从这部分开始 结合虫子的demo程序给大家分析下cuda的性能与可行性。
一。先概述下实现流程。
CUDA在执行的时候是让host里面的一个一个的kernel按照线程网格(Grid)的概念在显卡硬件(GPU)上执行。每一个线程网格又可以包含多个线程块(block),每一个线程块中又可以包含多个线程(thread)。
每一个kernel交给每一个Grid来完成。当要执行这些任务的时候,每一个Grid又把任务分成一部分一部分的block,block再分线程来完成。每个Grid中的任务是一定的。二维线程块的索引关系为如下:
unsigned int xIndex = blockDim.x * blockIdx.x + threadIdx.x;
unsigned int yIndex = blockDim.y * blockIdx.y + threadIdx.y;
block中的每个线程都有自己的寄存器和local memory,block中的所有线程共享一个shared memory,一个grid共享一个global memory。
每一个时钟周期内,warp(一个block里面一起运行的thread)包含的thread数量是有限的,现在的规定是32个。一个block中含有16个warp。所以一个block中最多含有512个线程
每次Device(就是显卡)只处理一个grid。
下面说明一下硬件的执行模型。
假如出于某种原因,公司的办公室被征用了搞活动。只留下一个小房间来给开发团队。每一个时钟周期内按照wrap(就理解为运行的时候,一个block里面一起运行的thread,例如block里面有512个thread,但是每次只有32个thread在运行,那么这32个thread就是一个运行的warp组 线程束)。每一个warp里面包含的thread数量是有限的,现在的规定是32个。将来不知道会不会有变化,这个只有CUDA开发人员知道了。每次Device(就是显卡)只处理一个grid(在未来支持directX11 的硬件中这一限制可能被解除)。假如我们一个部门有x个人,办公室内有N个桌子,每张桌子可以坐32个人。然后轮流来开发….。这里的桌子可以理解成multiprocessor(多处理器)。每个sm中包含8个标量流处理器(sp)。GPU所谓的多核中核的概念就是sp的数量。Cuda中的kernel函数实质上是以block为单位执行的。同一block中需要共享数据,因此他们必须在同一个SM中发射,而block中的每一个线程则被发射到sp上去执行。 疑点:既然有这样的线程簇限制、为何还要设置高于warp线程数的线程。
二。demo
安装部署方面driver、toolkit、sdk顺序安装好。Cuda的项目支持4种调试方式release、debug、emurelease、emudubug。前2个是需要gpu真正的支持cuda后者是cpu模拟gpu。至于你的电脑能否支持cuda 可以运行下deviceQuery.exe程序
计算框架(二) 实例相关-cuda并行计算举例">
图中我们关注一下几点就可以了,首先 有一个支持cuda的设备。计算能力1,局存储器的大小,核的数量,多处理器的数量,常量存储器的大小、每个block的共享存储器的大小、wrap的线程数等等。
想看cuda在图形领域的应用可以运行这个smokeParticles.exe程序哦。
在我的demo中,cpp文件主要是处理一些cpu端的处理、cu文件通常是与gpu核函数和cuda api的一些内容。其中My_kernel封装了具体的核函数实现方法。Cudatool项目就是cuda的应用程序,CudaProviders是我连接C#与cuda之间的驱动、CudaWeb就是我们平常的web项目。CUDAWinApp这个就是一些小的功能演示。
下面介绍下cuda的函数类型限定符。
__device__ 在设备上执行、只能在设备上调用。
__global__ 用于声明内核函数、在设备上执行只能从主机端调用。
__host__ 在主机端执行,只能从主机端调用,默认。
__device__与__global__不支持递归,函数体内不能声明静态变量、参数数目不可变化,不能对device取指针。__global__与__host__不能连用。__global__只能返回空,调用__global__函数必须声明其执行配置、__global__函数的调用是异步的、__global__参数的值目前是通过共享存储器传递,总的大小不能超过256byte。
变量类型限定符分为__device__(变量存在设备端上)、__constant__(存在常数存储器空间)、__share__(block的共享存储器)、volatile关键字 当线程间数据可能互相影响变换时使用。