[AMDGPU] Replace target feature for global fadd32

[gandhi21299]

Summary

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.

[rampitec]

You cannot just enable it on gfx908 which does not have return version of it.

[gandhi21299]

• gfx908 does not support return version of the builtin

[gandhi21299]

• removed builtins-amdgcn-gfx908.cl

[rampitec]

Can you please also add gfx90a and gfx940 tests?

Otherwise LGTM *if* @b-sumner has no objections.

[gandhi21299]

• Adding tests for gfx90a and gfx940

[rampitec]

LGTM. Please wait for @b-sumner.

[gandhi21299]

Sounds good, thanks for the review @rampitec

[b-sumner]

No objection here.

clang/include/clang/Basic/BuiltinsAMDGPU.def

 
 //===----------------------------------------------------------------------===//
 // GFX9+ only builtins.
 //===----------------------------------------------------------------------===//
 
 TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts")
 
 TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f64, "dd*1d", "t", "gfx90a-insts")
-TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f32, "ff*1f", "t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f32, "ff*1f", "t", "atomic-fadd-rtn-insts")
 TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2f16, "V2hV2h*1V2h", "t", "atomic-buffer-global-pk-add-f16-insts")
 TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmin_f64, "dd*1d", "t", "gfx90a-insts")
 TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmax_f64, "dd*1d", "t", "gfx90a-insts")
 
 TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f64, "dd*0d", "t", "gfx90a-insts")
 TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmin_f64, "dd*0d", "t", "gfx90a-insts")
 TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmax_f64, "dd*0d", "t", "gfx90a-insts")
 

clang/lib/Basic/Targets/AMDGPU.cpp

case GK_GFX1100:
       Features["dl-insts"] = true;
       Features["16-bit-insts"] = true;
       Features["dpp"] = true;
       Features["gfx8-insts"] = true;
       Features["gfx9-insts"] = true;
       Features["gfx10-insts"] = true;
       Features["gfx10-3-insts"] = true;
       Features["gfx11-insts"] = true;
+      Features["atomic-fadd-rtn-insts"] = true;
       break;
     case GK_GFX1036:
     case GK_GFX1035:
     case GK_GFX1034:
     case GK_GFX1033:
     case GK_GFX1032:
     case GK_GFX1031:
     case GK_GFX1030:
case GK_GFX940:
       Features["fp8-insts"] = true;
       Features["atomic-ds-pk-add-16-insts"] = true;
       Features["atomic-flat-pk-add-16-insts"] = true;
       Features["atomic-global-pk-add-bf16-inst"] = true;
       [[fallthrough]];
     case GK_GFX90A:
       Features["gfx90a-insts"] = true;
       Features["atomic-buffer-global-pk-add-f16-insts"] = true;
+      Features["atomic-fadd-rtn-insts"] = true;
       [[fallthrough]];
     case GK_GFX908:
       Features["dot3-insts"] = true;
       Features["dot4-insts"] = true;
       Features["dot5-insts"] = true;
       Features["dot6-insts"] = true;
       Features["mai-insts"] = true;
+      Features["atomic-fadd-rtn-insts"] = true;
       [[fallthrough]];
     case GK_GFX906:
       Features["dl-insts"] = true;
       Features["dot1-insts"] = true;
       Features["dot2-insts"] = true;
       Features["dot7-insts"] = true;
       Features["dot10-insts"] = true;
       [[fallthrough]];

clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx908-err.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx908 \
 // RUN:   -verify -S -o - %s
 
 typedef half __attribute__((ext_vector_type(2))) half2;
 
 void test_global_fadd(__global half2 *addrh2, __local half2 *addrh2l, half2 xh2,
                       __global float *addrf, float xf,
                       __global double *addr, double x) {
   half2 *half_rtn;
   float *fp_rtn;
   double *rtn;
   *half_rtn = __builtin_amdgcn_global_atomic_fadd_v2f16(addrh2, xh2); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_v2f16' needs target feature atomic-buffer-global-pk-add-f16-insts}}
-  *fp_rtn = __builtin_amdgcn_global_atomic_fadd_f32(addr, x); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_f32' needs target feature gfx90a-insts}}
   *rtn = __builtin_amdgcn_global_atomic_fadd_f64(addr, x); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_f64' needs target feature gfx90a-insts}}
   *rtn = __builtin_amdgcn_global_atomic_fmax_f64(addr, x); // expected-error{{'__builtin_amdgcn_global_atomic_fmax_f64' needs target feature gfx90a-insts}}
   *rtn = __builtin_amdgcn_global_atomic_fmin_f64(addr, x); // expected-error{{'__builtin_amdgcn_global_atomic_fmin_f64' needs target feature gfx90a-insts}}
   *rtn = __builtin_amdgcn_flat_atomic_fadd_f64(addr, x); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_f64' needs target feature gfx90a-insts}}
   *rtn = __builtin_amdgcn_flat_atomic_fmin_f64(addr, x); // expected-error{{'__builtin_amdgcn_flat_atomic_fmin_f64' needs target feature gfx90a-insts}}
   *rtn = __builtin_amdgcn_flat_atomic_fmax_f64(addr, x); // expected-error{{'__builtin_amdgcn_flat_atomic_fmax_f64' needs target feature gfx90a-insts}}
   __builtin_amdgcn_ds_atomic_fadd_v2f16(addrh2l, xh2); // expected-error{{'__builtin_amdgcn_ds_atomic_fadd_v2f16' needs target feature atomic-ds-pk-add-16-insts}}
 }

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl

void test_permlane64(global uint* out, uint a) {
   *out = __builtin_amdgcn_permlane64(a);
 }
 
 // CHECK-LABEL: @test_s_wait_event_export_ready
 // CHECK: call void @llvm.amdgcn.s.wait.event.export.ready
 void test_s_wait_event_export_ready() {
   __builtin_amdgcn_s_wait_event_export_ready();
 }
+
+// CHECK-LABEL: @test_global_add_f32
+// CHECK: call float @llvm.amdgcn.global.atomic.fadd.f32.p1.f32(ptr addrspace(1) %{{.*}}, float %{{.*}})
+void test_global_add_f32(float *rtn, global float *addr, float x) {
+  *rtn = __builtin_amdgcn_global_atomic_fadd_f32(addr, x);
+}

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx908.cl

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -S -emit-llvm -o - %s | FileCheck %s --check-prefix=IR
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -S -o - %s | FileCheck %s --check-prefix=GFX908
+
+// IR-LABEL: @test_global_add_f32
+// IR: call float @llvm.amdgcn.global.atomic.fadd.f32.p1.f32(ptr addrspace(1) %{{.*}}, float %{{.*}})
+// GFX908-LABEL: test_global_add_f32
+// GFX908: global_atomic_add_f32
+void test_global_add_f32(float *rtn, global float *addr, float x) {
+  *rtn = __builtin_amdgcn_global_atomic_fadd_f32(addr, x);
+}