[BUG] Cascaded algorithm correctness depends on compiler flag
Opened this issue · 2 comments
Describe the bug
On branch-2.2, compiling the Cascaded algorithm with the -G flag changes the static shared memory alignment behavior, causing misalignment errors.
Steps/Code to reproduce bug
- Modify CMakeLists.txt:
- Comment out line CMakeLists.txt:33
set(CMAKE_CUDA_FLAGS_DEBUG "${CMAKE_CUDA_FLAGS_DEBUG};-g")- Uncomment line CMakeLists.txt:32
set(CMAKE_CUDA_FLAGS_DEBUG "${CMAKE_CUDA_FLAGS_DEBUG};-G").- Run tests/test_cascaded.cpp with the following test case:
TEST_CASE("comp/decomp cascaded-small-uint64", "[nvcomp][small]").To compile successfully, I reduced the default_chunk_size in default_chunk_size from 4096 to 2048.
Expected behavior
The test should pass without errors.
Environment details (please complete the following information):
-
Environment location:
-
Ubuntu-22.04
-
Driver Version: 555.99
-
CUDA Version: 12.5
-
NVIDIA GeForce RTX 3080
-
Method of nvCOMP install: branch-2.2 source code
Additional context
After debugging, I explicitly declared alignment for shared memory allocation in the following files:
__shared__ __align__(sizeof(data_type)) uint8_t shmem[shmem_size];__shared__ __align__(sizeof(data_type)) uint32_t chunk_metadata[max_chunk_metadata_size / sizeof(uint32_t)];After making these changes, all tests in test_cascaded.cpp passed. I believe this dependency on compiler optimization for correctness is a bug.
This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.
Thanks for the bug report. As far as I can tell, this bug was most likely fixed in nvCOMP 2.6.0 or 2.6.1, in January 2023.