@@ -357,24 +357,24 @@ BUILTIN(__builtin_amdgcn_set_fpenv, "vWUi", "n")
357357
358358// ===----------------------------------------------------------------------===//
359359
360- BUILTIN(__builtin_amdgcn_wave_reduce_add_i32, " iii " , " nc" )
361- BUILTIN(__builtin_amdgcn_wave_reduce_sub_i32, " iii " , " nc" )
362- BUILTIN(__builtin_amdgcn_wave_reduce_min_i32, " iii " , " nc" )
363- BUILTIN(__builtin_amdgcn_wave_reduce_min_u32, " UiUii " , " nc" )
364- BUILTIN(__builtin_amdgcn_wave_reduce_max_i32, " iii " , " nc" )
365- BUILTIN(__builtin_amdgcn_wave_reduce_max_u32, " UiUii " , " nc" )
366- BUILTIN(__builtin_amdgcn_wave_reduce_and_b32, " iii " , " nc" )
367- BUILTIN(__builtin_amdgcn_wave_reduce_or_b32, " iii " , " nc" )
368- BUILTIN(__builtin_amdgcn_wave_reduce_xor_b32, " iii " , " nc" )
369- BUILTIN(__builtin_amdgcn_wave_reduce_add_i64, " WiWii " , " nc" )
370- BUILTIN(__builtin_amdgcn_wave_reduce_sub_i64, " WiWii " , " nc" )
371- BUILTIN(__builtin_amdgcn_wave_reduce_min_i64, " WiWii " , " nc" )
372- BUILTIN(__builtin_amdgcn_wave_reduce_min_u64, " WUiWUii " , " nc" )
373- BUILTIN(__builtin_amdgcn_wave_reduce_max_i64, " WiWii " , " nc" )
374- BUILTIN(__builtin_amdgcn_wave_reduce_max_u64, " WUiWUii " , " nc" )
375- BUILTIN(__builtin_amdgcn_wave_reduce_and_b64, " WiWii " , " nc" )
376- BUILTIN(__builtin_amdgcn_wave_reduce_or_b64, " WiWii " , " nc" )
377- BUILTIN(__builtin_amdgcn_wave_reduce_xor_b64, " WiWii " , " nc" )
360+ BUILTIN(__builtin_amdgcn_wave_reduce_add_i32, " ZiZiZi " , " nc" )
361+ BUILTIN(__builtin_amdgcn_wave_reduce_sub_i32, " ZiZiZi " , " nc" )
362+ BUILTIN(__builtin_amdgcn_wave_reduce_min_i32, " ZiZiZi " , " nc" )
363+ BUILTIN(__builtin_amdgcn_wave_reduce_min_u32, " ZUiZUiZi " , " nc" )
364+ BUILTIN(__builtin_amdgcn_wave_reduce_max_i32, " ZiZiZi " , " nc" )
365+ BUILTIN(__builtin_amdgcn_wave_reduce_max_u32, " ZUiZUiZi " , " nc" )
366+ BUILTIN(__builtin_amdgcn_wave_reduce_and_b32, " ZiZiZi " , " nc" )
367+ BUILTIN(__builtin_amdgcn_wave_reduce_or_b32, " ZiZiZi " , " nc" )
368+ BUILTIN(__builtin_amdgcn_wave_reduce_xor_b32, " ZiZiZi " , " nc" )
369+ BUILTIN(__builtin_amdgcn_wave_reduce_add_i64, " WiWiZi " , " nc" )
370+ BUILTIN(__builtin_amdgcn_wave_reduce_sub_i64, " WiWiZi " , " nc" )
371+ BUILTIN(__builtin_amdgcn_wave_reduce_min_i64, " WiWiZi " , " nc" )
372+ BUILTIN(__builtin_amdgcn_wave_reduce_min_u64, " WUiWUiZi " , " nc" )
373+ BUILTIN(__builtin_amdgcn_wave_reduce_max_i64, " WiWiZi " , " nc" )
374+ BUILTIN(__builtin_amdgcn_wave_reduce_max_u64, " WUiWUiZi " , " nc" )
375+ BUILTIN(__builtin_amdgcn_wave_reduce_and_b64, " WiWiZi " , " nc" )
376+ BUILTIN(__builtin_amdgcn_wave_reduce_or_b64, " WiWiZi " , " nc" )
377+ BUILTIN(__builtin_amdgcn_wave_reduce_xor_b64, " WiWiZi " , " nc" )
378378
379379// ===----------------------------------------------------------------------===//
380380// R600-NI only builtins.
0 commit comments