__AMDGCN_WAVEFRONT_SIZE__ and __AMDGCN_WAVEFRONT_SIZE to be deprecated in future ROCM release

[mrowan137]

As an FYI, the macros __AMDGCN_WAVEFRONT_SIZE__ and __AMDGCN_WAVEFRONT_SIZE are planned to be deprecated in a future release of ROCm (mentioned in ROCm 6.3 release notes). I am not aware of any suitable way for applications to detect the wavefront size at compile time, so for the time being there may be no better alternative than the current hardcoding.
(This check will become problematic at some point: https://github.com/AMReX-Codes/amrex/blob/b3f67385e62f387b548389222840486c0fffca57/Src/Base/AMReX_GpuDevice.cpp#L67C1-L73C7)
@WeiqunZhang

[WeiqunZhang]

@mrowan137 Thanks for letting us know. I guess we will do the static_assert only when the macros are defined.

[WeiqunZhang]

@mrowan137 https://rocm.docs.amd.com/projects/rocPRIM/en/latest/reference/intrinsics.html#_CPPv4N7rocprim9warp_sizeEv

It seems that we can use __device__ inline constexpr unsigned int rocprim::device_warp_size() at compile time on device and __host__ inline hipError_t rocprim::host_warp_size(const hipStream_t stream, unsigned int &warp_size) on host at runtime.

[WeiqunZhang]

Will they become deprecated as well?

[mrowan137]

Yes, those too will become deprecated. It remains to be seen whether they're replaced with something, but for the time being it may be safest to assume that any workaround based on a ROCm library would also be deprecated.