Skip to content

Commit f412ebc

Browse files
JacobfaibJacobfaib
committed
bugprone-misplaced-widening-cast
1 parent a4ffbcf commit f412ebc

75 files changed

Lines changed: 245 additions & 172 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎.clang-tidy‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,6 @@ Checks:
3232
- 'bugprone-*'
3333
- '-bugprone-forward-declaration-namespace'
3434
- '-bugprone-narrowing-conversions'
35-
- '-bugprone-misplaced-widening-cast'
3635
- '-bugprone-signed-char-misuse'
3736
- '-bugprone-suspicious-include'
3837
- '-bugprone-multi-level-implicit-pointer-conversion'

‎cub/cub/agent/agent_batch_memcpy.cuh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -188,7 +188,7 @@ GetAlignedPtrs(const void* in_begin, void* out_begin, ByteOffsetT num_bytes)
188188
// Bytes after `out_chars_aligned` to the last VectorT-aligned
189189
// address at (or before) `out_begin` + `num_bytes`
190190
uint32_t out_end_aligned{};
191-
if (in_end_padding_req + alignment_offset > num_bytes)
191+
if (in_end_padding_req + alignment_offset > num_bytes) // NOLINT(bugprone-misplaced-widening-cast)
192192
{
193193
out_end_aligned = out_start_aligned;
194194
}

‎cub/cub/agent/agent_for.cuh‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -41,7 +41,8 @@ struct agent_block_striped_t
4141
_CCCL_PRAGMA_UNROLL_FULL()
4242
for (int item = 0; item < items_per_thread; item++)
4343
{
44-
const auto idx = static_cast<OffsetT>(threads_per_block * item + threadIdx.x);
44+
const auto idx =
45+
static_cast<OffsetT>(threads_per_block * item + threadIdx.x); // NOLINT(bugprone-misplaced-widening-cast)
4546

4647
if (IsFullTile || idx < items_in_tile)
4748
{

‎cub/cub/agent/agent_merge.cuh‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -78,8 +78,9 @@ struct agent_t
7878
struct alignas(cub::detail::LoadToSharedBufferAlignBytes<ValueT>()) buffer_t
7979
{
8080
// Need extra bytes of padding for TMA because this static buffer has to hold the two dynamically sized buffers.
81-
static constexpr int bytes_needed = cub::detail::LoadToSharedBufferSizeBytes<ValueT>(items_per_tile + 1)
82-
+ (alignof(ValueT) < bl2sh_minimum_align ? 2 * bl2sh_minimum_align : 0);
81+
static constexpr int bytes_needed =
82+
cub::detail::LoadToSharedBufferSizeBytes<ValueT>(items_per_tile + 1) // NOLINT(bugprone-misplaced-widening-cast)
83+
+ (alignof(ValueT) < bl2sh_minimum_align ? 2 * bl2sh_minimum_align : 0);
8384

8485
char c_array[bytes_needed];
8586
};

‎cub/cub/agent/agent_radix_sort_downsweep.cuh‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -257,7 +257,9 @@ struct AgentRadixSortDownsweep
257257

258258
key = bit_ordered_conversion::from_bit_ordered(decomposer, key);
259259

260-
if (FULL_TILE || (static_cast<OffsetT>(threadIdx.x + (ITEM * BLOCK_THREADS)) < valid_items))
260+
if (FULL_TILE
261+
|| (static_cast<OffsetT>(threadIdx.x + (ITEM * BLOCK_THREADS)) // NOLINT(bugprone-misplaced-widening-cast)
262+
< valid_items))
261263
{
262264
d_keys_out[relative_bin_offsets[ITEM] + threadIdx.x + (ITEM * BLOCK_THREADS)] = key;
263265
}
@@ -291,7 +293,9 @@ struct AgentRadixSortDownsweep
291293
{
292294
ValueT value = exchange_values[threadIdx.x + (ITEM * BLOCK_THREADS)];
293295

294-
if (FULL_TILE || (static_cast<OffsetT>(threadIdx.x + (ITEM * BLOCK_THREADS)) < valid_items))
296+
if (FULL_TILE
297+
|| (static_cast<OffsetT>(threadIdx.x + (ITEM * BLOCK_THREADS)) // NOLINT(bugprone-misplaced-widening-cast)
298+
< valid_items))
295299
{
296300
d_values_out[relative_bin_offsets[ITEM] + threadIdx.x + (ITEM * BLOCK_THREADS)] = value;
297301
}

‎cub/cub/agent/agent_radix_sort_histogram.cuh‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -247,7 +247,8 @@ struct AgentRadixSortHistogram
247247
// Process the tiles.
248248
OffsetT portion_offset = portion * MAX_PORTION_SIZE;
249249
OffsetT portion_size = ::cuda::std::min(MAX_PORTION_SIZE, num_items - portion_offset);
250-
for (OffsetT offset = blockIdx.x * TILE_ITEMS; offset < portion_size; offset += TILE_ITEMS * gridDim.x)
250+
for (OffsetT offset = static_cast<OffsetT>(blockIdx.x) * TILE_ITEMS; offset < portion_size;
251+
offset += static_cast<OffsetT>(TILE_ITEMS) * gridDim.x)
251252
{
252253
OffsetT tile_offset = portion_offset + offset;
253254
bit_ordered_type keys[ITEMS_PER_THREAD];

‎cub/cub/agent/agent_radix_sort_onesweep.cuh‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -398,7 +398,7 @@ struct AgentRadixSortOnesweep
398398
{
399399
// gather and scatter the values
400400
ValueT values[ITEMS_PER_THREAD];
401-
LoadValues(block_idx * TILE_ITEMS, values);
401+
LoadValues(block_idx * TILE_ITEMS, values); // NOLINT(bugprone-misplaced-widening-cast)
402402
if (full_block)
403403
{
404404
StoreDirectWarpStriped(threadIdx.x, d_values_out + global_offset, values);
@@ -598,7 +598,7 @@ struct AgentRadixSortOnesweep
598598

599599
// load values
600600
ValueT values[ITEMS_PER_THREAD];
601-
LoadValues(block_idx * TILE_ITEMS, values);
601+
LoadValues(block_idx * TILE_ITEMS, values); // NOLINT(bugprone-misplaced-widening-cast)
602602

603603
// scatter values
604604
__syncthreads();
@@ -618,7 +618,7 @@ struct AgentRadixSortOnesweep
618618
// if warp1 < warp2, all elements of warp1 occur before those of warp2
619619
// in the source array
620620
bit_ordered_type keys[ITEMS_PER_THREAD];
621-
LoadKeys(block_idx * TILE_ITEMS, keys);
621+
LoadKeys(block_idx * TILE_ITEMS, keys); // NOLINT(bugprone-misplaced-widening-cast)
622622

623623
// rank keys
624624
int ranks[ITEMS_PER_THREAD];

‎cub/cub/agent/agent_reduce.cuh‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -331,14 +331,15 @@ struct AgentReduceImpl
331331
// Read first item
332332
if (IsFirstTile && (thread_offset < valid_items))
333333
{
334-
thread_aggregate = transform_op(d_wrapped_in[block_offset + thread_offset]);
334+
thread_aggregate =
335+
transform_op(d_wrapped_in[block_offset + thread_offset]); // NOLINT(bugprone-misplaced-widening-cast)
335336
thread_offset += NumThreads;
336337
}
337338

338339
// Continue reading items (block-striped)
339340
while (thread_offset < valid_items)
340341
{
341-
InputT item(d_wrapped_in[block_offset + thread_offset]);
342+
InputT item(d_wrapped_in[block_offset + thread_offset]); // NOLINT(bugprone-misplaced-widening-cast)
342343

343344
thread_aggregate = reduction_op(thread_aggregate, transform_op(item));
344345
thread_offset += NumThreads;

‎cub/cub/agent/agent_reduce_by_key.cuh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -432,8 +432,8 @@ struct AgentReduceByKey
432432
for (int item = threadIdx.x; item < num_tile_segments; item += BLOCK_THREADS)
433433
{
434434
KeyValuePairT pair = temp_storage.raw_exchange.Alias()[item];
435-
d_unique_out[num_tile_segments_prefix + item] = pair.key;
436-
d_aggregates_out[num_tile_segments_prefix + item] = pair.value;
435+
d_unique_out[num_tile_segments_prefix + item] = pair.key; // NOLINT(bugprone-misplaced-widening-cast)
436+
d_aggregates_out[num_tile_segments_prefix + item] = pair.value; // NOLINT(bugprone-misplaced-widening-cast)
437437
}
438438
}
439439

‎cub/cub/agent/agent_rle.cuh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -362,7 +362,7 @@ struct AgentRle
362362
T tile_successor_item;
363363
if (threadIdx.x == BLOCK_THREADS - 1)
364364
{
365-
tile_successor_item = d_in[tile_offset + TILE_ITEMS];
365+
tile_successor_item = d_in[tile_offset + TILE_ITEMS]; // NOLINT(bugprone-misplaced-widening-cast)
366366
}
367367

368368
BlockDiscontinuityT(temp_storage.aliasable.scan_storage.discontinuity)
@@ -388,7 +388,7 @@ struct AgentRle
388388
T tile_successor_item;
389389
if (threadIdx.x == BLOCK_THREADS - 1)
390390
{
391-
tile_successor_item = d_in[tile_offset + TILE_ITEMS];
391+
tile_successor_item = d_in[tile_offset + TILE_ITEMS]; // NOLINT(bugprone-misplaced-widening-cast)
392392
}
393393

394394
// Get the last item from the previous tile

0 commit comments

Comments
 (0)