[Clang] Remove 't' from __builtin_amdgcn_ds_atomic_fadd_f32/f64

[jmmartinez]

Allows for type checking depending on the builtin signature.

[llvmbot]

@llvm/pr-subscribers-clang

@llvm/pr-subscribers-backend-amdgpu

Author: Juan Manuel Martinez Caamaño (jmmartinez)

Changes

Allows for type checking depending on the builtin signature.

---

Full diff: https://github.com/llvm/llvm-project/pull/172293.diff

2 Files Affected:

• (modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+2-2)
• (added) clang/test/SemaHIP/amdgpu-ds-atomic-fadd.hip (+33)

diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index a867144d83928..88b306462a92c 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -276,8 +276,8 @@ TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f64, "dd*0d", "t", "gfx90a-inst
 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")
 
-TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f64, "dd*3d", "t", "gfx90a-insts")
-TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f32, "ff*3f", "t", "gfx8-insts")
+TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f64, "dd*3d", "", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f32, "ff*3f", "", "gfx8-insts")
 
 TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f32, "ff*0f", "t", "gfx940-insts")
 TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2f16, "V2hV2h*0V2h", "t", "atomic-flat-pk-add-16-insts")
diff --git a/clang/test/SemaHIP/amdgpu-ds-atomic-fadd.hip b/clang/test/SemaHIP/amdgpu-ds-atomic-fadd.hip
new file mode 100644
index 0000000000000..840b26e52cc25
--- /dev/null
+++ b/clang/test/SemaHIP/amdgpu-ds-atomic-fadd.hip
@@ -0,0 +1,33 @@
+// RUN: %clang_cc1 -fsyntax-only -triple amdgcn -target-cpu gfx90a -verify=device,both %s -fcuda-is-device
+// RUN: %clang_cc1 -fsyntax-only -triple x86_64 -aux-triple amdgcn -verify=host,both %s
+
+#define __device__ __attribute__((device))
+#define __shared__ __attribute__((shared))
+
+__device__ void test_ds_atomic_fadd_f32_valid(__shared__ float *lds_ptr, float val) {
+  float result;
+  result = __builtin_amdgcn_ds_atomic_fadd_f32(lds_ptr, val);
+}
+
+__device__ void test_ds_atomic_fadd_f32_errors(__shared__ float *lds_ptr, float val, 
+                                                __shared__ double *lds_ptr_d,
+                                                float *global_ptr) {
+  float result;
+  result = __builtin_amdgcn_ds_atomic_fadd_f32(lds_ptr, val, 0); // both-error{{too many arguments to function call, expected 2, have 3}}
+  result = __builtin_amdgcn_ds_atomic_fadd_f32(global_ptr, val);
+  result = __builtin_amdgcn_ds_atomic_fadd_f32(lds_ptr_d, val); // device-error{{cannot initialize a parameter of type '__shared__ float *' with an rvalue of type '__shared__ double *'}} host-error{{cannot initialize a parameter of type '__attribute__((address_space(3))) float *' with an rvalue of type '__attribute__((address_space(3))) double *'}}
+}
+
+__device__ void test_ds_atomic_fadd_f64_valid(__shared__ double *lds_ptr, double val) {
+  double result;
+  result = __builtin_amdgcn_ds_atomic_fadd_f64(lds_ptr, val);
+}
+
+__device__ void test_ds_atomic_fadd_f64_errors(__shared__ double *lds_ptr, double val,
+                                                __shared__ float *lds_ptr_f,
+                                                double *global_ptr) {
+  double result;
+  result = __builtin_amdgcn_ds_atomic_fadd_f64(lds_ptr, val, 0); // both-error{{too many arguments to function call, expected 2, have 3}}
+  result = __builtin_amdgcn_ds_atomic_fadd_f64(global_ptr, val);
+  result = __builtin_amdgcn_ds_atomic_fadd_f64(lds_ptr_f, val); // device-error{{cannot initialize a parameter of type '__shared__ double *' with an rvalue of type '__shared__ float *'}} host-error{{cannot initialize a parameter of type '__attribute__((address_space(3))) double *' with an rvalue of type '__attribute__((address_space(3))) float *'}}
+}

[arsenm]

I'm assuming we already have OpenCL tests for these?

[jmmartinez]

I'm assuming we already have OpenCL tests for these?

There are, but not for the type-checking (only code-gen and target feature support). I'll add some extra tests similar to the ones I've added for HIP.

[arsenm]

I'm assuming we already have OpenCL tests for these?

There are, but not for the type-checking (only code-gen and target feature support). I'll add some extra tests similar to the ones I've added for HIP.

codegen is probably good enough. We kind of need both for all builtins that use pointer arguments

[shiltian]

Great to see that we are removing more. :-D We should probably remove all of them.

[ranapratap55]

We use -err.cl suffix for error check test files.

[jmmartinez]

Great to see that we are removing more. :-D We should probably remove all of them.

There are some of them make some sense. There are __builtin_amdgcn_mov_dpp8/__builtin_amdgcn_mov_dpp/__builtin_amdgcn_update_dpp ones that take are polymorphic (in the tests they take several types, but fall into i32/i64, I haven't played with them).

There is also some bfloat16 ones that are awkward++: they take short. I'm not sure what the support for bfloat16 as a value type is for OpenCL (some extensions represent them as short while others as ushort).

[arsenm]

We have a bunch of intrinsics that could use the work to migrate from short to bfloat