[Clang][AMDGPU] Accept builtins in lambda declarations

[jmmartinez]

Sema::getCurFunctionDecl(AllowLambda = false) returns a nullptr when the lambda declaration is outside a function (for example, when assigning a lambda to a static constexpr variable).

This triggered an assertion in SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall.

Using Sema::getCurFunctionDecl(AllowLambda = true) returns the declaration of the enclosing lambda.

Stumbled with this issue when refactoring some code in CK.

[llvmbot]

@llvm/pr-subscribers-clang

@llvm/pr-subscribers-backend-amdgpu

Author: Juan Manuel Martinez Caamaño (jmmartinez)

Changes

Sema::getCurFunctionDecl(AllowLambda = false) returns a nullptr when the lambda declaration is outside a function (for example, when assigning a lambda to a static constexpr variable).

This triggered an assertion in SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall.

Using Sema::getCurFunctionDecl(AllowLambda = true) returns the declaration of the enclosing lambda.

Stumbled with this issue when refactoring some code in CK.

---

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

2 Files Affected:

• (modified) clang/lib/Sema/SemaAMDGPU.cpp (+1-1)
• (added) clang/test/SemaHIP/amdgpu-builtin-in-lambda.hip (+24)

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);
+}

[yxsamliu]

LGTM. Pls address other comments

[arsenm]

IIRC there were bugs when you try to use lambda in conjunction with the target attribute. Can you add a test where you try to use a builtin that depends on a subtarget feature, in a function marked with the target attribute (and not available in the main compile target), that we still error?

#60005 is the issue for OpenCL blocks, I assume lambdas have a similar issue

[jmmartinez]

IIRC there were bugs when you try to use lambda in conjunction with the target attribute. Can you add a test where you try to use a builtin that depends on a subtarget feature, in a function marked with the target attribute (and not available in the main compile target), that we still error?

I've added a separate test for this (the error is not triggered with -fsyntax-only but it is with -emit-llvm).

I've also added tests for local variables inside a function.