CUDA(Compute Unified Device Architecture)는 계산 병렬처리를 지원하는 NVIDIA GPU 아키텍처이며 프로그래밍 기술이다.
CUDA와 CUDnn은 운영체제 및 버전에 맞게 엔비디아 사이트에서 다운로드 받고 설치해야 한다.
만약, 아나콘다로 설치하려면 다음 명령을 입력한다.
설치 후 cudnn.h 파일 중 아래 부분을 수정한다.
아래와 같이 쉘 파일을 편집해 경로를 추가한다.
환경변수를 체크한다.
제대로 설치되었는 지 버전 등 체크한다.
예제를 아래 폴더에서 home 폴더로 복사한다.
cudnn 폴더 아래 mnist 예제 폴더안에서 다음 같이 메이크한다.
빌드된 결과를 실행하면 다음 결과를 얻을 수 있다.
CUDA는 다음 그림과 같이 쓰레드, 블록 개념을 이용해 병렬처리할 수 있다. 각 이미지 픽셀에 특정 필터 연산을 하는 등의 작업 시간을 크게 줄일 수 있다.
CUDA는 병렬로 단위연산이 가능한 실수형 연산에 특화된 계산유닛을 제어할 수 있는 아키텍처로 설계되어 있다. GPU내에는 여러 개의 멀티프로세서가 존재하는데, 하나의 멀티프로세서에서 작동되는 스레드의 모음을 블록이라고 한다. 각 유닛은 스래드로 불리고, 다차원 연산이 가능하도록 블럭으로 스래드를 묶여 관리한다. 즉, Block과 Thread는 GPU 코드를 병렬로 처리하기 위한 단위로서 1개 블록은 N개 쓰레드로 구성된다.
CUDA는 병렬처리 대상이 데이터이므로, 데이터를 행렬이나 다차원 텐서로 표현해, 각 컴포넌트를 병렬처리할 수 있도록 함수를 정의하는 방법을 제공한다. 그러므로, 함수는 공동으로 사용되며, 함수의 입출력 데이터가 보관되는 메모리만 호스트 메모리(PC. CPU)와 디바이스 메모리(GPU)로 복사 전환하는 함수를 제공해 GPU에서 병렬 연산시 GPU메모리에 있는 데이터를 입출력할 수 있도록 하였다.
그러므로, 프로그래밍 시 정의되는 함수는 호스트와 디바이스 모두 실행되거나 각 영역에서만 실행되는 표식(키워드)가 필요하며, 다음 그림과 같이 CPU 와 GPU 메모리 사이를 교환할 수 있는 메모리 관리 API가 제공되어야 한다. CUDA TOOLKIT 은 이를 제공하는 라이브러리와 컴파일러의 패키지이다.
GPU프로그래밍에서 kernel(커널)은 병렬 실행되는 명령 집합을 의미한다. 다음과 같음 커널 함수를 코딩하는 경우, 블록갯수 B와 스레드수 S를 지정해주도록 한다.
func<<<B, S>>>();
아울러, __global__ 키워드를 함수 앞에 붙여서, GPU 동작 함수임을 명시해줘야 한다. 이 키워드를 이용해 CUDA 컴파일러인 NVCC가 입력되는 소스코드에서 함수들을 구분해 CPU 기계어, GPU 기계어로 오브젝트 바이너리 파일을 생성할 수 있다.
NVIDIA는 게임 렌더링에 특화되어 발전되었으므로, VERTEX BUFFER와 같이 3차원 형상을 계산하고 렌더링할 때 필요한 자료구조와 함수를 추가적으로 제공한다.
CUDA TOOLKIT 을 설치하고, 해당 버전에 맞는 VISUAL STUDIO를 설치하면, 관련된 기능 개발이 가능하다.
쿠다 프로그래밍 예시
개발 툴킷을 설치하면, 다음과 같은 예제를 실행해 볼 수 있다.
다음은 병렬 연산에서 주로 많이 적용되는 벡터 더하기 계산 예제이다.
#include <stdio.h>
#include <cuda_runtime.h>
#include <helper_cuda.h>
__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements)
{
int i = blockDim.x * blockIdx.x + threadIdx.x; // 병렬처리 시 사용되는 유닛 인덱스 계산
if (i < numElements)
C[i] = A[i] + B[i]; // 해당 유닛마다 덧셈 실수 연산 처리
}
int main(void)
{
cudaError_t err = cudaSuccess;
int numElements = 50000; // 5만개 실수 연산
size_t size = numElements * sizeof(float);
printf("[Vector addition of %d elements]\n", numElements);
// 호스트 CPU 메모리 할당
float *h_A = (float *)malloc(size); // 입력1
float *h_B = (float *)malloc(size); // 입력2
float *h_C = (float *)malloc(size); // 출력
if (h_A == NULL || h_B == NULL || h_C == NULL)
{
fprintf(stderr, "Failed to allocate host vectors!\n");
exit(EXIT_FAILURE);
}
for (int i = 0; i < numElements; ++i) // 램덤값으로 초기화
{
h_A[i] = rand()/(float)RAND_MAX;
h_B[i] = rand()/(float)RAND_MAX;
}
float *d_A = NULL;
err = cudaMalloc((void **)&d_A, size); // CUDA GPU 메모리도 할당
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
float *d_B = NULL;
err = cudaMalloc((void **)&d_B, size);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to allocate device vector B (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
float *d_C = NULL;
err = cudaMalloc((void **)&d_C, size);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to allocate device vector C (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
printf("Copy input data from the host memory to the CUDA device\n");
err = cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); // CPU > GPU 메모리로 값 복사
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to copy vector A from host to device (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
err = cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to copy vector B from host to device (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// 병렬처리 커널함수 실행
int threadsPerBlock = 256;
int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;
printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid, threadsPerBlock);
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements); // 병렬처리 커널함수
err = cudaGetLastError();
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to launch vectorAdd kernel (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// 계산 결과를 GPU에서 CPU메모리로 복사
printf("Copy output data from the CUDA device to the host memory\n");
err = cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to copy vector C from device to host (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// 계산 결과 검증
for (int i = 0; i < numElements; ++i)
{
if (fabs(h_A[i] + h_B[i] - h_C[i]) > 1e-5) // 실수 연산이므로, 특정 소숫점 자리수 이하로는 차이 발생함
{
fprintf(stderr, "Result verification failed at element %d!\n", i);
exit(EXIT_FAILURE);
}
}
printf("Test PASSED\n");
// 메모리 해제
err = cudaFree(d_A);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to free device vector A (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
err = cudaFree(d_B);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to free device vector B (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
err = cudaFree(d_C);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to free device vector C (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Free host memory
free(h_A);
free(h_B);
free(h_C);
printf("Done\n");
return 0;
}
다음은 프로그램 실행 결과이다. CPU 사용해 계산하는 것 보다 훨씬 빠르게 계산하는 것을 알 수 있다.
다음은 쿠다 기반 OPEN GL 렌더링의 간단한 예제이다.
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#ifdef _WIN32
# define WINDOWS_LEAN_AND_MEAN
# define NOMINMAX
# include <windows.h>
#endif
// OpenGL Graphics includes
#include <helper_gl.h>
#include <GL/freeglut.h>
// includes, cuda
#include <cuda_runtime.h>
#include <cuda_gl_interop.h>
// Utilities and timing functions
#include <helper_functions.h> // includes cuda.h and cuda_runtime_api.h
// CUDA helper functions
#include <helper_cuda.h> // helper functions for CUDA error check
#include <vector_types.h>
#define MAX_EPSILON_ERROR 10.0f
#define THRESHOLD 0.30f
#define REFRESH_DELAY 10 //ms
// 해상도 변수
const unsigned int window_width = 512;
const unsigned int window_height = 512;
const unsigned int mesh_width = 256;
const unsigned int mesh_height = 256;
// 정점 버퍼 vbo 변수 정의
GLuint vbo;
struct cudaGraphicsResource *cuda_vbo_resource;
void *d_vbo_buffer = NULL;
float g_fAnim = 0.0;
// mouse controls
int mouse_old_x, mouse_old_y;
int mouse_buttons = 0;
float rotate_x = 0.0, rotate_y = 0.0;
float translate_z = -3.0;
StopWatchInterface *timer = NULL;
// Auto-Verification Code
int fpsCount = 0; // FPS count for averaging
int fpsLimit = 1; // FPS limit for sampling
int g_Index = 0;
float avgFPS = 0.0f;
unsigned int frameCount = 0;
unsigned int g_TotalErrors = 0;
bool g_bQAReadback = false;
int *pArgc = NULL;
char **pArgv = NULL;
#define MAX(a,b) ((a > b) ? a : b)
// 함수 선언
bool runTest(int argc, char **argv, char *ref_file);
void cleanup();
bool initGL(int *argc, char **argv);
void createVBO(GLuint *vbo, struct cudaGraphicsResource **vbo_res,
unsigned int vbo_res_flags);
void deleteVBO(GLuint *vbo, struct cudaGraphicsResource *vbo_res);
// OPEN GL 렌더링 콜백 함수
void display();
void keyboard(unsigned char key, int x, int y);
void mouse(int button, int state, int x, int y);
void motion(int x, int y);
void timerEvent(int value);
// Cuda functionality
void runCuda(struct cudaGraphicsResource **vbo_resource);
void runAutoTest(int devID, char **argv, char *ref_file);
void checkResultCuda(int argc, char **argv, const GLuint &vbo);
const char *sSDKsample = "simpleGL (VBO)";
// 각 정점에 대해, 정점 데이터를 처리하는 BLOCK, THREAD X, Y 위치를 입력받고,
// SIN, COS함수를 통과해, Y값인 W값을 계산해, POS에 저장함.
// 각 위치 데이터 계산을 병렬처리하므로, 전체 실수 연산 시간이 크게 줄어듬을 보여줌.
__global__ void simple_vbo_kernel(float4 *pos, unsigned int width, unsigned int height, float time)
{
unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
// calculate uv coordinates
float u = x / (float) width;
float v = y / (float) height;
u = u*2.0f - 1.0f;
v = v*2.0f - 1.0f;
// calculate simple sine wave pattern
float freq = 4.0f;
float w = sinf(u*freq + time) * cosf(v*freq + time) * 0.5f;
// write output vertex
pos[y*width+x] = make_float4(u, w, v, 1.0f);
}
// 매 시간에 대한 메쉬의 정점 높이값을 실수 계산, 병렬 처리.
void launch_kernel(float4 *pos, unsigned int mesh_width,
unsigned int mesh_height, float time)
{
// execute the kernel
dim3 block(8, 8, 1);
dim3 grid(mesh_width / block.x, mesh_height / block.y, 1);
simple_vbo_kernel<<< grid, block>>>(pos, mesh_width, mesh_height, time); // 블록과 스레드 유닛 병렬처리
}
// GPU 지원 여부 확인
bool checkHW(char *name, const char *gpuType, int dev)
{
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
strcpy(name, deviceProp.name);
if (!STRNCASECMP(deviceProp.name, gpuType, strlen(gpuType)))
{
return true;
}
else
{
return false;
}
}
// 메인 함수
int main(int argc, char **argv)
{
char *ref_file = NULL;
pArgc = &argc;
pArgv = argv;
printf("%s starting...\n", sSDKsample);
if (argc > 1)
{
if (checkCmdLineFlag(argc, (const char **)argv, "file"))
{
getCmdLineArgumentString(argc, (const char **)argv, "file", (char **)&ref_file);
}
}
printf("\n");
runTest(argc, argv, ref_file);
printf("%s completed, returned %s\n", sSDKsample, (g_TotalErrors == 0) ? "OK" : "ERROR!");
exit(g_TotalErrors == 0 ? EXIT_SUCCESS : EXIT_FAILURE);
}
void computeFPS()
{
frameCount++;
fpsCount++;
if (fpsCount == fpsLimit)
{
avgFPS = 1.f / (sdkGetAverageTimerValue(&timer) / 1000.f);
fpsCount = 0;
fpsLimit = (int)MAX(avgFPS, 1.f);
sdkResetTimer(&timer);
}
char fps[256];
sprintf(fps, "Cuda GL Interop (VBO): %3.1f fps (Max 100Hz)", avgFPS);
glutSetWindowTitle(fps);
}
// 일반적인 GL 렌더링 초기화 함수
bool initGL(int *argc, char **argv)
{
glutInit(argc, argv);
glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);
glutInitWindowSize(window_width, window_height);
glutCreateWindow("Cuda GL Interop (VBO)");
glutDisplayFunc(display);
glutKeyboardFunc(keyboard);
glutMotionFunc(motion);
glutTimerFunc(REFRESH_DELAY, timerEvent,0);
// initialize necessary OpenGL extensions
if (! isGLVersionSupported(2,0))
{
fprintf(stderr, "ERROR: Support for necessary OpenGL extensions missing.");
fflush(stderr);
return false;
}
// default initialization
glClearColor(0.0, 0.0, 0.0, 1.0);
glDisable(GL_DEPTH_TEST);
// viewport
glViewport(0, 0, window_width, window_height);
// projection
glMatrixMode(GL_PROJECTION);
glLoadIdentity();
gluPerspective(60.0, (GLfloat)window_width / (GLfloat) window_height, 0.1, 10.0);
SDK_CHECK_ERROR_GL();
return true;
}
// CUDA로 병렬 처리함
bool runTest(int argc, char **argv, char *ref_file)
{
sdkCreateTimer(&timer); // 타이머 생성
int devID = findCudaDevice(argc, (const char **)argv); // CUDA 장치 ID획득
if (ref_file != NULL)
{
checkCudaErrors(cudaMalloc((void **)&d_vbo_buffer, mesh_width*mesh_height*4*sizeof(float))); // GPU 메모리 할당
runAutoTest(devID, argv, ref_file);
checkResultCuda(argc, argv, vbo);
cudaFree(d_vbo_buffer);
d_vbo_buffer = NULL;
}
else
{
if (false == initGL(&argc, argv))
return false;
// register callbacks
glutDisplayFunc(display);
glutKeyboardFunc(keyboard);
glutMouseFunc(mouse);
glutMotionFunc(motion);
glutCloseFunc(cleanup);
createVBO(&vbo, &cuda_vbo_resource, cudaGraphicsMapFlagsWriteDiscard); // VBO 생성
runCuda(&cuda_vbo_resource); // 병렬 처리
glutMainLoop();
}
return true;
}
void runCuda(struct cudaGraphicsResource **vbo_resource)
{
// map OpenGL buffer object for writing from CUDA
float4 *dptr;
checkCudaErrors(cudaGraphicsMapResources(1, vbo_resource, 0));
size_t num_bytes;
checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void **)&dptr, &num_bytes,
*vbo_resource));
// execute the kernel
// dim3 block(8, 8, 1);
// dim3 grid(mesh_width / block.x, mesh_height / block.y, 1);
// kernel<<< grid, block>>>(dptr, mesh_width, mesh_height, g_fAnim);
launch_kernel(dptr, mesh_width, mesh_height, g_fAnim); // 병렬 처리 계산
// unmap buffer object
checkCudaErrors(cudaGraphicsUnmapResources(1, vbo_resource, 0));
}
#ifdef _WIN32
#ifndef FOPEN
#define FOPEN(fHandle,filename,mode) fopen_s(&fHandle, filename, mode)
#endif
#else
#ifndef FOPEN
#define FOPEN(fHandle,filename,mode) (fHandle = fopen(filename, mode))
#endif
#endif
void sdkDumpBin2(void *data, unsigned int bytes, const char *filename)
{
printf("sdkDumpBin: <%s>\n", filename);
FILE *fp;
FOPEN(fp, filename, "wb");
fwrite(data, bytes, 1, fp);
fflush(fp);
fclose(fp);
}
// CUDA 계산
void runAutoTest(int devID, char **argv, char *ref_file)
{
char *reference_file = NULL;
void *imageData = malloc(mesh_width*mesh_height*sizeof(float));
// execute the kernel
launch_kernel((float4 *)d_vbo_buffer, mesh_width, mesh_height, g_fAnim);
cudaDeviceSynchronize();
getLastCudaError("launch_kernel failed");
checkCudaErrors(cudaMemcpy(imageData, d_vbo_buffer, mesh_width*mesh_height*sizeof(float), cudaMemcpyDeviceToHost));
sdkDumpBin2(imageData, mesh_width*mesh_height*sizeof(float), "simpleGL.bin");
reference_file = sdkFindFilePath(ref_file, argv[0]);
if (reference_file &&
!sdkCompareBin2BinFloat("simpleGL.bin", reference_file,
mesh_width*mesh_height*sizeof(float),
MAX_EPSILON_ERROR, THRESHOLD, pArgv[0]))
{
g_TotalErrors++;
}
}
void createVBO(GLuint *vbo, struct cudaGraphicsResource **vbo_res,
unsigned int vbo_res_flags)
{
assert(vbo);
// create buffer object
glGenBuffers(1, vbo);
glBindBuffer(GL_ARRAY_BUFFER, *vbo);
// initialize buffer object
unsigned int size = mesh_width * mesh_height * 4 * sizeof(float);
glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);
glBindBuffer(GL_ARRAY_BUFFER, 0);
// register this buffer object with CUDA
checkCudaErrors(cudaGraphicsGLRegisterBuffer(vbo_res, *vbo, vbo_res_flags));
SDK_CHECK_ERROR_GL();
}
void deleteVBO(GLuint *vbo, struct cudaGraphicsResource *vbo_res)
{
// unregister this buffer object with CUDA
checkCudaErrors(cudaGraphicsUnregisterResource(vbo_res));
glBindBuffer(1, *vbo);
glDeleteBuffers(1, vbo);
*vbo = 0;
}
// CALL BACK 함수
void display()
{
sdkStartTimer(&timer);
// run CUDA kernel to generate vertex positions
runCuda(&cuda_vbo_resource);
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
// set view matrix
glMatrixMode(GL_MODELVIEW);
glLoadIdentity();
glTranslatef(0.0, 0.0, translate_z);
glRotatef(rotate_x, 1.0, 0.0, 0.0);
glRotatef(rotate_y, 0.0, 1.0, 0.0);
// render from the vbo
glBindBuffer(GL_ARRAY_BUFFER, vbo);
glVertexPointer(4, GL_FLOAT, 0, 0);
glEnableClientState(GL_VERTEX_ARRAY);
glColor3f(1.0, 0.0, 0.0);
glDrawArrays(GL_POINTS, 0, mesh_width * mesh_height);
glDisableClientState(GL_VERTEX_ARRAY);
glutSwapBuffers();
g_fAnim += 0.01f;
sdkStopTimer(&timer);
computeFPS();
}
void timerEvent(int value)
{
if (glutGetWindow())
{
glutPostRedisplay();
glutTimerFunc(REFRESH_DELAY, timerEvent,0);
}
}
void cleanup()
{
sdkDeleteTimer(&timer);
if (vbo)
{
deleteVBO(&vbo, cuda_vbo_resource);
}
}
void keyboard(unsigned char key, int /*x*/, int /*y*/)
{
switch (key)
{
case (27) :
#if defined(__APPLE__) || defined(MACOSX)
exit(EXIT_SUCCESS);
#else
glutDestroyWindow(glutGetWindow());
return;
#endif
}
}
void mouse(int button, int state, int x, int y)
{
if (state == GLUT_DOWN)
{
mouse_buttons |= 1<<button;
}
else if (state == GLUT_UP)
{
mouse_buttons = 0;
}
mouse_old_x = x;
mouse_old_y = y;
}
void motion(int x, int y)
{
float dx, dy;
dx = (float)(x - mouse_old_x);
dy = (float)(y - mouse_old_y);
if (mouse_buttons & 1)
{
rotate_x += dy * 0.2f;
rotate_y += dx * 0.2f;
}
else if (mouse_buttons & 4)
{
translate_z += dy * 0.01f;
}
mouse_old_x = x;
mouse_old_y = y;
}
void checkResultCuda(int argc, char **argv, const GLuint &vbo)
{
if (!d_vbo_buffer)
{
checkCudaErrors(cudaGraphicsUnregisterResource(cuda_vbo_resource));
// map buffer object
glBindBuffer(GL_ARRAY_BUFFER, vbo);
float *data = (float *) glMapBuffer(GL_ARRAY_BUFFER, GL_READ_ONLY);
// check result
if (checkCmdLineFlag(argc, (const char **) argv, "regression"))
{
// write file for regression test
sdkWriteFile<float>("./data/regression.dat",
data, mesh_width * mesh_height * 3, 0.0, false);
}
// unmap GL buffer object
if (!glUnmapBuffer(GL_ARRAY_BUFFER))
{
fprintf(stderr, "Unmap buffer failed.\n");
fflush(stderr);
}
checkCudaErrors(cudaGraphicsGLRegisterBuffer(&cuda_vbo_resource, vbo,
cudaGraphicsMapFlagsWriteDiscard));
SDK_CHECK_ERROR_GL();
}
}
실행 결과는 다음과 같다.
이외, 파티클 시뮬레이션 등을 실행해 보면, 병렬 계산되어 실시간으로 렌더링되는 것을 볼 수 있다.
