Skip to content

Commit a981654

Browse files
jirikrausjirikraus
committed
Added NVTX examples for blogpost
1 parent 1bedcc3 commit a981654

4 files changed

Lines changed: 315 additions & 0 deletions

File tree

‎posts/nvtx/Makefile‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
CC=gcc
2+
CXX=g++
3+
NVCC=nvcc
4+
5+
BINARIES=manual_nvtx compiler_inst_nvtx
6+
7+
all: $(BINARIES)
8+
9+
inst_nvtx.o: inst_nvtx.cpp Makefile
10+
g++ -fPIC -I${CUDA_ROOT}/include -c inst_nvtx.cpp
11+
12+
manual_nvtx: manual_nvtx.cu Makefile
13+
nvcc -DUSE_NVTX -arch=sm_20 -lnvToolsExt -o manual_nvtx manual_nvtx.cu
14+
15+
compiler_inst_nvtx: compiler_inst_nvtx.cu Makefile inst_nvtx.o
16+
nvcc -Xcompiler -fPIC -Xcompiler -finstrument-functions -arch=sm_20 inst_nvtx.o -ldl -lnvToolsExt -o compiler_inst_nvtx compiler_inst_nvtx.cu
17+
18+
clean:
19+
rm -f *.o $(BINARIES)
20+
21+
run: $(BINARIES)
22+
nvprof -o compiler_inst_nvtx.nvvp ./compiler_inst_nvtx
23+
nvprof -o manual_nvtx.nvvp ./manual_nvtx

‎posts/nvtx/compiler_inst_nvtx.cu‎

Lines changed: 99 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,99 @@
1+
#include <cstdio>
2+
3+
__global__ void init_data_kernel( int n, double* x)
4+
{
5+
int i = blockIdx.x * blockDim.x + threadIdx.x;
6+
if ( i < n )
7+
{
8+
x[i] = n - i;
9+
}
10+
}
11+
12+
13+
__global__ void daxpy_kernel(int n, double a, double * x, double * y)
14+
{
15+
int i = blockIdx.x*blockDim.x + threadIdx.x;
16+
if (i < n)
17+
{
18+
y[i] = a*x[i] + y[i];
19+
}
20+
}
21+
22+
__global__ void check_results_kernel( int n, double correctvalue, double * x )
23+
{
24+
int i = blockIdx.x*blockDim.x + threadIdx.x;
25+
if (i < n)
26+
{
27+
if ( x[i] != correctvalue )
28+
{
29+
printf("ERROR at index = %d, expected = %f, actual: %f\n",i,correctvalue,x[i]);
30+
}
31+
}
32+
}
33+
34+
void init_host_data( int n, double * x )
35+
{
36+
for (int i=0; i<n; ++i)
37+
{
38+
x[i] = i;
39+
}
40+
}
41+
42+
void init_data(int n, double* x, double* x_d, double* y_d)
43+
{
44+
cudaStream_t copy_stream;
45+
cudaStream_t compute_stream;
46+
cudaStreamCreate(&copy_stream);
47+
cudaStreamCreate(&compute_stream);
48+
49+
cudaMemcpyAsync( x_d, x, n*sizeof(double), cudaMemcpyDefault, copy_stream );
50+
init_data_kernel<<<ceil(n/256),256,0,compute_stream>>>(n, y_d);
51+
52+
cudaStreamSynchronize(copy_stream);
53+
cudaStreamSynchronize(compute_stream);
54+
55+
cudaStreamDestroy(compute_stream);
56+
cudaStreamDestroy(copy_stream);
57+
}
58+
59+
void daxpy(int n, double a, double* x_d, double* y_d)
60+
{
61+
daxpy_kernel<<<ceil(n/256),256>>>(n,a,x_d,y_d);
62+
cudaDeviceSynchronize();
63+
}
64+
65+
void check_results( int n, double correctvalue, double* x_d )
66+
{
67+
check_results_kernel<<<ceil(n/256),256>>>(n,correctvalue,x_d);
68+
}
69+
70+
void run_test(int n)
71+
{
72+
double* x;
73+
double* x_d;
74+
double* y_d;
75+
cudaSetDevice(0);
76+
cudaMallocHost((void**) &x, n*sizeof(double));
77+
cudaMalloc((void**)&x_d,n*sizeof(double));
78+
cudaMalloc((void**)&y_d,n*sizeof(double));
79+
80+
init_host_data(n, x);
81+
82+
init_data(n,x,x_d,y_d);
83+
84+
daxpy(n,1.0,x_d,y_d);
85+
86+
check_results(n, n, y_d);
87+
88+
cudaFree(y_d);
89+
cudaFree(x_d);
90+
cudaFreeHost(x);
91+
cudaDeviceSynchronize();
92+
}
93+
94+
int main()
95+
{
96+
int n = 1<<22;
97+
run_test(n);
98+
return 0;
99+
}

‎posts/nvtx/inst_nvtx.cpp‎

Lines changed: 60 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,60 @@
1+
#include <stdio.h>
2+
#include "nvToolsExt.h"
3+
#include <dlfcn.h>
4+
#include <cxxabi.h>
5+
6+
const char* const default_name = "Unknown";
7+
8+
const uint32_t colors[] = { 0x0000ff00, 0x000000ff, 0x00ffff00, 0x00ff00ff, 0x0000ffff, 0x00ff0000, 0x00ffffff };
9+
const int num_colors = sizeof(colors)/sizeof(uint32_t);
10+
static int color_id = 0;
11+
12+
extern "C" void __cyg_profile_func_enter(void *this_fn, void *call_site) __attribute__((no_instrument_function));
13+
extern "C" void __cyg_profile_func_exit(void *this_fn, void *call_site) __attribute__((no_instrument_function));
14+
15+
void rangePush(const char* const name ) __attribute__((no_instrument_function));
16+
void rangePop() __attribute__((no_instrument_function));
17+
18+
extern "C" void __cyg_profile_func_enter(void *this_fn, void *call_site)
19+
{
20+
Dl_info this_fn_info;
21+
if ( dladdr( this_fn, &this_fn_info ) )
22+
{
23+
int status = 0;
24+
rangePush(abi::__cxa_demangle(this_fn_info.dli_sname,0, 0, &status));
25+
}
26+
else
27+
{
28+
rangePush(default_name);
29+
}
30+
} /* __cyg_profile_func_enter */
31+
32+
extern "C" void __cyg_profile_func_exit(void *this_fn, void *call_site)
33+
{
34+
rangePop();
35+
} /* __cyg_profile_func_enter */
36+
37+
void rangePush(const char* const name )
38+
{
39+
nvtxEventAttributes_t eventAttrib = {0};
40+
eventAttrib.version = NVTX_VERSION;
41+
eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE;
42+
eventAttrib.colorType = NVTX_COLOR_ARGB;
43+
eventAttrib.color = colors[color_id];
44+
color_id = (color_id+1)%num_colors;
45+
eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII;
46+
if ( name != 0 )
47+
{
48+
eventAttrib.message.ascii = name;
49+
}
50+
else
51+
{
52+
eventAttrib.message.ascii = default_name;
53+
}
54+
nvtxRangePushEx(&eventAttrib);
55+
}
56+
57+
void rangePop()
58+
{
59+
nvtxRangePop();
60+
}

‎posts/nvtx/manual_nvtx.cu‎

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

0 commit comments

Comments
 (0)