3

共有メモリ内の配列の割り当てと割り当て解除のためのいくつかのルーチンを持つ CUDA アプリケーションを開発しています。

このアプリケーション (申し訳ありませんが、利用できません) には、メモリのチャンクを配列としてカプセル化するクラスがあります。このクラスにはcount、特定の値に一致する要素の数をカウントするメソッドがあります。

だから、（クラス全体の実際の部分である）のようなものを想像してください

template <class Type>
struct Array {
    // ...

    Type &operator[](int i) { return data_[i]; }
    Type operator[](int i) const { return data_[i]; }

    size_t count(const Type &val) const {
        size_t c = 0;
        for (size_t i = 0; i < len_; ++i)
            if (data_[i] == val)
                ++c;
        return c;
    }

    void print(const char *fmt, const char *sep, const char *end) const {
        for (size_t i = 0; i < len_ - 1; ++i) {
            printf(fmt, data_[i]);
            printf(sep);
        }
        printf(fmt, _data[len_ - 1]);
        printf(end);
    }
private:
    Type *data_;
    size_t len_;
};

私がアクセスしているメモリが正しく割り当てられていると仮定します (共有メモリは実行時に割り当てられ、次元をカーネルに渡します)、データを格納するのに十分な大きさであり、共有メモリ内の整列された (wrt ) 領域をdata_指します。私はこれを複数回確認しましたが、これらの仮定は有効である必要があります (ただし、さらに確認を求めてください)。Type

さて、コードをテストしているときに、非常に奇妙なことがわかりました。

  • を使用して明示的に値を代入し、 を使用operator[]して読み取る場合operator[] const、問題は発生しません。
  • を使用してデータを読み取る場合print、問題は発生しません。
  • を呼び出すとcount()、プログラムがクラッシュし、(x = sizeof(Type))Address ADDR is out of boundsが原因で cuda-memcheck によって報告されます。Invalid __global__ read of size xADDR は共有メモリ バッファ内にあるため、有効である必要があります。
  • 内部countを に置き換えるdata_[i]と(*this)[i]、プログラムは正常に実行され、クラッシュは発生しません。

今、私はこれが起こる可能性についてまったく知りませんし、舞台裏で何が起こっているかを確認するために何をチェックすればよいかわかりません.なぜ直接読むとクラッシュするのですか? なぜ使用operator[]しないのですか？そして、内部を (直接?) 読み取っprintてもクラッシュしないのはなぜですか?

この質問が難しいことは承知しており、コードに関するこの小さな情報を提供して申し訳ありません...しかし、詳細についてはお気軽にお尋ねください。できる限りお答えします。私が解決しようとしている日々であり、これは私が得ることができる限りであるため、どんなアイデアや提案も大歓迎です.

このコードをテストするために、2 つの異なる GPU を使用しています。1 つは機能 2.1 で、もう 1 つは 3.5 です (後者はこのクラッシュに関する詳細情報を提供しますが、最初のものは提供しません)。CUDA5.0

編集:このエラーが発生する最小限の例を見つけました。不思議なことに、sm_20 と sm_35 でコンパイルするとエラーが表示されますが、sm_30 では表示されません。私が使用している GPU の上限は 3.5 です。

/* Compile and run with:
  nvcc -g -G bug.cu -o bug -arch=sm_20 # bug!
  nvcc -g -G bug.cu -o bug -arch=sm_30 # no bug :|
  nvcc -g -G bug.cu -o bug -arch=sm_35 # bug!
  cuda-memcheck bug

Here's the output (skipping the initial rows) I get
Ctor for 0x3fffc10 w/o alloc, data 0x10000c8
Calling NON CONST []
Calling NON CONST []
Fill with [] ok
Fill with raw ok
Kernel launch failed with error:
        unspecified launch failure
========= Invalid __global__ write of size 8
=========     at 0x00000188 in /home/bio/are/AlgoCUDA/bug.cu:26:array<double>::fill(double const &)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x010000c8 is out of bounds
=========     Device Frame:/home/bio/are/AlgoCUDA/bug.cu:49:kernel_bug(unsigned long) (kernel_bug(unsigned long) : 0x8c0)
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/libcuda.so (cuLaunchKernel + 0x3dc) [0xc9edc]
=========     Host Frame:/opt/cuda-5.0/lib64/libcudart.so.5.0 [0x13324]
=========     Host Frame:/opt/cuda-5.0/lib64/libcudart.so.5.0 (cudaLaunch + 0x182) [0x3ac62]
=========     Host Frame:bug [0xbb8]
=========     Host Frame:bug [0xaa7]
=========     Host Frame:bug [0xac4]
=========     Host Frame:bug [0xa07]
=========     Host Frame:/lib/libc.so.6 (__libc_start_main + 0xfd) [0x1ec4d]
=========     Host Frame:bug [0x8c9]
=========
========= Program hit error 4 on CUDA API call to cudaDeviceSynchronize 
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame:/usr/lib/libcuda.so [0x26a180]
=========     Host Frame:/opt/cuda-5.0/lib64/libcudart.so.5.0 (cudaDeviceSynchronize + 0x1dd) [0x441fd]
=========     Host Frame:bug [0xa0c]
=========     Host Frame:/lib/libc.so.6 (__libc_start_main + 0xfd) [0x1ec4d]
=========     Host Frame:bug [0x8c9]
=========
========= ERROR SUMMARY: 2 errors


(cuda-gdb) set cuda memcheck on
(cuda-gdb) run
Starting program: /home/bio/are/AlgoCUDA/bug 
[Thread debugging using libthread_db enabled]
[New Thread 0x7ffff5c25700 (LWP 23793)]
[Context Create of context 0x625870 on Device 0]
[Launch of CUDA Kernel 0 (kernel_bug<<<(1,1,1),(1,1,1)>>>) on Device 0]
Memcheck detected an illegal access to address (@global)0x10000c8

Program received signal CUDA_EXCEPTION_1, Lane Illegal Address.
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 12, warp 0, lane 0]
0x0000000000881928 in array<double>::fill (this=0x3fffc10, v=0x3fffc08) at bug.cu:26
26                              data[i] = v;
*/

#include <stdio.h>

extern __shared__ char totalSharedMemory[];

template <class Type>
struct array {
    // Create an array using a specific buffer
    __device__ __host__ array(size_t len, Type *buffer):
        len(len),
        data(buffer) {
        printf("Ctor for %p w/o alloc, data %p\n", this, data);
    }
    __device__ __host__ Type operator[](int i) const {
        printf("Calling CONST []\n");
        return data[i];
    }
    __device__ __host__ Type &operator[](int i) {
        printf("Calling NON CONST []\n");
        return data[i];
    }
    __device__ __host__ void fill(const Type &v) {
        for (size_t i = 0; i < len; ++i) data[i] = v;
    }
    size_t len;
    Type *data;
};

__global__ void kernel_bug(size_t bytesPerBlock) {
    // This is a test writing to show that filling the memory
    // does not produce any error
    for (size_t i = 0; i < bytesPerBlock; ++i) {
        totalSharedMemory[i] = i % ('z' - 'a' + 1) + 'a';
        printf("[%p] %c\n", totalSharedMemory + i, totalSharedMemory[i]);
    }

    // 200 / 8 = 25 so should be aligned
    array<double> X(2, (double *)(totalSharedMemory + 200));
    X[0] = 2;
    X[1] = 4;
    printf("Fill with [] ok\n");
    X.data[0] = 1;
    X.data[1] = 0;
    printf("Fill with raw ok\n");
    X.fill(0); // Crash here
    printf("Fill with method ok\n");
}

int main(int argc, char **argv) {
    // Total memory required
    size_t bytesPerBlock = 686; // Big enough for 85 doubles
    kernel_bug<<<1, 1, bytesPerBlock>>>(bytesPerBlock);
    cudaError_t err = cudaDeviceSynchronize();
    if (err != cudaSuccess) {
        fprintf(stderr, "Kernel launch failed with error:\n\t%s\n", cudaGetErrorString(err));
        return 1;
    }
    return 0;
}

編集: CUDA 4.2 でもテストされていますが、問題は解決しません。

4

2 に答える 2

2

次の方法で問題を再現できました。

RHEL 5.5、ドライバー 304.54、CUDA 5.0、Quadro 5000 GPU。

次の問題を再現できませんでした。

RHEL 5.5、ドライバー 319.72、CUDA 5.5、Quadro 5000 GPU。

CUDA インストールを CUDA 5.5 に更新し、ドライバーを 319.72 以降に更新してください。

于 2014-02-09T04:22:46.310 に答える
0

クラッシュを特定しようとしている間は、呼び出しで 0 から 0.0 への暗黙的な変換を削除することをお勧めしX.fill(0);ます。これは有効な C++ ですが、CUDA では関数呼び出し演算子で一時変数を割り当てる際に問題が発生する可能性があります。確かに、彼らのドキュメントをざっと調べたところ、そのような一時ファイルがどこに割り当てられるかについての答えが見つかりませんでした-グローバルですか？デバイス？おそらくそれは問題ではありませんが、確かに.

于 2013-04-09T15:36:58.013 に答える