GPU program

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

GPU与CPU的区别

上图参考华盛顿大学Tianqi Chen的课件
CPU首先ALU的数目少于GPU,并且CPU在Fetch、Decode和Write back的负载是很高的。

CPU与GPU存储

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层面和内存层面

简单的GPU编程

更详细的指南可以参考
代码 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];
    }