GPU에서CUDA또는PTX코드실행하기- MATLAB和Simulink金宝app MathWorks한국

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

개요

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

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

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

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

% 1。创建CUDAKernel对象。k = parallel.gpu.CUDAKernel (“myfun.ptx”，“myfun.cu”，“entryPt1”）;% 2。设置对象的属性。k.GridSize = [8 1];k.ThreadBlockSize = [16 1];% 3。使用定义的输入调用feval。g1 = gpuArray (in);输入gpuArray %。g2 = gpuArray (in2);输入gpuArray %。结果=函数宏指令(k, g1, g2);

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

CUDAKernel객체만들기

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

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

nvcc -ptx myfun.cu

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

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

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

k = parallel.gpu.CUDAKernel (“myfun.ptx”，“myfun.cu”）;

참고

CUDAKernel객체에대해保存또는负载를수행할수없습니다。

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

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

k = parallel.gpu.CUDAKernel (“myfun.ptx”，'float *， const float *， float'）;

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

typedef ArgType浮动;__global__ void add3(ArgType * v1, const ArgType * v2) {int idx = threadadidx .x;[idx] v1 + v2 = [idx];｝

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

k = parallel.gpu.CUDAKernel (“test.ptx”，'float *， const float *'，“add3”）;

지원되는데이터형

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

浮动형 정수형 부울형및문자형

浮动형	정수형	부울형및문자형
`双`，`double2` `浮动`，`float2`	`短`，`无符号短`，`short2`，`ushort2` `int`，`无符号整型`，`int2`，`uint2` `长`，`无符号长`，`马龙`，`ulong2` `很久很久`，`无符号长时间长`，`longlong2`，`ulonglong2` `ptrdiff_t`，`size_t`	`保龄球` `字符`，`无符号字符`，`char2`，`uchar2`

双，double2

浮动，float2

短，无符号短，short2，ushort2

int，无符号整型，int2，uint2

长，无符号长，马龙，ulong2

很久很久，无符号长时间长，longlong2，ulonglong2

ptrdiff_t，size_t

保龄球

字符，无符号字符，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/走读生/ include / tmwtypes.h로제공됩니다。다음라인을사용하여프로그램에파일을포함합니다。

# include“tmwtypes.h”

인수제한사항

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

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

__global__ void kernel (input…)

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

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

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

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

CUDAKernel객체속성

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

k = parallel.gpu.CUDAKernel (“conv.ptx”，“conv.cu”）

k = parallel.gpu. cudakernel handlegpu属性: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파일이여러다른커널에대한여러진입점을포함할수있습니다。이러한각진입점마다고유한이름이있습니다。이러한이름은일반적으로변형(破坏)되어있습니다(C + +碾压과유사)。그러나学校网站로생성된PTX의이름에는항상铜파일의원래함수이름이포함됩니다。예를들어、铜파일이커널함수를다음과같이정의한다고가정합니다。

__global__ void simplestKernelEver(float * x, float val)

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

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

참고

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

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

template  __global__ void add4(T * v1, const T * v2) {int idx = threadadidx .x;[idx] v1 + v2 = [idx];｝

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

模板__global__ void add4(float *， const float *);模板__global__ void add4(double *， const double *);

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

_Z4add4IfEvPT_PKS0_ _Z4add4IdEvPT_PKS0_

浮动형버전에는진입점add4If를사용하고双형버전에는진입점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编程指南를참조하십시오。

CUDAKernel실행하기

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

작업공간변수사용하기

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

k = parallel.gpu.CUDAKernel (“conv.ptx”，“conv.cu”）;结果=函数宏指令(k,兰德(100 1),兰特(100 1));

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

gpuArray변수사용하기

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

k = parallel.gpu.CUDAKernel (“conv.ptx”，“conv.cu”）;i1 = gpuArray (rand(100年1“单一”））;i2 = gpuArray (rand(100年1“单一”））;result1 =函数宏指令编写此表达式(k, i1、i2);

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

result2 =函数宏指令(k, i1、i2);r1 =收集(result1)编写此表达式;r2 =收集(result2);

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

[out1, out2] = feval(内核，in1, in2, in3)을호출하면입력三机一体，in2，in3C함은铜파일내수의각입력인수에대응합니다。출력着干活과out2는C커널이실행된후C함수에대한첫번째와두번째의상수가아닌포인터입력인수값을저장합니다。

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

void reallySimple(float * pInOut, float c)

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

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

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

y =函数宏指令(k, x1, x2)

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

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

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

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

maxnumlhsararguments: 2 numrhsararguments: 3 ArgumentTypes: {'in single vector' 'inout single vector' 'inout single vector'}

이코드의커널(k)에函数宏指令을사용하려면다음구문처럼합니다。

(y1, y2) =函数宏指令(k, x1, x2, x3)

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

커널워크플로완료하기

2개의숫자추가하기

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

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

2개의벡터추가하기

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

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

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

N = 128;k.ThreadBlockSize = N;三机一体= 1 (N, 1,“gpuArray”）;in2 = 1 (N, 1“gpuArray”）;结果=函数宏指令(k,三机一体,in2);

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

CUDA를사용하는방법을보여주는예제를확인하고사용자가시도해볼수있는铜파일과PTX파일을받으려면说明GPU计算的三种方法:Mandelbrot集合항목을참조하십시오。

참고항목

mexcuda|CUDAKernel

并行计算工具箱문서

지원

仿真软件MATL金宝appAB,기타제품사용해보기

평가판신청