// 在设备上分配 count 字节的全局内存，并通过 devPtr 指针返回该内存的地址
cudaError_t cudaMalloc(void **devPtr, size_t count);

// 使用存储 value 的值（CPU 端）来初始化设备内存 devPtr 处开始的 count 字节
cudaError_t cudaMemset(void **devPtr, int value, size_t count);

// 销毁 devPtr 指向的全局内存，该内存必须在此前使用 cudaMalloc 分配过
cudaError_t cudaFree(void *devPtr);

cudaError_t cudaMemcpy(void *dst, void *src, size_t count, enum cudaMemcpyKind kind);

#include <iostream>
#include <cuda_runtime.h>

__device__ int deviceVar;

__global__ void addOneToDeviceVar() {
    deviceVar += 1;
}

int main() {
    int h_var = 100;
    int h_var_from_device = 0;

    // 将主机变量拷贝到设备符号变量
    cudaMemcpyToSymbol(deviceVar, &h_var, sizeof(int));

    // 调用核函数：1 个线程块，64 个线程
    addOneToDeviceVar<<<1, 64>>>();

    cudaDeviceSynchronize();

    // 将设备符号变量拷贝回主机
    cudaMemcpyFromSymbol(&h_var_from_device, deviceVar, sizeof(int));

    std::cout << "Value copied to device: " << h_var << std::endl;
    std::cout << "Value after adding one on device: "
              << h_var_from_device << std::endl;

    cudaDeviceReset();
    return 0;
}

#include <iostream>
#include <cuda_runtime.h>

__device__ int deviceVar;

__global__ void addOneToDeviceVar() {
    // 使用原子操作消除数据竞争
    atomicAdd(&deviceVar, 1);
}

int main() {
    int h_var = 100;
    int h_var_from_device = 0;

    cudaMemcpyToSymbol(deviceVar, &h_var, sizeof(int));

    // 64 个线程，每个线程为设备变量 +1
    addOneToDeviceVar<<<1, 64>>>();

    cudaDeviceSynchronize();

    cudaMemcpyFromSymbol(&h_var_from_device, deviceVar, sizeof(int));

    std::cout << "Value copied to device: " << h_var << std::endl;
    std::cout << "Value after adding one on device: "
              << h_var_from_device << std::endl;

    cudaDeviceReset();
    return 0;
}

Value copied to device: 100
Value after adding one on device: 164

#include <cuda_runtime.h>
#include <iostream>
#include <iomanip>

__global__ void copyData(const int *input, int *output) {
    // 静态分配共享内存数组（大小为 64，编译时分配）
    __shared__ int sharedData[64];

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

    // 从全局内存读取数据到共享内存
    sharedData[threadIdx.x] = input[tid] * 2;

    // 线程块内同步：确保所有线程完成写入
    __syncthreads();

    // 从共享内存写回全局内存
    output[tid] = sharedData[threadIdx.x];
}

int main() {
    int numElements = 64;
    size_t size = numElements * sizeof(int);

    // 在主机端分配输入和输出数组
    int *h_input = new int[numElements];
    int *h_output = new int[numElements];

    // 给 64 个数赋初值
    for (int i = 0; i < numElements; ++i) {
        h_input[i] = i;
    }

    // 在设备端分配输入和输出数组
    int *d_input, *d_output;
    cudaMalloc((void **)&d_input, size);
    cudaMalloc((void **)&d_output, size);

    // 从主机拷贝数据到设备
    cudaMemcpy(d_input, h_input, size, cudaMemcpyHostToDevice);

    // 调用 2 个 blocks，每个 block 有 32 个线程
    copyData<<<2, 32>>>(d_input, d_output);

    // 把显存数据从 GPU 拷贝到 CPU
    cudaMemcpy(h_output, d_output, size, cudaMemcpyDeviceToHost);

    for (int i = 0; i < numElements; ++i) {
        std::cout << std::setw(15) << h_input[i]
                  << std::setw(15) << h_output[i] << std::endl;
    }

    // 收尾：删除数组（主机端和设备端）
    delete[] h_input;
    delete[] h_output;
    cudaFree(d_input);
    cudaFree(d_output);

    return 0;
}

#include <iostream>
#include <cuda_runtime.h>

// GPU 全局内存里的设备变量
__device__ float d_sum = 0.0;

__global__ void dynamic_shared_memory_kernel(unsigned int ndim) {
    // 动态分配共享内存（大小由调用核函数时确定）
    // 使用 __shared__ 关键字，"动态"由 "extern" 关键字确定
    extern __shared__ float shared_array[];

    int tid = threadIdx.x;

    // 共享内存赋值
    shared_array[tid] = tid;

    // 线程块内的线程同步
    __syncthreads();

    // 原子操作：对共享数组进行求和，防止数据竞争
    atomicAdd(&d_sum, shared_array[tid]);
}

int main() {
    unsigned int ndim = 100;

    dim3 grid(1);
    dim3 block(ndim);

    float c_sum = 0;

    // 调用核函数：1 个 block 和 100 个线程
    // 注意第三个参数是共享内存大小（字节）
    dynamic_shared_memory_kernel<<<grid, block, ndim * sizeof(float)>>>(ndim);

    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        std::cout << "Kernel Error: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }

    // 从设备端 d_sum 拷贝到 CPU 端 c_sum 后输出
    cudaMemcpyFromSymbol(&c_sum, d_sum, sizeof(float));

    err = cudaGetLastError();
    if (err != cudaSuccess) {
        std::cout << "memcpy Error: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }

    cudaDeviceSynchronize();

    std::cout << "sum=" << c_sum << std::endl;

    return 0;
}

nvcc 09_smeam.cu

sum=4950

// 添加 __device__ 限定符的 GPU 函数
__device__ float add(float a, float b) {
    return a + b;
}

// 添加 __global__ 限定符的 GPU 函数
__global__ void vectorAdd(
    const float *A,
    const float *B,
    float *C,
    int numElements) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < numElements) {
        C[i] = add(A[i], B[i]);  // 调用设备函数
    }
}

// global_factor 存放在全局内存
__device__ double global_factor = 2.0;

// __constant__ 表明是一个常量内存里定义的变量
__constant__ double constant_factor = 2.0;

// GPU 的核函数：做矩阵加法
__global__ void add_matrices(
    double *a,
    double *b,
    double *c,
    const int n) {
    const int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index < n) {
        for (int i = 0; i < 100000; ++i) {
            // global_factor：每次循环，每个线程都要发起一次全局内存读取
            c[index] += (a[index] + b[index]) * global_factor;

            // constant_factor：同一线程束内所有线程读取同一个常量地址时，
            // 硬件只需从常量缓存读一次，然后广播给所有线程，延迟极低
            c[index] += (a[index] + b[index]) * constant_factor;
        }
    }
}

__texture__

// 首先，声明和分配 3 个托管数组，A, B 用于输入，数组 gpuRef 用于输出
float *A, *B, *gpuRef;

// cudaMallocManaged() 支持 Unified Memory 的 GPU 上分配内存
cudaMallocManaged((void **)&A, nBytes);
cudaMallocManaged((void **)&B, nBytes);
cudaMallocManaged((void **)&gpuRef, nBytes);

// 然后，使用指向托管内存的指针来初始化主机上的两个输入矩阵
initiaData(A, nxy);
initiaData(B, nxy);

// 最后，通过指向托管内存的指针调用矩阵加法核函数
sumMatrixGPU<<<grid, block>>>(A, B, gpuRef, nx, ny);
cudaDeviceSynchronize();

// GPU 核函数修改统一内存数据
kernel<<<grid, block>>>(managed_ptr);

// 内存栅栏：确保 GPU 完成所有内存操作
cudaDeviceSynchronize();

// 现在 CPU 可以安全读取最新数据
printf("%f\n", managed_ptr[0]);