OpenCL 과 병렬 프로그래밍
OpenCL 은 연산 가능한 디바이스(CPU, GPU, DSP) 를 이용해 이종 간 병렬 프로그래밍을 가능하게 하는 소프트웨어입니다. 그래픽 작업과 같이 여러 오브젝트, 또는 픽셀에 대해 같은 연산이 이루어져야 할 때 유용하게 사용할 수 있습니다. OpenCL API 는 C 로 작성하지만 디바이스 안에서 수행되는 연산에 대해서는 OpenCL C 언어를 사용합니다. OpenCL C 는 C99 에 기초한 언어이고 C 언어에 익숙하다면 어렵지 않게 작성할 수 있습니다. OpenCL C 로 작성한 소스코드는 실행 커널(함수와 같음)로 컴파일 된 후 각각의 디바이스에 대해 연산이 수행됩니다.
OpenCL
컴파일된 프로그램은 CPU 가 데이터를 연산하고 메모리에 읽고 쓰면서 실행됩니다. 기존의 연산은 모두 CPU 의 한 코어에서 이루어졌습니다. 병렬 연산은 멀티 쓰레드를 지원하는 GPU 또는 CPU 를 이용합니다. OpenCL 아키텍쳐에서는 프로그램의 전체 프로세스를 담당하는 곳을 호스트, 병렬 연산을 수행하는 곳을 디바이스라 부릅니다. 아래 그림은 호스트와 디바이스의 관계를 간단히 보여주고 있습니다.
호스트와 디바이스의 관계
디바이스의 메모리 입출력, 연산은 모두 OpenCL API 를 통해 가능합니다. 호스트는 디바이스의 연산 결과와 같은 크기의 버퍼를 할당해 둔 후 병렬 연산이 끝나면 결과 버퍼를 할당해 둔 버퍼에 복사해 프로그램의 나머지 프로세스에서 사용합니다. 예를 들면 1000 800 화면에 이미지를 렌더하기 위해 호스트에서는 화면 픽셀 사이즈와 같은 크기의 버퍼를 미리 할당해 둔 뒤 필요한 연산을 디바이스에 맡기고 반환된 결과를 할당해 둔 버퍼에 복사해 화면에 출력합니다. 호스트는 대개 CPU 가 담당하고 디바이스는 멀티 쓰레딩을 지원하는 CPU 또는 GPU 가 담당합니다.
디바이스의 상위 개념으로 플랫폼이 있습니다. 플랫폼은 디바이스들이 연결되어 있는 모체, 예를 들어 맥에서 플랫폼은 Apple (clGetPlatformInfo 에 CL_PLATFORM_NAME 매개변수를 넣은 결과) 입니다. 디바이스는 플랫폼의 하위 개념이므로 OpenCL 에서는 플랫폼을 쿼리한 뒤 플랫폼의 정보를 이용해 디바이스를 쿼리할 수 있습니다.
CPU 또는 GPU 디바이스에서 실제 연산은 유닛에 의해 수행됩니다. 아래 그림은 병렬 연산을 해내는 유닛의 계층을 보여주고 있습니다.
연산 유닛의 계층
전체 연산 유닛들이 자리한 최상위 계층을 grid (디바이스), 중간 계층을 block (work-group), 최하위 계층을 thread (work-item) 라 할 수 있습니다. 같은 계층의 유닛들은 디바이스 상 메모리를 공유합니다. 이 개념은 변수의 스코프에 비유할 수 있습니다. grid 에 속한 모든 유닛들은 global 메모리를 공유합니다. 같은 work-group 에 속한 유닛들은 local 메모리를 공유하고 다른 work-group 의 유닛들은 이에 접근할 수 없습니다. 각각의 유닛은 private 메모리를 사용하고 이를 다른 유닛에서는 접근할 수 없습니다.
호스트와 디바이스의 메모리 계층
메모리는 호스트와 디바이스의 관계, 유닛의 계층 구조를 이해했다면 어렵지 않습니다. 호스트 메모리는 정적 또는 동적으로 할당하는 프로그램 상의 메모리를 뜻하며 디바이스의 연산 결과를 복사해 결과로 출력되거나 병렬 연산을 위한 매개 변수로서 역할합니다. 위 그림의 아랫부분의 context 는 OpenCL 에서 수행되는 커널, 메모리 등의 관리를 담당하는 구조체입니다. 앞서 잠시 언급했지만 OpenCL 은 이종간 병렬 프로그래밍이 가능합니다. 즉 OpenCL 을 이용해 여러 플랫폼의 여러 디바이스를 모두 사용할 수 있고 이들의 관리를 context 구조체가 담당합니다. context 의 관리 하에 디바이스의 메모리는 global > local > private 과 같은 계층을 갖습니다.
OpenCL 코드 작성
유닛의 계층 구조 이해를 위해 간단한 예제 코드를 작성해 보겠습니다. 아래 예시는 하나의 플랫폼으로부터 하나의 디바이스(GPU)만을 사용합니다. API 에 관해선 OpenCL 공식문서를 참고하셔도 좋습니다.
1. 플랫폼과 디바이스 정보
GPU 디바이스를 사용하기 위해 플랫폼 정보를 얻습니다.
cl_int clGetPlatformIDs(
cl_uint num_entries, // 필요한 플랫폼의 개수
cl_platform_id *platforms, // 플랫폼 ID 를 담을 포인터
cl_uint *num_platforms // 플랫폼의 개수를 담을 포인터
)플랫폼의 개수를 얻기 위해 함수 clGetPlatformIDs 에 개수를 담을 포인터 변수만을 전달하고 이후 실제 필요한 구조체를 담을 포인터 변수를 전달합니다.
cl_uint num_platforms;
cl_platform_id *platforms;
clGetPlatformIDs(0, NULL, &num_platforms);
platforms = (cl_platform_id *)malloc(sizeof(cl_platform_id) * num_platforms);
clGetPlatformIDs(num_platforms, platforms, NULL);우리는 하나의 플랫폼으로부터 하나의 디바이스만 필요하므로 아래와 같이 작성하겠습니다.
cl_platform_id platform;
clGetPlatformIDs(1, &platform, NULL);플랫폼을 이용해 디바이스 ID 를 얻습니다.
cl_int clGetDeviceIDs(
cl_platform_id platform, // clGetPlatformIDs 로부터 얻은 플랫폼 ID
cl_device_type device_type, // 디바이스 유형
cl_uint num_entries, // 필요한 디바이스의 개수
cl_device_id *devices, // 디바이스 ID를 담을 포인터
cl_uint *num_devices // 디바이스의 개수를 담을 포인터
)함수 clGetPlatformIDs 와 같은 방식으로 디바이스 ID 를 얻겠습니다. 이 때 디바이스 유형은 사용자가 반드시 지정해주어야 합니다. CPU 를 사용하려면 CL_DEVICE_TYPE_CPU, GPU 를 사용하려면 CL_DEVICE_TYPE_GPU 매크로 상수를 매개변수로 전달하면 됩니다. 디바이스 유형의 모든 매크로 상수는 OpenCL 공식문서에서 확인할 수 있습니다. 우리는 GPU 한 개를 사용할 것이므로 다음과 같이 작성합니다.
cl_device_id device;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);2. context 와 커맨드 큐 생성
필요한 디바이스를 얻었으므로 OpenCL 의 커널, 메모리를 관리할 context 구조체를 만들겠습니다.
cl_context clCreateContext(
cl_context_properties *properties, // context 의 속성 리스트
cl_uint num_devices, // 디바이스의 개수
const cl_device_id *devices, // 디바이스 ID 를 담은 포인터
void *pfn_notify( // context 에서 에러가 발생했을 때 실행할 콜백 함수
const char *errinfo,
const void *private_info,
size_t cb,
void *user_data
),
void *user_data, // 콜백 함수의 매개변수 포인터
cl_int *errcode_ret // 에러 여부를 알려줄 결과값 포인터
)아래는 필요한 매개변수만 전달해 clCreateContext 를 호출한 코드입니다.
cl_context context;
cl_int ret;
context = clCreateContext(NULL, 1, &device, NULL, NULL, &ret);결과에 에러가 있는지 판단하기 위해 아래와 같이 조건문을 만들어도 좋습니다.
if (ret != CL_SUCCESS)
{
printf("Failed to create context\n");
...
}디바이스에서 연산 유닛에 의해 동시에 수행되는 병렬 연산을 하나의 함수라고 한다면 OpenCL 에서는 이 함수를 커널(kernel) 이라 정의하고 있습니다. 수행될 커널들은 우선 커맨드 큐(command queue) 에 넣습니다. 이 때 커맨드 큐는 디바이스마다 만들 수 있으며 아래 함수를 이용합니다.
cl_command_queue clCreateCommandQueue(
cl_context context, // context
cl_device_id device, // 디바이스 ID
cl_command_queue_properties properties, // 커맨드 큐의 속성을 담은 구조체
cl_int *errcode_ret // 에러 여부를 알려줄 결과값 포인터
)아래와 같이 함수를 호출해 커맨드 큐 구조체를 생성하겠습니다.
cl_command_queue command_queue;
cl_int ret;
command_queue = clCreateCommandQueue(context, device, 0, &ret);3. 커널 소스 컴파일
디바이스 ID 를 얻고 사용할 커맨드 큐를 만들어 병렬 연산을 위한 첫번째 준비를 마쳤습니다. 이번에는 커널 소스를 .cl 파일에 작성하고 이를 컴파일 하는 과정을 알아보겠습니다. 아래는 각각의 연산 유닛의 grid 상 인덱스를 출력 버퍼에 쓰는 커널입니다.
// example.cl
__kernel void get_unit_idx(__global int *out)
{
int idx = get_global_id(0);
out[idx] = idx;
}__kernel 은 작성한 함수가 OpenCL 의 커널임을 가리킵니다. __global 은 매개변수 int *out 가 다른 유닛들과 공유되는, 즉 grid 상 메모리임을 가리킵니다. 함수 get_global_id(0) 은 현재 유닛의 grid 상 인덱스를 반환합니다. 이 때 매개변수 0은 차원의 인덱스입니다. 차원은 커맨드 큐에 컴파일된 커널을 넣는 과정에서 알아볼 예정입니다.
커널 소스 작성은 .c 파일에 문자열로 작성할 수도 있습니다. 만약 위와 같이 .cl 파일에 따로 작성하셨다면 파일 포인터 또는 파일 디스크립터를 이용해 소스파일의 컨텐트를 하나의 문자열 버퍼로 저장해야 합니다. 이 과정은 어렵지 않게 할 수 있으니 이 글에서는 생략하겠습니다.
작성한 커널 소스는 프로그램 구조체를 만드는 함수의 매개변수로 사용됩니다.
cl_program clCreateProgramWithSource(
cl_context context, // context
cl_uint count, // 커널 소스를 담은 문자열 버퍼의 개수
const char **strings, // 커널 소스를 담은 문자열 버퍼의 포인터
const size_t *lengths, // 문자열 버퍼의 길이
cl_int *errcode_ret // 에러 여부를 알려줄 결과값 포인터
)커널 소스코드를 문자열 변수에 저장한 뒤 아래와 같이 함수를 호출해 프로그램 구조체를 만듭니다.
cl_program program;
const char *source_str = /* 커널 소스 코드 문자열 */;
cl_int ret;
program = clCreateProgramWithSource(context, 1, (const char **)&source_str, NULL, &ret);매개변수로서 문자열 버퍼 길이의 포인터를 NULL 로 전달하면 source_str 이 반드시 null 문자로 끝나야 합니다.
만들어진 프로그램 구조체는 함수 clBuildProgram 를 이용해 커널을 사용할 수 있도록 빌드합니다.
cl_int clBuildProgram(
cl_program program, // 프로그램 구조체
cl_uint num_devices, // 프로그램을 빌드할 디바이스 개수
const cl_device_id *device_list, // 프로그램을 빌드할 디바이스 리스트
const char *options, // 빌드 옵션
void (*pfn_notify)(cl_program, void *user_data), // 빌드 후 호출될 콜백 함수
void *user_data // 콜백 함수의 매개변수
)빌드 옵션과 콜백 함수를 제외하고 필요한 매개변수만을 넣어 함수를 호출하겠습니다.
clBuildProgram(program, 1, &device, NULL, NULL, NULL);만약 커널 소스에 컴파일 에러 요소가 있다면 clBuildProgram 은 CL_SUCCESS 가 아닌 값을 반환합니다.
빌드에 성공한 프로그램에서 원하는 커널을 실행시켜줄 커널 구조체를 만들어 보겠습니다.
cl_kernel clCreateKernel(
cl_program program, // 프로그램 구조체
const char *kernel_name, // 커널 이름
cl_int *errcode_ret // 에러 여부를 알려줄 결과값 포인터
)커널 이름은 작성한 커널 소스에 있는 함수의 이름을 말합니다. 예를 들어 우리가 작성한 소스의 커널 이름은 get_unit_idx 입니다. 하나의 소스에 여러 커널을 작성하고 각각의 커널에 대해 커널 구조체를 만들어 사용할 수 있습니다.
cl_kernel kernel;
cl_int ret;
kernel = clCreateKernel(program, "get_unit_idx", &ret);커널 구조체를 프로그램 빌드 후 곧바로 만들 필요는 없습니다. 프로세스 상 적당한 때에 커널 객체를 만들고 커맨드 큐에 넣어도 무방합니다.
4. 메모리 오브젝트 생성
host 의 메모리는 정적 또는 동적으로 어떻게 할당해도 좋습니다. 병렬 연산의 매개변수 또는 결과로 출력될 디바이스의 메모리는 아래 함수를 이용해 오브젝트를 생성합니다.
cl_mem clCreateBuffer(
cl_context context, // context
cl_mem_flags flags, // 메모리 사용 용도를 정의하는 플래그
size_t size, // 메모리 사이즈
void *host_ptr, // host 메모리의 포인터
cl_int *errcode_ret // 에러 여부를 알려줄 결과값 포인터
)매개변수 flags 를 통해 메모리의 사용 용도를 정합니다. 예를 들어 디바이스에서 메모리의 데이터를 읽기만 한다면 CL_MEM_READ_ONLY, 쓰기만 한다면 CL_MEM_WRITE_ONLY, 그 외에 host 의 메모리와 연관되어 있다면 해당 매크로 상수를 전달할 수 있습니다. 공식문서에 사용 가능한 매크로 상수와 그 용도가 자세히 기술되어 있습니다.
우리는 각각의 연산 유닛의 grid 상 인덱스 데이터가 쓰여질 CL_MEM_WRITE_ONLY 메모리가 필요합니다.
cl_mem dev_mem_obj;
size_t size = /* 메모리 오브젝트의 크기(윈도우에 출력될 이미지 버퍼의 크기 등) */
dev_mem_obj = clCreateBuffer(context, CL_MEM_WRITE_ONLY, size, NULL, &ret);플랫폼과 디바이스의 ID를 얻고 메모리 오브젝트를 생성하기까지는 병렬 연산을 위한 준비과정이라 할 수 있습니다. 다음 포스트에서는 준비된 커널을 커맨드 큐에 넣어 실행하고 반환된 메모리 오브젝트를 host 의 메모리 버퍼에 복사해 결과값을 출력해보겠습니다.
