CUDA에서 벡터 덧셈을 구현할 때, 데이터 개수와 할당된 스레드 수의 관계에 따라 접근 방식이 달라진다. 스레드가 충분한 경우와 부족한 경우를 각각 살펴다.

스레드가 충분할 때: 1:1 매핑

데이터 개수만큼 블록을 할당하면 각 요소를 독립적인 스레드가 처리한다.

#include <iostream>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#define TOTAL_ELEMS 10

__global__ void vecAddKernel(const int* __restrict__ srcA,
                             const int* __restrict__ srcB,
                             int* __restrict__ dst)
{
    int idx = blockIdx.x;
    dst[idx] = srcA[idx] + srcB[idx];
}

int main()
{
    int hostA[TOTAL_ELEMS], hostB[TOTAL_ELEMS], hostC[TOTAL_ELEMS];
    int *gpuA, *gpuB, *gpuC;

    cudaMalloc((void**)&gpuA, TOTAL_ELEMS * sizeof(int));
    cudaMalloc((void**)&gpuB, TOTAL_ELEMS * sizeof(int));
    cudaMalloc((void**)&gpuC, TOTAL_ELEMS * sizeof(int));

    for (int k = 0; k < TOTAL_ELEMS; k++) {
        hostA[k] = k;
        hostB[k] = k + 1;
    }

    cudaMemcpy(gpuA, hostA, TOTAL_ELEMS * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(gpuB, hostB, TOTAL_ELEMS * sizeof(int), cudaMemcpyHostToDevice);

    vecAddKernel<<<TOTAL_ELEMS, 1>>>(gpuA, gpuB, gpuC);

    cudaMemcpy(hostC, gpuC, TOTAL_ELEMS * sizeof(int), cudaMemcpyDeviceToHost);

    for (int k = 0; k < TOTAL_ELEMS; k++) {
        std::cout << hostC[k] << " ";
    }

    cudaFree(gpuA);
    cudaFree(gpuB);
    cudaFree(gpuC);

    return 0;
}

위 예시에서는 TOTAL_ELEMS개의 블록에 각각 1개의 스레드를 배치하여, 전체 요소를 동시에 처리한다.

스레드가 부족할 때: 순환 처리(Strided Loop)

32개의 요소를 더해야 하는데 GPU에 2개 블록 × 4개 스레드 = 8개 스레드만 할당된 상황을 가정한다. 8개 스레드가 32개 데이터를 번갈아가며 처리해야 한다.

각 스레드는 자신의 고유 인덱스에서 시작해, 전체 스레드 수(gridDim.x × blockDim.x)만큼 건너며 데이터를 처리한다.

#include <iostream>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#define DATA_COUNT 32

__global__ void stridedVecAdd(const int* __restrict__ inX,
                              const int* __restrict__ inY,
                              int* __restrict__ outZ)
{
    int pos = blockIdx.x * blockDim.x + threadIdx.x;
    int leap = gridDim.x * blockDim.x;

    for (; pos < DATA_COUNT; pos += leap) {
        outZ[pos] = inX[pos] + inY[pos];
    }
}

int main(int argc, char** argv)
{
    int cpuBufX[DATA_COUNT], cpuBufY[DATA_COUNT], cpuBufZ[DATA_COUNT];
    int *gpuBufX, *gpuBufY, *gpuBufZ;

    cudaMalloc((void**)&gpuBufX, DATA_COUNT * sizeof(int));
    cudaMalloc((void**)&gpuBufY, DATA_COUNT * sizeof(int));
    cudaMalloc((void**)&gpuBufZ, DATA_COUNT * sizeof(int));

    for (int j = 0; j < DATA_COUNT; j++) {
        cpuBufX[j] = j;
        cpuBufY[j] = j;
    }

    cudaMemcpy(gpuBufX, cpuBufX, DATA_COUNT * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(gpuBufY, cpuBufY, DATA_COUNT * sizeof(int), cudaMemcpyHostToDevice);

    stridedVecAdd<<<2, 4>>>(gpuBufX, gpuBufY, gpuBufZ);

    cudaMemcpy(cpuBufZ, gpuBufZ, DATA_COUNT * sizeof(int), cudaMemcpyDeviceToHost);

    for (int j = 0; j < DATA_COUNT; j++) {
        std::cout << cpuBufZ[j] << " ";
    }

    cudaFree(gpuBufX);
    cudaFree(gpuBufY);
    cudaFree(gpuBufZ);

    return 0;
}

0번 스레드는 0, 8, 16, 24번 인덱스를, 1번 스레드는 1, 9, 17, 25번 인덱스를 담당한다. 이 패턴은 leap 간격의 등차수열 형태로 전개되며, 제한된 스레드 자원으로 대용량 데이터를 처리할 수 있게 한다.