로직은
https://class.coursera.org/hetero-002
http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html
등등을 참조했습니다.
데이터가 많아지면(10만개 정도....ㅋㅋㅋ), CPU에서 prefix sum를 구하는 것보다 GPU에서 prefix sum를 구하는 게 더 빠르다는 군요.
(링크 참조)
보통은 쓸 일이 없겠지만, 데이터 많아지면(ㅋㅋ) 소용이 있겠네요.
--------이하 코딩-----------
//되도록이면 2의 자승으로 블럭 사이즈를 지정해주는 게 로직상 좋을 듯 싶습니다.
#define BLOCK_SIZE 512
__global__ void addsums(float * output, int len)
{
//블럭과 쓰레드의 인덱스 관련 변수입니다. 개인 취향입니다. 적절하게 사용하면 되겠습니다.
unsigned int t = threadIdx.x;
unsigned int bi = blockIdx.x;
unsigned int bd = blockDim.x;
unsigned tx = t + bi * bd;
//아마 이대로 코딩하면 에러날 겁니다. ㅋㅋㅋ main 함수에서 extern shared memory관련 작업을 해줘야 됩니다. 자세한 건 검색 ㄱㄱ.
extern __shared__ float sumvalue[];
//경계(boundry)를 정확하게 지정해주는 if문은 행여나 하는 마음에라도 꼭 써주는 게 좋겠습니다.
if (tx < len)
{
sumvalue[bi] = 0;
__syncthreads();
//밑의 scan kernel을 실행하면 로직에 의해서 output의 각 블럭 마지막 자리에 블럭의 합계가 위치하게 됩니다.
//그 값들을 shared memory에 더합니다. shared memory의 크기는 이 코드에서는 sizeof(float)만 해도 될 것 같습니다. 그런데 고치기 귀찮아서 그냥 extern으로 해놨습니다. 쓸데없이 크게 메모리를 잡아놨다면, 메모리 낭비겠죠.
for ( int i = 1; i < bi+1; i++ )
{
sumvalue[bi] += output[ bd * i - 1];
}
//for 문 안에 __syncthreads()를 넣으면 값이 예상과 다르게 나올 수 있습니다.
//각 블럭의 계산이 다 끝난 후에 __syncthreads()를 불러주는 게 좋을 것 같습니다.
//마치 한가지 일이 다 끝나기 전에 다른 일 주면 헷갈려서 엉망이 되는 것과 비슷한 것 같습니다.
__syncthreads();
}
//앞 블럭들의 합계값(scan이 끝난 후의 블럭의 맨 마지막 값)을 각 블럭의 값들에 더해줘야 prefix sum이 완료되겠습니다. 헉헉
if( tx < len && t < BLOCK_SIZE )
{
__syncthreads();
output[tx] += sumvalue[bi];
}
}
__global__ void scan(float * input, float * output, int len)
{
//@@ Modify the body of this function to complete the functionality of
//@@ the scan on the device
//@@ You may need multiple kernel calls; write your kernels before this
//@@ function and call them from here
//역시 개인 취향
unsigned int t = threadIdx.x;
unsigned int tx = t + blockIdx.x * blockDim.x;
//코딩하다 남은 찌꺼기
int ite;
//@@ int maxite;
//하라는 대로 그냥 해봅니다. 이 강좌의 교수님은 shared memory를 충분하게 주라는 뜻인지, 꼭 블럭사이즈의 2배를 해줍니다. 뭐 하라면 합니다.
__shared__ float array[2*BLOCK_SIZE];
//shared memory에 입력값을 넣습니다. shared memory가 DDR보다 빨라서 계산시에는 shared memory 쓰라는 것 같습니다.
if( tx < len)
{
__syncthreads();
array[t] = input[tx];
input[tx] = 0;
}
for(ite=1; ite <= BLOCK_SIZE; ite*=2)
{
int index = (t + 1) * ite * 2 - 1;
//for 문 안에는 __syncthreads()는 안 넣는 게 좋을 것 같은데, 어쨌든 성공했으니 냅둡니다.
//로직의 첫번째 과정을 실행합니다. 저는 index와 ite와의 관계를 수십번 써보고 이해했습니다. --;;
__syncthreads();
if( index < 2*BLOCK_SIZE )
array[index] +=array[index - ite];
}
//로직의 두번째를 실행합니다. 이 로직이 끝나면 일단 하나의 블럭안에 있는 shared memory 의 값들은 prefix sum 상태가 됩니다.
for(ite=BLOCK_SIZE/2; ite > 0; ite/=2)
{
int index = (t+1) * ite * 2 -1;
__syncthreads();
if(index+ite < 2*BLOCK_SIZE)
array[index+ite] += array[index];
}
//각 shared memory 값들을 output 메모리에 넣습니다.
if( tx < len )
{
__syncthreads();
output[tx] = array[t];
}
}
------------------------------
계속 삽질을 해보니까 이제 감이 좀 잡히는 것 같습니다. 병렬계산은 기존의 로직하고는 좀 다르게 접근해야될 필요가 있습니다. 가래떡을 자를 때, 칼 하나로 빠르게 하느냐, 아니면 칼이 여러개 달린 틀로 자르냐의 차이와 비슷한 것 같습니다.
그리고, 수학적으로 더 빠르다고 증명이 된 로직이라고 하더라도, 물리적인 이유로 CPU와 별 차이가 없을 수도 있다는 생각이 듭니다. CPU메모리에서 GPU메모리로 옮기고, 그걸 shared memory에 옮기고, 그걸 계산시키고, 그걸 다시 GPU메모리로 옮기고, 다시 CPU메모리로 옮기는 작업을 잘 해야될 것 같네요.
참고:
- http://celdee.tistory.com/686
- http://micol.tistory.com/m/post/view/id/227
- http://www.slideshare.net/dgtman/cuda-12213924
댓글 없음:
댓글 쓰기