[OpenCL] OpenCL 프로그래밍 - 1 OpenCL 호스트와 디바이스

1.1 OpenCL 프로그램 구성

• OpenCL로 프로그램을 개발할 때는 디바이스에서 동작하는 커널 프로그램과 호스트 프로그램을 따로 작성
• 커널 프로그램[코드 3-1] 은 OpenCL C언어로 작성
• 호스트 프로그램[코드 3-2] 은 C언어로 작성하되 OpenCL 런타임 API를 사용

1.2 커널 프로그램

__kernel void hello(__global char* string)

• OpenCL C언어 문법에 맞춰 작성한다.
• hello() 함수의 선언에 함수 수식자 __kernel을 지정 [C언어와 차이점 1]
• 함수 수식자 __kernel 은 hello() 함수가 디바이스에서 동작하며 호스트에서 호출할 수 있는 커널 함수임을 의미.
• __kernel 수식자를 사용하는 커널 함수는 다음 2가지 OpenCL 런타임 API 함수를 이용해 호스트에서 호출 가능
  - 태스크 호출 API 함수 : clEnqueueTask()
  - 데이터 병렬 호출 API 함수 : clEnqueueNDRangeKernel()
  - hello() 함수는 데이터 병렬을 위한 함수는 아님. 호스트 프로그램에서 clEnqueueTask 함수를 이용해 커널을 호출함.
• hello() 함수의 첫 번째 파라미터 string에 주소 공간 수식자 __global을 지정 [C언어와 차이점 2]
  - OpenCL 메모리 모델 중 글로벌 메모리에 존재한다는 것을 의미함.
  - 일반적으로 디바이스 프로그램은 OpenCL 메모리 모델로 정의된 글로벌 메모리, 컨스턴트 메모리, 로컬 메모리, 프라이빗 메모리 중 어느 한 가지에 속한다.
  - 메모리 영역에 읽기와 쓰기가 가능하며, 각각 __global, __constant, __local, __private 이라는 수식자를 정의
  - 주소 공간 수식자를 생략하면 __private으로 지정된다.

01 커널내에서 스레드가 몇번째 위치인지 확인하는 방법

• get_global_size(0) : 가로 크기를 반환한다.
• get_global_id(1) : 현재 스레드의 위치가 위에서부터 몇번째에 위치해 있는지를 반환한다.
• get_global_id(0) : 현재 스레드의 위치가 왼쪽에서부터 몇번째에 위치해 있는지를 반환한다.

1.3 호스트 프로그램

// OpenCL 커널 실행

ret = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);

• OpenCL C언어의 문법에 따라 개발한 커널 프로그램은 OpenCL 런타임 API를 이용해 개발한 호스트 프로그램에 의해 디바이스에서 실행된다.
• 커널 함수 hello()를 호스트에서 호출할 때는 단지 OpenCL 런타임 API 함수의 하나인 clEnqueueTask()를 호출할 뿐이다. 하지만 실제로는 이를 위해 여러 단계의 절차를 거친다.
• 전체 워크 플로우 

1. 플랫폼 지정
2. 디바이스 선택
3. 컨텍스트 생성 (선택된 Device를 이용해 Context 생성)
4. 커맨드 큐 생성
5. 메모리 오브젝트 생성
6. 커널 프로그램 파일 읽기
7. 프로그램 오브젝트 생성
8. 커널 컴파일
9. 커널 오브젝트 생성
10. 커널 파라미터 설정
11. 커널 실행
12. 메모리 오브젝트 로드 (Device로부터 연산이 완료된 Buffer Data 읽기)
13. 오브젝트 해제

01 플랫폼 지정

cl_platform_id platform_id = NULL;

...

cl_uint ret_num_platforms;

...

// 플랫폼, 디바이스 정보를 얻음

ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);

• OpenCL의 플랫폼 모델은 하나의 호스트와 호스트에 연결된 하나 이상의 디바이스로 구성된다.
• 따라서 호스트에서 OpenCL을 이용할 때는 가정 먼저 프로그램이 동작할 OpenCL 플랫폼을 지정해야 한다.
• clGetPlatformIDs()는 OpenCL 런타임 API에 정의된 함수이며, 시스템에 존재하는 플랫폼 ID를 얻는다.
  - 시스템에 OpenCL이 동작하는 플랫폼(하드웨어)이 존재하는지 찾아본 후,
  - 검색된 플랫폼이 있을 때 cl_platform_id 형의 파라미터인 platform_id에 플랫폼 ID를 반환한다.
  - 이후 호스트 프로그램은 platform_id를 지정해 해당 플랫폼을 이용할 수 있다.
• 첫 번째 파라미터는 호스트에 존재하는 애플리케이션이 원하는 플랫폼 수이며 보통 1을 지정한다.
• 세 번째 파라미터 num_pltform에는 시스템에서 실제로 이용 가능한 OpenCL 플랫폼 수를 반환한다.

[참고] 01-1 terasic hello example 에서는 아래와 같이 플랫폼 지정

static cl_platform_id platform = NULL;

...

// Get the OpenCL platform.

platform = findPlatform("Altera");

if (platform == NULL) {

    printf("ERROR: Unable to find Altera OpenCL platform.\n");

    return false;

}

[참고] 01-2 후로린 블로그 예제

// get all platforms

clGetPlatformIDs(0, NULL, &platformCount);

platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * platformCount);

clGetPlatformIDs(platformCount, platforms, NULL);

• platform이란 opencl 서비스를 제공할 수 있는 환경을 말한다.
• 현재 컴퓨터에 얼마나 많은 플랫폼이 있는지를 먼저 알아와야 한다.
  - clGetPlatformIDs()를 이용해 플랫폼의 개수를 platformCount에 먼저 받아온다.
• 이후 플랫폼의 개수만큼 메모리를 할당하고 플랫폼을 얻어오게 된다.
  - platform 변수 공간을 할당하고, clGetPlatformIDs()를 통해 플랫폼을 받아온다.

02 디바이스 선택

cl_device_id device_id = NULL;

...

cl_uint ret_num_devices;

...

// 플랫폼, 디바이스 정보를 얻음

ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);

• clGetDeviceIDs() 역시 OpenCL 런타임 API에 정의된 함수이며, 플랫폼에서 이용 가능한 디바이스를 얻는다.
• 첫번째 파라미터에는 디바이스를 선택하려는 플랫폼을 지정.
• 두 번째 파라미터에는 디바이스의 종류를 지정. 해당 예제는 플랫폼이 제공하는 표준 디바이스를 나타내는 CL_DEVICE_TYPE_DEFAULT 를 지정
• 세 번째 파라미터에는 얻어야 할 디바이스의 수를 지정.
• 네 번째 파라미터에는 device_id에 저장된 디바이스를 호스트 프로그램에서 이용하기 위한 cl_device_id 형의 디바이스 핸들로 반환한다. 
• 두 번째 파라미터에서 지정한 디바이스가 플랫폼에 존재하지 않을 때는 다섯 번째 파라미터 num_devices에 0을 반환한다.

[참고] 02-1 terasic hello example 에서는 아래와 같이 디바이스 선택

static cl_device_id device = NULL;

...

// Query the available OpenCL devices.

scoped_array<cl_device_id> devices;

cl_uint num_devices;

devices.reset(getDevices(platform,CL_DEVICE_TYPE_ALL, &num_devices));

// We'll just use the first device.

device = devices[0];

[참고] 02-2 후로린 블로그 예제

// get all devices

clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, NULL, &deviceCount);

devices = (cl_device_id*)malloc(sizeof(cl_device_id) * deviceCount);

clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, deviceCount, device, NULL);

• 디바이스는 계산을 수행할 수 있는 유닛의 집합이다.
• GPU로 설명하자면 GPU내의 수많은 코어들의 집합이라고 비유할 수 있다.
• 플랫폼 내의 디바이스를 찾아 실제로 연산을 진행하게 되며 platform을 받아오는 과정과 동일하다.

03 컨텍스트 생성

cl_context context = NULL;

...

// OpenCL 컨텍스트 생성

context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);

• 이용하려는 디바이스 핸들을 얻었다면 다음에는 OpenCL 실행 환경이 될 OpenCL 컨텍스트를 생성해야 한다. clCreateContext()는 컨텍스트를 생성하기 위한 함수이다.
• 보통 컨텍스트를 생성할 때는 커널을 실행할 디바이스를 반드시 하나 이상 지정해야 한다. 
  - 두 번째 파라미터에는 이용할 디바이스의 수를
  - 세 번째 파라미터에는 이용할 디바이스에 해당되는 디바이스 핸들의 리스트를 지정.
• 이후에 생성되는 각종 OpenCL 오브젝트는 모두 여기서 지정한 OpenCL의 컨텍스트에 속하게 된다. 같은 컨텍스트 안의 오브젝트를 통해서 각 디바이스를 제어한다.
• 즉, 컨텍스트란 하나 이상의 연산 디바이스를 이용할 수 있는 가상의 컴퓨터라고 생각하면 쉽다.

[참고] 03-1 terasic hello example 에서는 아래와 같이 컨텍스트 생성

static cl_context context = NULL;

...

// Create the context.

context = clCreateContext(NULL, 1, &device, NULL, NULL, &status);

checkError(status, "Failed to create context");

[참고] 03-2 후로린 블로그 예제

// create context

context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);

• CL 커널이 실행되는 환경으로, 동기화와 메모리 관리가 정의된다.
• OpenCL디바이스에서 실행할 OpenCL함수들을 포함하고 있다.
• clCreateContext()함수를 이용해 생성한다.

04 커맨드 큐 생성

cl_command_queue command_queue = NULL;

...

// 커맨드 큐 생성

command_queue = clCreateCommandQueue(context, device_id, 0, &ret);

• 디바이스를 제어하기 위한 준비 과정
• OpenCL에서는 호스트가 디바이스에게 무언가를 지시(커널 실행 명령, 혹은 호스트~디바이스 사이의 메모리 전송 명령)하려면 커맨드 큐 오브젝트 command_queue를 통한다.
• 각각의 디바이스에는 반드시 하나 이상의 커맨드 큐 오브젝트를 생성해야 한다.
• 함수는 cl_command_queue 형의 커맨드 큐 핸들을 반환하며 이후 커널을 실행하거나 호스트~디바이스 사이에서 메모리를 전송하려면 이 커맨드 큐 핸들을 지정한다.
• clCreateCommandQueue()는 커맨드 큐를 생성하기 위한 함수이다.
  - context : 커맨드 큐가 소속될 컨텍스트
  - device_id : 커맨드 큐의 명령을 실행할 디바이스를 지정

[참고] 04-1 terasic hello example 에서는 아래와 같이 커맨드 큐 생성

static cl_command_queue queue = NULL;

...

// Create the command queue.

queue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, &status);

checkError(status, "Failed to create command queue");

[참고] 04-2 후로린 블로그 예제

// create command queue

queue = clCreateCommandQueue(context, device, 0, NULL);

• 호스트에서 디바이스 별로 생성되는 것으로 하나의 디바이스에서 여러개의 커맨드 큐가 연결 가능하다.
• 커맨드 큐를 이용해 커널을 실행하고 메모리의 매핑, 언매핑, 동기화 등을 할 수 있다.

05 메모리 오브젝트 생성

cl_mem memobj = NULL;

...

// 메모리 버퍼 생성

memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(char), NULL, &ret);

• 디바이스에서 커널이 동작할 때 커널의 처리 대상인 배열 등의 데이터는 모두 디바이스 메모리에 미리 준비되어야 하며, 디바이스 메모리에 처리할 데이터를 준비하는 것은 호스트의 역할이다.
• 커널 자체는 디바이스의 외부에 있는 저장장치나 네트워크로부터 디바이스 메모리에 데이터를 가져올 수 없다. 따라서 호스트가 디바이스 메모리에 데이터를 준비할 때는 메모리 오브젝트를 생성해야 디바이스 메모리에 접근할 수 있다.
• clCreateBuffer()는 디바이스 상에 메모리를 확보하기 위한 함수이다. 여기서 확보된 메모리는 메모리 오브젝트 memobj를 통해 호스트에서 참조할 수 있다.
  - context : 메모리 오브젝트가 소속될 컨텍스트를 지정
  - CL_MEM_READ_WRITE : 메모리 속성, 커널이 디바이스 상에 확보한 메모리 안에 데이터 읽고 쓰기 가능
  - MEM_SIZE * sizeof(char) : 확보할 메모리 크기를 지정

[참고] terasic example 중 hello는 메모리 사용하지 않고, vector add는 아래와 같이 메모리 오브젝터 생성

scoped_array<cl_mem> input_a_buf; // num_devices elements

scoped_array<cl_mem> input_b_buf; // num_devices elements

scoped_array<cl_mem> output_buf; // num_devices elements

...

// Input buffers.

input_a_buf[i] = clCreateBuffer(context, CL_MEM_READ_ONLY, n_per_device[i] * sizeof(float), NULL, &status);

checkError(status, "Failed to create buffer for input A");

input_b_buf[i] = clCreateBuffer(context, CL_MEM_READ_ONLY, n_per_device[i] * sizeof(float), NULL, &status);

checkError(status, "Failed to create buffer for input B");

// Output buffer.

output_buf[i] = clCreateBuffer(context, CL_MEM_WRITE_ONLY, n_per_device[i] * sizeof(float), NULL, &status);

checkError(status, "Failed to create buffer for output");

06 커널 프로그램 파일 읽기

FILE *fp;

char fileName[] = "./hello.cl";

char *source_str;

size_t source_size;

// 커널을 포함한 소스 코드를 로드

fp = fopen(fileName, "r");

if (!fp)

{

        fprintf(stderr, "Failed to load kernel.\n");

        exit(1);

}

source_str = (char*)malloc(MAX_SOURCE_SIZE);

source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);

fclose(fp);

• 디바이스에서 커널을 실행하려면 먼저 커널 프로그램을 읽어들일 필요가 있으며 이 작업은 호스트 프로그램에서 이루어진다.
• 여기서 커널 프로그램은 OpenCL 컴파일러로 미리 컴파일된 바이너리 데이터 혹은 컴파일하기 전의 소스 코드 일 수도 있다.
• 커널 파일을 읽을 때는 OpenCL 런타임 API를 이용하지 않는다. 
• fopen과 fread를 통해 OpenCL C언어의 소스 코드를 커널 프로그램으로 읽어들인다.
• 커널 프로그램은 통상적으로 호스트의 저장장치에 파일로 존재하므로 다음처럼 hello.cl 을 호스트의 버퍼인 source에 읽어온다.

[참고] 06-1 후로린 블로그 예제

// 프로그램으로부터 커널 생성

simpleKernel = clCreateKernel(program, "simpleKernel", NULL);

• kernel의 집합인 program 내에서 특정 커널을 선택해 커널을 실행할 수 있도록 커널을 만들고 커널의 포인터를 반환하는 코드이다. 

07 프로그램 오브젝트 생성

cl_program program = NULL;

...

// 미리 로드한 소스 코드로 커널 프로그램을 생성

program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);

• 위와 같은 커널 프로그램의 소스 코드를 그대로 커널로 인식시킬 수 없다.
• 커널 프로그램에는 복수의 커널 함수가 포함될 수 있기 때문이다. 
• OpenCL에서는 커널 프로그램을 먼저 프로그램 오브젝트로 인식시켜야 한다.
• clCreateProgramWithSource()는 소스 코드로부터 프로그램 오브젝트를 생성하기 위한 함수다.
  - (const char **)&source_str : 커널 프로그램의 소스 코드
  - (const size_t *)&source_size : 소스 코드의 문자열 크기(바이트 단위)
• 해당 함수는 이름에서 알 수 있듯이 커널 프로그램의 소스 코드 문자열로부터 프로그램 오브젝트인 program를 생성한다.
• 미리 컴파일된 커널 프로그램의 바이너리 파일을 가지고 프로그램 오브젝트를 생성할 때는 clCreateProgramWithSource()함수 대신 clCreateProgramWithBinary()함수를 이용한다.  

[참고] 07-1 terasic hello example 에서는 아래와 같이 커널 프로그램 파일 읽고 프로그램 오브젝트 생성

cl_program program = NULL;

...

// Create the program.

std::string binary_file = getBoardBinaryFile("hello_world", device);

printf("Using AOCX: %s\n", binary_file.c_str());

program = createProgramFromBinary(context, binary_file.c_str(), device, num_devices);

[참고] 07-2 후로린 블로그 예제

program = clCreateProgramWithSource(context, 1, (const char **)&source, NULL, NULL);

cl_int build_status;

build_status = clBuildProgram(program, 1, &device, NULL, NULL, NULL);

• program이란 kernel의 집합이라고 보면 된다.
• 간단히 말해 프로그램 소스 혹은 빌드된 바이너리이다.

08 커널 컴파일

// 커널 프로그램 빌드

ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);

• 소스 코드로 등록한 프로그램 오브젝트를 OpenCL C 컴파일러와 링커를 사용해 빌드하는 과정이다.
• clBuildProgram() 함수는 OpenCL 컴파일러와 링커에 의해 빌드를 실행한다.
  - program : 빌드 대상이 되는 프로그램 오브젝트
  - 두 번째 파라미터 1은 세 번째 파라미터에 지정한 디바이스 수
  - &device_id : 프로그램 오브젝트에 연관된 디바이스의 리스트를 가리키는 포인터
  - 네 번째 파라미터에는 컴파일러에게 넘겨줄 파라미터 문자열을 지정한다.
• 단, 이 컴파일 절차는 clCreateProgramWithSource() 함수에 의해 소스 코드로부터 프로그램 오브젝트를 생성했을 때만 필요하다. clCreateProgramWithBinary() 함수를 이용해 컴파일 된 바이너리로부터 직접 프로그램 오브젝트를 생성할 때는 이 컴파일 절차가 필요없다.

[참고] terasic hello example 에서는 아래와 같이 커널 컴파일

// Build the program that was just created.

status = clBuildProgram(program, 0, NULL, "", NULL, NULL);

checkError(status, "Failed to build program");

09 커널 오브젝트 생성

cl_kernel kernel = NULL;

...

// OpenCL 커널 생성

kernel = clCreateKernel(program, "hello", &ret);

• 컴파일된 프로그램 오브젝트를 생성했다면 커널 오브젝트를 생성해야 한다.
• 하나의 커널 오브젝트는 하나의 커널 함수를 나타낸다. 따라서 커널 오브젝트를 생성할 때는 다음처럼 커널 함수 이름인 hello를 지정한다.
• clCreateKernel() 함수는 커널 오브젝트를 생성하기 위해 사용한다.
  - 첫 번째 파라미터 program : 커널을 포함한 프로그램 오브젝트
  - 두 번째 파라미터 "hello" : 커널 오브젝트로 지정한 커널 함수 이름
• 이번 예에서는 하나의 프로그램 오브젝트에 하나의 커널 함수만을 생성했지만 실제로는 복수의 커널 함수를 하나의 프로그램 오브젝트로서 생성할 수도 있다. 
• 이는 하나의 프로그램 오브젝트로부터 여러 개의 커널 오브젝트를 생성할 수 있다는 의미다. 단, 커널 오브젝트와 커널 함수는 1대1로 대응하기 때문에 clCreateKernel() 함수를 여러 번 호출해야 한다.

[참고] terasic hello example 에서는 아래와 같이 커널 오브젝트 생성

// Create the kernel - name passed in here must match kernel name in the

// original CL file, that was compiled into an AOCX file using the AOC tool

const char *kernel_name = "hello_world";  // Kernel name, as defined in the CL file

kernel = clCreateKernel(program, kernel_name, &status);

checkError(status, "Failed to create kernel");

10 커널 파라미터 설정

// OpenCL 커널 파라미터 설정

ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);

• 커널 함수가 얻는 파라미터에 데이터를 전달하는 과정이다.
• 이번 예제에서는 문자열 형식의 커널 파라미터인 string에 호스트에서 생성한 메모리 오브젝트 memobj를 전달한다. 
• 이처럼 호스트에서 확보한 메모리 오브젝트를 커널 파라미터로 전달할 때 커널 파라미터는 주소 공간 수식자인 __global을 지정해 선언한다.
• clSetKernelArg() 함수는 커널에 파라미터를 전달하기 위한 함수다.
  - 첫 번째 파라미터 kernel : 커널 파라미터를 전달할 커널 오브젝트
  - 두 번째 파라미터 0 : 넘겨줄 커널 파라미터의 인덱스로서 범위는 0 ~ 커널 파라미터수-1,
  - 이번 예제는 첫 번째 문자열 형식의 커널 파라미터인 string에 데이터를 전달하기 때문에 0을 지정함. 따라서 데이터를 전달할 커널 파라미터가 여럿 존재한다면 커널 파라미터 각각을 clSetKernelArg() 함수로 호출한다.
  - 세 번째 파라미터 sizeof(cl_mem) : 넘겨줄 파라미터의 크기를 지정
  - 네 번째 파라미터 (void *)&memobj : 전달할 데이터로 메모리 오브젝트 memobj를 지정
• 해당 예제에서는 메모리 오브젝트를 파라미터로 지정했다. 만약 호스트의 데이터를 커널 파라미터로 직접 전달하려면 아래와 같이 코드를 작성한다.

int a = 10;

clSetKernelArg(kernel, 0, sizeof(int), (void *)&a);

[참고] terasic hello example 에서는 아래와 같이 커널 파라미터 설정

// Set the kernel argument (argument 0)

status = clSetKernelArg(kernel, 0, sizeof(cl_int), (void*)&thread_id_to_output);

checkError(status, "Failed to set kernel arg 0");

[참고] terasic vector example 에서는 아래와 같이 커널 파라미터 설정

// Set kernel arguments.

unsigned argi = 0;

status = clSetKernelArg(kernel[i], argi++, sizeof(cl_mem), &input_a_buf[i]);

checkError(status, "Failed to set argument %d", argi - 1);

status = clSetKernelArg(kernel[i], argi++, sizeof(cl_mem), &input_b_buf[i]);

checkError(status, "Failed to set argument %d", argi - 1);

status = clSetKernelArg(kernel[i], argi++, sizeof(cl_mem), &output_buf[i]);

checkError(status, "Failed to set argument %d", argi - 1);

11 커널 실행

// OpenCL 커널 실행

ret = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);

• 위의 코드는 호스트에서 디바이스로 커널을 실행하는 과정, 디바이스에게 태스크를 맡기기 위해 커맨드 큐에 태스크를 전달하는 구문이다.
• hello라는 커널을 실행하는 명령은 clEnqueueTask() 함수에 의해 큐에 전달된 후 디바이스 상의 실행 가능한 연산 유닛에서 실행한다.
• clEnqueueTask() 함수는 커맨드 큐에 명령을 전달할 뿐이므로 호스트에서 호출한 clEnqueueTask() 함수의 실행이 끝난 시점에서는 hello라는 커널의 실행이 (디바이스에서) 종료되었다고 보장할 수 없다.
• clEnqueueTask() 함수가 정상 종료되었을 때는 hello의 실행을 의미하는 명령이 커맨드 큐에 제대로 전달되었다는 점만 보장할 뿐이다. 즉, 실제로 디바이스 상에서 hello의 실행이 종료될 때까지 기다리려면 다섯 번째 파라미터를 통해서 이벤트 오브젝트를 얻어야한다.
• clEnqueueTask() 함수는 태스크 병렬을 실행할 때 사용한다. 
• 데이터 병렬을 실행할 때는 인덱스 공산 설정을 필요로 하는 clEnqueueNDRangeKernel() 함수를 사용한다. -> 아래 예제 참고

[참고] terasic vector example 에서는 아래와 같이 커널 실행

// Launch the kernel

status = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, gSize, wgSize, 0, NULL, NULL);

checkError(status, "Failed to launch kernel");

12 메모리 오브젝트 로드

char string[MEM_SIZE];

...

// 실행 결과를 메모리 버퍼에서 얻음

ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, MEM_SIZE * sizeof(char), string, 0, NULL, NULL);

• 호스트에서 커널의 처리 결과를 확인하려면 디바이스의 메모리 영역에서 호스트의 메모리 영역으로 데이터를 전송해야 한다.
• clEnqueueReadBuffer()는 디바이스의 메모리에서 호스트의 메모리로 데이터를 전송하기 위한 함수다. 반대로 호스트 메모리에서 디바이스 메모리로 데이터를 전송할 때는 clEnqueueWriteBuffer() 함수를 이용한다.
• 함수 이름에 'Enqueue'가 포함된 것은 디바이스를 조작하는 함수인 것 !!
• clEnqueueReadBuffer() 함수는 커맨드 큐에 메모리 명령이 전달된 후 데이터 복사를 시작한다.
  - 두 번째 파라미터 memobj : 데이터를 읽어들일 디바이스 상의 메모리를 나타내는 메모리 오브젝트를 지정
  - 세 번째 파라미터 CL_TRUE : 동기 메모리 복사를 하겠다는 뜻, string에 복사가 끝날 때까지 clEnqueueReadBuffer() 함수는 종료되지 않는다. 
  -> 만약 CL_FALSE를 지정하면 비동기 메모리 복사를 하겠다는 뜻, 커맨드 큐에 복사 명령이 제대로 전달되면 실제 데이터 복사가 끝나지 않더라도 clEnqueueReadBuffer() 함수가 종료된다.
  - 다섯 번째 파라미터 MEM_SIZE * sizeof(char) : 복사할 데이터의 크기(바이트)를 지정
  - 여섯 번째 파라미터 string : 데이터를 저장할 호스트 상의 메모리 영역에 대한 포인터를 지정
• 호스트에서 clEnqueueReadBuffer() 함수에게 디바이스의 메모리를 읽으라는 지시를 내리기 전, hello 커널이 반드시 종료되었다고 할 수 있을까? 에 대한 질문
  - 호스트의 커널 실행 함수 clEnqueueTask()는 큐에 전달된 지점에서 빠져나온다.
  - 커널의 실행이 끝나기 전에 디바이스의 메모리를 읽는다면 문제가 생길 수 있지만 Enqueue 계열의 함수는 호스트에서 보면 비동기 호출이지만 커맨드 큐에 투입된 명령의 실행 순서는 커맨드 큐 생성시 특별한 지정을 하지 않는 이상 동기 실행(큐에 전달된 순서대로 명령을 실행)이므로 우려할 필요 없다.
  - 이번 예는 커맨드 큐에 비동기 실행 지시를 내리지 않았기 때문에 먼저 큐에 전달된 커널 실행 명령이 종료된 것을 확인한 후 메모리 복사 명령이 실행된다.

[참고] terasic hello example 에서는 읽을 것이 없음

[참고] terasic vector example 에서는 아래와 같이 메모리 오브젝트 로드

// Read the result. This the final operation.

status = clEnqueueReadBuffer(queue[i], output_buf[i], CL_FALSE, 0, n_per_device[i] * sizeof(float), output[i], 1, &kernel_event[i], &finish_event[i]);

13 오브젝트 해제

ret = clReleaseKernel(kernel);

ret = clReleaseProgram(program);

ret = clReleaseMemObject(memobj);

ret = clReleaseCommandQueue(command_queue);

ret = clReleaseContext(context);

• 프로그램의 안전한 종료를 위해 생성한 각종 오브젝트를 해제해야 한다. 
• 사용이 끝난 오브젝트는 재사용할 수 있기 때문에 생성과 해제를 매번 실시할 필요는 없다.
• 실제 프로그램에서는 커널 파라미터의 설정 혹은 입력 버퍼를 통해 메모리 복사 -> 커널 실행 -> 메모리 읽기를 반복하는 것이 기본이다.
• 오브젝트를 생성만 하고 해제시키지 않으면서 계속 사용하면 호스트의 각종 오브젝트 관리용 메모리 영역이 부족해진다. 이 때는 OpenCL 런타임이 메모리 부족 에러를 반환한다.

[참고] 13-1 terasic hello example 에서는 아래와 같이 오브젝트 해제

// Free the resources allocated during initialization

void cleanup() {

        if (kernel) {

               clReleaseKernel(kernel);

        }

        if (program) {

               clReleaseProgram(program);

        }

        if (queue) {

               clReleaseCommandQueue(queue);

        }

        if (context) {

               clReleaseContext(context);

        }

}