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