CUDA学习笔记(一)

说明：
  1.此乃本人学习CUDA过程中的笔记，不定期更新，如有错误，欢迎指出⁄(⁄ ⁄•⁄ω⁄•⁄ ⁄)⁄。
  2.由于我是初学，对CUDA的理解不多，排版不佳，还请原谅。（哎，高中作文都写不利索）
  3.每日一问：今天你买到显卡了吗[○･｀Д´･ ○]？

一.核函数
  作用：实现主机（CPU）对设备（GPU）的调用

  声明形式：

__global__ void hello_from_gpu()
{
    printf("Hello World from the GPU!\n");
}

   注：
    1.核函数必须被限定词__global__修饰，其中global前后是双下划线。也可以加上一些其他C++中的限定符，如static。限定符的次序可任意。
    2.核函数的返回必须是空类型，即void。
    3.在核函数中使用printf()需要包含头文件<stdio.h>或<cstdio>，核函数不支持iostream。
    4.函数名无特殊要求，而且支持C++中的重载。
    5.不支持可变数量的参数列表，即参数的个数必须确定。
    6.可以向核函数传递非指针变量，其内容对每个线程可见。
    7.除非使用统一内存编程机制，否则传给核函数的数组（指针）必须指向设备内存。
    8.核函数不可成为一个类的成员。通常的做法是用一个包装函数调用核函数，而将包装函数定义为类的成员。
    9.从计算能力3.5开始，引入了**动态并行（dynamic parallelism）**机制，在核函数内部可以调用其他核函数，甚至可以调用自己（递归函数）。

  调用格式：

hello_from_gpu<<<1, 1>>>();

   注：
    1.三对括号中的信息: 第一个数字是线程块的个数，第二个数字是每个线程块中的线程数。
    2.主机在调用一个核函数时，必须指明需要在设备中指派多少个线程，否则设备不知如何工作。

  CUDA线程组织：
    一个核函数的全部线程块构成一个网格（grid），而线程块的个数就记为网格大小（grid size）。每个线程块含有同样数目的线程，该数目成为线程块大小（block size）。

    <<<grid_size, block_size>>>

  grid_size和block_size，可以是一个结构体类型的变量，也可以是一个整型变量。
  此处最大支持三维网格，grid_size和block_size可以是类型为dim3的结构体变量。（dim3有x,y,z三个成员，在头文件vector_types.h中定义）
  核函数中的内建变量(build-in variable):
    gridDim:：相对应grid_size,类型为dim3
    blockDim：相对应block_size，类型为dim3
    blockIdx：相对应一个线程所在的线程块在一个网格中的线程块索引。blockIdx.x的取值范围是0到gridDim.x-1（对于y，z同理）。类型为uint3（uint3为包含x，y，z三个成员的结构体，在头文件vector_types.h中定义）
    threadIdx：一个线程在一个线程块中的线程索引，类型为uint3。
    wrapSize：线程束大小。一个线程束是同一个线程块中相邻的warpSize个线程。目前所有的GPU架构都是32。
  综上所述：一个线程块在网格中有三维坐标，一个线程在线程块中有三维坐标，定位线程需先找线程块，然后在线程块中找线程。

  网格与线程块大小的限制：
    网格大小在x、y和z这三个方向的最大允许值分别为$2^{31}-1$、65535和65535；线程块大小在x、y和z这三个方向的最大允许值分别为1024、1024和64。还要求线程块总的大小，即blockDim.x、blockDim.y和blockDim.z的乘积不能大于1024。

二.CUDA常用函数

  1.cudaDeviceSynchronize()
    作用：同步主机与设备，一般在核函数调用后使用。

  2.cudaError_t cudaMalloc(void **address, size_t size);
    作用：动态分布内存

address：待分配设备内存的指针。注意：因为内存（地址）本身就是一个指针，所以待分配设备内存的指针就是指针的指针，即双重指针。
size：待分配内存的字节数。
返回值是一个错误代号。如果调用成功，则返回cudaSuccess，否则返回一个代表某种错误的代号。

  3.cudaError_t cudaFree(void* address);
    作用：释放分配的设备内存

address：待释放的设备内存（不是双重指针）
返回值是一个错误代号，如果调用成功，返回cudaSuccess。

  4.cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)
    作用：主机和设备之间数据的传递

dst: 目标地址
src：源地址
count：复制数据的字节数
kind：枚举类型变量，标志数据传递方向，只能取如下几个值：
  1.cudaMemcpyHostToHost, 表示从主机复制到主机
  2.cudaMemcpyHostToDevice, 表示从主机复制到设备
  3.cudaMemcpyDeviceToHost, 表示从设备复制到主机
  4.cudaMemcpyDeviceToDevice, 表示从设备复制到设备
  5.cudaMemcpyDefault, 表示根据指针dst和src所指地址自动判断数据传输的方向。要求系统具有统一虚拟地址功能（要求64位的主机）。

  5.cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src, size_t count, size_t offset=0, cudaMemcpyKind kind = cudaMemcpyHostToDevice)

symbol：静态全局内存变量名
src：主机内存缓冲区指针
count：复制的字节数
offset：从symbol对应设备地址开始偏移的字节数
kind： 可选参数

cudaError_t cudaMemcpyFromSymbol(const void* dst, const void* symbol, size_t count, size_t offset=0, cudaMemcpyKind kind = cudaMemcpyDeviceToHost)

dst：主机内存缓冲区指针
symbol：静态全局内存变量名
count：复制的字节数
offset：从symbol对应设备地址开始偏移的字节数
kind： 可选参数

注：这两个函数的参数symbol可以是静态全局内存变量的变量名，也可以是常量内存变量的变量名。

三.CUDA中的头文件

  用nvcc编译器驱动编译.cu文件时，将自动包含必要的CUDA头文件，如<cuda.h>和<cuda_runtime.h>。
  device_launch_parameters.h——用以使用内建变量

四.设备函数

  定义：核函数可以调用不带执行配置的自定义函数，这样的自定义函数成为设备函数（device function）。它是在设备中执行，且在设备中调用的。与之相比，核函数是在设备中执行，但在主机端被调用的。

  CUDA函数执行空间标识符（前后均为双下划线）
    （1）__global__ 核函数，一般由主机调用，在设备中执行。如果使用动态并行，则也可以在核函数中调用自己或其他核函数。
    （2）__device__设备函数，只能被核函数或其他设备函数调用，在设备中执行。
    （3）__host__主机端的普通C++函数，在主机中调用，在主机中执行。有时可以用__host__和__device__同时修饰一个函数，使得该函数既是一个C++普通函数，又是一个设备加密手机。这样做可以减少代码冗余。编译器将针对主机和设备分别编译该函数。
    （4）不能同时用__global__和__device__修饰一个函数，也不能同时用__global__和__host__修饰同一个函数。
    （5）编译器决定把设备函数当作内联函数（inline function）或非内联函数，但可以用修饰符__noinline__建议一个设备函数为非内联函数（编译器不一定接收），也可以用修饰符__forceinline__建议一个设备函数为内联函数。
学习资料：
  CUDA编程:基础与实践