Efficient GEMM

Introduction

Matrix Multiplication은 다양한 과학 분야에서 널리 사용된다. Matrix Multiplication은 본질적으로 단순한 산술 연산의 집합으로 구성되어 있기 때문에, 상대적으로 더 많은 수의 코어(core)를 갖춘 GPU를 이용하여 효율적으로 수행된다. 그러나 GPU를 활용하여 Matrix Multiplication을 수행할 때, 빈번하게 발생하는 Memory Hierarchy 간 데이터 이동 문제가 성능 저하의 주요 원인이 된다.

일반적으로 Matrix Multiplication은 다음과 같이 정의된다.

\[c_{ij} = \sum_{k} a_{ik} \times b_{kj}\]

한 변의 길이가 \(O(N)\)인 행렬의 계산(computation) 복잡도는 \(O(N^3)\)이며, 계산을 수행하기 위해 GPU의 core에 로드되어야 하는 데이터의 크기는 \(O(N^2)\)이다. 데이터를 최소한으로 로드하면서 효율성을 극대화하기 위해서는 데이터가 한 번 로드될 때마다 \(O(N)\)의 연산(computation)을 수행해야 한다. 그러나 실제로 \(O(N^2)\) 크기의 데이터는 일반적으로 매우 크기 때문에 모든 데이터를 한 번에 캐시(cache)에 로드하는 것이 불가능하다. 따라서 연산량은 동일하지만, 데이터를 여러 번 반복적으로 로드해야 하므로, 이러한 데이터의 빈번한 로드 과정이 병목 현상(bottleneck)의 주요 원인이 된다.

모든 \(N\)에 대해 앞서 정의한 일반적인 형태인 \(c_{ij} = \sum_{k} a_{ik} \times b_{kj}\)의 계산 방식을 그대로 수행한다면, 빈번한 데이터 로드 문제는 해결할 수 없다. 이 문제를 해결하기 위해, Matrix Multiplication의 정의를 수학적으로 동일하지만 다르게 해석하여 계산한다. 즉, Matrix Multiplication을 Block Matrix Multiplication의 형태로 재정의하여 계산하는 것이다. Block Matrix Multiplication을 활용하여 Matrix Multiplication을 수행하면 데이터 로드 횟수를 줄일 수 있다.

더 나아가, 각 Block 내에서도 Block Matrix Multiplication을 다시 수행하도록 분해하여 구현할 수 있다. 추가적으로, Block Multiplication 내에서도 Outer Product 형태로 정의하여 구현함으로써, 메모리 접근 효율성을 더욱 향상할 수 있다.

BackGround

데이터의 반복적인 로드(load)를 피하기 위해 반드시 이해해야 할 중요한 개념들이 있습니다.
그 중 하나가 GPU의 Memory Hierarchy입니다.
GPU의 Memory Hierarchy는 데이터가 어떻게 로드되고, 연산이 어떤 순서로 이루어지는지를 이해하는 데 매우 유용한 정보를 제공합니다.
이를 통해 Matrix Multiplication을 보다 미세하게(fine-grained) 최적화하여 구현할 수 있습니다.

Arithmetic Intensity는 데이터가 전송되는 속도에 비해 얼마나 많은 연산이 수행되는지를 나타내는 지표입니다.
따라서, Arithmetic Intensity가 높을수록 동일한 데이터 전송량으로 더 많은 연산을 수행할 수 있기 때문에 연산 효율이 높아집니다.
이로 인해 Arithmetic Intensity를 높이는 것이 왜 중요한지, 그 이점을 깊이 이해할 필요가 있습니다.

또한, GPU는 Data Parallelism(데이터 병렬성)을 적용하여 하나의 명령어를 여러 thread가 동시에 실행합니다.
이러한 GPU의 특성은 동일한 데이터에 대한 반복적인 로드를 줄이고, 데이터 로드와 연산을 최적화하는 데 매우 중요한 역할을 합니다.
따라서, GPU의 Data Parallelism이 데이터 로드 최적화와 어떤 관계가 있는지 이해하는 것이 중요합니다.

그리고 Data Locality와 Memory Layout은 Processing Unit의 하드웨어적 특성에 해당합니다.
이러한 하드웨어적 특성을 잘 이해하고 적절히 활용할 경우, 데이터 접근 효율을 높이고 데이터 로드를 최적화하는 데 큰 도움이 됩니다.
따라서 연산 성능을 극대화하기 위해서는 Data Locality와 Memory Layout까지 고려한 접근 방식이 필요합니다.

또한, Matrix Multiplication의 정의(Definition)를 다양한 관점에서 해석할 수 있어야 합니다.
기존의 전통적인 방식으로 Matrix Multiplication을 수행하게 되면, 필연적으로 빈번한 데이터 로드가 발생하게 되어 메모리 접근 병목이 발생할 수 있습니다.
따라서 Matrix Multiplication을 효율적으로 구현하기 위해서는 그 정의를 여러 방식으로 재구성하고 전개할 수 있어야 하며,
이를 통해 중복 데이터 로드를 최소화하는 전략을 세울 수 있습니다.

Memory Hierarchy

GPU의 Memory Hierarchy는 DRAM, L2 Cache, L1 Cache/Shared Memory, Register로 구분됩니다.
GPU가 연산을 수행하기 위해 데이터에 접근할 때, 일반적으로 DRAM에서 데이터를 가져와 L2 Cache, L1 Cache/Shared Memory를 거쳐 Register에 저장한 후 연산을 수행합니다.
또한, **L1 Cache/Shared Memory로 직접 데이터를 로드하는 경우에는 L2 Cache를 거치지 않고 곧바로 L1 Cache/Shared Memory로 전달된 후 Register로 이동합니다.

GPU는 Streaming Multiprocessor (SM) 라는 연산 유닛을 포함하고 있으며, 프로그램이 실행될 때 SM 내에서 PTX라는 Instruction Set Architecture (ISA) 가 동작합니다. PTX 명령어는 opcode와 operand로 구성되며, opcode는 사전에 정의된 명령어 코드들이고, operand는 주로 Register에 저장된 값들을 사용합니다. 따라서, GPU가 실제로 연산을 수행하려면 필요한 데이터가 DRAM으로부터 최종적으로 Register까지 전송되어야 합니다.

GPU에서 Register는 가장 빠른 저장 장치이며, Register에 물리적으로 가까운 저장 장치일수록 데이터 전송 속도(bandwidth)가 빠릅니다.
최근의 최신 GPU의 경우 DRAM 용량은 최대 80GB, L2 Cache는 약 50MB, L1 Cache/Shared Memory는 최대 약 33MB 수준입니다.

특히 딥러닝과 같은 대규모 데이터 연산에서는 행렬(Matrix) 전체를 캐시에 로드할 수 없는 경우가 빈번히 발생하기 때문에, 여러 번 DRAM에 접근하여 필요한 데이터를 L2, L1/Shared Memory, Register로 가져오는 작업이 반복됩니다.
하지만 DRAM의 접근 속도와 대역폭은 L2 Cache, L1 Cache/Shared Memory, Register에 비해 매우 느리기 때문에, DRAM에서 Register로의 데이터 전송을 최소화하는 것이 GPU 연산 최적화의 핵심입니다.

Arthimetic Intensity

Register로의 데이터 전송을 최적화하게 되면 Arithmetic Intensity가 증가하게 됩니다.
Arithmetic Intensity는 다음과 같이 정의됩니다.

\[\text{Arithmetic Intensity} = \frac{\text{Arithmetic Operation Work}}{\text{Transferred Bytes per second}}\]

GPU에서 연산을 최적화하기 위해서는 DRAM에서 Register로 전송되는 데이터의 양을 최소화하는 것이 중요합니다.
이는 앞서 언급한 것처럼, Matrix Multiplication과 같은 연산을 고려할 때 전체 계산량(Arithmetic Operation Work)은 고정되어 있기 때문입니다.

특히, 행렬(Matrix)의 각 요소(element)가 연산에 사용되는 횟수는 결정되어 있기 때문에, 동일한 데이터를 DRAM에서 여러 번 가져오는 것을 줄이면 데이터 전송 시간이 감소하고, 결과적으로 Arithmetic Intensity가 증가하게 됩니다.

이러한 이유로 Arithmetic Intensity는 딥러닝 연산이 얼마나 효율적으로 수행되는지를 나타내는 중요한 지표로 사용됩니다.
또한, Arithmetic Intensity가 높다는 것은 동일한 데이터를 메모리로부터 로드했을 때 수행되는 연산의 양이 많다는 의미이며, 이는 메모리 대역폭 대비 계산량이 많은 구조로, 메모리 병목을 줄이고 GPU의 계산 성능을 극대화할 수 있음을 의미합니다.

Data Parallelism

GPU에서 프로그램을 실행할 때는 데이터를 분산하여 여러 개의 스레드(thread)가 동시에 같은 명령어(instruction)를 실행하도록 합니다.
이와 같은 병렬 실행 방식을 위해 GPU는 하나의 명령어를 여러 데이터에 동시에 적용하는 구조를 사용합니다.

NVIDIA GPU에서는 이를 SIMT (Single Instruction, Multiple Threads) 방식이라고 부르며,
AMD GPU에서는 SIMD (Single Instruction, Multiple Data) 방식이라는 용어를 사용합니다.

구체적으로, NVIDIA GPU에서는 하나의 명령어를 동시에 실행하는 스레드 묶음을 “warp”라고 부르며,
각 warp는 32개의 thread로 구성되어 있습니다.
반면, AMD GPU에서는 “wavefront”라는 단위를 사용하며,
각 wavefront는 64개의 thread로 구성됩니다.

특히, 여러 개의 thread가 동시에 load instruction을 수행할 때,
warp (NVIDIA)나 wavefront (AMD) 내에서 연속적인 데이터에 접근하게 되면,
여러 개의 데이터를 하나의 load instruction으로 여러 thread가 동시에 가져올 수 있습니다.

이러한 방식을 “coalesced” 되었다라고 표현하며,
메모리 접근이 병합되어(coalesced) 하나의 메모리 명령으로 처리되기 때문에, 메모리 instruction의 실행 횟수를 줄일 수 있습니다.
결과적으로, 데이터 로드 속도가 빨라지고 메모리 대역폭 활용이 최적화됩니다.

따라서, GPU에서 데이터 접근 패턴을 설계할 때, 가능한 한 warp나 wavefront 내의 thread들이 연속적인 메모리 주소에 접근하도록 하는 것이 매우 중요합니다.

NVIDIA와 AMD 모두 다수의 스레드가 동일한 명령어를 실행하는 방식이지만, 스레드 묶음의 크기에는 차이가 있으며,
이러한 차이는 병렬 처리 방식과 성능 최적화 전략에 중요한 영향을 미칩니다.

Data Locality & Memory Layout

거의 모든 Processing Unit(연산 장치) 에는 Cache Memory가 존재합니다.
Cache Memory는 데이터를 Register에서 사용하기 전에 메모리(RAM 또는 DRAM)로부터 데이터를 가져와 임시로 저장하는 공간입니다.
이때 Cache는 하나의 데이터를 로드할 때 해당 데이터의 주소와 인접한 데이터들도 함께 로드(prefetch) 하여 메모리 접근 횟수를 줄이는 역할을 합니다.
이는 대부분의 Processing Unit이 다음 연산에서 이전에 사용한 데이터 근처의 데이터를 사용할 확률이 높다는 경향성을 가지기 때문입니다.

따라서, Processing Unit을 설계할 때 이러한 경향성을 고려하여 주소상 인접한 데이터를 미리 Cache로 로드하도록 되어 있습니다.
그리고 이렇게 Cache에 로드된 데이터들 중에서 실제로 얼마나 연속적인 데이터가 사용되는가를 “Data Locality” 라고 합니다.
결과적으로, 연산을 수행할 때 가능한 한 연속적인 데이터를 접근하는 것이 Data Locality를 높이는 방법이 됩니다.

특히, 대부분의 Processing Unit에서는 행렬(Matrix)을 메모리에 저장할 때 행(row) 방향으로 연속된 형태로 저장합니다.
따라서 Matrix의 원소(element)를 접근할 때도 가능한 한 row 방향으로 접근하는 것이 Data Locality를 높이고, Cache 효율성을 극대화할 수 있는 좋은 방식입니다.

Matrix Mutliplication Definition

Matrix Multiplication은 다음과 같이 정의됩니다.

\[c_{ij} = \sum_{k} a_{ik} \times b_{kj}\]

하지만, 이러한 정의를 사용하면 Matrix Multiplication은 중첩된 for loop으로 구현됩니다.

for k K
  for i  M
    for j N
      c[i][j] += A[i][k] *B[k][j]

이때 C의 한 개의 원소를 계산하기 위해서는 A의 하나의 row와 B의 하나의 column 전체를 load해야 합니다.
따라서, 일반적인 Matrix Multiplication 정의를 그대로 사용하게 되면, A와 B 행렬의 각 원소를 C[i][j]를 계산하는 순간마다 메모리에서 가져와서 Register로 로드한 후 연산을 수행하게 됩니다.

결과적으로, 하나의 C[i][j]를 계산하기 위해 A와 B의 데이터를 여러 번 반복해서 load하게 되며, A와 B의 각 원소가 O(N)번씩 메모리로부터 로드되는 비효율이 발생합니다.

Block Multiplication

Matrix는 여러 개의 작은 Matrix Block으로 나누어 표현될 수 있습니다.
그리고 모든 Matrix Multiplication은 이러한 Matrix Block들의 곱으로 표현할 수 있습니다.
즉, 큰 행렬의 곱셈도 작은 블록 단위로 나누어 계산할 수 있습니다.

예를 들어, 다음과 같이 두 행렬을 Block 형태로 나눌 수 있습니다.

\[\begin{bmatrix} A_{11} & A_{12} \\ A_{21} & A_{22} \end{bmatrix} \begin{bmatrix} B_{11} \\ B_{21} \end{bmatrix} = \begin{bmatrix} A_{11} B_{11} + A_{12} B_{21} \\ A_{21} B_{11} + A_{22} B_{21} \end{bmatrix}\]

따라서, Matrix Multiplication은 Block 단위의 곱셈과 덧셈으로 분할하여 표현할 수 있으며, 이는 대규모 행렬 연산의 효율성을 높이기 위해 널리 사용되는 방식입니다.

Outer Product

Outer Product란 어떤 벡터 $v \ (m \times 1)$, $u \ (n \times 1)$ 가 있을 때,
$v$와 $u^T$의 외적(Outer Product)을 통해 $(m \times n)$ 크기의 행렬을 만들어내는 계산 방법입니다.

또한, 어떤 Matrix Multiplication도 여러 개의 Outer Product의 합으로 표현할 수 있습니다.
즉, 다음과 같이 나타낼 수 있습니다.

\[AB = \sum_k (a_{\text{col}_k} \otimes b_{\text{row}_k})\]

여기서 $a_{\text{col}k}$는 A의 열(column), $b{\text{row}_k}$는 B의 행(row)을 의미하며, $\otimes$는 Outer Product를 의미합니다.

이 방법의 장점은 각 행렬의 원소(element)에 대해 외적(Outer Product)을 한 번 수행한 이후에는 해당 element가 다시 필요 없다는 점입니다.
즉, 한 번 로드한 데이터를 재사용하지 않고 곧바로 모든 연산에 사용할 수 있어 메모리 접근을 최소화할 수 있습니다.

구체적으로, Register에 1차원 배열로 row 또는 column을 로드한 후, 한 번의 Outer Product 연산을 수행하면 이미 Register에 로드된 element들을 다시 메모리에서 로드할 필요가 없습니다.
따라서, 효율적인 데이터 사용과 최소한의 메모리 접근이 가능해지며, 이는 특히 GPU에서 연산 최적화를 위해 매우 중요한 기법입니다.

Data coallescing

Processor Unit에서 register에 load되는 data에서 가까운 data를 cache에 load하는 것이 구현되어 있습니다.따라서, 아래와 같이 load를 하게 되면 A를 row방향으로 load하게 되어 data locality가 높아지게 됩니다.

__global__ void sgemm_naive(int M, int N, int K, float alpha, const float *A,
                            const float *B, float beta, float *C) {
  // compute position in C that this thread is responsible for
  const uint x = blockIdx.x * blockDim.x + threadIdx.x;
  const uint y = blockIdx.y * blockDim.y + threadIdx.y;

  // `if` condition is necessary for when M or N aren't multiples of 32.
  if (x < M && y < N) {
    float tmp = 0.0;
    for (int i = 0; i < K; ++i) {
      tmp += A[x * K + i] * B[i * N + y];
    }
    // C = α*(A@B)+β*C
    C[x * N + y] = alpha * tmp + beta * C[x * N + y];
  }
}

각 thread는 행렬(Matrix)의 하나의 element를 계산합니다.
특히 행렬 A에 대해서는 각 thread가 각 row를 담당하기 때문에,
연속된 thread 그룹이 row의 element들을 순차적으로 load하게 됩니다.

하지만 이러한 방식에서는 warp가 instruction을 수행하기 위해 데이터를 load할 때,
load되는 element들이 메모리 상에서 연속적(consecutive)이지 않기 때문에
데이터 로드가 coalescing되지 못하는 문제가 발생합니다.
즉, 메모리 접근이 비효율적으로 이루어져 메모리 대역폭을 충분히 활용하지 못하는 상황이 됩니다.

반면, 행렬 B에 대해서는 모든 thread가 동일한 element를 사용하기 때문에,
각 thread가 B의 동일한 element를 공유(broadcasting) 하여 사용합니다.
이로 인해 B의 데이터는 여러 thread가 같은 값을 참조하게 되어 메모리 로드 측면에서 효율적입니다. .

따라서,아래와 같이 access pattern을 변경해줍니다.

const int x = blockIdx.x * BLOCKSIZE + (threadIdx.x / BLOCKSIZE);
const int y = blockIdx.y * BLOCKSIZE + (threadIdx.x % BLOCKSIZE);

if (x < M && y < N) {
  float tmp = 0.0;
  for (int i = 0; i < K; ++i) {
    tmp += A[x * K + i] * B[i * N + y];
  }
  C[x * N + y] = alpha * tmp + beta * C[x * N + y];
}

각 thread가 행렬 A로부터 데이터를 사용할 때,
모든 thread가 같은 데이터를 broadcasting 방식으로 공유하여 사용합니다.
즉, 각 thread가 동일한 A의 원소를 참조하면서 연산을 수행합니다.

반면, 행렬 B의 경우,
각 thread는 매 instruction마다 B의 column에 해당하는 데이터를 load하여 사용합니다.
이때 연속된 thread 그룹이 메모리 상에서 연속된 column 데이터를 순차적으로 load하기 때문에,
데이터 로드가 coalescing되어 효율적인 메모리 접근이 가능해집니다.
결과적으로, B의 데이터를 로드할 때는 메모리 대역폭을 최적화할 수 있는 장점이 있습니다.

Tiling

Tiling은 Matrix Multiplication의 정의를 Block Multiplication으로 재정의하여 구현하는 방식입니다.
아래와 같이 일반적인 Matrix Multiplication은 중첩된 for loop을 통해 계산됩니다.

일반적인 Matrix Multiplication 방식에서는 A, B, C의 각 포인터로부터 원소(element) 하나씩을 Register에 로드하여 연산합니다.
이때 GPU에서 L1 Cache 또는 Shared Memory에 A, B, C 전체를 한 번에 로드해서 사용하면 좋겠다고 생각할 수 있지만, 현실적으로는 A, B, C 전체를 캐시에 올릴 수 없습니다.

예를 들어, (4000, 4000) 크기의 행렬이 fp16 데이터 타입일 경우, L1 Cache나 Shared Memory의 크기를 초과하게 됩니다.
특히, 딥러닝에서는 행렬 크기가 (4000, 4000) 을 훨씬 넘어서는 경우가 많습니다.
또한, Training 시에는 weight에 대한 gradient까지 저장해야 하므로, weight와 동일한 크기의 추가 메모리가 필요하게 되어 사용 메모리는 더 증가합니다.

따라서, 이러한 이유로 인해 반복적인 데이터 로드를 피하기 위해서 기존의 단순한 Matrix Multiplication 정의를 그대로 사용할 수 없으며,
이를 해결하기 위한 접근 방식이 바로 Tiling 기법을 통해 Block 단위로 데이터를 나누어 처리하는 방식입니다.

Block Tiling

Matrix 전체를 한 번에 Cache에 로드하기는 현실적으로 어렵고,
필요할 때마다 개별 element를 매번 로드하는 방식은 매우 비효율적입니다.
따라서, Matrix의 전체가 아닌 일부(Block)를 Cache에 로드한 후, 해당 부분을 이용해 계산을 수행하는 방식이 사용됩니다.

이처럼 Matrix의 일부를 나누어 로드하고 계산하는 방식은 결국 Matrix Multiplication을 Block Multiplication으로 해석한 것이라 할 수 있습니다.
즉, 큰 행렬 연산을 작은 블록 단위로 나누어 처리함으로써, 데이터 로드와 메모리 사용을 최적화하는 기법입니다.

Output Matrix에서 Block을 설정하여 GPU의 각 코어(core)들이 해당 Block을 병렬적으로 계산하게 됩니다.
이렇게 Matrix를 Block 단위로 나누어 처리하는 Block Multiplication 방식은
상대적으로 작은 크기의 Matrix 연산으로 분할되기 때문에, 이 블록 데이터를 Cache Memory에 로드하여 사용할 수 있는 장점이 있습니다.

따라서, GPU Programming에서는 이러한 Block 단위 데이터를 직접 Shared Memory에 로드할 수 있도록 명시적인 구문(syntax)을 제공합니다.
이를 통해 데이터를 미리 Shared Memory에 저장해두고, 반복적인 Global Memory 접근을 줄이는 최적화가 가능해집니다.

이를 바탕으로 한 Pseudocode는 다음과 같습니다.

for bm M
  for bn N
    for bk K
      A_block(shared memory) = A block
      B_block(shared memory) =  B block
      C_block(shared memory) = C block
      for i bm
        for j bn
          c_block[i][j] += A_block[i][k] *B_block[k][j]

Block Multiplication에서는 Output Matrix의 각 Block을 계산하기 위해
행렬 A의 Block은 $N // b_n$ 번, 행렬 B의 Block은 $M // b_m$ 번 로드(load)하게 됩니다.

따라서, 각 element들이 A에서는 $N // b_n$ 번, B에서는 $M // b_m$ 번씩 DRAM으로부터 Shared Memory로 로드됩니다.
이는 기존 방식처럼 각 element를 연산마다 매번 DRAM에서 로드하는 것(O(N)번) 과는 다르게,
Block 단위로 묶어서 가져오기 때문에 각 데이터가 훨씬 적은 횟수로 DRAM에서 로드되는 것입니다.즉, Arthimetic Intensity 가 증가하게 됩니다.

결과적으로, Block Multiplication을 사용하면 데이터가 DRAM에서 불필요하게 여러 번 로드되는 것을 줄일 수 있으며,
메모리 대역폭 사용을 줄이고 성능을 향상시키는 데 큰 도움이 됩니다.

Warp Tiling

Block Tiling은 Block 단위로 데이터를 Shared Memory에 로드해서 사용하는 기법으로,
DRAM에서 불필요하게 여러 번 데이터를 로드하는 것을 줄이는 효과가 있습니다.

하지만, Shared Memory에 데이터를 저장하더라도 연산을 수행할 때마다 매번 Shared Memory로부터 Register로 데이터를 가져와야 하기 때문에 추가적인 오버헤드(overhead)가 발생합니다.
특히, 같은 element들이 A에서는 $N // b_n$ 번, B에서는 $M // b_m$ 번씩 반복적으로 Shared Memory에서 Register로 로드되어 사용되기 때문에
Shared Memory 내부에서도 반복적인 데이터 로드가 발생하는 비효율이 존재합니다.

이러한 문제를 해결하기 위해, Block 내에서 자주 사용되는 일부 element들을 Register에 미리 로드한 후, Register에서 바로 연산을 수행하고, 결과만 C에 저장하는 방식이 사용됩니다. shared memory에서 register로의 load를 줄였기에 Arthimetic Intensity 가 증가하게 됩니다.

이를 통해 Shared Memory로부터 동일한 데이터를 반복해서 로드하는 비용을 줄일 수 있으며,
결과적으로 메모리 접근을 최소화하고 연산 성능을 향상시킬 수 있습니다.

GPU에서 프로그램이 실행될 때는 연속된 thread 그룹(warp 또는 wavefront)이 같은 instruction을 서로 다른 데이터에 대해 동시에 실행합니다.
이를 통해 데이터 병렬성(Data Parallelism)을 극대화할 수 있습니다.

또한, Output Matrix의 영역을 Block 내에 할당된 warp 수만큼 균등하게 나누고,
각 warp 내의 thread들이 Outer Product 방식을 이용하여 결과를 병렬로 누적하도록 합니다.
이 방식은 각 thread가 담당하는 데이터 범위에 대해 독립적으로 연산을 수행하면서도, 결과는 효율적으로 통합될 수 있도록 설계됩니다.

특히, 기존의 Block Tiling 방식과 비교했을 때,
Outer Product를 사용하면 같은 데이터를 Shared Memory에서 Register로 가져오는 횟수를 줄일 수 있기 때문에,
Shared Memory 접근을 줄여 더 높은 성능 최적화가 가능합니다.

결론적으로, Warp 단위의 병렬 연산과 Outer Product 기반 접근법을 함께 사용하면,
메모리 접근 효율을 높이고, GPU의 병렬 연산 자원을 보다 효과적으로 활용할 수 있습니다.

Thread Tiling

기존의 Warp Tiling 기법은 Block Tiling에 비해 Shared Memory에서 Register로 데이터를 로드하는 횟수를 줄일 수 있는 장점이 있습니다.
즉, Warp 단위로 데이터를 나누어 처리함으로써 Block 단위로 처리할 때보다 더 효율적인 메모리 접근이 가능합니다.

하지만 그럼에도 불구하고, Warp Tiling 방식에서도 여전히 동일한 데이터를 여러 thread가 Register로 중복해서 로드하는 현상이 발생합니다.
따라서, 중복된 Register 로드를 완전히 피할 수는 없으며, 이러한 부분에서 추가적인 최적화 여지가 존재합니다.

결론적으로, Warp Tiling은 Block Tiling보다 메모리 접근 최적화가 개선되었지만,
중복된 Register 로드 문제를 해결하기 위한 추가적인 고려가 필요합니다.

Warp 내의 각 thread가 자신이 담당한 데이터만 Register로 로드하여 Output Product를 계산하고, 그 결과를 Output Matrix C에 누적하도록 합니다.
이 방식은 각 thread가 필요한 데이터만 직접 로드해서 사용하기 때문에,
기존의 Warp Tiling에서 발생했던 중복적인 Register 로드를 줄일 수 있습니다.

즉, Warp 단위의 협업(cooperation)을 통해 각 thread가 독립적으로 연산을 수행하면서도,
중복되지 않는 효율적인 데이터 로드로 메모리 접근 최적화를 달성할 수 있습니다.
결과적으로, Warp 내의 thread들이 최소한의 데이터 로드로 최대한의 연산을 수행할 수 있도록 개선된 방식입니다.

Conclusion

지금까지 살펴본 것처럼, 효율적인 Matrix Multiplication을 구현하기 위해서는 GPU의 메모리 계층 구조와 데이터 접근 방식에 대한 깊은 이해가 필수적입니다.
특히, DRAM에서 Register로의 데이터 전송을 최소화하고, 메모리 계층 내에서 반복적인 데이터 로드를 피하는 전략이 GPU 연산 성능을 극대화하는 핵심입니다.

이를 위해 Block Tiling과 Warp Tiling 같은 기법들이 활용되며,
Block Tiling은 DRAM 접근을 줄이는 장점이 있지만 Shared Memory에서 Register로의 반복적인 로드 문제가 존재합니다.
Warp Tiling은 이러한 문제를 일부 완화했지만, 여전히 Warp 내에서 중복된 Register 로드가 발생하는 한계가 있습니다.

이러한 한계를 극복하기 위해 Outer Product 기반의 접근 방식과 Warp 내 thread들이 자신이 담당하는 데이터만을 Register에 로드하여 연산하는 방식이 제안되었습니다.
이를 통해 Shared Memory와 Register 간의 데이터 이동을 최소화하고, 중복된 메모리 접근을 줄임으로써 메모리 대역폭을 보다 효율적으로 사용할 수 있습니다.

결국, 데이터 로드와 연산의 균형을 맞추고, Arithmetic Intensity를 높이며, Data Locality와 메모리 계층을 고려한 최적화가 고성능 GPU 연산을 위한 필수 전략임을 확인할 수 있습니다.

References

https://docs.nvidia.com/cuda/parallel-thread-execution/

https://product.kyobobook.co.kr/detail/S000003500906

https://developer.nvidia.com/blog/cutlass-linear-algebra-cuda/

https://siboehm.com/articles/22/CUDA-MMM