CUDA 큰 벡터의 합 - 공부하는 도비

2025.01.21 - [CUDA] - CUDA 스레드 레이아웃 설정 및 커널 호출<<<>>> - 공부하는 도비

CUDA 스레드 레이아웃 설정 및 커널 호출<<<>>> - 공부하는 도비

 

앞 전 피드에서 한 블록의 x차원 최대 길이는 1,024이며 한 블록이 가질 수 있는 최대 스레드 수는 1,024개라고 하였습니다. 

따라서 큰 수가 입력될 때는, 스레드 레이아웃을 적절하게 잘 설정하는 것이 중요합니다.

#include "cuda_runtime.h"   
#include "device_launch_parameters.h"

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

#define NUM_DATA 10240

__global__ void vectorAdd(int *d_a, int *d_b, int *d_c)
{
    int tID = blockIdx.x * blockDim.x + threadIdx.x;
    d_c[tID] = d_a[tID] + d_b[tID];
}

int main()
{
    
    int *a, *b, *c;
    int *d_a, *d_b, *d_c;

    int memSize = NUM_DATA * sizeof(int);

    printf("%d elements, memSize = %d bytes\n", NUM_DATA, memSize);

    a = new int[NUM_DATA]; memset(a, 0, memSize);
    b = new int[NUM_DATA]; memset(b, 0, memSize);
    c = new int[NUM_DATA]; memset(c, 0, memSize);

    for (int i = 0; i < NUM_DATA; i++) {
        a[i] = rand() %10;
        b[i] = rand() %10;
    }


    cudaMalloc(&d_a, memSize);
    cudaMalloc(&d_b, memSize);
    cudaMalloc(&d_c, memSize);

    // Copy data to device
    cudaMemcpy(d_a, a, memSize, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, memSize, cudaMemcpyHostToDevice);

    dim3 block(NUM_DATA);
    dim3 grid(1);
    
    // If the number of data is greater than 1024, we need to use multiple blocks
    // Because the maximum number of threads in a block is 1024
    if(NUM_DATA > 1024) {
        block.x = 1024;
        grid.x = NUM_DATA / 1024;
        if(NUM_DATA % 1024 != 0) {
            grid.x++;
        }
    } else {
        block.x = NUM_DATA;
    }
    // block.x is the number of threads in a block
    // grid.x is the number of blocks in a grid

    // To check the index of the thread
    printf("grid  : (%d, %d, %d)\n", grid.x, grid.y, grid.z);
    printf("block : (%d, %d, %d)\n", block.x, block.y, block.z);


    // Launch kernel
    vectorAdd<<<grid, block>>>(d_a, d_b, d_c);

    // Copy data back to host
    cudaMemcpy(c, d_c, memSize, cudaMemcpyDeviceToHost);

    // Check the result
    bool result = true;
    for (int i = 0; i < NUM_DATA; i++) {
        if (c[i] != a[i] + b[i]) {
            result = false;
            break;
        }
    }
    if(result) {
        printf("Test PASSED\n");
    } else {
        printf("Test FAILED\n");
    }

    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);   

    delete[] a; delete[] b; delete[] c; 


    return 0;
}

 

코드에서 NUM_DATA 가 10240으로 설정이 되어있습니다. 

그리고 단순하게 dim3를 써서 블록의 형태를 (10240, 1, 1), 그리드의 형태를 (1, 1, 1) 로 설정을 해두었습니다.

다만, 한 블록이 갖는 스레드의 수가 1,024를 넘지 못하기 때문에, 블록을 여러 개 사용해야 합니다.

즉, 그리드의 형태가 바뀌어야 하죠. 

NUM_DATA의 수가 1024가 넘으면, NUM_DATA를 1024로 나눈 몫 만큼 gird.x 에 저장해두어야 하고, 이는 블록의 개수를 의미합니다. 

만약 입력된 NUM_DATA 가 1024로 나누어 떨어지지 않을 땐, 블록이 하나 더 필요하기 때문에 코드 사에 grid.x ++ 로 처리를 해주었습니다. 

 

위의 전체 코드를 실행시키면 아래 결과가 출력이 됩니다. 

총 10240개의 요소가 입력되었기 때문에, int형 4byte만큼 할당을 해주었고, 

grid의 형태(크기) : (10, 1, 1)

block 의 형태(크기) : (1024, 1, 1) 로 잘 셋팅이 된 걸 볼 수 있습니다. 

 

 

실행 결과