Bit-reversal 재귀 트리를 따라 내려가면서 배열이 어떻게 찢어지는지 보면, 특이한 패턴을 발견할수 있다. N = 8인 배열이 재귀 함수를 타고 내려가는 과정을 보자 매 step마다 2로 나뉘어지고, 숫자표현이 이진수로 되어있으니 각 step 별로 짝홀 판단은 LSB로 시작해서 MSB에 결정됨을 쉽게 이해할수 있다. 따라서 우선 배열...

Butterfly 연산 다항식을 짝수항, 홀수항으로 분할 할수 있다는 것은 이해가 되는데 후에 병합될때가 잘 이해가 안된다. [P(x)=P_{even}(x^2)+xP_{odd}(x^2)] N개의 복소수 거듭제곱근을 각각 대입했을 때 나오는 point value들을 가지고 butterfly 연산을 수행하면 더 촘촘한 결과를 뽑을수 있다는 것인데 이...

Sound를 Frequency로 분해, 신호를 다항식으로 보고 그 다항식을 단위근에서 평가 한게 신호 x[n]을 계수로 가지는 다항식 뿐만 아니라, 불확정성원리, 리만 제타 함수와 소수, 미분 방정식등 여러 분야에서 활용됨 다항식 관점에서의 DFT (Convolution) Coefficient representation -> Point va...

Matrix addition을 stream을 사용해서 최적화 시키고 분석해보면서 배운것 기존 synchronous하게 matrix addition에서 DMA Controller 작업과 SM작업을 overlapping하서 최적화를 진행 하였다. 기존 default stream으로 synchronous하게 진행되는것에서 4개의 stream에 asyn...

Kernel을 실행할때 block수, thread수 정하기 Block당 thread수 정하기 Warp 단위에 맞춰서 32의 배수로 설정; 아니면 32의 나머지는 놀게 된다. (보통 256으로 설정) Block 하나의 너무 많은 thread를 배정하는 경우 SM의 shared memory, register의 공간이 부족해져...

GPU에서 thread hierachy가 왜 이리 헷갈리는지 정리하고 넘어가자 Overview Kernel 실행시, 새로운 Grid가 하나 생성되는 거야. Grid: 커널이 실행되는 공간 (kernel 실행시 Gird가 생성) Block: 여러개의 워프들의 그룹, 같은 block 내의 thread들은 Shared M...

Asynchronous Execution What is Asynchronous Concurrent Execution? Concurret Execution이 가능한 것들 Host에서의 연산 Device 연산 Data transfer Host -> Device Device -> Host ...

Basics of SIMT, Thread Hierarchy 모든 thread가 같은 kernel code를 실행하지만 고유한 thread ID가 있어서 서로 다른 데이터에 접근하여 작업을 수행 Grid > Block > Thread의 3단계 계층 구조를 갔는다. Thread: 고유한 register 존재 Block: Shared ...

Kernels Specifying kernels Code for kernel is specified using __global__ indicates to the compiler that this funciton will be compiled for the GPU Kernels are functions with a void return type. L...

FLOPS: FLoating point Operations Per Second GPU는 일반적으로 SIMD가 아닌 SIMT(Signle Instruction Multiple Thread)로 정의 된다. 한 쓰레드 그루부 내 스레드들을 하나의 제어 장치로 제어 각 쓰레드는 자신만의 제어 문맥을 가진다. 하나의 제어 장치가 여러개의 쓰레드를 제...