摘要：本文論述了使用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編程探索