[HIP][AMDGPU] Remove 't' from all __builtin_*_load_lds builtins

[jmmartinez]

Allows for type checking depending on the builtin signature.

Stacked on top of: #165387 and #165388

[llvmbot]

@llvm/pr-subscribers-backend-amdgpu

@llvm/pr-subscribers-clang

Author: Juan Manuel Martinez Caamaño (jmmartinez)

Changes

[HIP][AMDGPU] Remove 't' from all _builtin*_load_lds builtins

Allows for type checking depending on the builtin signature.

---

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

2 Files Affected:

• (modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+3-3)
• (modified) clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip (+15-15)

diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index f265d82efee75..0a4deb182d80f 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -187,8 +187,8 @@ TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32, "ffQbiiIi", "",
 TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64, "ddQbiiIi", "", "atomic-fmin-fmax-global-f64")
 TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64, "ddQbiiIi", "", "atomic-fmin-fmax-global-f64")
 
-TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_load_lds, "vQbv*3IUiiiIiIi", "t", "vmem-to-lds-load-insts")
-TARGET_BUILTIN(__builtin_amdgcn_struct_ptr_buffer_load_lds, "vQbv*3IUiiiiIiIi", "t", "vmem-to-lds-load-insts")
+TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_load_lds, "vQbv*3IUiiiIiIi", "", "vmem-to-lds-load-insts")
+TARGET_BUILTIN(__builtin_amdgcn_struct_ptr_buffer_load_lds, "vQbv*3IUiiiiIiIi", "", "vmem-to-lds-load-insts")
 
 //===----------------------------------------------------------------------===//
 // Ballot builtins.
@@ -286,7 +286,7 @@ TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2bf16, "V2sV2s*1V2s", "t", "
 TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s", "t", "atomic-ds-pk-add-16-insts")
 TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2hV2h*3V2h", "t", "atomic-ds-pk-add-16-insts")
 TARGET_BUILTIN(__builtin_amdgcn_load_to_lds, "vv*v*3IUiIiIUi", "", "vmem-to-lds-load-insts")
-TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "t", "vmem-to-lds-load-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "", "vmem-to-lds-load-insts")
 
 //===----------------------------------------------------------------------===//
 // Deep learning builtins.
diff --git a/clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip b/clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip
index ad8342b9fddb5..509906d8c87a8 100644
--- a/clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip
+++ b/clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip
@@ -59,17 +59,17 @@ __global__ void i_am_kernel(void* src, __amdgpu_buffer_rsrc_t rsrc, __shared__ v
 }
 
 __device__ void i_am_wrong(void* src, __amdgpu_buffer_rsrc_t rsrc, __shared__ void* dst, int vindex, int voffset, int soffset) {
-    __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 1, voffset, soffset, 0, 0, 4);
-    __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 2, voffset, soffset, 0, 0, 4);
-    __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 4, voffset, soffset, 0, 0, 4);
-    __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 12, voffset, soffset, 0, 0, 4);
-    __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 16, voffset, soffset, 0, 0, 4);
+    __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 1, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}
+    __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 2, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}
+    __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 4, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}
+    __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 12, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}
+    __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 16, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}
 
-    __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 1, vindex, voffset, soffset, 0, 0, 4);
-    __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 2, vindex, voffset, soffset, 0, 0, 4);
-    __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 4, vindex, voffset, soffset, 0, 0, 4);
-    __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 12, vindex, voffset, soffset, 0, 0, 4);
-    __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 16, vindex, voffset, soffset, 0, 0, 4);
+    __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 1, vindex, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}
+    __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 2, vindex, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}
+    __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 4, vindex, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}
+    __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 12, vindex, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}
+    __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 16, vindex, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}
 
     __builtin_amdgcn_load_to_lds(src, dst, 1, 0, 0, 4); // expected-error{{too many arguments to function call}}
     __builtin_amdgcn_load_to_lds(src, dst, 2, 0, 0, 4); // expected-error{{too many arguments to function call}}
@@ -77,9 +77,9 @@ __device__ void i_am_wrong(void* src, __amdgpu_buffer_rsrc_t rsrc, __shared__ vo
     __builtin_amdgcn_load_to_lds(src, dst, 12, 0, 0, 4); // expected-error{{too many arguments to function call}}
     __builtin_amdgcn_load_to_lds(src, dst, 16, 0, 0, 4); // expected-error{{too many arguments to function call}}
 
-    __builtin_amdgcn_global_load_lds(src, dst, 1, 0 , 0, 4);
-    __builtin_amdgcn_global_load_lds(src, dst, 2, 0 , 0, 4);
-    __builtin_amdgcn_global_load_lds(src, dst, 4, 0 , 0, 4);
-    __builtin_amdgcn_global_load_lds(src, dst, 12, 0 , 0, 4);
-    __builtin_amdgcn_global_load_lds(src, dst, 16, 0 , 0, 4);
+    __builtin_amdgcn_global_load_lds(src, dst, 1, 0 , 0, 4); // expected-error{{too many arguments to function call}}
+    __builtin_amdgcn_global_load_lds(src, dst, 2, 0 , 0, 4); // expected-error{{too many arguments to function call}}
+    __builtin_amdgcn_global_load_lds(src, dst, 4, 0 , 0, 4); // expected-error{{too many arguments to function call}}
+    __builtin_amdgcn_global_load_lds(src, dst, 12, 0 , 0, 4); // expected-error{{too many arguments to function call}}
+    __builtin_amdgcn_global_load_lds(src, dst, 16, 0 , 0, 4); // expected-error{{too many arguments to function call}}
 }