16 #define CK_TILE_S_CNT_MAX 0b1100'1111'0111'1111
17 #define CK_TILE_VMCNT(cnt) \
18 ([]() { static_assert(!((cnt) >> 6), "VMCNT only has 6 bits"); }(), \
19 ((cnt) & 0b1111) | (((cnt) & 0b110000) << 10))
20 #define CK_TILE_EXPCNT(cnt) \
21 ([]() { static_assert(!((cnt) >> 3), "EXP only has 3 bits"); }(), ((cnt) << 4))
22 #define CK_TILE_LGKMCNT(cnt) \
23 ([]() { static_assert(!((cnt) >> 4), "LGKM only has 4 bits"); }(), ((cnt) << 8))
27 template <
typename,
bool>
28 struct safe_underlying_type;
31 struct safe_underlying_type<T, true>
33 using type = std::underlying_type_t<T>;
37 struct safe_underlying_type<T, false>
65 #if defined(__GFX9__) || !defined(__HIP_DEVICE_COMPILE__)
74 hipDeviceProp_t props{};
76 auto status = hipGetDevice(&device);
77 if(status != hipSuccess)
81 status = hipGetDeviceProperties(&props, device);
82 if(status != hipSuccess)
86 return props.major > 9;
103 template <
bool ReturnSgpr = true>
107 if constexpr(ReturnSgpr)
124 asm volatile(
"s_wait_loadcnt %0 \n"
125 "s_barrier_signal -1 \n"
131 asm volatile(
"s_waitcnt vmcnt(%0) \n"
150 template <index_t cnt>
153 static_assert(cnt >= 0 && !(cnt >> 6),
"valid range is [0..63]");
154 return MAX & ((cnt & 0b1111) | ((cnt & 0b110000) << 10));
157 template <index_t cnt>
160 static_assert(cnt >= 0 && !(cnt >> 3),
"valid range is [0..7]");
161 return MAX & (cnt << 4);
164 template <index_t cnt>
167 static_assert(cnt >= 0 && !(cnt >> 4),
"valid range is [0..15]");
168 return MAX & (cnt << 8);
172 template <
index_t vmcnt = waitcnt_arg::kMaxVmCnt,
173 index_t expcnt = waitcnt_arg::kMaxExpCnt,
174 index_t lgkmcnt = waitcnt_arg::kMaxLgkmCnt>
177 __builtin_amdgcn_s_waitcnt(waitcnt_arg::from_vmcnt<vmcnt>() |
178 waitcnt_arg::from_expcnt<expcnt>() |
179 waitcnt_arg::from_lgkmcnt<lgkmcnt>());
182 template <
index_t vmcnt = waitcnt_arg::kMaxVmCnt,
183 index_t expcnt = waitcnt_arg::kMaxExpCnt,
184 index_t lgkmcnt = waitcnt_arg::kMaxLgkmCnt>
187 s_waitcnt<vmcnt, expcnt, lgkmcnt>();
188 __builtin_amdgcn_s_barrier();
191 template <index_t lgkmcnt = 0>
194 s_waitcnt_barrier<waitcnt_arg::kMaxVmCnt, waitcnt_arg::kMaxExpCnt, lgkmcnt>();
197 template <index_t vmcnt = 0>
200 s_waitcnt_barrier<vmcnt, waitcnt_arg::kMaxExpCnt, waitcnt_arg::kMaxLgkmCnt>();
206 asm volatile(
"s_nop %0" : :
"n"(cnt) :);
208 __builtin_amdgcn_sched_barrier(cnt);
212 #define CK_CONSTANT_ADDRESS_SPACE \
213 __attribute__((address_space( \
214 static_cast<safe_underlying_type_t<address_space_enum>>(address_space_enum::constant))))
216 template <
typename T>
221 #pragma clang diagnostic push
222 #pragma clang diagnostic ignored "-Wold-style-cast"
224 #pragma clang diagnostic pop
227 template <
typename T>
232 #pragma clang diagnostic push
233 #pragma clang diagnostic ignored "-Wold-style-cast"
235 #pragma clang diagnostic pop
240 #if defined(__gfx950__)
252 case address_space_enum::generic:
return "generic";
253 case address_space_enum::global:
return "global";
254 case address_space_enum::lds:
return "lds";
255 case address_space_enum::sgpr:
return "sgpr";
256 case address_space_enum::constant:
return "constant";
257 case address_space_enum::vgpr:
return "vgpr";
258 default:
return "unknown";
272 #if defined(__gfx11__)
#define CK_CONSTANT_ADDRESS_SPACE
Definition: arch.hpp:212
constexpr CK_TILE_HOST_DEVICE const char * address_space_to_string(address_space_enum addr_space)
Helper function to convert address space enum to string.
Definition: arch.hpp:248
constexpr CK_TILE_HOST_DEVICE index_t get_smem_capacity()
Definition: arch.hpp:238
CK_TILE_DEVICE void s_nop(index_t cnt=0)
Definition: arch.hpp:203
__device__ T * cast_pointer_to_generic_address_space(T CK_CONSTANT_ADDRESS_SPACE *p)
Definition: arch.hpp:217
CK_TILE_DEVICE void block_sync_lds_direct_load()
Definition: arch.hpp:198
CK_TILE_DEVICE void s_waitcnt_barrier()
Definition: arch.hpp:185
__host__ __device__ T CK_CONSTANT_ADDRESS_SPACE * cast_pointer_to_constant_address_space(T *p)
Definition: arch.hpp:228
CK_TILE_DEVICE void block_sync_lds()
Definition: arch.hpp:192
CK_TILE_DEVICE void s_waitcnt()
Definition: arch.hpp:175
#define CK_TILE_DEVICE
Definition: config.hpp:41
#define CK_TILE_HOST
Definition: config.hpp:40
#define CK_TILE_HOST_DEVICE
Definition: config.hpp:42
Definition: cluster_descriptor.hpp:13
CK_TILE_DEVICE void atomic_add(X *p_dst, const X &x)
__device__ uint32_t amd_wave_read_first_lane(uint16_t v)
Definition: amd_buffer_addressing.hpp:2834
CK_TILE_HOST_DEVICE T add(const T &a, const T &b)
Definition: generic_memory_space_atomic.hpp:16
int32_t index_t
Definition: integer.hpp:9
__device__ index_t get_grid_size()
Definition: get_id.hpp:49
__device__ index_t get_block_size()
Definition: get_id.hpp:51
__device__ index_t get_block_1d_id()
Definition: get_id.hpp:47
__device__ index_t get_thread_global_1d_id()
Definition: get_id.hpp:43
__device__ X atomic_max(X *p_dst, const X &x)
constexpr __device__ index_t get_warp_size()
Definition: get_id.hpp:10
__device__ index_t get_thread_local_1d_id()
Definition: get_id.hpp:41
const GenericPointer< typename T::ValueType > T2 value
Definition: pointer.h:1350
unsigned short uint16_t
Definition: stdint.h:125