CUDAKernel 对象属性

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

k = parallel.gpu.CUDAKernel("conv.ptx","conv.cu")

k = 
  parallel.gpu.CUDAKernel handle
  Package: parallel.gpu

  Properties:
     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 对象的属性控制它的一些执行行为。使用点符号来改变可以更改的属性。有关对象属性的描述，请参阅 CUDAKernel。修改可设置属性的典型原因是指定线程数，如下所述。

指定入口点

单个 PTX 文件可以包含不同内核的多个入口点。每个入口点都有一个唯一的名称。每个入口点的名称都被重组（如同 C++ 重组），但始终包含来自 CU 文件的原始函数名称。例如，如果 CU 文件将核函数定义为

__global__ void simplestKernelEver( float * x, float val )

那么 PTX 代码包含一个可能被称为 _Z18simplestKernelEverPff 的条目。

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

k = parallel.gpu.CUDAKernel("myfun.ptx","myfun.cu","myKernel1");

您可能无法控制原始条目名称，在这种情况下，您必须了解每个条目派生的唯一混乱。例如，考虑以下函数模板。

template <typename T>
__global__ void add4( T * v1, const T * v2 )
{
    int idx = threadIdx.x;
    v1[idx] += v2[idx];
}

当模板针对 float 和 double 进行扩展时，会产生两个入口点，都包含子字符串 add4。

template __global__ void add4<float>(float *, const float *);
template __global__ void add4<double>(double *, const double *);

PTX 有相应的条目：

_Z4add4IfEvPT_PKS0_
_Z4add4IdEvPT_PKS0_

对于浮点版本使用入口点 add4If，对于双精度版本使用入口点 add4Id。

k = parallel.gpu.CUDAKernel("test.ptx","double *, const double *","add4Id");

指定线程数

您可以通过设置 CUDAKernel 的两个对象属性来指定其计算线程的数量：

这两个属性的默认值都是 [1 1 1]，但假设您想使用 500 个线程对 500 个元素的向量并行运行逐元素操作。设置 ThreadBlockSize 使得其元素的乘积为 500。

k = parallel.gpu.CUDAKernel("myfun.ptx","myfun.cu");
k.ThreadBlockSize = [500,1,1];

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

使用 C 原型输入构造 CUDAKernel 对象.  如果您没有与 PTX 文件对应的 CU 文件，您可以为您的 C 内核指定 C 原型，而不是 CU 文件。例如：

k = parallel.gpu.CUDAKernel("myfun.ptx","float *, const float *, float");

C 原型输入的另一个用途是当您的源代码使用无法识别的受支持数据类型的重命名时。假设您的内核包含以下代码。

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

ArgType 本身不被识别为受支持的数据类型，因此在 MATLAB 中创建 CUDAKernel 对象时，包含它的 CU 文件不能直接用作输入。但是，add3 内核支持的输入类型可以指定为 CUDAKernel 构造函数的 C 原型输入。例如：

k = parallel.gpu.CUDAKernel("test.ptx","float *, const float *","add3");

支持的数据类型.  支持的 C/C++标准数据类型如下表所示。

此外，当您在程序中包含 tmwtypes.h 头文件时，以下整数类型也受支持。

头文件以 matlabroot/extern/include/tmwtypes.h 的形式提供。您可以使用以下行将文件包含到程序中：

#include "tmwtypes.h"

参量限制.  所有输入都可以是标量或指针，并且可以使用 const 标记为常量值。

内核的 C 声明始终采用以下形式：

__global__ void aKernel(inputs ...)

将 C 语言中的内核定义翻译成 MATLAB 时：

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

确定输入和输出对应关系

调用 [out1, out2] = feval(kernel,in1,in2,in3) 时，输入 in1、in2 和 in3 对应于 CU 文件内函数的每个输入参量。输出 out1 和 out2 将内核执行后的第一个和第二个非常量指针输入参量的值存储到函数中。

例如，如果 CU 文件中的内核具有以下签名：

void reallySimple( float * pInOut, float c )

MATLAB 中对应的内核对象（k）具有以下属性：

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

因此，要将此代码中的内核对象与 feval 一起使用，您需要为 feval 提供两个输入参量（除了内核对象），并且可以使用一个输出参量。

y = feval(k,x1,x2)

输入值 x1 和 x2 对应于函数原型中的 pInOut 和 c。输出参量 y 对应于内核执行后函数原型中 pInOut 的值。

下面是一个稍微复杂一点的例子，展示了常量指针和非常量指针的组合：

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

MATLAB 中对应的内核对象具有以下属性：

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

您可以使用三个输入参量和两个输出参量在此代码的内核（k）上使用 feval。

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

三个输入参量 x1、x2、x3，对应传递给函数的三个参量。输出参量 y1 和 y2 对应内核执行后的 pInOut1 和 pInOut2 的值。

两个数字相加

此示例在 GPU 中将两个双精度数相加。

两个向量相加

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

CU 和 PTX 文件示例

有关如何使用 CUDA 的示例，以及提供 CU 和 PTX 文件供您试验，请参阅 说明 GPU 计算的三种方法：曼德布洛特集合。