[NFC][HIP] Add __builtin_*_load_lds type check test cases

[jmmartinez]

This tests show how type-checking is performed for __builtin_amdgcn_load_to_lds,
but not for __builtin_amdgcn_raw_ptr_buffer_load_lds, __builtin_amdgcn_struct_ptr_buffer_load_lds and __builtin_amdgcn_global_load_lds since they are declared with the 't'
attribute.

Stacked on top of: #165387

[llvmbot]

@llvm/pr-subscribers-clang

@llvm/pr-subscribers-backend-amdgpu

Author: Juan Manuel Martinez Caamaño (jmmartinez)

Changes

[NFC][HIP] Add _builtin*_load_lds type check test cases

This tests show how typechecking is performed for
__builtin_amdgcn_load_to_lds, but not for
__builtin_amdgcn_raw_ptr_buffer_load_lds,
__builtin_amdgcn_struct_ptr_buffer_load_lds and
__builtin_amdgcn_global_load_lds since they are declared with the 't'
attribute.

---

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

1 Files Affected:

• (modified) clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip (+26-1)

diff --git a/clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip b/clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip
index b49c1866caa1c..ad8342b9fddb5 100644
--- a/clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip
+++ b/clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip
@@ -1,7 +1,6 @@
 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -fsyntax-only -triple amdgcn -target-cpu gfx950 -verify %s -fcuda-is-device
 // RUN: %clang_cc1 -fsyntax-only -triple x86_64 -aux-triple amdgcn -verify %s
-// expected-no-diagnostics
 
 #define __device__ __attribute__((device))
 #define __global__ __attribute__((global))
@@ -58,3 +57,29 @@ __global__ void i_am_kernel(void* src, __amdgpu_buffer_rsrc_t rsrc, __shared__ v
     __builtin_amdgcn_global_load_lds(src, dst, 12, 0 , 0);
     __builtin_amdgcn_global_load_lds(src, dst, 16, 0 , 0);
 }
+
+__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_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_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}}
+    __builtin_amdgcn_load_to_lds(src, dst, 4, 0, 0, 4); // expected-error{{too many arguments to function call}}
+    __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);
+}