CUDA文件无需c-mex使用，GPU与MATLAB混合编程

当计算机中没有安装并行计算工具箱时，可以采用c-mex文件在GPU上运行CUDA文件（.cu文件），这部分内容在第2章中已经介绍过。当有并行计算工具箱时，可以直接在GPU上运行CUDA文件，而不需要c-mex文件。在这种情况下，我们无需编译c-mex文件，只需对CUDA文件进行nvcc编译即可。

我们打算对AddVectors.cu的例子进行微调，该例在第2章中用于c-mex。这里是AddVectors_noCmex.cu文件。首先，AddVectors_noCmex.cu文件只有CUDA内核部分（__global__ void AddVectors_noCmex()），与第2章中AddVectors.cu不同，该文件具有CUDA存储器配置和与c-mex接口相关的存储复制。

在MATLAB中直接使用.cu文件的基本流程如下所示：

在shell（DOS或Linux shell）内，应该在NVIDIA CUDA工具箱内加 -ptx选项，使用nvcc编译cu文件生成ptx文件。也可以在MATLAB指令窗口内使用如下system指令运行这一shell指令：

虽然我们不编译c-mex文件，但是仍然需要C/C++编译器中的nvcc编译cu文件，这是因为cu文件具有C/C++格式，并且nvcc转发全部非CUDA编译步骤到C/C++编译器。如果你的系统没有指向C/C++编译器（例如cl.exe）的路径，接下来可以用 -ccbin选项确定路径。关于从NVIDIA网站安装nvcc可参考第2章。通常，成功安装NVIDIA CUDA工具箱会自动创建路径指向nvcc.exe位置（C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\bin）。不过，如果由于某些原因没有创建指向该位置的路径，那么需要在运行前添加addpath('C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\bin')指令：

使用相同文件名的ptx与cu文件，由parallel.gpu.CUDAKernel()函数生成CUDA核函数。当cu文件由多个函数组成时，可以按照下面的方法指定入口点名称：

通过CUDA对象AddVectors_noCmex的函数句柄myCu，有两个输入实参矩阵A和B）的feval函数运行。下面的nvcc_noCmex.m文件展示了编译cu文件、生成CUDA对象和使用测试输入运行的全部代码：

当运行这一代码时，有如下结果：(www.xing528.com)

通过MATLAB指令myCu.ThreadBlockSize = N（这里N=10为向量A的长度），ThreadBlockSize设为[10 1 1]。ThreadBlockSize控制AddVectors_noCmex.cu里的 i = threadIdx.x，在GPU中实现对10个元素进行并行加法，结果为10个数据元素均为11。如果把myCu.ThreadBlockSize的值都改为8，如下所示，最后两个元素没有正确求和：

除了使用c-mex接口，AddVectors.cu与AddVectors_noCmex.cu的主要区别是CUDA核函数实参的数目：

在第2章中的AddVectors.cu例子中，c-mex的CUDA核函数（addVectorsKernel()）在c-mex运行后被其他函数调用，因此用户可以自由地规定实参的数量和顺序。不过，MATLAB函数feval()通过parallel.gpu.CUDAKernel直接使用的CUDA核函数addVectorsKernel()），则应该遵循严格的规则。当调用[output1, output2] = feval(CUDA_Kernel, input1, input2, input3)时，input1、input2和 input3应该与CUDA_Kernel函数的输入实参一致。如果有两个以上输入实参，那么前两个输入实参用作输出实参；如果有两个或者更少的输入实参，那么仅第一个输入实参用作输出实参。可以通过CUDA核函数句柄来核实CUDA核函数的属性：

这里，可以看到ArgumentTypes里有两个inout的描述。inout表明了这两个实参既可以用于输入，也可以用于输出。下面看一个很简单的例子：

在这个简单的CUDA核函数（Simple_noCmex.cu）中，核心句柄（handleCu）性质显示一个inout实参和一个in实参。最后一行为核函数的返回类型，始终应该为void类型，而输出值应该存回第一个或第二个输入实参中。

[1]由于MATLAB分布式计算服务器产品可以让使用者在远程计算机集群上运行额外的worker，因此称单机多核worker为本地worker。简单起见，本书中只采用单机进行并行处理，所以这里的worker即为“本地worker”。