Main Content

이 번역 페이지는 최신 내용을 담고 있지 않습니다. 최신 내용을 영문으로 보려면 여기를 클릭하십시오.

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

CUDAKernel 워크플로 개요

이 페이지에서는 CUDA® C++ 소스 파일(CU) 파일에서 실행 가능한 커널을 만들고 MATLAB®의 GPU에서 해당 커널을 실행하는 방법을 설명합니다. 커널은 MATLAB에서 CUDAKernel 객체로 표현되며 이 객체는 호스트 메모리에 저장된 배열 또는 GPU 배열에 대해 연산을 수행할 수 있습니다.

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

  1. mexcuda를 사용하여 CU 파일에서 PTX(Parallel Thread Execution) 파일을 컴파일합니다. mexcuda를 사용하여 PTX 파일을 컴파일하는 데 CUDA 툴킷이 필요하지 않습니다.

    R2023a 이전: mexcuda 함수 대신 NVIDIA® CUDA 툴킷의 nvcc 컴파일러를 사용하여 PTX 파일을 컴파일합니다.

  2. parallel.gpu.CUDAKernel 함수를 사용하여 CU 파일과 PTX 파일에서 CUDAKernel 객체를 만듭니다. CUDAKernel은 GPU 실행 코드를 포함합니다.

  3. GPU에서의 실행을 제어할 수 있도록 CUDAKernel에 대한 속성을 설정합니다.

  4. CUDAKernel에 대해 필요한 입력값을 지정하고 feval을 호출하여 GPU에서 커널을 실행합니다.

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

% 1. Compile a PTX file.
mexcuda -ptx myfun.cu

% 2. Create CUDAKernel object.
k = parallel.gpu.CUDAKernel("myfun.ptx","myfun.cu");

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

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

result = feval(k,g1,g2);

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

CUDAKernel 객체 만들기

GPU에서 실행하려는 CU 파일이 있으면 먼저 CU 파일을 컴파일해서 PTX 파일을 만들어야 합니다. PTX 파일을 컴파일하려면 CU 파일을 -ptx 플래그와 함께 mexcuda에 전달합니다.

mexcuda -ptx myfun.cu

그러면 PTX 파일 myfun.ptx가 생성됩니다.

CU 파일과 PTX 파일을 사용하여 CUDAKernel 객체를 만듭니다.

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

참고

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

CUDAKernel 객체 속성

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 객체의 속성은 실행 동작의 일부를 제어합니다. 점 표기법을 사용하면 변경이 허용되는 속성을 변경할 수 있습니다. 객체 속성에 대한 설명은 CUDAKernel 항목을 참조하십시오. 설정 가능한 속성을 수정하는 일반적인 이유는 아래 설명된 대로 스레드 수를 지정하기 위해서입니다.

진입점 지정하기

하나의 PTX 파일이 여러 다른 커널에 대한 여러 진입점을 포함할 수 있습니다. 이러한 각 진입점마다 고유한 이름이 있습니다. 각 진입점의 이름은 C++ 변형(mangling)에서와 같이 변형되지만 항상 CU 파일의 원래 함수 이름을 포함합니다. 예를 들어, CU 파일이 커널 함수를 다음과 같이 정의한다고 가정합니다.

__global__ void simplestKernelEver( float * x, float val )

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

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

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

참고

parallel.gpu.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의 객체 속성 중 다음 두 가지를 설정하여 계산 스레드 수를 지정합니다.

  • GridSize — 3개 요소로 구성된 벡터로, 이들 요소를 곱한 값이 블록 수를 결정합니다.

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

이 두 속성의 디폴트 값은 [1 1 1]이지만 500개의 스레드를 사용하여 500개의 요소로 구성된 벡터에 대해 요소별 연산을 병렬로 실행한다고 가정합니다. 요소의 곱이 500이 되도록 ThreadBlockSize를 설정합니다.

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

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

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형정수형부울형 및 문자형

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

헤더 파일은 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 double * 등)은 MATLAB에서 스칼라이거나 행렬일 수 있습니다. 이러한 입력값은 올바른 유형으로 형변환되어 장치에 복사되고, 첫 번째 요소에 대한 포인터가 커널로 전달됩니다. 원래 크기에 대한 정보는 커널로 전달되지 않습니다. 이는 마치 커널이 mxArray에 대한 mxGetData 결과를 직접 받는 것과 같습니다.

  • C의 상수가 아닌 모든 포인터 입력값은 똑같이 상수가 아닌 포인터로 커널로 전송됩니다. 그러나 상수가 아닌 포인터는 커널에서 변경될 수 있기 때문에 커널의 출력값으로 간주됩니다.

  • MATLAB 작업 공간에서의 스칼라, 배열 입력값은 요청된 유형으로 형변환된 다음 커널로 전달됩니다. 그러나 gpuArray 입력값은 자동으로 형변환되지 않으므로 데이터형과 실수/복소수 여부가 예상과 정확하게 일치해야 합니다.

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

CUDAKernel 실행하기

GPU에서 CUDAKernel을 계산하려면 feval 함수를 사용합니다.

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

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

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

입력값이 그렇지 않더라도 출력값은 gpuArray입니다. 그러나 커널을 실행할 때 gpuArray 객체를 입력값으로 사용하는 것이 더 효율적일 수 있습니다.

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

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

result1 = feval(k,i1,i2);

출력값이 gpuArray이기 때문에 GPU 메모리와 호스트 메모리 간에 추가 전송 없이도 이 입력 데이터와 출력 데이터를 사용하여 다른 연산을 수행할 수 있습니다.

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

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

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

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)

입력값 x1x2는 함수 프로토타입의 pInOutc에 대응합니다. 출력 인수 y는 커널이 실행된 후 함수 프로토타입의 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'}

3개의 입력 인수와 2개의 출력 인수를 사용하여 이 코드의 커널(k)에서 feval을 사용할 수 있습니다.

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

3개의 입력 인수 x1, x2, x3은 함수로 전달되는 3개의 인수에 대응합니다. 출력 인수 y1y2은 커널이 실행된 후 pInOut1pInOut2의 값에 대응합니다.

커널 워크플로 완료하기

2개의 숫자 추가하기

이 예제에서는 GPU에서 2개의 double형을 함께 추가합니다.

  1. 이 작업을 수행하는 CU 코드는 다음과 같습니다.

    __global__ void add1( double * a, double b ) 
    {
        *a += b;
    }

    지시문 __global__은 이 코드가 커널에 대한 진입점임을 나타냅니다. 코드는 포인터를 사용하여 입력값이자 출력값인 a의 결과를 보냅니다. 이 코드를 현재 디렉터리의 test.cu라는 파일에 저장합니다.

  2. mexcuda를 사용하여 CU 코드를 컴파일하고 test.ptx라는 PTX 파일을 생성합니다.

    mexcuda -ptx test.cu
  3. MATLAB에서 커널을 만듭니다. 현재 이 PTX 파일에는 진입점이 하나뿐이므로 지정할 필요가 없습니다. PTX 파일에 둘 이상의 커널 진입점이 포함된 경우 add1을 진입점으로 지정합니다.

    k = parallel.gpu.CUDAKernel("test.ptx","test.cu");
  4. 2개의 숫자형 입력값으로 커널을 실행합니다. 기본적으로 커널은 하나의 스레드에서 실행됩니다.

    result = feval(k,2,3)
    result = 
        5
    

2개의 벡터 추가하기

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

  1. CU 코드는 바로 이전 예제와 약간 다릅니다. 두 입력값 모두 포인터이며 변경하지 않았기 때문에 하나는 상수입니다. 각 스레드는 이들 요소를 해당 스레드 인덱스에 추가하기만 할 것입니다. 스레드 인덱스는 이 스레드가 추가해야 할 요소를 파악해야 합니다. 이러한 스레드 관련 값과 블록 관련 값을 가져오는 것은 CUDA 프로그래밍에서 매우 일반적인 패턴입니다.

    __global__ void add2( double * v1, const double * v2 ) 
    {
        int idx = threadIdx.x;
        v1[idx] += v2[idx];
    }

    이 코드를 파일 test.cu에 저장합니다.

  2. mexcuda를 사용하여 CU 파일에서 test.ptx라는 PTX 파일을 컴파일합니다.

    mexcuda -ptx test.cu
  3. 이 코드를 동일한 CU 파일에 첫 번째 예제의 코드와 함께 추가했다면 이번에는 구분을 위해 진입점 이름을 지정해야 합니다.

    k = parallel.gpu.CUDAKernel("test.ptx","test.cu","add2");
    
  4. 커널을 실행하기 전에 추가할 벡터에 대한 스레드 수를 올바르게 설정합니다.

    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 항목을 참조하십시오.

참고 항목

| |

관련 항목