...
 
Commits (8)
......@@ -143,7 +143,7 @@ public:
invert_assign(t...);
}
__device__ __host__ inline grid_key_dx<dim,index_type> move(int i, int m)
__device__ __host__ inline grid_key_dx<dim,index_type> move(int i, int m) const
{
grid_key_dx<dim,index_type> tmp = *this;
......@@ -336,6 +336,33 @@ public:
return !this->operator==(key_t);
}
/*! \brief Check order of two keys
*
* \param key_t key to check
*
* \return true if this is lexicographically less than other key
*
*/
bool operator<(const grid_key_dx<dim> & key_t) const
{
// Check the two key index by index
for (long int i = dim-1 ; i >= 0; --i)
{
if (k[i] < key_t.k[i])
{
return true;
}
else if (k[i] > key_t.k[i])
{
return false;
}
}
// identical key
return false;
}
/*! \brief set the Key from a list of numbers
*
......
......@@ -31,12 +31,12 @@ struct cta_segscan_t {
int cta_mask = 0x7fffffff>> (31 - lane); // exclusive search.
// Build a head flag bitfield and store it into shared memory.
int warp_bits = __ballot(has_head_flag);
int warp_bits = __ballot_sync(0xFFFFFFFF,has_head_flag);
storage.delta[warp] = warp_bits;
__syncthreads();
if(tid < num_warps) {
int cta_bits = __ballot(0 != storage.delta[tid]);
int cta_bits = __ballot_sync(0xFFFFFFFF,0 != storage.delta[tid]);
int warp_segment = 31 - clz(cta_mask & cta_bits);
int start = (-1 != warp_segment) ?
(31 - clz(storage.delta[warp_segment]) + 32 * warp_segment) : 0;
......