2

私は非常に奇妙な状況を経験しています。私はこのテンプレート構造を持っています:

#ifdef __CUDACC__
#define __HOSTDEVICE __host__ __device__
#else
#define __HOSTDEVICE
#endif

template <typename T>
struct matrix
{
    T* ptr;
    int col_size, row_size;
    int stride;
    // some host & device methods
};

struct dummy1 {};
struct dummy2 : dummy1 {};

template <typename T>
struct a_functor : dummy2
{
    matriz<T> help_m;
    matrix<T> x, y;
    T *x_ptr, *y_ptr;
    int bsx, ind_thr;
    __HOSTDEVICE void operator()(T* __x, T* __y)
    {
        // functor code
    }
};

cpp ファイルと cu ファイルを分離するようにコードを構成したので、_functor オブジェクトが cpp ファイルで作成され、カーネル関数で使用されます。問題は、カーネル内で operator() を実行すると、コードを見ただけでは説明できないランダムな動作が見つかったことです。私の構造体がちょっと壊れていたようでした。したがって、a_functor オブジェクトで sizeof() を呼び出すと、次のことがわかりました。

  • CPU コード (カーネル外の .cpp および .cu): 64 バイト

  • GPU コード (カーネル内): 68 バイト

明らかに、全体を台無しにするある種のミスマッチがありました。さらに進んで、構造体パラメーター ポインターと構造体自体の間の距離を追跡し、生成されたメモリ レイアウトを検査しようとしましたが、次のような結果が得られました。

a_functor foo;
// CPU
(char*)(&foo.help_m)    - (char*)(&foo)   = 0
(char*)(&foo.x)         - (char*)(&foo)   = 16
(char*)(&foo.y)         - (char*)(&foo)   = 32
(char*)(&foo.x_ptr)     - (char*)(&foo)   = 48
(char*)(&foo.y_ptr)     - (char*)(&foo)   = 52
(char*)(&foo.bsx)       - (char*)(&foo)   = 56
(char*)(&foo.ind_thr)   - (char*)(&foo)   = 60

// GPU - inside a_functor::operator(), in-kernel
(char*)(&this->help_m)  - (char*)(this)   = 4
(char*)(&this->x)       - (char*)(this)   = 20
(char*)(&this->y)       - (char*)(this)   = 36
(char*)(&this->x_ptr)   - (char*)(this)   = 52
(char*)(&this->y_ptr)   - (char*)(this)   = 56
(char*)(&this->bsx)     - (char*)(this)   = 60
(char*)(&this->ind_thr) - (char*)(this)   = 64

なぜ nvcc が私の構造体のためにこのメモリ レイアウトを生成したのか、私には本当に理解できません (その 4 バイトは何をすべきか!?!)。アラインメントの問題である可能性があると考え、a_functor を明示的にアラインしようとしましたが、カーネルで値によって渡されるため、できません。

template <typename T, typename Str>
__global__ void mykernel(Str foo, T* src, T*dst);

コンパイルしようとすると、

エラー: win32 プラットフォームのグローバルルーチンに明示的なアラインメントが大きすぎるパラメーターを渡すことはできません

では、この奇妙な状況を解決するには (...これは nvcc のバグだと思います)、どうすればよいでしょうか? 私が考えることができる唯一のことは、前述のエラーを回避するために、アラインメントで遊んで、構造体をポインターでカーネルに渡すことです。しかし、私は本当に疑問に思っています.なぜそのメモリレイアウトが一致しないのですか?! 本当に意味不明…

詳細情報: Visual Studio 2008 を使用しており、Windows XP 32 ビット プラットフォームで MSVC を使用してコンパイルしています。最新の CUDA Toolkit 5.0.35 をインストールしました。私のカードは GeForce GTX 570 (計算能力 2.0) です。

4

1 に答える 1

3

コメントから、実際に実行しているコードと投稿したコードに違いがあるように見えるため、誰かが問題を再現できない限り、あいまいな回答以上のものを提供することは困難です. とはいえ、Windows では、構造体のレイアウトとサイズが CPU と GPU で異なる場合があります。これらはプログラミング ガイドに記載されています。

Windows では、CUDA コンパイラは、次の条件のいずれかを満たすクラス型 T の C++ オブジェクトに対して、ホストの Microsoft コンパイラとは異なるメモリ レイアウトを生成する場合があります。

  • T は仮想関数を持つか、仮想関数を持つ直接または間接の基本クラスから派生します。
  • T には、直接または間接の仮想基本クラスがあります。
  • T には、複数の直接または間接の空の基底クラスを持つ多重継承があります。

このようなオブジェクトのサイズも、ホスト コードとデバイス コードで異なる場合があります。タイプ T がホストまたはデバイス コードで排他的に使用されている限り、プログラムは正しく動作するはずです。ホスト コードとデバイス コードの間で型 T のオブジェクトを渡さないでください (たとえば、グローバル関数への引数として、または cudaMemcpy*() 呼び出しを介して)。

3 番目のケースは、空の基本クラスがある場合に適用される可能性があります。実際のコードに複数の継承がありますか?

于 2013-05-20T09:21:22.143 に答える