[clang] [OpenACC][CIR] Implement 'nohost' lowering. (PR #170369)

[Erich Keane via cfe-commits]

https://github.com/erichkeane created https://github.com/llvm/llvm-project/pull/170369

This clause is pretty small/trivial and is a simple 'set a bool' value on the IR node, so its implementation is quite simple.  We create the Operation with this as 'false', so the 'nohost' marks it as true always.

>From ca1ca1c8f639cb67dda298448abe898cfaa9ded0 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Tue, 2 Dec 2025 09:28:31 -0800
Subject: [PATCH] [OpenACC][CIR] Implement 'nohost' lowering.

This clause is pretty small/trivial and is a simple 'set a bool' value
on the IR node, so its implementation is quite simple.  We create the
Operation with this as 'false', so the 'nohost' marks it as true always.
---
 clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp       |  4 ++++
 clang/test/CIR/CodeGenOpenACC/routine-clauses.cpp | 14 +++++++-------
 2 files changed, 11 insertions(+), 7 deletions(-)

diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index 19def59db14ba..a5322ac4e1930 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -329,6 +329,10 @@ class OpenACCRoutineClauseEmitter final
   void VisitVectorClause(const OpenACCVectorClause &clause) {
     routineOp.addVector(builder.getContext(), lastDeviceTypeValues);
   }
+
+  void VisitNoHostClause(const OpenACCNoHostClause &clause) {
+    routineOp.setNohost(/*attrValue=*/true);
+  }
 };
 } // namespace
 
diff --git a/clang/test/CIR/CodeGenOpenACC/routine-clauses.cpp b/clang/test/CIR/CodeGenOpenACC/routine-clauses.cpp
index 7ad64f2d05158..81437e7e02ab1 100644
--- a/clang/test/CIR/CodeGenOpenACC/routine-clauses.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/routine-clauses.cpp
@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
 
-#pragma acc routine seq
+#pragma acc routine seq nohost
 void Func1() {}
 
 void Func2() {}
@@ -10,16 +10,16 @@ void Func2() {}
 void Func3() {}
 
 void Func4() {}
-#pragma acc routine(Func4) worker
+#pragma acc routine(Func4) worker nohost
 
-#pragma acc routine vector
+#pragma acc routine nohost vector
 void Func5() {}
 
 void Func6() {}
-#pragma acc routine(Func6) vector
+#pragma acc routine(Func6) nohost vector
 
 // CHECK: cir.func{{.*}} @[[F1_NAME:.*Func1[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F1_R_NAME:.*]]]>}
-// CHECK: acc.routine @[[F1_R_NAME]] func(@[[F1_NAME]]) seq
+// CHECK: acc.routine @[[F1_R_NAME]] func(@[[F1_NAME]]) seq nohost
 
 // CHECK: cir.func{{.*}} @[[F2_NAME:.*Func2[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F2_R_NAME:.*]]]>}
 
@@ -34,5 +34,5 @@ void Func6() {}
 // CHECK: cir.func{{.*}} @[[F6_NAME:.*Func6[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F6_R_NAME:.*]]]>}
 
 // CHECK: acc.routine @[[F2_R_NAME]] func(@[[F2_NAME]]) seq
-// CHECK: acc.routine @[[F4_R_NAME]] func(@[[F4_NAME]]) worker
-// CHECK: acc.routine @[[F6_R_NAME]] func(@[[F6_NAME]]) vector
+// CHECK: acc.routine @[[F4_R_NAME]] func(@[[F4_NAME]]) worker nohost
+// CHECK: acc.routine @[[F6_R_NAME]] func(@[[F6_NAME]]) vector nohost