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