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

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