CUDA C编程探索

摘要：本文论述了使用CUDA C编写Windows Console Application、动态链接库(DLL)、在 .NET 中使用CUDA C编写的DLL的基本方法。

1、 CUDA C编写Windows Console Application

下面我们从一个简单的例子开始学习CUDA C。

打开VS，新建一个CUDAWinApp项目，项目名称为Vector，解决方案名称为CUDADemo。依次点击“确定”，“下一步”，选择Empty project。点击“Finished”。这样一个CUDA的项目就建成了。

右键点击Vector项目，依次选择“添加”、“新建项”、“代码”、“CUDA”。在名称中输入要添加的文件名。如Vector.cu。然后点击添加。

下面在Vector.cu文件里实现两个向量相加的程序。

    //添加系统库
#include <stdio.h>
#include <stdlib.h>
//添加CUDA支持
#include <cuda.h>

__global__ void VecAdd(float *A, float *B, float *C);

__host__ void runVecAdd(int argc, char **argv);

int main(int argc, char **argv)
{
 runVecAdd(argc,argv);

 CUT_EXIT(argc,argv);

}

__host__ void runVecAdd(int argc,char **argv)

{//初始化host端内存数据

const unsigned int N = 8;//向量维数

const unsigned int memSize = sizeof(float)*N;//需要空间的字节数

float *h_A = (float*)malloc(memSize);

float *h_B = (float*)malloc(memSize);

float *h_C = (float*)malloc(memSize);

for (unsigned int i = 0; i < N; i++)

{h_A[i] = i;h_B[i] = i;}

//设备端显存空间

float *d_A, *d_B, *d_C;

//初始化Device

CUT_DEVICE_INIT(argc,argv);

CUDA_SAFE_CALL(cudaMalloc((void**)&d_A, memSize));

CUDA_SAFE_CALL(cudaMalloc((void**)&d_B, memSize));

CUDA_SAFE_CALL(cudaMalloc((void**)&d_C, memSize));

CUDA_SAFE_CALL(cudaMemcpy(d_A, h_A, memSize, cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaMemcpy(d_B, h_B, memSize, cudaMemcpyHostToDevice));

VecAdd<<<1,N,memSize>>>(d_A, d_B, d_C);

CUT_CHECK_ERROR("Kernel execution failed");

CUDA_SAFE_CALL(cudaMemcpy(h_C, d_C, memSize, cudaMemcpyDeviceToHost));

for (unsigned int i = 0; i < N; i++)

{ printf("%.0f ",h_C[i]); }

free(h_A);free(h_B);free(h_C);

CUDA_SAFE_CALL(cudaFree(d_A));

CUDA_SAFE_CALL(cudaFree(d_B));

CUDA_SAFE_CALL(cudaFree(d_C));

}

__global__ void VecAdd(float *A, float *B, float *C)

{

//分配shared memory

extern __shared__ float s_A[];

extern __shared__ float s_B[];

extern __shared__ float s_C[];

//从global memory拷贝到shared memory

const unsigned int i = threadIdx.x;

s_A[i] = A[i];

s_B[i] = B[i];

//计算

s_C[i] = s_A[i] + s_B[i];

//拷贝到global memory

C[i] = s_C[i];

}
  

由于这里不是讲CUDA编程的，关于它的编程模型已经超出了我要介绍的范围，您可以阅读《GPU高性能运算之CUDA》来获得CUDA编程模型的知识。

编译Vector项目，执行此项目后会得到图1如下输出：

图1 Vector项目执行结果

2、CUDA C编写DLL模块

更多情况下的您的软件可能只是使用CUDA来实现一段程序的加速，这种情况下我们可以使用CUDA C 编写DLL来提供接口。下面我们就将例1编译成DLL。

在刚才的CUDADemo解决方案目录下添加一个新的CUDA项目（当然您也可以重新建立一个解决方案）。项目名为VecAdd_dynamic。Application Type选为DLL，Additional Options选择Empty Project。

第一步，添加头文件，文件名最好与工程名同名，这样便于您的维护工作。这里我向项目中添加了VecAdd_dynamic.h，在此头文件中添加如下代码

    #ifndef _VECADD_DYNAMIC_H_

#define _VECADD_DYNAMIC_H_

//并行计算N维向量的加法

__declspec(dllexport) void VecAdd(float* h_A, float* h_B, float* h_C, int N);

#endif

第二步，添加cpp文件，文件名为VecAdd_dynamic.cpp，在此文件中添加如下代码

#include

#include "VecAdd_dynamic.h"

#ifdef _MANAGED

#pragma managed(push, off)

#endif

BOOL APIENTRY DllMain(HMODULE hModule,DWORD ul_reason_for_call,LPVOID lpReserved)

{

return TRUE;

}

#ifdef _MANAGED

#pragma managed(pop)

#endif
  

第三步，添加def文件，此文件的功能就是确保其它厂商的编译器能够调用此DLL里的函数。这一点非常关键，因为您的程序可能用到多个厂家的编译器。文件名为VecAdd_dynamic.def。向该文件中添加：

    EXPORTS

VecAdd
  

第四步，添加cu文件，文件名为VecAdd_dynamic.cu。注意此文件最好直接添加到项目目录下，不要添加到源文件选项卡或其它已有的选项卡下。如图2所示

图2 VecAdd_dynamic项目文件组织

在cu文件里添加如下代码，实现要导出的函数。

    #include

#include

#include

#if __DEVICE_EMULATION__

bool InitCUDA(void)

{ return true;}

#else

bool InitCUDA(void)

{

int count = 0;

int i = 0;

cudaGetDeviceCount(&count);

if(count == 0)

{

fprintf(stderr, "There is no device./n");

return false;

}

for(i = 0; i < count; i++)

{

cudaDeviceProp prop;

if(cudaGetDeviceProperties(&prop, i) == cudaSuccess)

{

if(prop.major >= 1)

{ break; }

}

}

if(i == count)

{

fprintf(stderr, "There is no device supporting CUDA./n");

return false;

}

cudaSetDevice(i);

printf("CUDA initialized./n");

return true;

}

#endif

__global__ void D_VecAdd(float *g_A, float *g_B, float *g_C, int N)

{

unsigned int i = threadIdx.x;

if (i < N)

{ g_C[i] = g_A[i] + g_B[i]; }

}

void VecAdd(float* h_A, float* h_B, float* h_C, int N)

{

if(!InitCUDA())

{ return; }

float *g_A, *g_B, *g_C;

unsigned int size = N * sizeof(float);

CUDA_SAFE_CALL(cudaMalloc((void**)&g_A, size));

CUDA_SAFE_CALL(cudaMalloc((void**)&g_B, size));

CUDA_SAFE_CALL(cudaMalloc((void**)&g_C, size));

CUDA_SAFE_CALL(cudaMemcpy(g_A, h_A, size, cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaMemcpy(g_B, h_B, size, cudaMemcpyHostToDevice));

D_VecAdd<<<1,N>>>(g_A, g_B, g_C, N);

CUDA_SAFE_CALL(cudaMemcpy(h_C, g_C, size, cudaMemcpyDeviceToHost));

cudaFree(g_A);cudaFree(g_B);cudaFree(g_C);

}
  

第五步，如果您已经正确完成了以上四步，那么剩下的就只有编译，只要您用过VS，这一步就不需要我介绍了吧。成功之后，在您的解决方案文件目录下的Debug文件夹下会有一个VecAdd_dynamic.dll文件。

3、 在 .NET 中使用CUDA C编写的DLL

下面介绍在托管程序中如何使用VecAdd_dynamic.dll。

第一步，在上面的解决方案CUDADemo下添加一个C++/CLR的Windows窗体应用程序，工程名为NETDemo（当然您也可以重新建一个解决方案，工程名也是随意的）。

第二步，在窗体上添加一个按钮，名字随意，我将它的现实文本改为“调用CUDA_DLL”，给这个按钮添加click事件。我们的代码将在这个事件里添加调用VecAdd()的程序。在窗体上添加一个文本框用来显示调用输出的结果。

第三步，代码实现。为工程NETDemo添加一个头文件，我将它命名为Win32.h，这个文件中主要是实现VecAdd()函数的导入。在此文件中添加如下代码

    #pragma once

namespace Win32

{

using namespace System::Runtime::InteropServices;

[DllImport("VecAdd_dynamic.dll",EntryPoint="VecAdd",CharSet=CharSet::Auto)]

extern "C" void VecAdd(float* h_A, float* h_B, float* h_C, int N);

}
  

在Form1.h中，#pragma once 之后 namespace NETDemo 之前添加以下代码。

    #include "Win32.h"

#include 
  

在button1_Click()中添加如下代码

    int N = 8;

float* h_A = (float*)malloc(N*sizeof(float));

float* h_B = (float*)malloc(N*sizeof(float));

float* h_C = (float*)malloc(N*sizeof(float));

for (int i = 0; i < N; i++)

{h_A[i] = i;h_B[i] = i;}

Win32::VecAdd(h_A, h_B, h_C,N);

String ^reslut;

for (int i = 0; i < N; i++)

{reslut += Convert::ToString(h_C[i]) + ", ";}

this->textBox1->Text = Convert::ToString(reslut);

free(h_A);free(h_B);free(h_C);
  

第四步、执行NETDemo项目。点击“调用CUDA_DLL”，您会看到图3所示的结果

图3 NETDemo运行结果

到现在为止您已经完全可以正确使用CUDA了。

参考

[1]Jeffrey Richter ， Christophe Nasarre . Windows 核心编程（第五版） [M]. 北京：清华大学出版社， 2008.

[2] 张舒，褶艳利 . GPU 高性能运算之 CUDA [M]. 北京：中国水利水电出版社， 2009.

CUDA C编程探索