20160331

1. 동시성 프로그램의 이해
– Blocking에서 블록의 크기를 결정할 때에는 아키텍쳐를 고려해야 한다. G80에서 SM은 물리적으로 최대 768개의 thread, 최대 8개의 thread block을 지원한다. 8X8 크기의 블록을 사용하면 12개의 thread block이 생기기 때문에, 한 번에 8개의 thread block 밖에 실행하지 못한다. 결과적으로 한 번에 512개의 thread밖에 실행하지 못한다. 이처럼 블록 크기를 결정할 때 아키텍쳐 측면도 함께 고려해야 한다.

– 같은 thread block에 있는 thread들은 __syncthreads() 함수를 사용해 동기화할 수 있다. 이 함수를 사용하면 thread block 내의 모든 threads의 실행을 동기화할 수 있다.

– GPU에서 실행되는 어플리케이션은 1) 높은 수준의 병렬성을 갖는다. 2) vector-like 연산을 수행한다(중요). GPU는 이러한 특징을 고려해 설계되었다. 병렬성을 사용하기 위해 multi-core, multi-threaded 설계되었다. G80에서는 한 사이클에 8개의 thread, 동시에 최대 768개의 threads를 실행할 수 있다. CUDA는 SPMD 모델을 따르지만, 실제로 아키텍쳐의 구현은 SIMD 모델을 보인다. 따라서 분기되는 if문에서는 두 개의 경로를 모두 실행해야 한다. 이러한 경우에는 성능이 떨어진다. 고성능을 위해서는 if문의 사용을 줄여야 한다. 이러한 아키텍쳐의 특징이 다시 CUDA, application에 영향을 준다. 이처럼 어플리케이션과 아키텍쳐는 상호 영향을 주며 발전한다.

– UltraSPARC T2는 한 개의 물리적인 코어가 8개의 thread를 지원함. 8개의 코어가 있으면 논리적으로는 64개의 코어가 보이게 된다. GPU에서는 32개의 thread가 한 개의 warp이며, SM은 한 번에 8개의 thread를 실행하므로 4 cycle마다 새로운 warp를 스케쥴링한다. GPU는 SIMD도 MIMD도 아니고 SPMD이다.

– CPU는 캐시가 대부분의 크기를 차지하고, control logic은 상대적으로 크기가 작은 편이다. GPU에서는 복잡한 OoO logic, predictor, prefetch를 제거한다. Single fetch, single decode unit을 가지면서도 execution unit의 수는 늘렸다. 그리고 context 저장을 위한 공간을 늘렸다.

– 기본적으로 processor에서는 pipeline을 사용하고 있음. 그리고 SMT에서는 이러한 pipeline을 여러 개의 thread가 공유함. 여러 thread가 공유하기 위해서는 icache와 register map에서만 관리해주면 됨. backend에서는 dependency만 잘 관리해주면 됨. GPU에서는 한 번에 8개의 코어에서 32개의 thread를 실행하며, 따라서 한 번에 4 cycles가 걸림. 4 cycles마다 한 개의 warp를 실행. CPU에서는 memory latency를 숨기고자 OoO 실행을 함. GPU에서는 in-order execution을 하므로, 대신에 여러 개의 thread를 동시에 실행해 latency hiding 달성함. 따라서 GPU에서는 thread를 얼마나 빨리 실행하느냐가 중요하지 않음. GPU에서는 전체 threads의 실행을 얼마나 빨리 끝내느냐가 중요함. 따라서 bandwidth가 중요하다.

– 한 개의 SM은 한 번에 768개의 threads를 생성할 수 있지만, thread context에서 저장해야 하는 register file의 크기에 제약된다. thread context에서 저장해야 하는 register file의 크기가 크다면 더 적은 threads를 생성할 수밖에 없다.

– Out of order execution은 thread 안에서의 instruction 재정렬을 통한 실행을 의미한다. GPU에서 warp내 thread 종료 시간은 다를 수 있지만, thread 내에서의 instruction 재정렬은 없으므로 OoO execution이 아님.

– GPU는 register, local memory, shared memory, global memory, constant memory를 가짐. Constant memory는 3D rendering 과정에서 쓰이는 상수를 저장하기 위해 사용된다. local 변수는 DRAM에 생성되나, shared 변수는 streaming multiprocessor마다 있는 SRAM에 생성됨. constant로 선언하면 DRAM에 저장되지만 constant cache에 캐시됨.

– GPU 쓰레드 사이에서 실행을 제어할 수 없으므로 동기화가 어렵다. 이를 해결하기 위해서는 atomic operation을 사용하면 됨.

– MXN 행렬 곱에서 M의 한 요소는 n번 읽게 됨. Locality를 잘 사용한다면 반복해서 읽을 필요가 없음. Tiling을 적용하지 않았을 때에는 모든 thread가 독립적으로 동작하므로, 같은 데이터를 matrix width만큼 반복해서 읽게 됨. Tiling을 적용했을 때에는 한 블록씩 읽은 뒤, thread들이 협동하여 실행함. Tiling이 없을 때에는 각 요소를 n번 읽었으나, tiling을 적용하면 tile width만큼만 읽으면 됨.

– 코드에서 __shared__ Mds, Nds가 선언되어 있으나 이는 쓰레드 사이에 공유되는 것. Shared memory에 한 번만 가져오고, 그 뒤에 여러 번 연산한다. Shared memory는 SRAM이며, 사용자가 제어할 수 있다. phase가 있으니 prefetch 가능한지 생각해보자.


2. 창현이 형의 git 사용법
git log –oneline –all –graph –decorate 명령을 사용해, 현재까지의 작업 흐름을 쉽게 볼 수 있다.

Advertisements
Tagged with: , , , , , , , , , , , , , , , , , , ,
Posted in 1) Memo

Leave a Reply

Fill in your details below or click an icon to log in:

WordPress.com Logo

You are commenting using your WordPress.com account. Log Out / Change )

Twitter picture

You are commenting using your Twitter account. Log Out / Change )

Facebook photo

You are commenting using your Facebook account. Log Out / Change )

Google+ photo

You are commenting using your Google+ account. Log Out / Change )

Connecting to %s

누적 방문자 수
  • 98,564 hits
%d bloggers like this: