development

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 기반 모델의 정확도를 높이기 위해서는 보다 깊은 네트워크를 구성해야한다고 생각했습니다. 실제로 당시의 논문들을 보면 보다 깊은 네트워크의 성능이 더 나은 경우가 많았습니다.이러한 가정으로부터 다음..

Translation InvarianceTranslation에 invariant 하다는 것은 입력이 바뀌어도 출력은 바뀌지 않는 것을 의미합니다. 입력 이미지가 고양이인지 아닌지 판별하는 함수 f 는 고양이의 위치가 바뀌어도 동일하게 고양이라고 판별해 줄 것입니다.Translation EquivarianceTranslation Equivariance는 입력이 바뀌면 출력도 바뀐다는 뜻입니다. 예를들어 고양이가 있는 영역에는 1을 칠하고, 아닌 영역에는 0을 칠하는 함수 f 가 있다고 합시다. 이때, 고양이의 위치를 이동시키는 S 함수를 이용하여 입력 이미지를 바꿔주었다면 출력도 동일하게 변경된 고양이의 위치에 칠해져 있을 것입니다. Convolutional layer 자체는 아래 그림과 같이, 입력이 변..