스레드 블록 아키텍처(Thread Block Architecture) - 그리드 블록 모델 -3

2.6 매트릭스 곱

이전 글에서 살펴본 매트릭스는 12 X 12 크기로 작상서 GPU가 최대 효율을 발휘하지 않는다. 그러나 CUDA 프로그래밍에 익숙해지고 2차원 스레드의 작업 분할에 대하여 알아보는 것이 목적이므로 작은 매트릭스 크기로 진행하는 것이 이해하기 쉽다. 행렬의 곱 M X N = P 는 다음과 같이 계산된다.

P( tx, ty) = M(0, ty) X N(tx, 0) + M(1, ty) X N(tx, 1) + M(2, ty) X N(tx, 2) + M(3, ty) X N(tx, 3) + ..... M(11, ty) X N(tx, 11);

위의 수식을 C코드로 구현하게 되면 12회 계산의 for 루프를 가지게 되고 매트릭스에 12 X 12 개의 연산이 있기 때문에 총 연산은 12 X 12 X 12 회가 된다.

| 매트릭스 곱 M X N = P

| code1

void MatrixMulC(int *M, int *N, int* P, int Width){

int col = 0;

int raw = 0;

nit index = 0;

int Destindex = 0;

for( col = 0; col < Width; col++)

{

for( raw = 0; raw < Width; raw++)

{

Destindex = col * Width + raw;

for(index = 0; index < Width; index++)

P[Destindex] += M[col * Width + index] * N[index * Width + raw];

}

}

}

C로 구현한 12 X 12 의 매트릭스의 연산 횟수는 3승수가 되어 12 X 12 X 12 = 1,728번의 계산을 하게 된다. 12 X 12 매트릭스의 크기는 수치해석에선 작은 크기로, 매트릭스의 크기가 증가함에 따라 기하급수적으로 늘어나게 된다. 30 X 30 매트릭스의곱을 계산하게 되면 27,000번의 연산을 수행하게 된다.

이와 같은 형태의 계산은 매니코어 GPU를 사용하는 CUDA프로그래밍에서 큰 장점이 있다. 위의 예제를 CUDA 프로그램으로 구현할 때 12 X 12 개의 스레드로 구성하여 작업을 분할하면 하나의 스레드는 위의 수식을 한번만 계산하면 된다. 대규모 스레드가 담당하는 부하는 C 코드와 비교하면 1 / 144로 줄어들게 된다.

| 매트릭스 크기에 따른 연산 횟수 그래프

| Code2 (CUDA로 구현한 정방형 매트릭스 곱 커널 함수)

__global__ void MatrixMul( int* M, int* N, int* P, int Width)

{

int tid, tx, ty;    

//2차원 작업 분할 인덱스 계산

tx = blockDim.x * blockIdx.x + threadIdx.x;

ty = blockDim.y * blockIdx.y + threadIdx.y;

tid = Width * ty + tx;

int Value = 0;

int MVal = 0;

int NVal = 0;

for(int i = 0; i < Width; i++)

{

MVal = M[ty * Width + i];

NVal = N[i * Width + tx];

Value += MVal * NVal;

}

P[tid] = Value;

}

code1 에서 수행한 3중 루프는 CUDA 프로그램에서 1회 루프로 변경되고 2회의 루프는 스레드로 분할되어 각각의 코어에 작업이 할당된다. code2를 수행하기 위한 메인 함수는 다음과 같다.

| main 함수

int main()

{

const int MatrixWidth = 12;

const int MatrixHeight = 12;

const int MatrixSize = MatrixWidth * MatrixHeight;

const int BufferSize = MatrixSize * sizeof(int);

int* M;

int* N;

int* P_cuda;

int* P_C;

//호스팅 메모리 할당

M = (int*)malloc(BufferSize);

N = (int*)malloc(BufferSIze);

P_cuda = (int*)malloc(BufferSize);

P_C = (int*)malloc(BufferSize);

int i = 0;

//데이터 입력

for(int i = 0; i < MatrixSize; i++)

{

M[i] = i;

N[i] = i;

P_cuda[i] = 0;

P_C[i] = 0;

}

int* dev_M;

int* dev_N;

int* dev_P;

//디바이스 메모리 할당

cudaMalloc((void**)&dev_M, BufferSize);

cudaMalloc((void**)&dev_N, BufferSize);

cudaMalloc((void**)&dev_P, BufferSize);

//호스트 디바이스 입력 데이터 전송

cudaMemcpy(dev_M, M, BufferSize, cudaMemcpyHostToDevice);

cudaMemcpy(dev_N, N, BufferSize, cudaMemcpyHostToDevice);

dim3 Dg(3, 4, 1);

dim3 Db(4, 3, 1);

//CUDA kernel 매트릭스 곱 계산

MatrixMul<<<Dg, Db>>>(dev_M, dev_N, dev_P, 12);

//디바이스 호스트 출력 데이터 전송

cudaMemcpy(P_cuda, dev_P, BufferSize, cudaMemcpyDeviceToHost);

//C 함수 매트릭스 곱 계산

MatrixMulC(M, N, P_C, 12);

bool ResultFlag = true;

//결과 출력

for( i = 0; i < MatrixSize; i++)

{

//printf("Result[%d] : %d, %d\n", i, P_cuda[i], P_C[i]);

if(P_cuda[i] != P_C[i]) ResultFlag = false;

}

if(ResultFlag == true) printf("MatrixMul Result OK!\n");

else printf("MatrixMul Result Error!\n);

cudaFree(dev_M);

cudaFree(dev_N);

cudaFree(dev_P);

free(M);

free(N);

free(P_cuda);

free(P_C);

return 0;

}

위 코드 프로그램의 진행은 다음과 같다.

① 호스트 M, N, P_cuda, P_C 12 X 12 X sizeof(int)힙 메모리 할당. M, N 매트릭스는 입력용, P_cuda는 cuda kernel 결과용P_C는C 함수 결과용으로 사용한다.

② 디바이스 dev_M, dev_N, dev_P에 GPU 메모리 할당

③ 매트릭스 M, N을 0~143까지 차례로 입력하여 초기화

④ 매트릭스 M, N을 디바이스 메모리로 전달한다.

⑤ 그리드(3, 4), 블록(4, 3)을 생성하여 3 X 4 X 4 X 3 = 144개의 스레드를 생성을 설정한다.

⑥ 144개의 스레드를 수행하느 CUDA 매트릭스 곱 커널 실행

⑦ C_cuda 호스트 메모리로 GPU 계산 결과 출력

⑧ C 함수를 이용하여 매트릭스 곱 계산, P_C 메모리로 결과를 얻는다.

⑨ CUDA 매트릭스 곱 결과 P_cuda와 C 매트릭스 곱 P_C 결과를 비교한다.

⑩ 결과를 출력하고 모든 메모리를 해제한다.

| 매트릭스 곱 프로그램 실행 결과

결과가 출력되는 부분의 주석을 삭제하면 매트릭스 요소의 값을 확인할 수 있다.

2.7 스레드의 3차원 구성

그리드는 현재 2차원까지 구성할 수 있으며, 블록은 3차원으로 구성할 수 있다. dim3형의 변수로 블록 안의 스레드 개수를 3차원적으로 지정한다. Db(x, y, z)로 각 차원을 지정하여 x는 1부터 512까지, y는 1부터 512까지, z는 1부터 64까지 사용할 수 있다. x X y X z 의 값은 512를 넘으면 안된다.

1<= x or y <=512, 1<= z <= 64, 1<= x * y * z <= 512

2.7 그리드

스레드가 모여 블록을 구성하고 블록이 모여 그리드를 구성한다. 그리드는 커널이 실행될 때 스레드를 구성하는 최상위 집합체로 디바이스에서 커널이 실행되는 모듈, 응용 프로그램을 나타낸다. 하지만 디바이스에서 실행되는 그리드는 동일한 시점에 하나만 존재할 수 있다. 복수의 디바이스에서는 개별적으로 그리드를 실행시킬 수 있으며 동일한 그리드를 실행시키는 것도 가능하다.

 

위 사진은 CUDA프로그램에서 차례로 커널 2개를 실행시킨는 과정을 나타내고 있다. 0번 커널에 대한 그리드를 생성할 때 블록과 스레드를 구성하고 완료시킨다. 프로그램이 진행되는 중에는 스레드를 구성하는 그리드의 구조를 변결할 수 없다. 0번 커널이 완료되고 1번 커널을 실행할 때는 새로운 스레드 구조로 그리드를 생성하여 실행할 수 있다.