CUDA 核函数参数传递示例：函数 foo 作为 foo2 的参数

本示例演示了如何将一个函数 foo 作为参数传递给 CUDA 核函数 foo2。函数 foo 用于将一个整数的值复制到另一个整数，并在核函数中并行执行，实现数组数据的快速复制。

代码示例：

__global__ void foo2(int* dst, int* src, int size) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < size) {
        foo(dst[tid], src[tid]);
    }
}

__device__ void foo(int& dst, int src) {
    dst = src;
}

int main() {
    int size = 100;
    int* src = new int[size];
    int* dst = new int[size];

    // 初始化src数组
    for (int i = 0; i < size; ++i) {
        src[i] = i;
    }

    // 将src数组拷贝到GPU的内存中
    int* d_src;
    cudaMalloc(&d_src, size * sizeof(int));
    cudaMemcpy(d_src, src, size * sizeof(int), cudaMemcpyHostToDevice);

    // 分配GPU内存，并调用核函数
    int* d_dst;
    cudaMalloc(&d_dst, size * sizeof(int));
    int threadsPerBlock = 256;
    int blocksPerGrid = (size + threadsPerBlock - 1) / threadsPerBlock;
    foo2<<<blocksPerGrid, threadsPerBlock>>>(d_dst, d_src, size);

    // 将结果拷贝回主机内存
    cudaMemcpy(dst, d_dst, size * sizeof(int), cudaMemcpyDeviceToHost);

    // 验证结果
    for (int i = 0; i < size; ++i) {
        assert(src[i] == dst[i]);
    }

    // 释放内存
    delete[] src;
    delete[] dst;
    cudaFree(d_src);
    cudaFree(d_dst);

    return 0;
}

代码解释：

__global__ void foo2(int* dst, int* src, int size): 定义了一个名为 foo2 的 CUDA 核函数，它接收三个参数：dst 和 src 分别指向目标数组和源数组，size 表示数组的长度。
__device__ void foo(int& dst, int src): 定义了一个名为 foo 的设备函数，它接收两个参数：dst 是目标整数的引用，src 是源整数。函数将 src 的值赋值给 dst。
int main() : 主函数，包含以下步骤：
- 初始化源数组 src。
- 将源数组 src 拷贝到 GPU 内存中，创建 d_src 指针指向该内存。
- 分配 GPU 内存，创建 d_dst 指针指向该内存。
- 调用核函数 foo2，将 d_dst，d_src 和 size 作为参数传递给核函数。
- 将结果从 GPU 内存拷贝回主机内存，使用 d_dst 指向的内存更新 dst 数组。
- 验证 dst 数组是否正确，确保每个元素的值与源数组 src 中的元素值相同。
- 释放内存。

注意：

__global__ 和 __device__ 分别表示全局函数和设备函数，这些函数的执行环境不同。
cudaMalloc 用于在 GPU 上分配内存。
cudaMemcpy 用于在主机内存和 GPU 内存之间复制数据。
<<<blocksPerGrid, threadsPerBlock>>> 用于指定核函数的执行方式，blocksPerGrid 表示网格中的块数，threadsPerBlock 表示每个块中的线程数。

示例代码展示了如何在 CUDA 核函数中传递一个函数作为参数，并使用该函数实现数据处理。这种方法可以提高代码的可重用性和灵活性，并允许在 GPU 上并行执行不同的数据处理逻辑。