OpenCL/문법

#!if top = 문서명1 != null ? 문서명1 : calleeTitle != null ? (i = calleeTitle.lastIndexOf("/")) != -1 ? calleeTitle.substr(0, i) : '상위 문서' : '상위 문서'
[[파일:상위 문서 아이콘.svg|align=left&width=21.6px]] {{{#!html  }}} 상위 문서: [[]]

1. 개요

2. 코드 도메인

2.1. Platform 코드

2.1.1. cl.h와 opencl.hpp의 차이

#!syntax cpp
#include <CL/CL.h>
#include <stdlib.h>

int main()
{
  cl_int clError = CL_SUCCESS;

  // Platform Query
  cl_uint         numPlatforms = 0;
  cl_platform_id *platforms    = NULL;

  clError = clGetPlatformIDs(0, NULL, &numPlatforms);
  cl_platform_id platforms =
      (cl_platform_id *)malloc(sizeof(cl_platform_id) * numPlatforms);
  clGetPlatformIDs(numPlatforms, platforms, NULL);

  cl_platform_id platform = platforms[0];

  // Device Query
  cl_device_id *devices    = NULL;
  cl_uint       numDevices = 0;

  clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, NULL, &numDevices);
  cl_device_id = (cl_device_id *)malloc(sizeof(cl_device_id) * numDevices);

  clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, numDevices, devices, NULL);

  cl_device_id device = devices[0];

  char deviceName[128];

  clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(deviceName)-1, deviceName, NULL);

  deviceName;
}

#!syntax cpp
#include <CL/OpenCL.hpp>
#include <vector>
#include <string>

int main()
{
    // Platform Query
    std::vector<cl::Platform> platforms;
    cl::Platform::get(&platforms);
    
    cl::Platform platform = platforms[0];

    // Device Query
    std::vector<cl::Device> devices;

    platform.getDevices(CL_DEVICE_TYPE_ALL, &devices);
    cl::Device device = devices[0];

    std::string deviceName = device.getInfo<CL_DEVICE_NAME>();
    deviceName;
  // ...
}

2.2. 커널

#!syntax cpp
__kernel void
arrAdd(global const int *src1, global const int *src2, global int *dest)
{
  const uint index = get_global_id(0);
  dest[index]      = src1[index] + src2[index];
  return;
}

2.2.1. 자료형

#!syntax cpp
(float4)(float, float, float, float); // 4개의 float을 1개의 float4 형으로
(float4)(float2, float, float); // 한개의 float2 벡터와 두개의 float 스칼라를 float 4 형으로
(float4)(float); // 한개의 float을 float4 로 브로드캐스팅

#!syntax cpp
float4 v = (float4)u + (float)f

#!latex
\begin{bmatrix} v_x \\ v_y \\ v_z \\ v_w \end{bmatrix} = \begin{bmatrix} u_x \\ u_y \\ u_z \\ u_w \end{bmatrix} + \begin{bmatrix} f \\ f \\ f \\ f \end{bmatrix}

#!latex
\begin{bmatrix} w.x \\ w.y \\ w.z \\ w.w \end{bmatrix} = \begin{bmatrix} v.x \\ v.y \\ v.z \\ v.w \end{bmatrix} + \begin{bmatrix} u.x \\ u.y \\ u.z \\ u.w \end{bmatrix} = \begin{bmatrix} v.x + u.x \\ v.y + u.y \\ v.z + u.z \\ v.w + u.w \end{bmatrix}

2.2.2. 벡터 접근자

#!syntax cpp
float4 myFloat4 = (float4)(0.1f, 0.2f, 0.3f, 0.4f);
myFloat4.x; // 0.1f
myFloat4.y; // 0.2f

myFloat4.xy; // (0.1f, 0.2f) -> float2;
myFloat4.zwz; // (0.3f, 0.4f, 0.3f) -> float3;

myFloat4.lo; // (0.1f, 0.2f) -> float2;
myFloat4.hi; // (0.3f, 0.4f) -> float2;

myFloat4.even; // myFloat4.xz - > float2;
myFloat4.odd; // myFloat4.yw -> float2;

myFloat4.0123; // myFloat4.xyzw -> float4;

2.2.3. 지정자 (qualifiers)

2.2.4. 연산자

2.2.5. 내장 함수

• 삼각함수와 초월함수
  acos, acosh, acospi,...
  
  내장 함수와 같은 역할을 하지만 하드웨어를 통한 근사값을 통해 속도를 높일 수 있는 함수들이 있다. 이 함수들은 native_ Prefix가 붙어 있으며 예를 들어 native_rsqrt(), rsqrt() 둘 다 역제곱근을 계산하지만, native의 경우 하드웨어에 내장된 유닛을 사용할 수도 있고, 룩업 테이블로 근사할 수도 있는 등 구현은 제조사 마음대로다. 대체로 계산 정밀도는 떨어지지만, 속도가 매우 빠르다.[2]
  Native 함수를 안 써도 unsafe math 옵션을 사용하면 컴파일러가 전부 바꿔버리는데 성능과 정밀도가 둘 다 중요하다면 선택적으로 적용해 내장 함수를 직접 사용하는 것이 권장된다.

• vload<>, vstore<>
  메모리 배열을 벡터 단위로 읽거나 쓴다. 예를 들어 메모리에 rgb 픽셀이 이어진 경우 vload3()을 통해 3개의 원소를 가진 값을 가져올 수 있다. 보통 최대의 효율을 위해서는 같은 메모리 접근이여도 스칼라 단위로 연산을 하는 것이 아니라 벡터 단위로 하는 것이 하드웨어의 효율을 최대한 이끌어 낼 수 있다. 스칼라로 작성해도 대충은 돌아 가긴 하고 컴파일러가 가능하면 알아서 잘 처리해 주긴 하지만 이는 매우 제한적이므로 명시적으로 벡터 단위로 접근하는것이 표준이다.

• printf
  C의 printf 맞다. 보통 커널에서 동작하는 코드는 호스트와 독립적으로 동작하기에 디버깅이 까다롭지만 printf를 통해 값을 출력할 수 있다. 당연하지만 많이 쓰면 printf로 인한 block이 생기므로 디버깅용으로만 사용하는 것이 좋다. OpenCL 1.2 부터 표준화가 되었는데, 그 이전까지는 벤더의 자체 확장 (cl_amd_printf, cl_intel_printf, cl_arm_printf 등)에 의존했다. 워낙 유용한 기능이라 너도나도 자체 확장을 만들어 쓰다 보니 표준으로 편입된 기능이 되었다.

• get_global_id, get_local_id
  GPGPU의 경우 데이터를 쪼개서 연산하게 되는데, 각 연산 코어가 위치한 순번을 조회할 때 사용한다. 예를 들어 32*32 크기의 데이터는 32개 유닛을 가진 프로세서에서 각 코어가 32개씩, 총 32코어가 1024개의 데이터를 병렬로 처리하는데, 코어가 위치한 그리드 도메인 내를 식별하기 위한 방법이다. 이게 없이는 메모리 접근이 거의 불가능하다고 보면 된다.

2.3. 메모리 모델