[Nvidia] OpenCL 기초 문법 및 병렬처리 관련 정리

[Nvidia] OpenCL 기초 문법 및 병렬처리 관련 정리

2022. 12. 31. 12:51ㆍDevelopers 공간 [Basic]/Embedded

728x90

딥러닝 등 다양한 Computationally Expensive한 어플리케이션의 등장으로 병렬처리에 대한 요구가 많이 늘어났습니다.

우리는 자동화된 병렬화를 지원하는 다양한 프레임워크 및 라이브러리(deepspeed, pytorch DP & DDP, AWS SMDDP horovod, tensorRT 등)을 통해 병렬처리를 하고 있지만, 근본적으로 병렬화에 필요한 특징이나 어떤 것을 유의 깊게 보아야 할지에 대한 이해를 바탕으로 사용한다면 더욱 최적화된 환경에서 작업이 가능할 것 같기도 합니다.

또한 때로는 Customized된 병렬처리를 위해 직접 내가 만든 로직을 병렬처리를 통해 가속화하고 싶은 경우도 있습니다.

보통 이럴 때 Nvidia에서 제공하는 Toolkit인 CUDA를 활용하는 경우가 대부분이지만, 먼저 OpenCL을 활용해 병렬 처리에 대해 이해하고 나서 CUDA를 살펴보려고 합니다.

<구성>
1. 병렬처리 기초 지식
a. 기본용어
b. 병렬 처리 용어
2. OpenCL 함수
a. Host 함수
b. Device 함수
3. 병렬처리 최적화 방법
a. 커널 성능 저하 요인
b. 최적화 방법 개요
c. Case1. Host ↔ Global Mem ↔ Kernel
d. Case2. Host ↔ Global Mem ↔ Local Mem ↔ Kernel

글효과 분류1 : 코드

글효과 분류2 : 폴더/파일

글효과 분류3 : 용어설명

글효과 분류4 : 글 내 참조

1. 병렬처리 기초 지식

먼저 병렬처리를 이해하기 위해 기본적인 용어들에 대해 살펴보려고 합니다.

하기의 내용을 바탕으로 "1-b. 병렬처리 용어"를 이해하고 작업하는 것이 시스템을 이해하는데 큰 도움이 되며, "1-a. 기본용어" 같은 경우는 완벽하게 이해하지 않아도 상관 없습니다.

a. 기본용어

병렬성
- Task 병렬성 : Thread별로 다른 task를 실행하는 것을 의미하며, 작은 수만 가능합니다.
- Data 병렬성(Loop-level 병렬성) : Thread별로 같은 task를 실행시 여러가지 data를 병렬처리하는 것을 의미합니다.
  - ex) SIMD(Single Instruction Multiple Data), SPMD(Single Program Multiple Data)

Process 와 Thread
- Process : 프로그램의 실행단위이며, 여러개의 thread로 구성됩니다.
  - User Process와 Kernel Process 가 있습니다.
- Thread : CPU 코어에 스케줄 될 수 있는 최소 단위로, thread간에는 자원이 공유가 가능합니다.

Data Dependency : 순서와 상관없이 병렬처리시 항상 같은 결과가 나오기 위한 의존성의 종류입니다.
- True Dependence : 없앨 수 없는 dependence
  - ex) Flow Dependence : Read After Write(RAW)
    R2 <- R1 + R3
    R4 <- R2 + R3
- False Dependence : 없앨 수 있는 dependence (변수 Renaming 으로 해결가능합니다)
  - ex) Anti Dependence : Write After Read (WAR)
    문제) R2 <- R1 + R3
    R3 <- R4 + R5
    해결) R2 <- R1 + Ra
    Rb <- R4 + R5
  - ex) Output Dependence : Write After Write (WAW)
    문제) R2 <- R1 + R3
    R2 <- R4 + R5
    해결) Ra <- R1 + R3
    Rb <- R4 + R5
- No Dependense : 실제 Dependense 는 아닙니다.
  - ex) Input Dependence : Read After Read(RAR)
- 추가적인 Dependence
  - ex) Loop-Independent Dependence : Loop와 상관없이 발생하는 dependence이며, Loop가 없어져도 존재
    <RAW예시>
    for(i=0;i<N;i++){
    A[i] = B[i];
    F[i+1] = A[i];
    }
  - ex) Loop-Carried Dependence : Loop를 통해 발생하는 dependence이며, Loop가 사라지면 사라집니다.
    <예시>
    for(i=0;i<N;i++){
    F[i+1] = A[i];
    A[i+1] = F[i];
    }

------------------------------------------------------------------------------------------------------
<참고> Pipeline Hazzard : 파이프라인에서 stall 이 발생/필요하는 경우
** stall : 단순히는 "기다리는 것"으로, Hazzrd 발생시 가장 기본적으로 해결할 수 있는 방법입니다.

Structural Dependency : 같은 시간에 같은 메모리/레지스터 등에 접근하는 경우 (구조적인 문제로)
- 해결책 : 메모리를 여러개로 나누어 해결합니다.
Control Dependency : Branch Hazzard 라고도 부르며, jump나 branch 때문에 쓸모 없는 instruction이 실행되는 경우
** jump : 새로운 명령어가 있는 메모리로 jump 이동해 실행하는 것을 의미합니다.
** branch(분기) : 단순히는 if 문에 의해 실행되거나, 실행되지 않는 상황을 의미합니다.
- 해결책 : branch prediction(예측)으로 해결합니다.
Data Dependency : Data 간의 의존성으로 인해 다른 결과가 발생하는 경우
- 해결책 : Operand Forwarding 으로 해결합니다.
  ** Operand Forwarding : 계산된 결과를 기존 pipeline과 다르게 forwarding 해주는 것입니다.

------------------------------------------------------------------------------------------------------

Data Race (데이터 레이스) : 동시에 같은 데이터를 접근 하는 것
- 이런 문제가 발생할 가능성이 있는 곳을 Critical Section이라고 합니다.
- 해결책 : mutual exclusive하게 실행(한번에 하나만 실행)하면 해결할 수 있습니다.
  - Lock/Unlock 혹은 BusyWait를 활용해 구현합니다

멀티 코어
- Homogeneous : 한 칩 안에 같은 여러개의 코어
  - 종류
    - 프로세서 자체를 여러개 두는 것
    - 프로세서 내부에 하나의 공유 메모리로, 여러개의 Thread를 사용하는 것
      (Shared Memory Parallel Programming)
    - 여러개의 프로세서가 하나의 공유 메모리에 연결되어 있는 것
      (Symmetric Multi-Processing(SMP))
      ex) NUMA(Non-Uniform Memory Access)
  - 단점
    - Cache Coherence : 프로세서 각각의 캐시 값이 서로 달라, 실제 값에 혼돈이 생기는 것
    - False Sharing : Cache Coherence를 해결하기 위해 모든 캐시를 최신화(Update)하지않고 사용할 수 없는것에 대한 flag를 달아두는 경우(Invalidation), 거짓 공유를 발생할 수 있습니다. 이 때, 오버헤드를 발생키거나 서로 간의 Ping-Pong이 발생합니다.
- Heterogeneous : 한 칩 안에 다른 여러개의 코어
  - 종류 : CPU + GPU, CPU + 가속기, Big-Little
    ** Big-Little : ARM에서 제안한 멀티 코어 구조로, 느리고 효율적인 Cortext-A7과 빠르고 비효율적인 Cortex-A15를 하나의 칩에 집적하는 기술
  - 플랫폼 병렬 모델 설계 : OpenCL, CUDA
    ex) host CPU(C,C++,Python) + 여러개의 GPU(OpenCL, Cuda)

Cache의 저장방법
- Direct-Mapped Cache
  - i번째 block은 (i%L) 번째 cache line에 저장합니다
  - 동일 위치에 Cache miss가 높아집니다.
- Fully Associative Cache
  - block을 모든 캐시 line에 저장 가능
  - 찾을 때 모든 tag를 확인해야합니다.
- Set-Associate cache : 위 두가지를 절충한 방법입니다

---------------------------------------------------------------------------------------------------------
word : 저장 장치로부터 프로세서의 레지스터에 옮겨놓을 수 있는 데이터 단위
ex) 32bits, 64bits
cash line : 캐시에 있는 임시 block 하나를 의미.
block : 메인 메모리에 있는 block을 의미.
cash line tag : 임시 block의 이름(tag)으로, 메인 메모리의 block 중 어떤 주소에 있는 것인지를 의미합니다.
2^n block(Index) : 1 block의 위치로, 메인 메모리의 크기가 2^n blocks 일 때, 2^n개의 block을 표시할 tag 혹은 K associative-way가 필요합니다.
2^m word(Offset) : 1 block의 크기로, 한 block의 크기가 2^m word일 때, m bits가 필요합니다.
---------------------------------------------------------------------------------------------------------

[directed mapped / K-way set-associative / fully-associative]

Cache Miss의 종류
- Cold Miss : 특정 Data block에 처음 접근시에 발생하는 Miss (어쩔 수 없음)
- Capacity Miss : Data Block의 양이 캐시의 양보다 너무 많아서 발생하는 Miss
- Conflict Miss : 캐시는 충분히 크지만, 위치에 의해 발생한 Miss

b. 병렬 처리 용어

커널(Kernel) : 디바이스에서 실행되는 코드의 기본 단위. 하나의 프로그램은 여러개의 커널 인스턴스로 이루어짐
컨텍스트(Context) : 커널이 실행되는 환경 및 커널 인스턴스들의 단위로, 컨텍스트 단위로 같은 메모리와 같은 디바이스 내에서 수행.
(OpenCL) Work-Group & Work-Item
(Nvidia) Thread Block & Thread
- Work Item & Thread : 커널 인스턴스 하나를 실행하는 단위
  - get_global_id() : 전체 NDRange에서의 work item의 index
  - get_local_id() : work group내에서의 work item의 index
- Work Group & Thread Block: 여러개 워크 아이템의 묶음으로, 하나의 CU(Compute Unit)에서 실행하는 단위. "Work Items per Work Group" 개수, 즉 local work size가 성능에 가장 중요한 요소입니다.
(OpenCL) CU(Compute Unit)
(Nvidia) SM(Streaming Multi-processor) : 하드웨어 상 연산 스케줄링 단위이며, 다수의 ALU가 명령어 stream을 실행하는 집단 단위입니다. 예를들어 AMD GPU의 경우 16/32/60/96...개 의 CU를 가지고 있는 경우가 있으며, Nvidia GPU의 경우 A100에는 108개의 SM이 포함되어있습니다.
** SM은 8개의 SP(Stream Processor), 혹은 32개의 CUDA Core로 이루어졌다고 볼 수도 있습니다.
(GPU마다 다를 수 있습니다.)
(Nvidia) Warp : Nvidia GPU에서 thread의 뭉치로 32개의 thread를 warp단위로 부릅니다. GPU마다 SM당 실행가능한 최대 Warp의 수가 권장됩니다.
ex) 하나의 work group에 256개의 thread가 있을 때, 하나의 CU에서 8개의 warp를 context switching하면서 실행합니다.
ex) A100의 경우 SM 마다 최대 64개의 Warp를 할당할 수 있습니다.
(AMD GPU) WaveFront : AMD GPU에서 thead의 뭉치로 64개의 thread를 wave front단위로 부릅니다.
(Nvidia) SM내에 Cuda Core 이외 유닛들
- SFU(Special Function Unit) : 지수함수, 삼각함수, 로그함수 등 초월함수를 빠르게 하기 위한 유닛
- Tensor Core : Low precision을 활용하며 4x4x4 matrix MAC연산을 위한 빠른 유닛
  - MAC연산을 빠르게 할 수 있다.
  - Mixed Precision을 활용할 수 있습니다.
    ** Mixed Precision : 중간에 16bitx16bits+32bits과 같은 연산을 가능하게 해서 손실되는 precision을 최소화 하는 방향으로

		OpenCL	CUDA(Nvidia)
코드 구현		Host 코드 : 실행 이전에 Compile (.c) Kernel 코드* : 실행 중 Compile (*.cl)	Host 코드 : 실행 이전에 Compile (.cu) Kernel 코드* : 실행 이전에 Compile (*.cu)
Logical	연산 단위	Work Item	Thread(1/32 Warp)
	연산 단위	Work Group	Thread Block
	차원 명칭	NDRange	Grid
Physical	메모리	Global 메모리
		Constant 메모리
		Local 메모리	Shared 메모리
		Private 메모리	Registers
	연산 단위	PE(Processing Element)	CUDA 코어
	연산 단위	CU(Compute Unit)	SM(Streaming Multi-processor)

2. OpenCL 함수

다음으로 OpenCL을 활용해 직접 프로그램을 작성할 때 위의 내용이 어떻게 적용되는지를 살펴보도록 합니다.

구현시에는 Host 코드와 Device코드로 나뉘며, Host 코드는 CPU에서 실행하기 위한 코드이고, Device코드는 GPU에서 실행될 우리의 로직입니다.

a. Host 함수

Host 코드의 경우 C/C++ 를 활용해 구현하는 것을 예로 들었습니다.

// Host Code (*.c)

#include <CL/cl.h>

#define CHECK_ERROR(err) \
	if (err != CL_SUCCESS){ \
    	printf("[%s:%d] OpenCL error %d\n", __FILE__, __LINE__, err);
        exit(EXIT_FAILURE);
    }
    
Cl_int err;

// 설정 ************************************************
cl_platform_id platform;
err = clGetPlatformIDs(1, &platform, NULL);
CHECK_ERROR(err);

cl_device_id device;
err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
CHECK_ERROR(err);

cl_context context;
context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
CHECK_ERROR(err);

cl_command_queue queue;
queue = clCreateCommandQueue(context, device, 0, &err);
CHECK_ERROR(err);

char* kernel_source;
size_t kernel_source_size;
kernel_source = get_source_code("kernel.cl", &kernel_source_size);

cl_program program;
program = clCreateProgramWithSource(context, 1, (const char**)&kernel_source, &kernel_source_size, &err);
CHECK_ERROR(err);

err = clBuildProgram(program, 1, &device, "", NULL, NULL);
CHECK_ERROR(err);

cl_kernel kernel;
kernel = clCreateKernel(program, "function_name", &err);

cl_mem buffer_input, buffer_output;
buffer_input = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(int)*100, NULL, &err);
buffer_output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int)*100, NULL, &err);
CHECK_ERROR(err);
// 수행 ************************************************
int *input = (int*)malloc(sizeof(int)*100);
int *output = (int*)malloc(sizeof(int)*100);

err = clEnqueueWriteBuffer(queue, buffer_input, CL_FALSE, 0, sizeof(int)*100, input, 0, NULL, NULL);
CHECK_ERROR(err);

err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer_input);
CHECK_ERROR(err);

size_t global_size = 1000;
size_t local_size = 50;
err = clEnqeueNDRangeKernel(queue, kernel, 1, NULL, &global_size, &local_size, 0, NULL, NULL);
CHECK_ERROR(err);

int C[100];
err = clEnqueueReadBuffer(queue, buffer_output, CL_TRUE, 0, sizeof(int)*100, C, 0, NULL, NULL);
CHECK_ERROR(err);

// 마무리 ************************************************
clReleaseMemObject(buffer_input);
clReleaseMemObject(buffer_output);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseCommandQueue(queue);
clReleaseContext(context);

설정
- clGetPlatformIDs() : 플랫폼 ID를 얻는 방법 (AMD, Intel, Nvidia...)
- clGetDeviceIDs() : 플랫폼 내 호스트인 CPU혹은 디바이스인 GPU의 ID를 얻는 방법
- clCreateContext() : 컨텍스트를 만듭니다.
- clCreateCommandQueue() : 커널 실행/ 데이터 전송/ 동기화 등을 위한 Command Queue를 만듭니다.
- get_source_code() : '2-b. Device 함수'에서 구현된 코드를 가져오기 위해 하기에 적힌 커스텀 함수입니다.
- clCreateProgramWithSource() : 프로그램 오브젝트 생성
- clBuildProgram() : 프로그램 오브젝트 빌드
- clCreateKernel() : 커널 오브젝트 생성
- clCreateBuffer() : 메모리의 데이터 전송하기 위한 버퍼(Command Queue) 오브젝트 생성
수행
- clEnqueueWriteBuffer() : 버퍼에 데이터를 전송하는 방법
  - blocking_write == CL_TRUE : 동기식 입력, Command Queue에 Wrtie하고 Buffer에 Write까지 완료되면 return
  - blocking_write == CL_FALSE : 비동기식 입력, Command Queue에 Write하고 return
- clSetKernelArg() : 커널의 argument에 해당하는 버퍼 및 데이터를 할당하는 방법
- clEnqeueNDRangeKernel() : 실행될 커널을 커맨드 큐에 추가!! 이때 실제 커널 수행을 명령합니다.
  - const size_t *global_work_size : 전체 work item 개수
  - const size_t *local_work_size : work group당 item개수
- clEnqueueReadBuffer() : 버퍼에서 데이터를 읽는 방법
기타
- clFinish(queue) : 특정 commandQueue의 모든작업이 끝날때까지 기다림
- clWaitForEvents(num_events, events_list) : 특정 event들이 모두 Complete상태일때까지 기다림

// get_source_code() function

#include <stdio.h>
#include <stdlib.h>

char *get_source_code(const char *file_name, size_t *len){
	char *source_code;
    size_t length;
    FILE *file = fopen(file_name, "r");
    if(file == NULL){
    	printf("[%s:%d] Failed to open %s\n", __FILE__, __LINE__, file_name);
        exit(EXIT_FAILURE);
    }
    
    fseek(file, 0, SEEK_END);
    length = (size_t)ftell(file);
    rewind(file);
    
    source_code = (char *) malloc (length+1);
    fread(source_code, length, 1, file);
    source_code[length] = '\0';
    
    fclose(file);
    
    *len = length;
    return source_code;
}

b. Device 함수

Device 코드의 경우, 우리는 OpenCL을 활용해 구현하고 있으므로, 확장자는 *.cl 으로된 파일을 만들 예정입니다.

// Device Code(*.cl))

__kernel void function(__global int *input1, __constant int *input2, __global int *output){   
	int global_id1 = get_global_id(0);
    int global_id2 = get_global_id(1);
	
	output[id2] = input1[id1] + input2[id1];
}

전체 정보 받기
- get_work_dim() : 차원을 얻기 위함
- get_global_size(차원) : global 사이즈를 얻기 위함
- get_num_groups(차원) : work-group의 개수를 얻기 위함
- get_local_size(차원) : work-group의 local 사이즈를 얻기 위함
연산 정보 받기
- get_global_id(차원): global work-item의 인덱스를 얻기 위함
- get_group_id(차원) : work-group의 인텍스를 얻기 위함
- get_local_id(차원): work-group 내의 local work-item의 인덱스를 얻기 위함
기타
- barrier(CLK_LOCAL_MEM_FENCE) : 워크 그룹내 작업중인 로컬 메모리에 대해 모두 작업이 끝날 때까지 busy wait
- barrier(CLK_GLOBAL_MEM_FENCE) : 워크 그룹내 작업중인 글로벌 메모리에 대해 모두 작업이 끝날 때까지 busy wait
- 아토믹 연산자(atomic_add, atomic_sub, atomic_inc ...) : 즉시 다른 모든 워크아이템에 적용 가능하도록 하는 것으로, 순서는 보장할 수 없으며, 자주 사용하면 성능이 대폭 하락합니다.

3. 병렬처리 최적화 방법

위의 방법으로 병렬 처리 코드를 작성하는 것은 사실 오픈소스를 활용해서도 구현하기 어렵지 않으실 거라는 생각이 듭니다.

다만, 이런 코드를 최적화하는 것이 더욱 어려운 부분일 것입니다.

우리는 어떤 부분을 참조해 병렬처리를 최적화하곤 하는지를 살펴보기 위해 아래 내용을 살펴보고자 합니다.

a. 커널 성능 저하 요인

1. GPU의 ALU 자원을 충분히 활용하지 못하는 경우
2. 메모리 Bandwidth를 비효율적으로 사용하는 경우
- Global (GPU 전체) ~ Local 혹은 Device (CU)
- Local 혹은 Device (CU) ~ Private 혹은 캐시 (PE)
3. 앞선 두개의 중첩 : ALU연산시간과 메모리 이동시간을 서로 숨길 수 있도록

b. 최적화 방법 개요

Work item의 전체 개수가 충분히 많아야합니다
- 최소개수
  - Physical : SM 개수 x SM 당 ALU의 개수
  - Logical : SM 개수 x SM 당 한번에 가능한 Warp수 * Warp사이즈 (ex. 32)
- 메모리 명령을 숨기고 tail effect 효과를 줄이기 위해, 최소 개수보다 조금더 많은 Work-Item을 만들어 주는 것이 좋습니다.
  ** tail effect : 전체 수행량을 병렬처리하고 나서, 남는 task로 인해 쉬는 유닛이 생기는 경우
- Coarse-grained vs Fine-grained : 적절한 방법을 취한다.
  - Coarse-grained : outermost loop를 병렬화 함으로써, 여러개 연산을 한번에 하는 방법입니다.
  - Fine-grained : intermost loop를 겹겹이 쌓는 등 연산을 잘게 쪼개는 방법입니다. 중복된 메모리 접근이 생길 수도 있으니, 중복된 연산 간의 동기화가 필요할 수도 있습니다.

조건문(If문)등을 할때, 같은 Warp내에서 work item간의 Branch Divergence를 최소화해야합니다.
- 즉, 데이터마다 If의 결과가 다른 경우가 많으면 divergence가 발생합니다
- 해결책1. Data Reordering : 데이터의 순서를 바꾸는 방법으로, Reordering에 추가적인 오버헤드가 발생하지만, 접근패턴을 최적화 가능합니다.
- 해결책2. Indirect Access : 데이터에 맞게, 해당하는 연산과의 mapping table을 만드는 방법

Loop Unrolling : Loop를 크게 쪼개서 (i +=1 → i+=N) 여러개의 연산을 한번에 실행하기
- 단, 하나의 item을 처리하기 위해, 한번에 실행하므로 많은 register가 필요할수도 있습니다.

Occupancy(Active Warp / Max Warp) 최적화
** Active Warp : 현재 하나의 SM에 할당된 warp의 개수
- 다음의 메모리 사용량을 적절히 줄여야 Active Warp를 늘릴 수 있습니다.
  - Local Memory(Shared Memory) 사용량 : 하나의 group을 위한 커널에 따라 컴파일러에 의해 + 런타임에 따라 결정됩니다.
  - private memory(Registers) 사용량 : 하나의 item을 위한 커널에 따라 컴파일러에 의해 결정됩니다.
  - 따라서 warp를 늘리는 방법은 global
- 위 두 메모리를 너무 안쓰고 Global Memory를 많이쓴다면, (computationally expensive한) 메모리 접근 중에 다른 연산이 동시에 진행되지 않을 수도 있다.
- 위 두 메모리를 너무 많이 쓴다면, SM당 최대 메모리 개수가 정해져있으므로, Thread개수 자체가 많이 생기지 않을 수도 있다.

Local한 Work group의 크기를 잘 정해야 한다.
- Warp 단위의 배수로 정하기
- 하드웨어마다 Max Work Group 크기가 정해져있으므로 참조한다.
- Private/Local Memory 사용량이 많다면, Work-group의 크기를 줄여야 Register Spilling을 줄일 수 있습니다.
  ** Register Spilling : 두개의 thread가 같은 register를 사용할 때 Overwrite를 방지하기 위해 stack에 미리 저장한 후 원래 자리의 레지스터로 복원시키는 솔루션.

c. Case1. Host ↔ Global Mem ↔ Kernel

__kernel void Matrix_multiply(__global float *A, __global float *B, __global float *C){
	int i = get_global_id(0);
    int j = get_global_id(1);

    if (i>=N || j>N) return;
    
    C = [i*N + j] = 0.0;
    int k;
   	float ans = 0.0;
    for (k=0;k<N;k++){
    	 ans += A[i*N + k] * B[k*N + j];
    }
    C[i*N + j] = ans;
}

위 예시에서..
- A,B는 global memory에 read
- C는 global memory에 write
- k 는 private memory

Cache Hit Rate을 높이기
- Capacity Miss를 줄이자
  - 하나의 warp에서 작업하는 data의 단위를 작게 잡거나, warp 전체 개수자체를 줄인다.
- Conflict Miss를 줄이자
  - 메모리 접근 패턴을 최적화한다.
    - 방법1. global memory에 접근시, get_global_id(0)를 기준으로 병렬화할 것인지, get_global_id(1)을 기준으로 병렬화할 것인지에 따라서, global memory에 대한 conflict miss를 줄일 수 있다.
      ** 뒤에서 예시로 자세히 설명해보겠습니다.
    - 방법2. 위에서 언급한 것과 같이 input 데이터 Layout을 Reordering하거나 Indirect Access 하여 spatial locality를 높여주거나 Padding하여 cache에 대한 접근용이성을 높입니다.
      ** Padding : 데이터 중간에 일부러 빈칸을 넣어 캐시라인이 어긋나게 만든 후 confict miss를 줄이는 방법.
  - 위에서 언급한 Data Reordering, Indirect Access 를 통해

Memory 접근 회수 자체를 최소화
- Global Memory 에의 접근 중 중복된 것은 제거
  - 위 예시중, A[i]와 ,B[i]를 여러번 read 하는 경우 private memory에 저장해두고 계산하기
  - 위 예시중, C[i]에 여러번 write 하는 경우 private memory에 저장해두고 넣기
- Memory Coalescing 하기 (같은 block에 접근하는 것들을 모아서 처리하기)
  ** Memory Coalscing : 해당 블록에 접근하는 것들을 모아서, 한번의 메모리 접근으로 처리하기
  - Vectorization : 위 예시중 A[i], A[i+1], A[i+2] … 를 처리하는 경우 연속 위치 접근시 벡터화해 for(i=0; i<N; i+=3) 형태로 한번에 A[i], A[i+1], A[i+2]를 모두 처리

예시) Matrix Multiplication : 코드는 위에 제시되었습니다.
- 기본 조건은 Warp의 사이즈는 32이며 Work Group은 16x16(8warp)인 상황이라고 가정합시다.
- 기본적으로 위에 코드는 연산시 위 그림과 같이 연산될 것입니다. 필요한 정보를 정리해보겠습니다.
  - 1. 먼저 각각의 thread는 C Matrix의 Output NxN'개를 연산하기 위한 thread(Work-item)가 필요합니다.
    ** 즉, 위에서 1개의 thread는 C[i][j]를 연산하기 위해 진행하며, k개의 loop를 가지게 됩니다.
  - 2. 필요한 thread의 개수를 work group단위로 나누면 총 work group개수가 정해지고, 수많은 work group들을 CU에 분배해야 합니다.
  - 3. CU마다 할당 가능한 Work Group의 개수는 3가지로 결정됩니다.
    - CU의 레지스터 개수 / 커널당 필요한 register 개수 = SM에 가능한 thread개수
    - CU의 shared memory / 커널당 필요한 shared 메모리 = SM에 가능한 thread개수
    - CU에 할당 가능한 최대 warp의 개수
    ** 같은 Work Group은 같은 CU에서 진행되므로, local memory가 있다면 같은 local memory를 공유합니다.
  - 4. 하나의 CU에 여러개의 work group이 할당되었고, 이 중에 warp단위로 context switch를 진행합니다.
    ** 하나의 warp가 돌아가는 동안, 다른 warp들은 register에 필요한 data를 채우는 등의 작업을 통해, context switch에 latency가 거의 들지않습니다. register에 필요한 data를 채우는 작업이 바로 Local Memory 혹은 Global Memory에서 데이터를 가지고 오는 작업입니다.
  - 5. 하나의 work group은 결국 C matrix의 어떤 16x16 output을 연산하기 위한 thread의 집합이고, 해당 work group내에 thread가 분배되는 방법은 [아래 그림]과 같이 Row-Major Order 혹은 Column-Major Order로 배치됩니다. 위의 Matrix Multiplication을 하는 그림의 경우는 thread들이 Column-major order로 분배되어 8개의 warp가 배치되는 방법을 보여준 것입니다.
    ** 데이터는 reordering등의 작업을 따로하지 않는 이상, 일반적으로 row-major order로 저장되어 있을 것입니다.
  - 6. 위 그림의 경우 warp 내의 thread들이 A matrix에 대해서 다른 메모리를 요구하므로 각각 cache line을 요구 할텐데, set-associative의 way가 충분하지 않으면, warp 내의 thread들이 각각의 cache line을 요구하므로 conflict miss가 일어나게 됩니다. 특히나, N이 K의 배수인 경우는 같은 set에 포함될 확률이 더 높아지므로, conflict miss가 굉장히 많아 질 것 입니다.
    ** memory divergence : 하나의 warp에서 여러 번의 memory access가 발생하는 경우
    ** conflict miss : 위에서 이미 설명했지만 같은 메모리 주소를 가지는 Cache set에 접근하는 여러개의 warp 혹은 여러개의 thread 때문에 cache에게 eviction을 강제로 하게 하는 miss 입니다.
  - 7. 이때 아래와 같이 get_global_id(0)와 get_global_id(1)을 바꾸는 등의 방법으로 thread가 배치되는 순서를 Row-Major Order로 바꾸게 되면, warp내의 thread들은 A matrix에 대해 서로 다른 메모리를 요구할 확률이 줄어들기 때문에 메모리에 대한 접근을 최적화할 수 있습니다.
  - 8. 최종적으로 아래 [결과 그림]과 같이 메모리 접근 최적화가 가능합니다.

[2D에서 Row-Major Order와 Column-Major Order]

__kernel void Matrix_multiply(__global float *A, __global float *B, __global float *C){
	int j = get_global_id(0); // Changed
    int i = get_global_id(1); // Changed

    if (i>=N || j>N) return;
    
    C = [i*N + j] = 0.0;
    int k;
   	float ans = 0.0;
    for (k=0;k<N;k++){
    	 ans += A[i*N + k] * B[k*N + j];
    }
    C[i*N + j] = ans;
}

d. Case2. Host ↔ Global Mem ↔ Local Mem ↔ Kernel

GPU에는 CPU와 다르게 Global Memory보다 빠른 Local Memory가 CU 단위로 있습니다. OpenCL에서 이를 사용하기 위해서는 아래와 같은 방법을 활용합니다.
- 방법1. 커널 함수 안에서 __local 변수 선언을 통해 접근 가능합니다.
- 방법2. Host로부터 __local 포인터를 받아 사용하기 (아래 예시)

__kernel void foo(__local int *A){
	__local float B[64];
}

Local Memory를 활용 전 고민요소
- SM 단위로 Active warp 의 개수를 줄여서, 하나의 warp 가 작업하는 데이터가 충분하도록 만들거나
- SM 단위로 Active warp 의 개수를 높이기 위해, 하나의 warp에서 작업하는 데이터를 최소화 하거나
Local Memory을 활용해 최적화 하려는 목표
- 결론적으로, 아래의 두가지 tradeoff를 결정하기 위해서는 연산 granularity를 잘 조절하는 것이 중요합니다.
  - 즉, locality를 보장하면서, 필요한 데이터를 충분히 사용하기 위해서입니다.
- (예로 한가지 방법) work group을 tiling하는 방법 : work item 하나가 연산할 개수와 work group 내의 work item의 배치형태를 결정해 타일 단위로 shared memory에 위치시켜 Memory Coalescing하는 것.

병렬화 증가하기 위해 working data를 줄이면... (Fine-grained)	병렬화 감소하기 위해 working data를 늘리면... (Coarse-grained)
적은 data를 캐싱	많은 data를 캐싱
Local 메모리 적게 사용해도 된다.	Local 메모리 많이 사용해야한다
(결과) Data Locality 좋지 않다. (여러번 접근해야 하므로)	(결과) Occupancy 좋지 않다. (Active warp 가 줄어서)