今回は、CUDA2.0で for を使ったときに気づいた点です。

---

// メモリをコピーする関数
__device__ void copy_n(const unsigned char* first, unsigned char* dest, unsigned int size)
{
    const unsigned char* last = first + size;
    for (; first != last; ++first, ++dest) {
        *dest = *first;
    }
}

// 実験コード
__global__ void test(unsigned char* buf, unsigned int n)
{
    copy_n(buf, buf + 10, n - 1);
}

// ホストから test<<< grid, threads >>>(d_buf, 4); のように呼び出す。単一スレッド。

---

※ 長時間止まったり、ハングすることがあるかもしれません。

copy_n 関数は、stl にありそうな、メモリをコピーするだけの関数です(実際はstlには無いです)。単純な関数ですが、コンパイルして動かしてみると、しばらくGPUで動いた後、「the launch timed out and was terminated.」とエラーが出ます。

いろいろ試してみたところ、動いたり動かなかったりします。以下、いろいろ変えてみた内容と結果です。

- copy_n(buf, buf + 10, n - 1)  の n - 1 を n - 2 にしてもダメ
- copy_n(buf, buf + 10, n - 1)  の n - 1 を n + 0 にするとOK
- copy_n(buf, buf + 10, n - 1)  の n - 1 を n + 1 にするとOK
- copy_n(buf, buf + 10, n - 1)  の n - 1 を n + 2 にするとOK
- copy_n(buf, buf + 10, n - 1)  の n - 1 を n / 4 にするとOK
- copy_n(buf, buf + 10, n - 1)  の n - 1 を n / 2 にするとOK
- copy_n(buf, buf + 10, n - 1)  の n - 1 を n * 2 にするとOK
- copy_n(buf, buf + 10, n - 1)  の n - 1 を 3 にするとOK
- copy_n(buf, buf + 10, n - 1)  の n - 1 を 4 にするとOK
- copy_n(buf, buf + 10, n - 1)  の n - 1 を 5 にするとOK
- copy_n(buf, buf + 10, n - 1)  の n - 1 を n / 2 - 1 にしてもダメ
- copy_n(buf, buf + 10, n - 1)  の n - 1 を n / 2 + 1 にするとOK
- copy_n(buf, buf + 10, n - 1)  の行の上に、n = 3; という行を入れると、OK
- copy_n(buf, buf + 10, n - 1)  の行の上に、n = 4; という行を入れると、OK
- copy_n(buf, buf + 10, n - 1)  の行の上に、n = 4; という行を入れると、OK
- copy_n(buf, buf + 10, n - 1)  の行の上に、--n; という行を入れてもダメ
- copy_n(buf, buf + 10, n - 1)  の行の上に、++n; という行を入れると、OK
- 引数 n の値を変えてもダメ
- *dest = *first; をコメントアウトするとOK
- for ループを while を用いて書き換えてもダメ
- for ループを、for (unsigned int i = 0; i < size; ++i, ++first, ++dest) { と書き直したらOK

※ 正常にコピーされているかはチェックしてません。

不思議な条件です。はっきりしたことは分からないのですが、(特にforの終了条件に)ポインタ比較を行うと危険だということでしょうか。


2008-12-10 追記:
この記事へ、unsigned の引き算が怪しいとのコメントがありました。バイト数の異なる符号無し整数が絡んだ引き算が原因ぽいです(ポインタはたぶん64ビット整数)。とにかくコンパイラが間違ったコードを吐いているようです。

詳しくは サンマヤのプログラミングの部屋 をご覧ください。

ご指摘ありがとうございました。


開発環境: WinXP SP3 x64, VC++2005 x64, CUDA2.0 x64, GeFource 8600GT