// 假设有一个结构体表示3D点
struct Point3D {
    float x;
    float y;
    float z;
};

// 使用AoS方式存储多个点
Point3D pointsAoS[N]; // N个点

// 访问某个点的x坐标。在GPU上，如果我们要对所有的x坐标进行某种操作，我们需要遍历整个数组，并且每次都要从一个不同的结构体中读取x字段。这可能导致内存访问的不连续，从而降低性能。
float xCoord = pointsAoS[i].x;


// 使用SoA方式存储多个点的x、y、z坐标
struct Points{
  float xCoords[N];
  float yCoords[N];
  float zCoords[N]; // N个点
}

// 访问某个点的x坐标。在GPU上，如果我们要对所有的x坐标进行某种操作，我们可以直接对整个xCoords数组进行操作，而不需要跳转到不同的内存位置。这种连续的内存访问模式可以显著提高性能。
float xCoord = Points.xCoords[i];

#include <cuda_runtime.h>  
  
// 定义块和线程的大小  
const int BLOCK_SIZE = 256;  
const int THREADS_PER_WARP = 32;  
  
// 假设我们有一个数组，每个元素都是一个float  
__global__ void coalescedAccessKernel(float* d_array, int n) {  
    int index = threadIdx.x + blockIdx.x * blockDim.x;  
    if (index < n) {  
        // 确保warp内的线程访问连续的内存位置  
        // 这里假设n是warp大小的整数倍  
        float value = d_array[index];  
        // ... 对value进行一些操作 ...  
        d_array[index] = value; // 假设我们修改了value的值  
    }  
}  
  
int main() {  
    // 分配设备上的内存  
    int n = BLOCK_SIZE * (BLOCK_SIZE / THREADS_PER_WARP); // 确保n是warp大小的整数倍  
    float* h_array = new float[n];  
    float* d_array;  
    cudaMalloc((void**)&d_array, n * sizeof(float));  
  
    // 初始化h_array...  
  
    // 将数据从主机复制到设备  
    cudaMemcpy(d_array, h_array, n * sizeof(float), cudaMemcpyHostToDevice);  
  
    // 调用内核函数  
    coalescedAccessKernel<<<n / BLOCK_SIZE, BLOCK_SIZE>>>(d_array, n);  
  
    // 等待内核完成  
    cudaDeviceSynchronize();  
  
    // 将数据从设备复制回主机  
    cudaMemcpy(h_array, d_array, n * sizeof(float), cudaMemcpyDeviceToHost);  
  
    // 释放内存...  
  
    return 0;  
}

__constant__ float constantData[64]; // 声明常量内存

__global__ void useConstantMemory() {
    int index = threadIdx.x;
    float value = constantData[index]; // 访问常量内存
    // ... 其他操作 ...
}

__shared__ float sharedData[32]; // 声明共享内存

__global__ void useSharedMemory(float* globalData, int n) {
    int tid = threadIdx.x;
    int bid = blockIdx.x;
    int index = tid + bid * blockDim.x;

    if (index < n) {
        // 从全局内存加载数据到共享内存
        sharedData[tid] = globalData[index];
        __syncthreads(); // 等待所有线程完成加载

        // 使用共享内存中的数据
        // ...

        // 将结果写回全局内存（如果需要）
        // ...
    }
}

#include <cuda_runtime.h>

int main() {
    // 分配页锁定内存
    void* pinnedHostMemory;
    cudaHostAlloc(&pinnedHostMemory, sizeof(float) * N, cudaHostAllocDefault);

    // ... 填充pinnedHostMemory ...

    // 分配设备内存
    float* deviceMemory;
    cudaMalloc(&deviceMemory, sizeof(float) * N);
 
    // 异步传输数据到设备
    cudaMemcpyAsync(deviceMemory, pinnedHostMemory, sizeof(float) * N, cudaMemcpyHostToDevice);

    // ... 在这里启动GPU内核 ...

    // 等待数据传输完成（如果需要）
    cudaDeviceSynchronize();

    // 释放内存
    cudaFree(deviceMemory);
    cudaFreeHost(pinnedHostMemory);

    return 0;
}

#include <cuda_runtime.h>

// 假设我们有一个简单的内核函数
__global__ void myKernel(int *data, int size) {
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    if (index < size) {
        // 对数组中的每个元素进行操作
        data[index] *= 2;
    }
}

int main() {
    // 主机上的数据
    int hostData[1024];
    // 初始化数据...

    // 分配设备上的内存
    int *devData;
    cudaMalloc(&devData, sizeof(int) * 1024);

    // 创建两个CUDA流
    cudaStream_t stream1, stream2;
    cudaStreamCreate(&stream1);
    cudaStreamCreate(&stream2);

    // 使用第一个流将数据从主机复制到设备
    cudaMemcpyAsync(devData, hostData, sizeof(int) * 1024, cudaMemcpyHostToDevice, stream1);

    // 在第二个流上启动一个内核
    myKernel<<<1, 1024, 0, stream2>>>(devData, 1024); // 注意指定了流stream2

    // 假设我们需要等待第一个流的数据传输完成，然后再进行其他操作
    cudaStreamSynchronize(stream1);

    // 此时，由于第二个流上的内核已经启动，并且与第一个流是并发的，
    // 因此我们可以继续执行其他操作，而不需要等待内核完成。

    // ... 在这里可以执行其他与流无关的操作，或者启动更多内核到不同的流 ...

    // 最后，等待第二个流上的内核完成
    cudaStreamSynchronize(stream2);

    // 如果需要，可以将结果从设备复制回主机（这里省略了代码）

    // 清理
    cudaFree(devData);
    cudaStreamDestroy(stream1);
    cudaStreamDestroy(stream2);

    return 0;
}