并行设计课笔记

CUDA硬件概述

• CPU

  • ALU计算能力强
  • 大容量Cache
  • 控制逻辑复杂

• GPU

  • Cache容量小
  • 控制逻辑简单
  • ALU高效节能
  • 通过大量线程来隐藏延迟

• CPU为了串行程序的高效执行进行了大量优化

• GPU专为大量数据的并行执行设计

• Nvidia的图形处理器产品类型

  • GeForce，面向消费者的图形处理产品
  • Quadro，工作站图形处理产品
  • Tegra，用于移动设备的芯片系列
  • Tesla，主要用于服务器高性能计算

CUDA程序的结构

• 获取GPU硬件信息
  • 得到可用的GPU卡的数量

  cudaError_t cudaGetDeviceCount(int *n)
  

  • 设置/获得当前使用的GPU卡的信息

  cudaError_t cudaGetDevice(int *d)
  cudaError_t cudaSetDevice(int d)
  

  • 查看GPU卡的特性（计算能力，内存等）

  cudaError_t cudaGetDeviceProperties(cudaDevProp* prop, int dev)
  struct cudaDeviceProp{
  	char name[256];
  	int major;
  	int minor;
  	...
  }
  

• CUDA函数调用错误的处理(当错误发生的时候，会给出错误发生的原因，文件名以及行号信息)

  cudaError_t error;
  error = cudaMalloc(gpuptr, N*sizeof(float));
  if(error != cudaSuccess){
  	printf("Error allocating memory\n");
  	exit(1);
  }
  

• 数据在GPU内存上的空间分配，释放
  • CPU和GPU拥有不同的内存空间（从CUDA6.0开始，可以使用统一的内存访问机制）
  • 分配和释放内存

  cudaMalloc( (void**) &pointer, size);
  cudaFree(pointer);
  

• 数据在CPU和GPU间的传输

  // dst_pointer 是数据的目的地指针
  // src_pointer 是数据的源地址指针
  // size 是要拷贝的数据的数量，以字节为单位
  // direction 表示数据的拷贝方向
  cudaMemcpy(dst_pointer, src_pointer, size, direction)
  

• GPU上内核函数的定义

  __host__, // 函数在主机上被调用和执行，类似普通c函数的编译
  __device__, // 函数在GPU上被调用和执行，由nvcc进行编译
  __global__, // 函数在主机上被调用，GPU上执行
  

• CPU端调用GPU上的内核函数
  • 调用内核函数的方式
    声明GPU上执行的内核函数的线程的组结构，线程网格和线程块
  • 线程网格
    • 二维结构
    • dim3 gridDim(2,4) 表示有8个线程块，按照2*4排列
  • 线程块
    • 三维结构
    • dim3 blockDim(32, 8, 1) 表示每个线程块有256个线程，按照3281排列
  • 调用内核函数

    kernelFunc<<<gridDim, blockDim>>>(param1, param2)
    

线程网格，线程块以及线程

• 线程的组织

  • 线程网格可以按照一维或者二维组织，gridDim.x，gridDim.y（y是行，x是列）表示线程网格在x和y方向上线程块的数量
  • 线程块可以按照一维，二维，三维进行组织，blockDim.x，blockDim.y，blockDim.z表示线程在x，y，z方向上线程的数量；blockIdx.x, blockIdx.y,blockIdx.z表示线程块在所属线程网格中的索引
  • threadIdx.x，threadIdx.y，threadIdx.z表示线程在所属线程块的索引

• 线程在执行中，以线程束为单位（每个GPU包含多个SM，每个SM有多个核，以线程束为单位进行调度）

• 线程块映射
  所有线程以线程块为单位分配给GPU的SM，实际分配的线程块数量和线程块内部的线程数目有关，并且受到最大可执行线程数限制

• 线程间的同步

  • 同一个线程块内部，使用__syncthreads()函数进行同步
  • 线程块间，需要使用全局内存进行同步