OpenCL 기본 실행 예 (1)
OpenCL은 실행을 위한 세팅 과정이 너무 많다. 따라서 본 포스팅을 통해 OpenCL을 실행하는 예시 코드를 정리해 둔다.
0 단계. 사용 가능한 플랫폼 개수 확인하기.
OpenCL을 사용할 수 있는 플랫폼(CPU와 GPU 등 디바이스 조합)이 총 몇 개인지를 확인하기 위해 clGetPlatformIDs 함수를 사용한다. 2 번째 매개변수인 cl_platform_id * platforms에 NULL을 대입하면 플랫폼의 총 개수가 반환된다. 이 때 플랫폼의 개수는 1 번째 매개변수인 cl_uint num_entries 이하로 반환한다. 그러므로 최대의 개수를 얻고자 하면 1 번째 매개변수에 -1을 대입하면 된다. 이 매개변수의 자료형은 부호 없는 정수(cl_uint)이기 때문에 -1은 0xFFFF, 0xFFFFFFFF, 0xFFFFFFFFFFFFFFFF 중에 하나이고 부호 없는 정수의 최댓값과 같다.
사용 가능한 플랫폼의 수는 3 번째 매개변수인 cl_uint * num_platforms으로 반환된다. 함수 자체가 반환하는 값은 함수의 실행 결과로서, 정상 작동했다면 CL_SUCCESS를 반환한다.
/* C source */
// 새로 추가된 변수
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
/* ... */
// 0 단계. 시스템으로부터 OpenCL이 사용 가능한 플랫폼의 수를 확인한다.
errNo = clGetPlatformIDs
(
(cl_uint)(-1), // cl_uint num_entries
NULL, // cl_platform_id * platforms
&nPlatformId // cl_uint * num_platforms
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clGetPlatformIDs = %d\n", errNo);
// 실패 시 작업
}
else if (nPlatformId <= 0)
{
printf("[FAILURE] clGetPlatformIDs: nPlatformId = %u", nPlatformId);
// 실패 시 작업
}
else
{
printf("[SUCCEED] clGetPlatformIDs: nPlatformId = %u\n", nPlatformId);
}
/* ... */
1 단계. 사용 가능한 플랫폼 목록 가져오기.
앞서 0 단계에서 플랫폼 수를 확인하였다면, 메모리를 동적할당하여 플랫폼 목록을 가져온다. 마찬가지로 clGetPlatformIDs를 사용한다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
/* ... */
// nPlatformId개의 구조체를 동적할당한다.
assert(lpPlatformId = (cl_platform_id *)calloc(nPlatformId, sizeof(cl_platform_id)));
// 1 단계. 플랫폼 목록을 가져온다.
errNo = clGetPlatformIDs
(
nPlatformId, // cl_uint num_entries
lpPlatformId, // cl_platform_id * platforms
NULL // cl_uint * num_platforms
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clGetPlatformIDs = %d\n", errNo);
// 실패 시 작업
}
else
{
printf("[SUCCEED] clGetPlatformIDs: lpPlatformId = %p\n", lpPlatformId);
}
/* ... */
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
2 단계. 컨텍스트 생성
컨텍스트context는 OpenCL에서 사용 가능한 CPU + GPU + 메모리 객체 등등을 아우르는 개념이다. OpenCL은 컨텍스트 속에서 작동하는 구조로 되어 있으므로 이를 얻는 것이 가장 중요하다.
본 포스팅에서는 컨텍스트를 얻기 위해서 clCreateContextFromType 함수를 사용하겠다.
2-1 단계. 컨텍스트 프로퍼티 구성
컨텍스트 프로퍼티context property는 clCreateContextFromType에서 요구하는 매개변수로서 다음과 같은 구조의 배열이다.
{ CL_CONTEXT_PLATFORM, 플랫폼 1, CL_CONTEXT_PLATFORM, 플랫폼 2, ..., 0 }
마지막 원소는 항상 0이어야 한다.
다음은 앞서 얻은 플랫폼 배열을 통해 컨텍스트 프로퍼티 배열을 구성하는 예이다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
// 새로 추가된 변수
size_t i, j;
cl_uint nContextProperties = 0;
cl_context_properties * lpContextProperties = NULL;
cl_context context;
/* ... */
assert((nContextProperties = nPlatformId * 2 + 1) > 1);
assert(lpContextProperties = (cl_context_properties *)calloc(nContextProperties, sizeof(cl_context_properties)));
for (i = 0, j = 0; i < nContextProperties;)
{
if (i < nContextProperties - 1)
{
// 마지막 원소가 아니라면...
// 플랫폼에 대해 질의한다.
lpContextProperties[i++] = (cl_context_properties)CL_CONTEXT_PLATFORM;
// 질의하고 싶은 플랫폼을 지정한다.
lpContextProperties[i++] = (cl_context_properties)lpPlatformId[j++];
}
else
{
// 마지막 원소는 항상 zero
lpContextProperties[i] = (cl_context_properties)0;
break;
}
}
/* ... */
free(lpContextProperties); // 컨텍스트 속성 배열을 해제한다.
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
2-2 단계. 컨텍스트 객체 생성
앞서 구성한 컨텍스트 프로퍼티 배열을 가지고 OpenCL 컨텍스트 객체를 생성한다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
// 새로 추가된 변수
size_t i, j;
cl_uint nContextProperties = 0;
cl_context_properties * lpContextProperties = NULL;
cl_context context;
/* ... */
context = clCreateContextFromType
(
lpContextProperties, // cl_context_properties * properties
CL_DEVICE_TYPE_ALL, // cl_device_type device_type
NULL, // void (*pfn_notify) (const char *errinfo, const void *private_info, size_t cb, void *user_data)
NULL, // void * user_data
&errNo // cl_int * errcode_ret
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clCreateContextFromType = %d\n", errNo);
free(lpContextProperties); // 컨텍스트 얻기에 사용된 배열 해제
// 실패 시 작업
}
else
{
printf("[SUCCEED] clCreateContextFromType: context = %p\n", context);
}
/* ... */
clReleaseContext(context); // 컨텍스트를 해제한다.
free(lpContextProperties); // 컨텍스트 속성 배열을 해제한다.
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
3 단계. 컨텍스트에 포함된 장치들 얻기
컨텍스트에는 CPU, GPU를 포함한 가속 장치들이 포함되어 있다. 현재 컨텍스트에 몇 개의 장치들이 있고, 각각의 장치가 무엇인지 확인하기 위해 다음과 같이 코드를 작성한다.
3-1. 장치 정보를 얻기 위한 버퍼의 크기 구하기
컨텍스트에 포함된 장치들의 목록을 얻기 위해 총 몇 바이트의 버퍼가 필요한지 확인하는 과정이다. clGetContextInfo를 사용하며 매개변수에 param_value_size = 0, param_value = NULL을 하고 param_value_size_ret =&변수를 하면 변수를 통해 메모리의 크기가 반환된다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
size_t i, j;
cl_uint nContextProperties = 0;
cl_context_properties * lpContextProperties = NULL;
cl_context context;
// 새로 추가된 변수
size_t cbDeviceBuffer = 0;
size_t nDeviceId = 0;
cl_device_id * lpDeviceId = NULL;
/* ... */
errNo = clGetContextInfo
(
context, // cl_context context
CL_CONTEXT_DEVICES, // cl_context_info param_name
0, // size_t param_value_size
NULL, // void * param_value
&cbDeviceBuffer // size_t * param_value_size_ret
);
if (errNo != CL_SUCCESS)
{
fprintf(stderr, "[FAILURE] clGetContextInfo = %d\n", errNo);
// 실패 시 작업
}
else if (cbDeviceBuffer <= 0)
{
fprintf(stderr, "[FAILURE] deviceBufferSize = %zu\n", cbDeviceBuffer);
// 실패 시 작업
}
else
{
fprintf(stderr, "[SUCCESS] clGetContextInfo: cbDeviceBuffer = %zu\n", cbDeviceBuffer);
}
/* ... */
clReleaseContext(context); // 컨텍스트를 해제한다.
free(lpContextProperties); // 컨텍스트 얻기에 사용된 배열 해제
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
3-2. 컨텍스트에 있는 장치들의 정보 가져오기
장치 목록들을 얻기 위해 몇 바이트가 필요한지를 위와 같이 clGetContextInfo 함수로 확인하였다. 다시 호출하여 장치 목록을 복사해온다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
size_t i, j;
cl_uint nContextProperties = 0;
cl_context_properties * lpContextProperties = NULL;
cl_context context;
// 새로 추가된 변수
size_t cbDeviceBuffer = 0;
size_t nDeviceId = 0;
cl_device_id * lpDeviceId = NULL;
/* ... */
assert(nDeviceId = (cl_uint)(cbDeviceBuffer / sizeof(cl_device_id)));
assert(lpDeviceId = (cl_device_id *)calloc(nDeviceId, sizeof(cl_device_id)));
errNo = clGetContextInfo
(
context, // cl_context context
CL_CONTEXT_DEVICES, // cl_context_info param_name
cbDeviceBuffer, // size_t param_value_size
lpDeviceId, // void * param_value
NULL // size_t * param_value_size_ret
);
if (errNo != CL_SUCCESS)
{
fprintf(stderr, "[FAILURE] clGetContextInfo = %d\n", errNo);
// 실패 시 작업
}
else
{
printf("[SUCCESS] clGetContextInfo: nDeviceId = %zu\n", nDeviceId);
}
/* ... */
free(lpDeviceId); // 장치 버퍼 메모리 해제
clReleaseContext(context); // 컨텍스트를 해제한다.
free(lpContextProperties); // 컨텍스트 속성 배열을 해제한다.
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
3-3. 컨텍스트에 있는 장치들에 대한 커맨드 큐 객체 생성
커맨드 큐command queue는 연산장치가 수행할 연산들이 대기하는 장소이다. clCreateCommandQueue 함수로 컨텍스트에 있는 각 장치들의 커맨트 큐 객체를 생성한다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
size_t i, j;
cl_uint nContextProperties = 0;
cl_context_properties * lpContextProperties = NULL;
cl_context context;
size_t cbDeviceBuffer = 0;
size_t nDeviceId = 0;
cl_device_id * lpDeviceId = NULL;
// 새로 추가된 변수
size_t nCommandQueue = 0;
cl_command_queue * lpCommandQueue = NULL;
/* ... */
assert((nCommandQueue = nDeviceId) > 0);
assert(lpCommandQueue = (cl_command_queue *)calloc(nCommandQueue, sizeof(cl_command_queue)));
for (i = 0; i < nCommandQueue; i++)
{
lpCommandQueue[i] = clCreateCommandQueue
(
context, //cl_context context
lpDeviceId[i], // cl_device_id device
0, // cl_command_queue_properties properties
&errNo // cl_int * errcode_ret
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clCreateCommandQueue = %d\n", errNo);
// 실패 시 작업
}
else if (lpCommandQueue[i] == NULL)
{
printf("[FAILURE] commandQuene == %p @ %zu\n", lpCommandQueue[i], i);
// 실패 시 작업
}
else
{
printf("[SUCCESS] clCreateCommandQueue: commandQueue = %p @ %zu\n", lpCommandQueue[i], i);
}
}
/* ... */
free(lpDeviceId); // 장치 버퍼 메모리 해제
clReleaseContext(context); // 컨텍스트를 해제한다.
free(lpContextProperties); // 컨텍스트 속성 배열을 해제한다.
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
4 단계. OpenCL 소스 파일을 읽어서 컴파일
OpenCL 소스는 기본적으로 텍스트 상태로 배포된다. 사용자의 다양한 아키텍처에 맞추어 즉석으로 컴파일하기 때문이다. 이 때 사용되는 함수는 clCreateProgramWithSource, clBuildProgram이다.
4-1 단계. OpenCL 소스 작성 후 읽기
OpenCL 소스를 작성한다. 파일 이름은 hello.cl로 정한다.
/* OpenCL source */
// ptr1 배열의 각 원소와 ptr2 배열의 각 원소를 더하여 out으로 리턴한다.
__kernel void hello(__global float * out, __global const float * ptr1, __global const float * ptr2)
{
int global_id = 0; // 배열의 몇 번째 원소를 참조할 것인지를 나타내는 인덱스 변수이다.
global_id = get_global_id(0);
out[global_id] = (ptr1[global_id] + ptr2[global_id]);
return;
}
위 소스 파일을 C로 읽어들인다.
/* C source */
FILE * fp = NULL;
size_t cbSource = 0;
char * lpSource = NULL;
if (fp = fopen("./hello.cl", "rt")) // 파일을 읽기 모드로 연다
{
fseek(fp, 0, SEEK_END); // 파일이 몇 바이트인지를 체크한다.
cbSource = ftell(fp) + 1;
lpSource = (char *)calloc(cbSource, sizeof(char));
fseek(fp, 0, SEEK_SET); // 파일을 처음부터 끝까지 읽는다.
fread(lpSource, sizeof(char), cbSource, fp);
}
fclose(fp);
4-2 단계. OpenCL 소스로부터 프로그램 객체를 생성
위의 단계에서 읽어들인 소스 파일을 컴파일한다. 컴파일에 앞서 소스 문자열로부터 프로그램 객체를 생성해야 한다. 소스 문자열로부터 프로그램 객체를 생성하는 함수는 clCreateProgramWithSource이다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
size_t i, j;
cl_uint nContextProperties = 0;
cl_context_properties * lpContextProperties = NULL;
cl_context context;
size_t cbDeviceBuffer = 0;
size_t nDeviceId = 0;
cl_device_id * lpDeviceId = NULL;
size_t nCommandQueue = 0;
cl_command_queue * lpCommandQueue = NULL;
// 새로 추가된 변수
cl_program program;
/* ... */
program = clCreateProgramWithSource
(
context, // cl_context context
1, // cl_uint count
(const char **)&sourceContent, // const char **strings
&sourceContentSize, // const size_t *lengths
&errNum // cl_int *errcode_ret
);
if (errNum != CL_SUCCESS)
{
printf("[FAILURE] clCreateProgramWithSource == %d\n", errNum);
// 실패 시 작업
}
else
{
printf("[SUCCESS] clCreateProgramWithSource\n");
}
/* ... */
clReleaseProgram(program); // 프로그램 객체 해제
free(lpDeviceId); // 장치 버퍼 메모리 해제
clReleaseContext(context); // 컨텍스트를 해제한다.
free(lpContextProperties); // 컨텍스트 속성 배열을 해제한다.
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
4-3 단계. OpenCL 소스 컴파일
앞서 생성한 프로그램 객체를 컴파일한다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
size_t i, j;
cl_uint nContextProperties = 0;
cl_context_properties * lpContextProperties = NULL;
cl_context context;
size_t cbDeviceBuffer = 0;
size_t nDeviceId = 0;
cl_device_id * lpDeviceId = NULL;
size_t nCommandQueue = 0;
cl_command_queue * lpCommandQueue = NULL;
cl_program program;
/* ... */
errNo = clBuildProgram
(
program, // cl_program program : 컴파일하고자 하는 프로그램 객체
0, // cl_uint num_devices : 컴파일 대상이 될 장치의 개수 (0개 = 컨텍스트에 있는 기본 장치들)
NULL, // const cl_device_id * device_list : 컴파일 대상이 될 장치 (NULL = 컨텍스트에 있는 기본 장치들)
NULL, // const char * options : 빌드 옵션
NULL, // void (*pfn_notify)(cl_program, void * user_data) : 콜백 함수
NULL // void * user_data : 콜백함수에 전달할 사용자 데이터
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clBuildProgram == %d\n", errNo);
// 실패 시 작업
}
else
{
printf("[SUCCESS] clBuildProgram\n");
}
/* ... */
clReleaseProgram(program); // 프로그램 객체 해제
free(lpDeviceId); // 장치 버퍼 메모리 해제
clReleaseContext(context); // 컨텍스트를 해제한다.
free(lpContextProperties); // 컨텍스트 속성 배열을 해제한다.
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
5 단계. OpenCL에서 선언한 함수 실행하기
커널kernel은 OpenCL 소스에서 선언한 각 함수들의 진입점에 해당하는 개념이다. 앞서 컴파일한 OpenCL 소스에서 특정 함수의 이름을 통해 그 커널을 가져와 호출한다. 함수가 매개변수를 필요로 한다면 메모리 객체를 별도로 구성해야 한다.
5-1 단계. 커널 얻기
hello.cl 소스 파일에서 선언한 한 함수 hello를 호출해보도록 하겠다. 이를 위해 커널 객체를 생성한다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
size_t i, j;
cl_uint nContextProperties = 0;
cl_context_properties * lpContextProperties = NULL;
cl_context context;
size_t cbDeviceBuffer = 0;
size_t nDeviceId = 0;
cl_device_id * lpDeviceId = NULL;
size_t nCommandQueue = 0;
cl_command_queue * lpCommandQueue = NULL;
cl_program program;
// 새로 추가된 변수
cl_kernel kernel;
/* ... */
kernel = clCreateKernel
(
program, // cl_program program
"hello", // const char * kernel_name
&errNo // cl_int * errcode_ret
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clCreateKernel = %d\n", errNo);
// 실패 시 작업
}
else if (kernel == NULL)
{
printf("[FAILURE] clCreateKernel: kernel = %p\n", kernel);
// 실패 시 작업
}
else
{
printf("[SUCCESS] clCreateKernel\n");
}
/* ... */
clReleaseKernel(kernel); // 커널 객체 해제
clReleaseProgram(program); // 프로그램 객체 해제
free(lpDeviceId); // 장치 버퍼 메모리 해제
clReleaseContext(context); // 컨텍스트를 해제한다.
free(lpContextProperties); // 컨텍스트 속성 배열을 해제한다.
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
5-2 단계. 매개변수 구성하기
OpenCL 소스에서 선언한 hello 함수에서는 float * 형식의 3개의 매개변수(out, ptr1, ptr2)가 필요하다. 이를 구성해보겠다.
먼저 다음과 같이 테스트 배열을 마련한다. 하나는 1부터 50까지이고 다른 하나는 51부터 100까지의 각 50개의 실수로 구성된 배열이다.
/* C source */
float out[50] = { 0.0f, };
float ptr1[50] = { 0.0f, };
float ptr2[50] = { 0.0f, };
size_t i = 0;
for (i = 0; i < 50; i++)
{
ptr1[i] = (1.0f + (float)i);
ptr2[i] = (51.0f + (float)i);
}
C의 배열을 OpenCL 소스에서 작성한 함수로 보내기 위해서는 메모리 객체를 생성한다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
size_t i, j;
cl_uint nContextProperties = 0;
cl_context_properties * lpContextProperties = NULL;
cl_context context;
size_t cbDeviceBuffer = 0;
size_t nDeviceId = 0;
cl_device_id * lpDeviceId = NULL;
size_t nCommandQueue = 0;
cl_command_queue * lpCommandQueue = NULL;
cl_program program;
cl_kernel kernel;
// 새로 추가된 변수
cl_mem memOut;
cl_mem memPtr1;
cl_mem memPtr2;
/* ... */
memOut = clCreateBuffer
(
context, // cl_context context : 배열을 사용할 컨텍스트
CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, // cl_mem_flags flags : 각종 속성
sizeof(out), // size_t size : 배열의 크기 (단위 = 바이트)
out, // void * host_ptr : 배열
&errNo // cl_int * errcode_ret : 오류 코드
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clCreateBuffer = %d", errNo);
// 실패 시 작업
}
else if (memOut == NULL)
{
printf("[FAILURE] clCreateBuffer: memOut = %p", memOut);
// 실패 시 작업
}
else
{
printf("[SUCCESS] clCreateBuffer\n");
}
memOut = clCreateBuffer
(
context, // cl_context context
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, // cl_mem_flags flags
sizeof(ptr1), // size_t size
ptr1, // void * host_ptr
&errNo // cl_int * errcode_ret
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clCreateBuffer = %d", errNo);
// 실패 시 작업
}
else if (memOut == NULL)
{
printf("[FAILURE] clCreateBuffer: memOut = %p", memOut);
// 실패 시 작업
}
else
{
printf("[SUCCESS] clCreateBuffer\n");
}
memOut = clCreateBuffer
(
context, // cl_context context
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, // cl_mem_flags flags
sizeof(ptr2), // size_t size
ptr2, // void * host_ptr
&errNo // cl_int * errcode_ret
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clCreateBuffer = %d", errNo);
// 실패 시 작업
}
else if (memOut == NULL)
{
printf("[FAILURE] clCreateBuffer: memOut = %p", memOut);
// 실패 시 작업
}
else
{
printf("[SUCCESS] clCreateBuffer\n");
}
/* ... */
clReleaseMemObject(memOut); // float out[50]에 대한 메모리 객체 해제
clReleaseMemObject(memPtr1); // float ptr1[50]에 대한 메모리 객체 해제
clReleaseMemObject(memPtr2); // float ptr2[50]에 대한 메모리 객체 해제
clReleaseKernel(kernel); // 커널 객체 해제
clReleaseProgram(program); // 프로그램 객체 해제
free(lpDeviceId); // 장치 버퍼 메모리 해제
clReleaseContext(context); // 컨텍스트를 해제한다.
free(lpContextProperties); // 컨텍스트 속성 배열을 해제한다.
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
5-3 단계. 커널에 매개변수 전달하기
함수 진입점을 '커널'의 형태로 얻었고, 함수에 전달할 매개변수도 '메모리'의 형태로 구성하였다. 이제 커널과 매개변수를 연계해보겠다. OpenCL 소스에서 (out, ptr1, ptr2)의 순서로 선언하였으므로 이 순서대로 메모리 객체를 커널에 연계한다. 이 때 쓰이는 함수는 clSetKernelArg이다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
size_t i, j;
cl_uint nContextProperties = 0;
cl_context_properties * lpContextProperties = NULL;
cl_context context;
size_t cbDeviceBuffer = 0;
size_t nDeviceId = 0;
cl_device_id * lpDeviceId = NULL;
size_t nCommandQueue = 0;
cl_command_queue * lpCommandQueue = NULL;
cl_program program;
cl_kernel kernel;
// 새로 추가된 변수
cl_mem memOut;
cl_mem memPtr1;
cl_mem memPtr2;
/* ... */
errNo = clSetKernelArg
(
kernel, // cl_kernel kernel : 실행하고 싶은 커널
0, // cl_uint arg_index : 커널에 선언된 매개변수 순서 (0번째 = out)
sizeof(cl_mem), // size_t arg_size : 메모리 객체의 크기
&memOut // const void * arg_value : 버퍼
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clSetKernelArg = %d\n", errNo);
}
else
{
printf("[SUCCESS] clSetKernelArg\n");
}
errNo = clSetKernelArg
(
kernel, // cl_kernel kernel
1, // cl_uint arg_index : 1번째 = ptr1
sizeof(cl_mem), // size_t arg_size
&memPtr1 // const void * arg_value
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clSetKernelArg = %d\n", errNo);
}
else
{
printf("[SUCCESS] clSetKernelArg\n");
}
errNo = clSetKernelArg
(
kernel, // cl_kernel kernel
2, // cl_uint arg_index : 2번째 = ptr2
sizeof(cl_mem), // size_t arg_size
&memPtr2 // const void * arg_value
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clSetKernelArg = %d\n", errNo);
}
else
{
printf("[SUCCESS] clSetKernelArg\n");
}
/* ... */
clReleaseMemObject(memOut); // float out[50]에 대한 메모리 객체 해제
clReleaseMemObject(memPtr1); // float ptr1[50]에 대한 메모리 객체 해제
clReleaseMemObject(memPtr2); // float ptr2[50]에 대한 메모리 객체 해제
clReleaseKernel(kernel); // 커널 객체 해제
clReleaseProgram(program); // 프로그램 객체 해제
free(lpDeviceId); // 장치 버퍼 메모리 해제
clReleaseContext(context); // 컨텍스트를 해제한다.
free(lpContextProperties); // 컨텍스트 속성 배열을 해제한다.
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
5-4 단계. 커널을 호출하기
이제 커널을 실행한다. 커맨드 큐에 커널을 추가하는 것으로 커널은 실행 준비 상태가 된다. 앞서 장치의 수 만큼 여러개의 커맨드 큐를 얻었는데 0번째 커맨드 큐로 커널을 추가해 보겠다. 이 때 사용되는 함수는 clEnqueueNDRangeKernel이다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
size_t i, j;
cl_uint nContextProperties = 0;
cl_context_properties * lpContextProperties = NULL;
cl_context context;
size_t cbDeviceBuffer = 0;
size_t nDeviceId = 0;
cl_device_id * lpDeviceId = NULL;
size_t nCommandQueue = 0;
cl_command_queue * lpCommandQueue = NULL;
cl_program program;
cl_kernel kernel;
cl_mem memOut;
cl_mem memPtr1;
cl_mem memPtr2;
// 새로 추가된 변수
// 커널이 수행할 총 작업량을 어떻게 분배할 것인지 그 크기를 지정한다. 이 분배는 최적화 방식에 따라 달라질 수 있다.
size_t globalWorkSize[1] = { 50 }; // 각 배열이 가지고 있는 50개의 원소
size_t localWorkSize[1] = { 1 }; // 하나의 디바이스에 몰아서 처리
/* ... */
errNo = clEnqueueNDRangeKernel
(
lpCommandQueue[0], // cl_command_queue command_queue
kernel, // cl_kernel kernel
1, // cl_uint work_dim
NULL, // const size_t * global_work_offset
globalWorkSize, // const size_t * global_work_size
localWorkSize, // const size_t * local_work_size
0, // cl_uint num_events_in_wait_list
NULL, // const cl_event * event_wait_list
NULL // cl_event * event
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clEnqueueNDRangeKernel = %d\n", errNo);
// 실패 시 작업
}
else
{
printf("[SUCCESS] clEnqueueNDRangeKernel\n");
}
/* ... */
clReleaseMemObject(memOut); // float out[50]에 대한 메모리 객체 해제
clReleaseMemObject(memPtr1); // float ptr1[50]에 대한 메모리 객체 해제
clReleaseMemObject(memPtr2); // float ptr2[50]에 대한 메모리 객체 해제
clReleaseKernel(kernel); // 커널 객체 해제
clReleaseProgram(program); // 프로그램 객체 해제
free(lpDeviceId); // 장치 버퍼 메모리 해제
clReleaseContext(context); // 컨텍스트를 해제한다.
free(lpContextProperties); // 컨텍스트 속성 배열을 해제한다.
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
6 단계. 커널의 실행 결과 읽기
커널에서는 out 매개 변수를 통해 값을 함수 밖으로 전달한다. 이 매개변수는 C 소스에서 cl_mem 형식으로 전달된다. cl_mem으로부터 배열을 읽어올 때 사용되는 함수는 clEnqueueReadBuffer이다.
/* C source */
cl_int errNo;
cl_uint nPlatformId = 0;
cl_platform_id * lpPlatformId = NULL;
size_t i, j;
cl_uint nContextProperties = 0;
cl_context_properties * lpContextProperties = NULL;
cl_context context;
size_t cbDeviceBuffer = 0;
size_t nDeviceId = 0;
cl_device_id * lpDeviceId = NULL;
size_t nCommandQueue = 0;
cl_command_queue * lpCommandQueue = NULL;
cl_program program;
cl_kernel kernel;
float out[50] = { 0.0f, };
float ptr1[50] = { 0.0f, };
float ptr2[50] = { 0.0f, };
cl_mem memOut;
cl_mem memPtr1;
cl_mem memPtr2;
size_t globalWorkSize[1] = { 50 };
size_t localWorkSize[1] = { 1 };
/* ... */
errNo = clEnqueueReadBuffer
(
lpCommandQueue[0], // cl_command_queue command_queue
memOut, // cl_mem buffer
CL_TRUE, // cl_bool blocking_read
0, // size_t offset
sizeof(out), // size_t cb
out, // void * ptr
0, // cl_uint num_events_in_wait_list
NULL, // const cl_event * event_wait_list
NULL // cl_event * event
);
if (errNo != CL_SUCCESS)
{
printf("[FAILURE] clEnqueueReadBuffer = %d\n", errNo);
// 실패 시 작업
}
else
{
printf("[SUCCESS] clEnqueueReadBuffer\n");
}
/* ... */
clReleaseMemObject(memOut); // float out[50]에 대한 메모리 객체 해제
clReleaseMemObject(memPtr1); // float ptr1[50]에 대한 메모리 객체 해제
clReleaseMemObject(memPtr2); // float ptr2[50]에 대한 메모리 객체 해제
clReleaseKernel(kernel); // 커널 객체 해제
clReleaseProgram(program); // 프로그램 객체 해제
free(lpDeviceId); // 장치 버퍼 메모리 해제
clReleaseContext(context); // 컨텍스트를 해제한다.
free(lpContextProperties); // 컨텍스트 속성 배열을 해제한다.
free(lpPlatformId); // 플랫폼 구조체 동적할당을 해제한다.
6-1. 결과 확인하기
OpenCL 커널의 연산 결과는 clEnqueueReadBuffer를 통해 float out[50];로 복사되었다. 출력해보면 OpenCL 소스에서 작성하였던 ptr1[x] * 100.0f + ptr2[x]의 결과가 출력됨을 확인할 수 있다.
/* C source */
/* ... */
for (i = 0; i < 50; i++)
{
printf("%6.1f ", out[i]);
if ((i + 1) % 8 == 0) printf("\n");
}
printf("\n");
/* ... */
