ืืื ืึทืจืืึทื ืืขืฉืืงื ืขืก ืืืืฃ Github .
ืขืก ืืื ืึท ืคึผืฉืื ืืคึผื ืืึทืฉ ืืืฉ ืืืึธืก ืืื ืืืืืืช ืฆื ืคึผืจืึทืกืขืกืื ื ืืื ืืขืจืืขืจ ืคืื ืืืืืึทื ื ืคืื ืื ืกืขืจืฅ ืคึผืขืจ ืกืขืงืื ืืข. ืืืืฃ ืืืื NVIDIA GTX 1060 ืืึทืคึผืืึทืคึผ, ืื ืงืึธื ืื ืกืขืจืฅ 64 ืืืืืึธื ืจืึทื ืืึทืืื ืืืฉืขื ืขืจืืืืึทื ืฉืืืกื-ืืืขืจื ืคึผืขืจื ืืื ืืืขืื 210 ืืื ืืื ืจืืืืืื 32 ืืืืืึธื ืคึผืขืจื ืืื ืืืขืื 64 ืืื.
ืึทื ืืื, ืื ืืืืงืืึทื ืืืืฃ ืึท ืืึทืคึผืืึทืคึผ ืืื ืืขืขืจืขื 300 ืืืืืึธื ืื ืกืขืจืฅ / ืกืขืง ืืื 500 ืืืืืึธื ืืืืืฅ / ืกืขืง.
ืืขืจ ืืืฉ ืืื ืืขืฉืจืืื ืืื CUDA, ืืึธืืฉ ืื ืืขืืืข ืืขืื ืืง ืงืขื ืขื ืืืื ืืขืืืขื ืื ืฆื HLSL ืึธืืขืจ GLSL. ืื ืืืคึผืืึทืืขื ืืืืฉืึทื ืืื ืขืืืขืืข ืืืืืืืืฉืึทื ื ืฆื ืขื ืฉืืจ ืืืื ืคืึธืจืฉืืขืืื ื ืืืืฃ ืึท ืืืืืขื ืงืึทืจืื:
- ืืืืื 32-ืืืกื ืฉืืืกืืขื ืืื ืื ืืขืืืข ืืืึทืืืขืก ืืขื ืขื ืคึผืจืึทืกืขืกื.
- ืื ืืึทืฉ ืืืฉ ืืื ืึท ืคืึทืจืคืขืกืืืงื ืืจืืืก.
- ืืื ืืขื ืืจืืืก ืืืื ืืืื ืืืืึทื ืฆื ืฆืืืื ืฆื ืื ืืึทืื.
ืคึฟืึทืจ ืฉืืืกืืขื ืืื ืืืึทืืืขืก, ืืืจ ืืึทืจืคึฟื ืฆื ืจืขืืขืจืืืืจื ืึท ืคึผืฉืื ืืขืืืืืืขืจ ืืึทืจืงืขืจ (ืืื ืื ืืืืื ืงืึธื, ืืึธืก ืืื 0xffffffff).
ืืึทืฉ ืืืฉ ืึธื ืืึทืงืก
ืื ืืึทืฉ ืืืฉ ื ืืฆื ืขืคืขื ืขื ืึทืืจืขืกืื ื ืืื , ืืึธืก ืืื, ืขืก ืืื ืคืฉืื ืึท ืืขื ืืข ืคืื โโืฉืืืกื-ืืืขืจื ืคึผืขืจื ืืืึธืก ืืื ืกืืึธืจื ืืื ืืึผืจืื ืืื ืืื ืืขืืขืจ ืงืึทืฉ ืคืึธืจืฉืืขืืื ื. ืืขืจ ืืขืืืืงืขืจ ืงืขื ืขื ื ืื ืืืื ืืขืืึธืื ืคึฟืึทืจ ืืฉืึทืื ืื ื, ืืืึธืก ืื ืืืึทืืืื ืืืื ืคึฟืึทืจ ืึท ืืืึทืื ืืื ืึท ืืื ืืงื ืจืฉืืื. ื ืืึทืฉ ืืืฉ ืืื ืึท ืคึผืฉืื ืืขื ืืข ืกืืึธืจืื ื ืขืืขืืขื ืื KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
ืื ืืจืืืก ืคืื ืื ืืืฉ ืืื ืึท ืืึทืื ืคืื ืฆืืืื, ื ืืฉื ืึท ืืืืคึผื ื ืืืขืจ, ืืืืึทื ืืืื ืฉื ืขื ืื ืกืืจืืงืืืึธื ืืื ืืขื ืื ืฆื ืฆืืืืืื ืื ืคึผืึธืื2 / AND ืืึทืกืงืข, ืึธืืขืจ ืื ืืึธืืืืืก ืึธืคึผืขืจืึทืืึธืจ ืืื ืคืื ืกืืึธืืขืจ. ืืึธืก ืืื ืืืืืืืง ืืื ืืขื ืคืึทื ืคืื ืืื ืขืึทืจ ืคึผืจืึธืืืื ื, ืืืืึทื ืืื ืึท ืืื ืขืึทืจ ืืืฉ ืืืงืึทืคึผ ืื ืฉืคึผืขืืื ืืื ืืขืงืก ืืืื ืืืื ืืื ืืขืืืืงืื ืืื ืืขืืขืจ ืฉืคึผืขืืื. ืืื ืืื ืึท ืจืขืืืืืึทื, ืื ืคึผืจืืึทื ืคืื ืื ืึธืคึผืขืจืึทืฆืืข ืืื ืืืกืืฃ ืืึธืืืืึธ ืืื ืืขืืขืจ ืฉืคึผืขืืื.
ืืขืจ ืืืฉ ืืืืื ืกืืึธืจื ืื ืฉืืืกื ืืื ืืืขืจื ืคึฟืึทืจ ืืขืืขืจ ืขืืขืืขื ื, ื ืืฉื ืึท ืืึทืฉ ืคืื ืื ืฉืืืกื. ืืื ื ืื ืืืฉ ืืืืื ืกืืึธืจื 32-ืืืกื ืฉืืืกืืขื, ืื ืืึทืฉ ืืื ืงืึทืืงืืึทืืืืืื ืืืืขืจ ืืขืฉืืืื ื. ืืขืจ ืงืึธื ืืืืื ื ืืฆื ืื Murmur3 ืืึทืฉ, ืืืึธืก ืืืืื ืคึผืขืจืคืึธืจืื ืึท ืืืกื ืฉืืคืฅ, XORs ืืื ืงืืืคื.
ืื ืืึทืฉ ืืืฉ ื ืืฆื ืืึทืงืื ื ืฉืืฅ ืืขืงื ืืงืก ืืืึธืก ืืขื ืขื ืคืจืืึท ืคืื ืืึผืจืื ืกืืจ. ืืคืืื ืืืื ืขืืืขืืข ืฉืจืืึทืื ืึทืคึผืขืจืืืฉืึทื ื ืฆืขืฉืืขืจื ืื ืกืืจ ืคืื ืื ืืขืจืข ืึทืืึท ืึทืคึผืขืจืืืฉืึทื ื, ืื ืืึทืฉ ืืืฉ ืืืขื ื ืึธื ืืึทืืื ืื ืจืืืืืง ืฉืืึทื. ืืืจ ืืืขืื ืจืขืื ืืืขืื ืืขื ืืื ืื. ืื ืืขืื ืืง ืึทืจืืขื ืืจืืืก ืืื ืืืืืขื ืงืึทืจืืก ืืืึธืก ืืืืคื ืืืืื ืืขืจ ืคืื ืคึฟืขืืขื ืงืึทื ืงืขืจืึทื ืืื.
ืื ืฉืืืกืืขื ืืื ืืืึทืืืขืก ืืื ืื ืืึทืฉ ืืืฉ ืืขื ืขื ืื ืืืืึทืืืืื ืฆื ืืืืืืง.
ืืขืจ ืงืึธื ืงืขื ืขื ืืืื ืืึทืืึทืคืืื ืฆื ืฉืขืคึผื 64-ืืืกื ืฉืืืกืืขื ืืื ืืืึทืืืขืก. ืงืื ืืึทืจืคื ืึทืืึธืืืฉืข ืืืืขื ืขื, ืฉืจืืึทืื ืืื ืคืึทืจืืืืึทืื-ืืื-ืืืกืืืึทืื ืึทืคึผืขืจืืืฉืึทื ื. ืืื ืืืึทืืืขืก ืืึทืจืคื ืึทืืึธืืืฉืข ืืืืขื ืขื ืืื ืฉืจืืึทืื ืึทืคึผืขืจืืืฉืึทื ื. ืฆืื ืืืืง, ืืื CUDA, ืืืืขื ืขื-ืฉืจืืึทืื ืึทืคึผืขืจืืืฉืึทื ื ืคึฟืึทืจ 32- ืืื 64-ืืืกื ืืืึทืืืขืก ืืขื ืขื ืึทืืึธืืืฉืข ืืื ืืึทื ื ืืื ืืื ืืขื ืขื ืืขืืืืื ืืืขื ืึทืืืื ื (ืืขื ืืืืืืขืจ). ), ืืื ืืึธืืขืจื ืืืืืขื ืงืึทืจืืก ืฉืืืฆื 64-ืืืกื ืึทืืึธืืืฉืข ืคืึทืจืืืืึทืื-ืืื-ืืืขืงืกื ืึทืคึผืขืจืืืฉืึทื ื. ืคืื ืงืืจืก, ืืืขื ืืึธืืืื ื ืฆื 64 ืืืฅ, ืคืึธืจืฉืืขืืื ื ืืืขื ืคืึทืจืืื ืขืจื ืึท ืืืกื.
ืืึทืฉ ืืืฉ ืฉืืึทื
ืืขืืขืจ ืฉืืืกื-ืืืขืจื ืคึผืึธืจ ืืื ืึท ืืึทืฉ ืืืฉ ืงืขื ืขื ืืึธืื ืืืื ืขืจ ืคืื ืคืืจ ืฉืืึทืื:
- ืฉืืืกื ืืื ืืืขืจื ืืขื ืขื ืืืืืืง. ืืื ืืขื ืฉืืึทื, ืื ืืึทืฉ ืืืฉ ืืื ืื ืืฉืึทืืืืื.
- ืืขืจ ืฉืืืกื ืืื ืึทืจืึธืคึผืืขืฉืจืืื, ืึธืืขืจ ืื ืืืขืจื ืืื ื ืึธื ื ืืฉื ืืขืฉืจืืื. ืืืื ืื ืื ืืขืจ ืคืึธืืขื ืืื ืืืฆื ืืืืขื ืขื ืืึทืื, ืขืก ืงืขืจื ืืืืืืง. ืืึธืก ืืื ื ืึธืจืืึทื, ืื ืืขืืืข ืืึทื ืืืึธืื ืคึผืึทืกืืจื ืืืื ืื ืื ืืขืจ ืคืึธืืขื ืคืื ืืืจืืคืืจืื ื ืืื ืืขืืจืืขื ืึท ืืืกื ืคืจืืขืจ, ืืื ืืืจ ืืขื ืขื ืืขืจืขืื ืืืขืื ืึท ืงืึทื ืงืขืจืึทื ื ืืึทืื ืกืืจืืงืืืจ.
- ืืืืืข ืื ืฉืืืกื ืืื ืื ืืืขืจื ืืขื ืขื ืจืขืงืึธืจืืขื.
- ืื ืืืขืจื ืืื ืืืจืขืืืืื ืคึฟืึทืจ ืื ืืขืจืข ืคึฟืขืืขื ืคืื ืืืจืืคืืจืื ื, ืึธืืขืจ ืืขืจ ืฉืืืกื ืืื ื ืืฉื ื ืึธื. ืืึธืก ืงืขื ืคึผืึทืกืืจื ืืืืึทื ืื CUDA ืคึผืจืึธืืจืึทืืืื ื ืืึธืืขื ืืื ืึท ืืืกืื ืึธืจืืขืจื ืืึผืจืื ืืึธืืขื. ืืึธืก ืืื ื ืึธืจืืึทื; ืืื ืงืืื ืคืึทื, ืืขืจ ืฉืืืกื ืืื ื ืึธื ืืืืืืง, ืืคืืื ืืืื ืื ืืืขืจื ืืื ื ืื ืืขืจ ืึทืืื.
ืึท ืืืืืืืง ื ืืึทื ืก ืืื ืึทื ืึทืืึธื ืืขืจ ืฉืืืกื ืืื ืืขืฉืจืืื ืฆื ืื ืฉืคึผืขืืื, ืขืก ืืื ื ืื ืืขืจ ืืืืืขืื - ืืคืืื ืืืื ืืขืจ ืฉืืืกื ืืื ืืืืกืืขืืขืงื, ืืืจ ืืืขืื ืจืขืื ืืืขืื ืืขื ืืื ืื.
ืืขืจ ืืึทืฉ ืืืฉ ืงืึธื ืืคืืื ืึทืจืืขื ืืื ืืืกืื ืึธืจืืขืจื ืืึผืจืื ืืึธืืขืืก ืืื ืืืึธืก ืื ืกืืจ ืืื ืืืึธืก ืืึผืจืื ืืื ืืืืขื ืขื ืืื ืืขืฉืจืืื ืืื ืืืืืึทืงืึทื ื. ืืืขื ืืืจ ืงืืงื ืืื ืื ืกืขืจืฉืึทื, ืืืงืึทืคึผ ืืื ืืืืืฉืึทื ืืื ืึท ืืึทืฉ ืืืฉ, ืืขืืขื ืงืขื ืึทื ืืขืืขืจ ืฉืืืกื-ืืืขืจื ืคึผืึธืจ ืืื ืืื ืืืื ืขืจ ืคืื ืื ืคืืจ ืฉืืึทืื ืืืกืงืจืืืื ืืืืื.
ืื ืกืขืจืืื ื ืืื ืึท ืืึทืฉ ืืืฉ
ืื CUDA ืคึฟืื ืงืฆืืข ืืืึธืก ืื ืกืขืจืฅ ืฉืืืกื-ืืืขืจื ืคึผืขืจื ืืื ืึท ืืึทืฉ ืืืฉ ืงืืงื ืืื ืืึธืก:
void gpu_hashtable_insert(KeyValue* hashtable, uint32_t key, uint32_t value)
{
uint32_t slot = hash(key);
while (true)
{
uint32_t prev = atomicCAS(&hashtable[slot].key, kEmpty, key);
if (prev == kEmpty || prev == key)
{
hashtable[slot].value = value;
break;
}
slot = (slot + 1) & (kHashTableCapacity-1);
}
}
ืฆื ืึทืจืืึทื ืืืืื ืึท ืฉืืืกื, ืื ืงืึธื ืืืขืจืืืฅ ืืืจื ืื ืืึทืฉ ืืืฉ ืืขื ืืข ืกืืึทืจืืื ื ืืื ืื ืืึทืฉ ืคืื ืื ืื ืกืขืจืืึทื ืฉืืืกื. ืืขืืขืจ ืฉืคึผืขืืื ืืื ืื ืืขื ืืข ืคึผืขืจืคืึธืจืื ืึทื ืึทืืึธืืืฉืข ืคืึทืจืืืืึทืื-ืืื-ืืืกืืืึทืื ืึธืคึผืขืจืึทืฆืืข ืืืึธืก ืงืึทืืคึผืขืจื ืื ืฉืืืกื ืืื ืืขื ืฉืคึผืขืืื ืฆื ืืืืืืง. ืืืื ืึท ืืืกืืึทืืฉ ืืื ืืืืขืงืืึทื, ืืขืจ ืฉืืืกื ืืื ืื ืฉืคึผืขืืื ืืื ืืขืจืืืึทื ืืืงื ืืื ืื ืื ืกืขืจืืึทื ืฉืืืกื, ืืื ืืขืจ ืึธืจืืืื ืขื ืฉืคึผืขืืื ืฉืืืกื ืืื ืืืืืขืงืขืจื. ืืืื ืืขืจ ืึธืจืืืื ืขื ืฉืืืกื ืืื ืืืืืืง ืึธืืขืจ ืืึทืืฉื ืื ืื ืกืขืจืืึทื ืฉืืืกื, ืืขืจ ืงืึธื ืืขืคึฟืื ืขื ืึท ืคึผืึทืกืืง ืฉืคึผืขืืื ืคึฟืึทืจ ืื ืกืขืจืฉืึทื ืืื ืื ืกืขืจืืึทื ืื ืื ืกืขืจืืึทื ืืืขืจื ืืื ืื ืฉืคึผืขืืื.
ืืืื ืืื ืืืื ืงืขืจื ืจืืคื gpu_hashtable_insert() ืขืก ืืขื ืขื ืงืืืคื ืขืืขืืขื ืื ืืื ืืขืจ ืืขืืืืงืขืจ ืฉืืืกื, ืืื ืงืืื ืคืื ืืืืขืจ ืืืึทืืืขืก ืงืขื ืขื ืืืื ืืขืฉืจืืื ืฆื ืื ืฉืืืกื ืฉืคึผืขืืื. ืืึธืก ืืื ืืขืจืขืื ื ืืื ื ืึธืจืืึทื: ืืืื ืขืจ ืคืื ืื ืฉืืืกื-ืืืขืจื ืฉืจืืืื ืืขืฉืึทืก ืื ืจืืคื ืืืขื ืืืื ืืขืจืึธืื, ืึธืืขืจ ืืื ื ืึทืืข ืืขื ืืึทืคึผืึทื ื ืืื ืคึผืึทืจืึทืืขื ืืื ืขืืืขืืข ืคึฟืขืืขื ืคืื ืืืจืืคืืจืื ื, ืืืจ ืงืขื ืขื ื ืืฉื ืคืึธืจืืืกืืึธืื ืืืึธืก ืืึผืจืื ืฉืจืืึทืื ืืืขื ืืืื ืื ืืขืฆืืข.
ืืึทืฉ ืืืฉ ืืืื
ืงืึธื ืคึฟืึทืจ ืืืื ืฉืืืกืืขื:
uint32_t gpu_hashtable_lookup(KeyValue* hashtable, uint32_t key)
{
uint32_t slot = hash(key);
while (true)
{
if (hashtable[slot].key == key)
{
return hashtable[slot].value;
}
if (hashtable[slot].key == kEmpty)
{
return kEmpty;
}
slot = (slot + 1) & (kHashTableCapacity - 1);
}
}
ืฆื ืืขืคึฟืื ืขื ืื ืืืขืจื ืคืื ืึท ืฉืืืกื ืกืืึธืจื ืืื ืึท ืืืฉ, ืืืจ ืืืขืจืืื ืืืจื ืื ืืขื ืืข ืกืืึทืจืืื ื ืืื ืื ืืึทืฉ ืคืื ืื ืฉืืืกื ืืืึธืก ืืืจ ืืขื ืขื ืงืืงื ืคึฟืึทืจ. ืืื ืืขืืขืจ ืฉืคึผืขืืื, ืืืจ ืงืึธื ืืจืึธืืืจื ืืืื ืืขืจ ืฉืืืกื ืืื ืืขืจ ืืืึธืก ืืืจ ืืขื ืขื ืงืืงื ืคึฟืึทืจ, ืืื ืืืื ืึทืืื, ืืืจ ืฆืืจืืงืงืืืขื ืืืื ืืืขืจื. ืืืจ ืืืื ืงืึธื ืืจืึธืืืจื ืืืื ืืขืจ ืฉืืืกื ืืื ืืืืืืง, ืืื ืืืื ืึทืืื, ืืืจ ืึทืืึธืจื ืื ืืืื.
ืืืื ืืืจ ืงืขื ืขื ื ืืฉื ืืขืคึฟืื ืขื ืืขื ืฉืืืกื, ืืขืจ ืงืึธื ืงืขืจื ืึท ืืืืืืง ืืืขืจื.
ืึทืืข ืื ืืืื ืึทืคึผืขืจืืืฉืึทื ื ืงืขื ืขื ืืืื ืืืจืืืขืงืึธืื ืงืึทื ืงืขืจืึทื ืืื ืืืจื ืื ืกืขืจืฉืึทื ื ืืื ืืืืืฉืึทื ื. ืืขืืขืจ ืคึผืึธืจ ืืื ืื ืืืฉ ืืืขื ืืึธืื ืืืื ืขืจ ืคืื ืื ืคืืจ ืฉืืึทืื ืืืกืงืจืืืื ืืืืื ืคึฟืึทืจ ืื ืืืืคื.
ืืืืืืื ื ืืื ืึท ืืึทืฉ ืืืฉ
ืงืึธื ืฆื ืืืกืืขืงื ืฉืืืกืืขื:
void gpu_hashtable_delete(KeyValue* hashtable, uint32_t key, uint32_t value)
{
uint32_t slot = hash(key);
while (true)
{
if (hashtable[slot].key == key)
{
hashtable[slot].value = kEmpty;
return;
}
if (hashtable[slot].key == kEmpty)
{
return;
}
slot = (slot + 1) & (kHashTableCapacity - 1);
}
}
ืืืืืืื ื ืึท ืฉืืืกื ืืื ืืขืืื ืืื ืึท ืืืืขืืืืื ืืืขื ืืืขื: ืืืจ ืืึธืื ืืขื ืฉืืืกื ืืื ืื ืืืฉ ืืื ืฆืืืื ืืืึทื ืืืขืจื (ื ืืฉื ืื ืฉืืืกื ืืื) ืืื ืืืืืืง. ืืขื ืงืึธื ืืื ืืืืขืจ ืขื ืืขื ืฆื lookup(), ืึทืืืฅ ืึทื ืืืขื ืึท ืืืืึทืื ืืื ืืขืคึฟืื ืขื ืืืืฃ ืึท ืฉืืืกื, ืขืก ืืืื ืืืื ืืืขืจื ืืืืืืง.
ืืื ืืขืจืืื ื ืืืืื, ืึทืืึธื ืึท ืฉืืืกื ืืื ืืขืฉืจืืื ืฆื ืึท ืฉืคึผืขืืื, ืขืก ืืื ื ืื ืืขืจ ืืจืืืขืจืืขืคืืจื. ืืคืืื ืืืขื ืึทื ืขืืขืืขื ื ืืื ืืืืกืืขืืขืงื ืคืื ืื ืืืฉ, ืืขืจ ืฉืืืกื ืืืืืื ืืื ืคึผืืึทืฅ, ืืืึทื ืืืขืจื ืคืฉืื ืืืขืจื ืืืืืืง. ืืืก ืืืื ื ืื ืืืจ ืืืจืคื ื ืืฉื ื ืืฆื ืื ืืืืืืฉืข ืฉืจืืื ืืคืขืจืืฆืืข ืคืืจ ืื ืกืืื ืืืขืจื, ืืืืื ืก'ืืื ื ืืฉื ืงืืื ืืืืืง ืฆื ืื ืืขืฆืืืืข ืืืขืจื ืืื ืืืืืืง ืืืขืจ ื ืืฉื - ืขืก ืืืขื ื ืื ืืืฅ ืืืขืจื ืืืืืืง.
ืจืขืกืืืื ื ืึท ืืึทืฉ ืืืฉ
ืืืจ ืงืขื ืขื ืืืืฉื ืื ืืจืืืก ืคืื ืึท ืืึทืฉ ืืืฉ ืืืจื ืฉืึทืคึฟื ืึท ืืจืขืกืขืจืข ืืืฉ ืืื ืื ืกืขืจืืื ื ื ืื-ืืืืืืง ืขืืขืืขื ืื ืคืื ืื ืึทืื ืืืฉ ืืื ืขืก. ืืื ืืื ื ืืฉื ืื ืกืืจืืืขื ื ืืขื ืคืึทื ืืงืฉืึทื ืึทืืืื ืืืืึทื ืืื ืืขืืืืื ืฆื ืืึทืืื ืื ืืืกืืขืจ ืงืึธื ืคึผืฉืื. ืืขืจืฆื, ืืื CUDA ืืืืื, ืืึผืจืื ืึทืืึทืงืืืฉืึทื ืืื ืึธืคื ืืืจืืืขืงืึธืื ืืื ืืขืจ ืืึทืืขืืึธืก ืงืึธื ืืื ืืื ืืื ืื CUDA ืงืขืจื.
ืืขืจ ืึทืจืืืงื ืืืฉืจืืืื ืืื ืฆื ืืึธืืืคืืฆืืจื ืึทืืึท ืึท ืฉืืึธืก-ืคึผืจืึธืืขืงืืขื ืืึทืื ืกืืจืืงืืืจ.
ืงืึทืืคึผืขืืืืืืื ืึทืก
ืืื ืื ืืืืื ืคืื ืงืฆืืข ืงืึธื ืกื ืืคึผืึทืฅ gpu_hashtable_insert(), _lookup() ะธ _delete() ืคึผืจืึธืฆืขืก ืืืื ืฉืืืกื-ืืืขืจื ืคึผืึธืจ ืืื ืึท ืฆืืึทื. ืืื ื ืืืขืจืืงืขืจ gpu_hashtable_insert(), _lookup() ะธ _delete() ืคึผืจืึธืฆืขืก ืึท ืืขื ืืข ืคืื โโืคึผืขืจื ืืื ืคึผืึทืจืึทืืขื, ืืขืืขืจ ืคึผืึธืจ ืืื ืึท ืืึทืืื ืืขืจ ืืคึผื ืืืจืืคืืจืื ื ืคืึธืืขื:
// CPU code to invoke the CUDA kernel on the GPU
uint32_t threadblocksize = 1024;
uint32_t gridsize = (numkvs + threadblocksize - 1) / threadblocksize;
gpu_hashtable_insert_kernel<<<gridsize, threadblocksize>>>(hashtable, kvs, numkvs);
// GPU code to process numkvs key/values in parallel
void gpu_hashtable_insert_kernel(KeyValue* hashtable, const KeyValue* kvs, unsigned int numkvs)
{
unsigned int threadid = blockIdx.x*blockDim.x + threadIdx.x;
if (threadid < numkvs)
{
gpu_hashtable_insert(hashtable, kvs[threadid].key, kvs[threadid].value);
}
}
ืื ืฉืืึธืก-ืงืขืื ืฉืืขืืืง ืืึทืฉ ืืืฉ ืฉืืืฆื ืงืึทื ืงืขืจืึทื ื ืื ืกืขืจืฅ, ืืืงืึทืคึผืก ืืื ืืืืืฉืึทื ื. ืืืืึทื ืฉืืืกื-ืืืขืจื ืคึผืขืจื ืืขื ืขื ืฉืืขื ืืืง ืืื ืืืื ืขืจ ืคืื ืคืืจ ืฉืืึทืื ืืื ืฉืืืกืืขื ืืึธื ื ืื ืจืืจื, ืื ืืืฉ ืืขืจืึทื ืืื ืงืขืจืขืงืื ืึทืก ืืคืืื ืืืขื ืคืึทืจืฉืืืขื ืข ืืืืคึผืก ืคืื ืึทืคึผืขืจืืืฉืึทื ื ืืขื ืขื ืืขื ืืฆื ืกืืืืึทืืืืื ืืึทืกืื.
ืึธืืขืจ, ืืืื ืืืจ ืคึผืจืึธืฆืขืก ืึท ืคึผืขืงื ืคืื ืื ืกืขืจืฉืึทื ื ืืื ืืืืืฉืึทื ื ืืื ืคึผืึทืจืึทืืขื, ืืื ืืืื ืื ืึทืจืืึทื ืฉืจืืึทื ืืขื ืืข ืคืื โโืคึผืขืจื ืึผืืื ืืืคึผืืืงืึทื ืฉืืืกืืขื, ืืืจ ืืืขืื ื ืืฉื ืงืขื ืขื ืคืึธืจืืืกืืึธืื ืืืึธืก ืคึผืขืจื ืืืขื "ืืขืืืื ืขื" - ืืืขื ืืืื ืืขืฉืจืืื ืฆื ืื ืืึทืฉ ืืืฉ ืืขืฆืืข. ืืื ืก ืืึธืื ืืืจ ืืขืจืืคื ืื ืื ืกืขืจืฉืึทื ืงืึธื ืืื ืึท ืึทืจืืึทื ืฉืจืืึทื ืืขื ืืข ืคืื โโืคึผืขืจื A/0 B/1 A/2 C/3 A/4. ืืืขื ืื ืงืึธื ืงืึทืืคึผืืืฅ, ืคึผืขืจื B/1 ะธ C/3 ืืขื ืขื ืืขืจืึทื ืืื ืฆื ืืืื ืคืึธืจืฉืืขืื ืืื ืื ืืืฉ, ืึธืืขืจ ืืื ืืขืจ ืืขืืืืงืขืจ ืฆืืื, ืงืืื ืคืื ืื ืคึผืขืจื ืืืขื ืืขืจืฉืืึทื ืขื ืืื ืขืก A/0, A/2 ืึธืืขืจ A/4. ืืึธืก ืงืขื ืึธืืขืจ ืงืขื ื ืืฉื ืืืื ืึท ืคึผืจืึธืืืขื - ืขืก ืึทืืข ืืขืคึผืขื ืืก ืืืืฃ ืื ืึทืคึผืืึทืงืืืฉืึทื. ืืืจ ืงืขื ืืืืกื ืืื ืฉืืืึทืื ืึทื ืขืก ืืขื ืขื ืงืืื ืืืคึผืืืงืึทื ืงืื ืืื ืื ืึทืจืืึทื ืฉืจืืึทื ืืขื ืืข, ืึธืืขืจ ืืืจ ืงืขื ื ืืฉื ืืึธืจืื ืืืึธืก ืืืขืจื ืืื ืืขืฉืจืืื ืืขืฆืืข.
ืืืื ืืึธืก ืืื ืึท ืคึผืจืึธืืืขื ืคึฟืึทืจ ืืืจ, ืืืจ ืืึทืจืคึฟื ืฆื ืฆืขืืืืื ืื ืืืคึผืืืงืึทื ืคึผืขืจื ืืื ืคืึทืจืฉืืืขื ืข CUDA ืกืืกืืขื ืงืึทืืืก. ืืื CUDA, ืืขืืขืจ ืึธืคึผืขืจืึทืฆืืข ืืืึธืก ืจืืคื ืื ืงืขืจื ืืื ืฉืืขื ืืืง ืงืึทืืคึผืืืืื ืืืืืขืจ ืื ืืืืึทืืขืจ ืงืขืจื ืจืืคื (ืืคึผืืืช ืืื ืืืื ืคืึธืืขื. ืืื ืคืึทืจืฉืืืขื ืข ืคึฟืขืืขื, ืงืขืจื ืึทืื ืืขื ืขื ืขืงืกืึทืงืืืืึทื ืืื ืคึผืึทืจืึทืืขื). ืืื ืืขื ืืืึทืฉืคึผืื ืืืืื, ืืืื ืืืจ ืจืืคื ืืืื ืงืขืจื ืืื A/0 B/1 A/2 C/3, ืืื ืื ืื ืืขืจืข ืืื A/4, ืืขืืึธืื ืืขืจ ืฉืืืกื A ืืืขื ืืึทืงืืืขื ืื ืืืขืจื 4.
ืืืฆื ืืึธืื ืก ืจืขืื ืืืขืื ืฆื ืคืึทื ืืงืฉืึทื ื ืืึธื lookup() ะธ delete() ื ืืฆื ืึท ืงืืึธืจ ืึธืืขืจ ืืืึทืืึทืืึทื ืืืึทืื ืฆื ืึท ืืขื ืืข ืคืื โโืคึผืขืจื ืืื ืื ืืึทืฉ ืืืฉ. ืฉืืึทืื ืึทื:
ืืขืจ ืงืึทืืคึผืืืืขืจ ืงืขื ืงืืืึทืื ืฆื ืึทืคึผืืึทืืืื ืืืืขื ืขื ืืื ืฉืจืืืื ืฆื ืืืืืืืข ืึธืืขืจ ืฉืขืจื ืืึผืจืื ... ืื ืึธืคึผืืืืืืึทืืืึธื ืก ืงืขื ืขื ืืืื ืคืึทืจืงืจืืคึผืื ื ืืฆื ืื ืงืืืืขืจื
volatile: ... ืงืืื ืจืขืคึฟืขืจืขื ืฅ ืฆื ืืขื ืืืึทืืขืืืืืง ืืื ืงืึธืืคึผืืืขื ืืื ืึท ืคืึทืงืืืฉ ืืึผืจืื ืืืืขื ืขื ืึธืืขืจ ืฉืจืืึทืื ืืืืขื.
ืงืขืจืขืงืื ืึทืก ืงืึทื ืกืืืขืจืืืฉืึทื ื ืืึธื ื ืื ืืึทืจืคื ืึทืคึผืืึทืงืืืฉืึทื volatile. ืืืื ืืขืจ ืืืจืืคืืจืื ื ืคืึธืืขื ื ืืฆื ืึท ืงืึทืฉื ืืืขืจื ืคืื ืึท ืคืจืืขืจ ืืืืขื ืขื ืึธืคึผืขืจืึทืฆืืข, ืขืก ืืืขื ืืืื ื ืืฆื ืึท ืืืกื ืึทืืืืืืืื ืืื ืคึฟืึธืจืืึทืฆืืข. ืึธืืขืจ ื ืึธื, ืืึธืก ืืื ืืื ืคึฟืึธืจืืึทืฆืืข ืคืื โโืื ืจืืืืืง ืฉืืึทื ืคืื ืื ืืึทืฉ ืืืฉ ืืื ืึท ืืืืขืจ ืืึธืืขื ื ืคืื ืื ืงืขืจื ืจืืคื. ืืืื ืืืจ ืืึทืจืคึฟื ืฆื ื ืืฆื ืื ืืขืฆืืข ืืื ืคึฟืึธืจืืึทืฆืืข, ืืืจ ืงืขื ืขื ื ืืฆื ืืขื ืืื ืืขืงืก volatile, ืึธืืขืจ ืืขืืึธืื ืื ืคืึธืจืฉืืขืืื ื ืืืขื ืคืึทืจืืื ืขืจื ืึท ืืืกื: ืืืื ืืืื ืืขืกืฅ, ืืืขื ืื ืืืกืืขืงื 32 ืืืืืึธื ืขืืขืืขื ืื, ืื ืืืืงืืึทื ืืืงืจืืกื ืคืื 500 ืืืืืึธื ืืืืืฉืึทื ื / ืกืขืง ืฆื 450 ืืืืืึธื ืืืืืฉืึทื ื / ืกืขืง.
ืคืึธืจืฉืืขืืื ื
ืืื ืื ืคึผืจืึธืืข ืคึฟืึทืจ ืื ืกืขืจืืื ื 64 ืืืืืึธื ืขืืขืืขื ืื ืืื ืืืกืืขืงื 32 ืืืืืึธื ืคืื ืืื, ืคืึทืจืืขืกื ืฆืืืืฉื std::unordered_map ืืื ืขืก ืืื ืืืขื ืงืืื ืืึทืฉ ืืืฉ ืคึฟืึทืจ ืื ืืคึผื:
std::unordered_map ืคืืจืืจืืื 70 ืืื ืื ืกืขืจืืื ื ืืื ืจืืืืืืื ื ืขืืขืืขื ืื ืืื ืืึทื ืคืจืื ืืื unordered_map (ืืึทืงืืืขื ืืึทืคืจืืึทืขื ืคืื ืืืืืึทื ื ืคืื ืขืืขืืขื ืื ื ืขืื ืึท ืคึผืืึทืฅ ืคืื ืฆืืื, ืืืืึทื ืื unordered_map ืงืืืคื ืืึผืจืื ืึทืืึทืงืืืฉืึทื ื ืืขื ืขื ืืขืืืื). ืขืจืืขื ืืขืจืขืื, std:unordered_map ืืึธืจ ืคืึทืจืฉืืืขื ืข ืจืืกืืจืืงืฉืึทื ื. ืขืก ืืื ืึท ืืืื ืงืคึผื ืคืึธืืขื ืคืื ืืืจืืคืืจืื ื, ืฉืืืฆื ืฉืืืกื ืืืึทืืืขืก ืคืื ืงืืื ืืจืืืก, ืคึผืขืจืคืึธืจืื ืืขืืื ื ืืื ืืืื ืืืืึทืืึทืืืืฉืึทื ืจืึทืืขืก ืืื ืืืืึทืื ืกืืึทืืื ืคืึธืจืฉืืขืืื ื ื ืึธื ืงืืืคื ืืืืืฉืึทื ื.
ืืขืจ ืืขืืืืขืจ ืคืื ืื ืืึทืฉ ืืืฉ ืคึฟืึทืจ ืื ืืคึผื ืืื ืื ืืขืจ-ืคึผืจืึธืืจืึทื ืงืึธืืื ืืงืึทืฆืืข ืืื ืืขืืืขื 984 ืืื. ืืึธืก ืืืื ืื ืฆืืื ืืืึธืก ืืื ืคืืจืืจืืื ืฆื ืฉืืขืื ืื ืืืฉ ืืื ืืึผืจืื ืืื ืืืกืืขืงื ืขืก (ืึทืืึธืงืึทืืื ื 1 ืืืืืืืื ืคืื ืืึผืจืื ืืืื ืืึธื, ืืืึธืก ื ืขืื ืขืืืขืืข ืืึธื ืืื CUDA), ืื ืกืขืจืืื ื ืืื ืืืกืืขืงื ืขืืขืืขื ืื ืืื ืืืขืจืืืืื ื ืืืืขืจ ืืื. ืึทืืข ืงืืคืืขืก ืฆื ืืื ืคืื ืื ืืืืืขื ืงืึทืจืื ืืึผืจืื ืืขื ืขื ืืืื ืืขื ืืืขื ืืื ืืฉืืื.
ืื ืืึทืฉ ืืืฉ ืืื ืืขื ืืืขื 271 ืืื ืฆื ืคืึทืจืขื ืืืงื. ืืึธืก ืืืื ืื ืฆืืื ืคืืจืืจืืื ืืืจื ืื ืืืืืขื ืงืึทืจืื ืื ืกืขืจืืื ื ืืื ืืืกืืขืงื ืขืืขืืขื ืื, ืืื ืืื ื ืืฉื ื ืขืืขื ืืื ืืฉืืื ืื ืฆืืื ืคืืจืืจืืื ืงืึทืคึผืืื ื ืืื ืืึผืจืื ืืื ืืืขืจืืืืื ื ืืืืขืจ ืื ืจืืืึทืืืื ื ืืืฉ. ืืืื ืื ืืคึผื ืืืฉ ืืขืื ืคึฟืึทืจ ืึท ืืึทื ื ืฆืืึทื, ืึธืืขืจ ืืืื ืื ืืึทืฉ ืืืฉ ืืื ืงืึทื ืืืื ื ืืื ืื ืืืงืึธืจื ืคืื ืื ืืืืืขื ืงืึทืจืื (ืืืฉื, ืฆื ืฉืึทืคึฟื ืึท ืืึทืฉ ืืืฉ ืืืึธืก ืืืขื ืืืื ืืขืืืืื ื ืืืจื ืื ืืขืจืข ืืคึผื ืงืึธื ืืื ื ืืฉื ืื ืืืืคื ืคึผืจืึทืกืขืกืขืจ), ืืขืืึธืื ืื ืคึผืจืึธืืข ืจืขืืืืืึทื ืืื ืืึทืืืึทืืืง.
ืื ืืึทืฉ ืืืฉ ืคึฟืึทืจ ืึท ืืืืืขื ืงืึทืจืื ืืขืืึทื ืกืืจืืืฅ ืืืื ืคืึธืจืฉืืขืืื ื ืจืขืื ืฆื ืืืื ืืจืืคึผืื ืืื ืึทืงืืืื ืคึผืึทืจืึทืืขืืืืึทืืืึธื.
ืืืืืืืืฉืึทื ื
ืื ืืึทืฉ ืืืฉ ืึทืจืงืึทืืขืงืืฉืขืจ ืืื ืึท ืืืกื ืืฉืื ืฆื ืืืื ืึทืืืขืจ ืคืื:
- ืืื ืขืึทืจ ืคึผืจืึธืืืื ื ืืื ืืึทืืคึผืขืจื ืืืจื ืงืืึทืกืืขืจืื ื, ืืืึธืก ื ืื ืฉืืืกืืขื ืืื ืื ืืืฉ ืฆื ืืืื ืืขืฉืืขืื ืืืืื ืืงืขืจ ืืื ืืืฉืืืืืขืก.
- ืงืื ืืขื ืขื ื ืืฉื ืึทืืืขืงืืขื ืืืขื ื ืืฆื ืื ืคึฟืื ืงืฆืืข
deleteืืื ืืื ืืขืจ ืฆืฒึทื ืฆืขืฉืคึผืจืืืื ืืื ืืขื ืืืฉ.
ืืื ืึท ืจืขืืืืืึทื, ืื ืคืึธืจืฉืืขืืื ื ืคืื ืึท ืืึทืฉ ืืืฉ ืงืขื ืขื ืืืกืืขืืืืืึทื ืืืืจืืื, ืกืคึผืขืฆืืขื ืืืื ืขืก ืืืืืกืฅ ืคึฟืึทืจ ืึท ืืึทื ื ืฆืืึทื ืืื ืืื ืคืืืข ืื ืกืขืจืฅ ืืื ืืืืืฅ. ืืืื ืืืขื ืฆื ืคืึทืจืืื ืขืจื ืื ืืืกืึทืืืืึทื ืืืืืฉืื ืืื ืฆื ืจืืืึทืคึผื ืืื ืึท ื ืืึทืข ืืืฉ ืืื ืึท ืืึทื ืฅ ื ืืืขืจืืง ืืืืึทืืึทืืืืฉืึทื ืงืืจืก ืืื ืคืืืืขืจ ืืืืก ืื ืึทืืืขืงืืขื ืืืขื ืฉืืืกืืขื ืืขืฉืึทืก ืื ืจืืืึทืฉืื ื.
ืฆื ืืืืืกืืจืืจื ืื ืืืกืงืจืืืื ืืฉืื, ืืื ืืืขื ื ืืฆื ืื ืืืืื ืงืึธื ืฆื ืฉืึทืคึฟื ืึท ืืืฉ ืืื 128 ืืืืืึธื ืขืืขืืขื ืื ืืื ืฉืืืืฃ ืืืจื 4 ืืืืืึธื ืขืืขืืขื ืื ืืื ืืื ืืึธืื ืึธื ืืขืคืืื 124 ืืืืืึธื ืกืืึธืฅ (ืืืืึทืืึทืืืืฉืึทื ืงืืจืก ืคืื ืืืขืื 0,96). ืืึธ ืืื ืืขืจ ืจืขืืืืืึทื ืืืฉ, ืืขืืขืจ ืจืืืขืจื ืืื ืึท CUDA ืงืขืจื ืจืืคื ืฆื ืึทืจืืึทื ืืืืื 4 ืืืืืึธื ื ืืึท ืขืืขืืขื ืื ืืื ืืืื ืืึทืฉ ืืืฉ:
ืืึทื ืืฅ ืงืืจืก
ืื ืกืขืจืฉืึทื ืืขืืืืขืจ 4 ืขืืขืืขื ืื
0,00
11,608448 ืืก (361,314798 ืืืืืึธื ืงืื / ืกืขืง.)
0,03
11,751424 ืืก (356,918799 ืืืืืึธื ืงืื / ืกืขืง.)
0,06
11,942592 ืืก (351,205515 ืืืืืึธื ืงืื / ืกืขืง.)
0,09
12,081120 ืืก (347,178429 ืืืืืึธื ืงืื / ืกืขืง.)
0,12
12,242560 ืืก (342,600233 ืืืืืึธื ืงืื / ืกืขืง.)
0,16
12,396448 ืืก (338,347235 ืืืืืึธื ืงืื / ืกืขืง.)
0,19
12,533024 ืืก (334,660176 ืืืืืึธื ืงืื / ืกืขืง.)
0,22
12,703328 ืืก (330,173626 ืืืืืึธื ืงืื / ืกืขืง.)
0,25
12,884512 ืืก (325,530693 ืืืืืึธื ืงืื / ืกืขืง.)
0,28
13,033472 ืืก (321,810182 ืืืืืึธื ืงืื / ืกืขืง.)
0,31
13,239296 ืืก (316,807174 ืืืืืึธื ืงืื / ืกืขืง.)
0,34
13,392448 ืืก (313,184256 ืืืืืึธื ืงืื / ืกืขืง.)
0,37
13,624000 ืืก (307,861434 ืืืืืึธื ืงืื / ืกืขืง.)
0,41
13,875520 ืืก (302,280855 ืืืืืึธื ืงืื / ืกืขืง.)
0,44
14,126528 ืืก (296,909756 ืืืืืึธื ืงืื / ืกืขืง.)
0,47
14,399328 ืืก (291,284699 ืืืืืึธื ืงืื / ืกืขืง.)
0,50
14,690304 ืืก (285,515123 ืืืืืึธื ืงืื / ืกืขืง.)
0,53
15,039136 ืืก (278,892623 ืืืืืึธื ืงืื / ืกืขืง.)
0,56
15,478656 ืืก (270,973402 ืืืืืึธื ืงืื / ืกืขืง.)
0,59
15,985664 ืืก (262,379092 ืืืืืึธื ืงืื / ืกืขืง.)
0,62
16,668673 ืืก (251,627968 ืืืืืึธื ืงืื / ืกืขืง.)
0,66
17,587200 ืืก (238,486174 ืืืืืึธื ืงืื / ืกืขืง.)
0,69
18,690048 ืืก (224,413765 ืืืืืึธื ืงืื / ืกืขืง.)
0,72
20,278816 ืืก (206,831789 ืืืืืึธื ืงืื / ืกืขืง.)
0,75
22,545408 ืืก (186,038058 ืืืืืึธื ืงืื / ืกืขืง.)
0,78
26,053312 ืืก (160,989275 ืืืืืึธื ืงืื / ืกืขืง.)
0,81
31,895008 ืืก (131,503463 ืืืืืึธื ืงืื / ืกืขืง.)
0,84
42,103294 ืืก (99,619378 ืืืืืึธื ืงืื / ืกืขืง.)
0,87
61,849056 ืืก (67,815164 ืืืืืึธื ืงืื / ืกืขืง.)
0,90
105,695999 ืืก (39,682713 ืืืืืึธื ืงืื / ืกืขืง.)
0,94
240,204636 ืืก (17,461378 ืืืืืึธื ืงืื / ืกืขืง.)
ืืื ืืืืึทืืึทืืืืฉืึทื ืื ืงืจืืกืื, ืคืึธืจืฉืืขืืื ื ืืืงืจืืกืึทื. ืืึธืก ืืื ื ืืฉื ืืืืืืจืึทืืึทื ืืื ืจืืึฟ ืงืึทืกืขืก. ืืืื ืึท ืึทืคึผืืึทืงืืืฉืึทื ืื ืกืขืจืฅ ืขืืขืืขื ืื ืืื ืึท ืืืฉ ืืื ืืขืืึธืื ืึทืืืขืงืืืึทืจืคื ืืื (ืืืฉื, ืืืขื ืงืึทืื ืืื ื ืืืขืจืืขืจ ืืื ืึท ืืื), ืืึธืก ืืื ื ืืฉื ืึท ืคึผืจืึธืืืขื. ืึธืืขืจ ืืืื ืื ืึทืคึผืืึทืงืืืฉืึทื ื ืืฆื ืึท ืืึทื ื-ืืขืืขืื ืืึทืฉ ืืืฉ (ืืืฉื, ืืื ืึท ืืจืึทืคืืงืก ืจืขืืึทืงืืึธืจ ืฆื ืงืจืึธื ื ืื-ืืืืืืง ืคึผืึทืจืฅ ืคืื ืืืืืขืจ ืืื ืืขืจ ืืึทื ืืฆืขืจ ืึธืคื ืื ืกืขืจืฅ ืืื ืืืกืืขืงื ืืื ืคึฟืึธืจืืึทืฆืืข), ืืขื ื ืึทืืืจ ืงืขื ืืืื ืคึผืจืึธืืืขืืึทืืืง.
ืืื ืืขืืืกืื ืื ืืึทืฉ ืืืฉ ืคึผืจืึธืืืื ื ืืืคืงืืึทื ื ืึธื 64 ืืืืืึธื ืื ืกืขืจืฅ (ืืืืึทืืึทืืืืฉืึทื ืคืึทืงืืึธืจ 0,5). ืื ืืืจืืฉื ืืืืขื ืืืคืงืืึทื ืืื ืืขืืืขื 0,4774, ืึทืืื ืจืืึฟ ืฉืืืกืืขื ืืขื ืขื ืืขืืืขื ืืื ืืขืจ ืืขืกืืขืจ ืืขืืืขื ืฉืคึผืขืืื ืึธืืขืจ ืืืื ืฉืคึผืขืืื ืึทืืืขืง ืคืื ืืขืจ ืืขืกืืขืจ ืฉืืขืืข. ืื ืืึทืงืกืืืื ืกืึทืื ืืื ื ืืืคืงืืึทื ืืื ืืขืืืขื 60.
ืืื ืืขืืึธืื ืืขืืืกืื ืื ืคึผืจืึธืืืื ื ืืืคืงืืึทื ืืืืฃ ืึท ืืืฉ ืืื 124 ืืืืืึธื ืื ืกืขืจืฅ (ืืืืึทืืึทืืืืฉืึทื ืคืึทืงืืึธืจ 0,97). ืื ืืืจืืฉื ืืืืขื ืืืคืงืืึทื ืืื ืฉืืื 10,1757, ืืื ืื ืืึทืงืกืืืื - 6474 (!!). ืืื ืขืึทืจ ืกืขื ืกืื ื ืคืึธืจืฉืืขืืื ื ืืจืืคื ืก ืืืืืืืืง ืืื ืืืื ืืืืึทืืึทืืืืฉืึทื ืจืืืฅ.
ืขืก ืืื ืืขืกืืขืจ ืฆื ืืึทืืื ืื ืืืืึทืืึทืืืืฉืึทื ืงืืจืก ืคืื ืืขื ืืึทืฉ ืืืฉ ื ืืืขืจืืง. ืึธืืขืจ ืืขืืึธืื ืืืจ ืคืึทืจืืจืขืกืขืจื ืคืึธืจืฉืืขืืื ื ืืื ืื ืงืึธืกื ืคืื ืืึผืจืื ืงืึทื ืกืึทืืฉืึทื. ืฆืื ืืืืง, ืืื ืื ืคืึทื ืคืื 32-ืืืกื ืฉืืืกืืขื ืืื ืืืึทืืืขืก, ืืึธืก ืงืขื ืืืื ืืขืจืขืืืคืืจืืืงื. ืืืื ืืื ืื ืืืืื ืืืึทืฉืคึผืื, ืืื ืึท ืืืฉ ืืื 128 ืืืืืึธื ืขืืขืืขื ืื, ืืืจ ืืึทืืื ืื ืืืืึทืืึทืืืืฉืึทื ืคืึทืงืืึธืจ ืคืื 0,25, ืืืจ ืงืขื ืขื ืฉืืขืื ื ืื ืืขืจ ืืื 32 ืืืืืึธื ืขืืขืืขื ืื ืืื ืขืก, ืืื ืื ืจืืขื 96 ืืืืืึธื ืกืืึธืฅ ืืืขื ืืืื ืคืึทืจืคืึทืื - 8 ืืืืขืก ืคึฟืึทืจ ืืขืืขืจ ืคึผืึธืจ , 768 ืื ืคืื ืคืึทืจืคืึทืื ืืึผืจืื.
ืืืืข ืืึธื ืึทื ืืืจ ืืขื ืขื ืืขืจืขืื ืืืขืื ืื ืึธื ืืืขืจ ืคืื ืืืืืขื ืงืึทืจืื ืืึผืจืื, ืืืึธืก ืืื ืึท ืืขืจ ืืืขืจืืคืื ืืืื ืืื ืกืืกืืขื ืืึผืจืื. ืืึธืืฉ ืจืืึฟ ืืึธืืขืจื ืืขืกืงืืึทืคึผ ืืจืึทืคืืงืก ืงืึทืจืืก ืืืึธืก ืฉืืืฆื CUDA ืืึธืื ืืคึผืืืช 4 ืืืืืืืื ืคืื ืืึผืจืื (ืืื ืืขืจ ืฆืืื ืคืื ืฉืจืืืื, ืื NVIDIA 2080 ืื ืืื 11 ืืืืืืืื), ืขืก ืืืึธืื ื ืึธื ื ืืฉื ืืืื ืื ืืืืืืึทืกื ืืึทืฉืืืก ืฆื ืคืึทืจืืืจื ืึทืืึท ืึทืืึทืื ืฅ.
ืฉืคึผืขืืขืจ ืืื ืืืขื ืฉืจืืึทืื ืืขืจ ืืืขืื ืงืจืืืืืื ื ืืึทืฉ ืืืฉื ืคึฟืึทืจ ืืืืืขื ืงืึทืจืืก ืืืึธืก ืืึธื ื ืื ืืึธืื ืคึผืจืึธืืืขืืก ืืื ืคึผืจืึธืืืื ื ืืืคืงืืึทื, ืืื ืืขืืื ื ืืื ืืืขืื ืฆื ืจืืืื ืืืืกืืขืืขืงื ืกืืึธืฅ.
ืกืึธืื ืืื ื ืืืฃ ืืขืืฉืขืจืืึทื ื
ืฆื ืืึทืฉืืืืขื ืื ืคึผืจืึธืืืื ื ืืืคืงืืึทื ืคืื ืึท ืฉืืืกื, ืืืจ ืงืขื ืขื ืขืงืกืืจืึทืงื ืื ืืึทืฉ ืคืื ืื ืฉืืืกื (ืืืึทื ืืืขืึทื ืืืฉ ืืื ืืขืงืก) ืคืื ืืืื ืคืึทืงืืืฉ ืืืฉ ืืื ืืขืงืก:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
ืืืืึทื ืคืื ืื ืืึทืืืฉ ืคืื ืฆืืืื ืก ืฆืืืื ืก ืืขืจืืึทื ื ืืืื ืขืจื ื ืืืขืจื ืืื ืื ืคืึทืงื ืึทื ืื ืงืึทืคึผืึทืฆืืืขื ืคืื ืื ืืึทืฉ ืืืฉ ืืื ืฆืืืื ืฆื ืื ืืึทืื ืคืื ืฆืืืื, ืืขื ืฆืืืึทื ื ืืืขื ืึทืจืืขืื ืืคืืื ืืืขื ืื ืฉืืืกื ืืื ืืขืงืก ืืื ืืจืืืขืจืืขืคืืจื ืฆื ืื ืึธื ืืืื ืคืื ืื ืืืฉ. ืืึธืืืจ ื ืขืืขื ืึท ืฉืืืกื ืืืึธืก ืืึทืฉืขื ืฆื 1, ืึธืืขืจ ืืื ืื ืกืขืจืืึทื ืืื ืฉืคึผืขืืื 3. ืืขืจื ืึธื ืคึฟืึทืจ ืึท ืืืฉ ืืื ืงืึทืคึผืึทืฆืืืขื 4 ืืืจ ืืึทืงืืืขื (3 โ 1) & 3, ืืืึธืก ืืื ืขืงืืืืืืึทืืขื ื ืฆื 2.
ืกืึธืฃ
ืืืื ืืืจ ืืึธื ืคึฟืจืืื ืึธืืขืจ ืืึทืืขืจืงืื ืืขื, ืืืืข ืืืืฆืคึผืึธืกื ืืืจ ืืืืฃ ืึธืืขืจ ืขืคึฟืขื ืขื ืึท ื ืืึทืข ืืขืืข ืืื.
ืืขืจ ืงืึธื ืืื ืืขืฉืจืืื ืืื ืืขืจ ืื ืกืคึผืืจืึทืฆืืข ืคืื โโืืืกืืขืฆืืืื ื ืึทืจืืืงืืขื:
ืืื ืืขืจ ืฆืืงืื ืคึฟื, ืืื ืืืขื ืคืึธืจืืขืฆื ืฆื ืฉืจืืึทืื ืืืขืื ืืึทืฉ ืืืฉ ืืืคึผืืึทืืึทื ืฅ ืคึฟืึทืจ ืืืืืขื ืงืึทืจืืก ืืื ืึทื ืึทืืืื ืืืืขืจ ืคืึธืจืฉืืขืืื ื. ืืืึทื ืคึผืืึทื ื ืึทืจืืึทื ื ืขืืขื ืืฉืึทืื ืื ื, ืจืึธืืื ืืื ืืึทืฉืื ื ืืื ืงืืงื ืืึทืฉืื ื ื ืืฆื ืึทืืึธืืืฉืข ืึทืคึผืขืจืืืฉืึทื ื ืืื ืืึทืื ืกืืจืึทืงืืฉืขืจื ืืืึธืก ืืขื ืขื ืืคึผื ืคืจืืึทื ืืืขื.
ืืงืืจ: www.habr.com