これはCUDAのバグですか？ （不正なメモリアクセスが発生しました）

Question

私は、次のCUDAカーネル使用しています

：それは、このコードを使用して起動さ

__global__ 
void sum_worker(int *data, int *sum_ptr) 
{ 
     __shared__ int block_sum; 
     int idx = threadIdx.x; 
     int thread_sum = 0; 

     if (threadIdx.x == 0) 
       block_sum = 2; 

     for (int i = idx; i < MAX_INDEX; i += blockDim.x) 
       thread_sum += data[i]; 

     __syncthreads(); 

     atomicAdd(&block_sum, thread_sum); 

     __syncthreads(); 

     if (threadIdx.x == 0) 
       *sum_ptr = block_sum; 
} 

：

sum_worker<<<1, 32>>>(primes_or_zeros, sum_buffer); 

そして、それが働いているファイン（無ランタイムエラーをして生成します正しい結果）。

Cuda error 'an illegal memory access was encountered' in primes_gpu.cu at line 97 

cuda-memcheckでカーネルを実行している：

========= Invalid __global__ read of size 4 
=========  at 0x00000108 in /home/clifford/Work/handicraft/2016/perfmeas/primes_gpu.cu:35:sum_worker(int*, int*) 
=========  by thread (31,0,0) in block (0,0,0) 
=========  Address 0x703b70d7c is out of bounds 
=========  Saved host backtrace up to driver entry point at kernel launch time 
=========  Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x472225] 
=========  Host Frame:/usr/lib/x86_64-linux-gnu/libcudart.so.7.5 [0x146ad] 
=========  Host Frame:/usr/lib/x86_64-linux-gnu/libcudart.so.7.5 (cudaLaunch + 0x143) [0x2ece3] 
=========  Host Frame:./perfmeas [0x17c7] 
=========  Host Frame:./perfmeas [0x16b7] 
=========  Host Frame:./perfmeas [0x16e2] 
=========  Host Frame:./perfmeas [0x153f] 
=========  Host Frame:./perfmeas [0xdcd] 
=========  Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xf0) [0x20830] 
=========  Host Frame:./perfmeas [0xf39] 
.... 

アドレス0x703b70d7cがdataの境界の外に確かにある：私はi += blockDim.xi += 32に変更した場合しかし、私はエラー私はcudaDeviceSynchronize()呼び出す次の時間を取得します配列は0x703b40000から始まり、MAX_INDEX要素を持ちます。このテストでMAX_INDEXは50000です。これはCUDAのバグです

for (int i = idx; i < MAX_INDEX; i += 32) { 
      if (i >= MAX_INDEX) 
        printf("WTF!\n"); 
      thread_sum += data[i]; 
    } 

または私はここに愚かな何かをやっている： - （0x703b70d7c 0x703b40000）/ 4 = 50015.

がi >= 50000のための追加のチェックを追加するには、魔法のように離れて行く問題になりますか？

私はUbuntu 2016.04でCUDA 7.5を使用しています。 nvcc --versionの出力：

nvcc: NVIDIA (R) Cuda compiler driver 
Copyright (c) 2005-2015 NVIDIA Corporation 
Built on Tue_Aug_11_14:27:32_CDT_2015 
Cuda compilation tools, release 7.5, V7.5.17 

このテストケースのための完全なソースコードはここで見つけることができます：
http://svn.clifford.at/handicraft/2016/perfmeas

（オプション-gxで実行します。このバージョンはi += blockDim.xを使用している再現するi += 32にそれを変更します。 ）

編集：@njuffaは、彼はリンクオフにリンクしたくない彼は "コンピュータが何かを捕まえるかもしれない"と恐怖を感じており、&のペーストを直接スタックオーバーフローからコピーできるというテストケースを好むからです。だからここには行く：

#include <string.h> 
#include <stdio.h> 
#include <stdbool.h> 
#include <math.h> 

#define MAX_PRIMES 100000 
#define MAX_INDEX (MAX_PRIMES/2) 

__global__ 
void primes_worker(int *data) 
{ 
    int idx = threadIdx.x + blockIdx.x * blockDim.x; 

    if (idx >= MAX_INDEX) 
     return; 

    int p = 2*idx+1; 
    for (int i = 3; i*i <= p; i += 2) { 
     if (p % i == 0) { 
      data[idx] = 0; 
      return; 
     } 
    } 

    data[idx] = idx ? p : 0; 
} 

__global__ 
void sum_worker(int *data, int *sum_ptr) 
{ 
    __shared__ int block_sum; 
    int idx = threadIdx.x; 
    int thread_sum = 0; 

    if (threadIdx.x == 0) 
     block_sum = 2; 

#ifdef ENABLE_BUG 
    for (int i = idx; i < MAX_INDEX; i += 32) 
     thread_sum += data[i]; 
#else 
    for (int i = idx; i < MAX_INDEX; i += blockDim.x) 
     thread_sum += data[i]; 
#endif 

    __syncthreads(); 

    atomicAdd(&block_sum, thread_sum); 

    __syncthreads(); 

    if (threadIdx.x == 0) 
     *sum_ptr = block_sum; 
} 

int *primes_or_zeros; 
int *sum_buffer; 

void primes_gpu_init() 
{ 
    cudaError_t err; 

    err = cudaMalloc((void**)&primes_or_zeros, sizeof(int)*MAX_INDEX); 

    if (err != cudaSuccess) 
     printf("Cuda error '%s' in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__); 

    err = cudaMallocHost((void**)&sum_buffer, sizeof(int)); 

    if (err != cudaSuccess) 
     printf("Cuda error '%s' in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__); 
} 

void primes_gpu_done() 
{ 
    cudaError_t err; 

    err = cudaFree(primes_or_zeros); 

    if (err != cudaSuccess) 
     printf("Cuda error '%s' in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__); 

    err = cudaFreeHost(sum_buffer); 

    if (err != cudaSuccess) 
     printf("Cuda error '%s' in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__); 
} 

int primes_gpu() 
{ 
    int num_blocks = (MAX_INDEX + 31)/32; 
    int num_treads = 32; 

    primes_worker<<<num_blocks, num_treads>>>(primes_or_zeros); 
    sum_worker<<<1, 32>>>(primes_or_zeros, sum_buffer); 
    cudaError_t err = cudaDeviceSynchronize(); 

    if (err != cudaSuccess) 
     printf("Cuda error '%s' in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__); 

    return *sum_buffer; 
} 

int main() 
{ 
    primes_gpu_init(); 

    int result = primes_gpu(); 
    printf("Result: %d\n", result); 

    if (result != 454396537) { 
     printf("Incorrect result!\n"); 
     return 1; 
    } 

    primes_gpu_done(); 
    return 0; 
} 

使用法：

$ nvcc -o demo demo.cu 
$ ./demo 
Result: 454396537 

$ nvcc -D ENABLE_BUG -o demo demo.cu 
$ ./demo 
Cuda error 'an illegal memory access was encountered' in demo.cu at line 99 
Result: 0 
Incorrect result! 

Answer 1

TL; DR：観測された動作が非常に高い、CUDA 7.5ツールチェーンのptxasコンポーネントのバグにより特異的にループが発生していますアンローラー。このバグは、公開されているCUDA 8.0 RCですでに修正されている可能性があります。

質問で報告された動作を、Quadro K2200 GPU（sm_50デバイス）を搭載した64ビットWindows 7プラットフォームで再現できました。 ENABLE_BUGで定義された生成されたマシンコード（SASS）の主な違いは、ループが4倍に展開されていることです。これは、ループの増分がvariabe（threadIdx.x）からコンパイル時定数32に変更されたことによる直接的な結果です。これによりコンパイラはコンパイル時にトリップ数を計算できます。

中間PTXレベルで、ループがあっても32の増分で圧延されることに注意することは興味深いことである。

BB7_4: 
ld.global.u32 %r12, [%rd10]; 
add.s32 %r16, %r12, %r16; 
add.s64 %rd10, %rd10, 128; 
add.s32 %r15, %r15, 32; 
setp.lt.s32  %p3, %r15, 50000; 
@%p3 bra BB7_4; 

ループをマシンコードに展開されるように、それはそれを適用ptxas unrollerなければなりません変換。

は私が期待通りのコードは動作しますが、nvccコマンドラインで-Xptxas -O1を指定することで、-O1からptxas最適化レベルを下げた場合。 sm_30のコード（sm_50デバイスで実行しているときにJITコンパイルを引き起こす）を作成すると、最新のドライバであるWindows 369.26を実行すると、コードが正常に動作します。これは、CUDA 7.5のptxasコンポーネントのアンローラーにバグがあることを強く示唆していますが、CUDAドライバのptxasコンポーネントはCUDA 7.5ツールチェーンのptxasコンポーネントよりもはるかに新しいため、すでに修正されています。この場合、展開がアンロールループが既にPTXレベルで存在する意味、コンパイラのnvvmコンポーネントによって実行されているので、ループの前に直接#pragma unroll 4を配置

はまた、問題を修正：

#if ENABLE_BUG 
#pragma unroll 4 
    for (int i = idx; i < MAX_INDEX; i += 32) 
     thread_sum += data[i]; 
#else 

結果のPTX：[、最小完全、かつ検証例]なし

BB7_5: 
.pragma "nounroll"; 
ld.global.u32 %r34, [%rd14]; 
add.s32 %r35, %r34, %r45; 
ld.global.u32 %r36, [%rd14+128]; 
add.s32 %r37, %r36, %r35; 
ld.global.u32 %r38, [%rd14+256]; 
add.s32 %r39, %r38, %r37; 
ld.global.u32 %r40, [%rd14+384]; 
add.s32 %r45, %r40, %r39; 
add.s64 %rd14, %rd14, 512; 
add.s32 %r44, %r44, 128; 
setp.lt.s32  %p5, %r44, %r3; 
@%p5 bra BB7_5;