OpenCL [1] concept

OpenCL [1] concept

Heterogeneous Platform을 위한 application 개발

- 디바이스(device) 별로 다른 프로그래밍 방법이 존재합니다. (CPU, GPU, FPGA, …)

- 개발 및 유지 보수의 어려움이 있습니다.

OpenCL

Open Computing Language

Heterogeneous Platform을 위한 표준 병렬 프로그래밍 모델입니다.

- 다양한 아키텍처의 디바이스를 위해 프로그래밍이 가능합니다.

- 추상화된 플랫폼 모델을 제공합니다.

Open source로 사용이 되며, 다양한 회사들과 연구 기관에서 표준 개발에 참여하고 있습니다.

- data/task parallel programming 모델 지원 (실제로는 데이터 병렬 프로그래밍 모델을 많이 사용)

- ANSI/ISO C99 기반의 언어와 C++ 지원이 추가

- 코드 이식성(source-code portability) : 다양한 장치에서 실행 가능

- 이식성 있고 효율적인 코드를 작성하기 원하는 프로그래머가 대상

- 많은 벤더(vendor)가 지원 : Altera, AMD, IBM, Intel, NVIDIA, Samsung 등

OpenCL Limit

성능 이식성(performance portability)의 부재

- target 장치에 따라 프로그래머의 최적화가 필요합니다.

- 프로그래밍의 복잡성 - 많은 요소를 직접 통제해야 하며, 최적화를 위해 많은 시간이 소요됩니다.

- 단일 OS 내의 디바이스들만 프로그래밍 가능합니다.

OpenCL VS CUDA

CUDA(Compute Unified Device Architecture)

- https://developer.nivida.com/cuda-zone

- NVIDIA에서 만든 GPU 병렬 프로그래밍 모델

OpenCL과 1:1로 대응되는 점이 많습니다.

하나만 배우면 쉽게 익힐 수 있습니다.

OpenCL standard

- OpenCL 아키텍처(architecture) : 플랫폼(platform) 모델, 실행 모델, 메모리 모델

- OpenCL API 함수

- OpenCL C 언어 명세

OpenCL Framework

OpenCL application을 개발하고 실행하기 위한 소프트웨어 시스템

Platform Layer + Runtime + Compiler

- Platform Layer : host와 1 이상의 device로 이루어진 (가상의) 플랫폼 제공

- Runtime : Host program에서 OpenCL API 함수를 사용할 수 있도록합니다. 각종 객체를 생성, 관리하고, command queue에 들어온 command를 처리합니다.

- OpenCL C Compiler : OpenCL 프로그램 소스 코드로부터 device에서 실행될 바이너리(binary)를 생성합니다. host 프로그램에서 프로그램 객체를 만들고 빌드(build)할 때 불리어 집니다(API 함수 사용)

Platform

OpenCL Framework 마다 하나 이상의 플랫폼(platform) 제공합니다.

- 시스템에 여러 벤더(vendor)의 OpenCL 프레임워크가 설치되어 있으면 platform도 여러 개입니다.

- 플랫폼(platform) = 호스트(host) + 하나 이상의 디바이스(device)

* host : CPU에 해당

* device : Multi-core CPU, GPU, FPGA, …

Device

global 메모리(memory)와  constant 메모리 있습니다.

디바이스(device)마다 하나씩 있어서, 해당 device 안에서 접근 가능합니다.

device는 constant memory를 읽을 수 만 있습니다.

device에서는 호스트(host) 메인 메모리에 바로 접근할 수 없습니다.

반면 host에서는 global와 constant memory에 읽기/쓰기가 가능합니다. 이에 따라 main memory와 global/constant memory 사이의 데이터 복사를 담당합니다.

Device는 하나 이상의 Compute Unit(CU)로 구성됩니다.

CU 마다 로컬 메모리(local memory)가 하나씩 있습니다.

- 해당 CU 안에서만 접근이 가능합니다. global/constant memory는 모든 CU가 공유합니다.

- 호스트에서는 접근할 수 없습니다.

각 CU는 하나 이상의 Processing Element(PE)로 구성됩니다.

PE에서는 실제로 코드를 실행하면서 계산을 수행합니다(SIMD, SPMD 방식)

PE마다 private memory가 하나씩 있습니다.

해당 PE에서만 접근이 가능합니다. local memory는 한 CU 안의 모든 PE가 공유합니다.

OpenCL Application

Host program + OpenCL program

- Host program은 호스트 프로세서에 실행됩니다.

- OpenCL program은 여러 커널(kernel)로 이루어져 있으며 device에서 실행됩니다.

Host program

- 일반적인 프로그래밍 언어(C, C++, Fortran, Java, Python, …)로 작성됩니다.

- OpenCL API 함수를 사용해 device를 관리합니다.

- kernel 단위, 계산(kernel 실행), 데이터 전송, 동기화(synchronization)

Kernel

- OpenCL C 언어로 작성

- device에서 실행되는 코드의 기본 단위

- 많은 instance가 동시에 device에서 실행됩니다. 이에 따라 data parallelism 이용

Host program과 Kernel은 병렬로 실행됩니다. 여러 device를 동시에 사용하는 것도 가능합니다.

OpenCL C

Kernel 작성을 위한 프로그래밍 언어

제약 조건으로는 기존의 표준 헤더 파일(stdio.h, stdlib.h) 지원이 안됩니다. 함수 포인터, 재귀 함수 지원 안됩니다.

확장으로는 다음과 같습니다.

- 벡터(vector) 타입 :  char2, int4, float8, double16

- 이미지(image) 타입 : image2d_t

- 주소 공간 한정자(qualifier) :  __global, __constant, __local, __private

- 동기화(synchronization)

- 빌트인 함수 : sin(), cos(), printf(), atomic_add(), get_global_id()

Kernel Execution

Host program에서 kernel을 실행할 때 N-차원(dimension) 인덱스(index) 공간 지정합니다.

N = 1,2,3 중 하나이며, 이는 1차원 벡터 또는 흑백 이미지, 색상이 들어간 이미지 처럼 각 dimension에 맞게 지정하도록 선택할 수 있습니다.

work-item

index 공간의 각 제일 작은 사각형마다 kernel instance가 하나씩 실행됩니다.

커널 함수 안에서 OpenCL C의 빌트인 함수(get_global_id 등)을 사용해 index 공간에서의 work-item id를 얻어 올 수 있습니다.

work-group

여러 work-item 들이 work-group 단위로 묶여 있습니다.

모든 work-group의 크기는 동일해야 하기 때문에 같은 사이즈(전체 크기의 약수)로 쪼개야 합니다.

하나의 work-group은 하나의 CU에서 실행됩니다. 같은 CU에 있는 PE는 local memory를 공유합니다.

work-group 안의 여러 work-item들이 CU 안의 여러 PE에서 나누어 실행됩니다.

work-group 크기가 성능에 큰 영향을 미칩니다.

OpenCL kernel 함수

__kernel void vec_add(__global float *A, __global float *B, __global float *C) { int i = get_global_id(0); C[i] = A[i] + B[i]; }

Data Parallel programming model

work-item이 하나의 데이터를 처리합니다.

- kernel 코드는 데이터에 적용되는 명령어의 순서를 정의합니다.

- index 공간은 work-item과 데이터의 맵핑을 결정합니다.

- 각 work-item은 어떤 데이터를 이용해 계산을 수행하는지 확인합니다.

Kernel code Execution object

Context

- kernel이 실행되는 환경

- 다른 객체(object)들을 관리하기 위한 최상위 객체

- context 단위로 명령어 간 동기화 및 메모리 관리를 수행

- 자원 할당(초기화) 및 해제

Command-queue

- device에 커맨드(커널 실행, 데이터 전송, 동기화)를 보내기 위해 필요

- Host program이 command-queue에 command를 넣으면, OpenCL 프레임워크의 런타임 시스템에서 빼내어 실행

- device가 4개면 command-queue도 4개가 있어야 합니다.

Program Object

OpenCL 프로그램의 바이너리(binary)

- IR(Intermediate Representation)일 수도 있고, 실제 실행 가능한 코드일 수 도 있습니다.

바이너리를 미리 만들어 두는 것이 아니라, 실행 중에 컴파일을 하는 경우가 많습니다. 아직까지 device 간 binary 호환성이 지원되지 않기 때문입니다. SPIR 표준으로 시도 중이나 디바이스에 호환이 아직 많이 되고 있지 않습니다.

Memory Object

device의 global/constant memory에 공간을 할당하고 데이터를 저장하기 위해 필요합니다.

host program에서 memory object에 데이터를 읽고 쓸 수 있습니다.

-kernel 함수에서 memory object를 인자로 받을 수 있습니다.

- buffer object: 일반적인 배열(array)과 동일

- image object: 1~3 차원의 데이터를 처리하기 위한 특수 object, 텍스처(texture), 프레임 버퍼(frame buffer), 이미지

OpenCL Runtime

Host 프로그램이 command-queue에 kernel execution 명령을 넣습니다. kernel 인덱스 공간을 지정합니다.

OpenCL 런타임이 queue에 든 kernel 커맨드를 꺼내어 target device에 이슈(issue) 됩니다.

Runtime이 work-group들을 target device 내의 CU(Compute Unit)로 분배합니다.

OpenCL Framework

AMD OpenCL :  http://www.amd.com/en-us/solutions/professional/hpc/opencl

Intel OpenCL : https://software.intel.com/en-us/intel-opencl

NVIDIA OpenCL : https://developer.nvidia.com/opencl