| 59 | // Tests if data needs further reduction, including across block boundaries |
| 60 | template<typename Tk> |
| 61 | __global__ void test_needs_reduction(int *needs_another_reduction, |
| 62 | int *needs_block_boundary_reduced, |
| 63 | CParam<Tk> keys_in, const int n) { |
| 64 | const int tid = blockIdx.x * blockDim.x + threadIdx.x; |
| 65 | Tk k; |
| 66 | |
| 67 | if (tid < n) { k = keys_in.ptr[tid]; } |
| 68 | |
| 69 | int update_key = (k == shfl_down_sync(k, 1)) && |
| 70 | (tid < (n - 1)) && ((threadIdx.x % 32) < 31); |
| 71 | int remaining_updates = any_sync(update_key); |
| 72 | |
| 73 | __syncthreads(); |
| 74 | |
| 75 | if (remaining_updates && (threadIdx.x % 32 == 0)) |
| 76 | atomicOr(needs_another_reduction, remaining_updates); |
| 77 | |
| 78 | // check across warp boundaries |
| 79 | update_key = |
| 80 | (((threadIdx.x % 32) == 31) // last thread in warp |
| 81 | && (threadIdx.x < (blockDim.x - 1)) // not last thread in block |
| 82 | // next value valid and equal |
| 83 | && ((tid + 1) < n) && (k == keys_in.ptr[tid + 1])); |
| 84 | remaining_updates = any_sync(update_key); |
| 85 | |
| 86 | // TODO: single per warp? change to assignment rather than atomicOr |
| 87 | if (remaining_updates) atomicOr(needs_another_reduction, remaining_updates); |
| 88 | |
| 89 | // last thread in each block checks if any inter-block keys need further |
| 90 | // reduction |
| 91 | if (tid == ((blockIdx.x + 1) * blockDim.x) - 1 && |
| 92 | blockIdx.x < gridDim.x - 1) { |
| 93 | int k0 = keys_in.ptr[tid]; |
| 94 | int k1 = keys_in.ptr[tid + 1]; |
| 95 | if (k0 == k1) { atomicOr(needs_block_boundary_reduced, 1); } |
| 96 | } |
| 97 | } |
| 98 | |
| 99 | // Compacts "incomplete" block-sized chunks of data in global memory |
| 100 | template<typename Tk, typename To> |
nothing calls this directly
no test coverage detected