Passing einen Funktionszeiger und seine Parameter als Schub :: tuple zu einer globalen Funktion

Question

ich folgendes zu tun:

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

Was habe ich versucht:

1. Entfernen eines der beiden Parameter führt zu einer funktionierenden Lösung natürlich.
2. Es funktioniert, wenn ich someFunction eine statische Funktion in einem struct mit Vorlage Parameter machen. Aber im ursprünglichen Code someFunction ist ein CUDA-Kernel, also kann ich das nicht tun. Irgendwelche weiteren Ideen?
3. Es funktioniert, wenn ich schub :: tuple zu std :: tuple ändere. Gibt es eine Möglichkeit, ein stuck :: tuple aus einem std :: tuple zu konstruieren?

---

EDIT:

es deutlicher zu machen: someFunction und otherFunction sind __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; 
} 

ich einen Compiler-Fehler: template argument deduction/substitution failed in beiden Beispielen.

Answer 1

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

So etwas sollte bearbeitbar sein:

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

Eine ähnliche Methode auch praktikabel sein sollte, wenn Sie df beabsichtigen, eine __global__ Funktion zu sein, die Sie gerade richtig für die dynamische Parallelität Fall berücksichtigen müssen . Ebenso sollte nur eine geringfügige Abweichung von oben es Ihnen erlauben, das Tupel direkt an die Kindfunktion zu übergeben (d. H. df, ob Gerätefunktion oder Kernel). Es ist mir nicht klar, warum Sie variable Template-Argumente brauchen, wenn Ihre Parameter in einem Schub-Tupel gut verpackt sind.

EDIT: Wenn Sie Ihr Tuple an den Kind-Kernel übergeben können (Ich sehe nicht, warum Sie nicht in der Lage sein würden, da das Tupel und der Kind-Kernel gemäß dem aktualisierten Beispiel das gleiche variadic Parameter-Paket teilen) , dann können Sie immer noch variadische Vorlagen mit diesem Ansatz vermeiden können:

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

in Bezug auf Funktionalität (in der Lage, mit unterschiedlicher Parameter-Pack mehr Kind-Kernel zu versenden) ich sehe keinen Unterschied in der Fähigkeit, wieder Angenommen, Ihre Parameter sind gut in einem Tupel verpackt.

Answer 2

A quick and dirty Lösung ist es, die Funktionszeiger zu werfen:

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

Ich bin offen für schönere Lösungen!