x1eU1rPz

1CUDA block synchronization differences between GTS 250 and Fermi devices
2__device__ void lock(int *mutex) {
3    while(atomicCAS(mutex, 0, 1) != 0);
4}
5
6__device__ void unlock(int *mutex) {
7    atomicExch(mutex, 0);
8}
9
10__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)
11{
12    // Find entry if it exists
13    unsigned int hashValue = hash(key, table.count);
14
15    lock(&globalHashLocks[hashValue]);
16
17    int bucketHead = table.entries[hashValue];
18    int currentLocation = bucketHead;
19
20    bool found = false;
21    Entry currentEntry;
22
23    while (currentLocation != -1 && !found) {
24        currentEntry = table.pool[currentLocation];
25        if (currentEntry.data.x == key) {
26            found = true;
27        } else {
28            currentLocation = currentEntry.next;
29        }
30    }
31
32    if (currentLocation == -1) {
33        // If entry does not exist, create entry
34        lock(globalFreeLock);
35        int newLocation = (*globalFirstFree)++;
36        __threadfence();
37        unlock(globalFreeLock);
38
39        Entry newEntry;
40        newEntry.data.x = key;
41        newEntry.data.y = count;
42        newEntry.data.z = sum;
43        newEntry.data.w = sumSquared;
44        newEntry.next = bucketHead;
45
46        // Add entry to table
47        table.pool[newLocation] = newEntry;
48        table.entries[hashValue] = newLocation;
49    } else {
50        currentEntry.data.y += count;
51        currentEntry.data.z += sum;
52        currentEntry.data.w += sumSquared;
53        table.pool[currentLocation] = currentEntry;
54    }
55
56    __threadfence();
57    unlock(&globalHashLocks[hashValue]);
58}