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