CUDA 核函数示例：使用函数指针传递自定义函数

本示例演示了如何在 CUDA 中定义和调用一个名为 mykernel 的核函数，该函数使用函数指针参数 funcFormat 来接收用户定义的函数 foo。foo 函数的功能是将 src 指向的值复制到 dst 指向的值，并演示了如何使用 mykernel 将 src 中的数据复制到 dst 中。

#include <stdio.h>
#include <cuda_runtime.h>

typedef void(*funcFormat)(int*, int*);

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

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

int main() {
    int size = 1024;
    int* dst, * src;
    cudaMalloc(&dst, size * sizeof(int));
    cudaMalloc(&src, size * sizeof(int));

    int* h_dst = (int*)malloc(size * sizeof(int));
    for (int i = 0; i < size; i++) {
        h_dst[i] = 0;
    }
    int* h_src = (int*)malloc(size * sizeof(int));
    for (int i = 0; i < size; i++) {
        h_src[i] = i;
    }

    cudaMemcpy(src, h_src, size * sizeof(int), cudaMemcpyHostToDevice);

    dim3 blockDim(256);
    dim3 gridDim((size + blockDim.x - 1) / blockDim.x);
    mykernel<<<gridDim, blockDim>>>(foo, dst, src, size);

    cudaMemcpy(h_dst, dst, size * sizeof(int), cudaMemcpyDeviceToHost);

    for (int i = 0; i < size; i++) {
        printf('%d ', h_dst[i]);
    }
    printf('
');

    free(h_dst);
    free(h_src);
    cudaFree(dst);
    cudaFree(src);

    return 0;
}

在上述代码中，我们首先定义了一个名为 foo 的函数，该函数的功能是将 src 指向的值赋值给 dst 指向的值。然后，我们在 mykernel 函数中调用 foo 函数，以将 src 中的值复制到 dst 中。最后，我们将 dst 中的值从设备内存复制到主机内存，并打印出来，以验证程序的正确性。

代码说明：

• funcFormat 是一个函数指针类型，它指向一个接受两个整型指针作为参数的函数。
• mykernel 是一个核函数，它接收一个 funcFormat 类型的函数指针作为参数。
• foo 是一个自定义函数，它接受两个整型指针作为参数，并将 src 指向的值复制到 dst 指向的值。
• 在 main 函数中，我们首先分配主机内存和设备内存。然后，我们初始化主机内存，并将数据从主机内存复制到设备内存。最后，我们调用 mykernel 核函数，并将 foo 函数指针作为参数传递给它。

总结：

本示例展示了如何使用函数指针在 CUDA 中定义和调用核函数，并使用自定义函数来执行特定的操作。这提供了一种灵活的方式，允许用户根据需要定义和使用不同的函数，而无需修改核函数本身。