建议先看看前言中关于存储器的介绍：点击打开链接
线程
首先介绍进程，进程是程序的一次执行，线程是进程内的一个相对独立的可执行的单元。若把进程称为任务的话，那么线程则是应用中的一个子任务的执行。举个简单的例子：一个人要做饭，食谱就是程序代码，做的过程就是执行程序，做好的饭就是程序运行的结果，而在这期间，需要炒菜，放盐，放油等等就是线程。
线程同步
调用__syncthreads 创建一个 barrier 栅栏
每个线程在调用点等待块内所有线程执行到这个地方，然后所有线程继续执行后续命令
Mds[i] = Md[j]; __syncthreads(); func(Mds[i],Mds[i+1]);
要求线程的执行时间尽量接近，因为执行时间差别过大，等待时间过长，效率低下。
只在块内同步，全局同步开销大
线程同步会导致死锁，例如：
if(someFunc()) { __syncthreads(); } else { __syncthreads(); }
线程调度
Warp——块内的一组线程，一般是32个线程为一组
线程调度的基本单位
运行于同一个 SM(流多处理器)
内部threadIdx值连续
warp内部线程不能沿不同分支执行
线程层次
Grid——一维或多维线程块(block)
一维或二维
Block——一组线程
一维，二维或三维
一个 Grid 里的每个 Block 的线程数是一样的
block内部的内个线程可以同步(synchronize)和访问共享存储器(shared memory)
线程组织模型
block下由线程束(warp) 组成，block 内部的线程共享“shared memory”
可以同步 “__syncthreads()”
一个Kernel具有大量线程，Kernel启动一个 “grid”，包含若干线程块，用户设定
下图是一个 grid 模型：
从图中可以看出这个 grid 中用户指定了 2 个线程块，每个线程块中有 2 个线程
可以看到CUDA内存传输是这样：
Device端
针对每个线程，可以读写 local memory 和 register
针对每个线程块，可以读写 shared memory
针对每个 grid，可以读写 global memory，只读 constant memory 和 texture memory
Host端
可以读写 global memory，constant memory 和 texture memory
回顾线程层次
首先看一个计算 N × N 矩阵的例子：
假设N = 32
先指定 Grid 里面有 2 × 2 个线程块 Block，表示为
blockIdx ([0,1],[0,1])
每个块有 16 × 16 个线程 (跟 N 无关)，表示为
threadIdx ([0,15],[0,15])
块大小这样表示blockDim = 16
则某个线程的真实索引为
i = [0,1] * 16 + [0,15]
线程索引： threadIdx
一维Block Thread ID == Thread Index
二维 Block (Dx,Dy)
Thread ID of index(x,y) == x + yDy
三维 Block (Dx,Dy,Dz)
Thread ID of index(x,y,z) == x + yDx + zDxDy
计算 N × N 矩阵，指定线程块为 1 块，大小为矩阵大小
__global__ void MatAdd(float A[N][N],float B[N][N],float C[N][N]) { int i = threadIdx.x; int j = threadIdx.y; C[i][j] = A[i][j] + B[i][j]; } int main() { int numBlocks = 1; //指定线程块为 1 dim3 threadsPerBlock(N,N); //指定大小为 N × N MatAdd<<<numBlocks,threadsPerBlock>>>(A,B,C); }
Thread Block 线程块
线程的集合
位于相同的处理器核(SM)
共享所在核的存储器
块索引：blockIdx
维度：blockDim
一维，二维，三维
计算 N × N 矩阵，指定线程块大小为 16 × 16，若干个线程块
__global__ void MatAdd(float A[N][N],float B[N][N],float C[N][N]) { int i = blockIdx.x * blockDim.x + threadIdx.x; int j = blockIdx.y * blockDim.y + threadIdx.y; if(i < N && j < N) C[i][j] = A[i][j] + B[i][j]; } int main() { dim3 threadsPerBlock(16,16); dim3 numBlocks(N / threadsPerBlock.x , N / threadsPerBlock.y); MatAdd<<<numBlocks,threadsPerBlock>>>(A,B,C); }
线程块之间彼此独立执行
任意顺序：并行或串行
被任意数量的处理器以任意顺序调度：处理器的数量具有可扩展性
一个块内部的线程
共享容量有限的低延迟存储器
同步执行：合并访存，__syncThreads()：Barrier ——块内线程一起等待所有线程都执行某处语句，轻量级
实例演示
先介绍几个函数：
cudaMalloc()
在设备端分配 global memory
cudaFree()
释放存储空间
cudaMemcpy()
内存传输
Host to host
Host to device
Device to host
Device to device
cudaMemcpy(Md,M,size,cudaMemcpyHostToDevice);
Md 目的地址
M 原地址
size 大小
CPU端代码：
void MatrixMulOnHost(float** M,float** N,float** P,int width) { for(int i = 0 ; i < width ; ++i) for(int j = 0 ; j < width ; ++j) { float sum = 0; for(int k = 0 ; k < width ; ++k) { float a = M[i][k]; float b = N[k][j]; sum += a * b; } P[i][j] = sum; } }
CUDA 算法框架
int main()
{
1.管理内存，为输入的原始数据，以及输出结果在GPU上分配内存
2.并行处理
3.结果拷贝回 CPU，释放内存
return 0;
}
1.
int size = Width * Width * sizeof(float); cudaMalloc(Md,size); cudaMemcpy(Md,M,size,cudaMemcpyHostToDevice); cudaMalloc(Nd,N,size,cudaMemcpyHostToDevice); cudaMalloc(Pd,size);
2.
__global__ void MatrixMulKernel(float* Md,float* Nd,float* Pd,int width) { int tx = threadIdx.x; int ty = threadIdx.y; float Pvalue = 0; for(int k = 0 ; k < width ; ++k) { float Mdelement = Md[ty * Md.width + k]; float Ndelement = Nd[k * Nd.width + tx]; Pvalue += Mdelement * Ndelement; } Pd[ty * width + tx] = Pvalue; }
3.
cudaMemcpy(P,Pd,size,cudaMemcpyDeviceToHost);
cudaFree(Md);
cudaFree(Nd);
cudaFree(Pd);在main()中调用 Kernel
dim3 dimBlock(WIDTH,WIDTH);
dim3 dimGrid(1,1);
MatrixMulKernel <<<dimGrid,dimBlock>>>(Md,Nd,Pd);完整GPU程序如下：
/* ============================================================================ Name : Martix.cu Author : wzn Version : Copyright : Your copyright notice Description : CUDA compute reciprocals ============================================================================ */ #include <iostream> #include <numeric> #include <stdlib.h> __global__ void MatrixMulKernel(float* Md,float* Nd,float* Pd,int Width) { int tx = threadIdx.x; int ty = threadIdx.y; float Pvalue = 0; for(int k = 0 ; k < Width ; ++k) { float Mdelement = Md[ty * Width + k]; float Ndelement = Nd[k * Width + tx]; Pvalue += Mdelement * Ndelement; } Pd[ty * Width + tx] = Pvalue; } int main(void) { int Width = 30; int size = Width * Width * sizeof(float); float M[Width][Width],N[Width][Width],P[Width][Width]; for(int i = 0;i<30;i++){ for(int j = 0;j<30;j++){ M[i][j] = j; N[i][j] = j; } } float *Md,*Nd,*Pd; cudaMalloc((void **)&Md,size); cudaMalloc((void **)&Nd,size); cudaMemcpy(Md,M,size,cudaMemcpyHostToDevice); cudaMemcpy(Nd,N,size,cudaMemcpyHostToDevice); cudaMalloc((void **)&Pd,size); dim3 dimBlock(Width,Width); dim3 dimGrid(1,1); MatrixMulKernel <<<dimGrid,dimBlock>>>(Md,Nd,Pd,Width); cudaMemcpy(P,Pd,size,cudaMemcpyDeviceToHost); cudaFree(Md); cudaFree(Nd); cudaFree(Pd); for(int i = 0;i<30;i++){ for(int j = 0;j<30;j++){ std::cout<<P[i][j]<<" "; } std::cout<<" "; } return 0; } CPU端程序: #include <iostream> using namespace std; void MatrixMulOnHost(float** M,float** N,float** P,int width) { for(int i = 0 ; i < width ; ++i) for(int j = 0 ; j < width ; ++j) { float sum = 0; for(int k = 0 ; k < width ; ++k) { float a = M[i][k]; float b = N[k][j]; sum += a * b; } P[i][j] = sum; } } int main() { int Width = 30; float **M,**N,**P; M = new float* [Width]; for(int i = 0 ;i<Width;i++) M[i] = new float [Width]; N = new float* [Width]; for(int i = 0 ;i<Width;i++) N[i] = new float [Width]; P = new float* [Width]; for(int i = 0 ;i<Width;i++) P[i] = new float [Width]; for(int i = 0; i<30; i++) { for(int j = 0; j<30; j++) { M[i][j] = j; N[i][j] = j; } } MatrixMulOnHost(M,N,P,Width); for(int i = 0; i<30; i++) { for(int j = 0; j<30; j++) { cout<<P[i][j]<<" "; } cout<<endl; } return 0; }
运行结果比较
