forked from NVIDIA-Korea/CUDA
-
Notifications
You must be signed in to change notification settings - Fork 4
/
Copy pathprofile_step2.cu
128 lines (105 loc) · 2.88 KB
/
profile_step2.cu
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
/* compile command
compile : nvcc -DUSE_NVTX -lnvToolsExt -arch=sm_35 a.cu -o a.out
profile : nvprof -o log1.o ./a.out
view : nvvp
*/
#include <cstdio>
#ifdef USE_NVTX
#include "nvToolsExt.h"
const uint32_t colors[] = { 0x0000ff00, 0x000000ff, 0x00ffff00, 0x00ff00ff, 0x0000ffff, 0x00ff0000, 0x00ffffff };
const int num_colors = sizeof(colors)/sizeof(uint32_t);
#define PUSH_RANGE(name,cid) { \
int color_id = cid; \
color_id = color_id%num_colors;\
nvtxEventAttributes_t eventAttrib = {0}; \
eventAttrib.version = NVTX_VERSION; \
eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE; \
eventAttrib.colorType = NVTX_COLOR_ARGB; \
eventAttrib.color = colors[color_id]; \
eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII; \
eventAttrib.message.ascii = name; \
nvtxRangePushEx(&eventAttrib); \
}
#define POP_RANGE nvtxRangePop();
#else
#define PUSH_RANGE(name,cid)
#define POP_RANGE
#endif
__global__ void init_data_kernel( int n, double* x){
int i = blockIdx.x * blockDim.x + threadIdx.x;
if ( i < n ) {
x[i] = n - i;
}
}
__global__ void daxpy_kernel(int n, double a, double * x, double * y){
int i = blockIdx.x*blockDim.x + threadIdx.x;
if (i < n) {
y[i] = a*x[i] + y[i];
}
}
__global__ void check_results_kernel( int n, double correctvalue, double * x ){
int i = blockIdx.x*blockDim.x + threadIdx.x;
if (i < n)
{
if ( x[i] != correctvalue )
{
printf("ERROR at index = %d, expected = %f, actual: %f\n",i,correctvalue,x[i]);
}
}
}
void init_host_data( int n, double * x ){
PUSH_RANGE("init_host_data",1)
for (int i=0; i<n; ++i) {
x[i] = i;
}
POP_RANGE
}
void init_data(int n, double* x, double* x_d, double* y_d){
PUSH_RANGE("init_data",2)
cudaStream_t copy_stream;
cudaStream_t compute_stream;
cudaStreamCreate(©_stream);
cudaStreamCreate(&compute_stream);
cudaMemcpyAsync( x_d, x, n*sizeof(double), cudaMemcpyDefault, copy_stream );
init_data_kernel<<<ceil(n/256),256,0,compute_stream>>>(n, y_d);
cudaStreamSynchronize(copy_stream);
cudaStreamSynchronize(compute_stream);
cudaStreamDestroy(compute_stream);
cudaStreamDestroy(copy_stream);
POP_RANGE
}
void daxpy(int n, double a, double* x_d, double* y_d){
PUSH_RANGE("daxpy",3)
daxpy_kernel<<<ceil(n/256),256>>>(n,a,x_d,y_d);
cudaDeviceSynchronize();
POP_RANGE
}
void check_results( int n, double correctvalue, double* x_d ){
PUSH_RANGE("check_results",4)
check_results_kernel<<<ceil(n/256),256>>>(n,correctvalue,x_d);
POP_RANGE
}
void run_test(int n){
PUSH_RANGE("run_test",0)
double* x;
double* x_d;
double* y_d;
cudaSetDevice(0);
cudaMallocHost((void**) &x, n*sizeof(double));
cudaMalloc((void**)&x_d,n*sizeof(double));
cudaMalloc((void**)&y_d,n*sizeof(double));
init_host_data(n, x);
init_data(n,x,x_d,y_d);
daxpy(n,1.0,x_d,y_d);
check_results(n, n, y_d);
cudaFree(y_d);
cudaFree(x_d);
cudaFreeHost(x);
cudaDeviceSynchronize();
POP_RANGE
}
int main(){
int n = 1<<22;
run_test(n);
return 0;
}