GPU는 계속 세대가 발전했고, 그에 따라서 feature들이 바뀌었다. LLM을 다룰것이라면 Ampere 부터의 GPU의 하드웨어적 특성을 알면 좋다. A100, H100 등 각 아키텍쳐의 대표적인 GPU를 통해 알아보자. 텐서코어의 TOPS 같은 성능수치는 다루지 않는다.A100: Ampere (SM80, 2020)2020년 발표된 GPU로, L1 bypass를 통해서 많은 과정이 생략되면서 DRAM의 값을 Shared Memory (SRAM)에 불러오는 것이 최적화되었다.cp.async ptx로 구성되는 이 메모리 복사는 비동기적으로 일어나기 때문에 다음과 같은 소프트웨어 파이프라인을 통해서 latency를 숨길 수 있다.소프트웨어 파이프라인은 연속된 명령어의 종속성을 제거해서 하드웨어를 fully..

Arithmetic Intensity (AI)Arithmetic intensity, 산술강도는 연산량/메모리량, ops/byte(mem) 으로 나타낸다. 즉 AI가 높을수록 동일한 메모리로 더 많은 연산을 할 수 있음을 의미한다. 이전 챕터에서는 SRAM (Shared memory of CUDA), 1d tiling 을 활용해서 성능을 끌어올렸다. 한개의 스레드에서 아래와 같이 여러개의 결과를 만들어낸다. 살펴본 경우와 더불어 확장된 알고리즘의 AI를 생각해보자.앞서 살펴본 커널에서, 한개의 결과만 만들어내는 경우는 17 load 가 필요하다. 반면 1d tiling을 하는것만으로도 11 load 로 줄어들게 되는데, 2d tiling을 하게 되면 9 load로 그보다 더 줄어든다. 이는 GEMM 연산의..

이 글은 포스트를 참고하며 직접 커널과 그림을 작성하며 진행한 공부이다.CUDA는 cuBLAS에서 최적화된 GEMM api를 제공한다. 직접 작성한 커널도 최적화를 통해서 충분히 cuBLAS 급의 성능을 낼 수 있다. 단계적으로 CUDA의 최적화 개념들을 적용하면서 따라가보자.A: (M, K), row-majorB: (K, N), row-majorC: (M, N), row-majorDRAM: Global memorySRAM: Shared memory구현은 다음과 같고, 결과를 먼저 보이면 아래와 같다.Naive implementation, DRAM coalescingSRAM cachingSRAM 1d tiling[BENCHMARK] CUBLAS GEMM │ 0.04533..

CUDA는 아래 그림과 같은 Grid - Block - Thread 의 논리적 계층구조를 사용한다. 블록 내부의 thread 는 3차원으로 구성할 수 있는데, 결론을 먼저 말하면 별다른 메모리 구조를 가지지 않는 한, 스레드 인덱스 (x, y, z) 는 (x + y Dx + z Dx Dy) 로 계산하라. 왜 그런지 알아보자.이때, xyz dimension 순서를 마음대로 바꿔도 성능에 영향이 없을까? 답은 NO 다. CUDA Programming guide 에 나와있는 내용으로, 블록을 3차원으로 했을때의 id는 다음과 같이 계산된다.이걸 무시했을 때 실제 커널에서 어떤 영향을 끼치는지 알아보자. 다음과 같은 커널 두개가 있다고 생각해보자. 1024개 원소의 vector 2개를 더하고 output에 저..

Original softmax$$\sigma_i(\mathbf{z}) = \frac{e^{z_i}}{\sum^K_{j=1}e^{z_j}}$$Softmax의 수식은 위와 같다. 입력 벡터에 대해서 load 2번, store 1번, 총 3번의 메모리 접근이 필요하다. 이 함수의 의미는 주어진 값들 속에서 각각의 원소의 확률을 계산해준다. 유용한 함수이지만, 실제 컴퓨터상에서 구현하게 될 때는 부동소수점으로 인한 문제가 발생하기 쉽다. 부동소수점은 소수를 컴퓨터로 표현하는 방법인데, 중요한 점은 표현할 수 있는 범위가 한정된다는 것이다. 그런데 softmax 연산에는 지수함수(\(e^z\))의 특성으로 인해서 값이 아주 커진다. 이 값들의 합연산을 하면 overflow가 발생하기 쉽다. 또한 반대로 음수방향으..

Python에서 triton 언어를 활용해서 triton 함수를 정의해준다. triton 함수는 `@triton.jit decorator` 형태로 정의됨.Triton의 컴파일 과정은 다른 포스트에서 따로 다루는걸로..@triton.jitdef _fwd_kernel( Q, K, V, Out, Lse, TMP, softmax_scale, batch, nheads, ... , EVEN_M, EVEN_N, EVEN_HEADDIM, IS_CAUSAL: tl.constexpr, BLOCK_HEADDIM: tl.constexpr, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr,):함수를 정의했다면, 이제 컴파일러를 호출할 차례다.expor..

CUDA에서 host에서 device로 memory를 복사하는 방법은 cudaMemcpy API 를 활용하는 방법이 있다. 기본적으로 별다른 작업 없이 선언된 host의 데이터는 pageable data로 동작한다. pageable data로부터 데이터를 device에 복사하기 위해서는 host 내부에서 pageable -> pinned memory로 한 번 옮겨가야 하기 때문에 속도가 느려지게 된다. cudaMallocHost 로 직접 memory pinning, non-pageable memory를 선언하게 되면, 해당 과정이 생략되기 때문에 복사가 더욱 빠르게 이루어진다. 1GB의 메모리를 device로 복사하는 경우의 두 방법의 속도 차이는 다음과 같다.$ ./pinned_memory Pinned..

Bank ConflictsCUDA를 공부하면 Shared Memory에 대한 개념을 접할 수 있습니다. 그러면서 Bank Conflicts 에 대한 개념을 접할 수 있는데요. 아래와 같은 그림 혹은 비슷한 그림들을 보셨을 겁니다. Bank Conflicts는 메모리 요청이 뱅크에 중복되게 일어날 때 발생하는 것을 나타냅니다.Shared memory는 32-bit 크기를 가진 메모리 모듈 32개로 이루어져있습니다. 64개의 fp32 데이터를 SRAM에 넣으면 아래와 같이 들어가게 됩니다. 즉, 32개의 bank에 순차적으로 속하게 됩니다. 각 스레드가 각각 다른 뱅크에 접근할 때(1, 2번)에는 bank conflict가 발생하지 않습니다. 3번 경우에는 bank conflicts가 "일어날 가능성이" ..

Motivation & GoalNVIDIA는 AI 시장에서 하드웨어와 그에 맞는 소프트웨어 스택이 압도적임CUDA는 하나의 큰 생태계가 되었고, 유저들은 많은 기법들을 이용해서 NVIDIA GPU를 fully utilize하고 있음Pytorch, vLLM과 같은 라이브러리들은 backend에 숨겨서 일반 유저들이 CUDA를 몰라도 연구를 할 수 있게 하였음그러나 GPU의 내부를 알게되는것은 충분한 강점이며, 보다 나은 성능을 위해서는 필수적임큰 틀에서의 CUDA 의 개념들을 훑어보기로 합시다Memory Hierarchy그렇다면 CUDA, NVIDIA GPU는 내부적으로 어떻게 동작하는지 알아보겠습니다. 우선 메모리 구조를 살펴보면 다음과 같습니다.Global memoryGPU의 메모리로, DRAM 영역..

나온지 꽤 된 논문인데요, ML 도메인에서는 아주아주 유명한 논문인 것 같습니다. ILSVRC'15 에서 당당하게 1위를 차지한 모델로, CVPR'16의 best paper 로 선정되었습니다. 마이크로소프트 북경 연구소에서 개발하였고, ResNet 으로 불리는 네트워크입니다. Shortcut/Skip connection 을 이용하여 residual (잔차) 을 학습시켜 성능을 향상시킨 논문으로, 이후 대부분의 모델들에서 ResNet 을 활용하게 되었습니다.Inspiration이 논문이 발표되기 이전까지, CNN 기반 모델의 정확도를 높이기 위해서는 보다 깊은 네트워크를 구성해야한다고 생각했습니다. 실제로 당시의 논문들을 보면 보다 깊은 네트워크의 성능이 더 나은 경우가 많았습니다.이러한 가정으로부터 다음..