GPU에서 CUDA 또는 PTX 코드 실행하기

CUDA 툴킷 요구 사항

CU 코드에서 CUDA^® 커널 객체를 생성하거나 GPU Coder™를 사용하여 CUDA 호환 소스 코드, 라이브러리 및 실행 파일을 컴파일하려면 CUDA 툴킷을 설치해야 합니다. CUDA 툴킷에는 컴파일하기 위한 CUDA 라이브러리와 툴이 포함되어 있습니다.

참고

GPU에서 MATLAB^® 함수를 실행하거나 CUDA 지원 MEX 함수를 생성하기 위해 툴킷이 필요하지는 않습니다.

작업 요구 사항

작업	요구 사항
`gpuArray` 및 GPU 지원 MATLAB 함수를 사용합니다. GPU Coder 또는 `mexcuda`를 사용하여 CUDA 지원 MEX 함수를 컴파일합니다.	NVIDIA 드라이버 다운로드에서 최신 그래픽스 드라이버를 다운로드합니다. CUDA 툴킷이 필요하지 않습니다.
CU 코드에서 CUDA 커널 객체를 만듭니다.* GPU Coder를 사용하여 CUDA 호환 소스 코드, 라이브러리 및 실행 가능 파일을 컴파일합니다.	MATLAB 릴리스에서 지원되는 CUDA 툴킷 버전을 설치합니다.

gpuArray 및 GPU 지원 MATLAB 함수를 사용합니다.
GPU Coder 또는 mexcuda를 사용하여 CUDA 지원 MEX 함수를 컴파일합니다.

NVIDIA 드라이버 다운로드에서 최신 그래픽스 드라이버를 다운로드합니다.

CUDA 툴킷이 필요하지 않습니다.

CU 코드에서 CUDA 커널 객체를 만듭니다.*
GPU Coder를 사용하여 CUDA 호환 소스 코드, 라이브러리 및 실행 가능 파일을 컴파일합니다.

MATLAB 릴리스에서 지원되는 CUDA 툴킷 버전을 설치합니다.

* MATLAB에서 CUDA 커널 객체를 만들려면 CU 파일과 해당 PTX 파일이 둘 다 있어야 합니다. CU 파일에서 PTX 파일을 컴파일하려면 CUDA 툴킷이 필요합니다. 해당 PTX 파일이 이미 있는 경우에는 툴킷이 필요하지 않습니다. CUDA 툴킷에서 지원되는 모든 컴파일러가 MATLAB에서 지원되는 것은 아닙니다.

필요한 툴킷 버전은 사용 중인 MATLAB 버전에 따라 달라집니다. 다음 표를 참고하여 현재 사용 중인 MATLAB 버전과 호환되는 툴킷 버전을 확인하십시오. NVIDIA^®의 모든 업데이트와 패치를 포함하여 지원되는 툴킷의 최신 버전을 사용하는 것이 가장 좋습니다.

MATLAB 릴리스	CUDA 툴킷 버전
R2022b	11.2
R2022a	11.2
R2021b	11.0
R2021a	11.0
R2020b	10.2
R2020a	10.1
R2019b	10.1
R2019a	10.0
R2018b	9.1
R2018a	9.0
R2017b	8.0
R2017a	8.0
R2016b	7.5
R2016a	7.5
R2015b	7.0
R2015a	6.5
R2014b	6.0
R2014a	5.5
R2013b	5.0
R2013a	5.0
R2012b	4.2
R2012a	4.0
R2011b	4.0

CUDA 툴킷에 대한 자세한 내용을 알아보고 지원되는 버전을 다운로드하려면 CUDA Toolkit Archive(NVIDIA)를 참조하십시오.

CUDAKernel 워크플로 개요

여기에서는 CU 파일이나 PTX(Parallel Thread Execution) 파일에서 실행 가능한 커널을 만들고 이 커널을 MATLAB의 GPU에서 실행하는 방법을 설명합니다. 커널은 MATLAB에서 CUDAKernel 객체로 표현되며 이 객체는 MATLAB 배열이나 gpuArray 변수를 기반으로 동작을 수행할 수 있습니다.

다음 단계는 CUDAKernel의 일반적인 워크플로를 설명하고 있습니다.

컴파일된 PTX 코드를 사용하여 GPU 실행 코드가 포함된 CUDAKernel 객체를 만듭니다.
GPU에서의 실행을 제어할 수 있도록 CUDAKernel 객체에 대한 속성을 설정합니다.
CUDAKernel에 대해 필요한 입력값을 지정하고 feval을 호출하여 GPU에서 커널을 실행합니다.

이러한 단계를 수행하는 MATLAB 코드는 다음과 같습니다.

% 1. Create CUDAKernel object.
k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu','entryPt1');

% 2. Set object properties.
k.GridSize = [8 1];
k.ThreadBlockSize = [16 1];

% 3. Call feval with defined inputs.
g1 = gpuArray(in1); % Input gpuArray.
g2 = gpuArray(in2); % Input gpuArray.

result = feval(k,g1,g2);

다음 섹션에서는 이러한 명령과 워크플로 단계에 대해 자세히 설명합니다.

CU 파일에서 PTX 파일 컴파일하기

GPU에서 실행하려는 CU 파일이 있으면 먼저 CU 파일을 컴파일해서 PTX 파일을 만들어야 합니다. 이 작업을 수행하는 한 가지 방법은 NVIDIA CUDA 툴킷의 nvcc 컴파일러를 사용하는 것입니다. 예를 들어, CU 파일이 myfun.cu인 경우 다음 셸 명령을 사용하여 컴파일된 PTX 파일을 만들 수 있습니다.

nvcc -ptx myfun.cu

이렇게 하면 myfun.ptx라는 파일이 생성됩니다.

CU 파일 입력값으로 CUDAKernel 객체 생성하기

.cu 파일과 .ptx 파일로 MATLAB에서 CUDAKernel 객체를 만들어서 이 객체를 커널을 실행하는 데 사용할 수 있습니다.

k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu');

참고

CUDAKernel 객체에 대해 save 또는 load를 수행할 수 없습니다.

C 프로토타입 입력값으로 CUDAKernel 객체 생성하기

PTX 파일에 대응되는 CU 파일이 없는 경우 CU 파일 대신 C 커널에 대해 C 프로토타입을 지정할 수 있습니다. 예를 들면 다음과 같습니다.

k = parallel.gpu.CUDAKernel('myfun.ptx','float *, const float *, float');

C 프로토타입 입력값을 사용하는 또 다른 경우는 지원되는 데이터형이 소스 코드가 인식할 수 없는 변경된 이름을 사용하는 경우입니다. 아래 지원되는 유형을 참조하십시오. 커널이 다음 코드로 구성된다고 가정하겠습니다.

typedef float ArgType;
__global__ void add3( ArgType * v1, const ArgType * v2 )
{
    int idx = threadIdx.x;
    v1[idx] += v2[idx];
}

ArgType 자체는 지원되는 데이터형으로 인식되지 않습니다. 따라서 MATLAB에서 CUDAKernel 객체를 만들 때 이 데이터형을 포함하는 CU 파일을 그대로 입력값으로 사용할 수 없습니다. 그러나 add3 커널에 지원되는 입력 유형을 C 프로토타입 입력값으로 CUDAKernel 생성자에 지정할 수 있습니다. 예를 들면 다음과 같습니다.

k = parallel.gpu.CUDAKernel('test.ptx','float *, const float *','add3');

지원되는 데이터형

지원되는 C/C++ 표준 데이터형이 아래 표에 나와 있습니다.

Float형 정수형 부울형 및 문자형

Float형	정수형	부울형 및 문자형
`double`, `double2` `float`, `float2`	`short`, `unsigned short`, `short2`, `ushort2` `int`, `unsigned int`, `int2`, `uint2` `long`, `unsigned long`, `long2`, `ulong2` `long long`, `unsigned long long`, `longlong2`, `ulonglong2` `ptrdiff_t`, `size_t`	`bool` `char`, `unsigned char`, `char2`, `uchar2`

double, double2

float, float2

short, unsigned short, short2, ushort2

int, unsigned int, int2, uint2

long, unsigned long, long2, ulong2

long long, unsigned long long, longlong2, ulonglong2

ptrdiff_t, size_t

bool

char, unsigned char, char2, uchar2

또한 다음 정수형은 프로그램에 tmwtypes.h 헤더 파일을 포함하는 경우에 지원됩니다.

정수형

정수형
`int8_T`, `int16_T`, `int32_T`, `int64_T` `uint8_T`, `uint16_T`, `uint32_T`, `uint64_T`

int8_T, int16_T, int32_T, int64_T

uint8_T, uint16_T, uint32_T, uint64_T

헤더 파일은 matlabroot/extern/include/tmwtypes.h로 제공됩니다. 다음 라인을 사용하여 프로그램에 파일을 포함합니다.

#include "tmwtypes.h"

인수 제한 사항

모든 입력값은 스칼라나 포인터가 될 수 있으며 const로 레이블이 지정될 수 있습니다.

커널의 C 선언은 항상 다음과 같은 형식입니다.

__global__ void aKernel(inputs ...)

커널은 아무 것도 반환하지 않아야 하며 커널의 입력 인수(스칼라 또는 포인터)에 대해서만 동작해야 합니다.
커널은 어떤 형식의 메모리도 할당할 수 없으므로 모든 출력값은 커널이 실행되기 전에 미리 할당되어야 합니다. 따라서 커널을 실행하기 전에 모든 출력값의 크기를 알고 있어야 합니다.
원칙적으로 커널로 전달되는 const가 아닌 모든 포인터는 출력 데이터를 포함할 수 있는데, 이는 커널의 다수의 스레드가 이 데이터를 수정할 수 있기 때문입니다.

C의 커널 정의를 MATLAB으로 변환할 때 다음에 유의합니다.

C의 모든 스칼라 입력값(double, float, int 등)은 MATLAB의 스칼라이거나 스칼라(즉, 단일 요소) gpuArray 변수여야 합니다.
C의 모든 const 포인터 입력값(const double * 등)은 MATLAB에서 스칼라이거나 행렬일 수 있습니다. 이러한 입력값은 올바른 유형으로 형변환되어 장치에 복사되고, 첫 번째 요소에 대한 포인터가 커널로 전달됩니다. 원래 크기에 대한 정보는 커널로 전달되지 않습니다. 이는 마치 커널이 mxArray에 대한 mxGetData 결과를 직접 받는 것과 같습니다.
C의 상수가 아닌 모든 포인터 입력값은 똑같이 상수가 아닌 포인터로 커널로 전송됩니다. 그러나 상수가 아닌 포인터는 커널에서 변경될 수 있기 때문에 커널의 출력값으로 간주됩니다.
MATLAB 작업 공간에서의 스칼라, 배열 입력값은 요청된 유형으로 형변환된 다음 커널로 전달됩니다. 그러나 gpuArray 입력값은 자동으로 형변환되지 않으므로 유형과 복잡도가 예상과 정확하게 일치해야 합니다.

이러한 규칙에는 몇 가지 의미가 있습니다. 가장 주목할 만한 것은 (GPU에서 메모리를 할당할 수 없어서) 입력값으로 출력값의 크기를 정의하므로 커널의 모든 출력값은 반드시 커널에 대한 입력값도 되어야 한다는 것입니다.

CUDAKernel 객체 속성

커널 객체를 종료 세미콜론 없이 만들거나 명령줄에 객체 변수를 입력하면 MATLAB에서 커널 객체 속성을 표시합니다. 예를 들면 다음과 같습니다.

k = parallel.gpu.CUDAKernel('conv.ptx','conv.cu')

k = 
  parallel.gpu.CUDAKernel handle
  Package: parallel.gpu

  Properties:
     ThreadBlockSize: [1 1 1]
  MaxThreadsPerBlock: 512
            GridSize: [1 1 1]
    SharedMemorySize: 0
          EntryPoint: '_Z8theEntryPf'
  MaxNumLHSArguments: 1
     NumRHSArguments: 2
       ArgumentTypes: {'in single vector'  'inout single vector'}

커널 객체의 속성은 실행 동작의 일부를 제어합니다. 점 표기법을 사용하면 변경이 허용되는 속성을 변경할 수 있습니다.

객체 속성에 대한 설명은 CUDAKernel 객체 참조 페이지를 참조하십시오. 설정 가능한 속성을 수정하는 일반적인 이유는 아래 설명된 대로 스레드 수를 지정하기 위해서입니다.

진입점 지정하기

PTX 파일에 여러 개의 진입점이 포함된 경우 myfun.ptx에서 커널 객체 k가 참조할 특정 커널을 식별할 수 있습니다.

k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu','myKernel1');

하나의 PTX 파일이 여러 다른 커널에 대한 여러 진입점을 포함할 수 있습니다. 이러한 각 진입점마다 고유한 이름이 있습니다. 이러한 이름은 일반적으로 변형(mangled)되어 있습니다(C++ mangling과 유사). 그러나 nvcc로 생성된 PTX의 이름에는 항상 CU 파일의 원래 함수 이름이 포함됩니다. 예를 들어, CU 파일이 커널 함수를 다음과 같이 정의한다고 가정합니다.

__global__ void simplestKernelEver( float * x, float val )

그러면 PTX 코드에는 _Z18simplestKernelEverPff라는 진입점이 포함됩니다.

진입점이 여러 개 있으면 CUDAKernel을 호출하여 커널을 생성할 때 특정 커널의 진입점을 지정합니다.

참고

CUDAKernel 함수는 PTX 파일에서 진입점 이름을 검색하여 부분문자열이 일치하는 경우를 모두 확인합니다. 따라서 진입점의 이름을 다른 진입점의 부분문자열을 사용해서 지정하면 안 됩니다.

원래 진입점 이름을 제어할 수 없을 수도 있습니다. 이런 경우에는 각각으로부터 파생하여 변형된(mangled) 고유 이름을 알고 있어야 합니다. 예를 들어, 다음과 같은 함수 템플릿이 있다고 가정하겠습니다.

template <typename T>
__global__ void add4( T * v1, const T * v2 )
{
    int idx = threadIdx.x;
    v1[idx] += v2[idx];
}

템플릿이 float형과 double형으로 확장되면 두 개의 진입점이 생성되고 둘 다 부분문자열 add4를 포함합니다.

template __global__ void add4<float>(float *, const float *);
template __global__ void add4<double>(double *, const double *);

PTX는 다음과 같이 대응하는 진입점을 가집니다.

_Z4add4IfEvPT_PKS0_
_Z4add4IdEvPT_PKS0_

float형 버전에는 진입점 add4If를 사용하고 double형 버전에는 진입점 add4Id를 사용합니다.

k = parallel.gpu.CUDAKernel('test.ptx','double *, const double *','add4Id');

스레드 수 지정하기

CUDAKernel의 객체 속성 중 다음 두 가지를 설정하여 CUDAKernel의 계산 스레드 수를 지정합니다.

GridSize — 3개 요소로 구성된 벡터로, 이들 요소를 곱한 값이 블록 수를 결정합니다.
ThreadBlockSize — 3개 요소로 구성된 벡터로, 이들 요소를 곱한 값이 블록당 스레드 수를 결정합니다. (이렇게 곱한 값이 속성 MaxThreadsPerBlock의 값을 초과할 수 없습니다.)

이 두 속성의 디폴트 값은 [1 1 1]이지만 500개의 스레드를 사용하여 500개의 요소로 구성된 벡터에 대해 요소별 연산을 병렬로 실행한다고 가정합니다. 이 작업을 수행하는 간단한 방법은 CUDAKernel을 만들고 속성을 적절하게 설정하는 것입니다.

k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu');
k.ThreadBlockSize = [500,1,1];

일반적으로 입력값의 크기를 기반으로 그리드 크기와 스레드 블록 크기를 설정합니다. 스레드 계층 구조 및 다차원 그리드와 다차원 블록에 대한 자세한 내용은 NVIDIA CUDA C Programming Guide를 참조하십시오.

CUDAKernel 실행하기

작업 공간 변수 사용하기
gpuArray 변수 사용하기
입력 및 출력의 대응 관계 결정하기

GPU에서 CUDAKernel을 실행하려면 feval 함수를 사용하십시오. 다음 예제에서는 커널을 MATLAB 작업 공간 변수를 사용해서 실행하는 방법과 gpuArray 변수를 사용해서 실행하는 방법을 보여줍니다.

작업 공간 변수 사용하기

네이티브 언어로 일부 커널을 이미 작성했으며 MATLAB에서 이 커널을 사용하여 GPU에서 실행하려 한다고 가정하겠습니다. 두 벡터에 대해 컨벌루션을 수행하는 커널이 있습니다. 두 개의 확률 벡터를 입력값으로 사용하여 이 커널을 불러와서 실행합니다.

k = parallel.gpu.CUDAKernel('conv.ptx','conv.cu');

result = feval(k,rand(100,1),rand(100,1));

MATLAB 작업 공간 데이터에 대한 입력값이 상수 또는 변수인 경우에도 출력값은 gpuArray입니다.

gpuArray 변수 사용하기

커널을 실행할 때 gpuArray 객체를 입력값으로 사용하는 것이 더 효율적일 수 있습니다.

k = parallel.gpu.CUDAKernel('conv.ptx','conv.cu');

i1 = gpuArray(rand(100,1,'single'));
i2 = gpuArray(rand(100,1,'single'));

result1 = feval(k,i1,i2);

출력값이 gpuArray이기 때문에 MATLAB 작업 공간과 GPU 간에 추가 전송 없이도 이 입력 데이터와 출력 데이터를 사용하여 다른 연산을 수행할 수 있습니다. 모든 GPU 연산이 완료되면 최종 결과 데이터를 MATLAB 작업 공간으로 수집하십시오.

result2 = feval(k,i1,i2);

r1 = gather(result1);
r2 = gather(result2);

입력 및 출력의 대응 관계 결정하기

[out1, out2] = feval(kernel, in1, in2, in3)을 호출하면 입력 in1, in2, in3은 CU 파일 내 C 함수의 각 입력 인수에 대응합니다. 출력 out1과 out2는 C 커널이 실행된 후 C 함수에 대한 첫 번째와 두 번째의 상수가 아닌 포인터 입력 인수 값을 저장합니다.

예를 들어, CU 파일 내 C 커널이 다음 시그니처를 갖는 경우,

void reallySimple( float * pInOut, float c )

대응하는 MATLAB의 커널 객체(k)는 다음 속성을 갖습니다.

MaxNumLHSArguments: 1
   NumRHSArguments: 2
     ArgumentTypes: {'inout single vector'  'in single scalar'}

따라서 feval을 사용하여 이 코드에서 커널 객체를 사용하려면 커널 객체 외에 2개의 feval 입력 인수를 제공해야 하며 하나의 출력 인수를 사용할 수 있습니다.

y = feval(k,x1,x2)

입력값 x1과 x2는 C 함수 프로토타입의 pInOut과 c에 대응합니다. 출력 인수 y는 C 커널이 실행된 후 C 함수 프로토타입의 pInOut 값에 대응합니다.

다음은 상수 포인터와 상수가 아닌 포인터의 조합을 보여주는 좀 더 복잡한 예제입니다.

void moreComplicated( const float * pIn, float * pInOut1, float * pInOut2 )

그러면 대응하는 MATLAB의 커널 객체는 다음 속성을 갖습니다.

MaxNumLHSArguments: 2
   NumRHSArguments: 3
     ArgumentTypes: {'in single vector'  'inout single vector'  'inout single vector'}

이 코드의 커널(k)에 feval을 사용하려면 다음 구문처럼 합니다.

[y1,y2] = feval(k,x1,x2,x3)

3개의 입력 인수 x1, x2 및 x3은 C 함수로 전달되는 3개의 인수에 대응합니다. 출력 인수 y1과 y2는 C 커널이 실행된 후 pInOut1과 pInOut2의 값에 대응합니다.

커널 워크플로 완료하기

2개의 숫자 추가하기
2개의 벡터 추가하기
CU 파일과 PTX 파일을 사용한 예제

2개의 숫자 추가하기

이 예제에서는 GPU에서 2개의 double형을 함께 추가합니다. NVIDIA CUDA 툴킷이 설치되어 있고 장치에 맞는 CUDA 지원 드라이버가 있어야 합니다.

이 작업을 수행하는 CU 코드는 다음과 같습니다.
```
__global__ void add1( double * pi, double c ) 
{
    *pi += c;
}
```
지시문 __global__은 이 코드가 커널에 대한 진입점임을 나타냅니다. 코드는 포인터를 사용하여 입력값이자 출력값인 pi의 결과를 보냅니다. 이 코드를 현재 디렉터리의 test.cu라는 파일에 넣습니다.
셸 명령줄에서 CU 코드를 컴파일하여 test.ptx라는 PTX 파일을 생성합니다.
```
nvcc -ptx test.cu
```
MATLAB에서 커널을 만듭니다. 현재 이 PTX 파일에는 진입점이 하나뿐이므로 지정할 필요가 없습니다. 커널을 더 추가하려면 add1을 진입점으로 지정해야 합니다.
```
k = parallel.gpu.CUDAKernel('test.ptx','test.cu');
```
2개의 숫자형 입력값으로 커널을 실행합니다. 기본적으로 커널은 하나의 스레드에서 실행됩니다.
```
result = feval(k,2,3)
```
```
result = 
    5
```

2개의 벡터 추가하기

이 예제에서는 이전 예제를 확장하여 2개의 벡터를 함께 추가합니다. 간단하게 벡터의 요소와 정확히 같은 수의 스레드가 있고 스레드 블록은 하나만 있다고 가정합니다.

CU 코드는 바로 이전 예제와 약간 다릅니다. 두 입력값 모두 포인터이며 변경하지 않았기 때문에 하나는 상수입니다. 각 스레드는 이들 요소를 해당 스레드 인덱스에 추가하기만 할 것입니다. 스레드 인덱스는 이 스레드가 추가해야 할 요소를 파악해야 합니다. (이러한 스레드 관련 값과 블록 관련 값을 가져오는 것은 CUDA 프로그래밍에서 매우 일반적인 패턴입니다.)
```
__global__ void add2( double * v1, const double * v2 ) 
{
    int idx = threadIdx.x;
    v1[idx] += v2[idx];
}
```
이 코드를 파일 test.cu에 저장합니다.
nvcc를 사용하여 위에서처럼 컴파일합니다.
```
nvcc -ptx test.cu
```
이 코드를 동일한 CU 파일에 첫 번째 예제의 코드와 함께 추가했다면 이번에는 구분을 위해 진입점 이름을 지정해야 합니다.
```
k = parallel.gpu.CUDAKernel('test.ptx','test.cu','add2');
```

커널을 실행하기 전에 추가할 벡터에 대한 스레드 수를 올바르게 설정합니다.

N = 128;
k.ThreadBlockSize = N;
in1 = ones(N,1,'gpuArray');
in2 = ones(N,1,'gpuArray');
result = feval(k,in1,in2);

CU 파일과 PTX 파일을 사용한 예제

CUDA를 사용하는 방법을 보여주는 예제를 확인하고 사용자가 시도해 볼 수 있는 CU 파일과 PTX 파일을 받으려면 Illustrating Three Approaches to GPU Computing: The Mandelbrot Set 항목을 참조하십시오.

참고 항목

mexcuda | CUDAKernel