Matrix add & 2-D 인덱싱

새 습관 둘: index 매핑이랑 2-D launch geometry

Vector add는 1-D grid launch. Matrix add는 데이터가 2-D니까 2-D grid launch. Launch shape이 데이터 shape 미러링, 커널 안 인덱싱이 thread 좌표를 데이터 좌표에 매핑.

매핑은 두 부분:

Thread 좌표 → 데이터 좌표: row = blockIdx.y * blockDim.y + threadIdx.y, col = blockIdx.x * blockDim.x + threadIdx.x.
데이터 좌표 → flat 메모리 offset: row-major면 idx = row * ld + col; column-major는 반대.

Block size 선택: 16×16 = 256 thread가 robust default. Warp size (32)의 배수, latency hiding용 ~8 warp 줘, occupancy vs register 압박 균형. 32×32 = 1024도 동작하는데 무거운 커널에선 register 한계 hit 가능.

Code

Matrix add — 2-D launch + 2-D 인덱싱 (CUDA)·cuda

__global__ void mat_add(const float* A, const float* B, float* C,
                        int rows, int cols, int ld) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < rows && col < cols) {
        size_t idx = (size_t)row * ld + col;     // row-major
        C[idx] = A[idx] + B[idx];
    }
}

// Launch:
//   dim3 block(16, 16);                      // 16×16 = 256 thread
//   dim3 grid((cols + 15)/16, (rows + 15)/16);
//   mat_add<<<grid, block>>>(dA, dB, dC, rows, cols, ld);

Metal에서 같은 커널·metal

#include <metal_stdlib>
using namespace metal;

kernel void mat_add(const device float *A [[buffer(0)]],
                    const device float *B [[buffer(1)]],
                    device       float *C [[buffer(2)]],
                    constant     uint2 &dims [[buffer(3)]],   // (rows, cols)
                    constant     uint  &ld   [[buffer(4)]],
                    uint2 gid [[thread_position_in_grid]])
{
    uint row = gid.y, col = gid.x;
    if (row < dims.x && col < dims.y) {
        uint idx = row * ld + col;
        C[idx] = A[idx] + B[idx];
    }
}

// Host: encoder.dispatchThreads(
//   MTLSize(width: cols, height: rows, depth: 1),
//   threadsPerThreadgroup: MTLSize(width: 16, height: 16, depth: 1))

Exercise

위 matrix add 커널 4096×4096 행렬에 작성하고 결과를 numpy(A + B)랑 검증. 그 다음 일부러 index swap: idx = col * ld + row (row-major로 read하면서 buffer를 column-major처럼 취급). 대부분 위치에서 잘못된 답 — layout이 가정 아니라 존중되어야 한다는 거 느끼는 데 유용.

새 습관 둘: index 매핑이랑 2-D launch geometry

Code

External links

Exercise

Progress

댓글 0