CUDA的全局内存、常量内存 、纹理内存、共享内存、寄存器

CUDA 编程模型中有几种不同类型的内存，每种内存都有其特定的用途和特性。以下是主要的几种 CUDA 内存类型：

1. 全局内存 (Global Memory)

• 描述：GPU 的主内存，所有线程都可以访问。
• 特性：
  • 大容量，但访问速度较慢。
  • 适合存储大量数据。

2. 常量内存 (Constant Memory)

• 描述：只读内存，适用于在内核执行期间不改变的值。
• 特性：
  • 访问速度较快（当所有线程访问相同地址时）。
  • 最大容量为 64KB。

3. 纹理内存 (Texture Memory)

• 描述：专门用于处理图像和纹理数据的只读内存。
• 特性：
  • 提供缓存机制，适合进行空间局部性访问。
  • 支持各种过滤和坐标变换。

4. 共享内存 (Shared Memory)

• 描述：线程块内部的共享内存，可以被同一线程块中的所有线程访问。
• 特性：
  • 访问速度非常快，适合线程间的数据共享。
  • 最大容量通常为 48KB（依赖于 GPU 架构）。

5. 寄存器 (Registers)

• 描述：每个线程的私有内存，存储线程的局部变量。
• 特性：
  • 访问速度最快，但数量有限。
  • 每个线程有其独立的寄存器。

6. 主机内存 (Host Memory)

• 描述：CPU 的内存，通常用于存储数据和指令。
• 特性：
  • 需要通过 CUDA API 进行数据传输到 GPU。
  • 访问速度较慢，相比于 GPU 内存。

#include <stdio.h>

__constant__ float constantArray[256]; // 常量内存声明

__global__ void kernel(float *globalArray) {
    // 使用寄存器
    float regValue = threadIdx.x * 1.0f;

    // 使用全局内存
    globalArray[threadIdx.x] = regValue;

    // 使用共享内存
    __shared__ float sharedValue[256]; // 共享内存声明
    sharedValue[threadIdx.x] = regValue;
    __syncthreads(); // 确保所有线程都写入共享内存

    // 使用常量内存
    float constantValue = constantArray[threadIdx.x];

    // 打印结果
    printf("Thread %d: Global = %f, Shared = %f, Constant = %f\n", 
           threadIdx.x, globalArray[threadIdx.x], sharedValue[threadIdx.x], constantValue);
}

int main() {
    const int arraySize = 256;
    float *d_globalArray;

    // 分配全局内存
    cudaMalloc((void**)&d_globalArray, arraySize * sizeof(float));

    // 初始化常量内存
    float h_constantArray[arraySize];
    for (int i = 0; i < arraySize; i++) {
        h_constantArray[i] = (float)i;
    }
    cudaMemcpyToSymbol(constantArray, h_constantArray, arraySize * sizeof(float));

    // 启动内核
    kernel<<<1, arraySize>>>(d_globalArray);

    // 等待内核完成
    cudaDeviceSynchronize();

    // 释放全局内存
    cudaFree(d_globalArray);

    return 0;
}