Skip to content

[OpenACC][CIR] Implement 'host_data' lowering, plus all clauses #143136

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Merged
merged 1 commit into from
Jun 6, 2025
Merged

[OpenACC][CIR] Implement 'host_data' lowering, plus all clauses #143136

merged 1 commit into from
Jun 6, 2025

Conversation

erichkeane
Copy link
Collaborator

'host_data' has its own Op kind, so this handles the lowering there, it looks exactly like the other ones we've done so far, so nothing novel here.

host_data takes 3 clauses, 1 of which is required.

'use_device' is required, and results in an acc.use_device operation,
which then feeds into the dataOperands list on acc.host_data.

'if_present' is a simple attribute on the operand.

'if' is a condition on the operand, identical to our other handling of 'if'.

This patch handles all of these.

@erichkeane
[OpenACC][CIR] Implement 'host_data' lowering, plus all clauses
8419c28 
'host_data' has its own Op kind, so this handles the lowering there, it
looks exactly like the other ones we've done so far, so nothing novel
here.

host_data takes 3 clauses, 1 of which is required.

'use_device' is required, and results in an acc.use_device operation,
  which then feeds into the dataOperands list on acc.host_data.

'if_present' is a simple attribute on the operand.

'if' is a condition on the operand, identical to our other handling of
'if'.

This patch handles all of these.
@llvmbot llvmbot added clang Clang issues not falling into any other category ClangIR Anything related to the ClangIR project labels Jun 6, 2025
@llvmbot
Copy link
Member

llvmbot commented Jun 6, 2025

@llvm/pr-subscribers-clangir

Author: Erich Keane (erichkeane)

Changes

'host_data' has its own Op kind, so this handles the lowering there, it looks exactly like the other ones we've done so far, so nothing novel here.

host_data takes 3 clauses, 1 of which is required.

'use_device' is required, and results in an acc.use_device operation,
which then feeds into the dataOperands list on acc.host_data.

'if_present' is a simple attribute on the operand.

'if' is a condition on the operand, identical to our other handling of 'if'.

This patch handles all of these.

Full diff: https://github.com/llvm/llvm-project/pull/143136.diff

3 Files Affected:

  • (modified) clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp (+40-1)
  • (modified) clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp (+11-5)
  • (added) clang/test/CIR/CodeGenOpenACC/host_data.c (+55)
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index f41f776225152..e3657e9014121 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -319,6 +319,21 @@ class OpenACCClauseCIREmitter final
     dataOperands.push_back(afterOp.getOperation());
   }
 
+  template <typename BeforeOpTy>
+  void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
+                      bool structured, bool implicit) {
+    DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand);
+    auto beforeOp =
+        builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured,
+                                   implicit, opInfo.name, opInfo.bounds);
+    operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
+
+    // Set the 'rest' of the info for the operation.
+    beforeOp.setDataClause(dataClause);
+    // Make sure we record these, so 'async' values can be updated later.
+    dataOperands.push_back(beforeOp.getOperation());
+  }
+
   // Helper function that covers for the fact that we don't have this function
   // on all operation types.
   mlir::ArrayAttr getAsyncOnlyAttr() {
@@ -550,7 +565,8 @@ class OpenACCClauseCIREmitter final
     if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
                                mlir::acc::KernelsOp, mlir::acc::InitOp,
                                mlir::acc::ShutdownOp, mlir::acc::SetOp,
-                               mlir::acc::DataOp, mlir::acc::WaitOp>) {
+                               mlir::acc::DataOp, mlir::acc::WaitOp,
+                               mlir::acc::HostDataOp>) {
       operation.getIfCondMutable().append(
           createCondition(clause.getConditionExpr()));
     } else if constexpr (isCombinedType<OpTy>) {
@@ -566,6 +582,17 @@ class OpenACCClauseCIREmitter final
     }
   }
 
+  void VisitIfPresentClause(const OpenACCIfPresentClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
+      operation.setIfPresent(true);
+    } else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
+      // Last unimplemented one here, so just put it in this way instead.
+      return clauseNotImplemented(clause);
+    } else {
+      llvm_unreachable("unknown construct kind in VisitIfPresentClause");
+    }
+  }
+
   void VisitDeviceNumClause(const OpenACCDeviceNumClause &clause) {
     if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp,
                                mlir::acc::SetOp>) {
@@ -791,6 +818,17 @@ class OpenACCClauseCIREmitter final
       return clauseNotImplemented(clause);
     }
   }
+
+  void VisitUseDeviceClause(const OpenACCUseDeviceClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
+      for (auto var : clause.getVarList())
+        addDataOperand<mlir::acc::UseDeviceOp>(
+            var, mlir::acc::DataClause::acc_use_device,
+            /*structured=*/true, /*implicit=*/false);
+    } else {
+      llvm_unreachable("Unknown construct kind in VisitUseDeviceClause");
+    }
+  }
 };
 
 template <typename OpTy>
@@ -826,6 +864,7 @@ EXPL_SPEC(mlir::acc::InitOp)
 EXPL_SPEC(mlir::acc::ShutdownOp)
 EXPL_SPEC(mlir::acc::SetOp)
 EXPL_SPEC(mlir::acc::WaitOp)
+EXPL_SPEC(mlir::acc::HostDataOp)
 #undef EXPL_SPEC
 
 template <typename ComputeOp, typename LoopOp>
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index d922ca0c74d5d..2aab9cecf93d8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -235,6 +235,17 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCCombinedConstruct(
     llvm_unreachable("invalid compute construct kind");
   }
 }
+
+mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(
+    const OpenACCHostDataConstruct &s) {
+  mlir::Location start = getLoc(s.getSourceRange().getBegin());
+  mlir::Location end = getLoc(s.getSourceRange().getEnd());
+
+  return emitOpenACCOpAssociatedStmt<HostDataOp, mlir::acc::TerminatorOp>(
+      start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
+      s.getStructuredBlock());
+}
+
 mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
     const OpenACCEnterDataConstruct &s) {
   cgm.errorNYI(s.getSourceRange(), "OpenACC EnterData Construct");
@@ -245,11 +256,6 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
   cgm.errorNYI(s.getSourceRange(), "OpenACC ExitData Construct");
   return mlir::failure();
 }
-mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(
-    const OpenACCHostDataConstruct &s) {
-  cgm.errorNYI(s.getSourceRange(), "OpenACC HostData Construct");
-  return mlir::failure();
-}
 mlir::LogicalResult
 CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
   cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct");
diff --git a/clang/test/CIR/CodeGenOpenACC/host_data.c b/clang/test/CIR/CodeGenOpenACC/host_data.c
new file mode 100644
index 0000000000000..4c3f7dd092a2f
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/host_data.c
@@ -0,0 +1,55 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_host_data(int cond, int var1, int var2) {
+  // CHECK: cir.func @acc_host_data(%[[ARG_COND:.*]]: !s32i {{.*}}, %[[ARG_V1:.*]]: !s32i {{.*}}, %[[ARG_V2:.*]]: !s32i {{.*}}) {
+  // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
+  // CHECK-NEXT: %[[V1:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["var1", init]
+  // CHECK-NEXT: %[[V2:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["var2", init]
+  // CHECK-NEXT: cir.store %[[ARG_COND]], %[[COND]] : !s32i, !cir.ptr<!s32i>
+  // CHECK-NEXT: cir.store %[[ARG_V1]], %[[V1]] : !s32i, !cir.ptr<!s32i>
+  // CHECK-NEXT: cir.store %[[ARG_V2]], %[[V2]] : !s32i, !cir.ptr<!s32i>
+
+#pragma acc host_data use_device(var1)
+  {}
+  // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+  // CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]] : !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+#pragma acc host_data use_device(var1, var2)
+  {}
+  // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+  // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
+  // CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+
+#pragma acc host_data use_device(var1, var2) if_present
+  {}
+  // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+  // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
+  // CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {ifPresent}
+
+#pragma acc host_data use_device(var1, var2) if(cond)
+  {}
+  // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+  // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
+  // CHECK-NEXT: %[[LOAD_COND:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[COND_BOOL:.*]] = cir.cast(int_to_bool, %[[LOAD_COND]] : !s32i), !cir.bool
+  // CHECK-NEXT: %[[COND_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_BOOL]] : !cir.bool to i1
+  // CHECK-NEXT: acc.host_data if(%[[COND_CAST]]) dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+
+#pragma acc host_data use_device(var1, var2) if(cond) if_present
+  {}
+  // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+  // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
+  // CHECK-NEXT: %[[LOAD_COND:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[COND_BOOL:.*]] = cir.cast(int_to_bool, %[[LOAD_COND]] : !s32i), !cir.bool
+  // CHECK-NEXT: %[[COND_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_BOOL]] : !cir.bool to i1
+  // CHECK-NEXT: acc.host_data if(%[[COND_CAST]]) dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {ifPresent}
+}

@llvmbot
Copy link
Member

llvmbot commented Jun 6, 2025

@llvm/pr-subscribers-clang

Author: Erich Keane (erichkeane)

Changes

'host_data' has its own Op kind, so this handles the lowering there, it looks exactly like the other ones we've done so far, so nothing novel here.

host_data takes 3 clauses, 1 of which is required.

'use_device' is required, and results in an acc.use_device operation,
which then feeds into the dataOperands list on acc.host_data.

'if_present' is a simple attribute on the operand.

'if' is a condition on the operand, identical to our other handling of 'if'.

This patch handles all of these.

Full diff: https://github.com/llvm/llvm-project/pull/143136.diff

3 Files Affected:

  • (modified) clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp (+40-1)
  • (modified) clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp (+11-5)
  • (added) clang/test/CIR/CodeGenOpenACC/host_data.c (+55)
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index f41f776225152..e3657e9014121 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -319,6 +319,21 @@ class OpenACCClauseCIREmitter final
     dataOperands.push_back(afterOp.getOperation());
   }
 
+  template <typename BeforeOpTy>
+  void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
+                      bool structured, bool implicit) {
+    DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand);
+    auto beforeOp =
+        builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured,
+                                   implicit, opInfo.name, opInfo.bounds);
+    operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
+
+    // Set the 'rest' of the info for the operation.
+    beforeOp.setDataClause(dataClause);
+    // Make sure we record these, so 'async' values can be updated later.
+    dataOperands.push_back(beforeOp.getOperation());
+  }
+
   // Helper function that covers for the fact that we don't have this function
   // on all operation types.
   mlir::ArrayAttr getAsyncOnlyAttr() {
@@ -550,7 +565,8 @@ class OpenACCClauseCIREmitter final
     if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
                                mlir::acc::KernelsOp, mlir::acc::InitOp,
                                mlir::acc::ShutdownOp, mlir::acc::SetOp,
-                               mlir::acc::DataOp, mlir::acc::WaitOp>) {
+                               mlir::acc::DataOp, mlir::acc::WaitOp,
+                               mlir::acc::HostDataOp>) {
       operation.getIfCondMutable().append(
           createCondition(clause.getConditionExpr()));
     } else if constexpr (isCombinedType<OpTy>) {
@@ -566,6 +582,17 @@ class OpenACCClauseCIREmitter final
     }
   }
 
+  void VisitIfPresentClause(const OpenACCIfPresentClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
+      operation.setIfPresent(true);
+    } else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
+      // Last unimplemented one here, so just put it in this way instead.
+      return clauseNotImplemented(clause);
+    } else {
+      llvm_unreachable("unknown construct kind in VisitIfPresentClause");
+    }
+  }
+
   void VisitDeviceNumClause(const OpenACCDeviceNumClause &clause) {
     if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp,
                                mlir::acc::SetOp>) {
@@ -791,6 +818,17 @@ class OpenACCClauseCIREmitter final
       return clauseNotImplemented(clause);
     }
   }
+
+  void VisitUseDeviceClause(const OpenACCUseDeviceClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
+      for (auto var : clause.getVarList())
+        addDataOperand<mlir::acc::UseDeviceOp>(
+            var, mlir::acc::DataClause::acc_use_device,
+            /*structured=*/true, /*implicit=*/false);
+    } else {
+      llvm_unreachable("Unknown construct kind in VisitUseDeviceClause");
+    }
+  }
 };
 
 template <typename OpTy>
@@ -826,6 +864,7 @@ EXPL_SPEC(mlir::acc::InitOp)
 EXPL_SPEC(mlir::acc::ShutdownOp)
 EXPL_SPEC(mlir::acc::SetOp)
 EXPL_SPEC(mlir::acc::WaitOp)
+EXPL_SPEC(mlir::acc::HostDataOp)
 #undef EXPL_SPEC
 
 template <typename ComputeOp, typename LoopOp>
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index d922ca0c74d5d..2aab9cecf93d8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -235,6 +235,17 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCCombinedConstruct(
     llvm_unreachable("invalid compute construct kind");
   }
 }
+
+mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(
+    const OpenACCHostDataConstruct &s) {
+  mlir::Location start = getLoc(s.getSourceRange().getBegin());
+  mlir::Location end = getLoc(s.getSourceRange().getEnd());
+
+  return emitOpenACCOpAssociatedStmt<HostDataOp, mlir::acc::TerminatorOp>(
+      start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
+      s.getStructuredBlock());
+}
+
 mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
     const OpenACCEnterDataConstruct &s) {
   cgm.errorNYI(s.getSourceRange(), "OpenACC EnterData Construct");
@@ -245,11 +256,6 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
   cgm.errorNYI(s.getSourceRange(), "OpenACC ExitData Construct");
   return mlir::failure();
 }
-mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(
-    const OpenACCHostDataConstruct &s) {
-  cgm.errorNYI(s.getSourceRange(), "OpenACC HostData Construct");
-  return mlir::failure();
-}
 mlir::LogicalResult
 CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
   cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct");
diff --git a/clang/test/CIR/CodeGenOpenACC/host_data.c b/clang/test/CIR/CodeGenOpenACC/host_data.c
new file mode 100644
index 0000000000000..4c3f7dd092a2f
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/host_data.c
@@ -0,0 +1,55 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_host_data(int cond, int var1, int var2) {
+  // CHECK: cir.func @acc_host_data(%[[ARG_COND:.*]]: !s32i {{.*}}, %[[ARG_V1:.*]]: !s32i {{.*}}, %[[ARG_V2:.*]]: !s32i {{.*}}) {
+  // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
+  // CHECK-NEXT: %[[V1:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["var1", init]
+  // CHECK-NEXT: %[[V2:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["var2", init]
+  // CHECK-NEXT: cir.store %[[ARG_COND]], %[[COND]] : !s32i, !cir.ptr<!s32i>
+  // CHECK-NEXT: cir.store %[[ARG_V1]], %[[V1]] : !s32i, !cir.ptr<!s32i>
+  // CHECK-NEXT: cir.store %[[ARG_V2]], %[[V2]] : !s32i, !cir.ptr<!s32i>
+
+#pragma acc host_data use_device(var1)
+  {}
+  // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+  // CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]] : !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+#pragma acc host_data use_device(var1, var2)
+  {}
+  // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+  // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
+  // CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+
+#pragma acc host_data use_device(var1, var2) if_present
+  {}
+  // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+  // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
+  // CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {ifPresent}
+
+#pragma acc host_data use_device(var1, var2) if(cond)
+  {}
+  // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+  // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
+  // CHECK-NEXT: %[[LOAD_COND:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[COND_BOOL:.*]] = cir.cast(int_to_bool, %[[LOAD_COND]] : !s32i), !cir.bool
+  // CHECK-NEXT: %[[COND_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_BOOL]] : !cir.bool to i1
+  // CHECK-NEXT: acc.host_data if(%[[COND_CAST]]) dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+
+#pragma acc host_data use_device(var1, var2) if(cond) if_present
+  {}
+  // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+  // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
+  // CHECK-NEXT: %[[LOAD_COND:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[COND_BOOL:.*]] = cir.cast(int_to_bool, %[[LOAD_COND]] : !s32i), !cir.bool
+  // CHECK-NEXT: %[[COND_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_BOOL]] : !cir.bool to i1
+  // CHECK-NEXT: acc.host_data if(%[[COND_CAST]]) dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {ifPresent}
+}

clementval
clementval approved these changes Jun 6, 2025
View reviewed changes
Copy link
Contributor

@clementval clementval left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM!

andykaylor
andykaylor approved these changes Jun 6, 2025
View reviewed changes
Copy link
Contributor

@andykaylor andykaylor left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

lgtm

@erichkeane erichkeane merged commit c02403e into llvm:main Jun 6, 2025
10 checks passed
rorth pushed a commit to rorth/llvm-project that referenced this pull request Jun 11, 2025
@erichkeane @rorth
[OpenACC][CIR] Implement 'host_data' lowering, plus all clauses (llvm…
ad069df 
…#143136)

'host_data' has its own Op kind, so this handles the lowering there, it
looks exactly like the other ones we've done so far, so nothing novel
here.

host_data takes 3 clauses, 1 of which is required.

'use_device' is required, and results in an acc.use_device operation,
  which then feeds into the dataOperands list on acc.host_data.

'if_present' is a simple attribute on the operand.

'if' is a condition on the operand, identical to our other handling of
'if'.

This patch handles all of these.
DhruvSrivastavaX pushed a commit to DhruvSrivastavaX/lldb-for-aix that referenced this pull request Jun 12, 2025
@erichkeane @DhruvSrivastavaX
[OpenACC][CIR] Implement 'host_data' lowering, plus all clauses (llvm…
9fc3c3a 
…#143136)

'host_data' has its own Op kind, so this handles the lowering there, it
looks exactly like the other ones we've done so far, so nothing novel
here.

host_data takes 3 clauses, 1 of which is required.

'use_device' is required, and results in an acc.use_device operation,
  which then feeds into the dataOperands list on acc.host_data.

'if_present' is a simple attribute on the operand.

'if' is a condition on the operand, identical to our other handling of
'if'.

This patch handles all of these.
tomtor pushed a commit to tomtor/llvm-project that referenced this pull request Jun 14, 2025
@erichkeane @tomtor
[OpenACC][CIR] Implement 'host_data' lowering, plus all clauses (llvm…
e17a5a0 
…#143136)

'host_data' has its own Op kind, so this handles the lowering there, it
looks exactly like the other ones we've done so far, so nothing novel
here.

host_data takes 3 clauses, 1 of which is required.

'use_device' is required, and results in an acc.use_device operation,
  which then feeds into the dataOperands list on acc.host_data.

'if_present' is a simple attribute on the operand.

'if' is a condition on the operand, identical to our other handling of
'if'.

This patch handles all of these.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@clementval clementval clementval approved these changes

@andykaylor andykaylor andykaylor approved these changes

@razvanlupusoru razvanlupusoru Awaiting requested review from razvanlupusoru

@lanza lanza Awaiting requested review from lanza lanza is a code owner

@bcardosolopes bcardosolopes Awaiting requested review from bcardosolopes bcardosolopes is a code owner

Assignees
No one assigned
Labels
clang Clang issues not falling into any other category ClangIR Anything related to the ClangIR project
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
4 participants
@erichkeane @llvmbot @clementval @andykaylor