/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.5.1/hipcub/include/hipcub/backend/rocprim/block/block_run_length_decode.hpp Source File#
block_run_length_decode.hpp
143 using RunOffsetScanT = BlockScan<DecodedOffsetT, BLOCK_DIM_X, BLOCK_SCAN_WARP_SCANS, BLOCK_DIM_Y, BLOCK_DIM_Z>;
191 : temp_storage(temp_storage.Alias()), linear_tid(RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z))
205 : temp_storage(temp_storage.Alias()), linear_tid(RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z))
217 : temp_storage(PrivateStorage()), linear_tid(RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z))
228 : temp_storage(PrivateStorage()), linear_tid(RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z))
273 RunOffsetT thread_dst_offset = static_cast<RunOffsetT>(linear_tid) * static_cast<RunOffsetT>(RUNS_PER_THREAD);
299 RunOffsetScanT(this->temp_storage.offset_scan).ExclusiveSum(run_offsets, run_offsets, decoded_size_aggregate);
302 // Ensure the prefix scan's temporary storage can be reused (may be superfluous, but depends on scan implementation)
327 HIPCUB_DEVICE __forceinline__ void RunLengthDecode(ItemT (&decoded_items)[DECODED_ITEMS_PER_THREAD],
332 DecodedOffsetT thread_decoded_offset = from_decoded_offset + linear_tid * DECODED_ITEMS_PER_THREAD;
335 // If this thread's <thread_decoded_offset> is already beyond the total decoded size, it will be assigned to the
338 StaticUpperBound<BLOCK_RUNS>(temp_storage.runs.run_offsets, BLOCK_RUNS, thread_decoded_offset) -
343 // If this thread is getting assigned the last run, we make sure it will not fetch any other run after this
357 // We make sure that a thread is not re-entering this conditional when being assigned to the last run already by
362 // If this thread is getting assigned the last run, we make sure it will not fetch any other run after this
363 assigned_run_end = (assigned_run == BLOCK_RUNS - 1) ? thread_decoded_offset + DECODED_ITEMS_PER_THREAD
383 HIPCUB_DEVICE __forceinline__ void RunLengthDecode(ItemT (&decoded_items)[DECODED_ITEMS_PER_THREAD],
The BlockRunLengthDecode class supports decoding a run-length encoded array of items....
Definition: block_run_length_decode.hpp:131
__device__ __forceinline__ void RunLengthDecode(ItemT(&decoded_items)[DECODED_ITEMS_PER_THREAD], RelativeOffsetT(&item_offsets)[DECODED_ITEMS_PER_THREAD], DecodedOffsetT from_decoded_offset=0)
Run-length decodes the runs previously passed via a call to Init(...) and returns the run-length deco...
Definition: block_run_length_decode.hpp:327
__device__ __forceinline__ BlockRunLengthDecode(TempStorage &temp_storage, ItemT(&run_values)[RUNS_PER_THREAD], RunLengthT(&run_lengths)[RUNS_PER_THREAD], TotalDecodedSizeT &total_decoded_size)
Constructor specialised for user-provided temporary storage, initializing using the runs' lengths....
Definition: block_run_length_decode.hpp:187
__device__ __forceinline__ BlockRunLengthDecode(TempStorage &temp_storage, ItemT(&run_values)[RUNS_PER_THREAD], UserRunOffsetT(&run_offsets)[RUNS_PER_THREAD])
Constructor specialised for user-provided temporary storage, initializing using the runs' offsets....
Definition: block_run_length_decode.hpp:202
__device__ __forceinline__ void RunLengthDecode(ItemT(&decoded_items)[DECODED_ITEMS_PER_THREAD], DecodedOffsetT from_decoded_offset=0)
Run-length decodes the runs previously passed via a call to Init(...) and returns the run-length deco...
Definition: block_run_length_decode.hpp:383
__device__ __forceinline__ BlockRunLengthDecode(ItemT(&run_values)[RUNS_PER_THREAD], UserRunOffsetT(&run_offsets)[RUNS_PER_THREAD])
Constructor specialised for static temporary storage, initializing using the runs' offsets.
Definition: block_run_length_decode.hpp:226
__device__ __forceinline__ BlockRunLengthDecode(ItemT(&run_values)[RUNS_PER_THREAD], RunLengthT(&run_lengths)[RUNS_PER_THREAD], TotalDecodedSizeT &total_decoded_size)
Constructor specialised for static temporary storage, initializing using the runs' lengths.
Definition: block_run_length_decode.hpp:214
Definition: block_scan.hpp:80
Definition: block_run_length_decode.hpp:174
A storage-backing wrapper that allows types with non-trivial constructors to be aliased in unions.
Definition: util_type.hpp:363