4 Commits

Author SHA1 Message Date
Anshil Gandhi
a955a31896 [AMDGPU] Replace target feature for global fadd32
Change target feature of __builtin_amdgcn_global_atomic_fadd_f32
to atomic-fadd-rtn-insts. Enable atomic-fadd-rtn-insts for gfx90a,
gfx940 and gfx1100 as they all support the return variant of
`global_atomic_add_f32`.

Fixes https://github.com/llvm/llvm-project/issues/61331.

Reviewed By: rampitec

Differential Revision: https://reviews.llvm.org/D146840
2023-03-28 15:58:30 -06:00
Mariusz Sikora
69061f9627 [AMDGPU] Add clang builtin for __builtin_amdgcn_ds_atomic_fadd_v2f16
Differential Revision: https://reviews.llvm.org/D146808
2023-03-24 16:27:44 +01:00
Mariusz Sikora
ea064ee2a3 [AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions
Introducing Subtarget Features for instructions:
- ds_pk_add_bf16
- ds_pk_add_f16
- ds_pk_add_rtn_bf16
- ds_pk_add_rtn_f16
- flat_atomic_pk_add_f16
- flat_atomic_pk_add_bf16
- global_atomic_pk_add_f16
- global_atomic_pk_add_bf16
- buffer_atomic_pk_add_f16

Differential Revision: https://reviews.llvm.org/D146701
2023-03-24 13:10:40 +01:00
Anshil Gandhi
39dac1f7f6 [clang] Add clang builtins support for gfx90a
Implement target builtins for gfx90a including fadd64, fadd32, add2h,
max and min on various global, flat and ds address spaces for which
intrinsics are implemented.

Differential Revision: https://reviews.llvm.org/D106909
2021-08-05 02:08:06 -06:00