cuda programming에 대해서 배울건데 programmming model, gpu hardware 왔다갔다 거릴 거임

programming model를 이해하려면 hardware을 알아야해서 왔다갔다 거리는 거임

GPU Frist Try

지난시간에 first try에 대해서 배웠고 cudaMallocManaged로 할당했음.

• let allowing cuda runtime automatically gpu memory and cpu memory

  → memory gpu and cpu 모두 관리해주고 싱크해줌

• cpu code로 접근하면 cpu메모리 접근, cuda kernel코드에서 접근하면 gpu 메모리 접근
• cuda는 cpu, gpu memory 모두 관리하기에 조금 느릴 수 있음

Real-world GPU programming

• 실질적으로는 cpu dram, gpu dram 다르게 할당하는 게 좋음

• host == CPU
• CPU에 input 준비하고 Output 할당 → GPU에 input, output 할당 → CPU input GPU에 복사 → GPU에서 계산(kernel 시작) → GPU output을 CPU에 복사
• 언제 gpu메모리를 접근해야할지 cpu 메모리를 접근해야할지 프로그래머가 아니까 빨라질 수 있는 것!
• but, cudaMallocManaged 명령어로 하면 unified memory로 관리하게 되는 거임

1. Setting up inputs and Outputs on Host(CPU)

• cpu memory에 할당 : new syntax 이용
• array content 역시 초기화해줘야 함

2. Allocate memory(input & output) on GPU

• 그 후에 cudaMalloc함수로 memory space in gpu
• d_a는 n * float 크기만큼 할당하고 d_a는 그 배열을 가리키게 됨

cf.

gpu에 있는 값을 가리키는 포인터는 주로 prefix로 device_인 d_를 사용

h_는 host_이기에 얘는 cpu를 가리키는 포인터를 의미

• cpu address가 의미하는 바는 저걸 의미함
• d_a는 실질적으로 cpu에 있지만 포인터는 gpu에 가리킴
• cpu는 d_a를 가리킬 수 있지만 d_a[0]는 가라킬 수 없음 가리킬 수는 있지만 cpu에 있는 어느 이상한 값을 가리키게 될 거임

3. Copy Input from Hosts to GPU

• cpu, gpu에 공간을 할당한 후에 cpu → gpu로 값을 복사
  • cudaMemcpy
  • host: cpu, device: gpu

4. Start GPU Kernel function

y가 아니라 d_a,d_b,d_c로 해서 d_c에 저장하도록 뭐 바꿔야 완전 같겠지..???

• kernal을 부른 후에 syncrhronize되길 기다림

5. Copy Output from GPU to Host

• 그 후에 gpu에 결과가 저장됙에 cudaMemcpy를 이용해서 gpu → cpu로 카피해줘야함

• 추가적으로 Exit, quit하면 구냥 다 알아서 없애주겠지만 다른 커널들이 있다면 free해줘야 함
• CPU memory는 new로 했으니까 free로 해제
• GPU memory는 cudaMalloc으로 했으니까 cudaFree로 해제

C extensions for gpu computing

• gpu prgramming에서 주로 cuda kernel을 .cu파일에 적고 해더는 .cuh에 저장됨(컨벤션)
• nvcc로 컴파일
  • g++, gcc로 컴파일 혹은 nvcc로 사실 해도 됨
• 양이 작다면 구냥 .cu에 넣고 컴파일해도 뭐 상관없대 뭐라니

#include <stdio.h>

__global__ void add(int n, float *x, float *y, float *z)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;
	for (int i = index; i< n; i+=stride)
	  z[i] = x[i] + y[i];
}

int main()
{
  float *h_a, *h_b, *h_c;
  int N = 10<<10;
  h_a = new float [N];
  h_b = new float [N];
  h_c = new float [N];

  for(int i=0;i<N;i++){
    h_a[i] = 1.0f;
    h_b[i] = 2.0f;
  }

  float *d_a, *d_b, *d_c;
  cudaMalloc((void**)&d_a, N*sizeof(float));
  cudaMalloc((void**)&d_b, N*sizeof(float));
  cudaMalloc((void**)&d_c, N*sizeof(float));

  cudaMemcpy(d_a, h_a, N*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_b, h_b, N*sizeof(float), cudaMemcpyHostToDevice);
  
  int blockSize = 256;
  int numBlocks = (N+blockSize - 1) / blockSize;
  add<<<numBlocks, blockSize>>>(N,d_a,d_b,d_c);

	/*이렇게 할 수도 있음
  dim3 grid(blockSize);
  dim3 block(numBlocks);
  add<<<grid, block>>>(N,d_a,d_b,d_c);
*/
  cudaDeviceSynchronize();

  cudaMemcpy(h_c, d_c, N*sizeof(float), cudaMemcpyDeviceToHost);
 
  for(int i=0;i<N;i++){
    printf("%.2f, ",h_c[i]);
  }

  free(h_a);
  free(h_a);
  free(h_a);
  cudaFree(d_a);
  cudaFree(d_b);
  cudaFree(d_c);
  return 0;
}

✅ Thread block organization Keywords (왠지 중요할 거 같은 느낌)

thread structure이 어떻게 organize되고 data에 맵핑되는지?

• thread는 three level이 존재
  • thread
  • group of thread: block
  • group of block: grid
• 1d,2d,3d dimension structure을 가짐
  • up to 3d structure
• thread는 basically cpu thread와 아주 비슷, 다른점이 있긴 하지만 그건 나중에..
• user는 group of thread를 block으로 해서 사용 가능
  • map to hardware computation unit
• group of bloock인 grid는 cuda kernel function이 끝나면 grid는 사라짐. 다른 커널이 생기면 grid이 또 생기고…구렇게 됨

• thread, thread block은 모두 3d structure을 가짐

  

  • block, thread는 3d structure을 가지기에 3d index를 가짐
  • 각 블럭에는 thread가 있는데 이거 역시 3d structure을 가짐

Q. 왜 저렇게 thread와 block 구조가 저렇게 되어있을까??

A. easy to map thread, block structure to data(often 1d,2d array)

Q. 왜 thread가 multilevel(thread → block → grid)이렇게 인걸까??

• 하드웨어 디자인 하기도 쉽고 잘맞음
• easy to design hardware scale
• thread와 블록을 하드웨어에 잘 맵핑되게 하기 위해. multi level로 해야 하드웨어 디자인 하기가 쉽기 때문
• thread block을 hardware unit(SM)에 맵핑할건데 thread block간에는 커뮤니케이션하지 않아도 디자인했기에 하드웨어도 그렇게 하는 거임…
  • 커뮤니케이션은 오버헤드가 커서 지운 거임
• 그래픽스는 그렇게 해도 되고 나중에 봤더니 다른 분야에서도 사실 그렇게 해도되더라~ 그런거임…

Q. global memory가 블락간에 공유가 되고 있지 않은가?

A. 그렇게 될 수도 있는데 스레드 간의 커뮤니케이션 가정이 없도록 디자인이 되는 것

(사실상 share memory가 있게 하려면 할 수 있지만 그냥 안한다라고 가정하고 짜는 게 맞다는 것!)

• 2048개들의 쓰레드를 묶어서 블럭을 했을 때 걔네들은 같은 stream multiplrocessor에서 실행되는데 얘네들 사이에는 커뮤니케이션 가정이 가능함. 근데 다른 쓰레드 블록끼리 있는거에는 가정을 못하도록!
• sm 간에 커뮤니케이션이 안된다고 가정

Q. block안에 shared memory에서 race condition이 있지 않낭??

• sm 사이의 데이터를 업데이트 될 때의 coherience라고 말함
• thread 사이에는 race dondition이 일어날 수도 있긴 함…
• atomic add 같은 것도 있지만 구냥 memory acccess 가져와서 global → local memory 에는 cohereicne가 없다고 함

⇒ thread block 노랑이는 stream multiplrocessor에 맵핑되고 다른 sm에 접근하는 애들은 서로 커뮤니케이션이 없음. block에 있는 애는 shared memory를 공유하는 형태

(정리)멀티레벨인게 hardware unit과 data structure에 잘 맞기 위해서라는 거만 알면 됨 구조는 정확히 알 필요는 없음

Block and grid dimension

• dim3 구조로 그리드 구조와 블럭 구조 선언 가능

Grid/block/thread visualized

• 3d structure map to data 1,2, 3d array
• grid에는 6개의 x,y thread block을 가짐
• thread block안에도 x,y thread 16개를 가짐

CUDA grid

9 thread block을 가지고 첫번째 블록을 보면 3 thread * 3 thread를 가짐

• index, dimension을 thread, block, grid level에서 가짐
• 가장 큰 grid level에는 gridDim을 가짐
  • block이 3개씩이니까 3
• 각 블럭에는 3개의 thread가 있기에 blockDim은 3
  • 만약에 2*2 thread block이 된다면 blockDim x,y 값은 2가 됨
• 각 블록의 idx는 threadidx.x, threadid.y
• 만약에 map to 9*9 array 그러면 어떻게 맵핑할 지 알 수 있을 것!
  • 3 dimension x * blockDim.x = number of elements
  • 얼마나 필요한지 알아낸 후 computation

⇒ 이제는 어떻게 Thread structure을 data에 map하는 거에 대해서 알아보자!

Inside a GPU

• thread block은 stream multiprocess(노랑이)에 맵핑됨
• 각 sm은 multiple processor(==core)과 레지스터를 가짐
  • 각 프로세서는 computation unit을 가짐
  • 각 프로세서는 own register를 가짐
• 하지만 single sm은 one shared memory를 가짐
• sm끼리에는 dram(Device memory)을 share함

• core(연두)끼리는 computation unit은 separate되어있지만 instruction unit은 공유

  == m개의 processor는 execute same instruction이라는 의미임!!!

  • 그래도 end up different computation임!
  • 왜냐면 they own register를 가지기에 add r1 r2를 하게 된다면 processor는 모두 저 동작을 하지만 r1, r2 값이 프로세서마다 다르기에 결과가 다를 수 있음

• fetch data from cpu dram보다는 fetch data from gpu dram이 빠름. PCI-e bus를 통해 가야 하기에 오래 걸림

  

• cuda kernel이 launch된다면 copy cpu memory → gpu memory되고 kernel은 thread block, grid에 대한 정보가 적혀있으니까 gpu hardware는 thread block 그만큼을 확보 → thread block scheduler로 each sm에 assign, sm은 computation을 실행함.

  

• 512개 쓰레드는 너무 많으니까 cuda는 저 512개를 다 실행하지 않을테니 small unit(warp)이라는 걸로 나눔

  

  • 32 thread이 warp이 있고 완전 동일한 연산을 진행
  • 싱글 블럭에는 여러 개의 warp들이 있을 거고 warp끼리에는 어떤 애는 다른 애보다 빨리 계산되지만 하나의 warp안에는 모두 완전히 같은 연산을 진행
  • warp은 32개로 고정됨, processor도 32개의 배수로 존재
• 512개는 너무 커서 32로 줄였지만 32 thread는 different execution이 될 수 있음.. branch가 있다면 어떤 애는 저렇게 어떤 애는 이렇게 이렇게 갈수 있음 : warp divergence
  • subst of thread within the same wrap인데 path가 여러개

cf.

Warp divergence(p35)

• differment thread면 idx가 달라지기에 a배열의 다른부분에 접근하게 될 거임
  • 비주얼화한다면 어떤 애는 연두로 어떤 애는 파랑이로 갈거임
• 만약에 화살표가 thread in wrap이라고 가정, warp에 들어있는 어떤 쓰레드 subset는 연두로 어떤 쓰레드 subset들은 블루로 감

Q. 어떻게 해결?

A. execute as number of path

• wrap에 8개의 스레드가 있다라고 가정한다면 path 두개 별로 다 실행시키는데 path a 시 해당하는 스레드만 enable시키고 path b시 해당하는 쓰레드만 enable시킨다
• 실제로는 표시되지 않는 화살표의 결과는 레지스터에 저장하지 않는 걸로 뭐 그렇게 한대

→ GPU programming에서는 if나 branch를 안쓰는게 좋다!

Streaming multiprocessor

• 싱글 sm 안의 모습을 보여주고 있음
• sm 안에 4개의 execution unit이 존재
  • 각 유닛에 32개의 core(computation unit)이 들어있음
  • LDIST, SFU 등등의 애들도 갖고 있음
  • 각각이 하나의 warp을 실행하니까 warp executor라고 부르기도 함
• 이 4개의 execution unit을 L1 cache를 공유(두개는 구냥 그림임)

cf.tex는 무시할 것