[CUDA] CUDA로 threading & shared memory

Parallel Blocks 나누기

•  cuda에서는 parallel로 실행되는 이 block들을 thread로 나눌 수 있는 방법을 제시한다. 
• 본 포스팅에서는 어떻게 thread로 block들을 나누어볼 수 있는지 splitting 방법에 대해서 다룬다. 
• 이전 포스팅에서 다음과 같은 kernel call 에 대해서 첫번째 인자가 block의 수 라고 했었다. 
  • kernel <<<number_of_parallel_blocks, threads per block>>> 
  • 두번째 인자는 block 당 돌리고 싶은 thread의 수를 의미한다. 
  • 예를 들어 add<<<1,N>>> (dev_a,dev_b,dev_c) 하게 되면 하나의 블록에 n개의 thread가 돌아가도록 설정한 것이다.

MaxThreadsPerBlock 

• 주로 block의 개수에는 한계가 있는데 일반적으로 한번의 실행에 65,535개 라고 한다.
• block에 실행할 수 있는 thread의 수도 512개로 한정이 되어 있는 경우가 많다. 
• 따라서 512보다 많은 양의 thread를 돌리고 싶은 경우는 block과 thread를 이용한 계산이 필요하다.

int tid= threadIdx.x +blockIdx.x*blockDim.x;

• blockDim : 블록의 크기 
• blockIdx: 블록의 index 번호
• threadIdx.x : 블록 내 thread의 번호 
• 이차원 배열의 개념으로 생각하면 쉽다. 

• 예시 
  • N개의 thread가 있다. 
  • 블록 당 128개의 thead를 실행하고 싶다면? 

add <<<(N+127)/128,128>>> (dev_a, dev_b,dev_c);

 

Shared Memory and Synchronization 

• cuda는 shared memory 가 있다!
• __shared__ 를 활용하여 shared memory 에 variable을 둘 수 있다. 
• block의 thread끼리는 공유할 수 있으나, 다른 block에서는 접근할 수 없다. 
• 물리적으로 gpu에 가깝기 때문에 접근 시간이 빠르다는 장점이 있다. 
• 그런데 여느 shared memory 와 같이 synchronization 에 주의해서 race condition이 발생하지 않도록 해야한다. 
• dot product 예시 

#define imin(a,b) (a<b?a:b) 

const int N=33*1024; 
cont int threadsPerBlock =256; 

__global__void dot(float* a, float* b, float* c){
	__shared__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x*blockDim.x; 
    int cacheIndex = threadIdx.x; 
    
    float temp =0 ; 
    
    //곱한 값 저장
    while(tid<N){
    	temp+=a[tid]*b[tid];
        tid+= blockDim.x*gridDim.x; 
        }
    cache[cacheIndex]= temp;
    
    __synchthread(); 
    
    //곱한 값들 더하기 
    int i = blockDim.x/2;
    while(i!=0) { 
    	if(cacheIndex <i) 
        	//reduction 방법으로 더하기를 반으로 나누어 가며 진행
        	cache[cacheIndex] += cache[cacheIndex+i];
            
        __synchthreads();
        i/=2;
   }
   
   //결과 값 저장 
   if(cacheIndex==0) {
	c[blockIdx.x]=cache[0];

• __shared__ 를 이용해서 shared memory 에 저장할 변수를 선언했다.
• __synchthreads() : 다음으로 넘어가기전 block의 모든 thread가 실행될때까지 기다리는 명령
• 결과는 global c variable에 저장