[General Purpose GPU] ch2. Programming model

 

 

*General - Purpose Graphics Processor Architecture의 chapter 2를 읽고 정리한 내용입니다.

---

 

 

non-graphic computing목적으로 GPU가 program되는 방법에 대해 알아볼 것이다.

이는 GPU 연구에서  hardware/software interface를 변경할 때 필수적이다.

 

현대 GPU는 data-level parallelim을 활용하기 위해서 wide SIMD hardware를 사용한다.

이 SIMD 하드웨어를 직접적으로 프로그래머에게 노출하기 보다 CUDA와 OpenCL은 MIMD 프로그래밍의 특징을 갖고 있어서 프로그래머가 많은 scalar thread를 GPU에 launch하도록 한다.

 

runtime 시 GPU hardware는 warp 단위로 실행된다.(NVIDIA에서는 32개 묶음의 thread를 warp로, AMD에서는 64개 묶음의 thread를 wavefront라 함). 또한 warp를 동시에 실행해서 규칙적이고 spatial locality를 활용하는 SIMT 방식을 사용한다.

 

 

 

2.1 Execution model

GPU computing application은 CPU에서 시작된다. 

 

discrete GPU에서 application의 CPU portion은

1. GPU에서 computation에 사용할 memory를 할당

2. input data를 GPU memory로 transfer

3. GPU에서 computation kernel을 시작

 

 

unified GPU

1. GPU에서 computation kernel을 시작

 

여기서 computation kernel은 수 천개의 thread로 구성된다, 각 thread는 동일한 프로그램을 실행하지만 computation 결과에 따라 프로그램을 통해 control flow가 달라질 수 있다. 

 

SAXPY(Single-precision A * X plus Y)

기본 선형대수 소프트웨어 라이브러리로, Gaussian elimination과 같은 higher level matrix operation을 구현하는 데 유용하다.

 

Y = a * X + Y

 

/* traditional code */

void saxpy_serial(int n, float a, float *x, float *y){
	for(int i = 0 ; i < n ; i++)
    {
    	y[i] = a * x[i] + y[i];
    }
}


main(){
	float *x, *y;
    int n;
    ...
    saxpy_serial(n, 2.0, x, y) // invoke serial kernel
    ...
}

 

위 코드는 매 iteration 마다 x, y의 element를 읽어서 연산을 수행한다.

또한 CPU에서 작동하는 코드이므로 serial 방식으로 구현되어 있다. 

 

/* CUDA code */

__global__ void saxpy(int n, float a, float *x, float *y)
{
	int i = blockIdx.x*blockDim.x + threadIdx.x;
	if(i<n)
	y[i] = a*x[i] + y[i];
}


int main() {
	float *h_x, *h_y;
	int n;
	...
	float *d_x, *d_y;
	int nblocks = (n + 255) / 256;
	cudaMalloc( &d_x, n * sizeof(float) );
	cudaMalloc( &d_y, n * sizeof(float) );
    	cudaMemcpy( d_x, h_x, n * sizeof(float), cudaMemcpyHostToDevice );
	cudaMemcpy( d_y, h_y, n * sizeof(float), cudaMemcpyHostToDevice );
	saxpy<<<nblocks, 256>>>(n, 2.0, d_x, d_y);
	cudaMemcpy( h_x, d_x, n * sizeof(float), cudaMemcpyDeviceToHost );

}

 

위 CUDA 프로그래밍 모델에서 kernel은 warp로 이루어진 thread block으로 구성된 grid 의 hierarchy로 organize된다.

각 thread는 operand가 scala value인 instruction을 수행한다.

GPU 하드웨어의 효율성을 높이기 위해 thread 그룹을 lock-step으로 실행한다.

따라서 GPU에서 작동하는 코드이기 때문에 parallel하게 구현되어 있다.

 

오늘날 mobile device system은 CPU와 GPU를 single chip으로 통합하고 있다. 

또한 unified memory를 사용해 GPU와 CPU 간 memory는 tranparent하게 업데이트하고 있다. 이를 통해 runtime에 하드웨어가 자동으로 처리하게 되므로  프로그래머가 memory copy를 직접 관리하지 않아도 되며 data 이동에 대한 걱정을 하지 않고 GPU와 CPU의 memory를 사용할 수 있다.

 

• h_ : CPU memory
• d_ : GPU memory
• cudaMalloc : CPU에서 call 되어 GPU driver가 GPU에 memory를 할당하도록 함
• cudaMemcpy : h_x가 가리키는 CPU memory의 data를 d_x가 가리키는 GPU memory로 복사. 혹은 역으로도 가

 

Parallel programming - GPU에서 thread 할당 방법

parallel programming은 data를 여러 부분으로 나누에 각각의 thread에 할당하는 방식을 사용한다.

GPU에서 각 thread는 grid 내 thread block 내에서 자신의 위치를 identify할 수 있다.

grid는 x, y, z dimension을 갖고 있다. 각 thread block은 grid 내에서 x, y, z 좌표를 갖는다.

thread block 역시 x, y, z dimension을 갖고 있다. 각 thread는 thread block 내에서 x, y, z 좌표를 갖는다.

따라서 GPU에서 각 thread는 실행 중에 3 dimension 공간에서 unique한 coordinate를 갖게 되고 해당 위치에 따라 작업을 수행하게 된다.

이러한 coordinate는 kernel configuration을 통해 세팅된다.

 

• threadIdx.x : thread block 내에서 특정 thread의 x 좌표
• blockIdx.x : grid 내에서 특정 thread block의 x 좌표
• blockDim.x : x dimension 상 thread의 최대 갯수
• blockDim.x + blockIdx.x +  threadIdx.x = array에 access할 때 offset을 나타냄

compiler와 hardware의 조합은 thread가 warp 안에서 lock-step 방식(=SIMD 방식)으로 실행되는 것에 대해, 각 thread가 warp 내에서 독립적으로 실행되는 것처럼 보이게 만들어, 프로그래머가 크게 신경쓰지 않도록 해준다. 

 

grid 내 모든 thread들의 computation이 끝나면, compute kernel은 CPU에게 control을 넘긴다. 

그리고 나면 CPU는 GPU driver로부터 array를 copy해서 CPU로 memory back 하도록 한다.

 

 

하나의 CTA(=thread block) 안에 있는 thread들은 SM 당 공유되는 scratchpad memory를 통해 communicate할 수 있다. 

NVIDIA에서는 이를 shared memory, AMD의 GCN architecture에서는 LDS(local data store)라 한다. 하나의 SM 당 하나의 shared memory를 갖고 있다. shared memory의 공간은 SM 내에서 실행 중인 모든 CTA들 간에 divide된다.

CUDA에서는 __shared__를 사용해서 shared memory를 할당하고, software적으로 manage되는 cache처럼 사용하도록 한다. 그밖에도 GPU는 hardware managed cache를 포함하는데, 이러한 cache를 통한 accessing은 cache miss를 빈번하게 일으킨다. 따라서 빈번하게 reuse되고 predictable한 방식으로 data가 사용된다면 프로그래머는 shared memory를 사용하면 된다. AMD의 GCN architecture는 또한 GDS(global data store)를 포함하고 있어 GPU 내 모든 core들이 data를 공유하도록 한다. shared memory는 graphics application이 다른 graphics shader 간 result를 pass하도록 하는 데 사용된다.

 

CTA 내 thread들은 hadware-supported barrier instruction을 사용해서 효율적으로 synchronize할 수 있다.

다른 CTA에 있는 thread들은 서로 global address space를 통해 communicate할 수 있다. 하지만 이러한 방식의 접근은 시간적인 측면, 에너지 측면에서 shared memory에 비해 더 많은 cost를 사용하게 한다.

 

NVIDIA Kepler architecture을 통 CDP(CUDA Dynamic Parallelism)가 소개되었다.

CDP는 GPU에서 data intensive irregular application을 실행 때 사용되는 thread 간 load imbalance로 인해 GPU hardware가 underutilize될 수 있다는 관찰을 기반으로 하고 있다.

 

 

 

 

 

 

2.2 GPU instruction set architectures

 

High-level language(CUDA, OpenCL)로 작성된 compute kernel을 GPU hardware에서 실행되는 assembly level language로 변환

GPU 아키텍쳐는 CPU와는 약간 다른 방식으로 instruction set을 지원하도록 발전해왔다. 

 

 

backward compatibility

이전 세대 아키텍쳐에 대해 프로그램 변경 없이 다음 세대 아키텍쳐에서 해당 instruction set이 실행될 수 있도록 하는 것

(x86 micro processor는 1976년 출시된 Intel 8086과 backwards compatible을 유지)

 

GPU에서는 compute kernel을 high-level language로 작성하고 이를 hardware에서 실행되는 assembly level로 변환하는 과정이 중요하다.  따라서 GPU의 instruction set이 변할 때 이전에 작성된 high-level language의 code가 새로운 instruction set에 대응하도록 하기 위해선 backward compatibility가 필요하다.

 

 

NVIDIA GPU INSTRUCTION SET ARCHITECTURES

 

PTX(Parallel Thread Execution ISA)

NVIDIA가 2007년 초기 CUDA를 출시했을 때 함께 release된 GPU computing을 위한 high level instruction set architecture

- PTX는 ARM, MIPS, SPARC, ALPHA와 같은 RISC instruction set architecutre와 많은 면에서 유사하고 optimizing compiler에서 사용되는 intermediate representation과도 비슷한 특성을 갖고 있다. virtual register 기능을 사용해 high-level abstraction을 제공한다.

 

 

SASS(Streaming ASSembler)

PTX code가 GPU에서 실행되기 전에 해당 하드웨어에서 지원하는 실제 instruction set architecture로 PTX를 컴파일해야하는 level

- PTX에서 SASS로 변환하는 프로세스는 GPU driver 혹은 CUDA toolkit에서 제공하는 ptxas라는 프로그램을 사용해서 수행될 수 있다.

 

/* PTX code */
.visible .entry _Z5saxpyifPfS_(
.param .u32 _Z5saxpyifPfS__param_0,
.param .f32 _Z5saxpyifPfS__param_1,
.param .u64 _Z5saxpyifPfS__param_2,
.param .u64 _Z5saxpyifPfS__param_3
)
{
.reg .pred %p<2>;
.reg .f32 %f<5>;
.reg .b32 %r<6>;
.reg .b64 %rd<8>;


ld.param.u32 %r2, [_Z5saxpyifPfS__param_0];
ld.param.f32 %f1, [_Z5saxpyifPfS__param_1];
ld.param.u64 %rd1, [_Z5saxpyifPfS__param_2];
ld.param.u64 %rd2, [_Z5saxpyifPfS__param_3];
mov.u32 %r3, %ctaid.x;
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %tid.x;
mad.lo.s32 %r1, %r4, %r3, %r5;
setp.ge.s32 %p1, %r1, %r2;
@%p1 bra BB0_2;

cvta.to.global.u64 %rd3, %rd2;
cvta.to.global.u64 %rd4, %rd1;
mul.wide.s32 %rd5, %r1, 4;
add.s64 %rd6, %rd4, %rd5;
ld.global.f32 %f2, [%rd6];
add.s64 %rd7, %rd3, %rd5;
ld.global.f32 %f3, [%rd7];
fma.rn.f32 %f4, %f2, %f1, %f3;
st.global.f32 [%rd7], %f4;

BB0_2:
ret;
}

해당 PTX 코드는 위의 cuda kernel code에 해당하는 PTX 코드이다.

 

/* NVIDIA Fermi architecture - SASS code */
Address 				Dissassembly 					Encoded Instruction
======== =============================================== ========================
/*0000*/ 		MOV R1, c[0x1][0x100]; 					/* 0x2800440400005de4 */
/*0008*/ 		S2R R0, SR_CTAID.X; 					/* 0x2c00000094001c04 */
/*0010*/ 		S2R R2, SR_TID.X; 						/* 0x2c00000084009c04 */
/*0018*/ 		IMAD R0, R0, c[0x0][0x8], R2; 			/* 0x2004400020001ca3 */
/*0020*/ 		ISETP.GE.AND P0, PT, R0, c[0x0][0x20], PT; /* 0x1b0e40008001dc23 */
/*0028*/	 @P0 BRA.U 0x78; 							/* 0x40000001200081e7 */
/*0030*/ 	 @!P0 MOV32I R5, 0x4; 						/* 0x18000000100161e2 */
/*0038*/ 	 @!P0 IMAD R2.CC, R0, R5, c[0x0][0x28]; 	/* 0x200b8000a000a0a3 */
/*0040*/ 	 @!P0 IMAD.HI.X R3, R0, R5, c[0x0][0x2c]; 	/* 0x208a8000b000e0e3 */
/*0048*/ 	 @!P0 IMAD R4.CC, R0, R5, c[0x0][0x30]; 	/* 0x200b8000c00120a3 */
/*0050*/	 @!P0 LD.E R2, [R2]; 						/* 0x840000000020a085 */
/*0058*/	 @!P0 IMAD.HI.X R5, R0, R5, c[0x0][0x34]; 	/* 0x208a8000d00160e3 */
/*0060*/	 @!P0 LD.E R0, [R4]; 						/* 0x8400000000402085 */
/*0068*/	 @!P0 FFMA R0, R2, c[0x0][0x24], R0; 		/* 0x3000400090202000 */
/*0070*/	 @!P0 ST.E [R4], R0; /* 0x9400000000402085 */
/*0078*/		 EXIT; 									/* 0x8000000000001de7 */

위 코드는 NVIDIA Fermi 아키텍쳐 기반 SASS 코드이다. 

 

PTX, SASS level에서 비슷한 점

- RISC 사용 -> memory access에 load store instruction 사용

- Prediction 사용 -> 조건에 따라 instruction을 실행하거나 제외

 

PTX, SASS level에서 다른 점

- register set 크기 : PTX는 infinite한 register set을 사용하며 각 definition이 일반적으로 새로운 register를 사용한다.

- kernel parameter 전달 : PTX에서는 parameter가 자체적인 별도의 parameter address에 할당되지만 SASS에서는 kernel parameter가 banked constant memory를 통해 전달되어 load/store instruction을 사용하지 않고도 액세스 할 수 있다.

 

 

 

 

/* NVIDIA Pascal architecture - SASS code */
Address Dissassembly Encoded Instruction
======== =============================================== ========================
/* 0x001c7c00e22007f6 */
/*0008*/ MOV R1, c[0x0][0x20]; /* 0x4c98078000870001 */
/*0010*/ S2R R0, SR_CTAID.X; /* 0xf0c8000002570000 */
/*0018*/ S2R R2, SR_TID.X; /* 0xf0c8000002170002 */
/* 0x001fd840fec20ff1 */
/*0028*/ XMAD.MRG R3, R0.reuse, c[0x0] [0x8].H1, RZ; /* 0x4f107f8000270003 */
/*0030*/ XMAD R2, R0.reuse, c[0x0] [0x8], R2; /* 0x4e00010000270002 */
/*0038*/ XMAD.PSL.CBCC R0, R0.H1, R3.H1, R2; /* 0x5b30011800370000 */
/* 0x081fc400ffa007ed */
/*0048*/ ISETP.GE.AND P0, PT, R0, c[0x0][0x140], PT; /* 0x4b6d038005070007 */
/*0050*/ @P0 EXIT; /* 0xe30000000000000f */
/*0058*/ SHL R2, R0.reuse, 0x2; /* 0x3848000000270002 */
/* 0x081fc440fec007f5 */
/*0068*/ SHR R0, R0, 0x1e; /* 0x3829000001e70000 */
/*0070*/ IADD R4.CC, R2.reuse, c[0x0][0x148]; /* 0x4c10800005270204 */
/*0078*/ IADD.X R5, R0.reuse, c[0x0][0x14c]; /* 0x4c10080005370005 */
/* 0x0001c800fe0007f6 */
/*0088*/ IADD R2.CC, R2, c[0x0][0x150]; /* 0x4c10800005470202 */
/*0090*/ IADD.X R3, R0, c[0x0][0x154]; /* 0x4c10080005570003 */
/*0098*/ LDG.E R0, [R4]; } /* 0xeed4200000070400 */
/* 0x0007c408fc400172 */
/*00a8*/ LDG.E R6, [R2]; /* 0xeed4200000070206 */
/*00b0*/ FFMA R0, R0, c[0x0][0x144], R6; /* 0x4980030005170000 */
/*00b8*/ STG.E [R2], R0; /* 0xeedc200000070200 */
/* 0x001f8000ffe007ff */
/*00c8*/ EXIT; /* 0xe30000000007000f */
/*00d0*/ BRA 0xd0; /* 0xe2400fffff87000f */
/*00d8*/ NOP; /* 0x50b0000000070f00 */
/* 0x001f8000fc0007e0 */
/*00e8*/ NOP; /* 0x50b0000000070f00 */
/*00f0*/ NOP; /* 0x50b0000000070f00 */
/*00f8*/ NOP; /* 0x50b0000000070f00 */

 

위 코드는 NVIDIA Pascal 아키텍쳐 기반 SASS 코드이다. 

NVIDIA Fermi 아키텍쳐 기반 SASS 코드와 비교해서 instruction encoding을 포함한 NVIDIA의 ISA가 상당히 변경되었다. 

- control instruction :  dependency 확인을 instruction scoreboard를 사용하여 없애며, Maxwell에서 64 bit control instruction은 stall count, hit flag, read, wait dependency barrier에 대한 정보를 인코딩하는 21bit 그룹을 포함한다.

- regular instruction에 대한 register reuse flag 사용 : Maxwell 아키텍쳐부터 추가된 operand reuse cache를 나타냄.

이는 주된 register file access 마다 register 값을 여러 번 읽을 수 있게 해주어 에너지 소비를 줄이거나 성능을 향상시키도록  함

 

 

 

 

AMD GRAPHICS CORE NEXT INSTRUCTION SET ARCHITECTURE

 

AMD의 Southern Islands는 AMD의 Graphics Core Next(GCN) 아키텍쳐의 첫 번째 세대이다.

AMD의 compile flow에는 Heterogeneous System Architecture(HSA)의 일부로 알려진 virtual instruction set architecture인 HSAIL이 포함되어 있다.

 

 

AMD GCN 아키텍쳐와 NVIDIA GPU의 주요한 차이점은 별도의 scala 및 vector instruction이다. 

AMD GCN 아키텍쳐에서 각 SIMT core는 scalar unit과 4개의 vector unit이 결합되어 있다.

• scalar instruction : scalar unit에서 실행되며 wavefront의 모든 thread에 대해 공유되는 single 32bit 값을 계산한다. control flow 처리와 관련 있으며 SIMT 프로그램의 특정 부분이 thread ID에 독립적으로 동일한 결과를 계산하는 경우가 있다면  유용하게 사용된다.

• vector instruction : vector unit에서 실행되며 wavefront의 각 개별 thread에 대해 다른 32bit 값을 계산한다.

exec : SIMT 실행을 위해 개별 vector lane의 실행을 조건부로 사용하는 특수 register

긴 latency 작업에 대해 data dependency를 해결하기 위해 S_WAITCHT instruction을 사용해 wavefront가 특정 유형의 미완료된 작업 갯수가 threashold 아래로 내려갈 때까지 기다리게 할 수 있다.

각 wavefront에는 vector memory count, loacal/global data storage count, register export count가 있다.