CUDA/문법

#!if 문서명2 != null
, [[]]

#!if 문서명3 != null
, [[]]

#!if 문서명4 != null
, [[]]

#!if 문서명5 != null
, [[]]

#!if 문서명6 != null
, [[]]

1. 개요

2. 편집 지침

#!syntax cpp 
__global__ void vecAdd(float* A, float* B, float* C)
{
   int workIndex = threadIdx.x + blockDim.x * blockIdx.x

   C[workIndex] = A[workIndex] + B[workIndex];
}

int main()
{
    ...
    vecAdd<<<4, 256>>>(A, B, C);
    ...
}

3. 시작하기

3.1. 프로그래밍 모델

3.2. 커널

#!syntax cpp
// 커널 정의
__global__ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x;

    C[i] = A[i] + B[i];
}

int main()
{
    ...
    // N개의 스레드를 사용한 커널 호출
    VecAdd<<<1, N>>>(A, B, C);
    ...
}

#!syntax cpp
__global__ void VecAdd(float* A, float* B, float* C, int N)
{
    int i = threadIdx.x + blockIdx.x * blockDim.x;

    if (i < N)
    {
        C[i] = A[i] + B[i];
    }
}

int main()
{
    ...
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    VecAdd<<<blocksPerGrid, threadsPerBlock>>>(A, B, C, N);
    ...
}

3.3. 메모리 모델

#!syntax cpp
float* d_A;

cudaMalloc(&d_A, N * sizeof(float));
cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice);

Kernel<<<grid, block>>>(d_A, N);

cudaMemcpy(h_A, d_A, N * sizeof(float), cudaMemcpyDeviceToHost);
cudaFree(d_A);

4. 기본 문법

4.1. 실행 지정자

#!syntax cpp
__host__ void HostFunc()
{
    // CPU에서 실행
}

__device__ float DeviceFunc(float x)
{
    // GPU에서 실행
    return x * 2.0f;
}

__global__ void Kernel(float* data)
{
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    data[i] = DeviceFunc(data[i]);
}

4.2. 메모리 지정자

#!syntax cpp
__device__ int deviceValue;

__constant__ float coeff[4];

__global__ void Kernel(float* data)
{
    __shared__ float cache[256];

    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    cache[tid] = data[i];

    __syncthreads();

    data[i] = cache[tid] * coeff[0] + deviceValue;
}

#!syntax cpp
__constant__ float coeff[4];

int main()
{
    float h_coeff[4] = {1.0f, 2.0f, 3.0f, 4.0f};

    cudaMemcpyToSymbol(coeff, h_coeff, sizeof(h_coeff));
}

4.3. 메모리 할당

#!syntax cpp
float* d_A = nullptr;
float* d_B = nullptr;
float* d_C = nullptr;

size_t size = N * sizeof(float);

cudaMalloc(&d_A, size);
cudaMalloc(&d_B, size);
cudaMalloc(&d_C, size);

cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);

4.4. 내장 자료형

#!syntax cpp
dim3 block(16, 16);
dim3 grid(
    (width + block.x - 1) / block.x,
    (height + block.y - 1) / block.y
);

Kernel2D<<<grid, block>>>(data, width, height);

#!syntax cpp
Kernel<<<16, 256>>>();

// 개념적으로는 다음과 유사하다.
dim3 grid(16, 1, 1);
dim3 block(256, 1, 1);
Kernel<<<grid, block>>>();

#!syntax cpp
float3 pos = make_float3(1.0f, 2.0f, 3.0f);

float x = pos.x;
float y = pos.y;
float z = pos.z;

4.5. 내장 변수

#!syntax cpp
int i = threadIdx.x + blockIdx.x * blockDim.x;

#!syntax cpp
__global__ void Kernel2D(float* data, int width, int height)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;

    if (x < width && y < height)
    {
        int i = x + y * width;
        data[i] = data[i] * 2.0f;
    }
}

4.6. 동기화

#!syntax cpp
Kernel<<<grid, block>>>(data);

cudaDeviceSynchronize();

#!syntax cpp
__global__ void SharedExample(float* data)
{
    __shared__ float cache[256];

    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    cache[tid] = data[i];

    __syncthreads();

    data[i] = cache[tid] * 2.0f;
}

4.7. 오류 처리

5. 문법 상세

5.1. 워프 함수

5.1.1. __shfl_sync()

#!syntax cpp
T __shfl_sync(unsigned mask, T var, int srcLane, int width = warpSize);

#!syntax cpp
__global__ void ShuffleExample(int* output)
{
    int lane = threadIdx.x % warpSize;
    int value = threadIdx.x;

    int fromLane0 = __shfl_sync(0xffffffff, value, 0);

    output[threadIdx.x] = fromLane0 + lane;
}

#!syntax cpp
__global__ void ShufflePermutationExample(int* output)
{
    int lane = threadIdx.x % warpSize;
    int value = lane;

    int srcLane = (lane + 1) % warpSize;
    int fromNextLane = __shfl_sync(0xffffffff, value, srcLane);

    output[threadIdx.x] = fromNextLane;
}

5.1.2. __shfl_down_sync()

#!syntax cpp
T __shfl_down_sync(unsigned mask, T var, unsigned int delta, int width = warpSize);

#!syntax cpp
__global__ void WarpReduceSum(const float* input, float* output)
{
    int lane = threadIdx.x % warpSize;
    int warpId = threadIdx.x / warpSize;
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    float value = input[i];

    for (int delta = warpSize / 2; delta > 0; delta /= 2)
    {
        value += __shfl_down_sync(0xffffffff, value, delta);
    }

    if (lane == 0)
    {
        output[blockIdx.x * (blockDim.x / warpSize) + warpId] = value;
    }
}

5.1.3. __shfl_up_sync()

#!syntax cpp
T __shfl_up_sync(unsigned mask, T var, unsigned int delta, int width = warpSize);

#!syntax cpp
__global__ void WarpPrefixSum(const int* input, int* output)
{
    int lane = threadIdx.x % warpSize;
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    int value = input[i];

    for (int delta = 1; delta < warpSize; delta *= 2)
    {
        int other = __shfl_up_sync(0xffffffff, value, delta);

        if (lane >= delta)
        {
            value += other;
        }
    }

    output[i] = value;
}

5.1.4. __shfl_xor_sync()

#!syntax cpp
T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width = warpSize);

5.2. 텐서 코어

6. 주의할 점

7. 최적화

8. 예제