Until now기존의 CUDA 프로그래밍은 CUDA C++을 바탕으로 하고 있었다. SIMT 기반으로 데이터를 블록과 스레드에 맵핑해줄 필요가 있다. 세밀한 설계를 통해 최선의 성능을 이끌어 낼 수 있으나, 모두가 손쉽게 할 수 있지는 않은 일이다. 또한 GPU 아키텍쳐가 발전함에 따라서 하드웨어 스펙도 달라지니 커널 개발자는 각 GPU에 맞는 최선의 커널을 다르게 개발해야했다.이런 어려움에 대해 떠오른 것이 triton-language다(NVIDIA Triton Inference Server와는 관계가 없다). Triton은 메모리관리, 동기화, 텐서코어 스케쥴링 등 블록 레벨의 추상화를 통해 접근성을 높였다. 이에 더해 MLIR을 통해 NVIDIA GPU뿐 아니라 AMD, NPU 등 다양한 하드웨어..

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..

이 글은 포스트를 참고하며 직접 커널과 그림을 작성하며 진행한 공부이다.CUDA는 cuBLAS에서 최적화된 GEMM api를 제공한다. 직접 작성한 커널도 최적화를 통해서 충분히 cuBLAS 급의 성능을 낼 수 있다. 단계적으로 CUDA의 최적화 개념들을 적용하면서 따라가보자.A: (M, K), row-majorB: (K, N), row-majorC: (M, N), row-majorDRAM: Global memorySRAM: Shared memory[BENCHMARK] CUBLAS GEMM │ 0.045334 ms (w:10 r:20)[BENCHMARK] GPU GEMM 0 NAIVE │ 3.943722 ms (w:10 r:20) [PASSED]..

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..

한참 찾았는데, vs code의 C++ extension은 더이상 cursor ide에 C++ intellisense 를 지원해주지 않는다고 한다. Cursor측에서는 자체 C++ extension 으로 따로 제작해서 clangd로 해결하였다. 하지만 나는 CUDA를 사용하는데, clangd에서 CUDA C/C++를 잘 인식하지는 못하는 것 같다. 난 야매 방법으로,1) vs code에서 구버전 C++ extension 폴더를 미리 백업2) cursor 에서 local directory로부터 extension을 설치하는 방법으로 우회해서 사용중이다.

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 영역..