GPU架构理解和CUDA编程优化

GPU架构及CUDA基础

理解cuda core，sm，sp

https://blog.csdn.net/asasasaababab/article/details/80447254

• 显存

显存（Global Memory）：显存是在GPU板卡上的DRAM，类似于CPU的内存，就是那堆DDR啊，GDDR5啊之类的。特点是容量大（可达16GB），速度慢，CPU和GPU都可以访问。

• 计算单元

计算单元（Streaming Multiprocessor）：执行计算的。每一个SM都有自己的控制单元（Control Unit），寄存器（Register），缓存（Cache），指令流水线（execution pipelines）

• 在GP100里，每一个SM有两个SM Processing Block（SMP），里边的绿色的就是CUDA Core，CUDA core也叫Streaming Processor（SP），这俩是一个意思。每一个SM有自己的指令缓存，L1缓存，共享内存。而每一个SMP有自己的Warp Scheduler、Register File等。要注意的是CUDA Core是Single Precision的，也就是计算float单精度的。双精度Double Precision是那个黄色的模块。所以一个SM里边由32个DP Unit，由64个CUDA Core，所以单精度双精度单元数量比是2:1。LD/ST 是load store unit，用来内存操作的。SFU是Special function unit，用来做cuda的intrinsic function的，类似于__cos()这种。
• SP：最基本的处理单元，streaming processor，也称为CUDA
  core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算，也就是很多个SP同时做处理。
• SM：多个SP加上其他的一些资源组成一个streaming multiprocessor。也叫GPU大核，其他资源如：warp scheduler，register，shared memory等。SM可以看做GPU的心脏（对比CPU核心），register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此，这些有限的资源就使每个SM中active warps有非常严格的限制，也就限制了并行能力。

理解threaad、block、grid、warp

https://blog.csdn.net/junparadox/article/details/50540602

• thread，block，grid，warp是CUDA编程上的概念，以方便程序员软件设计，组织线程，同样的我们给出一个示意图来表示。
• thread：一个CUDA的并行程序会被以许多个threads来执行。
• block：数个threads会被群组成一个block，同一个block中的threads可以同步，也可以通过shared
  memory通信。
• grid：多个blocks则会再构成grid。
• warp：GPU执行程序时的调度单位，目前cuda的warp的大小为32，同在一个warp的线程，以不同数据资源执行相同的指令,这就是所谓SIMT。

理解cuda stream

https://blog.csdn.net/dcrmg/article/details/55107518
CUDA流表示一个GPU操作队列，该队列中的操作将以添加到流中的先后顺序而依次执行。可以将一个流看做是GPU上的一个任务，不同任务可以并行执行。使用CUDA流，首先要选择一个支持设备重叠（Device Overlap）功能的设备，支持设备重叠功能的GPU能够在执行一个CUDA核函数的同时，还能在主机和设备之间执行复制数据操作。

理解Memory type

https://www.jianshu.com/p/3d4c9cc3a777
https://www.cnblogs.com/mtcnn/p/9411864.html
对CUDA架构而言，主机端的内存被分为两种，一种是可分页内存（pageable memroy）和页锁定内存（page-lock或 pinned）。可分页内存是由操作系统API malloc()在主机上分配的，页锁定内存是由CUDA函数cudaHostAlloc()在主机内存上分配的，页锁定内存的重要属性是主机的操作系统将不会对这块内存进行分页和交换操作，确保该内存始终驻留在物理内存中。

实践meory access–shared memory

矩阵乘以向量优化

https://m.baidu.com/sf_edu_wenku/view/f6d045ac185f312b3169a45177232f60dccce716

原始的方法

__global__ void cal_dis_floatgpunew(float *train_data, float *test_data, float *dis, int cntN)
{

    
    int tid = blockIdx.x*blockDim.x + threadIdx.x;// blockIdx.x;
    //printf("tid:%d，D:%d,%lf\n", tid, 0, *((float*)((char*)train_data + tid * pitch) + 0));
    if (tid < cntN)
    {
        float temp = 0;

        for (int i = 0; i < 512; i++)
        {
            temp += train_data[tid * 512 + i] * test_data[i];
 
        dis[tid] = temp; //Test
  
    }

缺点：必须按照矩阵维度来规划网格的尺寸，效率比较低。

优化1

一个线程处理 向量（5121）与矩阵（512cntN）的一行相乘，要求每个block的线程数大于512。对于矩阵的列维度没有特殊要求，原则上可以处理任意列维度的数据。

__global__ void cal_dis_floatgpu_wpf_mod1(float *train_data, float *test_data, float *dis, int cntN)
/*
//cntN 
//dis
//train_data
//test_data
//----- 
*/
{

	const int vector_dim=512;
    int tid = blockIdx.x*blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;
    __shared__ float test_data_shared[vector_dim];
 	
	
	if( threadIdx.x<=vector_dim)
	test_data_shared[ threadIdx.x]=test_data[ threadIdx.x];
    
	//printf("tid:%d，D:%d,%lf\n", tid, 0, *((float*)((char*)train_data + tid * pitch) + 0));
    for(int jj=tid;jj<cntN;jj+=stride)
    {
        float temp = 0;

        for (int i = 0; i < vector_dim; i++)
        {
            temp += train_data[i* cntN + jj] * test_data_shared[i];
            
        }
        dis[jj] = temp; //Test
       
    }
}

优化2

要求每个block的线程数量要大于等于512，每个block去处理一行矩阵数据和向量的点乘及加和。首先是在block中每个线程对512维向量做点乘，存储到temp中，然后对每个block中的temp存储到共享内存，对每个block中的共享内存进行二分求和，得到点乘的加和结果，然后按照block的id存储每个共享内存中的第一个结果。

__global__ void cal_dis_floatgpu_wpf_mod2(float *train_data, float *test_data, float *dis, int cntN)
/*
//cntN 
//dis
//train_data
//test_data
//----- 
*/
{

    // int tid = blockIdx.x;// *blockDim.x + threadIdx.x;// blockIdx.x;

        int tid = threadIdx.x;
        int bid = blockIdx.x;
        int stride = blockDim.x ;

        const int vector_dim=512;
        __shared__ float test_data_shared[vector_dim];
        __shared__ float dot_multiply_shared[vector_dim];

        if(tid<vector_dim)
        test_data_shared[tid]=test_data[tid];
        float temp=0.0;

        for (int j=bid;j<cntN;j+=gridDim.x)
        {
                temp=0.0;
                //for(int i=tid;i<vector_dim;i+=blockDim.x)
                if(tid<vector_dim)
                {
                        //temp=train_data[j*vector_dim+tid]*test_data_shared[tid];
                        temp=train_data[tid*cntN+j]*test_data_shared[tid];
                        dot_multiply_shared[tid]=temp;
                }
                //dot_multiply_shared[tid]=temp;
                __syncthreads();

                int jj=vector_dim/2;
                while(jj)
                {
                if(tid<jj)
                        dot_multiply_shared[tid] += dot_multiply_shared[tid + jj];
                        __syncthreads();
                        jj/=2;
                }
                if(tid==0)
                {
                        dis[j]=dot_multiply_shared[0];
                }

         __syncthreads();
        }
}

优化3

使用CUBLAS库
nvcc vector_multiply.cu -L /usr/local/cuda/lib64/ -lcublas

void cal_dis_floatgpu_wpf_mod3(float *train_data, float *test_data, float *dis, int cntN)
{
                cudaError_t cudastat;
                cublasStatus_t stat;
                cublasHandle_t handle;

                float alf=1.0;
                float beta=0;
                stat=cublasCreate(&handle);
                stat=cublasSgemv(handle,CUBLAS_OP_T,512,cntN,&alf,train_data,512,test_data,1,&beta,dis,1);//swap col and row
                cublasDestroy(handle);

                return ;

}

cublas
优点：快捷、方便；缺点：效率偏低