Detailed description
template<typename KeyT, typename HashT, uint32 CTA_SIZE, uint32 TABLE_SIZE, KeyT INVALID_KEY = 0xFFFFFFFF>
struct cugar::cuda::BlockHashMap< KeyT, HashT, CTA_SIZE, TABLE_SIZE, INVALID_KEY >
This class implements a cuda block-wide Hash Set, allowing arbitrary threads from the same CTA to add new entries at the same time.
__global__ void kernel()
{
typedef BlockHashMap<uint32,uint32,128,512> hash_map_type;
// allocate the temporary storage of the hash map in shared memory
__shared__ typename hash_map_type::TempStorage hash_map_storage;
hash_map_type hash_map( hash_map_storage );
hash_map.insert( threadIdx.x/2, hash(threadIdx.x/2) );
// allocate some storage for the hash map values
__shared__ uint32 hash_map_values[512];
// initialize the hash map values to zero
for (uint32 i = threadIdx.x; i < 512; i += 128)
hash_map_values[i] = 0;
__syncthreads();
// add 1 to each entry sharing the same slot in the hash map
const uint32 slot = hash_map.find( threadIdx.x/2 );
atomic_add( &hash_map_values[ slot ], 1u );
}
#include <hash.h>
Inheritance diagram for cugar::cuda::BlockHashMap< KeyT, HashT, CTA_SIZE, TABLE_SIZE, INVALID_KEY >:
Classes | |
| struct | TempStorage |
Public Methods | |
| CUGAR_DEVICE | BlockHashMap () |
| CUGAR_DEVICE | BlockHashMap (TempStorage &_storage) |
 Public Methods inherited from cugar::cuda::HashMap< KeyT, HashT, INVALID_KEY > | |
| CUGAR_DEVICE | HashMap () |
| CUGAR_DEVICE | HashMap (const uint32 _table_size, KeyT *_hash, KeyT *_unique, uint32 *_slots, uint32 *_count) |
| CUGAR_DEVICE void | insert (const KeyT key, const HashT hash_code) |
| CUGAR_DEVICE uint32 | find (const KeyT key, const HashT hash_code) |
| CUGAR_DEVICE uint32 | size () const |
| CUGAR_DEVICE KeyT | get_unique (const uint32 i) const |
Additional Inherited Members | |
| Public Members inherited from cugar::cuda::HashMap< KeyT, HashT, INVALID_KEY > | |
| uint32 | table_size |
| KeyT * | hash |
| KeyT * | unique |
| uint32 * | slots |
| uint32 * | count |
Constructor & Destructor Documentation
◆ BlockHashMap() [1/2]
template<typename KeyT , typename HashT , uint32 CTA_SIZE, uint32 TABLE_SIZE, KeyT INVALID_KEY = 0xFFFFFFFF>
|
inline |
empty constructor
◆ BlockHashMap() [2/2]
template<typename KeyT , typename HashT , uint32 CTA_SIZE, uint32 TABLE_SIZE, KeyT INVALID_KEY = 0xFFFFFFFF>
|
inline |
constructor
- Parameters
-
_storage the per-CTA storage backing this container
The documentation for this struct was generated from the following file:
- C:/p4research/research/jpantaleoni/Fermat-Public/contrib/cugar/basic/cuda/hash.h
