/* * Copyright 2016 Google Inc. * * Use of this source code is governed by a BSD-style license that can * be found in the LICENSE file. * */ // // // #ifdef __cplusplus extern "C" { #endif #include "common/cuda/assert_cuda.h" #include "common/macros.h" #include "common/util.h" #ifdef __cplusplus } #endif // // We want concurrent kernel execution to occur in a few places. // // The summary is: // // 1) If necessary, some max valued keys are written to the end of // the vin/vout buffers. // // 2) Blocks of slabs of keys are sorted. // // 3) If necesary, the blocks of slabs are merged until complete. // // 4) If requested, the slabs will be converted from slab ordering // to linear ordering. // // Below is the general "happens-before" relationship between HotSort // compute kernels. // // Note the diagram assumes vin and vout are different buffers. If // they're not, then the first merge doesn't include the pad_vout // event in the wait list. // // +----------+ +---------+ // | pad_vout | | pad_vin | // +----+-----+ +----+----+ // | | // | WAITFOR(pad_vin) // | | // | +-----v-----+ // | | | // | +----v----+ +----v----+ // | | bs_full | | bs_frac | // | +----+----+ +----+----+ // | | | // | +-----v-----+ // | | // | +------NO------JUST ONE BLOCK? // | / | // |/ YES // + | // | v // | END_WITH_EVENTS(bs_full,bs_frac) // | // | // WAITFOR(pad_vout,bs_full,bs_frac) >>> first iteration of loop <<< // | // | // +-----------<------------+ // | | // +-----v-----+ | // | | | // +----v----+ +----v----+ | // | fm_full | | fm_frac | | // +----+----+ +----+----+ | // | | ^ // +-----v-----+ | // | | // WAITFOR(fm_full,fm_frac) | // | | // v | // +--v--+ WAITFOR(bc) // | hm | | // +-----+ | // | | // WAITFOR(hm) | // | ^ // +--v--+ | // | bc | | // +-----+ | // | | // v | // MERGING COMPLETE?-------NO------+ // | // YES // | // v // END_WITH_EVENTS(bc) // // // NOTE: CUDA streams are in-order so a dependency isn't required for // kernels launched on the same stream. // // This is actually a more subtle problem than it appears. // // We'll take a different approach and declare the "happens before" // kernel relationships: // // concurrent (pad_vin,pad_vout) -> (pad_vin) happens_before (bs_full,bs_frac) // (pad_vout) happens_before (fm_full,fm_frac) // // concurrent (bs_full,bs_frac) -> (bs_full) happens_before (fm_full,fm_frac) // (bs_frac) happens_before (fm_full,fm_frac) // // concurrent (fm_full,fm_frac) -> (fm_full) happens_before (hm) // (fm_frac) happens_before (hm) // // concurrent (fm_full,fm_frac) -> (fm_full) happens_before (hm) // (fm_frac) happens_before (hm) // // launch (hm) -> (hm) happens_before (hm) // (hm) happens_before (bc) // // launch (bc) -> (bc) happens_before (fm_full,fm_frac) // // // We can go ahead and permanently map kernel launches to our 3 // streams. As an optimization, we'll dynamically assign each kernel // to the lowest available stream. This transforms the problem into // one that considers streams happening before streams -- which // kernels are involved doesn't matter. // // STREAM0 STREAM1 STREAM2 // ------- ------- ------- // // pad_vin pad_vout (pad_vin) happens_before (bs_full,bs_frac) // (pad_vout) happens_before (fm_full,fm_frac) // // bs_full bs_frac (bs_full) happens_before (fm_full,fm_frac) // (bs_frac) happens_before (fm_full,fm_frac) // // fm_full fm_frac (fm_full) happens_before (hm or bc) // (fm_frac) happens_before (hm or bc) // // hm (hm) happens_before (hm or bc) // // bc (bc) happens_before (fm_full,fm_frac) // // A single final kernel will always complete on stream 0. // // This simplifies reasoning about concurrency that's downstream of // hs_cuda_sort(). // typedef void (*hs_kernel_offset_bs_pfn)(HS_KEY_TYPE * const HS_RESTRICT vout, HS_KEY_TYPE const * const HS_RESTRICT vin, uint32_t const slab_offset); static hs_kernel_offset_bs_pfn const hs_kernels_offset_bs[] { #if HS_BS_SLABS_LOG2_RU >= 1 hs_kernel_bs_0, #endif #if HS_BS_SLABS_LOG2_RU >= 2 hs_kernel_bs_1, #endif #if HS_BS_SLABS_LOG2_RU >= 3 hs_kernel_bs_2, #endif #if HS_BS_SLABS_LOG2_RU >= 4 hs_kernel_bs_3, #endif #if HS_BS_SLABS_LOG2_RU >= 5 hs_kernel_bs_4, #endif #if HS_BS_SLABS_LOG2_RU >= 6 hs_kernel_bs_5, #endif #if HS_BS_SLABS_LOG2_RU >= 7 hs_kernel_bs_6, #endif #if HS_BS_SLABS_LOG2_RU >= 8 hs_kernel_bs_7, #endif }; // // // typedef void (*hs_kernel_bc_pfn)(HS_KEY_TYPE * const HS_RESTRICT vout); static hs_kernel_bc_pfn const hs_kernels_bc[] { hs_kernel_bc_0, #if HS_BC_SLABS_LOG2_MAX >= 1 hs_kernel_bc_1, #endif #if HS_BC_SLABS_LOG2_MAX >= 2 hs_kernel_bc_2, #endif #if HS_BC_SLABS_LOG2_MAX >= 3 hs_kernel_bc_3, #endif #if HS_BC_SLABS_LOG2_MAX >= 4 hs_kernel_bc_4, #endif #if HS_BC_SLABS_LOG2_MAX >= 5 hs_kernel_bc_5, #endif #if HS_BC_SLABS_LOG2_MAX >= 6 hs_kernel_bc_6, #endif #if HS_BC_SLABS_LOG2_MAX >= 7 hs_kernel_bc_7, #endif #if HS_BC_SLABS_LOG2_MAX >= 8 hs_kernel_bc_8, #endif }; // // // typedef void (*hs_kernel_hm_pfn)(HS_KEY_TYPE * const HS_RESTRICT vout); static hs_kernel_hm_pfn const hs_kernels_hm[] { #if (HS_HM_SCALE_MIN == 0) hs_kernel_hm_0, #endif #if (HS_HM_SCALE_MIN <= 1) && (1 <= HS_HM_SCALE_MAX) hs_kernel_hm_1, #endif #if (HS_HM_SCALE_MIN <= 2) && (2 <= HS_HM_SCALE_MAX) hs_kernel_hm_2, #endif }; // // // typedef void (*hs_kernel_fm_pfn)(HS_KEY_TYPE * const HS_RESTRICT vout); static hs_kernel_fm_pfn const hs_kernels_fm[] { #if (HS_FM_SCALE_MIN == 0) #if (HS_BS_SLABS_LOG2_RU == 1) hs_kernel_fm_0_0, #endif #if (HS_BS_SLABS_LOG2_RU == 2) hs_kernel_fm_0_1, #endif #if (HS_BS_SLABS_LOG2_RU == 3) hs_kernel_fm_0_2, #endif #if (HS_BS_SLABS_LOG2_RU == 4) hs_kernel_fm_0_3, #endif #if (HS_BS_SLABS_LOG2_RU == 5) hs_kernel_fm_0_4, #endif #if (HS_BS_SLABS_LOG2_RU == 6) hs_kernel_fm_0_5, #endif #if (HS_BS_SLABS_LOG2_RU == 7) hs_kernel_fm_0_6, #endif #endif #if (HS_FM_SCALE_MIN <= 1) && (1 <= HS_FM_SCALE_MAX) CONCAT_MACRO(hs_kernel_fm_1_,HS_BS_SLABS_LOG2_RU) #endif #if (HS_FM_SCALE_MIN <= 2) && (2 <= HS_FM_SCALE_MAX) #if (HS_BS_SLABS_LOG2_RU == 1) hs_kernel_fm_2_2, #endif #if (HS_BS_SLABS_LOG2_RU == 2) hs_kernel_fm_2_3, #endif #if (HS_BS_SLABS_LOG2_RU == 3) hs_kernel_fm_2_4, #endif #if (HS_BS_SLABS_LOG2_RU == 4) hs_kernel_fm_2_5, #endif #if (HS_BS_SLABS_LOG2_RU == 5) hs_kernel_fm_2_6, #endif #if (HS_BS_SLABS_LOG2_RU == 6) hs_kernel_fm_2_7, #endif #if (HS_BS_SLABS_LOG2_RU == 7) hs_kernel_fm_2_8, #endif #endif }; // // // typedef void (*hs_kernel_offset_fm_pfn)(HS_KEY_TYPE * const HS_RESTRICT vout, uint32_t const span_offset); #if (HS_FM_SCALE_MIN == 0) static hs_kernel_offset_fm_pfn const hs_kernels_offset_fm_0[] { #if (HS_BS_SLABS_LOG2_RU >= 2) hs_kernel_fm_0_0, #endif #if (HS_BS_SLABS_LOG2_RU >= 3) hs_kernel_fm_0_1, #endif #if (HS_BS_SLABS_LOG2_RU >= 4) hs_kernel_fm_0_2, #endif #if (HS_BS_SLABS_LOG2_RU >= 5) hs_kernel_fm_0_3, #endif #if (HS_BS_SLABS_LOG2_RU >= 6) hs_kernel_fm_0_4, #endif #if (HS_BS_SLABS_LOG2_RU >= 7) hs_kernel_fm_0_5, #endif }; #endif #if (HS_FM_SCALE_MIN <= 1) && (1 <= HS_FM_SCALE_MAX) static hs_kernel_offset_fm_pfn const hs_kernels_offset_fm_1[] { #if (HS_BS_SLABS_LOG2_RU >= 1) hs_kernel_fm_1_0, #endif #if (HS_BS_SLABS_LOG2_RU >= 2) hs_kernel_fm_1_1, #endif #if (HS_BS_SLABS_LOG2_RU >= 3) hs_kernel_fm_1_2, #endif #if (HS_BS_SLABS_LOG2_RU >= 4) hs_kernel_fm_1_3, #endif #if (HS_BS_SLABS_LOG2_RU >= 5) hs_kernel_fm_1_4, #endif #if (HS_BS_SLABS_LOG2_RU >= 6) hs_kernel_fm_1_5, #endif #if (HS_BS_SLABS_LOG2_RU >= 7) hs_kernel_fm_1_6, #endif }; #endif #if (HS_FM_SCALE_MIN <= 2) && (2 <= HS_FM_SCALE_MAX) static hs_kernel_offset_fm_pfn const hs_kernels_offset_fm_2[] { hs_kernel_fm_2_0, #if (HS_BS_SLABS_LOG2_RU >= 1) hs_kernel_fm_2_1, #endif #if (HS_BS_SLABS_LOG2_RU >= 2) hs_kernel_fm_2_2, #endif #if (HS_BS_SLABS_LOG2_RU >= 3) hs_kernel_fm_2_3, #endif #if (HS_BS_SLABS_LOG2_RU >= 4) hs_kernel_fm_2_4, #endif #if (HS_BS_SLABS_LOG2_RU >= 5) hs_kernel_fm_2_5, #endif #if (HS_BS_SLABS_LOG2_RU >= 6) hs_kernel_fm_2_6, #endif #if (HS_BS_SLABS_LOG2_RU >= 7) hs_kernel_fm_2_7, #endif }; #endif static hs_kernel_offset_fm_pfn const * const hs_kernels_offset_fm[] { #if (HS_FM_SCALE_MIN == 0) hs_kernels_offset_fm_0, #endif #if (HS_FM_SCALE_MIN <= 1) && (1 <= HS_FM_SCALE_MAX) hs_kernels_offset_fm_1, #endif #if (HS_FM_SCALE_MIN <= 2) && (2 <= HS_FM_SCALE_MAX) hs_kernels_offset_fm_2, #endif }; // // // typedef uint32_t hs_indices_t; // // // struct hs_state { // key buffers HS_KEY_TYPE * vin; HS_KEY_TYPE * vout; // can be vin cudaStream_t streams[3]; // pool of stream indices hs_indices_t pool; // bx_ru is number of rounded up warps in vin uint32_t bx_ru; }; // // // static uint32_t hs_indices_acquire(hs_indices_t * const indices) { // // FIXME -- an FFS intrinsic might be faster but there are so few // bits in this implementation that it might not matter. // if (*indices & 1) { *indices = *indices & ~1; return 0; } else if (*indices & 2) { *indices = *indices & ~2; return 1; } else // if (*indices & 4) { *indices = *indices & ~4; return 2; } } static uint32_t hs_state_acquire(struct hs_state * const state, hs_indices_t * const indices) { // // FIXME -- an FFS intrinsic might be faster but there are so few // bits in this implementation that it might not matter. // if (state->pool & 1) { state->pool &= ~1; *indices |= 1; return 0; } else if (state->pool & 2) { state->pool &= ~2; *indices |= 2; return 1; } else // (state->pool & 4) { state->pool &= ~4; *indices |= 4; return 2; } } static void hs_indices_merge(hs_indices_t * const to, hs_indices_t const from) { *to |= from; } static void hs_barrier_enqueue(cudaStream_t to, cudaStream_t from) { cudaEvent_t event_before; cuda(EventCreate(&event_before)); cuda(EventRecord(event_before,from)); cuda(StreamWaitEvent(to,event_before,0)); cuda(EventDestroy(event_before)); } static hs_indices_t hs_barrier(struct hs_state * const state, hs_indices_t const before, hs_indices_t * const after, uint32_t const count) // count is 1 or 2 { // return streams this stage depends on back into the pool hs_indices_merge(&state->pool,before); hs_indices_t indices = 0; // acquire 'count' stream indices for this stage for (uint32_t ii=0; iistreams[idx]; // // FIXME -- an FFS loop might be slower for so few bits. So // leave it as is for now. // if (wait & 1) hs_barrier_enqueue(to,state->streams[0]); if (wait & 2) hs_barrier_enqueue(to,state->streams[1]); if (wait & 4) hs_barrier_enqueue(to,state->streams[2]); } } hs_indices_merge(after,indices); return indices; } // // // #ifndef NDEBUG #include #define HS_STREAM_SYNCHRONIZE(s) \ cuda(StreamSynchronize(s)); \ fprintf(stderr,"%s\n",__func__); #else #define HS_STREAM_SYNCHRONIZE(s) #endif // // // static void hs_transpose(struct hs_state * const state) { HS_TRANSPOSE_KERNEL_NAME() <<bx_ru,HS_SLAB_THREADS,0,state->streams[0]>>> (state->vout); HS_STREAM_SYNCHRONIZE(state->streams[0]); } // // // static void hs_bc(struct hs_state * const state, hs_indices_t const hs_bc, hs_indices_t * const fm, uint32_t const down_slabs, uint32_t const clean_slabs_log2) { // enqueue any necessary barriers hs_indices_t indices = hs_barrier(state,hs_bc,fm,1); // block clean the minimal number of down_slabs_log2 spans uint32_t const frac_ru = (1u << clean_slabs_log2) - 1; uint32_t const full = (down_slabs + frac_ru) >> clean_slabs_log2; uint32_t const threads = HS_SLAB_THREADS << clean_slabs_log2; // stream will *always* be stream[0] cudaStream_t stream = state->streams[hs_indices_acquire(&indices)]; hs_kernels_bc[clean_slabs_log2] <<>> (state->vout); HS_STREAM_SYNCHRONIZE(stream); } // // // static uint32_t hs_hm(struct hs_state * const state, hs_indices_t const hs_bc, hs_indices_t * const hs_bc_tmp, uint32_t const down_slabs, uint32_t const clean_slabs_log2) { // enqueue any necessary barriers hs_indices_t indices = hs_barrier(state,hs_bc,hs_bc_tmp,1); // how many scaled half-merge spans are there? uint32_t const frac_ru = (1 << clean_slabs_log2) - 1; uint32_t const spans = (down_slabs + frac_ru) >> clean_slabs_log2; // for now, just clamp to the max uint32_t const log2_rem = clean_slabs_log2 - HS_BC_SLABS_LOG2_MAX; uint32_t const scale_log2 = MIN_MACRO(HS_HM_SCALE_MAX,log2_rem); uint32_t const log2_out = log2_rem - scale_log2; // // Size the grid // // The simplifying choices below limit the maximum keys that can be // sorted with this grid scheme to around ~2B. // // .x : slab height << clean_log2 -- this is the slab span // .y : [1...65535] -- this is the slab index // .z : ( this could also be used to further expand .y ) // // Note that OpenCL declares a grid in terms of global threads and // not grids and blocks // dim3 grid; grid.x = (HS_SLAB_HEIGHT / HS_HM_BLOCK_HEIGHT) << log2_out; grid.y = spans; grid.z = 1; cudaStream_t stream = state->streams[hs_indices_acquire(&indices)]; hs_kernels_hm[scale_log2-HS_HM_SCALE_MIN] <<>> (state->vout); HS_STREAM_SYNCHRONIZE(stream); return log2_out; } // // FIXME -- some of this logic can be skipped if BS is a power-of-two // static uint32_t hs_fm(struct hs_state * const state, hs_indices_t const fm, hs_indices_t * const hs_bc, uint32_t * const down_slabs, uint32_t const up_scale_log2) { // // FIXME OPTIMIZATION: in previous HotSort launchers it's sometimes // a performance win to bias toward launching the smaller flip merge // kernel in order to get more warps in flight (increased // occupancy). This is useful when merging small numbers of slabs. // // Note that HS_FM_SCALE_MIN will always be 0 or 1. // // So, for now, just clamp to the max until there is a reason to // restore the fancier and probably low-impact approach. // uint32_t const scale_log2 = MIN_MACRO(HS_FM_SCALE_MAX,up_scale_log2); uint32_t const clean_log2 = up_scale_log2 - scale_log2; // number of slabs in a full-sized scaled flip-merge span uint32_t const full_span_slabs = HS_BS_SLABS << up_scale_log2; // how many full-sized scaled flip-merge spans are there? uint32_t full_fm = state->bx_ru / full_span_slabs; uint32_t frac_fm = 0; // initialize down_slabs *down_slabs = full_fm * full_span_slabs; // how many half-size scaled + fractional scaled spans are there? uint32_t const span_rem = state->bx_ru - *down_slabs; uint32_t const half_span_slabs = full_span_slabs >> 1; // if we have over a half-span then fractionally merge it if (span_rem > half_span_slabs) { // the remaining slabs will be cleaned *down_slabs += span_rem; uint32_t const frac_rem = span_rem - half_span_slabs; uint32_t const frac_rem_pow2 = pow2_ru_u32(frac_rem); if (frac_rem_pow2 >= half_span_slabs) { // bump it up to a full span full_fm += 1; } else { // otherwise, add fractional frac_fm = MAX_MACRO(1,frac_rem_pow2 >> clean_log2); } } // enqueue any necessary barriers bool const both = (full_fm != 0) && (frac_fm != 0); hs_indices_t indices = hs_barrier(state,fm,hs_bc,both ? 2 : 1); // // Size the grid // // The simplifying choices below limit the maximum keys that can be // sorted with this grid scheme to around ~2B. // // .x : slab height << clean_log2 -- this is the slab span // .y : [1...65535] -- this is the slab index // .z : ( this could also be used to further expand .y ) // // Note that OpenCL declares a grid in terms of global threads and // not grids and blocks // dim3 grid; grid.x = (HS_SLAB_HEIGHT / HS_FM_BLOCK_HEIGHT) << clean_log2; grid.z = 1; if (full_fm > 0) { cudaStream_t stream = state->streams[hs_indices_acquire(&indices)]; grid.y = full_fm; hs_kernels_fm[scale_log2-HS_FM_SCALE_MIN] <<>> (state->vout); HS_STREAM_SYNCHRONIZE(stream); } if (frac_fm > 0) { cudaStream_t stream = state->streams[hs_indices_acquire(&indices)]; grid.y = 1; hs_kernels_offset_fm[scale_log2-HS_FM_SCALE_MIN][msb_idx_u32(frac_fm)] <<>> (state->vout,full_fm); HS_STREAM_SYNCHRONIZE(stream); } return clean_log2; } // // // static void hs_bs(struct hs_state * const state, hs_indices_t const bs, hs_indices_t * const fm, uint32_t const count_padded_in) { uint32_t const slabs_in = count_padded_in / HS_SLAB_KEYS; uint32_t const full_bs = slabs_in / HS_BS_SLABS; uint32_t const frac_bs = slabs_in - full_bs * HS_BS_SLABS; bool const both = (full_bs != 0) && (frac_bs != 0); // enqueue any necessary barriers hs_indices_t indices = hs_barrier(state,bs,fm,both ? 2 : 1); if (full_bs != 0) { cudaStream_t stream = state->streams[hs_indices_acquire(&indices)]; CONCAT_MACRO(hs_kernel_bs_,HS_BS_SLABS_LOG2_RU) <<>> (state->vout,state->vin); HS_STREAM_SYNCHRONIZE(stream); } if (frac_bs != 0) { cudaStream_t stream = state->streams[hs_indices_acquire(&indices)]; hs_kernels_offset_bs[msb_idx_u32(frac_bs)] <<<1,frac_bs*HS_SLAB_THREADS,0,stream>>> (state->vout,state->vin,full_bs*HS_BS_SLABS*HS_SLAB_THREADS); HS_STREAM_SYNCHRONIZE(stream); } } // // // static void hs_keyset_pre_merge(struct hs_state * const state, hs_indices_t * const fm, uint32_t const count_lo, uint32_t const count_hi) { uint32_t const vout_span = count_hi - count_lo; cudaStream_t stream = state->streams[hs_state_acquire(state,fm)]; cuda(MemsetAsync(state->vout + count_lo, 0xFF, vout_span * sizeof(HS_KEY_TYPE), stream)); } // // // static void hs_keyset_pre_sort(struct hs_state * const state, hs_indices_t * const bs, uint32_t const count, uint32_t const count_hi) { uint32_t const vin_span = count_hi - count; cudaStream_t stream = state->streams[hs_state_acquire(state,bs)]; cuda(MemsetAsync(state->vin + count, 0xFF, vin_span * sizeof(HS_KEY_TYPE), stream)); } // // // void CONCAT_MACRO(hs_cuda_sort_,HS_KEY_TYPE_PRETTY) (HS_KEY_TYPE * const vin, HS_KEY_TYPE * const vout, uint32_t const count, uint32_t const count_padded_in, uint32_t const count_padded_out, bool const linearize, cudaStream_t stream0, // primary stream cudaStream_t stream1, // auxilary cudaStream_t stream2) // auxilary { // is this sort in place? bool const is_in_place = (vout == NULL); // cq, buffers, wait list and slab count struct hs_state state; state.vin = vin; state.vout = is_in_place ? vin : vout; state.streams[0] = stream0; state.streams[1] = stream1; state.streams[2] = stream2; state.pool = 0x7; // 3 bits state.bx_ru = (count + HS_SLAB_KEYS - 1) / HS_SLAB_KEYS; // initialize vin uint32_t const count_hi = is_in_place ? count_padded_out : count_padded_in; bool const is_pre_sort_keyset_reqd = count_hi > count; bool const is_pre_merge_keyset_reqd = !is_in_place && (count_padded_out > count_padded_in); hs_indices_t bs = 0; // initialize any trailing keys in vin before sorting if (is_pre_sort_keyset_reqd) hs_keyset_pre_sort(&state,&bs,count,count_hi); hs_indices_t fm = 0; // concurrently initialize any trailing keys in vout before merging if (is_pre_merge_keyset_reqd) hs_keyset_pre_merge(&state,&fm,count_padded_in,count_padded_out); // immediately sort blocks of slabs hs_bs(&state,bs,&fm,count_padded_in); // // we're done if this was a single bs block... // // otherwise, merge sorted spans of slabs until done // if (state.bx_ru > HS_BS_SLABS) { int32_t up_scale_log2 = 1; while (true) { hs_indices_t hs_or_bc = 0; uint32_t down_slabs; // flip merge slabs -- return span of slabs that must be cleaned uint32_t clean_slabs_log2 = hs_fm(&state, fm, &hs_or_bc, &down_slabs, up_scale_log2); // if span is gt largest slab block cleaner then half merge while (clean_slabs_log2 > HS_BC_SLABS_LOG2_MAX) { hs_indices_t hs_or_bc_tmp; clean_slabs_log2 = hs_hm(&state, hs_or_bc, &hs_or_bc_tmp, down_slabs, clean_slabs_log2); hs_or_bc = hs_or_bc_tmp; } // reset fm fm = 0; // launch clean slab grid -- is it the final launch? hs_bc(&state, hs_or_bc, &fm, down_slabs, clean_slabs_log2); // was this the final block clean? if (((uint32_t)HS_BS_SLABS << up_scale_log2) >= state.bx_ru) break; // otherwise, merge twice as many slabs up_scale_log2 += 1; } } // slabs or linear? if (linearize) { // guaranteed to be on stream0 hs_transpose(&state); } } // // all grids will be computed as a function of the minimum number of slabs // void CONCAT_MACRO(hs_cuda_pad_,HS_KEY_TYPE_PRETTY) (uint32_t const count, uint32_t * const count_padded_in, uint32_t * const count_padded_out) { // // round up the count to slabs // uint32_t const slabs_ru = (count + HS_SLAB_KEYS - 1) / HS_SLAB_KEYS; uint32_t const blocks = slabs_ru / HS_BS_SLABS; uint32_t const block_slabs = blocks * HS_BS_SLABS; uint32_t const slabs_ru_rem = slabs_ru - block_slabs; uint32_t const slabs_ru_rem_ru = MIN_MACRO(pow2_ru_u32(slabs_ru_rem),HS_BS_SLABS); *count_padded_in = (block_slabs + slabs_ru_rem_ru) * HS_SLAB_KEYS; *count_padded_out = *count_padded_in; // // will merging be required? // if (slabs_ru > HS_BS_SLABS) { // more than one block uint32_t const blocks_lo = pow2_rd_u32(blocks); uint32_t const block_slabs_lo = blocks_lo * HS_BS_SLABS; uint32_t const block_slabs_rem = slabs_ru - block_slabs_lo; if (block_slabs_rem > 0) { uint32_t const block_slabs_rem_ru = pow2_ru_u32(block_slabs_rem); uint32_t const block_slabs_hi = MAX_MACRO(block_slabs_rem_ru, blocks_lo << (1 - HS_FM_SCALE_MIN)); uint32_t const block_slabs_padded_out = MIN_MACRO(block_slabs_lo+block_slabs_hi, block_slabs_lo*2); // clamp non-pow2 blocks *count_padded_out = block_slabs_padded_out * HS_SLAB_KEYS; } } } // // // void CONCAT_MACRO(hs_cuda_info_,HS_KEY_TYPE_PRETTY) (uint32_t * const key_words, uint32_t * const val_words, uint32_t * const slab_height, uint32_t * const slab_width_log2) { *key_words = HS_KEY_WORDS; *val_words = HS_VAL_WORDS; *slab_height = HS_SLAB_HEIGHT; *slab_width_log2 = HS_SLAB_WIDTH_LOG2; } // // //