Author: Tang Jiajun
Date: 2025-09-17T13:05:57+08:00
New Revision: 6d5c94203652a52b51b37ad4768bc1a7066f029c

URL: 
https://github.com/llvm/llvm-project/commit/6d5c94203652a52b51b37ad4768bc1a7066f029c
DIFF: 
https://github.com/llvm/llvm-project/commit/6d5c94203652a52b51b37ad4768bc1a7066f029c.diff

LOG: [clang] [CUDA] Support calling `consteval` function between different 
target. (#158688)

In CUDA, calling `consteval` functions cross excution space is allowed.
So the function with `consteval` attribute need be treated as a
`__host__ __device__` function.

Added: 
    clang/test/SemaCUDA/consteval-func.cu

Modified: 
    clang/docs/ReleaseNotes.rst
    clang/lib/Sema/SemaCUDA.cpp

Removed: 
    


################################################################################
diff  --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index 9af55951a1d82..5211d84be7a3b 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -448,6 +448,8 @@ CUDA/HIP Language Changes
 CUDA Support
 ^^^^^^^^^^^^
 
+Support calling `consteval` function between 
diff erent target.
+
 AIX Support
 ^^^^^^^^^^^
 

diff  --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 2e3cbb336a0c8..31735a0f5feb3 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -143,6 +143,9 @@ CUDAFunctionTarget SemaCUDA::IdentifyTarget(const 
FunctionDecl *D,
   if (D->hasAttr<CUDAGlobalAttr>())
     return CUDAFunctionTarget::Global;
 
+  if (D->isConsteval())
+    return CUDAFunctionTarget::HostDevice;
+
   if (hasAttr<CUDADeviceAttr>(D, IgnoreImplicitHDAttr)) {
     if (hasAttr<CUDAHostAttr>(D, IgnoreImplicitHDAttr))
       return CUDAFunctionTarget::HostDevice;

diff  --git a/clang/test/SemaCUDA/consteval-func.cu 
b/clang/test/SemaCUDA/consteval-func.cu
new file mode 100644
index 0000000000000..293c1ce85830a
--- /dev/null
+++ b/clang/test/SemaCUDA/consteval-func.cu
@@ -0,0 +1,8 @@
+// RUN: %clang_cc1 -std=c++20 -fsyntax-only -verify %s
+
+// expected-no-diagnostics
+
+#include "Inputs/cuda.h"
+
+__device__ consteval int f() { return 0; }
+int main() { return f(); }


        
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to