2014/02/17

CUDA 공부 (GPU 메모리 재활용하기)

기록용으로 적어 놓습니다.

CUDA를 공부(전 https://class.coursera.org/hetero-002 여기에서 공부하고 있습니다.)하다보니, device 메모리에 값 넣어놓고 cudafree 쓰기 전까지 메모리에 남아 있는 값을 재활용할 수 있는지, 그리고 거기에 새로운 계산을 할 수 있는지 궁금해지더군요.

그래서 간단한 더하기 프로그래밍 예제를 조금 수정해보았습니다.

(참조 : http://developer.download.nvidia.com/books/cuda-by-example/cuda-by-example-sample.pdf)


---------------------------------
#include <stdio.h>

#define N 10

static void HandleError( cudaError_t err,
                         const char *file,
                         int line ) {
    if (err != cudaSuccess) {
        printf( "%s in %s at line %d\n", cudaGetErrorString( err ),
                file, line );
    }
}

#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))

__global__ void minusone(int *c)
{
unsigned int t = threadIdx.x;
unsigned int tx = blockIdx.x * blockDim.x + t;

if (tx < N)
c[tx] -= 1;
}

__global__ void add( int *a, int *b, int *c )
{
int tid = blockIdx.x; // handle the data at this index
if (tid < N)
c[tid] = a[tid] + b[tid];
}

int main(void)
{
int a[N], b[N], c[N];
int *dev_a, *dev_b, *dev_c;

HANDLE_ERROR( cudaMalloc( (void**)&dev_a, N * sizeof(int) ) );
HANDLE_ERROR( cudaMalloc( (void**)&dev_b, N * sizeof(int) ) );
HANDLE_ERROR( cudaMalloc( (void**)&dev_c, N * sizeof(int) ) );

for (int i=0; i<N; i++)
{
a[i] = -i;
b[i] = i * i;
}

HANDLE_ERROR( cudaMemcpy( dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice ) );
HANDLE_ERROR( cudaMemcpy( dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice ) );

add<<<N,1>>>( dev_a, dev_b, dev_c );

HANDLE_ERROR( cudaMemcpy( c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost ) );

for (int i=0; i<N; i++)
{
printf( "%d + %d = %d\n", a[i], b[i], c[i] );
}

minusone<<<N,1>>>( dev_c );

HANDLE_ERROR( cudaMemcpy( c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost ) );

for (int i=0; i<N; i++)
{
printf( "%d + %d - 1 = %d\n", a[i], b[i], c[i] );
}

cudaFree( dev_a );
cudaFree( dev_b );
cudaFree( dev_c );

for (int i=0; i<N; i++)
{
printf( "%d + %d - 1 = %d\n", a[i], b[i], c[i] );
}

return 0;
}
------------------------

참조의 42쪽에 있는 코드에다가 굵은 글자만 넣어봤습니다. block과 dim 사이즈를 같이 하고, 앞의 cuda kernel에서 실행한 결과가 들어있는 메모리인 dev_c를 그대로 새로운 cuda kernel에 넣고, 새로운 cuda kernel에서는 각 값에 1을 빼줍니다.

그랬더니 되더군요.

kernel 함수 안 에서의 코딩 개념이 CPU랑 달라서 헷갈리는데, 그걸 이런 식으로 쪼개면 이해하기 쉽지 않을까 싶네요.

댓글 없음: