cuda编程例子

① 如何进行CUDA C程序核函数的调试

由于CUDA调试工具的不完善、CUDA调试工具上手难度较高，并行思想本身就难调试等因素，CUDA调试一直都是一件很蛋疼的事情。写CUDA也有三四年了，前段时间在群里见别人问CUDA调试的问题，突然有想法写个CUDA调试的博客。自己经验尚浅，希望各位大大看过后能够在评论里指点一二，共同完善这篇博客。
本博客只针对逻辑bug。

1 定位bug

出现bug的第一想法自然是定位bug。cuda比较奇特的地方在于，有时报错bug在500行，但500行出的代码没有错误，而是在1000行的地方逻辑错了，十分头疼。
下面介绍三种我总结的定位bug方法：

1.1 二分法

一半一半的注释代码，定位bug。比较笨拙和麻烦，但是十分好用。

1.2 输出定位法

将整体代码分为几个模块，正常的CUDA代码大概可以分为数据返纯初始化，内存申请，内存拷贝，核函数执行，结果拷贝等模块。在每个模块结束后输出标志，示例如图1。这样在调试时就可以根据输出快速定位bug大约在什么位置。如下图：

② CUDA编程基础——Grid、Block、Thread

本消尺余文主要介绍三者之间的关系。

三者之间关系如图所示，从中可以看出，三者存在包含关系。每个grid分为多个block，每困枣个block分为多个Thread，grid和block最多可拿滚以是三维的。

③ cuda程序无法传数据到显存

方法如下：
1、CPU和GPU之间
1）CPU->GPU
从CPU向GPU传输数据，最为人熟知的就是cudaMemcpy了。
默认情况下，数据是从系统的分页内存先到锁页内存，然后再到GPU显存。因此如果显式指定使用锁页内存，是可以加快数据传输速度的。
（锁页内存，在cuda编程里使用CudaHostMalloc分配。实质上和linux的mlock系统调用一样，就是给内存页打上标记，不让操作系统将其从物理内存交换到硬盘）
至于为什么cuda要这样设计，个人理解是为了实现的方便。因为操作系统已经处理了硬盘和物理内存间的页交换等情况，显卡驱动只需要实现物理内存到GPU显存这一种数据传输即可，不需要把操作系统内存管理的事情再做一遍。
2) GPU->CPU
GPU向CPU拷贝数据时，锁页内存同样比分页内存快
值得一提的是，适当使用pinned memory显然可以加快IO速度。但是并不是越多越好，因为锁页内存是完全独占住了物理内存，操作系统无法调度，可能会影响系统整体性能。
3)同一张GPU卡内部
同一张卡内两块显存对拷，实测P40上高达~285GB/s。也比较接近于GPU卡本身的访存速度
4）数据拷贝的overhead
在上面的测试数据中，可以看到传输数据量从1M->32M增长的过程中，测得的传输带宽是有逐渐增加的。
这是因为每次调用cuda api进行数据传输都有overhead，在数据量小的时候这个overhead在数据传输时间中的占比就显得很高。这也提示我们尽量合并小数据的传输
2、同机的GPU之间
一般可以通过cudaMemcpyPeer/cudaMemcpyPeerAsync函数进行显存拷贝
1）cudaMemcpyPeer withoutP2P
/********代码示例*******/

cudaSetDevice(1);

cudaMalloc((int**)&dest, bytes);

cudaSetDevice(2);

cudaMalloc((int**)&dsrc, bytes);

cudaMemcpyPeer(dest, 1, dsrc, 2, bytes);
通过nvprof+nvpp可以看到：禁用GPU P2P时，数据是先从GPU2拷贝到系统内存（DtoH），然后再从系统内存拷贝到GPU1（HtoD）

当然，这里是在一个进程内做GPU之间的数据拷贝。如果是2个进程分别运行在GPU1和GPU2上，那在CPU上这2个进程间可以通过共享内存或者socket通信来完成数据的拷贝。

2）cudaMemcpyPeer withP2P
/********代码示例*******/

cudaSetDevice(1);

cudaMalloc((int**)&dest, bytes);

cudaSetDevice(2);

cudaMalloc((int**)&dsrc, bytes);

cudaDeviceEnablePeerAccess(1,0);

cudaDeviceEnablePeerAccess(2,0);

cudaMemcpyPeer(dest, 1, dsrc, 2, bytes);
启用GPU P2P时，数据直接从GPU2拷贝到了GPU1，不再经过系统内存。

3)通过变量赋值方式传输数据
深度学习中，卡之间传递的数据其实很多都是参数数值，因此也可以直接用一个GPU内的变量给另一个GPU上的变量赋值来进行数据传输

/********代码示例*******/

(&numBlocks, &blockSize, p2p_float);

p2p_float<<>>(

(float *)dest, (float *)src, num_elems);

__global__ void p2p_float(float *__restrict__ dest, float const *__restrict__ src,

size_t num_elems) {undefined

size_t globalId = blockIdx.x * blockDim.x + threadIdx.x;

size_t gridSize = blockDim.x * gridDim.x;

#pragma unroll(5)

for (size_t i = globalId; i < num_elems; i += gridSize) {undefined

dest[i] = src[i];

}

}

④ 【CUDA 编程】bank 与bank冲突

以下内容摘笑塌抄于Nvida 官方教程碰则圆 https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#memory-hierarchy

Shared memory has 32 banks that are organized such that successive 32-bit words map to successive banks. Each bank has a bandwidth of 32 bits per clock cycle.
共享内存由连续的32bit单元映射到32个bank， 每个时钟周期内， 每个bank都有32bit的带宽

A shared memory request for a warp does not generate a bank conflict between two threads that access any address within the same 32-bit word (even though the two addresses fall in the same bank). In that case, for read accesses, the word is broadcast to the requesting threads and for write accesses, each address is written by only one of the threads (which thread performs the write is undefined).
一个wrap内的两个线程访问同一个在共享内存中的32bit数据， 这样并不会引起bank冲突 （怎么感觉和下面的图middle的描述不一样？？）盯桥 。对于读操作， 32bit数据会被广播给请求的线程； 对于写线程， 数据仅仅会被一个线程写(这样会产生未定义的行为)。

Figure 17 shows some examples of strided access.
Figure 18 shows some examples of memory read accesses that involve the broadcast mechanism.
Figure 17. Strided Shared Memory Accesses. Examples for devices of compute capability 3.x (in 32-bit mode) or compute capability 5.x and 6.x

Left Linear addressing with a stride of one 32-bit word (no bank conflict).
Middle Linear addressing with a stride of two 32-bit words (two-way bank conflict).
Right Linear addressing with a stride of three 32-bit words (no bank conflict).
Figure 18. Irregular Shared Memory Accesses. Examples for devices of compute capability 3.x, 5.x, or 6.x.

Left Conflict-free access via random permutation.
Middle Conflict-free access since threads 3, 4, 6, 7, and 9 access the same word within bank 5.
Right Conflict-free broadcast access (threads access the same word within a bank).

⑤ c语言 gpu

U越来越强大，GPU为显示图像做了优化之外，在计算上已经超越了通用的CPU。如此强大的芯片如果只是作为显卡就太浪费了，因此NVidia推出CUDA，让显卡可以用于图像计算以外的目的，也就是超于游戏，使得GPU能够发挥其强大的运算能力。

一年前NVIDIA发布CUDA，这是一种专门针对GPU的C语言开发工具。与以往采用图形API接口指挥GPU完成各种运算处理功能不同，CUDA的出现使研究人员和工程师可以在熟悉的C语言环境下，自由地输入代码调用GPU的并行处理架构。这使得原先需要花费数天数周才能出结果的运算大大缩短到数几小时，甚至几分钟之内。

CUDA是用于GPU计算的开发环境，它是一个全新的软硬件架构，可以将GPU视为一个并行数据计算的设备，对所进行的计算进行分配和管理。在CUDA的架构中，这些计算不再像过去所谓的GPGPU架构那样必须将计算映射到图形API（OpenGL和Direct 3D）中，因此对于开发者来说，CUDA的开发门槛大大降低了。CUDA的GPU编程语言基于标准的C语言，因此任何有C语言基础的用户都很容易地开发CUDA的应用程序。

那么，如何使得CPU与GPU之间很好的进行程序之间的衔接呢？以GPGPU的概念来看，显卡仍然需要以传统的DirectX和OpenGL这样的API来实现，对于编程人员来说，这样的方法非常繁琐，而CUDA正是以GPGPU这个概念衍生而来的新的应用程序接口，不过CUDA则提供了一个更加简便的方案——C语言。我们回顾一下CUDA的发展历史。

⑥ 如何使用CUDA 显卡编程

cuda是利用gpu编程。你需要先去下一个visual
studio,然后去cuda官网下一个现在版本的cuda7.0。全部安装好后就可以编程了。cuda是并行编程，利用线程组织架构。有很多教程，你可以去看看。最好的一本书是cuda
by
example。

⑦ 如何学习cuda c

1、CUDAC编写WindowsConsoleApplication

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

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

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

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

//添加系统库

#include

#include

//添加CUDA支持

#include

__global__voidVecAdd(float*A,float*B,float*C);

__host__voidrunVecAdd(intargc,char**argv);

intmain(intargc,char**argv)

{

runVecAdd(argc,argv);

CUT_EXIT(argc,argv);

}

__host__voidrunVecAdd(intargc,char**argv)

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

constunsignedintN=8;//向量维数

constunsignedintmemSize=sizeof(float)*N;//需要空间的字节数

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

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

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

for(unsignedinti=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("Kernelexecutionfailed");

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

for(unsignedinti=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__voidVecAdd(float*A,float*B,float*C)

{

//分配sharedmemory

extern__shared__floats_A[];

extern__shared__floats_B[];

extern__shared__floats_C[];

//从globalmemory拷贝到sharedmemory

constunsignedinti=threadIdx.x;

s_A[i]=A[i];

s_B[i]=B[i];

//计算

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

//拷贝到globalmemory

C[i]=s_C[i];

}

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

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

图1Vector项目执行结果

2、CUDAC编写DLL模块

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

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

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

#ifndef_VECADD_DYNAMIC_H_

#define_VECADD_DYNAMIC_H_

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

__declspec(dllexport)voidVecAdd(float*h_A,float*h_B,float*h_C,intN);

#endif

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

#include

#include"VecAdd_dynamic.h"

#ifdef_MANAGED

#pragmamanaged(push,off)

#endif

BOOLAPIENTRYDllMain(HMODULEhMole,DWORDul_reason_for_call,LPVOIDlpReserved)

{

returnTRUE;

}

#ifdef_MANAGED

#pragmamanaged(pop)

#endif

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

EXPORTS

VecAdd

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

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

#include

#include

#include

#if__DEVICE_EMULATION__

boolInitCUDA(void)

{returntrue;}

#else

boolInitCUDA(void)

{

intcount=0;

inti=0;

cudaGetDeviceCount(&count);

if(count==0)

{

fprintf(stderr,"Thereisnodevice./n");

returnfalse;

}

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

{

cudaDevicePropprop;

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

{

if(prop.major>=1)

{break;}

}

}

if(i==count)

{

fprintf(stderr,"ThereisnodevicesupportingCUDA./n");

returnfalse;

}

cudaSetDevice(i);

printf("CUDAinitialized./n");

returntrue;

}

#endif

__global__voidD_VecAdd(float*g_A,float*g_B,float*g_C,intN)

{

unsignedinti=threadIdx.x;

if(i<N)

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

}

voidVecAdd(float*h_A,float*h_B,float*h_C,intN)

{

if(!InitCUDA())

{return;}

float*g_A,*g_B,*g_C;

unsignedintsize=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中使用CUDAC编写的DLL

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

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

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

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

#pragmaonce

namespaceWin32

{

usingnamespaceSystem::Runtime::InteropServices;

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

extern"C"voidVecAdd(float*h_A,float*h_B,float*h_C,intN);

}

在Form1.h中，#pragmaonce之后namespaceNETDemo之前添加以下代码。

#include"Win32.h"

#include

在button1_Click()中添加如下代码

intN=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(inti=0;i<N;i++)

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

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

String^reslut;

for(inti=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所示的结果

图3NETDemo运行结果

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