https://github.com/jmmartinez updated https://github.com/llvm/llvm-project/pull/165388
From 8d08a6ef45d037b092fb7f542af768c617e4a08e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <[email protected]> Date: Tue, 28 Oct 2025 13:29:49 +0100 Subject: [PATCH] [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. stack-info: PR: https://github.com/llvm/llvm-project/pull/165388, branch: users/jmmartinez/fix/load_lds_typesignature/2 --- .../SemaHIP/amdgpu-gfx950-load-to-lds.hip | 27 ++++++++++++++++++- 1 file changed, 26 insertions(+), 1 deletion(-) 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); +} _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
