在GPU上运行CUDA或PTX代码- MATLAB和Simulink - Ma金宝appthWorks英国 - 金宝app,下载188bet金宝搏,金宝搏官方网站

运行CUDA或GPU上的PTX代码

概述

本主题解释如何从CU或PTX(并行线程执行)文件创建可执行内核，并在MATLAB的GPU上运行该内核^®．核在MATLAB中用a表示CUDAKernel对象，它可以对MATLAB数组或gpuArray变量进行操作。

下面的步骤描述了CUDAKernel的一般工作流程:

使用编译后的PTX代码创建一个CUDAKernel对象，其中包含GPU可执行代码。
设置CUDAKernel对象的属性来控制它在GPU上的执行。
调用函数宏指令在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(in1);%输入gpuArray。g2 = gpuArray(in2);%输入gpuArray。Result = feval(k,g1,g2);

以下部分提供了这些命令和工作流步骤的详细信息。

创建一个CUDAKernel对象

从CU文件编译一个PTX文件

如果你有一个CU文件想要在GPU上执行，你必须首先编译它来创建一个PTX文件。一种方法是用学校网站编译器在NVIDIA^®CUDA^®工具包。例如，如果您的CU文件被调用myfun.cu，你可以用shell命令创建一个编译好的PTX文件:

myfun.cu

这将生成命名为myfun.ptx．

用CU文件输入构造CUDAKernel对象

与一个.cu文件和.ptx文件可以创建CUDAKernel对象，然后你可以用它来计算内核:

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

请注意

你不能保存或负载CUDAKernel对象。

用C原型输入构造CUDAKernel对象

如果您没有与您的PTX文件相对应的CU文件，您可以为您的C内核指定C原型，而不是CU文件。例如:

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

C原型输入的另一个用途是当您的源代码使用无法识别的重命名所支持的数据类型时。金宝app(请参阅下面支持的金宝app类型。)假设您的内核包含以下代码。

ArgType;__global__ void add3(ArgType * v1, const ArgType * v2) {int idx = threadaddx .x;V1 [idx] += v2[idx];｝

ArgType它本身不能被识别为受支持的数据类型，因此在MATLAB中创建金宝appCUDAKernel对象时，包含它的CU文件不能直接用作输入。方法所支持的输入类型金宝appadd3kernel可以指定为CUDAKernel构造函数的C原型输入。例如:

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

金宝app支持的数据类型

支持的C金宝app/ c++标准数据类型如下表所示。

浮点数类型整数类型布尔和字符类型

浮点数类型	整数类型	布尔和字符类型
`双`，`double2` `浮动`，`float2`	`短`，`无符号短`，`short2`，`ushort2` `int`，`无符号整型`，`int2`，`uint2` `长`，`无符号长`，`马龙`，`ulong2` `很久很久`，`未签名的long long`，`longlong2`，`ulonglong2` `ptrdiff_t`，`size_t`	`保龄球` `字符`，`无符号字符`，`char2`，`uchar2`

双，double2

浮动，float2

短，无符号短，short2，ushort2

int，无符号整型，int2，uint2

长，无符号长，马龙，ulong2

很久很久，未签名的long long，longlong2，ulonglong2

ptrdiff_t，size_t

保龄球

字符，无符号字符，char2，uchar2

类型时，还支持以下整数类型金宝apptmwtypes.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 aKernel(输入…)

内核必须不返回任何内容，只对其输入参数(标量或指针)进行操作。
内核不能分配任何形式的内存，因此在执行内核之前必须预先分配所有输出。因此，在运行内核之前必须知道所有输出的大小。
原则上，所有传递到内核的指针都不是常量可以包含输出数据，因为内核的许多线程可以修改这些数据。

在将C语言中的核定义转换为MATLAB时:

所有标量输入在C (双，浮动，int等)必须是MATLAB中的标量，或标量(即单元素)gpuArray变量。
所有常量C语言的指针输入(Const double *等)在MATLAB中可以是标量或矩阵。它们被转换为正确的类型，复制到设备上，指向第一个元素的指针被传递给内核。没有关于原始大小的信息传递给内核。就好像内核已经直接接收到的结果mxGetData在一个mxArray．
C语言中所有的非常量指针输入都以非常量指针的形式传递给内核。然而，由于内核可以更改非常量指针，因此这将被视为来自内核的输出。
来自MATLAB工作空间标量和数组的输入被转换为所请求的类型，然后传递给内核。但是，gpuArray输入不会自动强制转换，因此它们的类型和复杂度必须与预期的完全匹配。

这些规则有一些含义。最值得注意的是，内核的每个输出都必须是内核的输入，因为输入允许用户定义输出的大小(这是由于无法在GPU上分配内存)。

对象属性

当您创建一个不带终止分号的内核对象时，或者当您在命令行中键入对象变量时，MATLAB将显示内核对象属性。例如:

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

k = parallel.gpu. cudakernel句柄包:parallel。gpu属性: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名称总是包含CU文件中的原始函数名。例如，如果CU文件将内核函数定义为

__global__ void simplestKernelEver(float * x, float val)

那么PTX代码包含一个可能被调用的条目_Z18simplestKernelEverPff．

当您有多个入口点时，请在调用时为特定内核指定入口名称CUDAKernel生成内核。

请注意

的CUDAKernel函数在PTX文件中搜索您的条目名称，并匹配任何出现的子字符串。因此，您不应该将任何条目命名为其他条目的子字符串。

您可能无法控制原始条目名称，在这种情况下，您必须知道每个条目的惟一扭曲派生。例如，考虑下面的函数模板。

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

当模板展开为float和double时，会产生两个入口点，这两个入口点都包含子字符串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的两个对象属性来指定计算线程的数量:

GridSize-三个元素的向量，它们的乘积决定了块的数量。
ThreadBlockSize-三个元素的向量，它们的乘积决定了每个块的线程数。(注意，产品不能超过属性的值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”）;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，“单一”));i2 = gpuArray(rand(100,1，“单一”));Result1 = feval(k,i1,i2);

因为输出是agpuArray，您现在可以使用此输入或输出数据执行其他操作，而无需在MATLAB工作区和GPU之间进行进一步传输。当所有GPU计算完成后，将最终结果数据收集到MATLAB工作空间中:

Result2 = feval(k,i1,i2);R1 = gather(result1);R2 = gather(result2);

确定输入和输出对应关系

当调用[out1, out2] = feval(kernel, in1, in2, in3)，输入三机一体，in2,in3对应于CU文件中C函数的每个输入参数。输出着干活而且out2在C内核执行之后，将第一个和第二个非const指针输入参数的值存储到C函数中。

例如，如果一个CU文件中的C内核具有以下签名:

void reallySimple(float * pInOut, float c)

对应的内核对象(k)在MATLAB中具有以下特性:

MaxNumLHSArguments: 1 NumRHSArguments: 2 ArgumentTypes: {'inout单个向量' '在单个标量'}

因此，使用内核对象从这段代码函数宏指令，你需要提供函数宏指令两个输入参数(除了内核对象)，你可以使用一个输出参数:

Y = feval(k,x1,x2)

输入值x1而且x2对应于引出线而且c在C函数原型中。输出参数y对应的值引出线在C内核执行后的C函数原型中。

下面是一个稍微复杂一点的例子，它显示了const指针和非const指针的组合:

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

在MATLAB中对应的内核对象具有如下属性:

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

你可以使用函数宏指令在此代码的内核上(k)，语法为:

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

三个输入参数x1，x2,x3，对应于传递给C函数的三个参数。输出参数日元而且y2，对应的值pInOut1而且pInOut2在C内核执行之后。

完整的内核工作流

两个数字相加

这个例子在GPU中将两个双精度值相加。您应该安装了NVIDIA CUDA工具包，并为您的设备安装了CUDA驱动程序。

这样做的CU代码如下。
```
__global__ void add1(double *pi, double c) {*pi += c;｝
```
该指令__global__指示这是内核的入口点。代码使用指针发送结果π，它既是输入也是输出。将这段代码放在一个名为test.cu在当前目录。
在shell命令行编译CU代码以生成一个名为test.ptx．
```
NVCC -ptx test.cu
```
用MATLAB创建内核。目前这个PTX文件只有一个条目，所以您不需要指定它。如果你要放入更多的核，你会指定add1作为入口。
```
k = parallel.gpu.CUDAKernel(“test.ptx”，“test.cu”）;
```
使用两个数字输入运行内核。默认情况下，内核运行在一个线程上。
```
结果= feval(k,2,3)
```
```
结果= 5
```

两个向量相加

这个例子扩展了前面的例子，将两个向量相加。为简单起见，假设线程数与向量中的元素数完全相同，并且只有一个线程块。

CU代码与上一个示例略有不同。两个输入都是指针，其中一个是常量，因为你没有改变它。每个线程将简单地在其线程索引处添加元素。线程索引必须计算出该线程应该添加哪个元素(在CUDA编程中，获取这些线程和块特定的值是一种非常常见的模式)。
```
__global__ void add2(double * v1, const double * v2) {int idx = threadaddx .x;V1 [idx] += v2[idx];｝
```
将此代码保存在文件中test.cu．
像以前一样编译学校网站．
```
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文件供您试验的示例，请参见阐述GPU计算的三种方法:Mandelbrot集．

另请参阅

CUDAKernel|mexcuda

并行计算工具箱文档

金宝app

试试MATLAB、Sim金宝appulink和其他产品下载188bet金宝搏

现在就去审判吧