[clang] c16297c - [CUDA][HIP] Fix host/device attribute of builtin (#138162)
via cfe-commits
cfe-commits at lists.llvm.org
Wed May 7 19:03:36 PDT 2025
Author: Yaxun (Sam) Liu
Date: 2025-05-07T22:03:33-04:00
New Revision: c16297cd3f0ed9d036e9cf16fb6885aa3c72d5d3
URL: https://github.com/llvm/llvm-project/commit/c16297cd3f0ed9d036e9cf16fb6885aa3c72d5d3
DIFF: https://github.com/llvm/llvm-project/commit/c16297cd3f0ed9d036e9cf16fb6885aa3c72d5d3.diff
LOG: [CUDA][HIP] Fix host/device attribute of builtin (#138162)
When a builtin function is passed a pointer with a different
address space, clang creates an overloaded
builtin function but does not copy the host/device attribute. This
causes
error when the builtin is called by device functions
since CUDA/HIP relies on the host/device attribute to treat
a builtin function as callable on both host and device
sides.
Fixed by copying the host/device attribute of the original
builtin function to the created overloaded builtin function.
Added:
clang/test/SemaCUDA/overloaded-builtin.cu
Modified:
clang/lib/Sema/SemaExpr.cpp
Removed:
################################################################################
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index c3ef5a70d5f6d..57135adf714ce 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -6362,6 +6362,14 @@ static FunctionDecl *rewriteBuiltinFunctionDecl(Sema *Sema, ASTContext &Context,
Params.push_back(Parm);
}
OverloadDecl->setParams(Params);
+ // We cannot merge host/device attributes of redeclarations. They have to
+ // be consistent when created.
+ if (Sema->LangOpts.CUDA) {
+ if (FDecl->hasAttr<CUDAHostAttr>())
+ OverloadDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
+ if (FDecl->hasAttr<CUDADeviceAttr>())
+ OverloadDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ }
Sema->mergeDeclAttributes(OverloadDecl, FDecl);
return OverloadDecl;
}
diff --git a/clang/test/SemaCUDA/overloaded-builtin.cu b/clang/test/SemaCUDA/overloaded-builtin.cu
new file mode 100644
index 0000000000000..c60c27e7f8627
--- /dev/null
+++ b/clang/test/SemaCUDA/overloaded-builtin.cu
@@ -0,0 +1,36 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -fsyntax-only -verify=host -xhip %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsyntax-only -fcuda-is-device -verify=dev -xhip %s
+
+// dev-no-diagnostics
+
+#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);
+}
+
+void hfun() {
+ __attribute__((address_space(0))) void *mem_ptr;
+ (void)__builtin_amdgcn_is_shared(mem_ptr); // host-error {{reference to __device__ function '__builtin_amdgcn_is_shared' in __host__ function}}
+}
+
+template<typename T>
+void template_hfun(T *p) {
+ __attribute__((address_space(0))) void *mem_ptr;
+ (void)__builtin_amdgcn_is_shared(mem_ptr); // host-error {{reference to __device__ function '__builtin_amdgcn_is_shared' in __host__ function}}
+}
+
+
+int main() {
+ int *p;
+ kernel<<<1,1>>>();
+ template_kernel<<<1,1>>>(p);
+ template_hfun(p); // host-note {{called by 'main'}}
+}
More information about the cfe-commits
mailing list