跑CUDA或PTX代码在GPU

概观

本主题说明如何创建一个从CU或PTX（并行线程执行）文件的可执行的内核，以及从MATLAB GPU上运行的内核函数®。内核在MATLAB由代表CUDAKernel对象，可以MATLAB阵列或gpuArray变量上操作。

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

  1. 使用编译PTX代码创建一个CUDAKernel对象，其中包含了GPU可执行代码。

  2. 在CUDAKernel对象上设置属性来控制其在GPU上执行。

  3. 呼叫feval与所需投入的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（平方英寸）;％输入gpuArray。导致= feval（K，G1，G2）;

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

创建CUDAKernel对象

从CU文件编译一个PTX文件

如果你要在GPU上执行CU文件，必须先编译它来创建一个PTX文件。这样做的一个方法是使用NVCC编译器在NVIDIA®CUDA®工具包。例如，如果你的CU文件名为myfun.cu，你可以创建一个shell命令编译PTX文件：

NVCC -ptx myfun.cu

这将生成文件名为myfun.ptx。

构建与CU文件输入CUDAKernel对象

随着.CU文件和.ptx文件可以创建一个CUDAKernel对象MATLAB，然后可以用它来评估的内核：

K = parallel.gpu.CUDAKernel（'myfun.ptx'，'myfun.cu'）;

注意

你不能保存要么加载CUDAKernel对象。

构建CUDAKernel对象与C语言输入

如果没有相应的你的PTX文件CU文件，你可以指定C原型你的C内核来代替CU文件。例如：

K = parallel.gpu.CUDAKernel（'myfun.ptx'，“浮动*，常量浮动*浮动”）;

对于C原型输入的另一个用途是当源代码使用一个支持的数据类型的未识别的重命名。万博1manbetx（请参阅下面的支持万博1manbetx的类型）。假设你的内核包括下面的代码。

的typedef浮动ArgType;__global__空隙ADD3（ArgType * V1，常量ArgType * V2）{INT IDX = threadIdx.x;V1 [IDX] + = V2 [IDX];}

ArgType本身不被识别为支持的数据类型，因此，其包括它不能创建在MAT万博1manbetxLAB中CUDAKernel对象时被直接用作输入的CU文件。但是，支持的输入类型的万博1manbetxADD3内核可被指定为C输入到CUDAKernel构造原型。例如：

K = parallel.gpu.CUDAKernel（'test.ptx'，“浮动*，常量浮动*”，'ADD3'）;

万博1manbetx支持的数据类型

所支持的万博1manbetxC / C ++标准数据类型列于下表。

float类型 整型 布尔和字符类型

双，double2

浮动，FLOAT2

短，无符号短，short2，ushort2

INT，无符号整型，INT2，UINT2

长，无符号长，long2，ulong2

很长很长，无符号很长很长，longlong2，ulonglong2

ptrdiff_t的，为size_t

布尔

烧焦，无符号的字符，CHAR2，uchar2

此外，当您包括以下整型都支持万博1manbetxtmwtypes.h在你的程序的头文件。

整型

中int8_t，int16_T，int32_T，的int64_t

uint8_T，uint16_T，uint32_t的，uint64_t中

头文件在装运时matlabroot/extern/include/tmwtypes.h。您在与行程序文件：

＃包括“tmwtypes.h”

参数限制

所有的输入可以是标量或指针，并且可被标记常量。

内核的C声明的形式总是：

__global__空隙aKernel（输入...）

  • 内核不能返回任何结果，只有在它的输入参数（标量或指针）操作。

  • 内核无法分配任何形式的存储器，因此，所有输出必须在执行前内核预先分配。因此，在运行内核前所有输出的尺寸必须是已知的。

  • 原则上，所有指针传递到了不是内核常量可以包含输出数据，因为内核的多个线程可能会修改该数据。

当翻译在C内核到MATLAB的定义：

  • 在C中的所有的标量输入（双，浮动，INT等）必须在MATLAB标量，或者标量（即，单元素）gpuArray变量。

  • 所有常量用C指针输入（常量双*等）可以是标量MATLAB或矩阵。它们被转换为正确的类型，复制到设备，以及一个指针的第一个元素被传递到内核。有关原始尺寸没有信息被传递给内核。这是因为虽然内核直接收到的结果mxGetData上mxArray。

  • 在C中的所有非恒定指针输入被精确地转移到内核非恒定指针。然而，由于非恒定的指针可以由内核来改变，这将被视为从内核的输出。

  • 从工作空间MATLAB标量和数组的输入被铸造成所要求的类型，然后传递到内核。然而，gpuArray输入没有自动转换，所以它们的类型和复杂性必须完全匹配那些预期。

这些规则有一定的影响。最值得注意的是，从一个内核每个输出必然也有一个输入到内核，由于输入允许用户定义输出（从暂时无法分配在GPU存储器下面）的大小。

CUDAKernel对象属性

当创建内核对象而无需终接分号，或当你在命令行键入对象变量，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：{”在单个载体”‘INOUT单个矢量’}

内核对象的属性控制它的一些执行行为的。使用点表示法来改变可以改变这些属性。

对于对象属性的描述，请参见CUDAKernel对象引用页面。用于修改可设置的属性的一个典型原因是指定的线程的数目，如下文所述。

指定入口点

如果您的PTX文件包含多个入口点，您可以识别特定内核myfun.ptx你想要的内核对象ķ引用：

K = parallel.gpu.CUDAKernel（'myfun.ptx'，'myfun.cu'，'myKernel1'）;

一个单一的PTX文件可以包含多个入口点，以不同的内核。各条目点都有一个唯一的名称。这些名称通常错位（如在C ++的mangling）。然而，当所产生的NVCC在PTX名称总是包含在CU文件中的原始函数名。例如，如果CU文件定义了核函数作为

__global__空隙simplestKernelEver（浮子* X，浮子val）的

那么PTX代码包含的信息可以被称为入口_Z18simplestKernelEverPff。

当你有多个入口点，打电话时指定特定内核条目名称CUDAKernel生成内核。

注意

该CUDAKernel在PTX文件的条目名称搜索功能，以及任何子串出现的比赛。所以，你不应该命名您的任何条目，任何其他的子串的。

您可能没有比原来的条目名称，在这种情况下，你必须意识到独特的错位得出每个控制。例如，考虑下面的函数模板。

模板 __global__空隙ADD4（T * V1，常量T * V2）{INT IDX = threadIdx.x;V1 [IDX] + = V2 [IDX];}

当模板被扩展出来float和double，它会导致两个入口点，它们都含有子ADD4。

模板__global__空隙ADD4 <浮子>（浮子*，常量浮子*）;模板__global__空隙ADD4 <双>（双*，常量双*）;

该PTX有相应的条目：

_Z4add4IfEvPT_PKS0_ _Z4add4IdEvPT_PKS0_

使用切入点add4If浮子版本，add4Id为双版本。

K = parallel.gpu.CUDAKernel（'test.ptx'，“双*，常量双*”，'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

使用feval功能评估在GPU上一个CUDAKernel。下面的例子说明了如何使用MATLAB的工作空间变量和gpuArray变量来执行内核。

使用工作区变量

假设你已经写了一些内核以本地语言和想使用它们在MATLAB中的GPU上执行。你有一个内核确实在两个向量的卷积;加载和具有两个随机输入向量运行：

K = parallel.gpu.CUDAKernel（'conv.ptx'，'conv.cu'）;导致= feval（K，兰特（100,1），兰特（100,1））;

即使输入用于MATLAB工作区数据常数或变量时，输出为gpuArray。

使用gpuArray变量

这可能是更有效地使用gpuArray运行内核对象时，输入：

K = parallel.gpu.CUDAKernel（'conv.ptx'，'conv.cu'）;I1 = gpuArray（兰特（100,1，'单'））;I2 = gpuArray（兰特（100,1，'单'））;RESULT1 = feval（K，I1，I2）;

因为输出是一个gpuArray，现在就可以执行使用此输入或输出数据，而不MATLAB工作空间和GPU之间进一步传送其他操作。当你所有的GPU计算是完整的，收集您的最终结果的数据到MATLAB工作区：

RESULT2 = feval（K，I1，I2）;R1 =聚集（RESULT1）;R 2 =聚（RESULT2）;

确定输入和输出对应

当调用[OUT1，OUT2] = feval（内核，IN1，IN2，IN3）中，输入端IN1，IN2和IN3对应于你的每一个CU文件中的输入参数的C函数。输出OUT1和OUT2的C内核已被执行后的第一和第二非const指针输入参数的值存储到C函数。

例如，如果一个CU文件中的C内核具有以下特征：

空隙reallySimple（浮子*的引脚配置，浮子c）中

相应的内核对象（ķ）在MATLAB具有以下性质：

MaxNumLHSArguments：1个NumRHSArguments：2个ArgumentTypes：{ 'INOUT单一载体' 在单个标 ''}

因此，使用内核对象，从这个代码feval，您需要提供feval两个输入参数（除了内核对象），并且可以使用一个输出参数：

Y = feval（K，X1，X2）

输入值X1和X2相当于引脚排列和C在C函数原型。输出参数ÿ对应于该值引脚排列在C函数原型后C-内核执行。

下面是一个稍微复杂的例子，示出了常数和非const指针的组合：

空隙moreComplicated（常量浮子*销，浮子* pInOut1，浮子* pInOut2）

在MATLAB相应的内核对象于是具有的属性：

MaxNumLHSArguments：2个NumRHSArguments：3个ArgumentTypes：{ '在单个载体' 'INOUT 'INOUT单个矢量' 单矢量'}

您可以使用feval这个代码的内核（ķ）的语法：

[Y1，Y2] = feval（K，X1，X2，X3）

这三个输入参数X1，X2和X3，对应于被传递到C函数的三个参数。输出参数Y1和Y2对应于的值pInOut1和pInOut2后C-内核执行。

整个内核工作流程

两个数字相加

这个例子在GPU将两个双打在一起。你应该有安装了NVIDIA CUDA工具包，并为您的设备支持CUDA的驱动程序。

  1. 该CU代码来做到这一点如下。

    __global__空隙ADD1（双* PI，双c）{* PI + = C;}

    该指令__全球__表明这是一个入口点的内核。代码使用指针发出的结果PI，这既是一个输入和一个输出。将这个代码在一个名为test.cu在当前目录。

  2. 在外壳命令行编译CU代码来生成称为PTX文件test.ptx。

    NVCC -ptx test.cu

  3. 创建MATLAB的内核。目前，该PTX文件只有一个入口，所以你并不需要指定它。如果你把更多的内核中，你会指定ADD1作为入门。

    K = parallel.gpu.CUDAKernel（'test.ptx'，'test.cu'）;

  4. 运行带有两个数字输入的内核。默认情况下，内核在一个线程中运行。

    导致= feval（K，2,3）
    导致= 5

两个向量

此示例扩展前一个添加两个矢量在一起。为简单起见，假定有相同数量的线程作为向量元素和只有一个线程块。

  1. 该CU代码是从最后一个例子稍有不同。两个输入都为指针，因为你不改变它一个是恒定的。每个线程将简单的螺纹指数在添加元素。螺纹指数必须制定出这个线程应该添加哪些元素。（获取这些thread-和块特定值是在CUDA编程一个非常普遍的图案）。

    __global__空隙ADD2（双* V1，常量双* V2）{INT IDX = threadIdx.x;V1 [IDX] + = V2 [IDX];}

    在文件中保存此代码test.cu。

  2. 作为编译使用前NVCC。

    NVCC -ptx test.cu

  3. 如果此代码被放在同一CU文件，第一个例子的代码一起，你需要指定入口点名称这段时间来区分它。

    K = parallel.gpu.CUDAKernel（'test.ptx'，'test.cu'，'ADD2'）;

  4. 之前运行的内核，你要添加的载体正确设置线程数。

    N = 128;k.ThreadBlockSize = N;IN1 =酮（N，1，'gpuArray'）;IN2 =酮（N，1，'gpuArray'）;导致= feval（K，IN1，IN2）;

例如与CU和PTX文件

举一个例子，说明如何使用CUDA的工作，并提供CU和PTX文件，为您进行实验，看龙虎斗三种方法来GPU计算：Mandelbrot集。

并行计算工具箱文档
万博1manbetx

尝试MATLAB，Sim万博1manbetxulink和其他产品s manbetx 845

现在就试用