https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#programming-model
Programming Guide :: CUDA Toolkit Documentation
Texture gather is a special texture fetch that is available for two-dimensional textures only. It is performed by the tex2Dgather() function, which has the same parameters as tex2D(), plus an additional comp parameter equal to 0, 1, 2, or 3 (see tex2Dgathe
docs.nvidia.com
커널(kernels)
CUDA C는 프로그래머가 커널(kernel)이라 불리는 C함수를 정의하도록 하여 C 언어를 확장한다. C 함수가 호출되어 한번 실행되는 데에 비해, 커널은 N개의 서로 다른 CUDA 스레드에 의해 N번 병렬적으로 수행된다.
커널은 __global__ 지정어를 통해 정의되고 커널을 실행할 CUDA 스레드의 수는 <<<...>>> 기호를 통해 특정된다. 커널을 실행하는 각 스레드는 유일한 thread ID를 부여받고, 내장변수(built-in variable)인 threadIdx를 통해 접근할 수 있다.
|
1
2
3
4
5
6
7
8
9
10
11
12
13
14
|
// Kernel definition of Vector addition (from CUDA c programming guide)
__global__ void VecAdd(float* A, float* B, float* C)
{
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main()
{
...
// Kernel invocation with N threads
VecAdd<<< 1, N >>>(A, B, C);
...
}
|
cs |
스레드 계층구조(Thread Hierarchy)
threadIdx 변수는 3차원 벡터이므로, 스레드는 각 차원에 따라 인덱스를 결정할 수 있고 또한 그 차원에 따라 thread block이 만들어진다.
스레드 인덱스와 스레드 ID는 직관적으로 연관되어 있다. (참고 : 첫 번째 스레드 블록에 대해서 생각한다면) 1 차원 스레드 블록의 스레드 인덱스는 스레드 ID와 동일하다. 크기가 (Dx, Dy)인 2차원 스레드 블록에서 인덱스가 (x, y)인 스레드의 스레드 ID는 ( x + y*Dx )이다. (참고: 그래픽에서 x, y는 행렬의 열(column)과 행(row)에 각각 해당된다는 점을 유의한다.) 마찬가지로 크기가 (Dx, Dy, Dz)인 3차원 스레드 블록에서 인덱스가 (x, y, z)인 스레드의 스레드 ID는 ( x + y*Dx + z*Dx*Dy )이다. N 정사각행렬의 두 행렬 곱으로 확장된 위의 코드는 다음과 같다.
|
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
|
// Kernel definition of Matrix addition (from CUDA c programming guide)
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if( i < N && j < N)
C[i][j] = A[i][j] + B[i][j];
}
int main()
{
...
// Kernel invocation with one block of (N * N * 1)threads
dim3 threadPerBlock(16,16);
dim3 numBlocks( N / threadPerBlock.x, N /threadPerBlock.y );
VecAdd<<< numBlocks , threadPerBlock >>>(A, B, C);
...
}
|
cs |
스레드 블록 안의 스레드 수는 제한이 있다. 한 스레드 블록 안의 모든 스레드들이 SM의 제한된 리소스를 함께 사용해야 하기 때문이다. 현재 GPU에서 하나의 스레드 블록은 최대 1024개의 스레드를 가질 수 있다.
그러나 커널은 동일한 형태를 가진 다수의 스레드 블록에 의해 실행될 수 있어서, 총 스레드의 수는 스레드 블록의 개수와 스레드 블록 당 스레드의 수를 곱한 것과 같다. 스레드 블록은 1/2/3차원 그리드 형태로 묶일 수 있다. 그리드 안의 스레드 블록 수는 대개 처리할 데이터의 수 또는 사용할 수 있는 프로세서 수에 의해 결정한다.
스레드 블록의 수와 블록 당 스레드 수는 <<<...>>> 안에서 int나 dim3 타입의 변수로 지정할 수 있다. 그리드 안에서 각 블록은 내장 변수인 blockIdx에 의해 접근될 수 있다. 또한 블록의 크기를 blockDim을 통해 커널이 알 수 있다. 위의 예제 코드는 다수의 스레드 블록을 사용하도록 다음과 같이 확장될 수 있다.
|
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
|
// Kernel definition of Matrix addition (from CUDA c programming guide)
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if( i < N && j < N)
C[i][j] = A[i][j] + B[i][j];
}
int main()
{
...
// Kernel invocation with one block of (N * N * 1)threads
dim3 threadPerBlock(16,16);
dim3 numBlocks( N / threadPerBlock.x, N /threadPerBlock.y );
VecAdd<<< numBlocks , threadPerBlock >>>(A, B, C);
...
}
|
cs |
스레드 블록은 독립적을 실행될 필요가 있다. 스레드 블록은 어떤 순서로도 실행시키는 것이 가능해야하고, 병렬 수행, 순차 실행이 가능하다. 이러한 독립성 덕분에 스레드블록은 수많은 코어 위에서 어떤 순서로도 스케줄이 될 수 있다.
블록 안의 스레드는 공유 메모리(shared memory)를 통해 데이터를 공유하고, 메모리 접근(memory access)을 함께 일어날 수 있도록 실행을 동기화하여서 협력적으로 동작할 수 있다. 사용자는 __syncthreads()를 호출하여서 커널 안에서 동기화 지점(synchronizing points)를 지정할 수 있다. __syncthreads( )는 블록 안의 모든 스레드가 다른 스레드의 동작이 끝나기를 기다려야 하는 배리어(barrier)로 동작한다. 효과적인 병렬수행을 위해서 공유 메모리는 (L1 캐시처럼) 프로세서에 인접한 low-latency 메모리이어야한다.
메모리 계층구조( Memory hierarachy)
CUDA 스레드는 실행과정에서 다양한 메모리 공간의 데이터에 접근할 수 있다. 각 스레드는 개별적인 로컬 메모리(private local memory)를 가질 수 있다. 각 스레드 블록은 블록 안의 모드 스레드가 접근할 수 있는 공유 메모리(shared memory)를 가질 수 있다. 공유 메모리 안의 데이터는 스레드 블록과 수명이 같다. 모든 스레드는 같은 글로벌 메모리(global memory)에 접근한다.
이외에도 스레드가 접근할 수 있는 두 가지 읽기전용(read-only) 메모리 공간이 있다. 상수메모리(constant memory)와 텍스처 메모리(texture memory), 글로벌 메모리(global memory)는 서로 다른 메모리 사용을 위해 최적화되어 있다.
이종 프로그래밍(Heterogenous Programming)
CUDA 프로그래밍 모델은 C 프로그램을 실행하는 호스트와 함께 실행되는 분리된 개별 디바이스 위에서 동작하는 CUDA 스레드를 가정한다. 예를 들어 커널은 GPU 위에서 실행되고, 프로그램의 나머지 부분은 CPU 위에서 동작하는 식이다. 또 호스트와 디바이스각 각각의 메모리 공간을 별도로 가지고 있다고 여긴다. (호스트 메모리(host memory), 디바이스 메모리(device memory))
따라서 프로그램은 커널이 볼 수 있는 글로벌, 상수, 텍스처 메모리 공간을 CUDA 런타임으로의 함수 호출로서 관리한다. 디바이스 메모리 할당과 해제, 호스트 메모리와 디바이스 메모리 간의 데이터 전송 등이 포함된다.
통합 메모리(Unified Memory)는 호스트와 디바이스 메모리를 연관지어주는 매니지드 메모리(managed memory)이다. 매니지드 메모리는 CPU와 GPU 모두에서 단일 주소공간을 가진 하나의 메모리로 접근되어진다(a single, coherent image with a common address space). 통합 메모리를 통해 디바이스 메모리의 용량보다 더 큰 데이터를 처리할 수 있게 된다(oversubscription).
연산 능력(Compute Capability)
디바이스의 연산능력(compute capability)는 버전 숫자로 표현되고, 때때로 SM version으로 표현되기도 한다. 이 버전 숫자를 통해 GPU 하드웨어에 의해 지원되는 기능을 알 수 있고, 애플리케이션이 어떤 하드웨어 기능을 활용할지 런타임에 결정할 수 있으며, 특정 명령어가 사용하는 GPU에서 실행 가능한 지를 알려준다. 연산능력은 메이저 버전 번호 X와 마이너 버전 번호Y를 조합하여 X.Y로 표시된다. 같은 메이저 버전 번호의 디바이스는 같은 코어 아키텍처를 가지고 있다. ( 7 - 볼타, 6-파스칼, 5-맥스웰, 3-케플러, 2-페르미, 1-테슬라
'CUDA 프로그래밍' 카테고리의 다른 글
| SAXPY(Single-precision A-X Plus Y) (0) | 2020.03.03 |
|---|---|
| MPS란? (0) | 2020.01.29 |
| Unified memory를 사용할 때, 공유변수의 동작 예제 (0) | 2020.01.08 |
| PTX 어셈블리를 활용하여 실행 중인 SM 확인하기 (1) | 2019.10.28 |
| Rodinia 벤치마크 - 개요 (0) | 2019.10.26 |
댓글