2017-07-04 67 views
1

クラスをカーネルに渡すC++ CUDAプログラムを作成したいと思います。クラスは、呼び出し演算子()を介してカーネル上の関数を評価するだけです。クラスの中で関数をハードワイヤリングすると、すべてがうまくいくように機能します。しかし、私はクラスにいくつかの柔軟性が必要なので、クラスを異なる関数でインスタンス化できるようにしたいと思います。ポインタ関数を渡すといいでしょう。ポインタ関数の実装を動作させることができません。私は、2つのクラス、機能(genericFunction)CUDA:クラス関数をクラスのメンバとする関数にポインタを渡す

//Functions.hh 
#include <iostream> 
#include <stdio.h> 

class fixedFunction{ 
public: 
__host__ fixedFunction() {} 
__host__ __device__ double operator()(double x) { 
    return x*x; 
} 
}; 

double f1(double x){ 
    return x*x; 
} 

typedef double (*pf) (double var); 

class genericFunction{ 
public: 
    __host__ genericFunction(double (*infunc)(double)) : func(infunc){} 
    __host__ __device__ double operator()(double x) { 
    return func(x); 
    } 
private: 
    pf func; 
}; 

__global__ void kernel1(fixedFunction* g1){ 
    unsigned int tid = blockIdx.x *blockDim.x + threadIdx.x; 
    printf("Func val is: %f\n", (*g1)(tid)); 
} 

__global__ void kernel2(genericFunction* g1){ 
    unsigned int tid = blockIdx.x *blockDim.x + threadIdx.x; 
    printf("Func val is: %f\n", (*g1)(tid)); 
} 

両方のクラスをインスタンス化し、ホスト作品にそれらを実行しているへのポインタを取ります(fixedFunction)定義された機能を持つものと別の定義の下に。私は、そのクラスは、ポインタ機能を呼び出すことkernel2は、私は、ポインタ機能で問題が任意のサイズにすることができます見ることができ、それがデバイスで会計処理されていない関連するカーネルに

#include "Functions.hh" 

int main(){ 

    fixedFunction h_g1; 
    fixedFunction* d_g1; 
    cudaMallocManaged(&d_g1, sizeof(h_g1)); 

    //Host call 
    std::cout << h_g1(2.0) << "\n"; 

    //device call 
    kernel1<<<1,32>>>(d_g1); 
    cudaDeviceSynchronize(); 

    genericFunction h_g2(f1); 
    genericFunction* d_g2; 
    cudaMallocManaged(&d_g2, sizeof(h_g2)); 

    //Host call 
    std::cout << h_g2(3.0) << "\n"; 

    //device call 
    kernel2<<<1,32>>>(d_g2); 
    cudaDeviceSynchronize(); 

を失敗を参照してください渡します。それで、ポインタ関数をクラスに渡してデバイス上で実行する方法はありますか?

おかげ

+1

f1は、デバイスの機能ではありません。関数ポインタの設定が正しいかどうかにかかわらず、おそらくsカーネルを使うことはできません – talonmies

答えて

3

これは約私はあなたが意図しに見えたとして、それはおよそ機能するために得るためにあなたのコードに作ることができ、変更の「最小」の数でした。また、CUDAの関数ポインタについては、this answerにいくつかのリンクがあります。

  1. f1__host__ __device__とを飾る。コンパイラにデバイスコール可能なルーチンを生成させるためには、これが必要です。それ以外の場合は、ホストコードのみが生成されます。

  2. 上記で作成したf1のデバイス呼び出し可能バージョンのデバイスエントリアドレスを取得する必要があります。これを行うにはいくつかの方法があります。私は別の__device__変数(f1_d)でそれを "静的に"キャプチャし、cudaMemcpyFromSymbolを使用してホストコードにプルします。

  3. genericFunctionクラスは、__host__と、目的の機能のために別の__device__エントリポイント(機能ポインタ)の両方を保持できるように変更されています。また、クラスのホストまたはデバイスのバージョン(__CUDA_ARCH__マクロ)をコンパイルするかどうかに基づいて、クラスを変更して、両方のエントリポイントを受け入れて割り当てるようにクラスコンストラクタを変更します。

  4. 最後に、デバイス上のd_g2オブジェクトも初期化する必要があります。 d_g1オブジェクトの場合、そのオブジェクトのクラスデータメンバーはありません。したがって、d_g1が指し示す「空の」オブジェクトを作成することができます。そのオブジェクトのクラスメンバー関数のエントリポイントデバイスコードではすでに認識されています。しかし、d_g2の場合、関数のそれぞれのホストとデバイスのバージョン(エントリポイント)へのポインタであるクラスデータメンバを介して間接的に関数にアクセスしています。したがって、h_g2オブジェクトをホストコードで初期化し、d_g2オブジェクトのストレージをデバイスコードに設定した後、d_g2の内容をとし、h_g2の内容をに初期化する必要があります。これらの変更により

、あなたのコードは、私のテストによると書かれたように動作します

$ cat t353.cu 
#include <iostream> 
#include <stdio.h> 

class fixedFunction{ 
public: 
__host__ fixedFunction() {} 
__host__ __device__ double operator()(double x) { 
    return x*x; 
} 
}; 

__host__ __device__ double f1(double x){ 
    return x*x; 
} 

typedef double (*pf) (double var); 

__device__ pf f1_d = f1; 

class genericFunction{ 
public: 
    __host__ genericFunction(double (*h_infunc)(double), double (*d_infunc)(double)) : h_func(h_infunc),d_func(d_infunc){} 
    __host__ __device__ double operator()(double x) { 
#ifdef __CUDA_ARCH__ 
    return d_func(x); 
#else 
    return h_func(x); 
#endif 
    } 
private: 
    pf h_func; 
    pf d_func; 
}; 

__global__ void kernel1(fixedFunction* g1){ 
    unsigned int tid = blockIdx.x *blockDim.x + threadIdx.x; 
    printf("Func val is: %f\n", (*g1)(tid)); 
} 

__global__ void kernel2(genericFunction* g1){ 
    unsigned int tid = blockIdx.x *blockDim.x + threadIdx.x; 
    printf("Func val is: %f\n", (*g1)(tid)); 
} 

int main(){ 

    fixedFunction h_g1; 
    fixedFunction* d_g1; 
    cudaMallocManaged(&d_g1, sizeof(h_g1)); 

    //Host call 
    std::cout << h_g1(2.0) << "\n"; 

    //device call 
    kernel1<<<1,32>>>(d_g1); 
    cudaDeviceSynchronize(); 
    pf d_f1; 
    cudaMemcpyFromSymbol(&d_f1, f1_d, sizeof(void*)); 
    genericFunction h_g2(f1, d_f1); 
    genericFunction* d_g2; 
    cudaMallocManaged(&d_g2, sizeof(h_g2)); 
    cudaMemcpy(d_g2, &h_g2, sizeof(h_g2), cudaMemcpyDefault); 
    //Host call 
    std::cout << h_g2(3.0) << "\n"; 

    //device call 
    kernel2<<<1,32>>>(d_g2); 
    cudaDeviceSynchronize(); 
} 
$ nvcc -arch=sm_61 -o t353 t353.cu 
$ cuda-memcheck ./t353 
========= CUDA-MEMCHECK 
4 
Func val is: 0.000000 
Func val is: 1.000000 
Func val is: 4.000000 
Func val is: 9.000000 
Func val is: 16.000000 
Func val is: 25.000000 
Func val is: 36.000000 
Func val is: 49.000000 
Func val is: 64.000000 
Func val is: 81.000000 
Func val is: 100.000000 
Func val is: 121.000000 
Func val is: 144.000000 
Func val is: 169.000000 
Func val is: 196.000000 
Func val is: 225.000000 
Func val is: 256.000000 
Func val is: 289.000000 
Func val is: 324.000000 
Func val is: 361.000000 
Func val is: 400.000000 
Func val is: 441.000000 
Func val is: 484.000000 
Func val is: 529.000000 
Func val is: 576.000000 
Func val is: 625.000000 
Func val is: 676.000000 
Func val is: 729.000000 
Func val is: 784.000000 
Func val is: 841.000000 
Func val is: 900.000000 
Func val is: 961.000000 
9 
Func val is: 0.000000 
Func val is: 1.000000 
Func val is: 4.000000 
Func val is: 9.000000 
Func val is: 16.000000 
Func val is: 25.000000 
Func val is: 36.000000 
Func val is: 49.000000 
Func val is: 64.000000 
Func val is: 81.000000 
Func val is: 100.000000 
Func val is: 121.000000 
Func val is: 144.000000 
Func val is: 169.000000 
Func val is: 196.000000 
Func val is: 225.000000 
Func val is: 256.000000 
Func val is: 289.000000 
Func val is: 324.000000 
Func val is: 361.000000 
Func val is: 400.000000 
Func val is: 441.000000 
Func val is: 484.000000 
Func val is: 529.000000 
Func val is: 576.000000 
Func val is: 625.000000 
Func val is: 676.000000 
Func val is: 729.000000 
Func val is: 784.000000 
Func val is: 841.000000 
Func val is: 900.000000 
Func val is: 961.000000 
========= ERROR SUMMARY: 0 errors 
$ 
+0

男、この多くのトラブルに払われることを願っています:-) – einpoklum

+0

Thanks Robert。あなたの助けは非常に高く評価されています! – photon

関連する問題