【IT168 技术】本文将集中讨论如何使用CUDA代码创建一个非托管DLL，并在C#程序中使用它，列举的例子将展示在数组上做计算的for()循环的托管、非托管和新的.NET 4并行版本之间的一些差异。
我将简要地介绍如何配置CUDA环境和运行示例程序，CUDA本身已经超出了本文的范围，我只会谈及CUDA内核执行速度和内存管理。
硬件开启CUDA和.NET 4(Visual Studio 2010 IDE或C# Express 2010)是成功运行示例代码的前提条件，Visual C++ Express 2008已经成为标准的CUDA C编辑器(2010版本修改了自定义生成规则功能，不支持CUDA SDK提供的生成规则)。
第1部分：为CUDA配置环境和工具
CUDA是由NVIDIA引入的一个通用目的并行计算架构，CUDA程序(内核)运行在GPU而不是CPU上，性能更好(数百个核心可以同时运行数千个计算线程)，它带有一个软件环境，允许开发人员使用C作为高级编程语言，这个计算技术可用于数学、科学、金融、建模、图像处理等领域。
针对开发的CUDA基础配置
下载并安装适配你操作系统的CUDA工具包，建议使用最新的版本进行设备模拟(如果你没有支持CUDA的设备，用模拟器是个不错的主意，虽然它可能有一些限制);
下载并安装相同操作系统，相同工具包版本的SDK;
在安装SDK时如果遇到问题，请更新显示驱动。
Visual C++ Express 2008(或Visual Studio 2008)配置
1、语法着色
从主窗口中打开“工具”*“选项”，选中“文本编辑器”*“文件扩展名”，添加“.cu”和“.cuh”扩展名;
将usertype.dat文件从“[sdk dir]\C\doc\syntax_highlighting\visual_studio_8\”拷贝到“Program Files\Microsoft Visual Studio 9.0\Common7\IDE\”文件夹;
重启Visual Studio。
2、新建项目，32位Windows XP版本
默认SDK位于“c:\Documents and Settings\All Users\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK\”;
默认工具包位于“C:\CUDA\”;
创建空白Win32控制台应用程序，添加.cu扩展名的源文件;
在“解决方案资源管理器”窗口选择新创建的项目，然后点击右键，选择“自定义生成规则”，使用“查找已有”按钮定位到“[sdk dir] \C\common\”文件夹中的Cuda.rules文件，添加它，并在可用规则文件列表上做上标记；
再次选中“项目”*“属性”，选择“发布配置”，然后在树形视图中依次选择“配置属性”*“链接程序”*“常规”*“附加库目录”，添加“C:\CUDA\lib;C:\Documents and Settings\All Users\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\common\lib”；
选择“配置属性”*“链接程序”*“输入”*“附加依赖”，输入cudart.lib；
为“调试配置”使用相同的设置；对于模拟器配置，使用“配置管理器”基于发布和调试添加新的配置，然后将它们分别命名为“EMU-Release”和“EMU-Debug”(或你自己想一个名字)；从“配置”组合框选择“EMU-Release”，将“cudart.lib”修改为“cudartemu.lib”，从树形视图中选择“CUDA生成规则(版本号)”，点击“常规”，将“仿真模式”设置为“是”；为“EMU-Debug”做相同的配置。
3、新建项目，64位Windows 7版本
Default SDK location is c:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\
默认SDK位置是“c:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\”；
Default toolkit location is C:\CUDA\
默认工具包位置是“C:\CUDA\”；
剩下的操作和32位Windows XP版本类似，但也有以下一些不同的地方：
1)使用配置管理器创建新解决方案配置AMD64_Release，AMD64_Debug和模拟器版本(如果需要的话，直接从Release和Debug配置拷贝设置)；
2)使用配置管理器添加新创建的x64平台解决方案，从32位版本拷贝设置;
3)在项目上下文窗口为所有AMD64配置选择x64平台；
4)依次点击“链接程序”*“常规”*“附加库目录”，增加“c:\CUDA\lib64; c:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\common\lib”；
5)再点击“链接程序”*“输入”*“附加依赖”，输入cudart.lib。
关于示例文件的重要注意事项
因为项目解决方案文件中的自定义生成规则条目，你必须做基本的配置才能打开连接的示例，如果该规则因路径错误不可访问，VC ++ 2008会返回错误，你也可以编辑cudalib.vcproj文件修复cuda.rules文件的路径。
如果这样做没有帮助，那么直接创建一个新项目，像前面配置小节描述的那样手动添加自定义生成规则，然后将示例中的.cu文件拷贝到新项目。
你也需要手工将cutil.h文件从“\NVIDIA GPU Computing SDK\C\common\inc\”目录拷贝到“C:\CUDA\include\”(这是修复导入路径最简单的方法)，并将cutil32.lib(或用于64位的cutil64.lib)添加到链接程序输入附加依赖，结果看起来是“cudart.lib cutil32.lib”(在示例项目中已经做好)。
CUDA工具包和SDK已经用于生成dll。
即使没有这一步，主要项目部分(用C#编写)也可以工作，因为我已经将生成dll库添加到“\bin\Debug”和“\bin\Release”文件夹了(如果你在dll部分做了任何修改，你都必须替换它)。
C#项目部分你需要VS 2010版本，因为它用到了.NET 4的功能。
第2部分：CUDA DLL
DLL部分代码已经用VC++ 2008 IDE写好，假设IDE配置正确，并用它创建了一个新的Win32控制台应用程序，将应用程序类型转换成“DLL”，并标记为“空白项目”(不需要预便于头和/或dllmain()函数)，添加新的资源文件并保存为扩展名为.cu的文件，语法着色应该会工作，记住添加CUDA自定义生成规则和链接程序依赖。
如果想运行CUDA内核(在GPU设备上执行的函数)，我们需要一些包装函数暴露给外部dll，内核安装和调用都将在这个函数内，非常棒的一个功能是，我们可以传递内核执行配置参数给它(Grid大小，块大小和贡献内存大小)，而不用给它们设置常量，这样我们就可以在目标机器上运行基准测试确定最合适的值。
extern "C" int __declspec(dllexport) __stdcall SomeCalculationsCU
(
float *a_h, // pointer to input array
const unsigned int N, // input array size
const unsigned int M, // kernel M parameter
const int cuBlockSize = 512, // kernel block size (max 512)
const int showErrors = 1 // show CUDA errors in console window
)
{
int tmp = PRINT_ERRORS;
PRINT_ERRORS = showErrors;
float *a_d; // pointer to device array
size_t size = N * sizeof(float);
int cuerr = 0; // no errors
unsigned int timer = 0;
cudaMalloc((void**)&a_d, size); // allocate array on device
cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
int n_blocks = N / cuBlockSize + (N % cuBlockSize == 0 ? 0 : 1);
cutCreateTimer(&timer); // from cutil.h
cutStartTimer(timer);
some_calculations <<<n_blocks, cuBlockSize>>> (a_d, N, M); // kernel invocation
cudaThreadSynchronize(); // by default kernel runs in parallel with CPU
code
cutStopTimer(timer);
cuerr = checkCUDAError("cuda kernel");
cudaMemcpy(a_h, a_d, size, cudaMemcpyDeviceToHost);
if(!cuerr) cuerr = checkCUDAError("cuda memcpy");
sExecutionTime = cutGetTimerValue(timer);
cudaFree(a_d);
if(!cuerr) cuerr = checkCUDAError("cuda free");
PRINT_ERRORS = tmp;
return cuerr;
}
该函数中最重要的部分是“extern "C" int _declspec(dllexport) _stdcall”部分，使得外部dll可见，必须指定调用约定(这里是_stdcall)，因为默认情况下，C函数使用“_cdecl”，.NET平台使用CallingConvention.Winapi调用(即_stdcall)。另一个重要事项是，为调用程序和被调用函数使用相同的约定。
这个函数也可以是控制台应用程序中的main()函数，因此，你可以修改它，并增加一些打印结果，例如：
{
float *a_h;
const unsigned int N = 2000;
const unsigned int M = 10;
const int cublocks = 256;
size_t size = N * sizeof(float);
a_h = (float*)malloc(size);
for(unsigned int i = 0; i < N; i++) a_h[i] = (float)i;
SomeCalculationsCU(a_h, N, M, cublocks, 1);
printf("exec time = %f ms\n", sExecutionTime);
}
如果你想以一个程序(不生成dll)运行它，在“项目属性”*“配置属性”*“常规”部分，将“配置类型修”改为“应用程序(.exe)”，这样就允许你使用CUDA Profiler(位于“C:\CUDA\cudaprof\bin\”)。
文件dllmain.cu也有变量可从外部dll访问：
extern "C"
{
float __declspec(dllexport) sExecutionTime = -1;
}
// variable wrapper function
extern "C" float __declspec(dllexport) __stdcall GetExecutionTime()
{
return sExecutionTime;
}
虽然可以直接访问变量，但使用包装函数方法更简单(就像get/set存取器)。
最终的内核函数示例：
__global__ void some_calculations(float *a, unsigned int N, unsigned int M)
{
unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N)
{
// note1: no need for shared memory here
// note2: global memory access is coalesced
// (no structs, float only used)
// do computations M times on each thread
// to extend processor's time
for(unsigned int i = 0; i < M; i++)
{
// some easy arithmetics
a[idx] = a[idx] * a[idx] * 0.1 - a[idx] - 10;
}
}
}
不同CPU版本执行时间对比：
(
float *a_h,
const unsigned int N,
const unsigned int M
)
unsigned int timer = 0;
cutCreateTimer(&timer);
cutStartTimer(timer);
for(unsigned int i = 0; i < N; i++)
for(unsigned int j = 0; j < M; j++)
*(a_h + i) = *(a_h + i) * *(a_h + i) * 0.1 - *(a_h + i) - 10;
cutStopTimer(timer);
sExecutionTime = cutGetTimerValue(timer);
}
