CUDA C Programming Guide
HTML
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
PDF
https://docs.nvidia.com/cuda/pdf/CUDA_C_Programming_Guide.pdf
p202
.reg .b32 r1, r2; // 32ºñÆ® ·¹Áö½ºÅÍ r1°ú r2 ¼±¾ð
.global .f32 array[N]; // array[N]Àº global·Î ¿ä¼Ò´Â f32
start: mov.b32 r1, %tid.x; // ½º·¹µå IdÀÇ x ¿ä¼Ò¸¦ r1À¸·Î º¹»ç
shl.b32 r1, r1, 2; // ½º·¹µå Id¸¦ 2ºñÆ® ¿ÞÂÊ ½ÃÇÁÆ®
ld.global.b32 r2, array[r1]; // array[tid]¸¦ ½º·¹µå tid·Î
add.f32 r2, r2, 0.5; // 0.5 ´õÇϱâ
p208
// ÀåÄ¡¿¡¼ ½ÇÇàÇÏ´Â Ä¿³Î
__global__ void VecAdd(float* A, float* B, float* C, int N)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < N)
C[i] = A[i] + B[i];
}
// È£½ºÆ®¿¡¼ ½ÇÇàÇÏ´Â ÄÚµå
int main()
{
int N = ...;
size_t size = N * sizeof(float);
// ÀÔ·Â º¤ÅÍ h_A¿Í h_B¸¦ È£½ºÆ® ¸Þ¸ð¸®¿¡ ÇÒ´ç
float* h_A = (float*)malloc(size);
float* h_B = (float*)malloc(size);
// ÀÔ·Â º¤ÅÍ ÃʱâÈ
...
// ÀÔÃâ·Â º¤ÅÍ ¿µ¿ªÀ» ÀåÄ¡ ¸Þ¸ð¸®¿¡ ÇÒ´ç
float* d_A;
cudaMalloc(&d_A, size);
float* d_B;
cudaMalloc(&d_B, size);
float* d_C;
cudaMalloc(&d_C, size);
// È£½ºÆ® ¸Þ¸ð¸®¿¡¼ ÀåÄ¡ ¸Þ¸ð¸®·Î ÀÔ·ÂÀ» º¹»ç
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// Ä¿³Î È£Ãâ
int threadsPerBlock = 256;
int blocksPerGrid =
(N + threadsPerBlock - 1) / threadsPerBlock;
VecAdd<<>>(d_A, d_B, d_C, N);
// °á°ú¸¦ ÀåÄ¡ ¸Þ¸ð¸®¿¡¼ È£½ºÆ® ¸Þ¸ð¸®·Î º¹»ç
// µ¡¼À °á°ú h_C¸¦ È£½ºÆ® ¸Þ¸ð¸®·Î º¹»ç
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
// »ç¿ëÀÌ ³¡³ ÀåÄ¡ ¸Þ¸ð¸® ÇØÁ¦
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
// È£½ºÆ® ¸Þ¸ð¸® ÇØÁ¦
...
}
p210
// Ä¿³Î Á¤ÀÇ
__global__ void MatMul(float A[N][N], float B[N][N], float C[N][N])
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
int j = blockDim.y * blockIdx.y + threadIdx.y;
for(int k=0; k>>(A, B, C);
...
}
p211
int main()
{
...
// ºí·ÏÀÇ ½º·¹µå ¼ö¸¦ (16, 16)À¸·Î ÇÏ¿© MalMul Ä¿³ÎÀ» ½ÇÇà
dim3 threadsPerBlock (16, 16);
dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
MatMul<<>>(A, B, C);
...
}
p215
// 2°³ÀÇ ½ºÆ®¸² »ý¼º
cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
cudaStreamCreate(&stream[i]);
float* hostPtr;
cudaMallocHost(&hostPtr, 2 * size);
// 2°³ÀÇ ½ºÆ®¸² °¢°¢¿¡ ¸Þ¸ð¸® Àü¼Û°ú Ä¿³Î ½ÇÇà ¸í·ÉÀ» ¹ßÇà
for (int i = 0; i < 2; ++i) {
// ÀÔ·Â µ¥ÀÌÅ͸¦ È£½ºÆ® ¸Þ¸ð¸®¿¡¼ ÀåÄ¡ ¸Þ¸ð¸®·Î ºñµ¿±â Àü¼Û
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
size, cudaMemcpyHostToDevice, stream[i]);
// MyKernel ½ÇÇà
MyKernel <<<100, 512, 0, stream[i]>>>
(outputDevPtr + i * size, inputDevPtr + i * size, size);
// ó¸® °á°ú¸¦ ÀåÄ¡ ¸Þ¸ð¸®¿¡¼ È£½ºÆ® ¸Þ¸ð¸®·Î ºñµ¿±â Àü¼Û
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
size, cudaMemcpyDeviceToHost, stream[i]);
}
p216
__device__ volatile int X = 1, Y = 2;
__device__ void writeXY( )
{
X = 10; // ¨ç
Y = 20; // ¨è
}
__device__ void readXY( )
{
int A = X; // ¨é
int B = Y; // ¨ê
}
p221
__global__ void AplusB( int *ret, int a, int b) { // Ä¿³Î AplusB Á¤ÀÇ
// a + b + threadIdx.x¸¦ °è»êÇÏ¿© ret¿¡ ¾²±â
ret[threadIdx.x] = a + b + threadIdx.x; }
int main( ) { // È£½ºÆ® ÂÊÀÇ main ÇÁ·Î±×·¥ ÀÛ¼º
int *ret; // ¹Ýȯ °ª ret¿¡ ´ëÇÑ Æ÷ÀÎÅÍ Á¤ÀÇ
cudaMalloc(&ret, 1000 * sizeof(int)); // retÀÇ ¿µ¿ªÀ» ÀåÄ¡ ¸Þ¸ð¸®¿¡ ÇÒ´ç
AplusB<<< 1, 1000 >>>(ret, 10, 100); // Ä¿³Î AplusB¸¦ 1,000¹ø È£Ãâ
// È£½ºÆ® ¸Þ¸ð¸®¿¡ °á°ú¸¦ ÀúÀåÇÒ ¿µ¿ª ÇÒ´ç
int *host_ret = (int *)malloc(1000 * sizeof(int));
// È£½ºÆ® ¸Þ¸ð¸®¿¡ º¹»ç
cudaMemcpy(host_ret, ret, 1000 * sizeof(int), cudaMemcpyDefault);
for(int i=0; i<1000; i++) // 1,000°³ÀÇ °á°ú¸¦ ¼ø¼´ë·Î
printf("%d: A+B = %d\n", i, host_ret[i]); // printf·Î Ãâ·Â
free(host_ret); // È£½ºÆ® ¸Þ¸ð¸® ÇØÁ¦
cudaFree(ret); // ÀåÄ¡ ¸Þ¸ð¸® ÇØÁ¦
return 0;
}
p222
__device__ __managed__ int ret[1000]; // ret¸¦ °ü¸® ¿µ¿ªÀ¸·Î Á¤ÀÇ
__global__ void AplusB(int a, int b) { // Ä¿³Î AplusB Á¤ÀÇ. ret´Â Àμö¿¡ ¾øÀ½
ret[threadIdx.x] = a + b + threadIdx.x; // a + b + threadIdx.x¸¦ °è»êÇÏ¿© ret¿¡ ÀúÀå
}
int main( ) {
AplusB<<< 1, 1000 >>>(10, 100); // Ä¿³Î AplusB È£Ãâ
cudaDeviceSynchronize( ); // ¸ðµç ½º·¹µåÀÇ Á¾·á¸¦ ±â´Ù¸²
for(int i=0; i<1000; i++) // 1,000°³ÀÇ ret °ªÀ»
printf("%d: A+B = %d\n", i, ret[i]); // printf·Î Ãâ·Â
return 0;
}
p224
size_t size = 1024 * sizeof(float);
cudaSetDevice(0); // ÀåÄ¡ 0 ¼±ÅÃ
float* p0;
cudaMalloc(&p0, size); // ÀåÄ¡ 0¿¡ ¸Þ¸ð¸® ÇÒ´ç
MyKernel<<<1000, 128>>>(p0); // MyKernelÀ» ÀåÄ¡ 0¿¡ ¹ßÇà
cudaSetDevice(1); // ÀåÄ¡ 1 ¼±ÅÃ
float* p1;
cudaMalloc(&p1, size); // ÀåÄ¡ 1¿¡ ¸Þ¸ð¸® ÇÒ´ç
MyKernel<<<1000, 128>>>(p1); // MyKernelÀ» ÀåÄ¡ 1¿¡ ¹ßÇà
p231
// ÃÖ¼Ò OpenCL ÇÁ·Î±×·¥
#include
#include
#define NWITEMS 512
// ¸Þ¸ð¸®¿¡ °ªÀ» ¾²±â¸¸ ÇÏ´Â °£´ÜÇÑ memset Ä¿³Î ÇÁ·Î±×·¥
// Ä¿³Î ÇÁ·Î±×·¥Àº ¹®ÀÚ¿·Î ¼Ò½º Äڵ带 ÀÔ·ÂÇصÒ
const char *source =
"__kernel void memset( __global uint *dst ) \n"
"{ \n"
" dst[get_global_id(0)] = get_global_id(0); \n"
"} \n";
int main(int argc, char ** argv)
{
// ¨çÇ÷§ÆûÀÇ ID ¾ò±â
// cl_platform_idÇüÀÇ º¯¼ö platformÀ» Á¤ÀÇÇÏ°í clGetPlatformIDs( ) ÇÔ¼ö È£Ãâ
cl_platform_id platform;
clGetPlatformIDs( 1, &platform, NULL );
// ¨èGPU ÀåÄ¡ ÇüÅ ȮÀÎ
// cl_device_idÇüÀÇ º¯¼ö device¸¦ ¸¸µé°í clGetDeviceIDs( ) ÇÔ¼ö È£Ãâ
cl_device_id device;
clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU,
1,
&device,
NULL);
// ¨éÄÁÅؽºÆ®¿Í ¸í·É Å¥ ¸¸µé±â
// clCreateContext( ) ÇÔ¼ö¸¦ È£ÃâÇÏ°í ¹Ýȯ °ªÀº context¿¡ ÀúÀå
// ÇØ´ç context¸¦ »ç¿ëÇÏ¿© clCreateCommandQueue( )¸¦ È£Ãâ, ¸í·É Å¥¸¦ ¸¸µê
cl_context context = clCreateContext( NULL,
1,
&device,
NULL, NULL, NULL);
cl_command_queue queue = clCreateCommandQueue( context,
device,
0, NULL );
// ¨ê¼Ò½º Äڵ带 ÄÄÆÄÀÏÇÏ°í Ä¿³ÎÀÇ ¿£Æ®¸® Æ÷ÀÎÆ®¸¦ ¾òÀ½
// clCreateProgramWithSource( )¸¦ È£ÃâÇÏ¿© Ä¿³Î ¼Ò½º¸¦ ÄÄÆÄÀÏ
// clBuildProgram( )°úclCreateKernel( ) ÇÔ¼ö¸¦ È£ÃâÇÏ¿© Ä¿³Î ¸¸µé±â
cl_program program = clCreateProgramWithSource( context,
1,
&source,
NULL, NULL );
clBuildProgram( program, 1, &device, NULL, NULL, NULL );
cl_kernel kernel = clCreateKernel( program, "memset", NULL );
// ¨ëµ¥ÀÌÅÍ ¹öÆÛ ¸¸µé±â
// clCreateBuffer( )¸¦ È£ÃâÇÏ¿© ¸Þ¸ð¸® ¿µ¿ª Buffer¸¦ È®º¸
cl_mem buffer = clCreateBuffer( context,
CL_MEM_WRITE_ONLY,
NMITEMS * sizeof(cl_uint),
NULL, NULL );
// ¨ìÄ¿³Î ½ÇÇà
// clSetKernelArg( )¸¦ È£ÃâÇÏ¿© Ä¿³Î°ú ¹öÆÛ¸¦ ¿¬°á
// clEnqueueNDRangeKernel( )À» È£ÃâÇÏ¿© Ä¿³Î ½ÇÇàÀ» ¸í·É Å¥¿¡ ³ÖÀ½
size_t global_work_size = NWITEMS;
clSetKernelArg(kernel, 0, sizeof(buffer), (void*) &buffer);
clEnqueueNDRangeKernel( queue,
kernel,
1,
NULL,
&global_work_size,
NULL, 0, NULL, NULL);
clFinish( queue );
// ¨í¹öÆÛ¸¦ ¸ÅÇÎÇÏ°í °á°ú¸¦ Àоî Ãâ·Â
// clEnqueueMapBuffer( )¸¦ È£ÃâÇÏ¿© Buffer¸¦ È£½ºÆ® ¿µ¿ª¿¡ ¸ÅÇÎÇÏ°í
// È£½ºÆ®¿¡¼ Àеµ·Ï ÇÔ
cl_uint *ptr;
ptr = (cl_uint *) clEnqueueMapBuffer( queue,
buffer,
CL_TRUE,
CL_MAP_READ,
0,
NWITEMS * sizeof(cl_uint),
0, NULL, NULL, NULL );
int i;
for(i=0; i < NWITEMS; i++)
printf("%d %d\n", i, ptr[i]);
return 0;
}
p237
if(cond) then {
func(param1, data);
} else {
func(param2, data);
}
p237
if(cond) then {
p=param1;
} else {
p=param2;
}
func(p, data);
p248
void VecAdd(int n, float *a, float *b, float *c)
{
#pragma acc kernels
for (int i=0; i