개발 환경 설정은 아래 링크를 참고하자.
https://velog.io/@pikamon/OpenCL-1

---

1. OpenCL이란?

Open Computing Language의 줄임말로, 이기종 플랫폼에 존재하는 여러 가지 프로세서들(예. CPU, GPU, DSP, FPGA 등)을 두루 활용하여 병렬 연산을 처리할 수 있도록 만들어진 병렬 처리용 프레임워크이다.

하나의 OpenCL 표준을 정해놓고, 프로세서 제조사들은 사용자들이 자신들의 프로세서에 접근하여 병렬 연산을 실행할 수 있도록 프로세서를 만들고, 사용자들은 어떤 프로세서든 동일한 OpenCL 인터페이스를 이용해 해당 프로세서를 이용해 연산을 수행할 수 있도록 만든 것이다.

어떤 문제의 계산량을 분할하여 여러 연산 장치에 분배함으로써 커다란 문제를 빠르게 해결할 수 있다는 장점이 있다. (Orchestration)

1. OpenCL 용어 정리

OpenCL에 나오는 용어들이 처음엔 낯설 수 있는데, 용어와 기본 개념을 정확히 알아야 프로그램을 설계할 수 있다.

1. 호스트

OpenCL 인터페이스를 수행하는 주체를 호스트라고 한다.

2. 플랫폼

프로세서 제조사를 플랫폼이라고 한다. 컴퓨터에 Intel CPU, Intel GPU, AMD GPU가 장착되어 있다면, 플랫폼은 Intel과 AMD를 말한다.
호스트에서 하나 또는 여러 개의 플랫폼을 선택하여 OpenCL을 실행할 수 있다.

clGetPlatformIDs 함수를 이용하여 현재 컴퓨터로부터 플랫폼을 가져올 수 있다.

// 2. 플랫폼 가져오기
clGetPlatformIDs(0, NULL, &platformCount); // 플랫폼 개수를 가져온다.
platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * platformCount); // 플랫폼 개수만큼 동적할당한다.
clGetPlatformIDs(platformCount, platforms, NULL); // 플랫폼을 가져온다.

3. 디바이스

직접적으로 연산을 수행하는 연산 장치를 디바이스라고 한다. CPU, GPU, FPGA 등이 디바이스에 해당한다.
하나의 플랫폼 안에 여러 개의 디바이스가 존재할 수 있다.

clGetDeviceIDs 함수를 이용하여 플랫폼에 속한 디바이스를 가져올 수 있다.

// 3. 디바이스 가져오기
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, devices, NULL); // 디바이스를 가져온다.

4. 컨텍스트

커널이 실행되는 환경을 컨텍스트라고 하며, 디바이스에 메모리 객체와 명령큐를 연결시켜주는 역할을 한다. 연산을 수행하는 디바이스, 피연산될 메모리 객체, 연산을 지닌 명령큐가 모두 컨텍스트 위에서 연결되어 상호작용하기 때문에, 컨텍스트는 일종의 랑데뷰라고 할 수 있다.

clCreateContext 함수를 이용하여 컨텍스트를 생성할 수 있다.

// 4. 컨텍스트 생성하기
context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL); // 컨텍스트를 생성한다.

5. 명령큐

디바이스에서 수행될 연산을 순서대로 담는 큐를 명령큐라고 한다. 컨텍스트와 디바이스를 명령큐로 연결한 후 명령큐에 커널을 집어넣으면 디바이스에서 커널을 꺼내서 수행한다.

clCreateCommandQueue 함수를 이용하여 명령큐를 생성할 수 있다.

// 5-1. 명령큐 생성하기
queue = clCreateCommandQueue(context, device, 0, NULL); // 명령큐를 생성한다.

clEnqueueNDRangeKernel 함수를 이용하여 명령큐에 커널을 삽입할 수 있다.

// 5-2. 명령큐에 커널 삽입하기
clEnqueueNDRangeKernel(queue, kernel, 2, NULL, globalSize, NULL, 0, NULL, NULL); // 명령큐에 커널을 삽입한다.

6. 프로그램

커널의 집합을 프로그램이라고 한다. .cl 파일에 있는 커널 소스를 컴파일하여 하나의 빌드된 형태로 존재한다. 디바이스에서는 이 프로그램 객체 내에서 특정 커널을 선택하여 실행할 수 있다.

clCreateProgramWithSource 함수를 이용하여 .cl 파일로부터 프로그램 객체를 생성할 수 있으며, clBuildProgram 함수를 이용하여 커널 소스 코드를 컴파일할 수 있다.

// 6. 프로그램 객체 생성하기
program = clCreateProgramWithSource(context, 1, (const char**)&source, NULL, NULL); // 프로그램 객체를 생성한다.
build_status = clBuildProgram(program, 1, &device, NULL, NULL, NULL); // 커널 소스 코드를 디바이스에서 실행 가능한 바이너리로 빌드한다.

.cl 파일은 파일 입출력을 이용해 해당 파일을 통째로 읽어오면 된다.

[C Style]
	FILE* fp = fopen("helloworld.cl", "rb");
	cl_int status = fseek(fp, 0, SEEK_END);
	long int size = ftell(fp);
	rewind(fp);
	char* source = (char*)malloc(sizeof(char) * (size + 1));
	fread(source, sizeof(char), size, fp);
	source[size] = '\0';

[C++ Style]
	std::ifstream file("helloworld.cl");
	char* source = std::string(std::istreambuf_iterator<char>(file), (std::istreambuf_iterator<char>())).c_str();

7. 커널

디바이스 위에서 실질적으로 수행되는 함수 형태의 단위 연산을 커널이라고 한다. 커널을 명령큐에 넣음으로써 디바이스에서 실행시킬 수 있다.

clCreateKernel 함수를 이용하여 프로그램 객체 내에서 특정 커널을 선택하여 커널 객체를 생성할 수 있다.

// 7-1. 커널 객체 생성하기
kernel = clCreateKernel(program, "myKernel", NULL); // 프로그램 객체로부터 커널 객체를 생성한다.

clSetKernelArg 함수를 이용하여 커널에 메모리 버퍼를 인자로 지정할 수 있다.

// 7-2. 커널에 인자 지정하기
clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufferA); // kernel의 0번째 인자에 bufferA를 지정한다.
clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufferB); // kernel의 1번째 인자에 bufferB를 지정한다.
clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufferC); // kernel의 2번째 인자에 bufferB를 지정한다.

8. 메모리 객체

호스트-디바이스 간 메모리 복사를 위해 사용되는 버퍼를 메모리 객체라고 한다. 메모리 버퍼라고도 하는데, 컨텍스트 위에서 입출력을 위한 메모리 버퍼를 정의하고 이를 디바이스에 전달함으로써 디바이스에 값을 입력하고 결과값을 받아올 수 있다.

clCreateBuffer 함수를 이용하여 메모리 객체를 생성할 수 있다.

// 8-1. 메모리 객체 생성하기
cl_mem bufferA = clCreateBuffer(context, CL_MEM_READ_WRITE, HEIGHT * WIDTH * sizeof(int), NULL, NULL); // HEIGHT * WIDTH * sizeof(int) 크기의 메모리 객체를 생성한다.

clEnqueueWriteBuffer 함수를 이용하여 메모리 버퍼를 명령큐에 삽입할 수 있다.

// 8-2. 명령큐에 메모리 객체 삽입하기
clEnqueueWriteBuffer(queue, bufferA, CL_TRUE, 0, HEIGHT * WIDTH * sizeof(int), arrayA, 0, NULL, NULL); // bufferA를 queue에 삽입한다.

3. 간단한 예제

예제를 통하여 위의 일련의 과정들을 수행해보자.

0. 프로젝트 생성

Visual Studio를 열어서 새 프로젝트를 생성한다.

새 프로젝트 만들기에서 Empty OpenCL Project for Windows 를 클릭한다.

생성하면 Source Files 안에 host.cpp 파일이 있다.

해당 파일을 열어서 아래의 예제 코드를 하나씩 따라서 작성해보자.

1. 디바이스 정보 출력
2. 커널 실행

1. 디바이스 정보 출력

clGetPlatformIDs와 clGetDeviceIDs를 이용하여 플랫폼과 디바이스를 가져와보자.
가져온 후 디바이스로부터 각종 정보를 가져와 출력해보자.

Source Files 폴더 안에 아래 파일을 만들어준다.

• host.cpp

// Add you host code

#include <cstdio>
#include <cstdlib>

#include "CL\cl.h"

int main(void)
{
	cl_uint platformCount;
	cl_platform_id* platforms;
	cl_uint deviceCount;
	cl_device_id* devices;
	size_t valueSize;
	char* value;
	cl_uint maxComputeUnits;

	// 1. 플랫폼 가져오기
	clGetPlatformIDs(0, NULL, &platformCount);
	printf("Total Platform Count: %d\n", platformCount);
	platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * platformCount);
	clGetPlatformIDs(platformCount, platforms, NULL);

	for (int i = 0; i < platformCount; i++)
	{
		printf("Platform %d\n", i);

		// 2. 디바이스 가져오기
		clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, NULL, &deviceCount);
		printf("\tTotal Device Count: %d\n", deviceCount);
		devices = (cl_device_id*)malloc(sizeof(cl_device_id) * deviceCount);
		clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, deviceCount, devices, NULL);

		// 3. 디바이스 정보 가져오기
		for (int j = 0; j < deviceCount; j++)
		{
			printf("\tDevice %d\n", j);

			// 3-1. 디바이스 이름 출력
			clGetDeviceInfo(devices[j], CL_DEVICE_NAME, 0, NULL, &valueSize);
			value = (char*)malloc(valueSize);
			clGetDeviceInfo(devices[j], CL_DEVICE_NAME, valueSize, value, NULL);
			printf("\t\tDevice Name: %s\n", value);
			free(value);

			// 3-2. 디바이스 버전 출력
			clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, 0, NULL, &valueSize);
			value = (char*)malloc(valueSize);
			clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, valueSize, value, NULL);
			printf("\t\tDevice Version: %s\n", value);
			free(value);

			// 3-3. 드라이버 버전 출력
			clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, 0, NULL, &valueSize);
			value = (char*)malloc(valueSize);
			clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, valueSize, value, NULL);
			printf("\t\tDriver Version: %s\n", value);
			free(value);

			// 3-4. OpenCL C 버전 출력
			clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, 0, NULL, &valueSize);
			value = (char*)malloc(valueSize);
			clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, valueSize, value, NULL);
			printf("\t\tOpenCL C Version: %s\n", value);
			free(value);

			// 3-5. 연산 유닛 수 출력
			clGetDeviceInfo(devices[j], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(maxComputeUnits), &maxComputeUnits, NULL);
			printf("\t\tMax Compute Units: %d\n", maxComputeUnits);
		}
	}

	system("pause");
	return 0;
}

작성 후 Ctrl + F5 를 입력하여 실행한다.

위와 같이 플랫폼 및 디바이스 정보가 출력되는 것을 볼 수 있다.

위 5가지 외에도 여러 가지 디바이스 정보를 가져올 수 있다.

디바이스 정보	설명
CL_DEVICE_NAME	디바이스 이름
CL_DEVICE_TYPE	디바이스 종류
CL_DEVICE_PLATFORM	소속된 플랫폼
CL_DEVICE_VENDOR	제조사 이름
CL_DEVICE_VENDOR_ID	제조사 ID
CL_DEVICE_AVAILABLE	사용 가능 여부
CL_DEVICE_MAX_CLOCK_FREQUENCY	최대 클럭 주파수 (MHz)
CL_DEVICE_GLOBAL_MEM_SIZE	Global Device Memory 크기
...	...

아래 OpenCL 공식 페이지에서 더 많은 디바이스 정보를 확인할 수 있다.

https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clGetDeviceInfo.html

2. 커널 실행

컨텍스트, 큐, 프로그램을 이용하여 디바이스 위에서 커널을 실행해보자.

아래 예제는 0번째 플랫폼의 0번째 디바이스를 선택하여 8x8 배열 A와 B를 더한 값을 배열 C에 출력하는 예제이다.

위에서 본 바, 필자의 0번째 디바이스가 인텔 내장 그래픽이었음을 참고한다.

Source Files 폴더 안에 아래 파일을 만들어준다.

• host.cpp

// Add you host code

#include <stdio.h>
#include <stdlib.h>
#include <string.h>

#include <CL/cl.h>

#define HEIGHT 8
#define WIDTH  8

int main(void)
{
	// 0. 사용자 데이터 정의
	int platformNum = 0;
	int deviceNum = 0;
	const char* sourceFile = "helloworld.cl";
	const char* kernelName = "arrayAdd";

	int arrayA[HEIGHT][WIDTH] = {
		{ 1, 1, 1, 1, 2, 2, 2, 2, },
		{ 1, 1, 1, 1, 2, 2, 2, 2, },
		{ 1, 1, 1, 1, 2, 2, 2, 2, },
		{ 1, 1, 1, 1, 2, 2, 2, 2, },
		{ 3, 3, 3, 3, 4, 4, 4, 4, },
		{ 3, 3, 3, 3, 4, 4, 4, 4, },
		{ 3, 3, 3, 3, 4, 4, 4, 4, },
		{ 3, 3, 3, 3, 4, 4, 4, 4, },
	};
	int arrayB[HEIGHT][WIDTH] = {
		{ 1, 1, 1, 1, 1, 1, 1, 1, },
		{ 1, 1, 1, 1, 1, 1, 1, 1, },
		{ 1, 1, 1, 1, 1, 1, 1, 1, },
		{ 1, 1, 1, 1, 1, 1, 1, 1, },
		{ 1, 1, 1, 1, 1, 1, 1, 1, },
		{ 1, 1, 1, 1, 1, 1, 1, 1, },
		{ 1, 1, 1, 1, 1, 1, 1, 1, },
		{ 1, 1, 1, 1, 1, 1, 1, 1, },
	};
	int arrayC[HEIGHT][WIDTH] = { 0, };

	// 입력 데이터 출력
	printf("Array A:\n");
	for (int i = 0; i < HEIGHT; i++)
	{
		for (int j = 0; j < WIDTH; j++)
		{
			printf("%d ", arrayA[i][j]);
		}
		printf("\n");
	}
	printf("\n");
	printf("Array B:\n");
	for (int i = 0; i < HEIGHT; i++)
	{
		for (int j = 0; j < WIDTH; j++)
		{
			printf("%d ", arrayB[i][j]);
		}
		printf("\n");
	}
	printf("\n");

	// 1. platform 가져오기
	cl_uint platformCount;
	clGetPlatformIDs(0, NULL, &platformCount);
	cl_platform_id* platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * platformCount);
	clGetPlatformIDs(platformCount, platforms, NULL);

	// 2. device 가져오기
	cl_uint deviceCount;
	clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, 0, NULL, &deviceCount);
	cl_device_id* devices = (cl_device_id*)malloc(sizeof(cl_device_id) * deviceCount);
	clGetDeviceIDs(platforms[platformNum], CL_DEVICE_TYPE_ALL, deviceCount, devices, NULL);
	cl_device_id device = devices[deviceNum];

	// 3. context 생성하기
	cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);

	// 4. command queue 생성하기
	cl_command_queue queue = clCreateCommandQueue(context, device, 0, NULL);

	// 5. source 가져오기
	FILE* fp = fopen(sourceFile, "rb");
	cl_int status = fseek(fp, 0, SEEK_END);
	long int size = ftell(fp);
	rewind(fp);
	char* source = (char*)malloc(sizeof(char) * (size + 1));
	fread(source, sizeof(char), size, fp);
	source[size] = '\0';

	// 6. program 빌드하기
	cl_program program = clCreateProgramWithSource(context, 1, (const char**)&source, NULL, NULL);
	cl_int build_status = clBuildProgram(program, 1, &device, NULL, NULL, NULL);

	// 7. kernel 생성하기
	cl_kernel kernel = clCreateKernel(program, kernelName, NULL);

	// 8. memory buffer 생성하기
	cl_mem bufferA = clCreateBuffer(context, CL_MEM_READ_WRITE, HEIGHT * WIDTH * sizeof(int), NULL, NULL);
	cl_mem bufferB = clCreateBuffer(context, CL_MEM_READ_WRITE, HEIGHT * WIDTH * sizeof(int), NULL, NULL);
	cl_mem bufferC = clCreateBuffer(context, CL_MEM_READ_WRITE, HEIGHT * WIDTH * sizeof(int), NULL, NULL);

	// 9. command queue에 memory buffer 삽입하기
	clEnqueueWriteBuffer(queue, bufferA, CL_TRUE, 0, HEIGHT * WIDTH * sizeof(int), arrayA, 0, NULL, NULL);
	clEnqueueWriteBuffer(queue, bufferB, CL_TRUE, 0, HEIGHT * WIDTH * sizeof(int), arrayB, 0, NULL, NULL);

	// 10. kernel argument 설정하기
	clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufferA);
	clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufferB);
	clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufferC);

	// 11. command queue에 kernel 삽입하기
	size_t globalSize[2] = { HEIGHT, WIDTH };
	clEnqueueNDRangeKernel(queue, kernel, 2, NULL, globalSize, NULL, 0, NULL, NULL);

	// 12. 연산 완료될 때까지 대기하기
	clFinish(queue);

	// 13. 출력 버퍼에 결과 반환하기
	clEnqueueReadBuffer(queue, bufferC, CL_TRUE, 0, HEIGHT * WIDTH * sizeof(int), arrayC, 0, NULL, NULL);

	// 결과 데이터 출력
	printf("Array C:\n");
	for (int i = 0; i < HEIGHT; i++)
	{
		for (int j = 0; j < WIDTH; j++)
		{
			printf("%d ", arrayC[i][j]);
		}
		printf("\n");
	}
	printf("\n");

	system("pause");

	// 할당한 역순으로 메모리 해제
	clReleaseMemObject(bufferA);
	clReleaseMemObject(bufferB);
	clReleaseMemObject(bufferC);
	clReleaseKernel(kernel);
	clReleaseProgram(program);
	free(source);
	fclose(fp);
	clReleaseCommandQueue(queue);
	clReleaseContext(context);
	clReleaseDevice(device);
	free(devices);
	free(platforms);

	return 0;
}

OpenCL Files 폴더 안에 아래 파일을 만들어준다.

• helloworld.cl

// TODO: Add OpenCL kernel code here.

__kernel void arrayAdd(__global int* pA, __global int* pB, __global int* pC)
{
    const int x     = get_global_id(0);
    const int y     = get_global_id(1);
    const int width = get_global_size(0);

    const int id = y * width + x;

    pC[id] = pA[id] + pB[id];
}

작성 후 Ctrl + F5 를 입력하여 실행한다.

실행하면 배열 C에 배열 A와 B를 더한 값이 들어가있는 것을 볼 수 있다.

여기까지 따라왔다면 여러분은 진짜 대단한 것이다!
스스로의 자존감을 고취시키도록 하자.

4. 솔루션 첨부

글을 쓰면서 직접 작성한 Visual Studio 솔루션을 깃허브에 푸시하였다.
https://github.com/pikamonvvs/OpenCL-Helloworld

필요하면 내려받아서 실행해보면 될 것 같다.

5. 참고 문헌

아래 글을 참고하여 작성하였다.

• OpenCL 프로그래밍 가이드 (저자: 아프탑 문시) - 책 링크
• 후로링의 프로그래밍 이야기 - https://hoororyn.tistory.com/3
• 후로링의 프로그래밍 이야기 - https://hoororyn.tistory.com/4