如何学习cuda c?

1、 CUDA C编写Windows Console Application

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

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

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

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

//添加系统库

#include

#include

//添加CUDA支持

#include

__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] = ih_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。注意此文件最好直接添加到项目目录下，不要添加到源文件选项卡或其它已有的选项卡下。

在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] = ih_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. CUDA out of memory

跑cuda 程序遇到下面错误：

RuntimeError: CUDA out of memory. Tried to allocate 588.00 MiB (GPU 011.00 GiB total capacity8.97 GiB already allocated190.44 MiB free9.00 GiB reserved in total by PyTorch)

运行程序之前，使用nvidia-smi 查看显存有没有被占用，如果有被占用5M以上，可能是显存没有被释放。通过如下命令来释放显存。

linux上使用命令行，云主机可以重启电脑

fuser -k /dev/nvidia* 或者 kill $(lsof -t /dev/nvidia*)

2. 减小batch size。

3. 更换更大显存的云主机。

有时候，我们需要在比较大的项目中调用CUDA，这就涉及到MFC+CUDA的环境配置问题，以矩阵相乘为例，在MFC中调用CUDA程序。我们参考罗振东iylzd@163.com（国防科学技术大学计算机学院）的方法。

环境： Windows 7 SP1

Microsoft Visual Studio 2010

CUDA 5.0

步骤：

1.首先建立一个空的名叫Matrix Multiplication_KahanMFC的“FCM应用程序”项目：

点击“确定”，这时d出如下窗口

我们需要对默认项目进行一些修改，点击“下一步”，我们设置一个空的MFC项目，选择“单个文档”和“MFC标准”：

点击“完成”。

2.创建CUDA的调用接口函数及其头文件

（1）头文件

“添加”-->“新建项”-->“Visual C++”-->“头文件（.h）”-->“名称”-->“CUDA_Transfer.h” -->“添加”，如下图：

在CUDA_Transfer.h中添加如下代码：

//CUDA_Transfer.h

#include

#include "math.h"

using namespace std

int run_cuda(float* GPU, float* CPU)

如下图所示：

（2）函数

按照和增加头文件相似的方法，添加函数。“添加”-->“新建项”-->“Visual C++”-->“C++文件（.cpp）” -->“名称”-->“CUDA_Transfer.cpp” -->“添加”，如下图：

在CUDA_Transfer.cpp中添加如下代码：

//CUDA_Transfer.cpp

#include "CUDA_Transfer.h"

#include "stdafx.h"

extern "C" int runtest(float* GPU, float* CPU)

int run_cuda(float* GPU, float* CPU)

{

runtest(GPU,CPU)

return 0

}

如下图所示：

需要注意的是在MFC的文件中是不能包含（include）.cu文件的，会报错，所以我们使用extern "C"的方式来实现函数的调用。

3. 创建存放cuda 代码的筛选器，名为CUDA

“添加”-->“新建筛选器”，重命名为CUDA

4. 在筛选器CUDA中创建一个CUDA源代码文件，kernel.cu。

我们直接把已经写好的矩阵相乘的程序kernel.cu复制到项目目录下，添加到CUDA筛选器中去。

添加”-->“现有项”-->“kernel.cu”-->“添加”：

把kernel.cu的int main()函数改为extern "C" int runtest(float* GPU, float* CPU)，两个参数用来获得GPU和CPU计算所使用的时间，单位为毫秒。

5. 右击项目-->“生成自定义”：

在d出的窗口中勾选CUDA 5.0(.target,.props)。如果使用其他版本的CUDA，就勾选对应的版本：

点击“确定”。

6. 修改 kernel.cu的编译链接设置

在解决方案资源管理器中右击kernel.cu文件-->“属性”，在d出窗口中-->“常规”-->“项类型”的下拉列表中选择

点击“应用”后，“常规”下方会出现一个“CUDA C/C++”的设置，没有特殊需求，不需要修改，点击“确定”。

7.修改工程设置。

工程设置需要修改“链接器”-->“输入”-->“附加依赖项”和“生成事件”-->“预先生成事件”-->“命令行”。需要设置的参数比较多，我们采用比较简单的方法。

我们新建一个空的CUDA项目，在这个空CUDA项目的项目属性中找到“链接器”-->“输入”-->“附加依赖项”，把“附加依赖项”中所包含的项复制到我们的MFC项目中：

按照同样的方法，设置“生成事件”-->“预先生成事件”-->“命令行”：

设置完成后，点击“确定”。

8.修改MFC文件，完成调用。

我们需要在MFC中调用CUDA程序，显示出GPU和CPU计算两个1024*1024矩阵相乘所消耗的时间。

在Matrix Multiplication_KahanMFCView.cpp中包含（include）"CUDA_Transfer.h"

文件；在CMatrixMultiplication_KahanMFCView::OnDraw(CDC* pDC)中添加如下代码：

float GPU

float CPU

run_cuda(&GPU, &CPU)

CString strGPU,strCPU

strGPU.Format(_T("GPU:%f \n"),GPU)

strCPU.Format(_T("CPU:%f \n"),CPU)

pDC->TextOut(0,0,strGPU)

pDC->TextOut(0,30,strCPU)

如图所示：

然后重新生成解决方案，运行。

计算要花费一些时间，需要等待，测试的时候可以把矩阵大小改小一些。因为把程序加到了OnDraw中，所以每当刷新窗口时候（例如调整窗口大小时），都会调用。由于计算耗时比较长，窗口看起来会像无响应一样，等计算完成就好了。

运行的结果如下：

在矩阵比较大的情况下，GPU的加速效果明显，GPU耗时只需要620ms，而CPU需要23438ms，要花费将近40倍的时间。

---