@@ -357,24 +357,24 @@ BUILTIN(__builtin_amdgcn_set_fpenv, "vWUi", "n")
357
357
358
358
// ===----------------------------------------------------------------------===//
359
359
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" )
378
378
379
379
// ===----------------------------------------------------------------------===//
380
380
// R600-NI only builtins.
0 commit comments