https://www.nvidia.cn/training/

课程示例：使用 CUDA C/C++ 加速应用程序

nvidia-smi (Systems Management Interface) 命令行命令查询GPU信息

GPU应用程序代码示例

<code>.cu 文件（.cu 是 CUDA 加速程序的文件扩展名）

void CPUFunction(){

printf("This function is defined to run on the CPU.\n");

}

__global__ void GPUFunction(){

printf("This function is defined to run on the GPU.\n");

}

int main(){

CPUFunction();

GPUFunction<<<1, 1>>>();

cudaDeviceSynchronize();

}

__global__ 关键字表明以下函数将在 GPU 上运行并可全局调用，而在此种情况下，则指由 CPU 或 GPU 调用。

通常，我们将在 CPU 上执行的代码称为主机代码，而将在 GPU 上运行的代码称为设备代码。

注意返回类型为 void。使用 __global__ 关键字定义的函数需要返回 void 类型。

通常，当调用要在 GPU 上运行的函数时，我们将此种函数称为已启动的核函数。

启动核函数时，我们必须提供执行配置，即在向核函数传递任何预期参数之前使用 <<< ... >>> 语法完成的配置。

在宏观层面，程序员可通过执行配置为核函数启动指定线程层次结构，从而定义线程组（称为线程块）的数量，以及要在每个线程块中执行的线程数量。

与许多 C/C++ 代码不同，核函数启动方式为异步：CPU 代码将继续执行而无需等待核函数完成启动。

调用 CUDA 运行时提供的函数 cudaDeviceSynchronize 将导致主机 (CPU) 代码暂作等待，直至设备 (GPU) 代码执行完成，才能在 CPU 上恢复执行。

编译并运行加速后的CUDA代码

CUDA 平台附带 NVIDIA CUDA 编译器 nvcc，可以编译 CUDA 加速应用程序，其中包含主机和设备代码。

http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html

nvcc -arch=sm_70 -o out some-CUDA.cu -run

arch 标志表示该文件必须编译为哪个架构类型。本示例中，sm_70 将用于专门针对本实验运行的 Volta GPU 进行编译，但有意深究的用户可以参阅有关 arch 标志、虚拟架构特性 和 GPU特性 的文档。

为方便起见，提供 run 标志将执行已成功编译的二进制文件。

CUDA的线程层次结构

<<< NUMBER_OF_BLOCKS, NUMBER_OF_THREADS_PER_BLOCK>>>

启动核函数时，核函数代码由每个已配置的线程块中的每个线程执行。

因此，如果假设已定义一个名为 someKernel 的核函数，则下列情况为真：

someKernel<<<1, 1>>() 配置为在具有单线程的单个线程块中运行后，将只运行一次。

someKernel<<<1, 10>>() 配置为在具有 10 线程的单个线程块中运行后，将运行 10 次。

someKernel<<<10, 1>>() 配置为在 10 个线程块（每个均具有单线程）中运行后，将运行 10 次。

someKernel<<<10, 10>>() 配置为在 10 个线程块（每个均具有 10 线程）中运行后，将运行 100 次。

分配将要在GPU和CPU上访问的内存

int a;

// Note the address of <code>a is passed as first argument.

cudaMallocManaged(&a, 10241024*2);

// Use a on the CPU and/or on any GPU in the accelerated system.

cudaFree(a);

如何处理块配置与所需线程数不匹配

鉴于 GPU 的硬件特性，所含线程的数量为 32 的倍数的线程块是最理想的选择，因其具备性能上的优势。假设我们要启动一些线程块且每个线程块中均包含 256 个线程（32 的倍数），并需运行 1000 个并行任务（此处使用极小的数量以便于说明），则任何数量的线程块均无法在网格中精确生成 1000 个总线程，因为没有任何整数值在乘以 32 后可以恰好等于 1000。

这个问题可以通过以下方式轻松地解决：

编写执行配置，使其创建的线程数超过执行分配工作所需的线程数。

将一个值作为参数传递到核函数 (N) 中，该值表示要处理的数据集总大小或完成工作所需的总线程数。

计算网格内的线程索引后（使用 threadIdx + blockIdx*blockDim），请检查该索引是否超过 N，并且只在不超过的情况下执行与核函数相关的工作。

以下是编写执行配置的惯用方法示例，适用于 N 和线程块中的线程数已知，但无法保证网格中的线程数和 N 之间完全匹配的情况。如此一来，便可确保网格中至少始终拥有 N 所需的线程数，且超出的线程数至多仅可相当于 1 个线程块的线程数量：

// Assume N is known

int N = 100000;

// Assume we have a desire to set threads_per_block exactly to 256

size_t threads_per_block = 256;

// Ensure there are at least N threads in the grid, but only 1 block’s worth extra

size_t number_of_blocks = (N + threads_per_block - 1) / threads_per_block;

some_kernel<<<number_of_blocks, threads_per_block>>>(N);

由于上述执行配置致使网格中的线程数超过 N，因此需要注意 some_kernel 定义中的内容，以确保 some_kernel 在由其中一个 ”额外的” 线程执行时不会尝试访问超出范围的数据元素：

global some_kernel(int N){

int idx = threadIdx.x + blockIdx.x * blockDim.x;

if (idx < N){ // Check to make sure idx maps to some value within N

// Only do work if it does

}

}

根据dataindex判断是否需要执行

数据集比网格大

在核函数中使用跨网格循环。

在跨网格循环中，每个线程将在网格内使用 threadIdx + blockIdx*blockDim 计算自身唯一的索引，并对数组内该索引的元素执行相应运算，然后将网格中的线程数添加到索引并重复此操作，直至超出数组范围。例如，对于包含 500 个元素的数组和包含 250 个线程的网格，网格中索引为 20 的线程将执行如下操作：

对包含 500 个元素的数组的元素 20 执行相应运算

将其索引增加 250，使网格的大小达到 270

对包含 500 个元素的数组的元素 270 执行相应运算

将其索引增加 250，使网格的大小达到 520

由于 520 现已超出数组范围，因此线程将停止工作

CUDA 提供一个可给出网格中线程块数的特殊变量：gridDim.x。然后计算网格中的总线程数，即网格中的线程块数乘以每个线程块中的线程数：gridDim.x * blockDim.x。带着这样的想法来看看以下核函数中网格跨度循环的详细示例：

__global void kernel(int *a, int N){

int indexWithinTheGrid = threadIdx.x + blockIdx.x * blockDim.x;

int gridStride = gridDim.x * blockDim.x;

for (int i = indexWithinTheGrid; i < N; i += gridStride) {

// do work on a[i];

}

}

错误处理

cudaError_t err = cudaMallocManaged(&a, N);// Assume the existence of <code>a and N.

if (err != cudaSuccess) { // cudaSuccess is provided by CUDA.

printf(“Error: %s\n”, cudaGetErrorString(err)); // cudaGetErrorString is provided by CUDA.

}

someKernel<<<1, -1>>>(); // -1 is not a valid number of threads.

cudaError_t err = cudaGetLastError(); // cudaGetLastError will return the error from above.

if (err != cudaSuccess){

printf(“Error: %s\n”, cudaGetErrorString(err));

}