にゃははー

はへらー

CUDA2.3のバグ?

Constant Memoryを使おうとして下のようなソースを書いてコンパイル通したら、通ってしまって動くんだけど、正しい結果が帰ってこなかった。

...(ry

__constant__ float d_data[ 10 ];

int main( void )
{
    float h_data[ 10 ];
...(ry

    cudaMemcpyToSymbol( d_data, h_data, 10 * sizeof( float ), cudaMemcpyHostToDevice );

...(ry
}

でろくに調べないでやろうとしていたからcudaMemcpyToSymbol()が受け取る引数が5つあって、本当は

cudaMemcpyToSymbol( d_data, h_data, 10 * sizeof( float ), 0, cudaMemcpyHostToDevice );

ってなってるべきだった。(こっちは正しい結果が帰ってきた

で本題に入ると、cudaMemcpyToSymbol()の引数が足りなくてもコンパイルが通ってしまう件。これ、結構危険ではないだろうか。

cudaMemcpyToSymbol()の第4引数を調べるとoffsetって書いてある。つまりoffset Byteずらしてゴニョるということらしい。
nvccは引数の数を計算していないのか、たまたま発生したものなのかは細かい実験をしてないから分からないけど、他のAPIについても引数の過不足は注意したほうが、いいかもしれない。

ちなみに毎回cudaMemcpyToSymbol()はcudaSuccessを返してくれてました。Device側メモリが破壊されるだけでなくスタックも破壊されるかも・・・

追記:
いろいろ調べたらバグではありませんでした。

include/cuda_runtime.hに以下のような・・・

static __inline__ __host__ cudaError_t cudaMemcpyToSymbol(
        char                *symbol,
  const void                *src,
        size_t               count,
        size_t               offset = 0,
        enum cudaMemcpyKind  kind   = cudaMemcpyHostToDevice
)
{
  return cudaMemcpyToSymbol((const char*)symbol, src, count, offset, kind );
}

template <class T>
__inline__ __host__ cudaError_t cudaMemcpyToSymbol(
        char                &symbol,
  const char                *src,
        size_t               count,
        size_t               offset = 0,
        enum cudaMemcpyKind  kind   = cudaMemcpyHostToDevice
)
{
  return cudaMemcpyToSymbol((const char*)&symbol, src, count, offset, kind );
}

!?

こんな定義、マニュアルにかいてねぇぞwwwwwwww

さらに見ていくとcudaMemcpyFromSymbol()も第4,5引数にそれぞれ0,cudaMemcpyDeviceToHostのデフォルト引数が与えられたり、テンプレートが作られたりしていた。

便利そうな定義を書くのはいいんだけど、マニュアルに書いてなかったから上の要なソースが通ってしまって非常に焦った。

デフォルト引数は折角の型チェックが無駄になるだけじゃなくて、マニュアルに書いてないとこういうことが起こって危険だと身をもって知った。

とりあえずNVIDIAさん。マニュアル更新しよう。