ಣ ᴄᴏᴍᴘᴜᴛᴇʀ ᴄɪᴛʏ ഒ

Search

SearchSearch

SHPC 13 - OpenCL

Dec 31, 202419 min read

OpenCL Application

  • Host processor와 OpenCL compute device 에서 동작하는 프로그램의 집합을 일컫는 말
    • Compute device: CPU, GPU, Xeon Phi 등
  • Host program + OpenCL programs
    • OpenCL program: kernel의 집합

OpenCL C Language

  • 커널작성에 사용됨
  • ISO C99에 기반을 두고 있음
    • 단 표준 C99 헤더, 함수 포인터, 재귀, 다변길이 array, bit field 등은 허용하지 않는다
    • Compute device의 하드웨어 제약 때문이다
      • e.g. GPU는 재귀가 불가능 (스택이 없기 때문)
  • Extension
    • Vector type
    • Image type
    • Synchronization
    • Address space qualifiers
    • Built in functions

OpenCL Application

  • Host program
    • 호스트에서 실행되며 디바이스들의 커널 실행을 관리한다
  • 커널
    • 컴퓨팅 디바이스에 존재하는 실행가능한 코드
    • 실행되면 많은 instance들이 생성된다
      • Data parallelism을 달성하기 위함
  • 호스트 프로그램과 커널은 병렬로 실행된다

OpenCL Framework

  • OpenCL application의 개발과 실행을 위한 소프트웨어
  • OpenCL platform layer
    • 플랫폼과 디바이스를 관리하는 역할 수행
      • 플랫폼: 벤더가 제공하는 실제 실행 환경
      • OpenCL API, 벤더별 드라이버, 디바이스 등으로 구성됨
    • OpenCL 플랫폼 및 디바이스 목록 쿼리
    • Context 생성 시 사용할 디바이스와 플랫폼을 초기화
  • OpenCL runtime
    • 커널 실행, 메모리 관리, 명령 대기열 처리 등을 담당
    • 커맨트 큐의 생성 및 관리, 메모리 관리, 커널 실행, 동기화 등을 담당함
  • OpenCL compiler (kernel compilation)
    • 커널 코드를 디바이스에서 실행 가능한 바이너리로 변환
    • 프로그램 빌드 시 에러 및 로그 정보 제공

OpenCL Platform Model

Overall Architecture

  • Host processor: Host 프로그램을 동작시키는 프로세서.
    • 메인 메모리, Constant memory, Global memory 등 모든 메모리에 Read/Write 접근 가능
    • 하나의 호스트에 여러 compute device가 붙을 수 있다
  • Main memory: Host processor용. Compute device는 접근할 수 없음
  • Constant memory: Compute device에게 read-only로 제공되는 메모리
  • Compute device 내에도 캐시가 존재하는데, Global/Constant memory와 latency 차이가 크게 나지 않음
    • Bandwidth를 exploit하기 위해 사용
    • 로컬 메모리가 L1 캐시라고 했을 때, L2 캐시 같은 존재
    • 반면 L1캐시에서 데이터를 읽는 것은 매우 빠름
      • L1-L2는 latency, L2-L3는 밴드위쓰를 취하기 위함이다

Compute Device

  • GPU 하나라고 볼 수 있음
  • Compute unit(CU) 여러개의 집합이다
    • GPU의 CU는 SM

Compute Unit

  • GPU의 SM
  • Processing Element(PE) 의 집합
    • SM의 ALU라고 볼 수 있음
  • PE들이 동시에 코드를 돌리면서 SIMD, SPMD 달성
    • PE 하나가 하나의 스레드를 실행하는 식
  • 각 PE는 자신의 private memory를 가진다
    • SM에서의 context (Register)
    • 다른 PE에서 접근할 수 없음
  • Local memory는 PE 사이에서 공유되는 메모리
    • SM에서의 shared context
    • 다른 CU에서는 접근할 수 없음

Kernel Index Space

  • N-dimensional index space(N=1, 2 or 3)
    • 3차원까지 만들어질 수 있음
  • 동시에 실행 가능한 Work-item의 수를 결정
    • Work-item == 커널 인스턴스 (CUDA에서는 쓰레드라 부름)
    • SPMD
  • 이렇게 work-item의 cartesian coordinate가 존재할 때, 이를 어떻게 compute unit에게 분배할 것인지가 중요
    • SM이 여러개 있을 때, 어느 SM에게 분배할 것인지
    • 아래와 같이 work-group이라는 그룹을 만들어 분배

  • 하나의 work group이 하나의 compute unit에서 실행
    • GPU로 치면 SM 하나가 하나의 work group을 할당받는 것
    • 여러 쓰레드를 돌리면서 work group을 실행
  • 크기는 최적화하여 프로그래머가 선택
  • Work group 내에서 work item의 index를 가지게 된다.
    • 오타가 있는데 work group (1, 1) 이 아니라 (1, 2)이고, work-item (2, 1) 도 (1, 2) 이다. work-item도 (10, 12)이다.

Kernel Example

  • 각각 다른 data element에 적용해서 for loop를 벗겨내고 kernel code를 만들 수 있다.
  • __global 은 글로벌 영역에 저장된 정보임을 나타낸다
  • __kernel은 커널 function임을 나타낸다
  • get_global_id : 전체 work item에서의 index
  • 그 결과 아래 그림과 같은 동작을 수행하는 코드가 만들어진다
    • SPMD
    • index space를 이용해 어떤 데이터에 매핑되었는지 확인하는 방식

Vector Addition Example

1. Kernel source code 준비

  • 커널에서 동작시킬 소스코드를 정의

2. Initializing Vectors

  • A, B, C 데이터를 호스트에 할당

3. Obtaining OpenCL Platforms and Devices

  • 플랫폼 하나를 받아와서 해당 플랫폼에서 사용 가능한 GPU의 device id를 가져온다
  • cl_platform_id: OpenCL platform을 선언
  • cl_int clGetPlatformIDs(cl_uint num_entries, cl_platform_id *platforms, cl_uint *num_platforms)
    • OpenCL에서 사용 가능한 플랫폼의 목록을 가져오는 함수
    • num_entries: platforms 배열에 저장할 수 있는 플랫폼의 갯수를 전달
    • platforms: cl_platform_id 배열을 리턴 (NULL인 경우 num_platform에 사용가능한 플랫폼 수만 저장됨)
    • num_platforms: 사용 가능한 전체 플랫폼의 갯수를 리턴
    • 성공 시 CL_SUCCESS 를 반환하고, 실패 시 오류 코드를 반환한다
  • cl_int clGetDeviceIDs(cl_platform_id platform, cl_device_type device_type, cl_uint num_entries, cl_device_id *devices, cl_uint *num_devices);
    • 특정 플랫폼의 디바이스 목록을 가져오는 함수
    • platform: 디바이스를 조회할 OpenCL 플랫폼을 가리키는 cl_platform_id
    • device_type: 조회하려는 디바이스의 유형을 전달
      • CL_DEVICE_TYPE_CPU, CL_DEVICE_TYPE_GPU, CL_DEVICE_TYPE_ACCELERATOR, CL_DEVICE_TYPE_ALL 등등
    • num_entries: devices 배열에 저장할 수 있는 디바이스의 최대 갯수
      • 디바이스의 개수만 알고 싶을 때는 0으로 설정하고, device를 NULL로 설정한다
    • devices: cl_device_id 배열을 리턴
    • num_devices: 사용할 수 있는 디바이스의 총 개수를 저장할 변수 포인터
    • 성공 시 CL_SUCCESS 를 반환하고, 실패 시 오류 코드를 반환한다

4. Creating OpenCL Context

  • OpenCL Context
    • Kernel이 실행되는 환경
    • OpenCL 리소스 관리, 디바이스와 호스트 간의 상호작용을 조정하는 역할
    • Synchronization, Memory management 등이 정의되는 도메인
  • cl_context clCreateContext(const cl_context_properties *properties, cl_uint num_devices, const cl_device_id *devices, void (*pfn_notify)(const char *errinfo, const void *private_info, size_t cb, void *user_data), void *user_data, cl_int *errcode_ret);
    • properties: 컨텍스트의 속성을 지정 (보통 NULL)
    • num_devices: 컨택스트 생성 시 사용할 디바이스의 갯수
    • devices: 컨택스트에서 사용할 디바이스의 목록
    • pfn_notify: 오류 발생시 콜백
    • user_data: 콜백에 전달할 데이터
    • errcode_ret: 함수 호출 결과를 저장할 변수 포인터 (CL_SUCCESS)

5. Creating Command-queues

  • Command queue
    • 특정 compute device에서 실행될 command를 담고 있음
      • 하나의 command queue가 Host program에 의해 하나의 compute device에 할당됨
      • in-order 혹은 out-of-order로 issue됨
  • Commands
    • Command queue에 제출되는 OpenCL operation
      • Kernel execution command
      • Memory command
      • Synchronization command
  • cl_command_queue clCreateCommandQueue(cl_context context, cl_device_id device, cl_command_queue_properties properties, cl_int *errcode_ret);
    • context: 명령 큐를 생성할 OpenCL 컨택스트
      • 컨택스트가 명령 큐가 사용 가능한 리소스의 유효 범위를 설정하기 때문에 필요
    • device: 명령 큐가 생성될 대상 디바이스 (Context에 포함된 디바이스 중 하나)
    • properties: 명령 큐의 특성을 지정
      • CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, …
    • errcode_ret: 함수 호출 결과를 저장할 변수 포인터 (CL_SUCCESS)

6. Allocating Memory Objects

  • OpenCL memory object를 할당
    • Buffer objects
      • 1차원 데이터 (array): Scalar type, vector type, user defined type
      • Accessed by a pointer in the kernel
    • Image objects
      • Two- or three-dimensional data
      • Cannot be accessed by a pointer
  • cl_mem clCreateBuffer(cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_int *errcode_ret);
    • OpenCL에서 사용할 메모리 버퍼 객체를 생성하는 함수
    • context: 버퍼를 생성할 OpenCL 컨텍스트
      • 컨텍스트 내부의 디바이스들이 해당 버퍼에 접근 가능
    • flags: 버퍼의 메모리 접근 속성을 지정하는 플래그
      • CL_MEM_READ_WRITE/CL_MEM_READ_ONLY …
    • size: 생성할 버퍼의 크기
    • host_ptr: 호스트 메모리 포인터
      • NULL인 경우 새로 할당
    • errcode_ret: 함수 호출의 결과 코드가 저장될 변수 포인터 (CL_SUCCESS)
    • OpenCL 메모리 객체 반환

7. Compiling and Building the OpenCL Program

  • cl_program clCreateProgramWithSource(cl_context context, cl_uint count, const char **strings, const size_t *lengths, cl_int *errcode_ret);
    • context: 프로그램을 생성할 OpenCL 컨텍스트
    • count: 소스코드 문자열의 갯수 (위 예시에서는 하나의 문자열이니 1)
    • strings: 소스코드 문자열의 배열 (const char * 로 저장)
    • lengths: 각 문자열의 길이를 나타내는 배열
    • errcode_ret: 함수 호출 결과 코드가 저장될 포인터 (CL_SUCCESS)
    • OpenCL 프로그램 객체 cl_program 반환
  • cl_int clBuildProgram(cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (*pfn_notify)(cl_program program, void *user_data), void *user_data);
    • program: clCreateProgramWithSource로 생성된 cl_program, 빌드 대상
    • num_devices: 빌드 대상 디바이스의 수
    • device_list: 빌드할 디바이스의 아이디 배열
    • options: 컴파일 옵션
    • pfn_notify: 비동기 빌드 완료 알림을 위한 콜백 함수
    • user_data: 콜백 함수에 전달할 사용자 데이터

7. Creating Kernel Objects

  • cl_kernel clCreateKernel(cl_program program, const char *kernel_name, cl_int *errcode_ret);
    • Kernel Object 생성
      • Device에서 실행되며 프로그램 실행 중 메모리 접근 및 컨텍스트를 관리
      • Device가 이를 기반으로 Compute Unit(SM)에게 워크 그룹을 분배
      • Compute Unit은 가지고 있는 PE를 이용해 해당 워크로드를 각자 실행하게 됨
  • cl_int clEnqueueWriteBuffer(cl_command_queue queue, cl_mem buffer, cl_bool blocking_write, size_t offset, size_t size, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
    • 커맨드 큐에 버퍼 쓰기 명령을 enqueue
    • 호스트 메모리(hostA)를 디바이스 메모리(bufferA)로 복사
      • Global/Constant 여부는 버퍼를 만들 때 명시
        • Global: cl_mem buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, NULL, &err);
        • Constant: cl_mem const_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY, buffer_size, NULL, &err);
      • 버퍼가 clCreateBuffer로 만들어지면 이는 OpenCL에 의해 관리되며 디바이스 메모리와 연동되기에 프로그래머 입장에서는 사실상 디바이스 메모리에 있다고 봐도 무방함

8. Launching the Kernel

  • cl_int clEnqueueNDRangeKernel(cl_command_queue queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
    • 커맨드 큐에 커널 객체의 실행을 enqueue
    • global: 전체 work item의 갯수
    • local: 하나의 work group에 속하는 work item의 갯수

9. Obtaining the Result from the Device

  • clFinish(command_queue)
    • command_queue의 모든 명령이 완료되기 전에 리턴하지 않음
    • Out-of-order인 경우 큐에 다른 명령어가 끝나지 않고 남아있을 수 있기 때문에 기다려야 한다
  • clEnqueueReadBuffer(command_queue, bufferC, CL_TRUE, 0, sizeC, hostC, 0, NULL, NULL);
    • blocking_read를 CL_TRUE로 설정하여 데이터 복사가 완료된 후에 리턴하게 함

What the Host does for Kernel Execution

  1. OepnCL 커널 코드 작성 및 컴파일
  2. 디바이스 메모리에 데이터 복사
  3. 커널 인자 입력
  4. 글로벌 및 로컬 워크 그룹 크기 정보 입력
  5. 커널에 명령어 enqueue
  6. (호스트는 다른 일을 하는 중) OpenCL runtime이 디바이스에 명령어 issue, 디바이스는 실행
  7. 디바이스 메모리로부터 결과값을 읽어옴

OpenCL Runtime

  • Host program이 커널 command를 command queue에 전달
  • 런타임은 타겟 디바이스에게 command를 issue
    • in-order 혹은 out-of-order
    • 명령어 간 의존성을 확인하고, 의존성이 해결된 명령, 즉 실행 가능한 명령을 전달
  • 런타임은 커널 워크로드를 워크 그룹 단위로 Computing unit에게 분배

(정리) Runtime의 역할

  • Host program이 command queue에 issue한 커널 명령어를 타겟 디바이스에게 Issue한다
  • 타겟 디바이스의 Computing Unit(SM)에게 워크 그룹 단위로 워크로드를 분배한다

Synchronization

Work-group Sync: Barrier

  • Work group 내에서의 synchronization 방법
    • 여러개의 work group 간의 barrier는 없음
    • 쿠다도 마찬가지
  • • void barrier(cl_mem_fence_flags flags)
    • 아래와 같이 사용
    • CLK_LOCAL_MEM_FENCE: 로컬 메모리 동기화에 사용
    • CLK_GLOBAL_MEM_FENCE: 글로벌 메모리 동기화에 사용

Command-queue Synchronization: Barrier, Event

  • Command queue 내에서 명령어 간 동기화
    • Out-of-order 방식의 command queue운영에서 특히 중요

Barrier

  • Barrier에서 이전 명령어가 완료되기까지 기다린 후, 다음 명령어를 실행하는 방식
    • 위 그림에서 Kernel A와 Kernel B는 Out of order로 병렬 실행되나, kernel C는 kernel A와 kernel B가 끝난 다음에야 수행됨

Event Synchronization

  • 모든 OpenCL 명령어는 실행 결과를 추적하기위해 Event 객체를 반환함
  • 이 Event를 이용해 명령어 간 동기화를 수행할 수 있음
    • 서로 다른 command queue에서도 수행 가능
  • 특정한 Event 객체가 주어졌을 때(특정 명령어가 실행되었을 떄) 실행 가능한 것

Memory Consistency Model

  • OpenCL은 Relaxed memory consistency model 을 사용
  • 즉, 동기화 지점에서만 일관성이 보장됨
  • Local/Global Memory의 경우
    • 같은 워크 그룹 내에서, 워크 그룹 배리어에서만 동기화가 이뤄짐
    • 다른 워크 그룹간의 동기화는 없음
  • Barrier 이후에서 메모리 업데이트가 워크 그룹 내의 워크 그룹 아이템들에게 보이게 됨

동기화 방법 정리

  • clFinish
  • Work-group barrier
  • Command queue barrier
  • Event sync

Vector Types

  • 뭐 이런 것들이 있는데, 솔직히 별로 중요하지는 않은 듯 하다
  • 이유는 GPU의 Processing Unit(ALU) 단위에서는 CPU와 같은 벡터 연산 하드웨어가 없기 때문!
  • 따라서 GPU에서 vector instruction을 쓰면, 사실상 PE 하나에서 루프 돌리듯이 실행되기 때문에 성능이 좋아지지 않는다
  • (쓰면 망한다! 라고 말씀하셨음)
  • 대신 Local memory로 데이터를 복사할 때 vector type을 쓰는 것은 조금 더 빠를 수 있다
    • 한 번의 읽기/쓰기 명령으로 여러개의 element를 로드할 수 있기 때문임

Limitations

  • 현재 OpenCL은 하나의 OS 인스턴스에서 여러 디바이스를 연산하는 것에 목적을 둠
    • 하나의 컴퓨터 내에서 CPU, GPU를 병렬로 활용하는 것
  • 따라서 서로 다른 이기종 클러스터에서는 사용하기 어려움
    • 서로 다른 노드가 각자의 운영체제를 실행하게 되고, 노드 간 작업 분배 및 통신을 위해서는 MPI 같은 또 다른 라이브러리를 사용해야 함
      • OpenCL이 각 노드의 로컬 디바이스만 관리하기 때문
    • 따라서 보통 MPI + OpenCL, MPI + CUDA 이런식으로 조합을 해서 사용
    • 이렇게 조합하면 이식성, 유지보수성이 떨어지고 코드가 복잡해짐

OpenCL ICD

  • 하나의 운영체제에서 여러 OpenCL 구현을 공존하게 지원
    • 인텔용, AMD용, NVIDIA용 등 서로 다른 벤더의 OpenCL 드라이버를 하나의 시스템(OS)에서 사용 가능
    • 애플리케이션이 실행될 때 어떤 플랫폼을 사용할지 지정

Limitations

  • 사용자가 항상 명시적으로 어떤 프레임워크를 쓸 것인지 지정해야 함
  • buffer, event 같은 메모리 오브젝트를 공유할 수 없고, 카피를 해야 한다는 문제점이 있음

Graph View

Backlinks

  • No backlinks found