1 CUDA程序基本框架
单源文件的CUDA程序，基本框架如下 头文件包含
#include <stdio.h>
#include <cuda_runtime.h>
常量定义（宏函数）
#define xxx
const double EPSILON = 1.0e-10;
C++自定义函数和CUDA函数声明
// 核函数。
__global__ void add(const double *x, const double *y, double *z, const int N);
主函数
int main()
{
// 申请主机和device内存（cudaMalloc）
// 初始化主机数据
// 从主机复制数据到设备（cudaMemcpy）
// 调用核函数在设备中执行计算
add<<<grid_size, block_size>>>(d_x, d_y, d_z, N);
// 从设备复制数据到主机（cudaMemcpy）
cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
// 释放主机内存。
// free(h_x)
// 释放设备内存。
cudaFree(d_z);
}
C++自定义函数和CUDA函数实现
__global__ void add(const double *x, const double *y, double *z, const int N) {
xxx
}
2 CUDA内存模型
- registers
- shared mem
- local mem
- constant mem
- texture mem
- global mem
CUDA内存模型分为以下层次：
thread: 每个线程独有registers（寄存器）、local mem（本地内存）； block：每个block（线程块）独有shared mem（共享内存），所有线程块内的线程共享shared mem； grid：每个grid独有global mem（全局内存）、constant mem（常量内存）、texture memory（纹理内存），不同线程块中线程可使用。
registers（寄存器） 每个线程私有，速度快。存放核函数中定义的、不加任何限定符的变量。
local mem（本地内存） 每个线程私有；本地内存是全局内存一部分，速度慢。线程使用过多寄存器、或声明大型结构体或数组、编译器无法确定数组大小等情况，线程私有数据会分配到local mem。
shared mem（共享内存） __shared__修饰符修饰的变量。block中线程共有，速度与registers一样快。 使用共享内存必须调用如下函数进行同步：
void __syncthreads()
constant mem（常量内存） __constant__修饰符进行修饰。只读；有缓存；空间小。
texture mem（纹理内存） 访问二维数据的线程可以达到最优性能。具有纹理缓存，只读。
global mem（全局内存） GPU中容量最大、延迟最高的内存空间。所有线程都可以访问；没有缓存
参考：
blog.csdn.net/QLeelq/arti…
3 栅栏机制
CUDA中的栅栏（Barrier）机制是一种同步机制，用于确保在多个线程块Block中的线程都完成某个任务。
__syncthreads() 实现一个线程块中所有线程按照代码出现的顺序执行指令，但是不同线程块之间依然是独立、异步的。
在某些情况，线程Block中的线程需要共享数据、协作，线程之间的执行速度可能有差异。可以通过插入栅栏，确保所有线程都完成上一步操作，避免数据竞争和不确定行为。
4 规约运算（Reduction）
规约运算是一种并行计算模式。将一个数组或向量的元素进行累积操作，得到一个单一结果。例如计算数组的最大值、最小值。
使用全局内存计算
// real *d_x, real *d_y为全局内存
void __global__ reduce_global(real *d_x, real *d_y)
{
const int tid = threadIdx.x;
// 定义一个寄存器指针变量来作为临时的缓存，指向每个线程块的起始地址
real *x = d_x + blockDim.x * blockIdx.x; // blockDim.x要是2的指数次方的整数
// 折半归约，用位运算代替/2，在和函数中更高效
for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
{
if (tid < offset)
{
x[tid] += x[tid + offset];
}
__syncthreads();
}
// 寄存器内存生命周期在和函数里，将值保存到全局内存里。
// 保证一个线程块中，仅执行一次。将寄存器缓存里的每个线程块里的第一个元素值赋值给全局内存，后面主机内存进行所有线程块首元素值相加即是数组归约值。
if (tid == 0) // 保证一个线程块中，仅执行一次
{
d_y[blockIdx.x] = x[0];
}
}
每个线程块内独立的对其中的数据进行归约。同步函数在每个线程块执行之后使用。每个线程块之间的计算执行不是顺序的，但这不影响结果的正确性。因为在核函数中，每个线程块是独立的处理不同的数据，相互之间没有依赖。
使用shared内存计算 全局内存不够高效，寄存器内存仅对单个线程可见，使用对整个线程块可见的共享内存来提高性能
void __global__ reduce_shared(real *d_x, real *d_y)
{
const int tid = threadIdx.x;
const int bid = blockIdx.x;
const int n = bid * blockDim.x + tid;
// 定义共享内存
__shared__ real s_y[128];
// 初始化
s_y[tid] = (n < N) ? d_x[n] : 0.0;
// 使用之前要使用线程块同步函数
__syncthreads();
for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
{
if (tid < offset)
{
s_y[tid] += s_y[tid + offset];
}
__syncthreads();
}
// 共享内存生命周期在和函数里，将值保存到全局内存里
if (tid == 0)
{
d_y[bid] = s_y[0];
}
}
参考：
blog.csdn.net/weixin_4131…
5 原子函数
原子函数（Atomic Functions）是一组特殊函数，用于在多个线程同时访问和修改共享内存时提供原子操作支持。原子函数确保多个线程按照预期顺序执行读取、修改和写入共享内存操作，不会导致数据竞争或不一致。
- atomicAdd(int *addr, int v)
- atomicSub(int *addr, int v)
- atomicMax(int *addr, int v)
- atomicExch(int *addr, int v)
