Cuda gpu 프로그래밍 공유 메모리 이용한 행렬의 곱셈

음 쿠다에서 공유 메모리랑 글로벌 메모리를 사용할 수 있는데 공유 메모리를 사용하는 편이 온칩으로 훨씬 빠르므로 공유 메모리를 사용해 보겠다. 공유 메모리는 각 블록 안에 있다. 그리드 > 블록 > 스레드 순으로 포함 관계를 갖는다. 그러기 위해서 수업에서는 공유 메모리의 메모리 뱅크를 잘 고려해서 각 스레드가 서로 다른 메모리 뱅크들에 접근하면 뱅크 콘플릭트를 피해서 빠르게 행렬 곱셈을 할 수 있다는 소리. 지피유의 글로벌 메모리에 접근하는 속도가 공유 메모리에 접근하는 속도 보다 100배는 느리다고 한다. 내가 해보려는건 공유 메모리를 사용했을때랑 글로벌 메모리 사용했을때 차이 그리고 cpu에서 했을때 3경우를 비교해 보자.

 

지금 공유 메모리를 사용하는 코드에서 계산하는 방법은 음 잘 몰라서 이 글을 쓰면서 정리해 보는 것이다.

 

공유 메모리는 커널 내부에서 초기화 해주어야 한다. 그래서 디바이스 함수에서 초기화 되도록 하고 __syncthreads();를 통해서 공유 메모리에 그 값을 모두가 공유하도록 해주어야 한다.

 

나는 512x512 행렬의 곱셈을 하려 한다. 그리고 각 스레드에서 한개의 행과 열을 곱해서 결과 행렬의 한 값을 구하도록 한다. 그러면 스레드가 총 512x512개 필요한 셈인듯. 

 

그리드 전체에 행렬 곱을 하는 셈이니까 그리드와 블록에 대한 그림으로 좌표를 이해해 보자.

 

https://www.3dgep.com/cuda-thread-execution-model/

y축이 row고 x축이 col이다.

 

 

각 스레드에서 그러니까 해야되는 일은 , 자신이 전체 512x512 행렬에서 몇행 몇열 스레드인지 알아보고 그곳에 올바른 값을 저장하면 된다.

 

그 스레드의 row = blockDim.y* blockIdx.y + threadIdx.y; 로 몇번째 블록인지 알아야 현재 스레드 위에 몇개의 스레드가 지나왔는지 계산하고, 현재 블록 내부에서 몇 행째에 있는 스레드인지 알면 된다. col 도 마찬가지. 

 

그리고 공유 sa, sb의 선언이 float형이므로 각 행렬 값은 4바이트로 bank width 가 보통 32비트니까 일치한다고 보면 된다.

bank width는 결국 그냥 구조체만 아니고 각 스레드가 한 영역만 간다면 콘플릭트가 일어날 일은 별로 없는듯...?

Row*numAColumns는 현재 위치 위의 스레드들이다. threadIdx.x는 현재 블록안에서 스레드의 x값이다. k*32는 그 행의 블록들을 하나씩 옮겨 가면서 계산한다. 즉 결국 sA는 현재 스레드의 같은 행의 여러 블록들을 한번씩 계산하는 셈이다......??반복문안에 왜 있지. 사실 잘 모르겠다. k의 역할을......