From e0b3d402f6e2a18c50ca3e3b96931653e6505be7 Mon Sep 17 00:00:00 2001 From: Paul Mattione Date: Mon, 7 Oct 2024 11:58:18 -0400 Subject: [PATCH] undo loop unroll, increased reg count --- cpp/src/io/parquet/decode_fixed.cu | 16 +++------------- 1 file changed, 3 insertions(+), 13 deletions(-) diff --git a/cpp/src/io/parquet/decode_fixed.cu b/cpp/src/io/parquet/decode_fixed.cu index 42f90880fe9..4522ea7fe56 100644 --- a/cpp/src/io/parquet/decode_fixed.cu +++ b/cpp/src/io/parquet/decode_fixed.cu @@ -71,19 +71,9 @@ __device__ static void scan_block_exclusive_sum(uint32_t warp_bits, // Compute block-wide results results.block_count = 0; results.thread_count_within_block = results.thread_count_within_warp; - if constexpr ((num_warps == 4) || (num_warps == 8)) { - results.block_count = warp_counts[0] + warp_counts[1] + warp_counts[2] + warp_counts[3]; - if constexpr (num_warps == 8) { - results.block_count += warp_counts[4] + warp_counts[5] + warp_counts[6] + warp_counts[7]; - } - for (int warp_idx = 0; warp_idx < warp_index; ++warp_idx) { - results.thread_count_within_block += warp_counts[warp_idx]; - } - } else { - for (int warp_idx = 0; warp_idx < num_warps; ++warp_idx) { - results.block_count += warp_counts[warp_idx]; - if (warp_idx < warp_index) { results.thread_count_within_block += warp_counts[warp_idx]; } - } + for (int warp_idx = 0; warp_idx < num_warps; ++warp_idx) { + results.block_count += warp_counts[warp_idx]; + if (warp_idx < warp_index) { results.thread_count_within_block += warp_counts[warp_idx]; } } }