Skip to content

Commit 6d5c942

Browse files
authored
[clang] [CUDA] Support calling consteval function between different target. (llvm#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.
1 parent c286a42 commit 6d5c942

File tree

3 files changed

+13
-0
lines changed

3 files changed

+13
-0
lines changed

clang/docs/ReleaseNotes.rst

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -448,6 +448,8 @@ CUDA/HIP Language Changes
448448
CUDA Support
449449
^^^^^^^^^^^^
450450

451+
Support calling `consteval` function between different target.
452+
451453
AIX Support
452454
^^^^^^^^^^^
453455

clang/lib/Sema/SemaCUDA.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -143,6 +143,9 @@ CUDAFunctionTarget SemaCUDA::IdentifyTarget(const FunctionDecl *D,
143143
if (D->hasAttr<CUDAGlobalAttr>())
144144
return CUDAFunctionTarget::Global;
145145

146+
if (D->isConsteval())
147+
return CUDAFunctionTarget::HostDevice;
148+
146149
if (hasAttr<CUDADeviceAttr>(D, IgnoreImplicitHDAttr)) {
147150
if (hasAttr<CUDAHostAttr>(D, IgnoreImplicitHDAttr))
148151
return CUDAFunctionTarget::HostDevice;
Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,8 @@
1+
// RUN: %clang_cc1 -std=c++20 -fsyntax-only -verify %s
2+
3+
// expected-no-diagnostics
4+
5+
#include "Inputs/cuda.h"
6+
7+
__device__ consteval int f() { return 0; }
8+
int main() { return f(); }

0 commit comments

Comments
 (0)