https://github.com/yxsamliu created
https://github.com/llvm/llvm-project/pull/138162
When a builtin function with generic pointer parameter is passed a pointer with
address space, clang creates an overloaded builtin function but does not make
it implicit. This causes error when the builtin is called by device functions
since CUDA/HIP relies on the implicit attribute to treat a builtin function as
callable on both host and device sides.
Fixed by making the created overloaded builtin functions implicit.
>From 9c73c14588c64623bfa6bcb4fd11e6e12dc0745e Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun....@amd.com>
Date: Thu, 1 May 2025 12:08:05 -0400
Subject: [PATCH] [CUDA][HIP] Fix implicit attribute of builtin
When a builtin function with generic pointer parameter is passed
a pointer with address space, clang creates an overloaded
builtin function but does not make it implicit. This causes
error when the builtin is called by device functions
since CUDA/HIP relies on the implicit attribute to treat
a builtin function as callable on both host and device
sides.
Fixed by making the created overloaded builtin functions implicit.
---
clang/lib/Sema/SemaExpr.cpp | 1 +
clang/test/SemaCUDA/overloaded-builtin.cu | 23 +++++++++++++++++++++++
2 files changed, 24 insertions(+)
create mode 100644 clang/test/SemaCUDA/overloaded-builtin.cu
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 0cd86dc54b0ab..d9eccb31e6d1e 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -6358,6 +6358,7 @@ static FunctionDecl *rewriteBuiltinFunctionDecl(Sema
*Sema, ASTContext &Context,
}
OverloadDecl->setParams(Params);
Sema->mergeDeclAttributes(OverloadDecl, FDecl);
+ OverloadDecl->setImplicit(true);
return OverloadDecl;
}
diff --git a/clang/test/SemaCUDA/overloaded-builtin.cu
b/clang/test/SemaCUDA/overloaded-builtin.cu
new file mode 100644
index 0000000000000..719bfea4aef2f
--- /dev/null
+++ b/clang/test/SemaCUDA/overloaded-builtin.cu
@@ -0,0 +1,23 @@
+// expected-no-diagnostics
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -aux-triple
amdgcn-amd-amdhsa -fsyntax-only -verify -xhip %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsyntax-only -fcuda-is-device
-verify -xhip %s
+
+#include "Inputs/cuda.h"
+
+__global__ void kernel() {
+ __attribute__((address_space(0))) void *mem_ptr;
+ (void)__builtin_amdgcn_is_shared(mem_ptr);
+}
+
+template<typename T>
+__global__ void template_kernel(T *p) {
+ __attribute__((address_space(0))) void *mem_ptr;
+ (void)__builtin_amdgcn_is_shared(mem_ptr);
+}
+
+int main() {
+ int *p;
+ kernel<<<1,1>>>();
+ template_kernel<<<1,1>>>(p);
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
