関数ポインタと推力としてそのパラメータを渡す::タプルをグローバル関数

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. 私はthrust :: tupleをstd :: tupleに変更すると動作します。 std :: tupleからthrust :: 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 
$ 

同様の方法論はまた、あなたが__global__機能するdfを予定であれば、あなただけのダイナミック並列処理の場合のために適切に考慮する必要があります実行可能でなければなりません。同様に、上記のわずかなバリエーションでは、タプルを子関数に直接渡すことができます（つまり、デバイス関数かカーネルかにかかわらず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; 
} 

私はよりよい解決策に開いています！