编译CUDA动态链接库及使用

除了直接在项目中使用cu或cuh来编写CUDA代码之外，还可以将CUDA相关操作代码放在一个DLL项目中，将项目编译成动态链接库DLL，然后在需要使用的项目中引用这个DLL并调用其内部函数即可。
现在新建一个DLL项目，项目名称为Test00302，如下图所示：



现在在项目中新建一个名为Test.cu的文件，如下图所示：



然后设置项目的生成自定义方式为，如下图所示：





然后设置Test.cu的属性中“项类型”为“CUDA C/C++”，如下图所示：



最后，在项目属性中链接器的附加依赖项里添加cudart.lib，如下图所示：



在Test.cu中添加如下代码：

#include <stdio.h>
#include <iostream>
#include <cuda_runtime.h> //CUDA运行时库头文件
using namespace std;

extern "C"
{
#ifdef TEST00302_EXPORTS
#define TEST00302_API __declspec(dllexport)
#else
#define TEST00302_API __declspec(dllimport)
#endif

TEST00302_API void ShowDeviceProp(void); //要导出的函数
}

//显示设备信息
void ShowDeviceProp(void)
{
int i,count;
cudaDeviceProp prop;
cudaError_t cudaStatus=cudaGetDeviceCount(&count);
if(cudaStatus == cudaSuccess) {
cout<<"共有设备数目："<<count<<"\n";
if(count>0)
{
for(i=0;i<count;i++)
{
cudaGetDeviceProperties(&prop,i);//获取设备的属性信息
cout<<"\n第"<<i+1<<"个设备信息：\n";
cout<<"设备名称："<<prop.name<<"\n";
cout<<"总内存："<<prop.totalGlobalMem/1048576<<"M\n";
cout<<"常量内存："<<prop.totalConstMem<<"字节\n";
cout<<"设备中处理器数目："<<prop.multiProcessorCount<<"个\n";
cout<<"每个线程块最多包含线程数目："<<prop.maxThreadsPerBlock<<"个\n";
cout<<"一个线程格中可包含的线程块数目：I="<< prop.maxGridSize[0]
<<" J="<<prop.maxGridSize[1]<<" K="<<prop.maxGridSize[2]<<"\n";
cout<<"多维线程块中可以包含的最大线程数目：I="<< prop.maxThreadsDim[0]
<<" J="<<prop.maxThreadsDim[1]<<" K="<<prop.maxThreadsDim[2]<<"\n";
}
}
}
else
{
cout<<"没有获取到设备信息！请检查计算机是否具有支持CUDA的显卡设备以及CUDA驱动程序版本是否需要更新！\n";
}
}

值得注意的是前面的extern “C”这段代码，它用于导出cu文件中定义的CUDA函数。编译项目，在debug文件中将会生成Test00302.dll文件，如下图所示：



现在新建一个VS项目，项目名称为Test00303，为控制台应用程序，在Test00303.cpp文件中添加动态引用dll的代码并执行dll中的ShowDeviceProp ()函数，如下：

// Test00303.cpp : 定义控制台应用程序的入口点。
//

#include "stdafx.h"
#include <stdio.h>
#include <stdlib.h>
#include <Windows.h>

//动态加载链接测试
typedef void(*DLLFUNC)(void);//声明需要从dll中调用的函数原型的函数指针

int _tmain(int argc, _TCHAR* argv[])
{
HINSTANCE myCudaDll=LoadLibrary(__T("Test00302.dll"));//动态地加载dll文件
if (myCudaDll)
{
DLLFUNC dllFun=(DLLFUNC)GetProcAddress(myCudaDll,"ShowDeviceProp");//获得函数指针
if (dllFun)
{
dllFun();//执行函数
}
else
{
printf("在DLL中不存在该函数");//可能由于函数名错误
}
FreeLibrary(myCudaDll);//动态地卸载Dll文件
}
else
{
printf("加载DLL失败!");
}
system("pause");
return 0;
}

然后将项目Test00302中生成的动态链接库文件Test00302.dll复制到项目Test00303中debug文件夹下，如下图所示：



运行项目Test00303，其结果如下图所示：



项目Test00303为普通的控制台应用程序，Test00302.dll为编译的包含CUDA函数的动态链接库，通过动态引用，即可在普通的应用程序里面加载CUDA程序。
下面再用调用CUDA内核函数的进行测试，先在项目Test00302中新建一个名为Test2.cu的文件，在Test2.cu中添加GPU中执行的核函数addKernel()，然后添加用于向量相加的函数vectorAdd()，在函数vectorAdd()中选择用于执行的GPU设备、在设备上分配内存、复制主机内存数据到设备内存、启动核函数、调用cudaDeviceSynchronize()监听核函数执行、复制设备内存数据到主机内存、重置CUDA设备、释放设备内存。Test2.cu中的代码如下：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

extern "C"
{
#ifdef TEST00302_EXPORTS
#define TEST00302_API __declspec(dllexport)
#else
#define TEST00302_API __declspec(dllimport)
#endif

TEST00302_API int vectorAdd(int c[], int a[], int b[],int size);
}

//CUDA核函数
__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
c[i] = a[i] + b[i];
}
//向量相加
int vectorAdd(int c[], int a[], int b[],int size)
{
int result=-1;
int *dev_a = 0;
int *dev_b = 0;
int *dev_c = 0;
cudaError_t cudaStatus;

// 选择用于运行的GPU
cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess) {
result=1;
goto Error;
}

// 在GPU中为变量dev_a、dev_b、dev_c分配内存空间.
cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
result=2;
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
result=3;
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
result=4;
goto Error;
}

// 从主机内存复制数据到GPU内存中.
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
result=5;
goto Error;
}
cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
result=6;
goto Error;
}

// 启动GPU内核函数
addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

// 采用cudaDeviceSynchronize等待GPU内核函数执行完成并且返回遇到的任何错误信息
cudaStatus = cudaDeviceSynchronize();
if (cudaStatus != cudaSuccess) {
result=7;
goto Error;
}

// 从GPU内存中复制数据到主机内存中
cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess) {
result=8;
goto Error;
}

result=0;

// 重置CUDA设备，在退出之前必须调用cudaDeviceReset
cudaStatus = cudaDeviceReset();
if (cudaStatus != cudaSuccess) {
return 9;
}
Error:
//释放设备中变量所占内存
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);

return result;
}

编译项目Test00302，在生成的动态链接库文件Test00302.dll复制到项目Test00303的debug文件夹下，然后修改项目Test00303中的文件Test00303.cpp内容如下：

// Test00303.cpp : 定义控制台应用程序的入口点。
#include "stdafx.h"
#include <stdio.h>
#include <stdlib.h>
#include <Windows.h>

typedef int(*DLLFUNC)(void);//声明需要从dll中调用的函数原型的函数指针
typedef int(*DLLFUNC2)(int *c, int *a, int *b,int size);//声明需要从dll中调用的函数原型的函数指针

int _tmain(int argc, _TCHAR* argv[])
{
HINSTANCE myCudaDll=LoadLibrary(__T("Test00302.dll"));//动态地加载dll文件
if (myCudaDll)
{
//调用DLL中的ShowDeviceProp函数显示设备属性信息
//DLLFUNC dllFun=(DLLFUNC)GetProcAddress(myCudaDll,"ShowDeviceProp");//获得函数指针
//if (dllFun)
//{
//	dllFun();//执行函数
//}
//else
//{
//	printf("在DLL中不存在该函数");//可能由于函数名错误
//}

//调用DLL中的vectorAdd函数
DLLFUNC2 dllFun2=(DLLFUNC2)GetProcAddress(myCudaDll,"vectorAdd");//获得函数指针
if (dllFun2)
{
const int arraySize = 5;
int a[arraySize] = { 1, 2, 3, 4, 5 };
int b[arraySize] = { 10, 20, 30, 40, 50 };
int c[arraySize] = { 0 };

int r=dllFun2(c,a,b,arraySize);//执行函数
if(r==0)
{
printf("计算成功\n计算结果：\n");
printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
c[0], c[1], c[2], c[3], c[4]);
}
else
printf("计算失败\n");
}
else
{
printf("在DLL中不存在该函数");//可能由于函数名错误
}
FreeLibrary(myCudaDll);//动态地卸载Dll文件
}
else
{
printf("加载DLL失败!");
}
system("pause");
return 0;
}

由于dll中的函数vectorAdd 定义如下：int vectorAdd(int c[], int a[], int b[],int size);所以在声明函数的原型指针时也需要定义为：
typedef int(*DLLFUNC2)(int *c, int *a, int *b,int size);
或：
typedef int(*DLLFUNC2)(int c[], int a[], int b[],int size);运行项目Test00303，其结果如下图所示：



虽然这个例子比较简单，但是向量相加的计算过程是由GPU完成的。