OpenCL Architecture

OpenCL Architecture

OpenCL 이란

- Open Computing Language

- 이종 플랫폼을 위한 표준 병렬 프로그래밍 모델

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

- CPU, GPU, FPGA 등

- 추상화된 플랫폼 모델 제공

- 저작권료를 지불하지 않아도 되는 공개 표준

OpenCL의 구성

- 병렬 프로그래밍을 위한 프레임워크

- OpenCL C

- 커널을 작성하기 위한 프로그래밍 언어

- 컴파일러

- 호스트 API

- 라이브러리

- 런타임 시스템

- 표준 문서의 구성

- 기본 표준

- OpenCL 구현이 지원해야할 필수 사항

- 임베디드 프로파일

- 모파일/임베디드 디바이스용

- 확장 표준

- 추후 표준에 추가될 수 있는 사항

OpenCL의 특징

- 기본 특징

- 데이터/태스크 병렬 프로그래밍 모델 지원

- ISO C99 기반의 언어 사용

- IEEE 754에 기반

- OpenGL 등과의 상호운용 가능

- 코드 이식성

- 한 번 작성한 코드는 다양한 디바이스에서 실행 가능

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

- 라이브러리 개발자, 미들웨어 벤더, 성능이 중요한 애플리케이션 개발자

- 애플리케이션 개발을 위해 성능과 관련된 많은 요소를 직접 제어

현재 OpenCL의 한계

- 성능 이식성의 부재

- 타깃 디바이스에 따라 프로그래머의 최적화가 필요

- 하지만 다른 프로그래밍 모델도 성능 이식성이 부재

- 프로그래밍의 복잡성

- 프로그래머가 많은 요소를 직접 통제해야 함

- 디바이스 별 최적화를 위해 많은 시간이 소요

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

- 이종 클러스터에서의 프로그래밍

- MPI + OpenCL 또는 MPI + CUDA

- 프로그래밍의 어려움 증가

- 코드 이식성 감소

- 유지 관리의 어려움

OpenCL Architecture

- 플랫폼 모델

- 실행 모델

- 메모리 모델

- 프로그래밍 모델

- OpenCL 프레임워크

플랫폼 모델(Platform Model)

- 호스트(host) + 하나 이상의 계산 디바이스

- 계산 디바이스(compute device)

- OpenCL 디바이스

- 하나 이상의 계산 유닛으로 구성

- 계산 유닛(compute unit, CU)

- 하나 이상의 Processing Element로 구성

- Processing Element (PE) (Core와 비슷)

- 디바이스 내에서 실제 계산을 수행

- 가상 스칼라 프로세서

OpenCL 애플리케이션

- 호스트 프로그램 + OpenCL 프로그램

- 호스트 프로그램(host program)

- 호스트 프로세서에서 실행

- OpenCL 프로그램의 실행을 조정

- 일반 프로그래밍 언어로 작성

- C, C++ 등

- 디바이스의 실행을 관리

- 커맨드 단위

- 메모리, 계산, 동기화 등

- OpenCL 프로그램

- OpenCL 커널의 집합

- OpenCL 디바이스에서 실행

- 디바이스 : CPU, GPU 등

- 커널(kernel)

- OpenCL C로 작성

- 디바이스에서 실행되는 코드의 기본 단위

- 많은 인스턴스가 디바이스에서 실행됨

- 데이터 병렬성 이용

- 호스트와 디바이스의 병렬 실행

- 호스트 프로그램과 커널은 병렬로 실행 가능

- 여러 디바이스를 동시에 사용 가능

OpenCL C

- 커널 작성을 위한 프로그래밍 언어

- 별도의 #include 사용 X

- 제약

- 표준 C99 헤더파일 지원 X

- 함수 포티인터, 재귀함수, 가변길이 배열(ex : int a[N]), 비트 필드 지원 X

- 확장

- 벡터 타입

- ex) char2, int4, float8

- 이미지 타입

- ex) image2d_t

- 주소 공간 한정자

- ex) __global, __constant, __local

- 동기화

- 많은 빌트인 함수

실행 모델(Execution Model)

- 호스트 프로그램이 커널 실행을 위해 "N-차원의 인덱스 공간"을 지정

- 인덱스 공간은 같은 크기의 하나 이상의 그룹으로 분할됨

- 그룹은 디바이스의 CU에서 실행됨

- 커널은 인덱스 공간의 각 지점에서 실행됨

- 각 지점은 하나 이상의 PE에서 실행됨

커널 인덱스 공간(Kernel Index Space)

- NDRange

- N 차원의 인덱스 공간(N : 1,2,3,)

- 전체 워크 아이템(work-item)의 수를 지정

- N 차원 튜플로 지정

- 워크 아이템은 커널 인스턴스

- 워크 아이템은 CUDA에서 스레드로 불림

- 각 워크 아이템은 인덱스 공간의 각 지점에서 실행

- 모든 워크 아이템은 병렬로 실행(모두 동시에 실행)

워크 그룹(Work-group)

- 워크 아이템의 집합

- 인덱스 공간을 분할하는 그룹

- CU에서 실행됨

- CPU의 경우

- 하나의 코어가 각 워크 아이템을 루프 형태로 실행

- GPU의 경우

- 하나의 SM이 여러 워크 그룹을 문맥 교환하며 실행

- 모양 및 크

- 모든 워크 그룹은 동일해야 함

- 성능에 큰 영향을 미침

- 알고리즘에 최적의 워크 그룹 차원과 크기를 찾는 것이 중요

- 선택 방법

- 프로그래머가 직접 코드에 지정

- 디바이스의 아키텍처와 커널 코드의 특성 고려

- OpenCL 런타임이 자동으로 선택

- 구현에 따라 다름

워크 그룹과 워크 아이템 식별

- 인덱스 공간 : (Gx, Gy)

- 워크 그룹

- 크기 : (Sx, Sy)

- 워크 그룹 ID : (wx, wy)

- 워크 그룹 개수

- (Wx, Wy) = (Gx/Sx, Gy/Sy)

- 워크 아이템

- 글로벌 ID : (gx, gy)

- 인덱스 공간 내의 전역 ID

- 글로벌 ID 오프셋 : (Fx, Fy)

- 로컬 ID : (sx, sy)

- 워크 그룹 내에서의 ID

- (gx, gy) = (wx * Sx + sx + Fx, wy * Sy + sy + Fy)

커널의 종류

- OpenCL 커널

- OpenCL C로 작성하고 OpenCL 컴파일러로 컴파일한 것

- OpenCL 표준에서 필수

- 네이티브 커널(Native kernel)

- 어플리케이션 코드나 라이브러리의 함수

- 호스트에서 함수 포인터를 이용해 접근

- OpenCl 표준에서 옵션

커널 실행

- 커널 코드가 커널 인텍스 공간의 모든 지점에서 실행됨

- SPMD

- 인덱스 공간

- 데이터가 워크 아이템에 어떻게 매핑 되는지 정의하는 역할

컨텍스트(Context)

- OpenCL 컨텍스트

- 커널이 실행되는 환경

- 동기화와 메모리 관리가 정의되는 도메인

- 호스트 프로그램이 API clCreateContext()를 이용하여 생성

- 포함하는 리소스

- 디바이스

- 호스트가 사용할 OpenCL 디바이스 목록

- 커널

- OpenCL 디바이스에서 실행할 OpenCL 함수들

- 프로그램 오브젝트

- 프로그램 소스 및 실행 오브젝트

- 메모리 오브젝트

- 호스트와 OpenCL 디바이스에서 접근 가능한 메모리 오브젝트 집합

- 커널은 계산 과정에서 메모리 오브젝트 사용

커맨트 큐(Command Queue)

- 호스트가 디바이스별로 생성

- 하나의 디바이스에서 여러 개의 커맨드 큐 연결 가능

- 호스트가 디바이스에서 실행될 커맨드를 enqueue

- 커맨드의 종류

- 커널 실행 커맨드

- 메모리 커맨드

- 데이터 이동, 호스트 메모리에 메모리 오브젝트의 map/unmap

- 동기화 커맨드

- 커맨드 실행 모드

- In-order Execution

- 삽입 순서대로

- 기본 실행 모드

- Out-of-order Execution

- 커맨드 큐 생성시 지정 가능

메모리 모델(Memory Model)

- 4개의 메모리 영역

- 글로벌 메모리

- 컨스탄트 메모리

- 로컬 메모리

- 프라이빗 메모리

- 글로벌 메모리(global memory)

- 모든 워크 그룹의 모든 워크 아이템이 접근 가능한 메모리

- 워크 아이템은 메모리의 임의의 위치에 읽기/쓰기 가능

- 글로벌 메모리 캐싱

- 디바이스의 캐시 메모리 유무에 따라

- 컨스탄트 메모리(constant memory)

- 커널 실행동안 변하지 않는 값(상수)으로 남아있는 글로벌 메모리 영역

- 호스트가 할당하고 초기화

- 로컬 메모리(local memory)

- 워크 그룹별로 각각 존재하는 메모리

- 한 워크 그룹 내의 워크 아이템들만 접근 가능

- 워크 그룹 내의 워크 아이템들간 공유를 위해 사용

- 프라이빗 메모리(private memory)

- 워크 아이템별로 갖는 메모리

- 한 워크 아이템의 프라이빗 메모리는 다른 워크 아이템이 접근할 수 없음

메모리 할당 및 접근

메모리 오브젝트

- 버퍼 오브젝트

- 일차원 데이터

- 스칼라 타입, 벡터 타입, 사용자 정의 타입

- 메모리에 순차적으로 데이터 저장

    - 커널에서 포인터로 접근 가능

    - 서브 버퍼

- 버퍼 오브젝트의 일부 영역으로부터 생성

- 이미지 오브젝트

- 2차원또는 3차원의 데이터

- 텍스처, 프레임 버퍼, 이미지

- 포인터로 직접 접근할 수 없음

메모리 컨시스턴시

- 릴렉스드 메모리 컨시스턴시 모델 사용

- 메모리의 내용이 서로 다른 워크 아이템에게 항상 같게 보일 필요는 없음

- 동기화 시 같게 보임

- 워크 아이템 하나가 사용하는 메모리

- 로컬 메모리

- 워크 그룹 배리어에서 워크 그룹 내의 모든 워크 아이템간에 컨시스턴시 보장

- 글로벌 메모리

- 워크 그룹 배리어에서 워크 그룹 내의 모든 워크 아이템간에 컨시스턴시 보장

- 서로 다른 워크 그룹 사이에는 메모리 컨시스턴시를 보장하지 않음

- 커맨드 간 공유되는 메모리

- 동기화 포인트에서 메모리 컨시스턴시 보장

- clFinish()

  - 워크 그룹 배리어

  - 이벤트 동기화

프로그래밍 모델

- 데이터병렬 프로그래밍 모델

- OpenCL의 기본 프로그래밍 모델

- 메모리 오브젝트의 데이터에 적용되는 명령어의 순서를 정의

- SPMD

- 인덱스 공간은 워크 아이템과 데이터의 매핑을 정의

- 워크 아이템이 어떤 데이터를 이용해 계산을 수행하는지 정의

- clEnqueueNDRangeKernel

- 태스크 병렬 프로그래밍 모델

- 인덱스 공간에 관계없이 하나의 커널 인스턴스만 실행

- 하나의 워크 아이템만 갖는 워크 그룹 한개를 디바이스에서 실행하는 것과 같은 의미

- 사용할 수 있는 병렬성

- 커널에서 벡터 타입 사용

- 여러 태스크를 동시에 실행

동기화

- 워크 아이템간의 동기화

- 하나의 워크 그룹 내에서만 가능

- 워크 그룹 배리어 이용

- 워크 그룹 간의 동기화

- 지원하지 않음

- 커널의 시작과 종료로는 가능

- 커맨드 간의 동기화

- 하나의 컨텍스트 내에서 커맨드 큐 내의 커맨드 사이에서 가능

- 하나 이상의 커맨드 큐

- 커맨드 큐 배리어 또는 이벤트 동기화 이용

워크 그룹 배리어

- 커널에서 barrier()이용

- barrier는 빌트인 함수

- 워크 그룹 내의 모든 워크 아이템은 배리어를 동시에 통과

- 모든 워크 아이템이 배리어까지 코드를 실행 후 그 다음 코드 실행

- 배리어 이후 워크 아이템 간의 메모리 일관성 보장

- 모든 워크 아이템이 barrer()를 실행하거나 전체가 실행하지 않아야 함

- 일부의 워크 아이템만 barrier()를 실행하는 것은 잘못된 커널

커맨드 큐 배리어

- 하나의 커맨드 큐에서 커맨드 사이의 동기화 방법

- 호스트가 커맨드 큐에 커맨드 큐 배리어를 enqueue

- 배리어 이전의 모든 커맨드의 실행이 완료됨을 보장

- 배리어 이전의 메모리 오브젝트의 갱신이 완료됨을 보장

- 배리어 이후 커맨드는 최신의 메모리 상태를 볼 수 있음

이벤트 동기화

- 하나 이상의 커맨드 큐에서 커맨드 간의 동기화 방법

- 커맨드를 enqueue하는 모든 API함수는 이벤트를 리턴

- 커맨드와 메모리 오브젝트 등을 식별

- 이벤트와 관련된 커맨드가 종료되기를 기다림

- 메모리 오브젝트의 갱신이 완료됨을 보장

opencl 프레임 워크

- opencl 어플리케이션 개발과 실행을 지원하는 소프트에어 시스템

- 호스트와 하나 이상의 opencl 디바이스를 단일 시스템처럼 사용

- 구성

- opencl 플랫폼 레이어

- opencl 디바이스 및 디바이스 정보 제공

- opencl 런타임

- 컨텍스트 조작 및 opencl 프로그램 실행

- opencl 컴파일

- opencl C로 작성된 커널 코드 컴파일