1、全局内存

特点：容量最大，访问延时最大，所有线程都可以访问。 线性内存。
cudaMalloc() 动态分配内存
cudaFree() 释放内存
_device_ int buff[N] 使用固定长度的静态全局内存
动态内存使用cudaMemcpy 在host 和 device之间传输
静态内存 host 和 device 之间传输数据：
cudaMemcpyToSymbol() host 传到device
cudaMemcpyFromSymbol() device传到host

#include <cuda.h>
#include <cuda_runtime.h>
__device__ int d_x = 1;
__device__ int d_y[2];
__global__ void cudaOut(void) {
	d_y[0]+= d_x;
	d_y[1]+= d_x;
}
int main(void) {
	int h_y[2] = { 10,20 };
	CHECK(cudaMemcpyToSymbol(d_y, h_y, sizeof(int) * 2));
	cudaOut<<<1,1>>>();
	CHECK(cudaDeviceSynchronize());
	CHECK(cudaMemcpyFromSymbol(h_y, d_y, sizeof(int) * 2));
	return 0;
}

2、局部内存

特点：存储线程私有的局部变量，在寄存器内存不足时使用。

3、共享内存

特点：片上内存，访问速度快，线程块内共享。
使用：存储线程块中的共享数据，加速线程间的数据处理
每个SM的共享内存数量是一定的，也就是说，如果在单个线程块中分配过度的共享内存，将会限制活跃线程束的数量；合适分配单个线程块的共享内存，使得SM的使用率最大化，起到加速的作用。

例如：blockSize = 128，一个SM有2048个线程，那么一个SM能同时处理16个block。如果SM有96K的共享内存，每个block则分配96 / 16 = 6K，太大其他block无法获得使用。

__syncthreads 通常用于协调同一块中线程间的通信。

__global__ void kernel_function(parameters) {
     // 假设算法的特殊设计导致对变量 data 的访问不得不是非合并的
     // 1. 定义共享内存
     __shared__ float s_data [data_size];
     // 2. 将 data 复制到共享内存
     s_data[copy_index] = data[copy_index];
     // 3. 等待线程块中所有线程完成复制操作
     __syncthreads();
     // 4. 进行操作（包含非合并内存访问）
     operations(data);
}

3.1 静态共享内存

// 静态共享内存
__global__ void staticReverse(int* d, int n){
   	__shared__ int s[64];
    int t = threadIdx.x;
    s[t] = d[t];
    __syncthreads();
}
int main(void){
staticReverse << <1, n >> > (d_d, n);
}

3.2 动态共享内存

调用时指定共享内存大小
如果一个共享内存的大小在编译时是未知的, 则需要添加 extern修饰，在内核调用时动态分配。

__global__ void dynamicReverse(int* d, int n){
    extern __shared__ int s[];
    int t = threadIdx.x;
    s[t] = d[t];
}
int main(void){
dynamicReverse << <1, n, n * sizeof(int) >> > (d_d, n);
}

4、纹理内存

特点：优化二维数据的访问，具有缓存加速
使用：适用图像处理和大规模数据访问

5、常量内存

特点：存储只读数据，访问速度快，广播式访问。数量有限，最多64KB
使用：频繁访问的常量数据，所有线程块都能访问，全局可见。
使用_constant_修饰
只能通过cudaMemcpyToSymbol() 或**cudaMemcpyToSymbolAsync()**进行数据传输

const int N = 4;
// 定义结构体
struct ConstStruct {
    float array[N];
    float singleValue;
};
// 常量内存变量声明
__constant__ float d_constArray[N];    // 数组
__constant__ float d_singleValue;      // 单个变量
__constant__ ConstStruct d_constStruct; // 结构体
// 定义核函数
__global__ void kernelFunction(float* d_result, float* d_result_s) {
    int idx = threadIdx.x;
    if (idx < N) {
        // 从常量内存中读取常量数组和单个值，将结果存入 d_result
        d_result[idx] = d_constArray[idx] + d_singleValue;
        // 从常量内存中的结构体读取数据，将数组中的值和单个值相加，并存入 d_result_s
        d_result_s[idx] = d_constStruct.array[idx] + d_constStruct.singleValue;
    }
}
int main() {
    // 将数据从主机复制到常量内存
    CHECK(cudaMemcpyToSymbol(d_constArray, h_array, N * sizeof(float)));
    CHECK(cudaMemcpyToSymbol(d_singleValue, &h_singleValue, sizeof(float)));
    CHECK(cudaMemcpyToSymbol(d_constStruct, &h_constStruct, sizeof(ConstStruct)));
	kernelFunction << <1, N >> > (d_result, d_result_s);
}

6、寄存器内存

特点：片上内存，访问速度最快
使用：局部变量

7、用CUDA运行时API函数查询设备

该段介绍用 CUDA 运行时 API 函数查询所用 GPU 的规格 ，可以通过以下代码查看显卡的信息：

#include <iostream>
#include <cuda_runtime.h>
#include "error_check.cuh"
int main(int argc, char* argv[]) {
	int device_id = 0;//如果你不止一个显卡，可以切换ID，输出不同显卡的信息
	if (argc > 1) device_id = atoi(argv[1]);
	CHECK(cudaSetDevice(device_id));
	cudaDeviceProp prop;
	CHECK(cudaGetDeviceProperties(&prop, device_id));
	printf("Device id: %d\n", device_id);
	printf("Device name: %s\n", prop.name);
	printf("Compute capability: %d.%d\n", prop.major, prop.minor);
	printf("Amount of global memory: %g GB\n", prop.totalGlobalMem / (1024.0 * 1024 * 1024));
	printf("Amount of constant memory: %g KB\n", prop.totalConstMem / 1024.0);
	printf("Maximum grid size: %d %d %d\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
	printf("Maximum block size: %d %d %d\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
	printf("Number of SMs: %d\n", prop.multiProcessorCount);
	printf("Maximum amount of shared memory per block: %g KB\n", prop.sharedMemPerBlock / 1024.0);//每个线程块可以使用的最大共享内存
	printf("Maximum amount of shared memory per SM: %g KB\n", prop.sharedMemPerMultiprocessor / 1024.0);//个SM可以分配的最大共享内存总量
	printf("Maximum number of registers per block: %d K\n", prop.regsPerBlock / 1024);//每个线程块可以使用的最大寄存器数量
	printf("Maximum number of registers per SM: %d K\n", prop.regsPerMultiprocessor / 1024);//每个SM可以分配的最大寄存器总量
	printf("Maximum number of threads per block: %d\n", prop.maxThreadsPerBlock);
	printf("Maximum number of threads per SM: %d\n", prop.maxThreadsPerMultiProcessor);//每个SM可以同时运行的最大线程数量
	return 0;
}

CUDA 错误检测

#define CHECK(call)                 \    
do                                  \
{                                   \
    const cudaError_t error_code = call;                \
    if (error_code != cudaSuccess)                  \
    {                                   \
    printf("CUDA ERROR:\n");                    \
        printf("    FILE:   %s\n", __FILE__);           \
        printf("    LINE:   %d\n", __LINE__);           \
        printf("    ERROR CODE: %d\n", error_code);         \
        printf("    ERROR TEXT: %s\n", cudaGetErrorString(error_code)); \
    exit(1);                            \
    }                                   \
} while(0)