이어지는 OpenCL 코드 작성
5. 커널 실행
이전 포스트에서 프로그램을 빌드하고 clCreateKernel 함수로 커널 구조체를 생성했습니다. 커널을 실행하기 위해서는 작성한 커널의 매개변수를 설정해 주어야 합니다. 예를 들어 이전 포스트에서 작성한 아래 커널 소스는 매개변수로 메모리 오브젝트를 필요로 합니다.
// example.cl
__kernel void get_unit_idx(__global int *out)
{
int idx = get_global_id(0);
out[idx] = idx;
}커널의 매개변수를 설정하는 함수 clSetKernelArg 는 다음과 같습니다.
cl_int clSetKernelArg(
cl_kernel kernel, // 커널 구조체
cl_uint arg_index, // 매개변수의 인덱스
size_t arg_size, // 매개변수의 크기
const void *arg_value // 매개변수의 포인터
)커널 get_unit_idx 의 매개변수로서 메모리 오브젝트를 전달한다면 다음과 같이 쓸 수 있습니다.
clSetKernelArg(kernel, 0, sizeof(dev_mem_obj), &dev_mem_obj);메모리 오브젝트가 아닌 데이터를 호스트 프로그램에서 직접 매개변수로 전달할 수 도 있습니다. 예를 들어 결과값이 될 메모리 각각에 호스트 프로그램에서 전달한 값을 저장 한다면 다음과 같이 소스를 작성하고 두 개의 매개변수를 전달합니다.
// example.cl
__kernel void set_x(__global int *out, int x)
{
int idx = get_global_id(0);
out[idx] = x;
}int x = 1;
clSetKernelArg(kernel, 0, sizeof(dev_mem_obj), &dev_mem_obj);
clSetKernelArg(kernel, 1, sizeof(x), &x);
매개변수 설정이 완료되면 실행될 커널을 커맨드 큐에 추가합니다.
cl_int clEnqueueNDRangeKernel(
cl_command_queue command_queue, // 커맨드 큐 구조체
cl_kernel kernel, // 커널 구조체
cl_uint work_dim, // 연산 유닛들이 자리한 grid 의 차원
const size_t *global_work_offset, // 오프셋(현재는 지원하지 않음)
const size_t *global_work_size, // 연산 유닛의 사이즈
const size_t *local_work_size, // 연산 유닛의 local 그룹별 사이즈
cl_uint num_events_in_wait_list, // 커맨드가 실행되기 전의 이벤트 커맨드 개수
const cl_event *event_wait_list,// 커맨드가 실행되기 전의 이벤트 커맨드 리스트
cl_event *event // 이벤트 오브젝트
)함수 clEnqueueNDRangeKernel 의 ND 는 N Dimension 을 뜻합니다. 커널 연산을 수행하는 유닛들은 global 계층에서의 인덱스, local 계층에서의 인덱스가 각각의 차원에서 정해집니다. 예를 들어 아래 그림의 연산 유닛 k는 global 계층에서 가로축에 대해 3, 세로축에 대해 2의 인덱스를 갖고 local 계층에서는 가로축에 대해 3, 세로축에 대해 0의 인덱스를 갖습니다.
위 그림처럼 연산 유닛들의 grid 를 형성하자면 다음과 같이 함수 clEnqueueNDRangeKernel 의 매개변수를 전달합니다.
size_t global_work_size = {10, 10};
size_t local_work_size = {5, 2};
clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
커널 get_unit_idx 의 get_global_id(dim) 함수는 매개변수 dim 을 차원의 인덱스로서 전달받습니다. 예를 들어 get_global_id(1) 은 연산 유닛의 global 계층상 두번째 차원에서의 인덱스를 반환합니다.
우리가 사용할 커널 get_unit_idx 은 첫번째 차원에서의 인덱스를 메모리에 저장하는 간단한 커널이므로 아래와 같이 1차원의 연산 유닛 grid 를 형성하겠습니다.
#define NUM_ELEMENTS 256
size_t global_work_size = NUM_ELEMENTS;
size_t local_work_size = 1;
clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);
clSetKernelArg 에 의해 커널의 매개변수로서 전달된 메모리 오브젝트는 연산이 끝난 후 호스트 프로그램에서 쓸 수 있는 메모리로 반환되어야 합니다. 이를 위해 커맨드 큐에 메모리 오브젝트를 호스트의 버퍼로 복사하는 커맨드를 추가합니다.
cl_int clEnqueueReadBuffer(
cl_command_queue command_queue, // 커맨드 큐 구조체
cl_mem buffer, // 읽어낼 메모리 오브젝트
cl_bool blocking_read,
/*
** TRUE : 메모리를 읽는 프로세스가 끝날 때 까지 함수가 결과를 반환하지 않음
** FALSE : 메모리를 읽는 프로세스가 끝날 때 까지 ptr 을 사용할 수 없음
*/
size_t offset, // 메모리 오브젝트를 읽어 들일 오프셋
size_t cb, // 읽어 들일 사이즈
void *ptr, // 메모리 오브젝트의 데이터가 복사될 호스트의 버퍼
cl_uint num_events_in_wait_list, // 커맨드가 실행되기 전의 이벤트 커맨드 개수
const cl_event *event_wait_list,// 커맨드가 실행되기 전의 이벤트 커맨드 리스트
cl_event *event // 이벤트 오브젝트
)데이터를 복사할 호스트의 버퍼를 만들어 메모리 오브젝트를 읽어 냅니다.
int *host_buf = (int *)malloc(sizeof(int) * NUM_ELEMENTS);
clEnqueueReadBuffer(command_queue, dev_mem_obj, CL_TRUE, 0,\
sizeof(*host_buf) * NUM_ELEMENTS, host_buf, 0, NULL, NULL);
커맨드 큐에 담긴 모든 커맨드는 아래 두 함수를 이용해 실행합니다.
cl_int clFlush(cl_command_queue command_queue)
cl_int clFinish(cl_command_queue command_queue)clFlush 는 큐의 모든 커맨드를 디바이스에 전달합니다. 디바이스에 전달된 커맨드들은 clFinish 함수에 의해 실행이 완료됩니다.
clFlush(command_queue);
clFinish(command_queue);
6. 각각의 구조체 해제
context, 커널 구조체, 프로그램 구조체, 메모리 오브젝트, 커맨드 큐 구조체는 병렬 연산이 끝나면 메모리를 해제해 주어야 합니다. 메모리 해제 함수는 clRelease... 형식을 갖고 아래와 같이 호출합니다.
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseMemObject(dev_mem_obj);
clReleaseCommandQueue(command_queue);
clReleaseContext(context);이전 포스트부터 지금까지 OpenCL 을 이용해 연산 유닛의 grid 상 인덱스를 출력하기까지의 과정을 알아보았습니다. 아래는 커널 소스와 지금까지의 과정을 다시 한번 정리한 코드입니다.
example.c
#include <OpenCL/cl.h>
#include <stdio.h>
#include <fcntl.h>
#define NUM_ELEMENTS 256
void example1(void)
{
int fd;
int *host_buf;
char *source_str;
cl_int ret;
cl_platform_id platform;
cl_device_id device;
cl_context context;
cl_command_queue command_queue;
cl_program program;
cl_kernel kernel;
cl_mem dev_mem_obj;
size_t global_work_size = NUM_ELEMENTS;
size_t local_work_size = 1;
host_buf = (int *)malloc(sizeof(int) * NUM_ELEMENTS);
/* 병렬 연산을 수행할 디바이스 ID 쿼리 */
clGetPlatformIDs(1, &platform, NULL);
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
/* context 와 커맨드 큐 구조체 생성 */
context = clCreateContext(NULL, 1, &device, NULL, NULL, &ret);
command_queue = clCreateCommandQueue(context, device, 0, &ret);
/*
** 아래는 .cl 파일의 소스를 문자열로 복사하는 코드입니다.
** get_file_content 와 같은 기능을 하는 함수는 어렵지 않게 만드실 수 있습니다.
*/
fd = open("kernels/setidx.cl", O_RDONLY);
source_str = get_file_content(fd);
close(fd);
/* 프로그램 구조체 생성 */
program = clCreateProgramWithSource(context, 1, (const char **)&source_str, NULL, &ret);
clBuildProgram(program, 1, &device, NULL, NULL, NULL);
/* 커널 구조체 생성 */
kernel = clCreateKernel(program, "get_unit_idx", &ret);
/* 메모리 오브젝트 생성 */
dev_mem_obj = clCreateBuffer(context, CL_MEM_WRITE_ONLY,\
sizeof(*host_buf) * NUM_ELEMENTS, NULL, &ret);
/* 커널 매개변수 설정 */
clSetKernelArg(kernel, 0, sizeof(dev_mem_obj), &dev_mem_obj);
/* 커맨드 큐에 커널과 메모리 읽는 커맨드 추가 */
clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,\
&global_work_size, &local_work_size, 0, NULL, NULL);
clEnqueueReadBuffer(command_queue, dev_mem_obj, CL_TRUE, 0,\
sizeof(*host_buf) * NUM_ELEMENTS, host_buf, 0, NULL, NULL);
/* 커맨드 큐에 추가된 커맨드 실행 */
clFlush(command_queue);
clFinish(command_queue);
/* 필요한 구조체 메모리 해제 */
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseMemObject(dev_mem_obj);
clReleaseCommandQueue(command_queue);
clReleaseContext(context);
for (size_t i=0; i < NUM_ELEMENTS; i++)
printf("%d ", host_buf[i]);
printf("\n");
free(host_buf);
}kernels/setidx.cl
__kernel void get_unit_idx(__global int *out)
{
int idx = get_global_id(0);
out[idx] = idx;
}main 함수에서 example1 함수를 호출하면 다음과 같이 출력됩니다.
0 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 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 123 124 125 126 127 128 129 130 131 132 133 134 135 136 137 138 139 140 141 142 143 144 145 146 147 148 149 150 151 152 153 154 155 156 157 158 159 160 161 162 163 164 165 166 167 168 169 170 171 172 173 174 175 176 177 178 179 180 181 182 183 184 185 186 187 188 189 190 191 192 193 194 195 196 197 198 199 200 201 202 203 204 205 206 207 208 209 210 211 212 213 214 215 216 217 218 219 220 221 222 223 224 225 226 227 228 229 230 231 232 233 234 235 236 237 238 239 240 241 242 243 244 245 246 247 248 249 250 251 252 253 254 255
이번에는 연산 유닛의 grid, local 상 인덱스를 쉽게 파악할 수 있는 커널을 만들어 보겠습니다. 각각의 차원의 인덱스를 이용해 아래와 같은 출력 결과를 얻고자 합니다.
2개의 차원을 사용하므로 아래와 같이 매크로 상수를 설정해 놓겠습니다.
#define NUM_ELEMENTS_X 16
#define NUM_ELEMENTS_Y 16
#define NUM_ELEMENTS (NUM_ELEMENTS_X * NUM_ELEMENTS_Y)clEnqueueNDRangeKernel 의 매개변수로 사용할 연산 유닛의 grid 크기와 local 그룹별 크기는 각각의 차원에 따라 크기를 정하기 위해 배열로 선언합니다.
size_t global_work_size[2] = {NUM_ELEMENTS_X, NUM_ELEMENTS_Y};
size_t local_work_size[2] = {4, 4};
이전의 예시와 마찬가지로 하나의 플랫폼으로부터 하나의 디바이스(GPU)를 사용할 예정입니다. 따라서 수정할 코드는 clEnqueueNDRangeKernel 의 매개변수와 커널 소스 파일의 이름뿐입니다.
/*
** 커널 소스를 문자열 변수에 직접 할당했거나
** 파일 디스크립터가 아닌 파일 포인터를 사용하셨다면
** 아래 코드의 수정 사항은 다를 수 있습니다.
*/
fd = open("kernels/set_group_index.cl", O_RDONLY);
...
/*
** 차원의 개수를 나타내는 매개변수를 변경합니다.
** 변수 global_work_size 와 local_work_size 가 배열이므로 포인터 연산자 & 을 지웁니다.
*/
clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL,\
global_work_size, local_work_size, 0, NULL, NULL);유닛 연산의 결과를 출력할 코드는 아래와 같이 수정합니다.
for (int i=0; i < NUM_ELEMENTS_Y; i++)
{
for (int j=0; j < NUM_ELEMENTS_X; j++)
printf("%3d ", host_buf[NUM_ELEMENTS_X * i + j]);
printf("\n");
}
이전에 하나의 차원에서 grid 상 인덱스를 얻기 위해 get_global_id 를 사용했다면 이번에는 각각의 차원과 local 그룹에서의 인덱스를 얻기 위해 다음과 같은 함수를 사용합니다.
get_group_id(dim) // dim 차원에서 local 그룹의 grid 계층상 인덱스
get_local_id(dim) // dim 차원에서 local 계층상 유닛의 인덱스
get_local_size(dim) // dim 차원에서 하나의 local 그룹의 유닛 개수
get_num_groups(dim) // dim 차원에서 local 그룹의 개수
위 함수들을 이용해 우리가 원하는 출력결과를 얻으려면 아래와 같이 커널 소스를 작성할 수 있습니다.
__kernel void set_group_index(__global int *out)
{
int idx_x = get_group_id(0) * get_local_size(0) + get_local_id(0);
int idx_y = get_group_id(1) * get_local_size(1) + get_local_id(1);
int row_size = get_num_groups(0) * get_local_size(0);
int idx = idx_y * row_size + idx_x;
int result = get_local_size(0) * get_local_size(1) *
(get_num_groups(0) * get_group_id(1) + get_group_id(0)) +
(get_local_id(1) * get_local_size(0) + get_local_id(0));
out[idx] = result;
}출력 결과는 아래와 같습니다.
0 1 2 3 16 17 18 19 32 33 34 35 48 49 50 51
4 5 6 7 20 21 22 23 36 37 38 39 52 53 54 55
8 9 10 11 24 25 26 27 40 41 42 43 56 57 58 59
12 13 14 15 28 29 30 31 44 45 46 47 60 61 62 63
64 65 66 67 80 81 82 83 96 97 98 99 112 113 114 115
68 69 70 71 84 85 86 87 100 101 102 103 116 117 118 119
72 73 74 75 88 89 90 91 104 105 106 107 120 121 122 123
76 77 78 79 92 93 94 95 108 109 110 111 124 125 126 127
128 129 130 131 144 145 146 147 160 161 162 163 176 177 178 179
132 133 134 135 148 149 150 151 164 165 166 167 180 181 182 183
136 137 138 139 152 153 154 155 168 169 170 171 184 185 186 187
140 141 142 143 156 157 158 159 172 173 174 175 188 189 190 191
192 193 194 195 208 209 210 211 224 225 226 227 240 241 242 243
196 197 198 199 212 213 214 215 228 229 230 231 244 245 246 247
200 201 202 203 216 217 218 219 232 233 234 235 248 249 250 251
204 205 206 207 220 221 222 223 236 237 238 239 252 253 254 255
OpenCL 의 API 를 이용해 연산 유닛의 계층 구조까지 알아보았습니다. 지금까지의 예시는 병렬 연산의 장점을 보여주지 않았습니다. 다음 포스트에서는 반복적으로 많은 연산을 필요로 하는 프랙탈 그리기 예제를 OpenCL API 를 이용해 구현해 보겠습니다.
