主要内容

跑步CUDA.或GPU上的PTX代码

概述

本主题介绍如何从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(IN2);%输入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对象

与A..cu.文件和A..ptx.文件可以创建一个Cudakernel.然后,MATLAB中的对象可以使用来评估内核:

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原型输入的用途是源代码使用无法识别的数据类型的重命名时。万博1manbetx(请参阅下面的支持万博1manbetx类型。)假设您的内核包含以下代码。

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

argtype.本身未被识别为支持的数据类型,因此在MATLAB中创建Cud万博1manbetxakernel对象时,CU文件不能直接用作输入。但是,支持的输入类型万博1manbetxAdd3.内核可以指定为Cudakernel构造函数的C原型输入。例如:

k = parallel.gpu.cudakernel('test.ptx''float *,const float *''Add3');

万博1manbetx支持的数据类型

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

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

双倍的双重

漂浮float2.

短的毫无符号短短期2.USHORT2.

unsigned int.INT2.uint2.

毫无符号朗2Ulong2.

长长毫无符号长龙龙2ulonglong2.

ptrdiff_t.size_t.

BOOL.

char无符号字符CHAR2.UCHAR2.

此外,当您包含时,支持以下整数类型万博1manbetxtmwtypes.h.程序中的标题文件。

整数类型

INT8_T.int16_t.INT32_T.INT64_T.

uint8_t.uint16_t.UINT32_T.UINT64_T.

标题文件已作为matlabroot./交换/ include/tmwtypes.h.。您将文件中的文件包含在内:

#include“tmwtypes.h”

论证限制

所有输入都可以是标量或指针,并且可以标记const

内核的C声明始终是表单:

__global__ void akernel(输入......)
  • 内核必须返回任何内容,并仅在其输入参数(标量或指针)上运行。

  • 内核无法分配任何形式的内存,因此必须在执行内核之前预先分配所有输出。因此,必须在运行内核之前已知所有输出的大小。

  • 原则上,所有指针都传递到没有的内核const可以包含输出数据,因为内核的许多线程可以修改该数据。

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

  • C的所有标量输入(双倍的漂浮等)必须是MATLAB中的标量,或标量(即单个元素)GPUARRAY变量。

  • 全部const指针输入在c(const double *等)可以是matlab中的标量或矩阵。它们被投射到正确的类型,复制到设备上,并将指向第一元素的指针传递给内核。没有关于原始大小的信息将传递给内核。就好像内核直接收到了结果mxgetdata.在AN.mxarray.

  • C中的所有非合作指针输入完全将C输入到内核中。但是,由于内核可以更改非合作指针,所以这将被视为来自内核的输出。

  • 来自MATLAB工作区标量和数组的输入将投入请求类型,然后传递给内核。但是,GPUARRAY输入不会自动演绎,因此它们的类型和复杂性必须与预期的类型完全匹配。

这些规则有一些含义。最值得注意的是,来自内核的每个输出必须必须是内核的输入,因为输入允许用户定义输出的大小(从不能在GPU上分配存储器)。

Cudakernel对象属性

在没有终止分号的情况下创建内核对象,或者在命令行键入对象变量时,MATLAB会显示内核对象属性。例如:

k = parallel.gpu.cudakernel('conv.ptx''conv.cu'
K = parallel.gpu.cudakernel句柄包在单载体'''inout单载体'}

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

有关对象属性的描述,请参阅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派生。例如,考虑以下功能模板。

模板 __global__ void add4(t * v1,const t * v2){int idx = threadidx.x;v1 [idx] + = v2 [idx];}

当模板扩展为float和double时,它会导致两个入口点,两者都包含子字符串Add4.

模板__Global__ void Add4 (float *,const float *);模板__Global__ void Add4 (双*,const double *);

PTX具有相应的条目:

_z4add4ifevpt_pks0__z4add4idevpt_pks0_

使用入口点Add4.IF.对于float版本,和add4id.对于双版本。

k = parallel.gpu.cudakernel('test.ptx''双人*,const double *''add4id');

指定线程数

通过设置其两个对象属性,您可以指定Cudakernel的计算线程数:

  • 网格化- 三个元素的向量,其产品确定块的数量。

  • threadblocksize.- 三个元素的向量,其产品确定每块的线程数。(请注意,该产品不能超过物业的价值maxthreadsperblock.。)

这两个属性的默认值是[1 1],但假设您要使用500个线程在500个元素的Vectors上运行Element-Wise操作并行。实现这一目标的简单方法是创建您的Cudakernel并相应地设置其属性:

k = parallel.gpu.cudakernel('myfun.ptx','myfun.cu');K.Threadblocksize = [500,1,1];

通常,您根据输入的大小设置网格和线程块大小。有关线程层次结构的信息,以及多维网格和块,请参阅NVIDIA CUDA C编程指南。

跑一个扶手

使用Feval.在GPU上评估Cudakernel的功能。以下示例显示如何使用MATLAB工作区变量和GPUARRAY变量执行内核。

使用工作区变量

假设您已经以母语编写了一些内核,并且希望在Matlab中使用它们来在GPU上执行。您有一个内核,这对两个向量进行了卷积;使用两个随机输入向量加载并运行它:

k = parallel.gpu.cudakernel('conv.ptx''conv.cu');结果= 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,'单身的'));结果1 = feval(k,i1,i2);

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

结果2 = Feval(k,i1,i2);R1 =聚集(结果1);r2 =聚集(结果2);

确定输入和输出对应

在呼唤时[OUT1,OUT2] = FEVAL(内核,IN1,IN2,IN3),输入In1In2., 和In3.对应于CU文件中的C函数的每个输入参数。产出OUT1.OUT2.在执行C内核之后将第一个和第二个非Const指针输入参数的值存储到C函数。

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

void soundsimple(float * pinout,float c)

相应的内核对象(K.)在Matlab中有以下属性:

maxnumlhsarguments:1 numrhsarguments:2 arearaltypes:{''inout single矢量''单scalar'}

因此,要使用此代码中的内核对象Feval.,你需要提供Feval.两个输入参数(除了内核对象),您可以使用一个输出参数:

y = feval(k,x1,x2)

输入值X1X2相当于引脚C在C函数原型中。输出参数y对应于价值引脚在C内核执行后的C函数原型中。

以下是一个稍微复杂的例子,显示Const和非Const指针的组合:

void morecomplicated(const float * pin,float * pinout1,float * pinout2)

然后,MATLAB中的相应内核对象具有属性:

maxnumlhsarguments:2 numrhsarguments:3 ArgumentTypes:{'在单载体'''Inout单vector'''Inout单vector'}

您可以使用Feval.在此代码的内核(K.)使用语法:

[Y1,Y2] = FeVal(k,x1,x2,x3)

三个输入参数X1X2, 和X3,对应于传递到C函数的三个参数。输出参数y1.Y2.,对应于值PINOUT1.PINOUT2.在C内核执行后。

完整的内核工作流程

添加两个数字

此示例在GPU中添加了两个双打。您应该安装NVIDIA CUDA Toolkit,并为您的设备提供CUDA的驱动程序。

  1. 执行此操作的CU代码如下。

    __global__ void add1(双* pi,double c){* pi + = c;}

    指令__全球的__表示这是核心的一个条目点。代码使用指针发送结果PI.,这是一个输入和输出。将此代码放在一个调用的文件中test.cu.在当前目录中。

  2. 在shell命令行上编译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代码与最后一个例子略有不同。两个输入都是指针,一个是常量,因为你没有改变它。每个线程都将简单地在其线程索引处添加元素。线程索引必须解决此线程应该添加的元素。(获取这些线程和块特定的值是CUDA编程中的一个非常常见的模式。)

    __global__ void add2(双* v1,const double * 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 Set

也可以看看

|

相关话题