GTS 250 和 Fermi 器件之间的 CUDA 块同步差异

Question

所以我一直在研究一个在全局内存中创建哈希表的程序。该代码在 GTS250（Compute 1.1 设备）上完全可用（尽管速度较慢）。但是，在 Compute 2.0 设备（C2050 或 C2070）上，哈希表已损坏（数据不正确，指针有时错误）。

基本上，当仅使用一个块（两个设备）时，代码可以正常工作。但是，当使用 2 个或更多块时，它仅适用于 GTS250，不适用于任何 Fermi 器件。

我知道两个平台之间的扭曲调度和内存架构是不同的，我在开发代码时考虑到了这一点。根据我的理解，使用 __theadfence() 应该确保任何全局写入都已提交并且对其他块可见，但是，从损坏的哈希表来看，它们似乎并非如此。

我还在 NVIDIA CUDA 开发者论坛上发布了该问题，可以在此处。

相关代码如下：

__device__ void lock(int *mutex) {
    while(atomicCAS(mutex, 0, 1) != 0);
}

__device__ void unlock(int *mutex) {
    atomicExch(mutex, 0);
}

__device__ void add_to_global_hash_table(unsigned int key, unsigned int count, unsigned int sum, unsigned int sumSquared, Table table, int *globalHashLocks, int *globalFreeLock, int *globalFirstFree)
{
    // Find entry if it exists
    unsigned int hashValue = hash(key, table.count);

    lock(&globalHashLocks[hashValue]);

    int bucketHead = table.entries[hashValue];
    int currentLocation = bucketHead;

    bool found = false;
    Entry currentEntry;

    while (currentLocation != -1 && !found) {
        currentEntry = table.pool[currentLocation];
        if (currentEntry.data.x == key) {
            found = true;
        } else {
            currentLocation = currentEntry.next;
        }
    }

    if (currentLocation == -1) {
        // If entry does not exist, create entry
        lock(globalFreeLock);
        int newLocation = (*globalFirstFree)++;
        __threadfence();
        unlock(globalFreeLock);

        Entry newEntry;
        newEntry.data.x = key;
        newEntry.data.y = count;
        newEntry.data.z = sum;
        newEntry.data.w = sumSquared;
        newEntry.next = bucketHead;

        // Add entry to table
        table.pool[newLocation] = newEntry;
        table.entries[hashValue] = newLocation;
    } else {
        currentEntry.data.y += count;
        currentEntry.data.z += sum;
        currentEntry.data.w += sumSquared;
        table.pool[currentLocation] = currentEntry;
    }

    __threadfence();
    unlock(&globalHashLocks[hashValue]);
}

Answer 1

正如 LSChien 在此 post，问题在于 L1 缓存一致性。虽然使用 __threadfence() 将保证共享和全局内存写入对其他线程可见，因为它不是原子的，块 1 中的线程 x可能会达到缓存的内存值，直到块0中的线程y执行到threadfence指令。相反，LSChien 在他的帖子中建议使用 atomicCAS() 强制线程从全局内存中读取而不是缓存的值。执行此操作的正确方法是将内存声明为易失性，要求对该内存的每次写入都立即对网格中的所有其他线程可见。

Answer 2

__threadfence 保证在返回之前对全局内存的写入对于当前块中的其他线程可见。这和“全局内存写操作完成”不一样！考虑每个多核上的缓存。