[Issue]: rocPRIM -Wdeprecated-declarations warnings for raw_storage

[ScottTodd]

Problem Description

I'm seeing almost 1GB of warning logs from rocPRIM after updating amd-llvm in TheRock in ROCm/TheRock#725. This is up from 23MB of logs before updating amd-llvm.

Logs: https://therock-artifacts.s3.us-east-2.amazonaws.com/15304904194-windows/logs/gfx110X-dgpu/rocPRIM_build.log

96.2	In file included from D:/projects/TheRock/math-libs/hipCUB/test/hipcub/test_hipcub_block_radix_rank.cpp:36:
96.2	In file included from D:/projects/TheRock/math-libs/hipCUB/hipcub/include\hipcub/block/block_radix_rank.hpp:38:
96.2	D:/projects/TheRock/math-libs/hipCUB/hipcub/include\hipcub/block\../backend/rocprim/block/block_radix_rank.hpp:141:45: warning: 'storage_type' is deprecated: To store non default-constructible types in local memory, use rocprim::uninitialized_array instead [-Wdeprecated-declarations]
96.2	  141 |     using TempStorage = typename base_type::storage_type;
96.2	      |                                             ^
96.2	D:/projects/TheRock/math-libs/hipCUB/test/hipcub/test_hipcub_block_radix_rank.cpp:146:18: note: in instantiation of template class 'hipcub::BlockRadixRank<37, 4, false>' requested here
96.2	  146 |         typename RankType::TempStorage         rank;
96.2	      |                  ^
96.2	D:/projects/TheRock/math-libs/hipCUB/test/hipcub/test_hipcub_block_radix_rank.cpp:143:16: note: in instantiation of member class '(unnamed union at D:\projects\TheRock\math-libs\hipCUB\test\hipcub\test_hipcub_block_radix_rank.cpp:143:16)' requested here
96.2	  143 |     __shared__ union
96.2	      |                ^
96.2	D:/projects/TheRock/math-libs/hipCUB/test/hipcub/test_hipcub_block_radix_rank.cpp:278:44: note: in instantiation of function template specialization 'rank_kernel<37U, 1U, 4U, false, RadixRankAlgorithm::RADIX_RANK_BASIC, float>' requested here
96.2	  278 |         hipLaunchKernelGGL(HIP_KERNEL_NAME(rank_kernel<block_size,
96.2	      |                                            ^
96.2	D:/projects/TheRock/build/math-libs/rocPRIM/dist/include\hip/amd_detail/amd_hip_runtime.h:192:30: note: expanded from macro 'HIP_KERNEL_NAME'
96.2	  192 | #define HIP_KERNEL_NAME(...) __VA_ARGS__
96.2	      |                              ^
96.2	D:/projects/TheRock/build/math-libs/rocPRIM/dist/include\hip/amd_detail/amd_hip_runtime.h:244:74: note: expanded from macro 'hipLaunchKernelGGL'
96.2	  244 | #define hipLaunchKernelGGL(kernelName, ...)  hipLaunchKernelGGLInternal((kernelName), __VA_ARGS__)
96.2	      |                                                                          ^
96.2	D:/projects/TheRock/build/math-libs/rocPRIM/dist/include\hip/amd_detail/amd_hip_runtime.h:241:9: note: expanded from macro 'hipLaunchKernelGGLInternal'
96.2	  241 |         kernelName<<<(numBlocks), (numThreads), (memPerBlock), (streamId)>>>(__VA_ARGS__);         \
96.2	      |         ^
96.2	D:/projects/TheRock/math-libs/hipCUB/test/hipcub/test_hipcub_block_radix_rank.cpp:314:5: note: in instantiation of function template specialization 'test_radix_rank<HipcubBlockRadixRank<params<float, 37, 1>>, RadixRankAlgorithm::RADIX_RANK_BASIC>' requested here
96.2	  314 |     test_radix_rank<TestFixture, RadixRankAlgorithm::RADIX_RANK_BASIC>();
96.2	      |     ^
96.2	D:/projects/TheRock/build/third-party/googletest/dist/include\gtest/internal/gtest-internal.h:456:44: note: in instantiation of member function 'HipcubBlockRadixRank_BlockRadixRankBasic_Test<params<float, 37, 1>>::TestBody' requested here
96.2	  456 |   Test* CreateTest() override { return new TestClass; }
96.2	      |                                            ^
96.2	D:/projects/TheRock/build/third-party/googletest/dist/include\gtest/internal/gtest-internal.h:723:13: note: (skipping 3 contexts in backtrace; use -ftemplate-backtrace-limit=0 to see all)
96.2	  723 |         new TestFactoryImpl<TestClass>);
96.2	      |             ^
96.2	D:/projects/TheRock/build/third-party/googletest/dist/include\gtest/internal/gtest-internal.h:727:9: note: in instantiation of member function 'testing::internal::TypeParameterizedTest<HipcubBlockRadixRank, testing::internal::TemplateSel<HipcubBlockRadixRank_BlockRadixRankBasic_Test>, testing::internal::Types<params<unsigned short, 512, 1, true>, params<double, 65, 1>, params<float, 37, 1>, params<long long, 510, 1, true>, params<unsigned int, 162, 1>, params<unsigned char, 255, 1>, params<unsigned long long, 64, 2, true>, params<int, 128, 4>, params<unsigned short, 256, 7>, params<float, 512, 2>, params<double, 33, 5>, params<char, 464, 2, true>, params<unsigned short, 100, 3>, params<__half, 234, 9>, params<hip_bfloat16, 234, 9>, params<unsigned long long, 64, 1, false, 8, 5>, params<unsigned short, 102, 3, true, 4, 3>, params<float, 60, 1, true, 8, 3>, params<unsigned int, 162, 2, true, 3, 6, 2>, params<__half, 193, 2, true, 1, 4, 3>, params<hip_bfloat16, 193, 2, true, 1, 4, 3>>>::Register' requested here
96.2	  727 |         Register(prefix, std::move(code_location), case_name, test_names,
96.2	      |         ^
96.2	D:/projects/TheRock/build/third-party/googletest/dist/include\gtest/internal/gtest-internal.h:727:9: note: in instantiation of member function 'testing::internal::TypeParameterizedTest<HipcubBlockRadixRank, testing::internal::TemplateSel<HipcubBlockRadixRank_BlockRadixRankBasic_Test>, testing::internal::Types<params<float, 256, 1>, params<unsigned short, 512, 1, true>, params<double, 65, 1>, params<float, 37, 1>, params<long long, 510, 1, true>, params<unsigned int, 162, 1>, params<unsigned char, 255, 1>, params<unsigned long long, 64, 2, true>, params<int, 128, 4>, params<unsigned short, 256, 7>, params<float, 512, 2>, params<double, 33, 5>, params<char, 464, 2, true>, params<unsigned short, 100, 3>, params<__half, 234, 9>, params<hip_bfloat16, 234, 9>, params<unsigned long long, 64, 1, false, 8, 5>, params<unsigned short, 102, 3, true, 4, 3>, params<float, 60, 1, true, 8, 3>, params<unsigned int, 162, 2, true, 3, 6, 2>, params<__half, 193, 2, true, 1, 4, 3>, params<hip_bfloat16, 193, 2, true, 1, 4, 3>>>::Register' requested here
96.2	D:/projects/TheRock/build/third-party/googletest/dist/include\gtest/internal/gtest-internal.h:727:9: note: in instantiation of member function 'testing::internal::TypeParameterizedTest<HipcubBlockRadixRank, testing::internal::TemplateSel<HipcubBlockRadixRank_BlockRadixRankBasic_Test>, testing::internal::Types<params<hip_bfloat16, 128, 1>, params<float, 256, 1>, params<unsigned short, 512, 1, true>, params<double, 65, 1>, params<float, 37, 1>, params<long long, 510, 1, true>, params<unsigned int, 162, 1>, params<unsigned char, 255, 1>, params<unsigned long long, 64, 2, true>, params<int, 128, 4>, params<unsigned short, 256, 7>, params<float, 512, 2>, params<double, 33, 5>, params<char, 464, 2, true>, params<unsigned short, 100, 3>, params<__half, 234, 9>, params<hip_bfloat16, 234, 9>, params<unsigned long long, 64, 1, false, 8, 5>, params<unsigned short, 102, 3, true, 4, 3>, params<float, 60, 1, true, 8, 3>, params<unsigned int, 162, 2, true, 3, 6, 2>, params<__half, 193, 2, true, 1, 4, 3>, params<hip_bfloat16, 193, 2, true, 1, 4, 3>>>::Register' requested here
96.2	D:/projects/TheRock/build/third-party/googletest/dist/include\gtest/internal/gtest-internal.h:727:9: note: in instantiation of member function 'testing::internal::TypeParameterizedTest<HipcubBlockRadixRank, testing::internal::TemplateSel<HipcubBlockRadixRank_BlockRadixRankBasic_Test>, testing::internal::Types<params<__half, 128, 1>, params<hip_bfloat16, 128, 1>, params<float, 256, 1>, params<unsigned short, 512, 1, true>, params<double, 65, 1>, params<float, 37, 1>, params<long long, 510, 1, true>, params<unsigned int, 162, 1>, params<unsigned char, 255, 1>, params<unsigned long long, 64, 2, true>, params<int, 128, 4>, params<unsigned short, 256, 7>, params<float, 512, 2>, params<double, 33, 5>, params<char, 464, 2, true>, params<unsigned short, 100, 3>, params<__half, 234, 9>, params<hip_bfloat16, 234, 9>, params<unsigned long long, 64, 1, false, 8, 5>, params<unsigned short, 102, 3, true, 4, 3>, params<float, 60, 1, true, 8, 3>, params<unsigned int, 162, 2, true, 3, 6, 2>, params<__half, 193, 2, true, 1, 4, 3>, params<hip_bfloat16, 193, 2, true, 1, 4, 3>>>::Register' requested here
96.2	D:/projects/TheRock/math-libs/hipCUB/test/hipcub/test_hipcub_block_radix_rank.cpp:312:1: note: in instantiation of member function 'testing::internal::TypeParameterizedTest<HipcubBlockRadixRank, testing::internal::TemplateSel<HipcubBlockRadixRank_BlockRadixRankBasic_Test>, testing::internal::Types<params<unsigned int, 64, 1>, params<__half, 128, 1>, params<hip_bfloat16, 128, 1>, params<float, 256, 1>, params<unsigned short, 512, 1, true>, params<double, 65, 1>, params<float, 37, 1>, params<long long, 510, 1, true>, params<unsigned int, 162, 1>, params<unsigned char, 255, 1>, params<unsigned long long, 64, 2, true>, params<int, 128, 4>, params<unsigned short, 256, 7>, params<float, 512, 2>, params<double, 33, 5>, params<char, 464, 2, true>, params<unsigned short, 100, 3>, params<__half, 234, 9>, params<hip_bfloat16, 234, 9>, params<unsigned long long, 64, 1, false, 8, 5>, params<unsigned short, 102, 3, true, 4, 3>, params<float, 60, 1, true, 8, 3>, params<unsigned int, 162, 2, true, 3, 6, 2>, params<__half, 193, 2, true, 1, 4, 3>, params<hip_bfloat16, 193, 2, true, 1, 4, 3>>>::Register' requested here
96.2	  312 | TYPED_TEST(HipcubBlockRadixRank, BlockRadixRankBasic)
96.2	      | ^
96.2	D:/projects/TheRock/build/third-party/googletest/dist/include\gtest/gtest-typed-test.h:215:31: note: expanded from macro 'TYPED_TEST'
96.2	  215 |                   CaseName)>::Register("",                                   \
96.2	      |                               ^
96.2	D:/projects/TheRock/build/math-libs/rocPRIM/dist/include\rocprim/device\../detail/various.hpp:195:10: note: 'raw_storage<rocprim::detail::block_radix_rank<37, 4>::storage_type_>' has been explicitly marked deprecated here
96.2	  195 | struct [[deprecated("To store non default-constructible types in local memory, use "
96.2	      |          ^

Operating System

Linux and Windows

CPU

N/A

GPU

N/A

ROCm Version

6.5 (built from TheRock with recent amd-llvm)

ROCm Component

rocPRIM

Steps to Reproduce

1. Checkout code from around Bump amd-llvm and HIPIFY to track amd-staging. TheRock#725, remove any -Wno-deprecated-pragma -Wno-deprecated-declarations code
2. Build rocPRIM
3. Watch logs in rocPRIM_build.log

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

No response

Additional Information

https://github.com/search?q=repo%3AROCm%2Frocm-libraries%20raw_storage&type=code shows this deprecated struct:

rocm-libraries/projects/rocprim/rocprim/include/rocprim/detail/various.hpp

Lines 193 to 197 in c925588

	// A storage-backing wrapper that allows types with non-trivial constructors to be aliased in unions
	template<typename T>
	struct [[deprecated("To store non default-constructible types in local memory, use "
	"rocprim::uninitialized_array instead")]] raw_storage
	{

and multiple uses within the same rocprim project:

rocm-libraries/projects/rocprim/rocprim/include/rocprim/block/detail/block_sort_merge.hpp

Lines 56 to 63 in c925588

	template<bool with_values>
	union storage_type_
	{
	typename warp_sort_type::storage_type warp_sort;
	ROCPRIM_DETAIL_SUPPRESS_DEPRECATION_WITH_PUSH
	detail::raw_storage<Key[ItemsPerBlock]> keys;
	ROCPRIM_DETAIL_SUPPRESS_DEPRECATION_POP
	};

I would recommend

1. Migrating from the deprecated struct to the recommended alternative
2. Never commiting a deprecation warning to a project without also migrating at least that current project
3. Ideally (now in the monorepo) roll out such migrations as an atomic commit across all projects, skipping the deprecation step entirely
4. Building with -Werror and enforcing that no new warnings as errors are introduced as a pre-merge check

[ScottTodd]

#192 is expected to fix this. Thanks @Naraenda and @stanleytsang-amd !