CUDA C 확장 키워드(CUDA C Extension)

 CUDA C extension, 즉 CUDA C에서 확장된 키워드들에 대하여 소개하고자 합니다. 예제 코드를 보면 __global__과 같은 키워드들을 발견할 수 있을 것입니다. 이러한 키워드들이 어떤 의미이며 무슨 역할을 하는지 알아야 보다 효율적인 프로그래밍을 할 수 있을 것입니다.



 1. 함수의 수식어

 함수의 수식어들은 어디서 호출할 수 있느냐와 어디서 실행되느냐에 따라 나뉩니다. __global__, __device__, __host__, __device__ __host__ 이렇게 총 4가지의 경우가 가능합니다.

• __global__

 디바이스에서 실행되는 함수를 뜻합니다. 여기서 device란 이전 포스팅에서도 언급했듯이 GPU를 뜻합니다. __global__로 수식된 함수는 host에서 호출할 수는 있어도 device에서 호출할 수는 없습니다. 대신 device로 실행하는 커널 함수 지정에 사용할 수 있습니다.

 다음은 __global__로 수식한 함수의 간단한 예시입니다.

?

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23

#include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio .h>   __global__ void add(int a, int b, int *c) {  *c = a + b; }   int main() {  int c;  int *dev_c;    cudaMalloc((void**)&dev_c, sizeof(int));  add<<<1, 1>>>(2, 7, dev_c);    cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost);     printf("2 + 7 = %d\n", c);    cudaFree(dev_c);    return 0; }

 위와 같이 리턴값은 항상 void 형이어야 하며, 다른 리턴값을 가질 수 없도록 되어 있습니다. 그 이유는 __global__ 함수는 host에서 호출하게 되는데, host와 device 간의 메모리 읽고 쓰기는 API 함수를 통해 일어나기 때문입니다. 함수 호출 시에 <<< block의 갯수 , thread의 갯수 >>>를 이용하여 block과 thread의 갯수를 지정해 줄 수 있습니다. __global__로 수식된 함수는 device에서 실행되는 함수이지만 device에서 호출할 수 없습니다. 즉, 재귀호출이 불가능합니다.

 또한 함수 내에 static 변수를 가질 수 없으며, 가변형 인수를 가질 수 없는 등의 제약사항이 존재합니다. 가변형 인수를 가질 수 없다는 것은 다음과 같은 식으로 함수를 호출하는 코드는 불가능하다는 것을 뜻합니다.

?

1	add<<<1, 1>>>(int a, int b, dev_c);

 또한 __global__ __host__와 같은 용법으로 쓰일 수 없고, 공유 메모리를 이용하며 256 바이트까지의 인수를 사용할 수 있습니다.

• __device__

 위의 __global__과 마찬가지로 디바이스에서 실행되는 함수를 뜻합니다. 하지만 __global__과는 다르게 host에서는 호출이 불가능하고 device에서만 호출이 가능하도록 되어 있습니다. 즉, 디바이스 코드 중에 작성하여 디바이스 내에서 실행되는 서브함수로 사용됩니다. device에서 실행되고 device에서 호출되기 때문에 재귀호출이 가능하지 않느냐고 생각할 수도 있지만 재귀호출은 할 수 없습니다.

?

1 2 3 4 5 6 7

__device__ int subAdd(int a, int b) {  return a + b; }   __global__ void add(int a, int b, int *c) {  *c = subAdd(a, b); }

 __global__에서 예시로 들었던 add 함수의 코드를 조금만 바꾼 __device__ 함수 예시입니다. 실행시켜보면 똑같은 결과값이 나오는 것을 알 수 있습니다. __global__ 함수는 device 내에서 실행되는 함수이기 때문에 __device__ 함수를 호출할 수 있습니다. 하지만 host에서는 호출할 수 없기 때문에 다음과 같은 호출은 불가능합니다.

?

1 2 3 4 5 6 7 8 9 10 11 12 13 14

#include <stdio.h>   __device__ int subAdd(int a, int b) {  return a + b; }   int main() {  int c;  c = subAdd(2, 7);    printf("2 + 7 = %d\n", c);    return 0; }

 위 코드를 실행시키면 error : calling a __device__ function("subAdd") from a __host__ function("main") is not allowed라는 에러가 뜹니다. 즉, __host__ 함수인 main 함수에서 __device__ 함수인 subAdd 함수를 호출할 수 없다는 것입니다.

 __device__ 함수 역시 __global__ 함수와 마찬가지로 static 변수를 함수 내에 가질 수 없고, 가변형 인수를 가질 수 없습니다.

• __host__

 __host__ 함수는 위에서 언급했던 __global__이나 __device__와는 실행되는 위치부터가 다릅니다. host에서 실행되며, host에서만 호출할 수 있고, device에서는 호출할 수 없습니다. main 함수가 그 대표적인 예입니다. main 함수를 통해서 알 수 있듯이, __global__, __device__, __host__ 등이 지정되지 않은 경우에는 __host__를 지정한 것과 동일한 효과를 지닙니다.

?

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25

#include <stdio.h>   __device__ int subAdd(int a, int b) {  return a + b; }   __global__ void add(int a, int b, int *c) {  *c = subAdd(a, b); }   __host__ int main() {  int c;  int *dev_c;    cudaMalloc((void**)&dev_c, sizeof(int));  add<<<1, 1>>>(2, 7, dev_c);    cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost);    printf("2 + 7 = %d\n", c);    cudaFree(dev_c);    return 0; }

 위의 코드처럼 main 함수를 __host__로 지정해주어도 아무런 문제 없이 잘 실행이 됩니다. 이는 main 함수가 __host__ 함수이기 때문이며, 어떤 함수인지 지정해주지 않았을 때는 default로 __host__로 지정되기 때문입니다.

 __host__ 수식어는 __global__ 수식어와는 동시에 사용할 수 없지만, __device__ 수식어와는 함께 사용할 수 있습니다. 그것에 대한 내용은 아래에서 따로 설명하도록 하겠습니다.

• __device__ __host__

 __host__ 수식어와 __device__ 수식어를 동시에 사용한 경우입니다. 이 경우 host와 device 양쪽에서 모두 사용할 수 있는 함수로 작성할 수 있습니다.

?

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15

#include <stdio.h>   __device__ __host__ int subAdd(int a, int b) {  return a + b; }   int main() {  int c;    c = subAdd(2, 7);    printf("2 + 7 = %d\n", c);    return 0; }

 위 코드를 실행시키면 역시나 아무런 문제 없이 잘 동작하게 됩니다. 이는 subAdd 함수가 host와 device 모두에서 사용할 수 있는 함수이기 때문입니다. 이러한 수식어를 사용하는 것이 무척이나 편리한 경우가 생기는데, 그 때문에 device와 host 모두에서 사용 가능하도록 만들어진 것입니다.



 2. 변수의 수식어

 변수의 수식어들은 메모리 영역에 따라서 구분되어 집니다. 즉, 변수가 위치하는 메모리의 위치가 어디냐에 따라 __device__, __constant__, __shared__ 세 가지로 나누어집니다. 메모리에 대해서는 이후 포스팅에서 따로 자세히 설명하도록 하고, 이번 포스팅에서는 간단하게 설명하도록 하겠습니다.

• __device__

 함수의 수식어인 __device__와는 다르게 __device__ 변수는 글로벌 메모리 영역에 할당되어 프로그램이 종료될 때까지 유효하게 됩니다. __device__ 변수에는 모든 thread가 접근할 수 있고, host에서는 API 함수를 통해 읽기와 쓰기가 가능합니다.

?

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21

#include <stdio.h>   __device__ int d_sum;   __global__ void add(int a, int b) {  d_sum = a + b; }   __host__ int main() {  int h_sum = 0;    cudaMemset(&d_sum, 0, sizeof(int));    add<<<1, 1>>>(2, 7);    cudaMemcpyFromSymbol(&h_sum, d_sum, sizeof(int), 0, cudaMemcpyDeviceToHost);    printf("2 + 7 = %d\n", h_sum);    return 0; }

 __device__ 변수는 위와 같이 cudaMemcpyFromSymbol이라는 함수를 통해 host 메모리에 불러올 수 있습니다. 여기서 주의하실 점은 CUDA 5.0부터 바뀐 점으로, 이전 버전들과는 symbol을 다른 방식으로 사용하여야 한다는 점입니다.

?

1	cudaMemcpyFromSymbol(&h_sum, "d_sum", sizeof(int), 0, cudaMemcpyDeviceToHost);

 이전 방식으로 cudaMemcpyFromSymbol 함수를 호출하면 위와 같이 호출하여야 합니다. 즉, 이전 버전들에서는 symbol을 character string으로 사용하였다면, CUDA 5.0부터는 symbol을 direct로 사용할 수 있도록 바뀌었습니다. 대신 이전 버전들과 같은 방법으로 사용할 수 없습니다.

 하지만 CUDA 5.0에서도 symbol을 direct로 사용하였을 때, 빨간 밑줄이 그이며 다음과 같은 Error가 발생한다고 합니다만, 실제로 실행시켰을 때는 아무런 문제가 없습니다. 추후에 새롭게 업데이트가 되면서 사라질 문제가 될 것 같습니다.

• __constant__

 __constant__ 변수는 그 이름과 같게 상수 메모리, 즉 constant memory 영역에 할당되어 프로그램이 종료될 때까지 유효하게 됩니다. 모든 thread에서 접근이 가능하지만, __device__ 변수와는 다르게 __constant__ 변수는 읽기만 가능합니다. 대신 host에서 cudaMemcpyToSymbol 함수를 통해 값을 쓸 수 있도록 되어 있습니다.

?

1 2 3 4 5 6 7 8 9 10 11 12 13 14

#include <stdio.h>   __constant__ int d_sum = 0;   int main() {  int h_sum1 = 9;  int h_sum2 = 0;    cudaMemcpyToSymbol(d_sum, &h_sum1, sizeof(int), 0, cudaMemcpyHostToDevice);  cudaMemcpyFromSymbol(&h_sum2, d_sum, sizeof(int), 0, cudaMemcpyDeviceToHost);    printf("h_sum2 = %d\n", h_sum2);  return 0; }

 여기서도 주의하셔야 할 점은 CUDA 5.0부터 cudaMemcpyToSymbol의 쓰임 역시 바뀌었다는 점입니다. 위에서 __device__ 변수를 설명하면서 언급하였던 cudaMemcpyFromSymbol 함수와 마찬가지로 symbol을 direct로 사용하도록 바뀌었습니다. 즉, 이전 버전들에서는 다음과 같이 사용하였지만, CUDA 5.0부터는 사용할 수 없습니다.

?

1	cudaMemcpyToSymbol("d_sum", &h_sum1, sizeof(int), 0, cudaMemcpyHostToDevice);

 즉, 위와 같이 symbol을 character string으로 사용하게 되면 값이 바뀌지 않는 것을 확인할 수 있습니다.

• __shared__

 __shared__ 변수 역시 그 이름과 같게 공유 메모리 영역에 할당됩니다. 다만 다른 변수들과는 달리 실행 중인 thread block 상에서만 유효합니다. __device__ 변수나 __constant__ 변수가 프로그램이 종료될 때까지 유효한 것과는 조금 다릅니다. 또, __shared__ 변수는 block 내의 thread는 접근하여 읽고 쓰는 것이 가능하도록 되어 있습니다.

?

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21

__global__ void add(int a, int b, int *c) {  __shared__ int sum;  sum = a + b;  *c = sum; }   int main() {  int c;  int *dev_c;    cudaMalloc((void**)&dev_c, sizeof(int));  add<<<1, 1>>>(2, 7, dev_c);    cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost);    printf("2 + 7 = %d\n", c);    cudaFree(dev_c);    return 0; }

 위는 __shared__ 변수를 사용한 예시입니다. 위는 단순하게 하나의 block에 하나의 thread만이 실행되는 코드이지만, 많은 thread가 실행되는 코드라면 __shared__ 변수를 유용하게 사용할 수 있을 것입니다.

 이것으로 CUDA C extension에 대하여 설명을 끝마치도록 하겠습니다. 메모리에 대한 상세한 설명이나 함수의 수식어들에 대하여도 이후 포스팅에서 더욱 자세하게 설명하도록 하겠습니다.