Различия в синхронизации блоков CUDA между устройствами GTS 250 и Fermi - PullRequest
0 голосов
/ 04 апреля 2011

Итак, я работал над программой, в которой я создаю хеш-таблицу в глобальной памяти. Код полностью функционален (хотя и медленнее) на 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]);
}

Ответы [ 2 ]

0 голосов
/ 08 апреля 2011

Как указано LSChien в этой публикации , проблема заключается в когерентности кэша L1.Хотя использование __threadfence() гарантирует, что записи в общую и глобальную память будут видны другим потокам, поскольку она не является атомарной, thread x в block 1 может достигать значения в кешируемой памяти до тех пор, пока thread y в block 0 не будет выполнена для потокаинструкция.Вместо этого LSChien предложил взломать в своем посте использование atomicCAS(), чтобы заставить поток читать из глобальной памяти вместо кэшированного значения.Правильный способ сделать это - объявить память как volatile, требуя, чтобы каждая запись в эту память была видна всем остальным потокам в сетке.

0 голосов
/ 04 апреля 2011

__ threadfence гарантирует, что записи в глобальную память видны другим потокам в текущем блоке перед возвратом. Это не то же самое, что «операция записи в глобальную память завершена»! Подумайте, кэширование на каждом многоядерном.

...