제가 CUDA관련한 글을 쓰는건 David박사가 했던 대학강의와 NVIDIA에서 나온 CUDA 문서를 참고로 하고 있습니다. 그래서 문서에 나온 내용도 정리 하겠습니다. 그나저나 NVIDIA 코리아에서는 CUDA 그렇게 밀면서 문서 한글화는 안해주네요.ㅡㅡ; 그것만 해줘도 CUDA하기 수월할텐데 말이죠.
CUDA 프로세서 구조는 SIMT(single-instruction, multiple-thread) 구조입니다. 기존에 있던 병렬프로그래밍으로 따지자면 OpenMP와 같은 성격인 것입니다. OpenMP는 CPU갯수만큼의 병렬화를 지원하지만 CUDA는 GPU안의 계산 유닛갯수만큼을 지원하는 것이 차이입니다.
SIMT 유닛은 32개의 병렬 스레드의 그룹을 만드는데 이것들을 warps 라고 부릅니다. device에서 계산을 시작하게 되면 SIMT 유닛은 warps안에 각각의 스레드를 분해해서 넣습니다. 여기서 첫번째 warp는 Thread ID 0을 가지게 되는 겁니다.
그리고 각 Multiprocessor는 아래와 같은 구조를 가지면서 실행되게 됩니다.
위의 그림에 나와있는 각 Multiprocessor들은 다음의 4가지 형태의 메모리들을 가지게 됩니다.
1. Processor당 32비트 register들의 세트 하나. 2. Multiprocessor안의 모든 processor 코어들이 공유하는 Shared memory 3. 모든 Processor 코어들에 공유되고 constant memory space의 읽기속도를 향상시켜주는, 읽기전용인 constant cache 4. 모든 Processor 코어들에 공유되고 texture memory space의 읽기속도를 향상시켜주는, 읽기전용인 texture chche. 각각의 Multiprocessor들은 texture unit을 통해서 texture cache를 액세스합니다.
위에 내용들이 중요한 내용들입니다. 앞 포스팅에 있던 Grid, Block, Thread와 위에 있는 메모리들이 어떻게 매치되는지를 지정해주어야 에러없이 프로그램이 실행됩니다. 이게 제대로 관리안되면 에러가 나게 되는 겁니다. 여기서부터는 상세한 설명이 없어서 스스로의 많은 삽질이 필요하게 되는 부분입니다^^;
그리고 SLI를 이용해서 GPU를 여러개 꼽아서 사용할수도 있다고 나와있군요. 하나도 제대로 쓰기 어렵구만..;;
여기서 Thread라고 되어 있는 것들이 CUDA연산에서 가장 작은 단위인 것입니다. 이 Thread들은 Block의 하위에 있고 Block들은 다시 Grid의 하위에 있습니다. 그래서 프로그래머가 연산을 하려면
'몇번Grid의 몇번Block의 몇번Thread를 사용하여 계산하겠다'
를 지정해 주어야 합니다.
여기서 각 Thread는 자신만의 local memory를 가지게 됩니다. 각 Thread 간에 data 공유를 위한 shared memory는 따로 있습니다. 그리고 Block들의 집합인 Grid간에 data 공유를 위한 Global memory가 있습니다. 구조는 아래와 같습니다.
이 memory들을 어떻게 사용하느냐에 따라서 같은 연산을 하는 CUDA프로그램이라 하더라도 큰 성능차이를 내게 됩니다. memory간의 data를 이동시키는 비용이 꽤 크기 때문에 적절한 알고리즘을 작성하지 않는다면 계산하는데서 단축한 시간을 data 이동시키는데서 다 까먹을 수 있게 됩니다. 각 memory들의 속도 또한 차이가 있어서 이런 여러가지 요소를 고려하기 시작하면 복잡해지기 시작하는 겁니다. 결국 CUDA 프로그래밍의 핵심은 각 memory들의 특성을 파악해서 최대한 효율적인 알고리즘을 만드는데 있습니다.
아래 그림은 Grid를 이용해 계산한다는 것이 코드로는 어떻게 표현되는지를 나타냅니다.
결국 template project에서 보았던 kernel을 실행시키는 부분이 kernel에 해당하는 Grid를 실행시킨다는 것을 알 수 있습니다.
정리 해놓으신 포스트들 많은 도움이 되었습니다.
Visual Studio 2005에서 CUDA를 이용해 새로운 프로젝트를 생성해서 작업을 하려고 합니다. 혹시 release note에 있는 template프로젝트를 이용하는 방법 말고, 새로운 프로젝트를 생성하는 방법을 알고 계시면 도움을 주시면 감사하겠습니다. Visual studio 프로젝트 환경 설정에서 애를 먹고 있어서,,;;;;
그리고 원래 생성되어 있는 프로젝트를 CUDA환경으로 변환(?)할 수 있는 방법을 알고 계신지,,,;;;;
코딩에서 에러가 있어서 진행이 안되면 그나마 디버깅이라도 해볼 수 있겠는데,, 프로젝트 환경설정이 안되서 애를 먹고 있으니 너무 답답하네요,,;;;;
조언좀 해주시면 감사하겠습니다.
아,, 제 메일 주소는 stylidy@metq.com입니다.
아리수님 블로그 번창하시길 바랍니다.
제가 원래 CUDA프로그래밍을 리눅스에서 많이 했었기 때문에 질문주신 윈도우에서 새로운 프로젝트 환경 설정하는 방법은 해본적이 없습니다.
그리고 CUDA 디버깅은 Visual Studio 로는 어려운 부분이 있습니다. Visual Studio는 메인보드와 CPU상에서 일어나는 부분만 디버깅 가능하기 때문에 GPU상에서 어떻게 연산이 진행되는지에 관한 정보는 알기가 어렵습니다.
오랜만에 쓰네요. 여기는 CUDA관련해서 들어오시는 분들이 제일 많군요..;; 잠깐 검색해봤는데 해가 지났는데 아직 CUDA관련한 한글자료는 잘 없는 편이군요. 국내에 어느정도 보급은 된거 같은데 정보공유는 어떻게들 하시는지 모르겠습니다. nvidia싸이트 가봐도 그다지 활성화되어 있는거 같지는 않구요.
이번 글은 CUDA 메모리 구조 관련한 정리 입니다. 2.0은 안봐서 모르겠는데 이전버전에서의 CUDA는 메모리 구조를 숙지하고 있었어야 했기 때문에 그냥 그럴거라 생각하고 메모리구조에 대한 정리부터 이어갑니다.
우리가 CUDA를 쓰려고하는 이유는 GPU를 사용하여 연산속도의 향상을 얻기 위함입니다. GPU가 CPU연산보다 빠른 이유가 CPU보다 많은 ALU를 보유하고 있기 때문입니다.
위 그림에서 보시면 GPU가 CPU보다 월등히 많은 ALU를 가지고 있다는 것을 알 수 있습니다. 저 ALU를 어떻게 잘 사용하느냐가 속도향상의 관건입니다. CPU에서는 특정 ALU를 이용하는걸 프로그래머가 지정해주지 않지만 CUDA에서는 저 ALU하나하나를 모두 프로그래머가 제어해 주어야 합니다. 여기에 CUDA프로그래밍의 어려움이 있는 것입니다. 저 ALU구성은 큰 구조는 같지만, 하드웨어에 따라서 전체 ALU의 개수나 ALU에서 사용하는 GPU의 내장메모리 용량이 틀려지기 때문에 그러한 사항들을 감안해서 프로그래밍을 해주어야 속도향상이라는 결과를 얻을 수 있습니다. 그렇기 때문에 GPU 메모리 구조에 대한 이해가 수반되어야 하는 겁니다.
여기부터는 template_kernel.cu 에 관한 내용입니다. 첫번째 포스트쓰고 이래저래 딴짓하느라 그 사이 CUDA2.0 이 나왔네요.
template_kernel.cu는 실제 GPU연산을 담당하는 함수가 들어가 있습니다.
__global__ void testKernel( float* g_idata, float* g_odata) { // shared memory // the size is determined by the host application extern __shared__ float sdata[];
// access thread id const unsigned int tid = threadIdx.x; // access number of threads in this block const unsigned int num_threads = blockDim.x;
// read in input data from global memory // use the bank checker macro to check for bank conflicts during host // emulation SDATA(tid) = g_idata[tid]; __syncthreads();
// write data to global memory g_odata[tid] = SDATA(tid); }
여기서 void라는 것은 다른 프로그래밍 언어를 보아왔다면 많이들 보셨을 반환값이 없다는 그 뜻입니다. 그런데 여기에 __global__라는 구문이 붙어 있는데 이건 그래픽카드에 있는 여러가지 메모리종류중에서 global 메모리를 사용하겠다는 뜻입니다. CUDA프로그래밍 중에서는 상황에 따라서 어떤 메모리를 사용하느냐가 속도에 상당히 큰 영향을 미칩니다. 우선은 __global__이 들어가는 부분이 어떤종류의 메모리를 사용한다를 선언한다는 정도만 알고 넘어갑니다.
// shared memory // the size is determined by the host application extern __shared__ float sdata[]; 여기서도 마찬가지로 __shared__는 어떤 메모리를 사용할까를 지정하는 것입니다.
// access thread id const unsigned int tid = threadIdx.x; // access number of threads in this block const unsigned int num_threads = blockDim.x; 스레드아이디와 블럭아이디를 얻어옵니다. 이건 CUDA 병렬화와 관계가 있는데 CUDA는 데이터를 병렬로 처리하는데서 계산속도의 큰 향상을 얻습니다. 그러기 위해서 물리적으로 여러개의 계산 유닛이 그래픽카드안에 구조를 가지고 자리를 잡게 됩니다. 이러한 구조도 나중에 살펴보겠지만 위에 나오는 구문은 그 구조중에서 어떤부분을 사용할지를 얻어온다고 생각하면 되겠습니다. 예를 들어 for문을 돌려서 루프를 계산할때 1부터 10까지의 계산을 순차적으로 한다고 하면 1번 계산하고나서 2번 계산하고 그 다음에 3번 계산하고 이런식으로 진행이 되지만 CUDA에서는 그 각각의 계산들이 한꺼번에 동시에 진행이 되기 때문에 계산속도가 10분의 1로 줄어들게 됩니다. 물론, 그렇게 계산을 수행하기 위해서는 for문 안에서 돌아가는 계산들이 각각 의존적이지 않고 독립적으로 돌아가야 정상적으로 계산을 수행하게 됩니다.
// read in input data from global memory // use the bank checker macro to check for bank conflicts during host // emulation SDATA(tid) = g_idata[tid]; __syncthreads(); GPU메모리에 있는 데이터를 읽어와서 bank를 체크한다는데 이부분은 아직 메뉴얼을 봐도 정확히 잘 모르겠네요. nvidia의 데이비드 아저씨가 강의한 파워포인트 자료를 봐도 이부분은 헷갈리는 부분입니다. 영어가 짧아서 듣기는 안되고 강의자료만 보려니 그렇네요..ㅡㅡ;
// perform some computations SDATA(tid) = (float) num_threads * SDATA( tid); __syncthreads(); GPU에서 실제 계산을 수행합니다. 여기서 __syncthreads(); 이 부분이 있어야 전체 데이터들이 엉키지 않고 정상적으로 처리가 됩니다. 병렬화를 위해서 쪼개어져 계산되었던 값들이 모두 계산을 끝낼때까지 기다린다음에 값을 취합하는 겁니다. 만약 이 부분이 없다면 계산이 모두 끝나지 않았는데도 다음 구문으로 넘어가게 되니 계산 결과에 심각한 에러를 초래할 수 있습니다.
// write data to global memory g_odata[tid] = SDATA(tid); 그 다음에 사용했던 데이터를 global 메모리로 옮깁니다. 이건 GPU상에 존재하는 여러가지 메모리중 global로 값을 옮겨서 계산이 끝난 결과값에 CPU에서 접근이 가능하도록 만드는 겁니다
이 정도면 대략적으로 CUDA에 포함되어 있는 template project의 소스는 살펴본 것 같습니다. 지루한 개념공부보다는 일단 뭐라도 코드를 한번 돌려보는게 어떨까 싶어서 이걸 제일 먼저 다뤘습니다. 다음부터는 CUDA 프로그래밍에 필요한 개념정리와 디버깅 방법에 대해 정리해 볼까 합니다.
cuda를 설치하고 나면 기본적으로 제공되는 예제들 중에 template에 관해서 살펴보겠습니다.
뭐 시작하기전에 메모리 종류라던가 커널이라던가 이것저것 알아야 할 것들이 많지만 일단 돌려보고 그런건 차차 하면서 알아 나가는게 좋지 않을까 합니다.
이 카테고리는 제가 공부하고 정리하는 의미가 있기 때문에 크게 순서에 얽매이지 않고 그냥 생각나는대로 쓰고 있음을 밝혀둡니다.;;
그럼 우선 한번 실행 시켜보도록 하죠. Visual C++을 열고 C:\Program Files\NVIDIA Corporation\NVIDIA CUDA SDK\projects 폴더로 가면 예제들이 쭉 있습니다. 한 50개 좀 넘는군요.. 이 중에서 template라는 폴더로 가서 솔루션을 엽니다.
그 다음 F5를 눌러 프로그램을 실행시켜 봅니다. 별 문제 없이 CUDA환경이 설정되었다면 아래와 같은 화면이 나올 겁니다.
시간은 머신마다 틀리게 나올거고 저렇게 나오면 정상적으로 돌아간 겁니다.
그럼 이제 저 화면이 어떻게 출력되었나를 template프로젝트의 소스를 보면서 알아보도록 하겠습니다.
우선 visual C++의 솔루션탐색기에서 src를 보면 3개의 파일이 존재함을 알 수 있습니다. template.cu, template_gold.cpp, template_kernel.cu 이것들이 CUDA 프로그래밍을 하기 위한 기본적인 파일들인 것입니다.
이 파일들의 관계를 제대로 알아야지 CUDA프로그래밍을 할 수 있습니다. NVIDIA에서는 쉽다고 말하면서 예제로 matrix연산 하는걸 메뉴얼에 올려놨는데 쉽기는...ㅡㅡ; 메뉴얼만 보면 쉽네? 하고 생각할 수 있겠지만 저같은 평민은 그 뒤에 감춰진 kernel이라던가 cu파일들 보면 어렵기만 하더군요.
우선적으로 template_gold.cpp를 보면 computeGold라는 함수에서 reference를 계산하는걸 볼 수 있습니다.
void computeGold( float* reference, float* idata, const unsigned int len) { const float f_len = static_cast<float>( len); for( unsigned int i = 0; i < len; ++i) { reference[i] = idata[i] * f_len; } }
reference를 계산하는 목적은 GPU를 이용해서 계산한 값이 CPU와 이용한 값과 같은지를 비교하기 위해서 CPU를 이용해서 참조용 결과를 만드는 겁니다. 물론, 이건 예제이기에 저런 부분이 들어가는 것이고 GPU결과 검증을 위해서 필요한 부분이지만 실제로 프로그램을 작성할때는 저 부분은 빼주어도 cuda프로그래밍에 문제가 되지는 않습니다.
실제 CUDA프로그램은 template.cu, template_kernel.cu 라는 두 파일을 통해서 이루어 집니다. 이중에 template_kernel.cu는 실제 GPU상에서 연산이 이루어지는 내용을 포함합니다. 그리고 template.cu는 template_kernel.cu에 있는 testKernel이라는 함수를 돌리기 위한 준비과정과 돌리고 난후의 후처리를 담당합니다. 나중에 되면 굳이 파일을 2개로 분리할 필요없이 template.cu와 template_kernel.cu를 합쳐서 돌려도 별 문제없이 잘 돌아가는걸 확인할 수 있습니다. 일단, 개념 정립을 위해서는 template.cu는 GPU연산을 위한 준비와 후처리, kernel은 실제 GPU연산 수행 정도로 알아두면 될 것 같습니다.
// check if kernel execution generated and error CUT_CHECK_ERROR("Kernel execution failed");
// allocate mem for the result on host side float* h_odata = (float*) malloc( mem_size); // copy result from device to host CUDA_SAFE_CALL( cudaMemcpy( h_odata, d_odata, sizeof( float) * num_threads, cudaMemcpyDeviceToHost) );
// check result if( cutCheckCmdLineFlag( argc, (const char**) argv, "regression")) { // write file for regression test CUT_SAFE_CALL( cutWriteFilef( "./data/regression.dat", h_odata, num_threads, 0.0)); } else { // custom output handling when no regression test running // in this case check if the result is equivalent to the expected soluion CUTBoolean res = cutComparef( reference, h_odata, num_threads); printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED"); }
우선 CUT_DEVICE_INIT(); 를 호출하여 GPU를 사용할 준비를 합니다. 그 다음라인에서 수행시간을 측정하기 위한 변수를 할당하고 초기화 합니다. unsigned int timer = 0; CUT_SAFE_CALL( cutCreateTimer( &timer)); CUT_SAFE_CALL( cutStartTimer( timer));
그리고 몇 개의 쓰레드를 돌릴 것인지 지정하고 거기에 맞는 메모리 크기를 지정합니다. 여기서 중요한게 쓰레드의 갯수를 몇개로 지정할 것인가인데 나중에 이부분이 계산 수행속도에 큰 영향을 미치게 됩니다. unsigned int num_threads = 32; unsigned int mem_size = sizeof( float) * num_threads;
그 다음 호스트와 디바이스에 각각 메모리를 할당하는 부분이 나오는데 호스트는 일반적인 CPU가 사용하는 ram을 말하는 것이고 디바이스 메모리할당은 그래픽카드에 있는 메모리를 사용하기 위해 할당하는 것입니다. 이때, cudaMalloc라는 함수를 이용해서 디바이스 메모리를 할당하고 cudaMemcpy라는 함수를 이용하여 호스트메모리에 있는 내용을 디바이스 메모리로 가져옵니다. // allocate host memory float* h_idata = (float*) malloc( mem_size); // initalize the memory for( unsigned int i = 0; i < num_threads; ++i) { h_idata[i] = (float) i; }
그리고 계산결과를 저장할 디바이스 메모리를 할당합니다. // allocate device memory for result float* d_odata; CUDA_SAFE_CALL( cudaMalloc( (void**) &d_odata, mem_size));
그 다음 CUDA에 존재하는 built-in 변수인 dim3형으로 그리드와 쓰레드를 지정해 줍니다. 이 부분은 nVidia 그래픽카드의 메모리구조를 알아야 이해할 수 있는 부분인데 그냥 아까 지정했던 갯수만큼 사용할 쓰레드를 지정한다고 생각하면 될 듯 합니다.(이 부분에 대한 정확한 내용은 아직 잘 모르겠습니다. 좀 더 공부 필요..근데 이런게 쉽다고하다니...메뉴얼 어디에도 설명이 없는데..-_-+) // setup execution parameters dim3 grid( 1, 1, 1); dim3 threads( num_threads, 1, 1);
그 다음 드디어 kernel을 실행시킵니다. // execute the kernel testKernel<<< grid, threads, mem_size >>>( d_idata, d_odata);
커널을 실행한 다음에 문제없이 잘 실행됐는지 확인을 합니다. // check if kernel execution generated and error CUT_CHECK_ERROR("Kernel execution failed");
호스트메모리에 디바이스에 저장되어 있는 결과를 가져올 메모리를 할당한 다음 디바이스에 있는 메모리의 결과정보를 호스트 메모리로 가져옵니다. // allocate mem for the result on host side float* h_odata = (float*) malloc( mem_size); // copy result from device to host CUDA_SAFE_CALL( cudaMemcpy( h_odata, d_odata, sizeof( float) * num_threads, cudaMemcpyDeviceToHost) );
그 다음 수행시간을 측정해서 화면에 출력합니다. 제일처음에 보았던 시간 출력이 이 부분입니다. CUT_SAFE_CALL( cutStopTimer( timer)); printf( "Processing time: %f (ms)\n", cutGetTimerValue( timer)); CUT_SAFE_CALL( cutDeleteTimer( timer));
reference용으로 쓸 결과를 CPU에서 계산 합니다. // compute reference solution float* reference = (float*) malloc( mem_size); computeGold( reference, h_idata, num_threads);
GPU에서 계산되어진 결과와 CPU에서 계산되어진 결과가 일치하는지 비교합니다. // check result if( cutCheckCmdLineFlag( argc, (const char**) argv, "regression")) { // write file for regression test CUT_SAFE_CALL( cutWriteFilef( "./data/regression.dat", h_odata, num_threads, 0.0)); } else { // custom output handling when no regression test running // in this case check if the result is equivalent to the expected soluion CUTBoolean res = cutComparef( reference, h_odata, num_threads); printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED"); }
* G80 시리즈 중에서 8400은 현재 4만원이면 구입가능. 나머진 너무 비쌈..ㅡ_ㅡ; CUDA에서 에뮬레이션 모드를 지원하나 실제 머신에서 돌리는 것과 에뮬레이션 모드에서 돌릴때의 차이가 있기 때문에 에뮬레이션보다는 실제머신에서 돌려보는 것을 추천. ex) double precision같은 경우 현재 출시되어 있는 GPU들은 지원하지 않음. 그래서 실제 하드웨어 상에서 double을 쓰면 문제 발생. 그러나 에뮬레이션 모드에서는 이런 문제가 발생하지 않음. 필요 software windows xp visual c++ 2005 express edition
CUDA ("Compute Unified Device Architecture"),는 그래픽 처리장치(GPU)에서 수행하는 알고리즘을 코딩하는데 있어서 C 프로그래밍 언어를 사용할 수 있도록 하는 GPGPU 기술이다. CUDA는 엔비디아(Nvidia)에 의해서 개발되어져왔고 이 아키텍쳐의 사용하기위해선 Nvidia GPU와 특별한 스트림 처리 드라이버가 필요하다. CUDA는 G8X GPUs로 구성된 GeForce 8시리즈에서 동작할 수 있다.; Nvidia는 GeForce 8 시리즈상에 작성된 프로그램은 앞으로 개발될 Ncidia 비디오 카드에서도 프로그램 수정없이 작동할 것이라고 선언했다. CUDA는 CUDA GPUs 안의 명령셋과 대용량 병렬처리 메모리를 접근할 수 있도록 해준다.
최초의 CUDA SDK는 2007년 2월 15일에 공개되었다. CUDA의 컴파일러는 Open64에 기반을 두었다.
