本主题介绍如何从CU或PTX(并行线程执行)文件中创建可执行内核,并在MATLAB运行该GPU上的内核®。内核在matlab中表示Cudakernel.
对象,可以在MATLAB阵列或GPUARRAY变量上运行。
以下步骤描述了Cudakernel常规工作流程:
使用编译的PTX代码来创建一个Cudakernel对象,其中包含GPU可执行代码。
在Cudakernel对象上设置属性以控制其在GPU上的执行。
称呼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);
以下部分提供了这些命令和工作流程的详细信息。
如果您要在GPU上执行CU文件,则必须首先编译它以创建PTX文件。这样做的一种方法是NVCC.
编译器在nvidia®CUDA.®工具包。例如,如果调用CU文件myfun.cu.
,您可以使用shell命令创建编译的ptx文件:
nvcc -ptx myfun.cu.
这会生成命名的文件myfun.ptx.
。
与A..cu.
文件和A..ptx.
文件可以创建一个Cudakernel.
然后,MATLAB中的对象可以使用来评估内核:
k = parallel.gpu.cudakernel('myfun.ptx'那'myfun.cu');
笔记
你不能节省
要么加载
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');
支持的C万博1manbetx / C ++标准数据类型列于下表中。
浮动类型 | 整数类型 | 布尔和字符类型 |
---|---|---|
|
|
|
此外,当您包含时,支持以下整数类型万博1manbetxtmwtypes.h.
程序中的标题文件。
整数类型 |
---|
|
标题文件已作为
。您将文件中的文件包含在内: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上分配存储器)。
在没有终止分号的情况下创建内核对象,或者在命令行键入对象变量时,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.
运行内核时的对象是输入:
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)
,输入In1
那In2.
, 和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)
输入值X1
和X2
相当于引脚
和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)
三个输入参数X1
那X2
, 和X3
,对应于传递到C函数的三个参数。输出参数y1.
和Y2.
,对应于值PINOUT1.
和PINOUT2.
在C内核执行后。
此示例在GPU中添加了两个双打。您应该安装NVIDIA CUDA Toolkit,并为您的设备提供CUDA的驱动程序。
执行此操作的CU代码如下。
__global__ void add1(双* pi,double c){* pi + = c;}
指令__全球的__
表示这是核心的一个条目点。代码使用指针发送结果PI.
,这是一个输入和输出。将此代码放在一个调用的文件中test.cu.
在当前目录中。
在shell命令行上编译cu代码以生成调用ptx文件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(双* v1,const double * v2){int idx = threadidx.x;v1 [idx] + = v2 [idx];}
在文件中保存此代码test.cu.
。
使用之前编译NVCC.
。
nvcc -ptx test.cu.
如果此代码与第一个示例的代码一起放入同一CU文件中,则需要在此时指定输入点名称以区分其。
k = parallel.gpu.cudakernel('test.ptx'那'test.cu'那'Add2');
在运行内核之前,请正确设置要添加的向量的线程数。
n = 128;k.threadblocksize = n;In1 =那些(n,1,'gpuarray');In2 =那些(n,1,'gpuarray');结果= Feval(k,In1,In2);
有关展示如何使用CUDA的示例,并为您提供CU和PTX文件进行实验,请参阅说明GPU计算的三种方法:Mandelbrot Set。