c++ - 将函数指针及其参数作为 thrust::tuple 传递给全局函数

Question

我想做以下事情:

#include <thrust/tuple.h>
#include <tuple>

template<typename... Args>
void someFunction(void (*fp)(Args...), thrust::tuple<Args...> params) {
}

void otherFunction(int n) {
}

int main(int argc, char **argv) {
    //// template argument deduction/substitution failed ////
    someFunction<int>(&otherFunction, thrust::make_tuple(1));
    return 0;
}

我尝试过的:

1. 删除两个参数中的一个当然会产生一个可行的解决方案。
2. 当我使用模板参数在 struct 中将 someFunction 设为静态函数时，它会起作用。但是在原始代码中 someFunction 是一个 CUDA 内核，所以我不能那样做。还有其他想法吗？
3. 当我将 thrust::tuple 更改为 std::tuple 时它起作用了。有没有办法从 std::tuple 构建 thrust::tuple？

---

编辑:

更清楚一点:someFunction 和 otherFunction 是 __global__!

#include <thrust/tuple.h>
#include <tuple>

template<typename... Args>
__global__ void someFunction(void (*fp)(Args...), thrust::tuple<Args...> params) {
}

__global__ void otherFunction(int n) {
}
__constant__ void (*kfp)(int) = &otherFunction;

int testPassMain(int argc, char **argv) {
    void (*h_kfp)(int);
    cudaMemcpyFromSymbol(&h_kfp, kfp, sizeof(void *), 0, cudaMemcpyDeviceToHost);
    someFunction<int><<<1,1>>>(h_kfp, thrust::make_tuple(1));
    return 0;
}

我收到编译器错误:template argument deduction/substitution failed in both examples.

Answer 1

Passing a function pointer and its parameters as a thrust::tuple to a global function

这样的事情应该是可行的:

$ cat t1161.cu
#include <thrust/tuple.h>
#include <stdio.h>

template <typename T, typename T1>
__global__ void kernel(void (*fp)(T1), T params){ // "someFunction"

  fp(thrust::get<0>(params));
  fp(thrust::get<1>(params));
}

__device__ void df(int n){                        // "otherFunction"

  printf("parameter = %d\n", n);
}

__device__ void (*ddf)(int) = df;

int main(){

  void (*hdf)(int);
  thrust::tuple<int, int> my_tuple = thrust::make_tuple(1,2);
  cudaMemcpyFromSymbol(&hdf, ddf, sizeof(void *));
  kernel<<<1,1>>>(hdf, my_tuple);
  cudaDeviceSynchronize();
}


$ nvcc -o t1161 t1161.cu
$ cuda-memcheck ./t1161
========= CUDA-MEMCHECK
parameter = 1
parameter = 2
========= ERROR SUMMARY: 0 errors
$

如果您希望 df 成为一个 __global__ 函数，类似的方法也应该可行，您只需要正确考虑动态并行情况。同样，只要对上面稍加改动就可以将元组直接传递给子函数(即 df，无论是设备函数还是内核)。如果您的参数很好地打包在推力元组中，我不清楚为什么需要可变参数模板参数。

编辑:如果您可以将您的元组传递给子内核(我不明白为什么您不能这样做，因为根据您更新的示例，元组和子内核共享相同的可变参数包)，那么您仍然可以使用这种方法避免可变参数模板:

$ cat t1162.cu
#include <thrust/tuple.h>
#include <stdio.h>

template<typename T>
__global__ void someFunction(void (*fp)(T), T params) {
  fp<<<1,1>>>(params);
  cudaDeviceSynchronize();
}

__global__ void otherFunction(thrust::tuple<int> t) {
  printf("param 0 = %d\n", thrust::get<0>(t));
}

__global__ void otherFunction2(thrust::tuple<float, float> t) {
  printf("param 1 = %f\n", thrust::get<1>(t));
}
__device__ void (*kfp)(thrust::tuple<int>) = &otherFunction;
__device__ void (*kfp2)(thrust::tuple<float, float>) = &otherFunction2;

int main(int argc, char **argv) {
    void (*h_kfp)(thrust::tuple<int>);
    void (*h_kfp2)(thrust::tuple<float, float>);
    cudaMemcpyFromSymbol(&h_kfp, kfp, sizeof(void *), 0, cudaMemcpyDeviceToHost);
    someFunction<<<1,1>>>(h_kfp, thrust::make_tuple(1));
    cudaDeviceSynchronize();
    cudaMemcpyFromSymbol(&h_kfp2, kfp2, sizeof(void *), 0, cudaMemcpyDeviceToHost);
    someFunction<<<1,1>>>(h_kfp2, thrust::make_tuple(0.5f, 1.5f));
    cudaDeviceSynchronize();
    return 0;
}
$ nvcc -arch=sm_35 -rdc=true -o t1162 t1162.cu -lcudadevrt
$ CUDA_VISIBLE_DEVICES="1" cuda-memcheck ./t1162
========= CUDA-MEMCHECK
param 0 = 1
param 1 = 1.500000
========= ERROR SUMMARY: 0 errors
$

就功能而言(能够分派(dispatch)具有不同参数包的多个子内核)我没有看到任何功能差异，再次假设您的参数很好地打包在一个元组中。