OpenCL

역사

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은 맥 OS X 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 구현은 LLVM 및 Clang 컴파일러를 기반으로 한 것으로 알려졌다.
AMD는 OpenCL 및 다이렉트엑스 11을 지원하는 대신 AMD 고유의 스트림 프레임워크 내 "Close to Metal"을 포기하기로 결정했다.RapidMind는 OpenCL 채택을 공식 선언하여 자신들의 개발 플랫폼의 밑단에 쓰기로 하였는데, 여러 제조업체의 GPU를 단일 인터페이스를 통해 지원하기 위해서였다.엔비디아는 2008년 12월 9일 자사의 GPU 컴퓨팅 툴킷에서 OpenCL 1.0을 완벽히 지원한다고 발표하였다. OpenCL과 쿠다를 비교한다면 두개의 컴퓨터 언어를 비교하는 것과 비슷하다는 입장이다.
OpenCL 명세서는 크로노스에서 개발 중이며, 관심 있는 어떤 회사에라도 개방되어 있다.

예제

다음은 고속 푸리에 변환을 하는 예제이다. 호스트 프로그램은 다음과 같다.
// GPU 장치와 계산 context를 생성한다context = clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL);// 작업 대기열을 생성한다clGetDeviceIDs(NULL, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);queue = clCreateCommandQueue(context, device_id, 0, NULL);// 버퍼 메모리 객체를 생성한다memobjs[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float)*2*num_entries, srcA, NULL);memobjs[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float)*2*num_entries, NULL, NULL);// 계산 프로그램을 생성한다program = clCreateProgramFromSource(context, 1, &fft1D_1024_kernel_src, NULL, NULL);// 계산 프로그램 실행 코드를 생성한다clBuildProgram(program, 0, NULL, NULL, NULL, NULL);// 계산 커널을 생성한다kernel = clCreateKernel(program, "fft1D_1024", NULL);// args 값을 설정한다clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobjs[0]);clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobjs[1]);clSetKernelArg(kernel, 2, sizeof(float)*(local_work_size[0]+1)*16, NULL);clSetKernelArg(kernel, 3, sizeof(float)*(local_work_size[0]+1)*16, NULL);// 커널을 실행시킨다global_work_size[0] = num_entries;local_work_size[0] = 64;clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
실제 계산이 이루어지는 커널 함수는 다음과 같다.
// 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);}