[clang] [CIR] Implement device-wide nounwind attribute (PR #227016)
Steffen Larsen via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 1 23:06:06 PDT 2026
https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/227016
>From 5a1e273e4043b56b608b314e038bdb83a2c9af6f Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Mon, 28 Sep 2026 09:54:35 -0500
Subject: [PATCH 1/3] [CIR] Implement device-wide nounwind attribute
This commit implements the missing application of nounwind on device
functions and device calls. This leaves out OpenMP, similar to how OGCG
currently does not support it for OpenMP device functions.
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/lib/CIR/CodeGen/CIRGenCall.cpp | 12 +++-
.../test/CIR/CodeGen/offload-nounwind-attr.cu | 49 +++++++++++++
clang/test/CIR/CodeGenCUDA/nounwind.cu | 68 +++++++++++++++++++
clang/test/CIR/CodeGenHIP/nounwind.hip | 68 +++++++++++++++++++
clang/test/CIR/CodeGenOpenCL/nounwind.cl | 25 +++++++
.../CodeGenSYCL/kernel-caller-attributes.cpp | 2 +-
clang/test/CIR/CodeGenSYCL/nounwind.cpp | 42 ++++++++++++
7 files changed, 263 insertions(+), 3 deletions(-)
create mode 100644 clang/test/CIR/CodeGen/offload-nounwind-attr.cu
create mode 100644 clang/test/CIR/CodeGenCUDA/nounwind.cu
create mode 100644 clang/test/CIR/CodeGenHIP/nounwind.hip
create mode 100644 clang/test/CIR/CodeGenOpenCL/nounwind.cl
create mode 100644 clang/test/CIR/CodeGenSYCL/nounwind.cpp
diff --git a/clang/lib/CIR/CodeGen/CIRGenCall.cpp b/clang/lib/CIR/CodeGen/CIRGenCall.cpp
index 226bc0a0ed0d5..dbeebd68b218d 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCall.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCall.cpp
@@ -268,8 +268,16 @@ static void addTrivialDefaultFunctionAttributes(
mlir::UnitAttr::get(mlirCtx));
}
- // TODO(cir): Classic codegen adds 'nounwind' here in a bunch of offload
- // targets.
+ // Device code cannot unwind. 'nothrow' keeps calls from getting an unwind
+ // edge, so 'nounwind' becomes LLVM's nounwind.
+ // TODO: OpenMP offload is not covered.
+ if ((langOpts.CUDA && langOpts.CUDAIsDevice) || langOpts.OpenCL ||
+ langOpts.SYCLIsDevice) {
+ attrs.set(cir::CIRDialect::getNoThrowAttrName(),
+ mlir::UnitAttr::get(mlirCtx));
+ attrs.set(cir::CIRDialect::getNoUnwindAttrName(),
+ mlir::UnitAttr::get(mlirCtx));
+ }
if (codeGenOpts.SaveRegParams && !attrOnCallSite)
attrs.set(cir::CIRDialect::getSaveRegParamsAttrName(),
diff --git a/clang/test/CIR/CodeGen/offload-nounwind-attr.cu b/clang/test/CIR/CodeGen/offload-nounwind-attr.cu
new file mode 100644
index 0000000000000..68eb982c8328c
--- /dev/null
+++ b/clang/test/CIR/CodeGen/offload-nounwind-attr.cu
@@ -0,0 +1,49 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fcuda-is-device -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s -check-prefix=CIR
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fcuda-is-device -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s -check-prefix=LLVM
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fcuda-is-device -emit-llvm %s -o - \
+// RUN: | FileCheck %s -check-prefix=LLVM
+
+// Device code cannot unwind even when exceptions are enabled, so calls inside
+// an EH cleanup scope must stay plain calls.
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fcuda-is-device -fcxx-exceptions -fexceptions -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s -check-prefix=CIR
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fcuda-is-device -fcxx-exceptions -fexceptions -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s -check-prefix=LLVM
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fcuda-is-device -fcxx-exceptions -fexceptions -emit-llvm %s -o - \
+// RUN: | FileCheck %s -check-prefix=LLVM
+
+struct D {
+ __attribute__((device)) ~D();
+};
+
+extern "C" {
+__attribute__((device)) int ext(int);
+
+__attribute__((device)) int caller(int x) {
+ D d;
+ return ext(x);
+}
+}
+
+// CIR: cir.func{{.*}}@caller(
+// CIR-SAME: nothrow, nounwind
+// CIR-NOT: cir.try_call
+// CIR: cir.call @ext(%{{.*}}) nothrow nounwind
+// CIR-NOT: cir.try_call
+// CIR: cir.call @_ZN1DD1Ev(%{{.*}}) nothrow nounwind
+// CIR-NOT: cir.try_call
+// CIR: cir.func private @ext(
+// CIR-SAME: nothrow, nounwind
+
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: define {{.*}}@caller({{.*}}){{.*}} #{{[0-9]+}} {
+// LLVM-NOT: invoke
+// LLVM: call {{.*}}@ext({{.*}}) #[[CALL_ATTR:[0-9]+]]
+// LLVM-NOT: invoke
+// LLVM: call void @_ZN1DD1Ev({{.*}}) #[[CALL_ATTR]]
+// LLVM-NOT: landingpad
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: declare {{.*}}@ext(
+// LLVM: attributes #[[CALL_ATTR]] = {{{.*}}nounwind
diff --git a/clang/test/CIR/CodeGenCUDA/nounwind.cu b/clang/test/CIR/CodeGenCUDA/nounwind.cu
new file mode 100644
index 0000000000000..3a4d21b608338
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/nounwind.cu
@@ -0,0 +1,68 @@
+#include "Inputs/cuda.h"
+
+// REQUIRES: nvptx-registered-target
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device \
+// RUN: -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device \
+// RUN: -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device \
+// RUN: -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+
+// Device code cannot unwind even when exceptions are enabled, so calls inside
+// an EH cleanup scope must stay plain calls.
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device \
+// RUN: -fcxx-exceptions -fexceptions -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device \
+// RUN: -fcxx-exceptions -fexceptions -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device \
+// RUN: -fcxx-exceptions -fexceptions -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+
+struct D {
+ __device__ ~D();
+};
+
+extern "C" {
+__device__ int ext(int);
+
+__device__ int caller(int x) {
+ D d;
+ return ext(x);
+}
+
+__global__ void kernel(int *p) { *p = caller(*p); }
+}
+
+// CIR: cir.func{{.*}}@caller(
+// CIR-SAME: nothrow, nounwind
+// CIR-NOT: cir.try_call
+// CIR: cir.call @ext(%{{.*}}) nothrow nounwind
+// CIR-NOT: cir.try_call
+// CIR: cir.call @_ZN1DD1Ev(%{{.*}}) nothrow nounwind
+// CIR-NOT: cir.try_call
+// CIR: cir.func private @ext(
+// CIR-SAME: nothrow, nounwind
+// CIR: cir.func{{.*}}@kernel(
+// CIR-SAME: nothrow, nounwind
+// CIR: cir.call @caller(%{{.*}}) nothrow nounwind
+
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: define {{.*}}@caller({{.*}}){{.*}} #{{[0-9]+}} {
+// LLVM-NOT: invoke
+// LLVM: call {{.*}}@ext({{.*}}) #[[CALL_ATTR:[0-9]+]]
+// LLVM-NOT: invoke
+// LLVM: call void @_ZN1DD1Ev({{.*}}) #[[CALL_ATTR]]
+// LLVM-NOT: landingpad
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: declare {{.*}}@ext(
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: declare {{.*}}@_ZN1DD1Ev(
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: define {{.*}}ptx_kernel void @kernel({{.*}}){{.*}} #{{[0-9]+}} {
+// LLVM: call {{.*}}@caller({{.*}}) #[[CALL_ATTR]]
+// LLVM: attributes #[[CALL_ATTR]] = {{{.*}}nounwind
diff --git a/clang/test/CIR/CodeGenHIP/nounwind.hip b/clang/test/CIR/CodeGenHIP/nounwind.hip
new file mode 100644
index 0000000000000..6523f8a5a30c6
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/nounwind.hip
@@ -0,0 +1,68 @@
+#include "../CodeGenCUDA/Inputs/cuda.h"
+
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+
+// Device code cannot unwind even when exceptions are enabled, so calls inside
+// an EH cleanup scope must stay plain calls.
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -fcxx-exceptions -fexceptions -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -fcxx-exceptions -fexceptions -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -fcxx-exceptions -fexceptions -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+
+struct D {
+ __device__ ~D();
+};
+
+extern "C" {
+__device__ int ext(int);
+
+__device__ int caller(int x) {
+ D d;
+ return ext(x);
+}
+
+__global__ void kernel(int *p) { *p = caller(*p); }
+}
+
+// CIR: cir.func{{.*}}@caller(
+// CIR-SAME: nothrow, nounwind
+// CIR-NOT: cir.try_call
+// CIR: cir.call @ext(%{{.*}}) nothrow nounwind
+// CIR-NOT: cir.try_call
+// CIR: cir.call @_ZN1DD1Ev(%{{.*}}) nothrow nounwind
+// CIR-NOT: cir.try_call
+// CIR: cir.func private @ext(
+// CIR-SAME: nothrow, nounwind
+// CIR: cir.func{{.*}}@kernel(
+// CIR-SAME: nothrow, nounwind
+// CIR: cir.call @caller(%{{.*}}) nothrow nounwind
+
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: define {{.*}}@caller({{.*}}){{.*}} #{{[0-9]+}} {
+// LLVM-NOT: invoke
+// LLVM: call {{.*}}@ext({{.*}}) #[[CALL_ATTR:[0-9]+]]
+// LLVM-NOT: invoke
+// LLVM: call void @_ZN1DD1Ev({{.*}}) #[[CALL_ATTR]]
+// LLVM-NOT: landingpad
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: declare {{.*}}@ext(
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: declare {{.*}}@_ZN1DD1Ev(
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: define {{.*}}amdgpu_kernel void @kernel({{.*}}){{.*}} #{{[0-9]+}} {
+// LLVM: call {{.*}}@caller({{.*}}) #[[CALL_ATTR]]
+// LLVM: attributes #[[CALL_ATTR]] = {{{.*}}nounwind
diff --git a/clang/test/CIR/CodeGenOpenCL/nounwind.cl b/clang/test/CIR/CodeGenOpenCL/nounwind.cl
new file mode 100644
index 0000000000000..a3b35afaad301
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenCL/nounwind.cl
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 %s -fclangir -emit-cir -triple spirv64-unknown-unknown -o - \
+// RUN: | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 %s -fclangir -emit-llvm -triple spirv64-unknown-unknown -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 %s -emit-llvm -triple spirv64-unknown-unknown -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+
+// OpenCL has no exceptions, so every function and call is nounwind.
+
+int ext(int);
+
+int caller(int x) { return ext(x); }
+
+// CIR: cir.func{{.*}}@caller(
+// CIR-SAME: nothrow, nounwind
+// CIR: cir.call @ext(%{{.*}}) nothrow nounwind
+// CIR: cir.func private @ext(
+// CIR-SAME: nothrow, nounwind
+
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: define {{.*}}@caller(
+// LLVM: call {{.*}}@ext({{.*}}) #[[CALL_ATTR:[0-9]+]]
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: declare {{.*}}@ext(
+// LLVM: attributes #[[CALL_ATTR]] = {{{.*}}nounwind
diff --git a/clang/test/CIR/CodeGenSYCL/kernel-caller-attributes.cpp b/clang/test/CIR/CodeGenSYCL/kernel-caller-attributes.cpp
index 54b168b2929f7..6ba946d9428c7 100644
--- a/clang/test/CIR/CodeGenSYCL/kernel-caller-attributes.cpp
+++ b/clang/test/CIR/CodeGenSYCL/kernel-caller-attributes.cpp
@@ -44,4 +44,4 @@ void test(int *p) {
// caller must not carry mustprogress; norecurse and sycl-module-id remain. The
// exact attribute set (which omits mustprogress) is verified here.
// LLVM-NOMP: define spir_kernel void @_ZTS2KN({{.*}}) #[[KATTR:[0-9]+]]
-// LLVM-NOMP: attributes #[[KATTR]] = { convergent noinline norecurse "sycl-module-id"="{{.*}}kernel-caller-attributes.cpp" }
+// LLVM-NOMP: attributes #[[KATTR]] = { convergent noinline norecurse nounwind "sycl-module-id"="{{.*}}kernel-caller-attributes.cpp" }
diff --git a/clang/test/CIR/CodeGenSYCL/nounwind.cpp b/clang/test/CIR/CodeGenSYCL/nounwind.cpp
new file mode 100644
index 0000000000000..6dee9f3f7e6e1
--- /dev/null
+++ b/clang/test/CIR/CodeGenSYCL/nounwind.cpp
@@ -0,0 +1,42 @@
+// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple spirv64-unknown-unknown -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s -check-prefix=CIR
+// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple spirv64-unknown-unknown -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s -check-prefix=LLVM
+// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple spirv64-unknown-unknown -emit-llvm %s -o - \
+// RUN: | FileCheck %s -check-prefix=LLVM
+
+// SYCL device code has no exceptions, so every function and call is nounwind.
+
+template <typename KernelName, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
+
+template <typename KernelName, typename KernelType>
+[[clang::sycl_kernel_entry_point(KernelName)]]
+void kernel_single_task(KernelType kf) { kf(); }
+
+struct KN;
+
+int ext(int);
+
+void test(int *p) {
+ kernel_single_task<KN>([p]() { *p = ext(*p); });
+}
+
+// CIR: cir.func{{.*}}@_ZTS2KN(
+// CIR-SAME: nothrow, nounwind
+// CIR: cir.call @_ZZ4testPiENKUlvE_clEv(%{{.*}}) nothrow nounwind
+// CIR: cir.func{{.*}}@_ZZ4testPiENKUlvE_clEv(
+// CIR-SAME: nothrow, nounwind
+// CIR: cir.call @_Z3exti(%{{.*}}) nothrow nounwind
+// CIR: cir.func private @_Z3exti(
+// CIR-SAME: nothrow, nounwind
+
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: define {{.*}}@_ZTS2KN(
+// LLVM: call {{.*}}@_ZZ4testPiENKUlvE_clEv({{.*}}) #[[CALL_ATTR:[0-9]+]]
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: define {{.*}}@_ZZ4testPiENKUlvE_clEv(
+// LLVM: call {{.*}}@_Z3exti({{.*}}) #[[CALL_ATTR]]
+// LLVM: ; Function Attrs: {{.*}}nounwind
+// LLVM-NEXT: declare {{.*}}@_Z3exti(
+// LLVM: attributes #[[CALL_ATTR]] = {{{.*}}nounwind
>From ea5997a3a1aab3baa11462c9a0a551ad76733971 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Wed, 30 Sep 2026 01:24:57 -0500
Subject: [PATCH 2/3] Fix triple
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/test/CIR/CodeGenHIP/nounwind.hip | 12 ++++++------
1 file changed, 6 insertions(+), 6 deletions(-)
diff --git a/clang/test/CIR/CodeGenHIP/nounwind.hip b/clang/test/CIR/CodeGenHIP/nounwind.hip
index 6523f8a5a30c6..4169f7e42f834 100644
--- a/clang/test/CIR/CodeGenHIP/nounwind.hip
+++ b/clang/test/CIR/CodeGenHIP/nounwind.hip
@@ -1,25 +1,25 @@
#include "../CodeGenCUDA/Inputs/cuda.h"
// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
// RUN: -fclangir -emit-cir %s -o - \
// RUN: | FileCheck %s --check-prefix=CIR
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
// RUN: -fclangir -emit-llvm %s -o - \
// RUN: | FileCheck %s --check-prefix=LLVM
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
// RUN: -emit-llvm %s -o - \
// RUN: | FileCheck %s --check-prefix=LLVM
// Device code cannot unwind even when exceptions are enabled, so calls inside
// an EH cleanup scope must stay plain calls.
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
// RUN: -fcxx-exceptions -fexceptions -fclangir -emit-cir %s -o - \
// RUN: | FileCheck %s --check-prefix=CIR
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
// RUN: -fcxx-exceptions -fexceptions -fclangir -emit-llvm %s -o - \
// RUN: | FileCheck %s --check-prefix=LLVM
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
// RUN: -fcxx-exceptions -fexceptions -emit-llvm %s -o - \
// RUN: | FileCheck %s --check-prefix=LLVM
>From 01a4381e941cbe0db6ba8546c0cbb8c01bdc8886 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Wed, 30 Sep 2026 07:45:41 -0500
Subject: [PATCH 3/3] Unify tests
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
.../test/CIR/CodeGen/offload-nounwind-attr.cu | 49 -------------
clang/test/CIR/CodeGenCUDA/nounwind.cu | 22 +++++-
clang/test/CIR/CodeGenHIP/nounwind.hip | 68 -------------------
3 files changed, 20 insertions(+), 119 deletions(-)
delete mode 100644 clang/test/CIR/CodeGen/offload-nounwind-attr.cu
delete mode 100644 clang/test/CIR/CodeGenHIP/nounwind.hip
diff --git a/clang/test/CIR/CodeGen/offload-nounwind-attr.cu b/clang/test/CIR/CodeGen/offload-nounwind-attr.cu
deleted file mode 100644
index 68eb982c8328c..0000000000000
--- a/clang/test/CIR/CodeGen/offload-nounwind-attr.cu
+++ /dev/null
@@ -1,49 +0,0 @@
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fcuda-is-device -fclangir -emit-cir %s -o - \
-// RUN: | FileCheck %s -check-prefix=CIR
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fcuda-is-device -fclangir -emit-llvm %s -o - \
-// RUN: | FileCheck %s -check-prefix=LLVM
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fcuda-is-device -emit-llvm %s -o - \
-// RUN: | FileCheck %s -check-prefix=LLVM
-
-// Device code cannot unwind even when exceptions are enabled, so calls inside
-// an EH cleanup scope must stay plain calls.
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fcuda-is-device -fcxx-exceptions -fexceptions -fclangir -emit-cir %s -o - \
-// RUN: | FileCheck %s -check-prefix=CIR
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fcuda-is-device -fcxx-exceptions -fexceptions -fclangir -emit-llvm %s -o - \
-// RUN: | FileCheck %s -check-prefix=LLVM
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fcuda-is-device -fcxx-exceptions -fexceptions -emit-llvm %s -o - \
-// RUN: | FileCheck %s -check-prefix=LLVM
-
-struct D {
- __attribute__((device)) ~D();
-};
-
-extern "C" {
-__attribute__((device)) int ext(int);
-
-__attribute__((device)) int caller(int x) {
- D d;
- return ext(x);
-}
-}
-
-// CIR: cir.func{{.*}}@caller(
-// CIR-SAME: nothrow, nounwind
-// CIR-NOT: cir.try_call
-// CIR: cir.call @ext(%{{.*}}) nothrow nounwind
-// CIR-NOT: cir.try_call
-// CIR: cir.call @_ZN1DD1Ev(%{{.*}}) nothrow nounwind
-// CIR-NOT: cir.try_call
-// CIR: cir.func private @ext(
-// CIR-SAME: nothrow, nounwind
-
-// LLVM: ; Function Attrs: {{.*}}nounwind
-// LLVM-NEXT: define {{.*}}@caller({{.*}}){{.*}} #{{[0-9]+}} {
-// LLVM-NOT: invoke
-// LLVM: call {{.*}}@ext({{.*}}) #[[CALL_ATTR:[0-9]+]]
-// LLVM-NOT: invoke
-// LLVM: call void @_ZN1DD1Ev({{.*}}) #[[CALL_ATTR]]
-// LLVM-NOT: landingpad
-// LLVM: ; Function Attrs: {{.*}}nounwind
-// LLVM-NEXT: declare {{.*}}@ext(
-// LLVM: attributes #[[CALL_ATTR]] = {{{.*}}nounwind
diff --git a/clang/test/CIR/CodeGenCUDA/nounwind.cu b/clang/test/CIR/CodeGenCUDA/nounwind.cu
index 3a4d21b608338..a66366834ddd8 100644
--- a/clang/test/CIR/CodeGenCUDA/nounwind.cu
+++ b/clang/test/CIR/CodeGenCUDA/nounwind.cu
@@ -1,6 +1,6 @@
#include "Inputs/cuda.h"
-// REQUIRES: nvptx-registered-target
+// Device code is nounwind for both CUDA and HIP, independent of the target.
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device \
// RUN: -fclangir -emit-cir %s -o - \
// RUN: | FileCheck %s --check-prefix=CIR
@@ -10,6 +10,15 @@
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device \
// RUN: -emit-llvm %s -o - \
// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
// Device code cannot unwind even when exceptions are enabled, so calls inside
// an EH cleanup scope must stay plain calls.
@@ -22,6 +31,15 @@
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device \
// RUN: -fcxx-exceptions -fexceptions -emit-llvm %s -o - \
// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -fcxx-exceptions -fexceptions -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -fcxx-exceptions -fexceptions -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -fcxx-exceptions -fexceptions -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
struct D {
__device__ ~D();
@@ -63,6 +81,6 @@ __global__ void kernel(int *p) { *p = caller(*p); }
// LLVM: ; Function Attrs: {{.*}}nounwind
// LLVM-NEXT: declare {{.*}}@_ZN1DD1Ev(
// LLVM: ; Function Attrs: {{.*}}nounwind
-// LLVM-NEXT: define {{.*}}ptx_kernel void @kernel({{.*}}){{.*}} #{{[0-9]+}} {
+// LLVM-NEXT: define {{.*}}void @kernel({{.*}}){{.*}} #{{[0-9]+}} {
// LLVM: call {{.*}}@caller({{.*}}) #[[CALL_ATTR]]
// LLVM: attributes #[[CALL_ATTR]] = {{{.*}}nounwind
diff --git a/clang/test/CIR/CodeGenHIP/nounwind.hip b/clang/test/CIR/CodeGenHIP/nounwind.hip
deleted file mode 100644
index 4169f7e42f834..0000000000000
--- a/clang/test/CIR/CodeGenHIP/nounwind.hip
+++ /dev/null
@@ -1,68 +0,0 @@
-#include "../CodeGenCUDA/Inputs/cuda.h"
-
-// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
-// RUN: -fclangir -emit-cir %s -o - \
-// RUN: | FileCheck %s --check-prefix=CIR
-// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
-// RUN: -fclangir -emit-llvm %s -o - \
-// RUN: | FileCheck %s --check-prefix=LLVM
-// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
-// RUN: -emit-llvm %s -o - \
-// RUN: | FileCheck %s --check-prefix=LLVM
-
-// Device code cannot unwind even when exceptions are enabled, so calls inside
-// an EH cleanup scope must stay plain calls.
-// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
-// RUN: -fcxx-exceptions -fexceptions -fclangir -emit-cir %s -o - \
-// RUN: | FileCheck %s --check-prefix=CIR
-// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
-// RUN: -fcxx-exceptions -fexceptions -fclangir -emit-llvm %s -o - \
-// RUN: | FileCheck %s --check-prefix=LLVM
-// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device \
-// RUN: -fcxx-exceptions -fexceptions -emit-llvm %s -o - \
-// RUN: | FileCheck %s --check-prefix=LLVM
-
-struct D {
- __device__ ~D();
-};
-
-extern "C" {
-__device__ int ext(int);
-
-__device__ int caller(int x) {
- D d;
- return ext(x);
-}
-
-__global__ void kernel(int *p) { *p = caller(*p); }
-}
-
-// CIR: cir.func{{.*}}@caller(
-// CIR-SAME: nothrow, nounwind
-// CIR-NOT: cir.try_call
-// CIR: cir.call @ext(%{{.*}}) nothrow nounwind
-// CIR-NOT: cir.try_call
-// CIR: cir.call @_ZN1DD1Ev(%{{.*}}) nothrow nounwind
-// CIR-NOT: cir.try_call
-// CIR: cir.func private @ext(
-// CIR-SAME: nothrow, nounwind
-// CIR: cir.func{{.*}}@kernel(
-// CIR-SAME: nothrow, nounwind
-// CIR: cir.call @caller(%{{.*}}) nothrow nounwind
-
-// LLVM: ; Function Attrs: {{.*}}nounwind
-// LLVM-NEXT: define {{.*}}@caller({{.*}}){{.*}} #{{[0-9]+}} {
-// LLVM-NOT: invoke
-// LLVM: call {{.*}}@ext({{.*}}) #[[CALL_ATTR:[0-9]+]]
-// LLVM-NOT: invoke
-// LLVM: call void @_ZN1DD1Ev({{.*}}) #[[CALL_ATTR]]
-// LLVM-NOT: landingpad
-// LLVM: ; Function Attrs: {{.*}}nounwind
-// LLVM-NEXT: declare {{.*}}@ext(
-// LLVM: ; Function Attrs: {{.*}}nounwind
-// LLVM-NEXT: declare {{.*}}@_ZN1DD1Ev(
-// LLVM: ; Function Attrs: {{.*}}nounwind
-// LLVM-NEXT: define {{.*}}amdgpu_kernel void @kernel({{.*}}){{.*}} #{{[0-9]+}} {
-// LLVM: call {{.*}}@caller({{.*}}) #[[CALL_ATTR]]
-// LLVM: attributes #[[CALL_ATTR]] = {{{.*}}nounwind
More information about the cfe-commits
mailing list