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_resultd_result[idx] = d_constArray[idx] + d_singleValue;// 從常量內存中的結構體讀取數據，將數組中的值和單個值相加，并存入 d_result_sd_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)