ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-02 21:27:45 +00:00
Code Issues Packages Projects Releases Wiki Activity

SWDEV-535598 - make ck::get_warp_size() constexpr and provide an implementation when used in constexpr contexts

Browse Source
This commit is contained in:
Gerardo Hernandez
2025-06-02 12:27:23 +01:00
parent 7dcaf46158
commit 05598cecee
1 changed files with 14 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

19
include/ck/utility/get_id.hpp
View File

@@ -9,13 +9,22 @@ namespace ck {
__host__ __device__ constexpr index_t get_warp_size()
{
hipDeviceProp_t props;
hipError_t status = hipGetDeviceProperties(&props, get_device_id());
if (std::is_constant_evaluated()) {
// this is only for cases where the use of constexpr cannot be avoided
#if defined(__GFX9__)
return 64;
#else
return 32;
#endif
} else {
hipDeviceProp_t props;
hipError_t status = hipGetDeviceProperties(&props, get_device_id());
if(status != hipSuccess)
throw std::runtime_error("Failed to get device properties when trying to get warp size");
if(status != hipSuccess)
throw std::runtime_error("Failed to get device properties when trying to get warp size");
return props.warpSize;
return props.warpSize;
}
}
__device__ index_t get_thread_local_1d_id() { return threadIdx.x; }