建议先看看前言中关于存储器的介绍：

 

线程

首先介绍进程，进程是程序的一次执行，线程是进程内的一个相对独立的可执行的单元。若把进程称为任务的话，那么线程则是应用中的一个子任务的执行。举个简单的例子：一个人要做饭，食谱就是程序代码，做的过程就是执行程序，做好的饭就是程序运行的结果，而在这期间，需要炒菜，放盐，放油等等就是线程。

 

 

线程同步

调用__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<<
      
       >>(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<<
      
       >>(A,B,C);}
      

 

 

线程块之间彼此独立执行

任意顺序：并行或串行

被任意数量的处理器以任意顺序调度：处理器的数量具有可扩展性

 

一个块内部的线程

共享容量有限的低延迟存储器

同步执行：合并访存，__syncThreads()：Barrier ——块内线程一起等待所有线程都执行某处语句，轻量级

 

 

实例演示

编写一个矩阵 M，N 相乘的程序,维度一致。

先介绍几个函数：

 

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 <<
    
     >>(Md,Nd,Pd);
    

完整GPU程序如下：

 

/* ============================================================================ Name        : Martix.cu Author      : wzn Version     : Copyright   : Your copyright notice Description : CUDA compute reciprocals ============================================================================ */#include 
     
      #include 
      
       #include 
       
        __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 <<
        
         >>(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<
         
          <<" ";              }              std::cout<<"\n";          }    return 0;}CPU端程序:#include 
          
           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
          
         
        
       
      
     

 

运行结果比较