CUDA代码不起作用，为什么？

Question

我有这个代码，但有时它的工作，有时不（写printf（“ERR：％d \ n”，id））。 我在CUDA 4.1下工作，并具有计算能力2.1的GTS450。

可是没有更高的目标代码，荫只是想找到为什么它不工作，因为我的心告诉我，对不对：]

如果你想运行它，也许你需要执行几次，当出现“错误”或更改网格大小时！

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

这似乎发生，因为这 代码

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 !!! 

包含竞争条件，如果一些块写信给d_mans [男人] .m_id，但还是有不写信给d_ids [id]。可能编译器交换指令“设置为非零”和“下次保存新的ID”或缓存只是不及时更新。

实际上，问题在于你的分配器 - 最好记住上次使用的'man'的索引而不是查找它的索引。