행렬 곱에 기반한 grad_input, grad_weight, shared memory 데이터 문제, tiling 알고리즘 사용 중에서 발생한 오류

CUDA 에서의 행렬 곱 연산의 경우,

메모리 접근 패턴과 캐시 효율성을 위해 transpose 된 행렬을 미리 만들어 사용하여 coalesced memory access 를 이용, shared memory tiling 을 효율적을 사용했음

 

그런데 CUDA 에서 제공하는 라이버르리 기반의 전치를 사용했을 때 전치 기반 행렬 곱 연산의 오류가 발생했음 ( 잘못된 메모리 접근으로 인한 오류 값 저장)

 

때문에 전치의 직접 구현이 필요했음

#include <cuda_runtime.h>
#include <iostream>
#include "transpose.cuh"

#define TILE_WIDTH 16

__global__ void transpose_kernel(const float* __restrict__ input, float* __restrict__ output, int rows, int cols) {
    int row = blockIdx.y * TILE_WIDTH + threadIdx.y;
    int col = blockIdx.x * TILE_WIDTH + threadIdx.x;

    if (row < rows && col < cols) {
        float val = input[row * cols + col];
        output[col * rows + row] = val;
    }
}

void launch_transpose(const float* input, float* output, int rows, int cols) {
    dim3 blockDim(TILE_WIDTH, TILE_WIDTH);
    dim3 gridDim((cols + TILE_WIDTH - 1) / TILE_WIDTH, (rows + TILE_WIDTH - 1) / TILE_WIDTH);
    transpose_kernel<<<gridDim, blockDim>>>(input, output, rows, cols);
    cudaDeviceSynchronize();
}