Enable more warning flags. by alliepiper · Pull Request #249 · NVIDIA/cub

alliepiper · 2020-12-11T02:44:33Z

Fixes #228.

brycelelbach · 2021-02-08T20:21:10Z

cmake/CubBuildCompilerTargets.cmake

+    append_option_if_available("-Wgnu" cxx_compile_options)
+    # Calling a variadic macro with zero args is a GNU extension until C++20,
+    # but the THRUST_PP_ARITY macro is used with zero args. Need to see if this
+    # is a real problem worth fixing.


It shouldn't be a problem as every compiler under the sun supports the extension, including NVRTC I believe.

brycelelbach · 2021-02-08T20:28:50Z

You should re-run this with CUB in -DDEBUG debugging mode - you'll hit a LOT more of the assignment in conditional warning.

brycelelbach · 2021-02-08T20:31:12Z

The whole if (CubDebug(error = ...)) break pattern should probably be encapsulated in a macro.

brycelelbach · 2021-02-08T20:34:48Z

cub/grid/grid_even_share.cuh

+        this->block_end             = num_items_;    // Initialize past-the-end
+        this->num_items             = num_items_;
+        this->total_tiles           = (num_items_ + tile_items - 1) / tile_items;
+        this->grid_size             = CUB_MIN(static_cast<int>(total_tiles), max_grid_size);


This could overflow, because total_tiles is OffsetT which could be 64bit.

This may be related to #221.

I think I'd prefer us to either make sure all the static casts don't overflow, or move them to a separate PR.

Fixed this and similar issues in the thrust/cub kernel dispatch code with a new cub::DivideAndRoundUp method that avoids these (n + d - 1) / d overflows.

brycelelbach · 2021-02-08T20:41:54Z

cub/grid/grid_even_share.cuh

        OffsetT avg_tiles_per_block = total_tiles / grid_size;
-        this->big_shares            = total_tiles - (avg_tiles_per_block * grid_size);        // leftover grains go to big blocks
+        // leftover grains go to big blocks
+        this->big_shares = static_cast<int>(total_tiles - (avg_tiles_per_block * grid_size));


Is there any reason big_share can't be an OffsetT?

I took a closer look and if anything, total_tiles should be changed to a 32-bit int since it's used as a block dimension and those can't exceed INT32_MAX. The only usages of the variable assume it is a 32-bit int.

I'll change this so that total_tiles is appropriately sized, which will eliminate the need for this (and a few other) static_casts.

brycelelbach · 2021-02-08T20:42:25Z

cub/iterator/tex_ref_input_iterator.cuh

+#pragma GCC diagnostic push
+#pragma GCC diagnostic ignored "-Wdeprecated-declarations"
+#endif
+


I would deprecate or remove this entirely.

brycelelbach · 2021-02-08T20:44:55Z

test/half.h

                    ir |= ia >> (24 - 11);
                    ia = ia << (32 - (24 - 11));
-                    ir = ir + ((14 + shift) << 10);
+                    ir = static_cast<uint16_t>(ir + ((14 + shift) << 10));


This is probably okay because it's just in the test code, but I'm pretty sure this is actually a bug due to potential overflow.

brycelelbach · 2021-02-08T20:46:21Z

test/test_device_run_length_encode.cu

-            TestDispatch<T, OffsetT, LengthT>(num_items);
+            unsigned int num;
+            RandomBits(num);
+            num = (unsigned int) ((double(num) * double(10000000)) / double(max_int));


Since we're changing this already, maybe static_cast here instead?

brycelelbach · 2021-02-08T20:46:56Z

test/test_device_scan.cu

-            TestOp<InputT>(num_items,  identity, initial_value);
+            unsigned int num;
+            RandomBits(num);
+            num = (unsigned int) ((double(num) * double(10000000)) / double(max_int));


Maybe static_cast here as a drive-by fix?

brycelelbach

This looks fine, I'd prefer to split the static_cast into a separate PR, or make sure they don't overflow, maybe add a comment to each explaining why they're fine.

alliepiper · 2021-02-09T23:54:55Z

You should re-run this with CUB in -DDEBUG debugging mode - you'll hit a LOT more of the assignment in conditional warning.

Updated #1175 with this recommendation.

alliepiper · 2021-02-10T00:24:39Z

This looks fine, I'd prefer to split the static_cast into a separate PR, or make sure they don't overflow, maybe add a comment to each explaining why they're fine.

After fixing #221, I reviewed the remaining usages and all are safe. Added comments to any non-obviously-safe ones.

Anonymous structs are C features. In C++, they're non-portable compiler extensions. These only seemed to pop up in CUB-style `TempStorage` objects, I just picked some reasonable sounding names for them.

We need to deprecate this class since the underlying CUDA APIs are deprecated. This suppression is a temporary workaround. Tracked by NVIDIA#191.

Changing to a `static_cast` fixes this warning.

The `cuda_std_17` compile feature is broken for MSVC when CMake < 3.18.3.

Users have been reporting that device algorithms return invalid `temp_storage_bytes` values when `num_items` is close to -- but not over -- INT32_MAX. This is caused by an overflow in the numerator of the pattern `num_tiles = (num_items + items_per_tile - 1) / items_per_tile`. The new function implements the same calculation but protects against overflow. Fixes NVIDIA#221. Bug 3075796

This value will always be representable with an int, and all usages of it treat it as a 32-bit int. Changing the type avoids some casts at usage sites.

alliepiper marked this pull request as draft December 11, 2020 02:44

alliepiper added this to the 1.12.0 milestone Dec 11, 2020

alliepiper force-pushed the enh/pedantic_flags/gh.cub228 branch from c065b3a to 94d78b3 Compare December 11, 2020 13:48

alliepiper mentioned this pull request Dec 11, 2020

Enable more warning flags. NVIDIA/thrust#1359

Merged

alliepiper mentioned this pull request Dec 24, 2020

Remove NPP dependency from test_device_histogram.cu. #254

Merged

alliepiper mentioned this pull request Jan 9, 2021

Warning when using GridEvenShare with unsigned offsets #257

Closed

alliepiper linked an issue Jan 9, 2021 that may be closed by this pull request

Warning when using GridEvenShare with unsigned offsets #257

Closed

alliepiper force-pushed the enh/pedantic_flags/gh.cub228 branch 2 times, most recently from 06366b2 to 8983abe Compare January 23, 2021 02:55

alliepiper requested a review from brycelelbach January 27, 2021 19:25

alliepiper marked this pull request as ready for review January 27, 2021 19:26

alliepiper force-pushed the enh/pedantic_flags/gh.cub228 branch from 8983abe to d588658 Compare January 27, 2021 22:22

brycelelbach reviewed Feb 8, 2021

View reviewed changes

brycelelbach approved these changes Feb 8, 2021

View reviewed changes

alliepiper mentioned this pull request Feb 9, 2021

NVBug: 3075796 [PyTorch] temp_storage_bytes overflows in InclusiveScan for size_cub value close to int32 max #221

Closed

alliepiper added 5 commits February 9, 2021 21:50

Name unnamed structs.

39a2172

Anonymous structs are C features. In C++, they're non-portable compiler extensions. These only seemed to pop up in CUB-style `TempStorage` objects, I just picked some reasonable sounding names for them.

Fix assorted shadowed variable warnings.

4b59207

warning C4706: assignment within conditional expression

d182eef

Replace deprecated cudaThreadSynchronize() calls.

f61bc46

Suppress deprecation warnings for TexRefInputIterator.

e4591c1

We need to deprecate this class since the underlying CUDA APIs are deprecated. This suppression is a temporary workaround. Tracked by NVIDIA#191.

alliepiper added 4 commits February 9, 2021 21:50

warning C4310: cast truncates constant value.

541916b

Changing to a `static_cast` fixes this warning.

Fix C++17 + NVCC + MSVC + CMake.

e9d77ab

The `cuda_std_17` compile feature is broken for MSVC when CMake < 3.18.3.

Fix numeric conversion warnings.

c353fce

Enable more compiler warning flags.

b7207a2

alliepiper force-pushed the enh/pedantic_flags/gh.cub228 branch from d588658 to 91e9ac7 Compare February 10, 2021 02:56

alliepiper linked an issue Feb 10, 2021 that may be closed by this pull request

NVBug: 3075796 [PyTorch] temp_storage_bytes overflows in InclusiveScan for size_cub value close to int32 max #221

Closed

alliepiper force-pushed the enh/pedantic_flags/gh.cub228 branch from 91e9ac7 to 0fb8228 Compare February 10, 2021 03:37

alliepiper added 2 commits February 9, 2021 22:45

Refactor GridEvenShare::total_tiles to int32.

e0a6736

This value will always be representable with an int, and all usages of it treat it as a 32-bit int. Changing the type avoids some casts at usage sites.

alliepiper force-pushed the enh/pedantic_flags/gh.cub228 branch from 0fb8228 to e0a6736 Compare February 10, 2021 03:46

alliepiper merged commit b229817 into NVIDIA:main Feb 16, 2021

alliepiper deleted the enh/pedantic_flags/gh.cub228 branch February 16, 2021 18:47

This was referenced Feb 19, 2021

[REVIEW] Reverting change that removed unsigned in dispatch_spmv_orig.cuh #196

Closed

cub.device_csrmv is corrupting sparse arrays cupy/cupy#3822

Closed

Conversation

alliepiper commented Dec 11, 2020

Uh oh!

Choose a reason for hiding this comment

Uh oh!

brycelelbach commented Feb 8, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

brycelelbach commented Feb 8, 2021

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

brycelelbach left a comment

Choose a reason for hiding this comment

Uh oh!

alliepiper commented Feb 9, 2021

Uh oh!

alliepiper commented Feb 10, 2021

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

brycelelbach commented Feb 8, 2021 •

edited

Loading