https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/228508
>From e7ab0e4278bd077bccd93b28a8f30203256e2d18 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Fri, 2 Oct 2026 18:28:59 +0200 Subject: [PATCH 1/2] [CIR] Add norecurse to main and GPU/shader functions --- clang/lib/CIR/CodeGen/CIRGenFunction.cpp | 9 ++++++++ clang/test/CIR/CodeGen/main-norecurse.cpp | 19 ++++++++++++++++ clang/test/CIR/CodeGenCUDA/norecurse.cu | 27 +++++++++++++++++++++++ clang/test/CIR/CodeGenHLSL/norecurse.hlsl | 13 +++++++++++ clang/test/CIR/CodeGenOpenCL/norecurse.cl | 13 +++++++++++ 5 files changed, 81 insertions(+) create mode 100644 clang/test/CIR/CodeGen/main-norecurse.cpp create mode 100644 clang/test/CIR/CodeGenCUDA/norecurse.cu create mode 100644 clang/test/CIR/CodeGenHLSL/norecurse.hlsl create mode 100644 clang/test/CIR/CodeGenOpenCL/norecurse.cl diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp index ca4cf4be63de3..69deb0fec3d6f 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp @@ -517,6 +517,15 @@ void CIRGenFunction::startFunction(GlobalDecl gd, QualType returnType, const auto *fd = dyn_cast_or_null<FunctionDecl>(d); curFuncDecl = (d ? d->getNonClosureContext() : nullptr); + // Recursion is disallowed for C++ main, OpenCL, HLSL, SYCL device code and + // CUDA kernels. + if (fd && + ((getLangOpts().CPlusPlus && fd->isMain()) || getLangOpts().OpenCL || + getLangOpts().HLSL || getLangOpts().SYCLIsDevice || + (getLangOpts().CUDA && fd->hasAttr<CUDAGlobalAttr>()))) + fn->setAttr(cir::CIRDialect::getNoRecurseAttrName(), + mlir::UnitAttr::get(fn.getContext())); + // This is an artifact of the legacy handling of constrained floating-point // modes. The rounding mode and exception behavior tracked in // clang::LangOptions don't correspond directly to the representation we diff --git a/clang/test/CIR/CodeGen/main-norecurse.cpp b/clang/test/CIR/CodeGen/main-norecurse.cpp new file mode 100644 index 0000000000000..4b544e6a8d055 --- /dev/null +++ b/clang/test/CIR/CodeGen/main-norecurse.cpp @@ -0,0 +1,19 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM + +int f() { return 0; } + +int main(int argc, char **argv) { return f(); } + +// CIR: cir.func {{.*}} @_Z1fv( +// CIR-NOT: norecurse +// CIR: cir.func {{.*}} @main({{.*}} attributes {{.*}}norecurse + +// LLVM: define {{.*}} @_Z1fv() #[[F_ATTR:[0-9]+]] +// LLVM: define {{.*}} @main({{.*}}) #[[MAIN_ATTR:[0-9]+]] +// LLVM-NOT: attributes #[[F_ATTR]] = {{.*}}norecurse +// LLVM: attributes #[[MAIN_ATTR]] = {{.*}}norecurse diff --git a/clang/test/CIR/CodeGenCUDA/norecurse.cu b/clang/test/CIR/CodeGenCUDA/norecurse.cu new file mode 100644 index 0000000000000..7dae9fdb0ad1a --- /dev/null +++ b/clang/test/CIR/CodeGenCUDA/norecurse.cu @@ -0,0 +1,27 @@ +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fclangir \ +// RUN: -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fclangir \ +// RUN: -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fclangir \ +// RUN: -emit-llvm -x hip %s -o %t-hip-cir.ll +// RUN: FileCheck --input-file=%t-hip-cir.ll %s -check-prefix=LLVM +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \ +// RUN: -emit-llvm %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM + +#include "Inputs/cuda.h" + +__device__ void f() {} + +__global__ void kernel1(int a) { f(); } + +// CIR: cir.func {{.*}} @_Z1fv( +// CIR-NOT: norecurse +// CIR: cir.func {{.*}} @_Z7kernel1i({{.*}} attributes {{.*}}norecurse + +// LLVM: define {{.*}} @_Z1fv() #[[F_ATTR:[0-9]+]] +// LLVM: define {{.*}} @_Z7kernel1i({{.*}}) #[[K_ATTR:[0-9]+]] +// LLVM-NOT: attributes #[[F_ATTR]] = {{.*}}norecurse +// LLVM: attributes #[[K_ATTR]] = {{.*}}norecurse diff --git a/clang/test/CIR/CodeGenHLSL/norecurse.hlsl b/clang/test/CIR/CodeGenHLSL/norecurse.hlsl new file mode 100644 index 0000000000000..43e48a24384a3 --- /dev/null +++ b/clang/test/CIR/CodeGenHLSL/norecurse.hlsl @@ -0,0 +1,13 @@ +// RUN: %clang_cc1 -x hlsl -triple spirv-unknown-vulkan-library -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR +// RUN: %clang_cc1 -x hlsl -triple spirv-unknown-vulkan-library -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM +// RUN: %clang_cc1 -x hlsl -triple spirv-unknown-vulkan-library -emit-llvm -disable-llvm-passes %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM + +export void f() {} + +// CIR: cir.func {{.*}} @_Z1fv() {{.*}} attributes {{.*}}norecurse + +// LLVM: define {{.*}} @_Z1fv() #[[ATTR:[0-9]+]] +// LLVM: attributes #[[ATTR]] = {{.*}}norecurse diff --git a/clang/test/CIR/CodeGenOpenCL/norecurse.cl b/clang/test/CIR/CodeGenOpenCL/norecurse.cl new file mode 100644 index 0000000000000..6534152b4647a --- /dev/null +++ b/clang/test/CIR/CodeGenOpenCL/norecurse.cl @@ -0,0 +1,13 @@ +// RUN: %clang_cc1 -triple spirv64-unknown-unknown -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR +// RUN: %clang_cc1 -triple spirv64-unknown-unknown -fclangir -emit-llvm -O0 %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM +// RUN: %clang_cc1 -triple spirv64-unknown-unknown -emit-llvm -O0 %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM + +kernel void kernel1(int a) {} + +// CIR: cir.func{{.*}} @kernel1({{.*}} attributes {{.*}}norecurse + +// LLVM: define {{.*}} @kernel1({{.*}}) #[[ATTR:[0-9]+]] +// LLVM: attributes #[[ATTR]] = {{.*}}norecurse >From d82bdf33639aa9cfbed82e40dcf700ad0c28d0d0 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 5 Oct 2026 17:59:51 +0200 Subject: [PATCH 2/2] Update clang/lib/CIR/CodeGen/CIRGenFunction.cpp Co-authored-by: Steffen Larsen <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenFunction.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp index 69deb0fec3d6f..c987e96af6716 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp @@ -518,7 +518,7 @@ void CIRGenFunction::startFunction(GlobalDecl gd, QualType returnType, curFuncDecl = (d ? d->getNonClosureContext() : nullptr); // Recursion is disallowed for C++ main, OpenCL, HLSL, SYCL device code and - // CUDA kernels. + // CUDA/HIP kernels. if (fd && ((getLangOpts().CPlusPlus && fd->isMain()) || getLangOpts().OpenCL || getLangOpts().HLSL || getLangOpts().SYCLIsDevice || _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
