将函数指针及其参数作为推力::元组传递给全局函数

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 构造一个推力::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 在两个示例中。

Answer 1

将函数指针及其参数作为推力::元组传递给全局函数

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

$ 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
$

在功能方面（能够调度具有不同参数包的多个子内核）我看不出任何能力差异，再次假设您的参数很好地打包在一个元组中。

Answer 2

一个快速而肮脏的解决方案是强制转换函数指针：

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

template<typename... Args>
__global__ void someFunction(void (*fp)(), thrust::tuple<Args...> params) {
    void (*kfp)(Args...) = (void (*)(Args...)) fp;
    kfp<<<1,1>>>(thrust::get<0>(params));
}

__global__ void otherFunction(int n) {
    printf("n = %d\n", n);
}
__constant__ void (*kfp)(int) = &otherFunction;

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

我愿意接受更好的解决方案！