CUDA 프로그래밍 강의를 들은건 작년 1학기 였는데 상당히 매력적인 과목이었다.
예제로 나온 CPU 코드와 GPU 코드를 돌려서 걸리는 시간을 확인했을 때의 충격이 엄청났다.
아무튼, 이번엔 CUDA 과제에 대한 질문을 받았다.
//block과 thread의 개수를 고정 (BLOCK_WIDTH와 TILE_WIDTH 값을 변경하지 말 것)
#define BLOCK_WIDTH 32
#define TILE_WIDTH 32
dim3 dimGrid(BLOCK_WIDTH, BLOCK_WIDTH);//32*32개 블록
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH);//하나의블록당 32 * 32개의 쓰레드Matrix 간의 병렬곱을 하는 것이 과제였는데, BLOCK_WIDTH와 TILE_WIDTH 값이 고정된 채로 진행하는 것이 과제였다.
dimGrid 안에 들어갈 Block_Width 값을 임의로 변경했었다.
dim3 dimGrid(1+((width-1)/TILE_WIDTH), 1+((width-1)/TILE_WIDTH));
dim3 dimBlock(32, 32);
CUDA를 이용한 병렬 프로그래밍은 사용할 블록의 수와 쓰레드 수에 따라 성능이 결정된다.
상황에 따라 블록 수, 쓰레드 수를 늘리거나 줄일 수도 있다.
고성능 GPU가 사용할 수 있는 블록의 수와 쓰레드 수는 저성능 GPU보다 많다.
동일한 성능을 설정하기 위해서는 블록의 수와 그리드 수를 임의로 설정해야 한다.
위와 같은 이유 때문에 다음과 같은 과제를 제출해준 것이라고 생각했다. 개인적인 추측이다.
처음 교수님이 제출해준 예제 코드는 다음과 같다.
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
// block과 thread의 개수를 고정 (BLOCK_WIDTH와 TILE_WIDTH 값을 변경하지 말 것)
#define BLOCK_WIDTH 32
#define TILE_WIDTH 32
__global__ void MultMatGPU(float *P, float *M, float *N, int width);
void main()
{
// width: 행렬의 한 변의 크기
// width를 어떤 값으로 바꾸더라도 작동해야 함
//int width = TILE_WIDTH * BLOCK_WIDTH;
//int width = TILE_WIDTH * BLOCK_WIDTH * 2;
int width = TILE_WIDTH * BLOCK_WIDTH * 3;
//int width = TILE_WIDTH * BLOCK_WIDTH * 4;
//int width = TILE_WIDTH * BLOCK_WIDTH * 5;
float *P = new float[width * width];
float *M = new float[width * width];
float *N = new float[width * width];
for (int i = 0; i < width * width; ++i)
{
M[i] = 1.0f;
N[i] = 1.0f;
P[i] = 0.0f;
}
cudaError_t cudaStatus = cudaSetDevice(0);
float *dev_P, *dev_M, *dev_N;
cudaStatus = cudaMalloc((void **)&dev_P, width * width * sizeof(float));
cudaStatus = cudaMalloc((void **)&dev_M, width * width * sizeof(float));
cudaStatus = cudaMalloc((void **)&dev_N, width * width * sizeof(float));
cudaStatus = cudaMemcpy(dev_M, M, width * width * sizeof(float), cudaMemcpyHostToDevice);
cudaStatus = cudaMemcpy(dev_N, N, width * width * sizeof(float), cudaMemcpyHostToDevice);
// block과 thread의 개수를 고정 (dimGrid와 dimBlock 변수를 변경하지 말 것)
dim3 dimGrid(BLOCK_WIDTH, BLOCK_WIDTH);
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH);
clock_t st = clock();
// 이 함수에 입력하는 dimGrid, dimBlock도 변경하지 말 것
MultMatGPU<<<dimGrid, dimBlock>>>(dev_P, dev_M, dev_N, width);
cudaDeviceSynchronize();
clock_t ed = clock();
printf("Elapsed time = %u ms\n", ed - st);
cudaStatus = cudaMemcpy(P, dev_P, width * width * sizeof(float), cudaMemcpyDeviceToHost);
cudaStatus = cudaDeviceReset();
// 검증 코드
for(int i = 0; i < width * width; ++i)
{
if(P[i] != width)
{
printf_s("error!\n");
break;
}
}
delete[] P;
delete[] M;
delete[] N;
cudaFree(dev_P);
cudaFree(dev_M);
cudaFree(dev_N);
}
__global__ void MultMatGPU(float *P, float *M, float *N, int width)
{
int i = blockIdx.y * TILE_WIDTH + threadIdx.y;
int j = blockIdx.x * TILE_WIDTH + threadIdx.x;
if(i < width && j < width)
{
float sum = 0.0;
for (int k = 0; k < width; ++k)
{
float a = M[i * width + k];
float b = N[k * width + j];
sum += a * b;
}
P[i * width + j] = sum;
}
}
위의 코드는 1 x 1에서 1024 x 1024 크기(최대 BLOCK_WIDTH * TILE_WIDTH)의 배열에 대한 행렬곱을 지원한다.
위의 코드를 수정하여 1024*n이 width인 Matrix에 대한 행렬곱이 지원되는 코드를 작성하는 것이다.
오래되어서 맞는지 모르겠지만
dimGrid는 Grid 내 블록의 개수에 대한 행렬.
dimBlock은 Block 내 쓰레드의 개수에 대한 행렬.
이라고 알고 있다.
일단 MultMatGPU() 함수를 보자.
각 쓰레드는 자신의 위치에서 가로, 세로로 Width 길이만큼의 행렬곱을 수행한다.
다만, BLOCK_WIDTH와 TILE_WIDTH가 크기가 제한되어져 있기 때문에 결과는 일부분만 계산이 된채로 나타난다.
2048 x 2048 행렬이라면 위와 같이 왼쪽 1/4 만큼의 공간만 계산이 될 것이다.
그럼 다음에 해야할 행동은 계산 범위의 변경이다.
파란색 네모가 현재 계산된 부분이라면 위의 사진처럼 구역을 나눠가면서 계산을 하면 전체 행렬이 계산이 된다.
int n = width / 1024;
if (width % 1024 != 0)
++n;
for (int i = 0; i < n; ++i){
for (int j = 0; j < n; ++j){
MultMatGPU <<<dimGrid, dimBlock>>> (dev_P, dev_M, dev_N, width, i, j);
cudaDeviceSynchronize();
}
}한 번에 계산할 수 있는 영역이 1024 x 1024 크기 이므로 width를 1024로 나누어 영역이 얼마나 필요한지 확인한다.
width가 1024보다 작거나 1024의 배수가 아닐 경우를 가정하여 나머지 값이 0이 아니면 n의 값을 1 더 올린다.
나누어진 영역의 크기만큼 for문을 돌려 MultMatGPU 함수를 실행한다.
__global__ void MultMatGPU(float *P, float *M, float *N, int width, int v, int h){
int i = blockIdx.y * TILE_WIDTH + threadIdx.y + v * 1024;
int j = blockIdx.x * TILE_WIDTH + threadIdx.x + h * 1024;
if (i < width && j < width) {
float sum = 0.0;
for (int k = 0; k < width; ++k){
float a = M[i * width + k];
float b = N[k * width + j];
sum += a * b;
}
P[i * width + j] = sum;
}
}함수 내에서는 영역의 위치에 따라 인덱스 값을 바꿔준다.
이를 통해 특정 영역 내에서의 행렬곱을 수행함으로써 전체 행렬곱이 가능하게 해준다.
다른 사람이 한 내용을 보았더니 다른 방식으로 작성도 가능하다.
이 경우는 MultMatGPU 내에서 계산을 여러번 실행하는 것으로 함수 내에서 width에 대하여 여러 위치에 대한 계산을 한다.
한 쓰레드에 들어갔을 때, 1024 x 1024 내에서의 해당 위치와 1024*m x 1024*n에서 동일한 위치를 같이 계산한다.
이렇게 할 경우, MultMatGPU 함수를 dimGrid와 dimBlock에 맞춰진 횟수만큼 실행한다.
효율성 면에서는 앞의 코드보다 떨어지지만 main 함수를 수정할 수 없도록 제한된 상황이라면
다음과 같은 코드가 해답이 될 수 있다.
