@@ -34,6 +34,20 @@ BUILTIN(__builtin_amdgcn_workgroup_id_x, "Ui", "nc")
3434BUILTIN(__builtin_amdgcn_workgroup_id_y, " Ui" , " nc" )
3535BUILTIN(__builtin_amdgcn_workgroup_id_z, " Ui" , " nc" )
3636
37+ TARGET_BUILTIN(__builtin_amdgcn_cluster_id_x, " Ui" , " nc" , " gfx1250-insts" )
38+ TARGET_BUILTIN(__builtin_amdgcn_cluster_id_y, " Ui" , " nc" , " gfx1250-insts" )
39+ TARGET_BUILTIN(__builtin_amdgcn_cluster_id_z, " Ui" , " nc" , " gfx1250-insts" )
40+
41+ TARGET_BUILTIN(__builtin_amdgcn_cluster_workgroup_id_x, " Ui" , " nc" , " gfx1250-insts" )
42+ TARGET_BUILTIN(__builtin_amdgcn_cluster_workgroup_id_y, " Ui" , " nc" , " gfx1250-insts" )
43+ TARGET_BUILTIN(__builtin_amdgcn_cluster_workgroup_id_z, " Ui" , " nc" , " gfx1250-insts" )
44+ TARGET_BUILTIN(__builtin_amdgcn_cluster_workgroup_flat_id, " Ui" , " nc" , " gfx1250-insts" )
45+
46+ TARGET_BUILTIN(__builtin_amdgcn_cluster_workgroup_max_id_x, " Ui" , " nc" , " gfx1250-insts" )
47+ TARGET_BUILTIN(__builtin_amdgcn_cluster_workgroup_max_id_y, " Ui" , " nc" , " gfx1250-insts" )
48+ TARGET_BUILTIN(__builtin_amdgcn_cluster_workgroup_max_id_z, " Ui" , " nc" , " gfx1250-insts" )
49+ TARGET_BUILTIN(__builtin_amdgcn_cluster_workgroup_max_flat_id, " Ui" , " nc" , " gfx1250-insts" )
50+
3751BUILTIN(__builtin_amdgcn_workitem_id_x, " Ui" , " nc" )
3852BUILTIN(__builtin_amdgcn_workitem_id_y, " Ui" , " nc" )
3953BUILTIN(__builtin_amdgcn_workitem_id_z, " Ui" , " nc" )
@@ -364,6 +378,31 @@ BUILTIN(__builtin_amdgcn_endpgm, "v", "nr")
364378BUILTIN(__builtin_amdgcn_get_fpenv, " WUi" , " n" )
365379BUILTIN(__builtin_amdgcn_set_fpenv, " vWUi" , " n" )
366380
381+ // ===----------------------------------------------------------------------===//
382+
383+ // Wave Reduction builtins.
384+
385+ // ===----------------------------------------------------------------------===//
386+
387+ BUILTIN(__builtin_amdgcn_wave_reduce_add_u32, " ZUiZUiZi" , " nc" )
388+ BUILTIN(__builtin_amdgcn_wave_reduce_sub_u32, " ZUiZUiZi" , " nc" )
389+ BUILTIN(__builtin_amdgcn_wave_reduce_min_i32, " ZiZiZi" , " nc" )
390+ BUILTIN(__builtin_amdgcn_wave_reduce_min_u32, " ZUiZUiZi" , " nc" )
391+ BUILTIN(__builtin_amdgcn_wave_reduce_max_i32, " ZiZiZi" , " nc" )
392+ BUILTIN(__builtin_amdgcn_wave_reduce_max_u32, " ZUiZUiZi" , " nc" )
393+ BUILTIN(__builtin_amdgcn_wave_reduce_and_b32, " ZiZiZi" , " nc" )
394+ BUILTIN(__builtin_amdgcn_wave_reduce_or_b32, " ZiZiZi" , " nc" )
395+ BUILTIN(__builtin_amdgcn_wave_reduce_xor_b32, " ZiZiZi" , " nc" )
396+ BUILTIN(__builtin_amdgcn_wave_reduce_add_u64, " WUiWUiZi" , " nc" )
397+ BUILTIN(__builtin_amdgcn_wave_reduce_sub_u64, " WUiWUiZi" , " nc" )
398+ BUILTIN(__builtin_amdgcn_wave_reduce_min_i64, " WiWiZi" , " nc" )
399+ BUILTIN(__builtin_amdgcn_wave_reduce_min_u64, " WUiWUiZi" , " nc" )
400+ BUILTIN(__builtin_amdgcn_wave_reduce_max_i64, " WiWiZi" , " nc" )
401+ BUILTIN(__builtin_amdgcn_wave_reduce_max_u64, " WUiWUiZi" , " nc" )
402+ BUILTIN(__builtin_amdgcn_wave_reduce_and_b64, " WiWiZi" , " nc" )
403+ BUILTIN(__builtin_amdgcn_wave_reduce_or_b64, " WiWiZi" , " nc" )
404+ BUILTIN(__builtin_amdgcn_wave_reduce_xor_b64, " WiWiZi" , " nc" )
405+
367406// ===----------------------------------------------------------------------===//
368407// R600-NI only builtins.
369408// ===----------------------------------------------------------------------===//
0 commit comments