https://github.com/jmmartinez updated
https://github.com/llvm/llvm-project/pull/135027
From a22e56d063e91df667d89ed5238bd886ce10bd0a 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] [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 +-
...n-in-lambda-with-unsupported-attribute.hip | 34 ++++++++++++
.../test/SemaHIP/amdgpu-builtin-in-lambda.hip | 53 +++++++++++++++++++
3 files changed, 88 insertions(+), 1 deletion(-)
create mode 100644
clang/test/SemaHIP/amdgpu-builtin-in-lambda-with-unsupported-attribute.hip
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..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();
+ 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);
diff --git
a/clang/test/SemaHIP/amdgpu-builtin-in-lambda-with-unsupported-attribute.hip
b/clang/test/SemaHIP/amdgpu-builtin-in-lambda-with-unsupported-attribute.hip
new file mode 100644
index 0000000000000..5b9223f2eaa3a
--- /dev/null
+++ b/clang/test/SemaHIP/amdgpu-builtin-in-lambda-with-unsupported-attribute.hip
@@ -0,0 +1,34 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu tahiti -emit-llvm
-fcuda-is-device -verify=no-memrealtime -o - %s
+// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu gfx950 -emit-llvm
-fcuda-is-device -o - %s
+
+#define __device__ __attribute__((device))
+#define __shared__ __attribute__((shared))
+
+struct S {
+ static constexpr auto memrealtime_lambda = []() {
+ __builtin_amdgcn_s_memrealtime(); //
no-memrealtime-error{{'__builtin_amdgcn_s_memrealtime' needs target feature
s-memrealtime}}
+ };
+};
+
+__attribute__((target("s-memrealtime")))
+__device__ void test_target_dependant_builtin_attr_fail() {
+ S::memrealtime_lambda();
+}
+
+constexpr auto memrealtime_lambda = []() {
+ __builtin_amdgcn_s_memrealtime(); //
no-memrealtime-error{{'__builtin_amdgcn_s_memrealtime' needs target feature
s-memrealtime}}
+};
+
+__attribute__((target("s-memrealtime")))
+__device__ void global_test_target_dependant_builtin_attr_fail() {
+ memrealtime_lambda();
+}
+
+__attribute__((target("s-memrealtime")))
+__device__ void local_test_target_dependant_builtin_attr_fail() {
+ static constexpr auto f = []() {
+ __builtin_amdgcn_s_memrealtime(); //
no-memrealtime-error{{'__builtin_amdgcn_s_memrealtime' needs target feature
s-memrealtime}}
+ };
+ f();
+}
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..8f0b14b7379d2
--- /dev/null
+++ b/clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip
@@ -0,0 +1,53 @@
+// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu gfx90a -fsyntax-only
-fcuda-is-device -verify=gfx90a -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); //
gfx90a-error{{invalid size value}} gfx90a-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);
+}
+
+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); //
gfx90a-error{{invalid size value}} gfx90a-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 global_test_target_dependant_builtin(void *src, __shared__
void *dst) {
+ global_load_lds_lambda(src, dst);
+}
+
+__device__ __amdgpu_buffer_rsrc_t local_test_simple_builtin(void *p, short
stride, int num, int flags) {
+ constexpr auto f = [](void *p, short stride, int num, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);
+ };
+ return f(p, stride, num, flags);
+}
+
+__device__ void local_test_target_dependant_builtin(void *src, __shared__ void
*dst) {
+ constexpr auto f = [](void* src, __shared__ void *dst) {
+ __builtin_amdgcn_global_load_lds(src, dst, 16, 0, 0); //
gfx90a-error{{invalid size value}} gfx90a-note{{size must be 1, 2, or 4}}
+ };
+ f(src, dst);
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
