CUDA的程序，自己写的，求高人指教

__global__ static void sumOfSquares(int *gpudata, int* result) 你把static去掉！另外我不明白你的global函数里面为什么没有定义线程的索引啊。那你那些数据运算是在哪运算的啊。我觉得应该定义int bx=blockIdx.x之类的变量啊，你这程序就是在一个线程中执行，没有并行。

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了。

---