Introductory Tutorial to OpenCL
Skip Navigation LinksHome > GPU Tools > ATI Stream SDK > Introductory Tutorial to OpenCL

Benedict R. Gaster, AMD Architect, OpenCL

AMD 는 CPU 에서 동작하는 OpenCL 의 퍼블릭 베타를 발표했습니다. 조만간 AMD 최신 GPU 에서도 구현될 것입니다. OpenCL 은 최신 기술입니다. 기술 규격에 대해서는 홈페이지(www.khronos.org/registry/cl/)에서 소개하고 있습니다만, 예제를 곁들인 기본 설명을 담고 있는 문서는 아직 얼마 되지 않습니다. 이 글은 OpenCL 에 대한 이해와 응용을 돕기 위해 만들어졌습니다.

먼저:

* 저는 AMD 에서 일하고 있으며, 제가 소개하는 예제들은 윈도우와 리눅스 환경에서 테스트할 것입니다. 하지만 저는 기종에 관계없이 OpenGL 의 활용에 중점을 둘 것입니다. 모든 예제들은 순전히 OpenGL 로 구현되며 따라서 기종에 관계없이 동작될 것입니다.
* 제가 제공하는 모든 예제들은 AMD 의 구현이 아니어도 동작될 수 있게끔 하였습니다만, AMD 에서 구현하지 않은 환경에서 테스트하지는 않았습니다. 따라서 해당 시스템에서 정상 동작되지 않을 수 있습니다. 그럴 경우 Stream Computing 포럼에 알려주시면 제가 코드를 다시 정리하여 업데이트 하도록 하겠습니다.

다음의 “Hello World” 예제는 OpenGL 을 간단히 소개하는 프로그램입니다. 이번 예제를 비롯해서 몇 가지 예제를 통해서 다음과 같은 내용을 다룰 것입니다:

* OpenCL 을 위한 플랫폼과 디바이스 레이어
* 프로그램 컴파일과 커널 객체
* 버퍼 관리
* 커널 실행
* 커널 프로그래밍 – 기초
* 커널 프로그래밍 – 동기화
* 행렬 곱셈 – 예제
* 커널 프로그래밍 – 빌트인

= The “Hello World” program in OpenCL

OpenCL 예제 작성시 고려했던 내용입니다:

* OpenCL 이 지정한 호스트 API 는 C89 호환이며 C++ 혹은 여타 다른 언어와의 연동은 언급하지 않고 있습니다. 현재 다른 언어와의 연동을 연구중이며 (이 글 마지막에 소개된 링크를 보세요) 특별히 C++ 과의 연동을 활발히 연구중입니다. 본 예제에서는 C++ 에서 OpenCL 활용법을 보여드리겠습니다. OpenCL 1.0 규격에서 해당되는 C API 를 참조하세요. 소스에서 어떻게 OpenCL 함수를 호출하고 있으며 C++ 에서 인수 전달 방법 등을 볼 수 있습니다.
* OpenCL 은 계산 장치 프로그램을 작성하기 위하여 C 형태의 언어를 사용합니다. 이 프로그램은 OpenCL 런타임에 char * 인수를 받는 API 호출로 전달됩니다. 보통 이 프로그램은 다른 소스 파일에 저장하는 것이 편리합니다. 다음의 예제들에서 디바이스 프로그램은 name_kernels.cl 과 같은 이름을 가진 파일에 저장합니다. name 은 예제 내용에 따라 바뀌지만 _kernels.cl 은 동일합니다. 해당되는 디바이스 프로그램은 실행시 읽어서 OpenCL API 로 전달됩니다. 그 외에도 다른 방법들이 있습니다만 이 방법이 이해하기 쉽기 때문에 선택했습니다.

첫 번째 OpenGL 프로그램 작성을 위하여 먼저 호스트 프로그램 소스부터 시작해 보겠습니다.

= Header files

다른 C++ 외부 API 처럼 OpenCL API 를 사용하려면 헤더를 include 해야 합니다. 보통 기본 include 디렉토리의 CL 디렉토리에 저장되어 있습니다. C++ 에서 사용을 위해서 (C API 인 cl.h 를 대신한)

#include
#define __NO_STD_VECTOR // Use cl::vector and cl::string and
#define __NO_STD_STRING // not STL versions, more on this later
#include

예제에서 사용하기 위한 C++ 헤더를 몇 개 추가합니다.

#include
#include
#include
#include
#include
#include

OpenCL 장치로부터 동적으로 “Hello World
” 문자열을 리턴할 것이므로 상수로 선언을 하겠습니다.

const std::string hw("Hello World
");

= Errors

OpenCL API 대부분 공통적으로 오류가 발생했을 때 함수의 리턴값으로 (cl_int 형) 에러 코드를 돌려주거나, 발생한 에러를 저장하여 유저에게 인수로 전달합니다. 에러 코드는 API 에서 오류가 발생했을 때 상황을 알아낼 수 있으므로 중요합니다. 하지만 지금은 문제를 간단히 하기 위해서 checkEr() 함수를 선언하여 어떤 작업이 무사히 끝났는지만 검사하고 넘어갑니다. 에러가 없으면 OpenCL 은 CL_SUCCESS 를 리턴합니다. 에러가 발생했을 경우 메세지를 발생하고 종료하거나 리턴하게 됩니다.

inline void
checkErr(cl_int err, const char * name)
{
if (err != CL_SUCCESS) {
std::cerr << "ERROR: " << name
<< " (" << err << ")" << std::endl;
exit(EXIT_FAILURE);
}

C++ 에서 에러를 처리하는 기본 방식은 exception 을 이용하는 것입니다. OpenCL C++ 바인딩에서는 유사한 인터페이스를 제공합니다. 뒷부분에 가서 exception 과 C++ 특성에 맞는 몇 가지 옵션 기능을 다루도록 하겠습니다. 여기서는 마지막으로 우리 첫 번째 OpenCL 프로그램에 필요한 main 함수를 보겠습니다.

= OpenCL Contexts

OpenCL 을 초기화하여 사용하는 첫 번째 단계는 컨텍스트를 만드는 것입니다. 나머지 작업들 (디바이스와 메모리를 만들고 컴파일 하고 실행하는 것) 은 컨텍스트 내에서 이루어집니다. 컨텍스트에는 사용 가능한 디바이스의 갯수 (예를 들어 CPU 와 GPU), 그리고 컨텍스트 내에는 장치 간 메모리를 유지합니다. 자세한 사항은 뒤에 다루겠습니다. 지금은 CPU 장치를 위한 CL_DEVICE_TYPE_CPU 디바이스 하나만 사용합니다. 만약 OpenCL 에서 지원한다면 CL_DEVICE_TYPE_GPU 등을 쓸 수도 있겠지요. 생성자 cl::Context 는 반드시 성공할 것이므로 변수 err 에는 CL_SUCCESS 가 들어있을 것입니다.

int
main(void)
{
cl_int err;

cl::Context context(
CL_DEVICE_TYPE_CPU,
NULL,
NULL,
NULL,
&err);
checkErr(err, "Conext::Context()");

실제 작업이 이루어지는 계산 장치를 파고들기 전에 먼저 장치 내에서 작업을 수행하는 커널의 결과값을 저장할 OpenCL 버퍼를 할당합니다. 여기서는 “Hello World
” 문자열이겠지요. 지금은 간단히 호스트에서 메모리를 할당하여 OpenCL 이 직접 메모리를 사용하도록, 버퍼를 생성할 때 CL_MEM_USE_HOST_PTR 플래그를 같이 전달합니다.

char * outH = new char[hw.length()+1];
cl::Buffer outCL(
context,
CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
hw.length()+1,
outH,
&err);
checkErr(err, "Buffer::Buffer()");

= OpenCL Devices

OpenCL 의 대부분 작업은 주어진 컨텍스트 내에서 수행됩니다. 예를 들어, 버퍼 (1D 메모리 공간) 와 이미지 (2D 혹은 3D 메모리 공간) 할당은 모두 컨텍스트 작업입니다. 하지만 특정 장치 작업도 있습니다. 예를 들어 프로그램 컴파일과 커널 실행은 장치 기준이고, 이를 위해서는 특정 장치 핸들이 필요합니다. 특정 장치의 핸들은 어떻게 얻어낼까요? 장치 컨텍스트를 요구하면 됩니다. OpenCL 은 특정 장치의 정보를 얻어내는 방법을 제공합니다. C++ API 에서는 object.getInfo() 와 같은 형식이 됩니다. 컨텍스트로부터 장치 정보를 얻어오는 경우는:

cl::vector devices;
devices = context.getInfo();
checkErr(
devices.size() > 0 ? CL_SUCCESS : -1, "devices.size() > 0");

이제 컨텍스트에 속해있는 장치 리스트를 얻었습니다. 지금 경우는 CPU 장치겠지요, 이제 계산 프로그램 (장치 내에서 작동되는 프로그램) 을 읽어서 빌드하는 작업이 필요합니다. 다음 코드는 OpenCL 디바이스 프로그램을 디스크로부터 읽어서 문자열로 치환한 뒤 helper constructor 를 이용하여 cl::Program::Sources 를 생성하겠습니다. 이 객체는 컨텍스트 내에서 특정한 디바이스를 위해 만들어진 것입니다.

std::ifstream file("lesson1_kernels.cl");
checkErr(file.is_open() ? CL_SUCCESS:-1, "lesson1_kernel.cl");

std::string prog(
std::istreambuf_iterator(file),
(std::istreambuf_iterator()));

cl::Program::Sources source(

1,
std::make_pair(prog.c_str(), prog.length()+1));

cl::Program program(context, source);
err = program.build(devices,"");
checkErr(file.is_open() ? CL_SUCCESS : -1, "Program::build()");

프로그램에는 커널이라고 하는 많은 엔트리 포인트가 있고, 커널 객체로 빌드해야 합니다. 커널은 문자열 형식의 커널 이름으로 일대일 대응되어 있고, 계산 프로그램 내에서는 __kernel 형식으로 정의된 함수입니다. 여기서는 cl::kernel 형의 객체 kernel 을 빌드합니다. 커널 인수는, 특정 인수를 인덱스와 값으로 받아들이는 kernel.setArg() 라는 C++ API 를 이용하여 전달합니다.

cl::Kernel kernel(program, "hello", &err);
checkErr(err, "Kernel::Kernel()");

err = kernel.setArg(0, outCL);
checkErr(err, "Kernel::setArg()");

이제 코드가 준비되었고, 결과를 계산합니다 (여기서는 “Hello World
” 문자열을 버퍼에 넣습니다). 모든 디바이스 계산은 커맨드 큐를 이용하는데, 사용하려는 장치의 가상 인터페이스입니다. 각 커맨드 큐는 특정 장치와 일대일 대응되며, 해당 컨텍스트에서 클래스의 생성자 cl::CommandQueue 를 호출함으로써 생성됩니다. 커널은 queue.enqueuNDRangeKernel 을 이용하여 해당 장치로부터 커널을 실행합니다. 커널은 주어진 자원 내에서 1, 2, 3 차원 도메인 인덱스를 병렬로 수행합니다. 실행된 도메인 내 인덱스의 총 갯수를 글로벌 워크 사이즈 라고 하며, 워크 아이템이라고 하는 것의 갯수입니다. 워크 아이템끼리 정보 교환이 필요할 경우 워크 아이템을 워크 그룹으로 묶습니다. 워크 그룹은 서브 인덱스 함수(로컬 워크 사이즈) 로 정의되며, 글로벌 실행 도메인에서 지정한 차원에 대한 각 차원의 크기를 지정합니다. 그 외에도 커널 실행에 대한 여러 가지 고려할 점이 있으며, 향후 다른 튜토리얼에서 다루겠습니다. 지금은 Hello World 에서 각각의 워크 아이템은 결과 문자열의 글자를 계산하며, , 프로그램 앞에서 const std::string 으로 선언했던 hw 의 길이hw.length()+1 만 실행해도 충분합니다. 나머지 워크 아이템은 NULL 을 넣습니다.

cl::CommandQueue queue(context, devices[0], 0, &err);
checkErr(err, "CommandQueue::CommandQueue()");


cl::Event event;
err = queue.enqueueNDRangeKernel(
kernel,
cl::NullRange,
cl::NDRange(hw.length()+1),
cl::NDRange(1, 1),
NULL,
&event);
checkErr(err, "ComamndQueue::enqueueNDRangeKernel()");

위의 마지막 enqueueNDRangeKernel 의 마지막 인수는 cl::Event 객체였는데, 해당 커맨드의 상태를 저장할 때 쓰입니다. (예를 들어 작업 종료 같은) 커맨드가 종료될 때까지 대기하는 wait() 를 지원합니다. 이 함수는 queue.enqueueReadBuffer() 를 이용해여 결과물을 호스트 메모리로 읽어들이기 전에 커널 작업이 종료되었는지를 확실히 할 때 유용합니다. 결과가 호스트 메모리로 돌아온 다음에는 stdout 으로 결과물을 출력하고 프로그램을 종료합니다.

event.wait();
err = queue.enqueueReadBuffer(
outCL,
CL_TRUE,
0,
hw.length()+1,
outH);
checkErr(err, "ComamndQueue::enqueueReadBuffer()");
std::cout << outH;
return EXIT_SUCCESS;
}

마지막으로 디바이스 프로그램 lesson1_kernels.cl 을 완성하기 위해서 가외의 엔트리 포인트 hello 가 필요합니다. 이 커널 프로그램은 간단합니다. get_global_id() 를 이용하여 실행 도메인의 인덱스를 얻어냅니다. 이것을 문자열 hw 의 인덱스로 삼아서 그 값을 출력 배열 out 에 넣습니다.

__constant char hw[] = "Hello World
";
__kernel void hello(__global char * out)
{
size_t tid = get_global_id(0);
out[tid] = hw[tid];
}

안정성을 위해서 쓰레드 ID (tid) 가 hw 의 범위를 벗어나지 않는지 검사하는 것이 좋습니다만, 여기서는 queue.enqueueNDRangeKernel() 호출이 정확하다고 가정하겠습니다.

= Building and running

리눅스에서는 OpenCL 프로그램을 빌드하려면 명령 한 줄이면 됩니다. 예를 들어
gcc –o hello_world –Ipath-opencl-incude –Lpath-opencl-libdir lesson1.cpp –lOpenCL

실행할 때에는:
LD_LIBRARY_PATH=path-opencl-libdir ./hello_world

윈도우 비쥬얼 스튜디오 커맨드 윈도우에서의 예는:
cl /Fehello_world.exe /Ipath-opencl-include lesson.cpp path-opencl-libdir/OpenCL.lib

OepnCL.dll 이 path 에 들어있다고 가정하면 실행할 때는:
.hello_world

다음과 같은 문자열을 출력합니다 pm stdout:
Hello World

OpenCL 입문 튜토리얼이 끝났습니다. 궁금한 점이나 덧붙이고 싶은 점이 있으시면 Stream 포럼을 방문해 주세요.

= Useful Links

다음 목록에는 OpenCL 을 C 이외 언어에서 사용하는 방법에 대한 내용입니다. 모든 내용을 다 테스트해본 것은 아닙니다만, 유용한 정보이길 바랍니다:

* OpenCL specification and headers:
http://www.khronos.org/registry/cl/
* OpenCL technical forum:
http://www.khronos.org/message_boards/viewforum.php?f=28
* The C++ bindings used in this tutorial can be found on the OpenCL web page at Khronos, along with complete documentation:
http://www.khronos.org/registry/cl/
* Python bindings can be found here:
http://pyopencl.next-touch.com/
* C# bindings can be found here:
http://www.khronos.org/message_boards/viewtopic.php?f=28&t=1932
* An Introduction to OpenCL:
http://ati.amd.com/technology/streamcomputing/intro_opencl.html