Передача указателя функции и ее параметры как тяги :: кортеж глобальной функции

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. Он работает, когда я делаю someFunction статической функцией в struct с параметром шаблона. Но в исходном коде someFunction есть ядро ​​CUDA, поэтому я не могу этого сделать. Какие-нибудь дальнейшие идеи?
3. Он работает, когда я изменяю тягу :: кортеж в std :: кортеж. Есть ли способ построить трюк :: кортеж из std :: tuple?

---

EDIT:

Чтобы было понятнее: 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

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, будь то функция устройства или ядро). Мне непонятно, почему вам нужны переменные аргументы шаблона, если ваши параметры хорошо упакованы в кортеж.

EDIT: Если вы можете передать свой кортеж дочернему ядру (я не понимаю, почему вы не сможете этого сделать, поскольку в соответствии с вашим обновленным примером кортеж и дочернее ядро ​​используют один и тот же пакет параметров) , то вы все равно можете быть в состоянии избежать VARIADIC шаблонов с использованием этого подхода:

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

Я открыт для более хороших решений!