^(코딩캣)^ = @"코딩"하는 고양이;

OpenCL 기본 실행 예 (1)

API/OpenCL
2019. 2. 2. 22:25

OpenCL 기본 실행 예 (1)


OpenCL은 실행을 위한 세팅 과정이 너무 많다. 따라서 본 포스팅을 통해 OpenCL을 실행하는 예시 코드를 정리해 둔다.

 

0 단계. 사용 가능한 플랫폼 개수 확인하기.


OpenCL을 사용할 수 있는 플랫폼(CPU와 GPU 등 디바이스 조합)이 총 몇 개인지를 확인하기 위해 clGetPlatformIDs 함수를 사용한다. 2 번째 매개변수인 cl_platform_id * platformsNULL을 대입하면 플랫폼의 총 개수가 반환된다. 이 때 플랫폼의 개수는 1 번째 매개변수인 cl_uint num_entries 이하로 반환한다. 그러므로 최대의 개수를 얻고자 하면 1 번째 매개변수에 -1을 대입하면 된다. 이 매개변수의 자료형은 부호 없는 정수(cl_uint)이기 때문에 -10xFFFF, 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개)를 확인할 수 있다.

 

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); // 플랫폼 구조체 동적할당을 해제한다.

 

이 PC에서는 컨텍스트에 2개의 장치가 있음을 확인 가능하다.

 

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); // 플랫폼 구조체 동적할당을 해제한다.

 

OpenCL 소스 컴파일 결과

 

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");
/* ... */

 

OpenCL 소스 코드의 실행 결과