GPU program

客官°小女子只卖身不卖艺 2023-02-13 11:45 1阅读 0赞

### 文章目录 ###

*   *  GPU与CPU的区别
     *  CPU与GPU存储
     *  程序运行模型
     *   *  thread层面和内存层面
         *  简单的GPU编程

参考华盛顿大学\[Tianqi Chen的课件\]  
http://dlsys.cs.washington.edu/pdf/lecture5.pdf

## GPU与CPU的区别 ##

![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70]  
上图参考华盛顿大学[Tianqi Chen的课件][Tianqi Chen]  
CPU首先ALU的数目少于GPU,并且CPU在Fetch、Decode和Write back的负载是很高的。  
![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70 1]

## CPU与GPU存储 ##

![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70 2]

GPU的L1/SM的存储大于CPU的L1/core的存储。GPU的ShareMem可以被porgammer 使用。

## 程序运行模型 ##

SIMT : Single Instruction , Multiple Threads  
程序在C语言层面是单个线程，但是执行的代码是多个线程。

（2）多个线程在一个block里面  
（3）多个block在一个组 grid 里面

执行：a grid of blocks of threads

一个block 只能被一个SM执行，一个SM可以执行多个block。

在kernel 内部控制threads是warp, 一个warp包含32个threads，warp是kernel内部最基础的调度单元。一个block包含32个warps。  
在每一个cycle, warp scheduler 选择一个ready的warps，然后 dispatches warps到GPU的cores上面执行。

### thread层面和内存层面 ###

![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70 3]

### 简单的GPU编程 ###

更详细的指南可以[参考][Link 1]  
代码 main.cu：

/* * *__global__ type func() 声明可以从主机调用并在设备 * 上运行的代码 * * 通过 __device__ type func() 声明只能被其它的 * __devide__ 函数调用或者 __global__ 函数调用 * * * func<<<N, 1>>> 调用需要在GPU中执行的代码，其中参数 * N代表设备在执行改函数线程块的个数，参数1代表每个线程块 * 创建线程的个数 * * * cudaMalloc 分配的是global memoy 全局内存。 * int * int_dev; * cudaMalloc(&int_dev, N*sizeof(int)); * 不能在主机代码中使用cudaMalloc()分配的指针进行内存读/写操作 * * cudaMemecpy 在CPU和GPU之间分配内存 * 将cpu数据b复制到int_dev * cudaMemcpy(int_dev, b, N*sizeof(int), cudaMemcpyHostToDevice); * * 将GPU的数据int_dev 复制到b * cudaMemcpy(b, int_dev, N* sizeof(int), cudaMemcpyDeviceToHost); * * * cudaFree释放在GPU中分配的内存 * cudaFree(int_dev) * * * cudaDeviceProp 结构体包含了当前CUDA设备的所有属性 * cudaGetDeviceCount(&count); 获取当前计算机的cuda设备的数量 * * cudaGetDevice(&dev) 获取当前使用设备的ID,ID从0开始编号 * * cudaGetDeviceProperties(&prop, i) 获取设备i的属性 * * cudaChooseDevice(&dev, &prop) 获取符合要求的设备ID到dev * * cudaSetDevice(dev) 使用ID为dev的设备 * 参考：http://docs.nvidia.com/cuda */
     
    #include <stdio.h>
    #include <cuda_runtime.h>
    #include "cuda.h"
    
     __global__
     void saxpy(int n, float a, float *x, float *y)
     { 
       int i = blockIdx.x*blockDim.x + threadIdx.x;
       if (i < n) y[i] = a*x[i] + y[i];
     }
    
     int main(void)
     { 
       int N = 1<<20;
       float *x, *y, *d_x, *d_y;
       x = (float*)malloc(N*sizeof(float));
       y = (float*)malloc(N*sizeof(float));
    
       cudaMalloc(&d_x, N*sizeof(float));
       cudaMalloc(&d_y, N*sizeof(float));
    
       for (int i = 0; i < N; i++) { 
         x[i] = 1.0f;
         y[i] = 2.0f;
       }
    
       cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
       cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);
    
       // Perform SAXPY on 1M elements
       saxpy<<<(N+255)/256, 256>>>(N, 2.0f, d_x, d_y);
    
       cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);
    
       float maxError = 0.0f;
       for (int i = 0; i < N; i++)
         maxError = max(maxError, abs(y[i]-4.0f));
       printf("Max error: %f\n", maxError);
    
       cudaFree(d_x);
       cudaFree(d_y);
       free(x);
       free(y);
     }

编译执行

nvcc main.cu -o main

// 点积
    const int N = 33*1024;
    const  int threadsPerBlock = 256;
    
    const int blocks = min(32, (N + threadsPerBlock -1)/256);
    
    // 未来确保性能，blocks的数目少于32
    __global__ void dot(float*a, float*b, float*c) { 
        __shared__ float cache[threadsPerBlock];
    
        int thread_id = blockDim.x * blockIdx.x + ThreadIdx.x;
    
        float tmp = 0;
    
        while (thread_id < N) { 
            tmp += a[thread_id]*b[thread_id];
            thread_id +=  * gridDim.x;
        }
    
        int cache_index =  ThreadIdx.x;
        cache[cache_index] =  tmp;
        __syncthreads();
    
        int i = blockDim.x /2;
        while (i !=0) { 
            if (cache_index < i) { 
                cache[cache_index] += cache_index[cache_index + i]
            }
            __syncthreads();
            i /=2;
        }
        if (cache_index ==0) { 
            c[blockIdx.x] = cache[0];
        }

[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70]: https://img-blog.csdnimg.cn/20200529090642354.png?x-oss-process=image/watermark,type_ZmFuZ3poZW5naGVpdGk,shadow_10,text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ==,size_16,color_FFFFFF,t_70
[Tianqi Chen]: http://dlsys.cs.washington.edu/pdf/lecture5.pdf
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70 1]: https://img-blog.csdnimg.cn/20200529090926799.png?x-oss-process=image/watermark,type_ZmFuZ3poZW5naGVpdGk,shadow_10,text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ==,size_16,color_FFFFFF,t_70
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70 2]: https://img-blog.csdnimg.cn/20200529093056348.png?x-oss-process=image/watermark,type_ZmFuZ3poZW5naGVpdGk,shadow_10,text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ==,size_16,color_FFFFFF,t_70
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70 3]: https://img-blog.csdnimg.cn/20200529102054149.png?x-oss-process=image/watermark,type_ZmFuZ3poZW5naGVpdGk,shadow_10,text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ==,size_16,color_FFFFFF,t_70
[Link 1]: http://notes.maxwi.com/2015/06/11/CUDA-study-notes/