来自共享内存中的数据结构的竞态检查错误

Posted 2023-03-23

【中文标题】来自共享内存中的数据结构的竞态检查错误

【英文标题】：racecheck error from a data structure in shared memory

【发布时间】：2022-01-18 14:40:09

【问题描述】：

我有一个数据结构哈希表，它具有线性探测哈希方案，并被设计为使用 CAS 无锁。

哈希表

constexpr uint64_t HASH_EMPTY = 0xffffffffffffffff;

struct OnceLock 

    static const unsigned LOCK_FRESH   = 0;
    static const unsigned LOCK_WORKING = 1;
    static const unsigned LOCK_DONE    = 2;

    volatile unsigned lock;

    __device__ void init() 
        lock = LOCK_FRESH;
     

    __device__ bool enter() 
        unsigned lockState = atomicCAS ( (unsigned*) &lock, LOCK_FRESH, LOCK_WORKING );
        return lockState == LOCK_FRESH;
    

    __device__ void done() 
        __threadfence();
        lock = LOCK_DONE;
        __threadfence();
    

    __device__ void wait() 
        while ( lock != LOCK_DONE );
    
;
template <typename T>
struct agg_ht 
    OnceLock lock;
    uint64_t hash;
    T payload;
;


template <typename T>
__global__ void initAggHT ( agg_ht<T>* ht, int32_t num ) 
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < num; i += blockDim.x * gridDim.x) 
    ht[i].lock.init();
    ht[i].hash = HASH_EMPTY;
    



// returns candidate bucket
template <typename T>
__device__ int hashAggregateGetBucket ( agg_ht<T>* ht, int32_t ht_size, uint64_t grouphash, int& numLookups, T* payl ) 
    int location=-1;
    bool done=false;
    while ( !done ) 
        location = ( grouphash + numLookups ) % ht_size;
        agg_ht<T>& entry = ht [ location ];
        numLookups++;
        if ( entry.lock.enter() ) 
            entry.payload = *payl;
            entry.hash = grouphash;
            entry.lock.done();
        
        entry.lock.wait();
        done = (entry.hash == grouphash);
        if ( numLookups == ht_size ) 
            printf ( "agg_ht hash table full at threadIdx %d & blockIdx %d \n", threadIdx.x, blockIdx.x );
            break;
        
    
    return location;

然后我有一个最小的内核以及主函数，只是为了让哈希表运行。重要的是哈希表用__shared__注解，它被分配在一个SM的共享内存中以便快速访问。 （我没有添加任何带有cudaMalloc 的输入数据以保持示例最小化。）

#include <cstdint>
#include <cstdio>

/**hash table implementation**/

constexpr int HT_SIZE = 1024;

__global__ void kernel()  
    __shared__ agg_ht<int> aht2[HT_SIZE]; 
    
        int ht_index;
        unsigned loopVar = threadIdx.x;
        unsigned step = blockDim.x;
        while(loopVar < HT_SIZE) 
            ht_index = loopVar;
            aht2[ht_index].lock.init();
            aht2[ht_index].hash = HASH_EMPTY;
            loopVar += step;
        
    

    int key = 1;
    int value = threadIdx.x;

    __syncthreads();

    int bucket = -1;
    int bucketFound = 0;
    int numLookups = 0;
    while(!(bucketFound)) 
        bucket = hashAggregateGetBucket ( aht2, HT_SIZE, key, numLookups, &(value));
        int probepayl = aht2[bucket].payload;
        bucketFound = 1;
        bucketFound &= ((value == probepayl));
    


int main() 
    kernel<<<1, 128>>>();
    cudaDeviceSynchronize();
    return 0;

标准编译方式，如果文件名为test.cu： $ nvcc -G test.cu -o test

我不得不说，这个哈希表会总是在大量输入下的并发插入期间给我正确的答案。 但是，当我在上面运行racecheck 时，我到处都看到了错误：

$ compute-sanitizer --tool racecheck ./test
========= COMPUTE-SANITIZER
========= Error: Race reported between Write access at 0xd20 in /tmp/test.cu:61:int hashAggregateGetBucket<int>(agg_ht<T1> *, int, unsigned long, int &, T1 *)
=========     and Read access at 0xe50 in /tmp/test.cu:65:int hashAggregateGetBucket<int>(agg_ht<T1> *, int, unsigned long, int &, T1 *) [1016 hazards]
========= 
========= Error: Race reported between Write access at 0x180 in /tmp/test.cu:25:OnceLock::done()
=========     and Read access at 0xd0 in /tmp/test.cu:30:OnceLock::wait() [992 hazards]
========= 
========= Error: Race reported between Write access at 0xcb0 in /tmp/test.cu:60:int hashAggregateGetBucket<int>(agg_ht<T1> *, int, unsigned long, int &, T1 *)
=========     and Read access at 0x1070 in /tmp/test.cu:103:kernel() [508 hazards]
========= 
========= RACECHECK SUMMARY: 3 hazards displayed (3 errors, 0 warnings)

我很困惑，我相信这个线性探测哈希表可以通过我的单元测试，但到处都有数据竞争的危险。我想这些危险与正确性无关。 (?)

经过一段时间的调试，我仍然无法消除危险错误。我坚信volatile 是原因。我希望有人能够对此有所了解，并帮助我解决这些烦人的危险。

我也希望这个问题可以反映一些关于该主题的设计思想：共享内存上的数据结构。在 *** 上搜索时，我看到的只是共享内存中的普通原始数组。

【参考方案1】：

我认为这些危险与正确性无关。 (?)

我不会尝试证明您的应用程序或算法的“正确性”。如果这就是你要找的，请忽略我的回答。

希望有人能解释一下

当一个线程写入共享内存中的某个位置，而另一个线程从该位置读取，并且代码中没有干预同步以确保写入发生在读取之前（或者可能，更准确地说，写入的值对读取线程是可见的）。这不是careful, exhaustive definition，但对于我们在这里处理的内容来说已经足够了。

就该定义而言，您的代码中肯定有该活动。一个被标记的具体案例是一个线程写在这里：

        entry.hash = grouphash;

另一个线程在这里读取相同的位置：

    done = (entry.hash == grouphash);

检查您的代码，我们可以看到这两个代码位置之间没有__syncthreads() 语句。此外，由于包含该活动的循环，与此相关的危害不止一种（有两种）。

另一个被标记的交互是一个线程在此处写入lock：

        entry.lock.done();

另一个线程在此处读取相同的lock 位置：

    entry.lock.wait();

这里报告的危险实际上是针对其他代码行报告的，因为它们都是函数调用。同样，没有干预同步。

我承认，由于您的应用程序的循环性质，我不确定是否需要尽早采用这些线程间通信路径中的任何一个来确保“正确性”。但是，我没有仔细研究过您的申请，也无意说明正确性。

帮我解决那些恼人的危险。

事实上，这两种交互都在代码的一小部分中，因此根据我的测试，我们可以通过以下添加来消除这 3 个危害：

    __syncthreads();  // add this line
    entry.lock.wait();
    done = (entry.hash == grouphash);
    __syncthreads();  // add this line

第一个同步与我已经指出的行之间明显的读写连接相交。由于此时代码的循环性质，需要第二次同步。

还请注意，proper usage 或 __syncthreads() 使得线程块中的所有线程都可以到达该同步点。快速浏览您在此处所拥有的内容并没有向我暗示需要仔细处理上述行/添加，但您应该确认这一点并注意一般应用程序/用法。 while bucketFound 循环可能会在此处创建一个应该以不同方式处理的情况，但是 compute-sanitizer --tool synccheck 没有报告任何问题，在 V100 上运行，并添加了我在此处建议的添加。

【讨论】：

嗨罗伯特。谢谢（你的）信息。由于它是一个哈希表，我们不能假设一个块中的所有线程都具有相同的执行路径并到达__syncthreads( ) 的同一点。如果内核更复杂，哈希表插入只是内核的一小部分，那么这个假设就不能总是成立。在这种复杂的情况下，您是否也有有效解决方案的提示？

是的，修复相当机械。 1.不允许任何线程退出。 2. 适当调节只有少数线程处于活动状态的区域。 3. 不要以__syncthreads() 语句为条件。因此，只有您选择的线程在工作，但所有线程都参与__syncthreads() 语句。 This 给出了总体思路。它允许您的代码具有条件字符，同时也满足__syncthreads() 的要求。