커널 호출(워크 그룹, 워크 아이템, 데이터 병렬, 태스트 병렬)

< 데이터 병렬과 태스크 병렬 >

• 프로세서 별로 같은 커널을 실행하면 데이터 병렬, 다른 커널을 실행하면 태스크 병렬임
• GPU는 굉장히 많은 프로세서를 가지고 있지만 서로 각기 다른 프로그램을 실행하기 적합하지 않음

 

< 데이터 병렬/태스크 병렬 커널 구조 >

• 데이터 병렬과 태스크 병렬용 호스트 함수가 따로 있음
• 태스크 병렬용 함수는 병렬로 돌릴 커널의 개수만큼 함수를 호출해줘야함

구현하려는 배열 연산

// 4행 4열의 float형 배열 A와 B가 있다고 할 때 사칙연산 동시에 해보기

// 데이터 병렬
__kernel void dataParallel(__global float* A, __global float* B, __flobal float* C)
{
	int base_id = 4*get_global_id(0); // NDRangeKernel의 global item size = 4
	
    C[base_id+0] = A[base_id+0] + B[base_id+0];
    C[base_id+1] = A[base_id+1] - B[base_id+1];
    C[base_id+2] = A[base_id+2] * B[base_id+2];
    C[base_id+3] = A[base_id+3] / B[base_id+3];
}

// 태스크 병렬
__kernel void taskParallelAdd(__global float* A, __global float* B, __global float* C)
{
	int base_id = 0; // 병렬적으로 수행하는게 아니라 직렬적으로 배열 값 계산
    
    C[base_id+0] = A[base_id+0] + B[base_id+0];
    C[base_id+4] = A[base_id+4] + B[base_id+4];
    C[base_id+8] = A[base_id+8] + B[base_id+8];
    C[base_id+12] = A[base_id+12] + B[base_id+12];
}

__kernel void taskParallelSub(__global float* A, __global float* B, __global float* C)
{
	int base_id = 1;
    
    C[base_id+0] = A[base_id+0] - B[base_id+0];
    C[base_id+4] = A[base_id+4] - B[base_id+4];
    C[base_id+8] = A[base_id+8] - B[base_id+8];
    C[base_id+12] = A[base_id+12] - B[base_id+12];
}

__kernel void taskParallelMul(__global float* A, __global float* B, __global float* C)
{
	int base_id = 2;
    
    C[base_id+0] = A[base_id+0] * B[base_id+0];
    C[base_id+4] = A[base_id+4] * B[base_id+4];
    C[base_id+8] = A[base_id+8] * B[base_id+8];
    C[base_id+12] = A[base_id+12] * B[base_id+12];
}

__kernel void taskParallelDiv(__global float* A, __global float* B, __global float* C)
{
	int base_id = 3;
    
    C[base_id+0] = A[base_id+0] / B[base_id+0];
    C[base_id+4] = A[base_id+4] / B[base_id+4];
    C[base_id+8] = A[base_id+8] / B[base_id+8];
    C[base_id+12] = A[base_id+12] / B[base_id+12];
}

 

< 워크 그룹 >

• 위에서 설명한 데이터 병렬은 워크 아이템 단위로 진행되었음
• 워크 아이템을 일정 수로 묶으면 워크 그룹
• 워크 그룹을 구성하는 워크 아이템의 수는 하나 이상이며 상한선은 플랫폼에 따라 상이
• 워크 그룹안의 워크 아이템은 동기처리나 로컬 메모리 공유가 가능함
• 데이터 병렬 방식의 커널을 실행할 때는 워크 아이템 뿐만 아니라 워크 그룹의 수도 입력을 지정해야하는데 cl::NullRange를 쓰면 자동으로 지정됨
• 모든 워크 그룹이 가지는 워크 아이템의 수는 동일하게 지정됨
• 이때 중요한게 전체 워크 아이템 수가 워크 그룹에 속하는 워크 아이템의 수로 나누어 떨어지지 않으면 안됨 -> 만약에 나누어 떨어지지 않으면 CL_INVALID_WORK_GROUP_SIZE를 반환
• 전체 워크 아이템에 대한 ID만 쓸 수 있는게 아니라 워크 그룹 안의 워크 아이템도 ID가 있어서 이 두개를 써서 ID를 나타낼 수 있음
• 디바이스 별로 인덱스 공간의 차원 수 상한선을 확인하려면 clGetDeviceInfo()함수를 사용해 CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS의 값을 얻으면 됨 -> cl_uint형식으로 반환
• 하나의 워크 그룹에 속하는 워크 아이템 수의 상한선을 확인하려면 CL_DEVICE_IMAGE_SUPPORT 값을 얻으면 됨   -> size_t 형태로 반환(NDRange에 입력하는 것도 size_t)

int id_global = get_global_id(); // 글로벌 워크 아이템 ID
int id_group = get_group_id(); // 워크 그룹 ID
int id_local = get_local_id(); // 로컬 워크 아이템 ID

인덱스 공간

 

< 태스크 병렬과 이벤트 오브젝트 >

• 커맨드 큐에 태스크를 전달하면 기본적으로는 병렬로 실행됨
• 그런데 의존 관계가 있는 태스크는 병렬로 실행하면 문제가 생길 수 있음
• 그래서 존재하는게 이벤트 오브젝트임
• 데이터 병렬에서도 쓰기는 쓰는데 그거는 커널 처리시간 측정용 플래그로 사용하기 위함임