https://github.com/jmmartinez updated https://github.com/llvm/llvm-project/pull/135027
From 4e5736d1b90ae477642ea03dcf39231f5f0a8e59 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <juama...@amd.com> Date: Wed, 9 Apr 2025 15:11:19 +0200 Subject: [PATCH 1/4] [Clang][AMDGPU] Accept builtins in lambda declarations Sema::getCurFunctionDecl(AllowLambda = false) returns a nullptr when the lambda declaration is outside a function (for example, when used in an assignment). Using Sema::getCurFunctionDecl(AllowLambda = true) returns the declaration of the enclosing lambda. --- clang/lib/Sema/SemaAMDGPU.cpp | 2 +- .../test/SemaHIP/amdgpu-builtin-in-lambda.hip | 24 +++++++++++++++++++ 2 files changed, 25 insertions(+), 1 deletion(-) create mode 100644 clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 7fec099374152..af139e36dd2ed 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -27,7 +27,7 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, // position of memory order and scope arguments in the builtin unsigned OrderIndex, ScopeIndex; - const auto *FD = SemaRef.getCurFunctionDecl(); + const auto *FD = SemaRef.getCurFunctionDecl(true); assert(FD && "AMDGPU builtins should not be used outside of a function"); llvm::StringMap<bool> CallerFeatureMap; getASTContext().getFunctionFeatureMap(CallerFeatureMap, FD); diff --git a/clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip b/clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip new file mode 100644 index 0000000000000..7a47fbd39817e --- /dev/null +++ b/clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip @@ -0,0 +1,24 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu tahiti -fsyntax-only -fcuda-is-device -verify=expected -o - %s +// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu gfx950 -fsyntax-only -fcuda-is-device -o - %s + +#define __device__ __attribute__((device)) +#define __shared__ __attribute__((shared)) + +struct S { + static constexpr auto make_buffer_rsrc_lambda = [](void *p, short stride, int num, int flags) { + return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags); + }; + + static constexpr auto global_load_lds_lambda = [](void* src, __shared__ void *dst) { + __builtin_amdgcn_global_load_lds(src, dst, 16, 0, 0); // expected-error{{invalid size value}} expected-note{{size must be 1, 2, or 4}} + }; +}; + +__device__ __amdgpu_buffer_rsrc_t test_simple_builtin(void *p, short stride, int num, int flags) { + return S::make_buffer_rsrc_lambda(p, stride, num, flags); +} + +__device__ void test_target_dependant_builtin(void *src, __shared__ void *dst) { + S::global_load_lds_lambda(src, dst); +} From ac85b98a3097efdcb241a73c617ebe3167dd5bb7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <juama...@amd.com> Date: Thu, 10 Apr 2025 09:51:38 +0200 Subject: [PATCH 2/4] [review] AllowLambda --- clang/lib/Sema/SemaAMDGPU.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index af139e36dd2ed..a6366aceec2a6 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -27,7 +27,7 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, // position of memory order and scope arguments in the builtin unsigned OrderIndex, ScopeIndex; - const auto *FD = SemaRef.getCurFunctionDecl(true); + const auto *FD = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true); assert(FD && "AMDGPU builtins should not be used outside of a function"); llvm::StringMap<bool> CallerFeatureMap; getASTContext().getFunctionFeatureMap(CallerFeatureMap, FD); From 1eecbd4928dfc7b659c8f4f6dd524008c1fc59ff Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <juama...@amd.com> Date: Thu, 10 Apr 2025 09:56:17 +0200 Subject: [PATCH 3/4] [review] remove requires --- clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip b/clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip index 7a47fbd39817e..1e7bb44975a69 100644 --- a/clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip +++ b/clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip @@ -1,4 +1,3 @@ -// REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu tahiti -fsyntax-only -fcuda-is-device -verify=expected -o - %s // RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu gfx950 -fsyntax-only -fcuda-is-device -o - %s From a4683c4b3ac04ff365da2d34b1cad0b92636dc3c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <juama...@amd.com> Date: Thu, 10 Apr 2025 10:00:08 +0200 Subject: [PATCH 4/4] Add tests case outside of struct --- clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip b/clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip index 1e7bb44975a69..bb465825b7eff 100644 --- a/clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip +++ b/clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip @@ -21,3 +21,19 @@ __device__ __amdgpu_buffer_rsrc_t test_simple_builtin(void *p, short stride, int __device__ void test_target_dependant_builtin(void *src, __shared__ void *dst) { S::global_load_lds_lambda(src, dst); } + +constexpr auto make_buffer_rsrc_lambda = [](void *p, short stride, int num, int flags) { + return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags); +}; + +constexpr auto global_load_lds_lambda = [](void* src, __shared__ void *dst) { + __builtin_amdgcn_global_load_lds(src, dst, 16, 0, 0); // expected-error{{invalid size value}} expected-note{{size must be 1, 2, or 4}} +}; + +__device__ __amdgpu_buffer_rsrc_t global_test_simple_builtin(void *p, short stride, int num, int flags) { + return make_buffer_rsrc_lambda(p, stride, num, flags); +} + +__device__ void lobal_test_target_dependant_builtin(void *src, __shared__ void *dst) { + global_load_lds_lambda(src, dst); +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits