@@ -95,8 +95,11 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
95
95
"cubema" => "__builtin_amdgcn_cubema",
96
96
"cubesc" => "__builtin_amdgcn_cubesc",
97
97
"cubetc" => "__builtin_amdgcn_cubetc",
98
+ "cvt.f16.bf8" => "__builtin_amdgcn_cvt_f16_bf8",
99
+ "cvt.f16.fp8" => "__builtin_amdgcn_cvt_f16_fp8",
98
100
"cvt.f32.bf8" => "__builtin_amdgcn_cvt_f32_bf8",
99
101
"cvt.f32.fp8" => "__builtin_amdgcn_cvt_f32_fp8",
102
+ "cvt.f32.fp8.e5m3" => "__builtin_amdgcn_cvt_f32_fp8_e5m3",
100
103
"cvt.off.f32.i4" => "__builtin_amdgcn_cvt_off_f32_i4",
101
104
"cvt.pk.bf8.f32" => "__builtin_amdgcn_cvt_pk_bf8_f32",
102
105
"cvt.pk.f16.bf8" => "__builtin_amdgcn_cvt_pk_f16_bf8",
@@ -181,6 +184,12 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
181
184
"dot4.f32.fp8.bf8" => "__builtin_amdgcn_dot4_f32_fp8_bf8",
182
185
"dot4.f32.fp8.fp8" => "__builtin_amdgcn_dot4_f32_fp8_fp8",
183
186
"ds.add.gs.reg.rtn" => "__builtin_amdgcn_ds_add_gs_reg_rtn",
187
+ "ds.atomic.async.barrier.arrive.b64" => {
188
+ "__builtin_amdgcn_ds_atomic_async_barrier_arrive_b64"
189
+ }
190
+ "ds.atomic.barrier.arrive.rtn.b64" => {
191
+ "__builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64"
192
+ }
184
193
"ds.bpermute" => "__builtin_amdgcn_ds_bpermute",
185
194
"ds.bpermute.fi.b32" => "__builtin_amdgcn_ds_bpermute_fi_b32",
186
195
"ds.gws.barrier" => "__builtin_amdgcn_ds_gws_barrier",
@@ -198,8 +207,32 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
198
207
"fdot2.f16.f16" => "__builtin_amdgcn_fdot2_f16_f16",
199
208
"fdot2.f32.bf16" => "__builtin_amdgcn_fdot2_f32_bf16",
200
209
"fdot2c.f32.bf16" => "__builtin_amdgcn_fdot2c_f32_bf16",
210
+ "flat.prefetch" => "__builtin_amdgcn_flat_prefetch",
201
211
"fmul.legacy" => "__builtin_amdgcn_fmul_legacy",
212
+ "global.load.async.to.lds.b128" => {
213
+ "__builtin_amdgcn_global_load_async_to_lds_b128"
214
+ }
215
+ "global.load.async.to.lds.b32" => {
216
+ "__builtin_amdgcn_global_load_async_to_lds_b32"
217
+ }
218
+ "global.load.async.to.lds.b64" => {
219
+ "__builtin_amdgcn_global_load_async_to_lds_b64"
220
+ }
221
+ "global.load.async.to.lds.b8" => "__builtin_amdgcn_global_load_async_to_lds_b8",
202
222
"global.load.lds" => "__builtin_amdgcn_global_load_lds",
223
+ "global.prefetch" => "__builtin_amdgcn_global_prefetch",
224
+ "global.store.async.from.lds.b128" => {
225
+ "__builtin_amdgcn_global_store_async_from_lds_b128"
226
+ }
227
+ "global.store.async.from.lds.b32" => {
228
+ "__builtin_amdgcn_global_store_async_from_lds_b32"
229
+ }
230
+ "global.store.async.from.lds.b64" => {
231
+ "__builtin_amdgcn_global_store_async_from_lds_b64"
232
+ }
233
+ "global.store.async.from.lds.b8" => {
234
+ "__builtin_amdgcn_global_store_async_from_lds_b8"
235
+ }
203
236
"groupstaticsize" => "__builtin_amdgcn_groupstaticsize",
204
237
"iglp.opt" => "__builtin_amdgcn_iglp_opt",
205
238
"implicit.buffer.ptr" => "__builtin_amdgcn_implicit_buffer_ptr",
@@ -291,6 +324,7 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
291
324
"s.incperflevel" => "__builtin_amdgcn_s_incperflevel",
292
325
"s.memrealtime" => "__builtin_amdgcn_s_memrealtime",
293
326
"s.memtime" => "__builtin_amdgcn_s_memtime",
327
+ "s.monitor.sleep" => "__builtin_amdgcn_s_monitor_sleep",
294
328
"s.sendmsg" => "__builtin_amdgcn_s_sendmsg",
295
329
"s.sendmsghalt" => "__builtin_amdgcn_s_sendmsghalt",
296
330
"s.setprio" => "__builtin_amdgcn_s_setprio",
@@ -300,11 +334,15 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
300
334
"s.sleep.var" => "__builtin_amdgcn_s_sleep_var",
301
335
"s.ttracedata" => "__builtin_amdgcn_s_ttracedata",
302
336
"s.ttracedata.imm" => "__builtin_amdgcn_s_ttracedata_imm",
337
+ "s.wait.asynccnt" => "__builtin_amdgcn_s_wait_asynccnt",
303
338
"s.wait.event.export.ready" => "__builtin_amdgcn_s_wait_event_export_ready",
339
+ "s.wait.tensorcnt" => "__builtin_amdgcn_s_wait_tensorcnt",
304
340
"s.waitcnt" => "__builtin_amdgcn_s_waitcnt",
305
341
"sad.hi.u8" => "__builtin_amdgcn_sad_hi_u8",
306
342
"sad.u16" => "__builtin_amdgcn_sad_u16",
307
343
"sad.u8" => "__builtin_amdgcn_sad_u8",
344
+ "sat.pk4.i4.i8" => "__builtin_amdgcn_sat_pk4_i4_i8",
345
+ "sat.pk4.u4.u8" => "__builtin_amdgcn_sat_pk4_u4_u8",
308
346
"sched.barrier" => "__builtin_amdgcn_sched_barrier",
309
347
"sched.group.barrier" => "__builtin_amdgcn_sched_group_barrier",
310
348
"sdot2" => "__builtin_amdgcn_sdot2",
@@ -346,8 +384,13 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
346
384
"smfmac.i32.16x16x64.i8" => "__builtin_amdgcn_smfmac_i32_16x16x64_i8",
347
385
"smfmac.i32.32x32x32.i8" => "__builtin_amdgcn_smfmac_i32_32x32x32_i8",
348
386
"smfmac.i32.32x32x64.i8" => "__builtin_amdgcn_smfmac_i32_32x32x64_i8",
387
+ "struct.ptr.buffer.load.lds" => "__builtin_amdgcn_struct_ptr_buffer_load_lds",
349
388
"sudot4" => "__builtin_amdgcn_sudot4",
350
389
"sudot8" => "__builtin_amdgcn_sudot8",
390
+ "tensor.load.to.lds" => "__builtin_amdgcn_tensor_load_to_lds",
391
+ "tensor.load.to.lds.d2" => "__builtin_amdgcn_tensor_load_to_lds_d2",
392
+ "tensor.store.from.lds" => "__builtin_amdgcn_tensor_store_from_lds",
393
+ "tensor.store.from.lds.d2" => "__builtin_amdgcn_tensor_store_from_lds_d2",
351
394
"udot2" => "__builtin_amdgcn_udot2",
352
395
"udot4" => "__builtin_amdgcn_udot4",
353
396
"udot8" => "__builtin_amdgcn_udot8",
@@ -6326,6 +6369,23 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
6326
6369
}
6327
6370
s390(name, full_name)
6328
6371
}
6372
+ "spv" => {
6373
+ #[allow(non_snake_case)]
6374
+ fn spv(name: &str, full_name: &str) -> &'static str {
6375
+ match name {
6376
+ // spv
6377
+ "num.subgroups" => "__builtin_spirv_num_subgroups",
6378
+ "subgroup.id" => "__builtin_spirv_subgroup_id",
6379
+ "subgroup.local.invocation.id" => {
6380
+ "__builtin_spirv_subgroup_local_invocation_id"
6381
+ }
6382
+ "subgroup.max.size" => "__builtin_spirv_subgroup_max_size",
6383
+ "subgroup.size" => "__builtin_spirv_subgroup_size",
6384
+ _ => unimplemented!("***** unsupported LLVM intrinsic {full_name}"),
6385
+ }
6386
+ }
6387
+ spv(name, full_name)
6388
+ }
6329
6389
"ve" => {
6330
6390
#[allow(non_snake_case)]
6331
6391
fn ve(name: &str, full_name: &str) -> &'static str {
0 commit comments