CUDAコードが動作しません、なぜですか？

Question

私はこのコードを持っていますが、ときどき動作します（printf（ "ERR：％d \ n"、id））。 私はCUDA 4.1で動作し、コンピューティング機能2.1であるGTS450を持っています。

コードdoesntのは、私の心はそれは右、私に言うので、IAMのはただ、そのが働いていない理由を見つけようとし、高い目的を持っている：]

あなたがそれを実行したい場合は、多分あなたは、数回実行する必要があります「エラー」が表示されたり、グリッドサイズが変更された場合

PS：here you can download exe file for win64 - you need to have cuda4.1 driver

class MAN 
{ 
public: 
    int m_id; 
    int m_use; 

    __device__ 
    MAN() 
    { 
     m_id = -1; 
     m_use = 0; 
    } 
}; 

__device__ int* d_ids = NULL; 
__device__ int d_last_ids = 0; 

__device__ MAN* d_mans = NULL; 


__global__ void init() 
{ 
    d_mans = new MAN[500]; //note: 500 is more than enough! 
    d_ids = new int[500]; 

    for(int i=0; i < 500; i++) 
     d_ids[i] = 0; 
} 


__device__ int getMAN() //every block get unique number, so at one moment all running blocks has different id 
{ 
    while(true) 
    { 
     for(int i=0; i < 500; i++) 
      if(atomicCAS(&(d_mans[i].m_use), 0, 1)==0) 
       return i; 
    } 
} 
__device__ void returnMAN(int id) 
{ 
    int s = atomicExch(&(d_mans[id].m_use), 0); 
} 



__global__ void testIt() 
{ 
    if(threadIdx.x==0) 
    { 
     int man = getMAN(); 

     int id = d_mans[man].m_id; 
     if(id == -1) //If It never works with this "id", its creating new 
     { 
      id = atomicAdd(&d_last_ids, 2); 

      d_ids[id] = 10; //set to non-zero 
      d_mans[man].m_id = id; //save new id for next time 

      printf("ADD:%d\n", id); 
     } 

     if(d_ids[id]==0) 
      printf("ERR:%d\n", id); //THIS SHOULD NEVER HAPPEN, BUT BECOMES !!! 

     returnMAN(man); 
    } 
} 



int main() 
{ 
    init<<<1, 1>>>(); 
    printf("init() err: %d\n", cudaDeviceSynchronize()); 

    testIt<<<20000, 512>>>(); 
    printf("testIt() err: %d\n", cudaDeviceSynchronize()); 

    getchar(); 
    return 0; 
} 

Answer 1

私は、この変更されました：これまで

__device__ int* d_ids = NULL; 

：

__device__ volatile int* d_ids = NULL; 

をし、それがうまく動作！

さらに、__threadfence（）は必要ありません。

Answer 2

いくつかのブロックがd_mansに書いた場合 このコード

int id = d_mans[man].m_id; 
    if(id == -1) //If It never works with this "id", its creating new 
    { 
     id = atomicAdd(&d_last_ids, 2); 

     d_ids[id] = 10; //set to non-zero 
     d_mans[man].m_id = id; //save new id for next time 

     printf("ADD:%d\n", id); 
    } 

    if(d_ids[id]==0) 
     printf("ERR:%d\n", id); //THIS SHOULD NEVER HAPPEN, BUT BECOMES !!! 

は競合状態が含まれていますので、これは、発生するようです[男]は.m_id、まだ持っていませんd_ids [id]に書きました。おそらく、コンパイラは命令を「非ゼロに設定」し、「次回のために新しいIDを保存する」か、またはキャッシュはただちに更新されないでしょう。

実際には、問題はアロケータにあります。最後に使用された 'man'のインデックスを探す方が良いでしょう。