OpenCL

이 문서는 프레임워크에 관한 것입니다. OpenCL이라 불렸던 암호화 라이브러리에 대해서는 Botan 문서를 참조하십시오.

OpenCL

개발자	크로노스 그룹
최근 버전	1.2
운영 체제	크로스 플랫폼
종류	API
라이선스	Royalty Free
웹사이트	www.khronos.org/opencl

OpenCL(Open Computing Language)은 개방형 범용 병렬 컴퓨팅 프레임워크이다. 여러 CPU, GPU 및 기타 프로세서 등으로 이루어진 다중 플랫폼에서 구동해야 하는 데이터 및 태스크 병렬성을 갖는 프로그램을 작성할 수 있게 해 주는 프레임워크이다. OpenCL은 C99에 기반한 새로운 언어를 포함하고 있다. 이 언어를 가지고 혼성 플랫폼을 정의하고 제어할 수 있는 커널, API 등을 작성할 수 있다. OpenCL은 임무 기반 병렬 또는 데이터 기반 병렬 컴퓨팅을 제공한다.

OpenCL이 만들어진 이유는 OpenGL이나 OpenAL이 만들어진 이유와 비슷하다. OpenGL과 OpenAL은 각각 3차원 컴퓨터 그래픽스 및 컴퓨터 오디오에 대한 산업계의 개방형 표준이다.

OpenCL은 그래픽 처리 장치(GPU)의 역량을 그래픽 처리 이외의 영역(범용 GPU)으로 확장한다. OpenCL은 비영리 기술 컨소시엄인 크로노스 그룹(Khronos Group)에서 관리하고 있다.

역사

OpenCL은 애플이 최초로 개발했으며 OpenCL의 상표권도 애플이 가지고 있다. 그 후 AMD, 인텔, 엔비디아 등과 함께 애플은 문서를 다듬어 최초의 제안서(proposal)를 크로노스 그룹에 제출하였다. 2008년 6월 16일 크로노스 컴퓨트 워킹 그룹(Khronos Compute Working Group)이 결성되었다. 크로노스 컴퓨트 워킹 그룹에는 CPU, GPU, 임베디드 프로세서, 소프트웨어 업체들이 참여하였다. 2008년 6월부터 5개월 동안 이 그룹은 OpenCL 1.0을 만들기 위해 작업하였다. 2008년 11월 18일 마침내 기술 규격 정보를 담은 OpenCL 1.0 명세서(specification)가 완성되었다. 이 기술 명세서는 크로노스 그룹 그룹원들이 검토(review)하였다. 그 후, 2008년 12월 8일 공식적으로 발표되었다.

OpenCL은 맥 오에스 텐 10.6 레퍼드 (스노 레퍼드)에서부터 지원된다. 프레스 릴리즈에서 다음과 같이 말하고 있다.

(비공식 번역임.) 스노 레오파드(MAC OS 10.6)는 오픈 컴퓨팅 언어(OpenCL)로 최신의 하드웨어에 대한 지원을 확장하였습니다. 이전에는 GPU의 방대한 기가플롭스 계산 능력을 그래픽 애플리케이션에만 사용해 왔지만, OpenCL을 통하여 이제 어떠한 응용 프로그램에서도 끌어와 쓸 수 있습니다. OpenCL은 C 프로그래밍 언어에 기반하고 있으며, 개방형 표준으로 제안되었습니다.

Snow Leopard further extends support for modern hardware with Open Computing Language (OpenCL), which lets any application tap into the vast gigaflops of GPU computing power previously available only to graphics applications. OpenCL is based on the C programming language and has been proposed as an open standard.

최초의 OpenCL 구현은 낮은 수준의 가상 머신(Low Level Virtual Machine, LLVM) 및 Clang 컴파일러 기술에 기반하여 구현된 것으로 알려졌다.

AMD는 OpenCL 및 다이렉트엑스 11을 지원하는 대신 AMD 고유의 스트림 프레임워크 내 "Close to Metal"을 포기하기로 결정했다. RapidMind는 OpenCL 채택을 공식 선언하여 자신들의 개발 플랫폼의 밑단에 쓰기로 하였는데, 여러 제조업체의 GPU를 단일 인터페이스를 통해 지원하기 위해서였다. 엔비디아는 2008년 12월 9일 자사의 GPU 컴퓨팅 툴킷에서 OpenCL 1.0을 완벽히 지원한다고 발표하였다. OpenCL과 쿠다를 비교한다면 두개의 컴퓨터 언어를 비교하는 것과 비슷하다는 입장이다.

OpenCL 명세서는 크로노스에서 개발 중이며, 관심 있는 어떤 회사에라도 개방되어 있다.

구현

2008/12/10 AMD와 엔비디아는 OpenCL 최초의 대중 시연을 실시하여 75분짜리 발표를 Siggraph Asia 2008에서 선보였다. AMD는 CPU 가속 OpenCL 시연으로 OpenCL의 코어 갯수에 대한 규모가변성을 보였고, 엔비디아는 GPU-가속 시범을 보였다.

2009/03/26 GDC 2009에서는 AMD와 하복이 최초로 AMD 라데온 HD 4000 시리즈 GPU 상에서 OpenCL을 이용하여 하복 클로스(Havok Cloth)를 가속시키는 시연을 실시하였다.

2009/04/20, 엔비디아가 OpenCL 조기 접근 프로그램 참가 개발자 대상으로 자신의 OpenCL 드라이버와 SDK를 배포한다고 발표하였다.

2009년 9월 애플에서 맥 오에스 텐 스노 레퍼드 버전에 OpenCL이 구현되었다. 스노 레퍼드에 포함되는 OpenCL은 초기에는 다음 GPU를 지원할 예정이다:

• NVIDIA Geforce 8600M GT, GeForce 8800 GT, GeForce 8800 GTS, Geforce 9400M, GeForce 9600M GT, GeForce GT 120, GeForce GT 130,
• ATI Radeon 4850, Radeon 4870.

예제

다음은 고속 푸리에 변환을 하는 예제이다.

// GPU 장치와 계산 context를 생성한다
context = clCreateContextFromType(CL_DEVICE_TYPE_GPU);
 
// 작업 대기열을 생성한다
queue = clCreateWorkQueue(context, NULL, NULL, 0);
 
// 버퍼 메모리 객채를 생성한다
memobjs[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float)*2*num_entries, srcA);
memobjs[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float)*2*num_entries, NULL);
 
// 계산 프로그램을 생성한다
program = clCreateProgramFromSource(context, 1, &fft1D_1024_kernel_src, NULL);
 
// 계산 프로그램 실행 코드를 생성한다
clBuildProgramExecutable(program, false, NULL, NULL);
 
// 계산 커널을 생성한다
kernel = clCreateKernel(program, "fft1D_1024");
 
// 작업 대상 칫수? 크기? 차원?을 포함하는 N-차원 범위 객체를 생성한다
// create N-D range object with work-item dimensions
global_work_size[0] = n;
local_work_size[0] = 64;
range = clCreateNDRangeContainer(context, 0, 1, global_work_size, local_work_size);
 
// args 값을 설정한다
clSetKernelArg(kernel, 0, (void *)&memobjs[0], sizeof(cl_mem), NULL);
clSetKernelArg(kernel, 1, (void *)&memobjs[1], sizeof(cl_mem), NULL);
clSetKernelArg(kernel, 2, NULL, sizeof(float)*(local_work_size[0]+1)*16, NULL);
clSetKernelArg(kernel, 3, NULL, sizeof(float)*(local_work_size[0]+1)*16, NULL);
 
// 커널을 실행시킨다
clExecuteKernel(queue, kernel, NULL, range, NULL, 0, NULL);

실제 계산 부분: (Fitting FFT onto the G80 Architecture에 기반함)

// This kernel computes FFT of length 1024. The 1024 length FFT is decomposed into 
// calls to a radix 16 function, another radix 16 function and then a radix 4 function 
 
__kernel void fft1D_1024 (__global float2 *in, __global float2 *out, 
                          __local float *sMemx, __local float *sMemy) { 
  int tid = get_local_id(0); 
  int blockIdx = get_group_id(0) * 1024 + tid; 
  float2 data[16]; 
 
  // 전역 메모리 입출력 영역 시작 주소 
  in = in + blockIdx;  out = out + blockIdx; 
 
  globalLoads(data, in, 64); // 한 덩어리로 전역 메모리 읽기
  fftRadix16Pass(data);       // 자리 변경 없이 radix-16 처리
  twiddleFactorMul(data, tid, 1024, 0);
 
  // 지역 메모리를 이용한 지역 shuffle 
  localShuffle(data, sMemx, sMemy, tid, (((tid & 15) * 65) + (tid >> 4))); 
  fftRadix16Pass(data);               // 자리 변경 없이 radix-16 처리 
  twiddleFactorMul(data, tid, 64, 4); // 회전 인수 곱셈
 
  localShuffle(data, sMemx, sMemy, tid, (((tid >> 4) * 64) + (tid & 15))); 
 
  // radix-4 함수 호출 4회 
  fftRadix4Pass(data); 
  fftRadix4Pass(data + 4); 
  fftRadix4Pass(data + 8);
  fftRadix4Pass(data + 12); 
 
  //한덩어리로 전역 메모리에 기록
  globalStores(data, out, 64); 
}

같이 보기

• GPU
• GPGPU
• CUDA
• Close to Metal
• BrookGPU
• Lib Sh
• 인텔 라라비

바깥 고리

• (영어) OpenCL 홈페이지
• (영어) OpenCL 기술 명세서
• (영어) OpenCL: What you need to know - article published in Macworld, August 2008
• (영어) HPCWire: OpenCL on the Fast Track
• (영어) The Khronos Group Releases OpenCL 1.0 Specification