[Mlir-commits] [flang] [mlir] [mlir][acc] Add OffloadTargetVerifier pass (PR #176467)

[llvmlistbot at llvm.org]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-mlir-openacc

Author: Razvan Lupusoru (razvanlupusoru)

<details>
<summary>Changes</summary>

Add a verification pass that checks live-in values and symbol references within offload regions are legal for the target execution model.

When code is offloaded to a device (e.g., GPU), not all values and symbols from the host context are directly accessible. Data must be explicitly mapped via OpenACC data clauses (copyin, create, present etc.), declared with device attributes, or be trivial scalars that can be passed by value. Similarly, symbol references to globals must have proper `declare` attributes or device-resident data attributes.

This pass walks operations implementing `OffloadRegionOpInterface`, which includes OpenACC compute constructs (`acc.parallel`, `acc.kernels`, `acc.serial`) as well as GPU operations like `gpu.launch`. For each region, it uses liveness analysis to identify values flowing into the region and checks their validity using the `OpenACCSupport` analysis.

Key features:
- Validates live-in values against OpenACC data mapping requirements
- Validates symbol references for device accessibility
- Supports soft-check mode for diagnostic-only verification
- Configurable device_type for target-specific behavior

---

Patch is 32.52 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/176467.diff


8 Files Affected:

- (modified) flang/include/flang/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.h (+3) 
- (modified) flang/lib/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.cpp (+16) 
- (added) flang/test/Transforms/OpenACC/offload-target-verifier.fir (+313) 
- (modified) mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td (+45) 
- (modified) mlir/lib/Dialect/OpenACC/Transforms/CMakeLists.txt (+1) 
- (added) mlir/lib/Dialect/OpenACC/Transforms/OffloadTargetVerifier.cpp (+236) 
- (modified) mlir/lib/Dialect/OpenACC/Utils/OpenACCUtils.cpp (+7) 
- (added) mlir/test/Dialect/OpenACC/offload-target-verifier.mlir (+231) 


``````````diff
diff --git a/flang/include/flang/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.h b/flang/include/flang/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.h
index c798681306c10..f5d44c7968b1d 100644
--- a/flang/include/flang/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.h
+++ b/flang/include/flang/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.h
@@ -14,6 +14,7 @@
 #define FORTRAN_OPTIMIZER_OPENACC_ANALYSIS_FIROPENACCSUPPORTANALYSIS_H
 
 #include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/Region.h"
 #include "mlir/IR/Value.h"
 #include <string>
 
@@ -43,6 +44,8 @@ class FIROpenACCSupportAnalysis {
 
   mlir::InFlightDiagnostic emitNYI(mlir::Location loc,
                                    const mlir::Twine &message);
+
+  bool isValidValueUse(mlir::Value v, mlir::Region &region);
 };
 
 } // namespace acc
diff --git a/flang/lib/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.cpp b/flang/lib/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.cpp
index 8cdbe1d5b170e..3ad3188314fbc 100644
--- a/flang/lib/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.cpp
+++ b/flang/lib/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.cpp
@@ -11,8 +11,11 @@
 //===----------------------------------------------------------------------===//
 
 #include "flang/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.h"
+
 #include "flang/Optimizer/Builder/Todo.h"
+#include "flang/Optimizer/Dialect/FIRType.h"
 #include "flang/Optimizer/OpenACC/Support/FIROpenACCUtils.h"
+#include "mlir/Dialect/OpenACC/OpenACCUtils.h"
 
 using namespace mlir;
 
@@ -36,5 +39,18 @@ FIROpenACCSupportAnalysis::emitNYI(Location loc, const Twine &message) {
   return mlir::emitError(loc, "not yet implemented: " + message.str());
 }
 
+bool FIROpenACCSupportAnalysis::isValidValueUse(Value v, Region &region) {
+  // First check using the base utility.
+  if (mlir::acc::isValidValueUse(v, region))
+    return true;
+
+  // FIR-specific: fir.logical is a trivial scalar type that can be
+  // passed by value.
+  if (mlir::isa<fir::LogicalType>(v.getType()))
+    return true;
+
+  return false;
+}
+
 } // namespace acc
 } // namespace fir
diff --git a/flang/test/Transforms/OpenACC/offload-target-verifier.fir b/flang/test/Transforms/OpenACC/offload-target-verifier.fir
new file mode 100644
index 0000000000000..3056503ed1b1c
--- /dev/null
+++ b/flang/test/Transforms/OpenACC/offload-target-verifier.fir
@@ -0,0 +1,313 @@
+// RUN: fir-opt %s --pass-pipeline="builtin.module(acc-initialize-fir-analyses,func.func(offload-target-verifier{soft-check=true}))" --verify-diagnostics -split-input-file
+
+// Test scalar i32 live-in value - should pass (scalars can be passed by value)
+func.func @test_scalar_i32() {
+  %alloca = fir.alloca i32
+  %livein = fir.load %alloca : !fir.ref<i32>
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %accalloca = fir.alloca i32
+    fir.store %livein to %accalloca : !fir.ref<i32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test fir.logical scalar - should pass (scalars can be passed by value)
+func.func @test_fir_logical_scalar() {
+  %alloca = fir.alloca !fir.logical<4>
+  %livein = fir.load %alloca : !fir.ref<!fir.logical<4>>
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %accalloca = fir.alloca !fir.logical<4>
+    fir.store %livein to %accalloca : !fir.ref<!fir.logical<4>>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test fir.ref live-in without data clause - should fail
+func.func @test_fir_ref() {
+  // expected-note @below {{value}}
+  %livein = fir.alloca f32
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  acc.serial {
+    %load = fir.load %livein : !fir.ref<f32>
+    %accalloca = fir.alloca f32
+    fir.store %load to %accalloca : !fir.ref<f32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test fir.ref with copyin data clause - should pass
+func.func @test_fir_ref_copyin() {
+  %alloca = fir.alloca f32
+  %livein = acc.copyin varPtr(%alloca : !fir.ref<f32>) -> !fir.ref<f32>
+  // expected-remark @below {{passed validity check}}
+  acc.serial dataOperands(%livein : !fir.ref<f32>) {
+    %load = fir.load %livein : !fir.ref<f32>
+    %accalloca = fir.alloca f32
+    fir.store %load to %accalloca : !fir.ref<f32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test fir.ref with private clause - should pass
+acc.private.recipe @privatization_ref_f32 : !fir.ref<f32> init {
+^bb0(%arg0: !fir.ref<f32>):
+  %0 = fir.alloca f32
+  acc.yield %0 : !fir.ref<f32>
+}
+
+func.func @test_fir_ref_private() {
+  %livein = fir.alloca f32
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %private = acc.private varPtr(%livein : !fir.ref<f32>) recipe(@privatization_ref_f32) -> !fir.ref<f32>
+    %load = fir.load %private : !fir.ref<f32>
+    %accalloca = fir.alloca f32
+    fir.store %load to %accalloca : !fir.ref<f32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test fir.global without declare attribute - should fail
+fir.global @_global_array : !fir.array<10xf32> {
+  %0 = fir.zero_bits !fir.array<10xf32>
+  fir.has_value %0 : !fir.array<10xf32>
+}
+
+func.func @test_fir_global_no_declare() {
+  // expected-warning @below {{illegal symbol(s): _global_array}}
+  acc.serial {
+    %liveinsym = fir.address_of(@_global_array) : !fir.ref<!fir.array<10xf32>>
+    %loaded = fir.load %liveinsym : !fir.ref<!fir.array<10xf32>>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test fir.global with acc.declare attribute - should pass
+fir.global @_global_array_declared {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.array<10xf32> {
+  %0 = fir.zero_bits !fir.array<10xf32>
+  fir.has_value %0 : !fir.array<10xf32>
+}
+
+func.func @test_fir_global_with_declare() {
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %liveinsym = fir.address_of(@_global_array_declared) : !fir.ref<!fir.array<10xf32>>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test fir.global with CUDA device attribute - should pass
+fir.global @_cuda_global_array {data_attr = #cuf.cuda<device>} : !fir.array<10xf32> {
+  %0 = fir.zero_bits !fir.array<10xf32>
+  fir.has_value %0 : !fir.array<10xf32>
+}
+
+func.func @test_fir_cuda_global() {
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %liveinsym = fir.address_of(@_cuda_global_array) : !fir.ref<!fir.array<10xf32>>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test fir.declare with CUDA device attribute - should pass
+func.func @test_fir_declare_cuda() {
+  %c10 = arith.constant 10 : index
+  %0 = fir.alloca !fir.array<10xf32>
+  %1 = fir.shape %c10 : (index) -> !fir.shape<1>
+  %2 = fir.declare %0(%1) {data_attr = #cuf.cuda<device>, uniq_name = "cuda_array"} : (!fir.ref<!fir.array<10xf32>>, !fir.shape<1>) -> !fir.ref<!fir.array<10xf32>>
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %3 = fir.load %2 : !fir.ref<!fir.array<10xf32>>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test hlfir.declare live-in without data clause - should fail
+func.func @test_hlfir_declare(%arg0: !fir.ref<f32> {fir.bindc_name = "var"}) {
+  %0 = fir.dummy_scope : !fir.dscope
+  // expected-note @below {{value}}
+  %1:2 = hlfir.declare %arg0 dummy_scope %0 {uniq_name = "_QEvar"} : (!fir.ref<f32>, !fir.dscope) -> (!fir.ref<f32>, !fir.ref<f32>)
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  acc.serial {
+    %cst = arith.constant 1.000000e+00 : f32
+    hlfir.assign %cst to %1#0 : f32, !fir.ref<f32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test acc.parallel region
+func.func @test_acc_parallel() {
+  // expected-note @below {{value}}
+  %alloca = fir.alloca f32
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  acc.parallel {
+    %load = fir.load %alloca : !fir.ref<f32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test acc.kernels region
+func.func @test_acc_kernels() {
+  // expected-note @below {{value}}
+  %alloca = fir.alloca f32
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  acc.kernels {
+    %load = fir.load %alloca : !fir.ref<f32>
+    acc.terminator
+  }
+  return
+}
+
+// -----
+
+// Test cuf.kernel region with invalid live-in - should fail
+func.func @test_cuf_kernel_invalid() {
+  %c1 = arith.constant 1 : index
+  %c1_i32 = arith.constant 1 : i32
+  // expected-note @below {{value}}
+  %alloca = fir.alloca f32
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  cuf.kernel<<<%c1_i32, %c1_i32>>> (%arg0 : index) = (%c1 : index) to (%c1 : index) step (%c1 : index) {
+    %load = fir.load %alloca : !fir.ref<f32>
+    "fir.end"() : () -> ()
+  }
+  return
+}
+
+// -----
+
+// Test cuf.kernel region with CUDA device scalar - should pass
+func.func @test_cuf_kernel_cuda_device() {
+  %c1 = arith.constant 1 : index
+  %c1_i32 = arith.constant 1 : i32
+  %alloca = fir.alloca f32
+  %decl = fir.declare %alloca {data_attr = #cuf.cuda<device>, uniq_name = "cuda_scalar"} : (!fir.ref<f32>) -> !fir.ref<f32>
+  // expected-remark @below {{passed validity check}}
+  cuf.kernel<<<%c1_i32, %c1_i32>>> (%arg0 : index) = (%c1 : index) to (%c1 : index) step (%c1 : index) {
+    %load = fir.load %decl : !fir.ref<f32>
+    "fir.end"() : () -> ()
+  }
+  return
+}
+
+// -----
+
+// Test that fir.shape live-in to cuf.kernel is illegal
+func.func @test_cuf_kernel_shape_illegal() {
+  %c1 = arith.constant 1 : index
+  %c1_i32 = arith.constant 1 : i32
+  %c10 = arith.constant 10 : index
+  %alloca = fir.alloca !fir.array<10xf32>
+  // expected-note @below {{value}}
+  %shape = fir.shape %c10 : (index) -> !fir.shape<1>
+  %decl = fir.declare %alloca(%shape) {data_attr = #cuf.cuda<device>, uniq_name = "cuda_arr"} : (!fir.ref<!fir.array<10xf32>>, !fir.shape<1>) -> !fir.ref<!fir.array<10xf32>>
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  cuf.kernel<<<%c1_i32, %c1_i32>>> (%arg0 : index) = (%c1 : index) to (%c10 : index) step (%c1 : index) {
+    %coor = fir.array_coor %decl(%shape) %arg0 : (!fir.ref<!fir.array<10xf32>>, !fir.shape<1>, index) -> !fir.ref<f32>
+    %load = fir.load %coor : !fir.ref<f32>
+    "fir.end"() : () -> ()
+  }
+  return
+}
+
+// -----
+
+// Test cuf.kernel region with cuf.alloc device data - should pass
+func.func @test_cuf_kernel_cuf_alloc() {
+  %c1 = arith.constant 1 : index
+  %c1_i32 = arith.constant 1 : i32
+  %alloca = cuf.alloc f32 {data_attr = #cuf.cuda<device>} -> !fir.ref<f32>
+  // expected-remark @below {{passed validity check}}
+  cuf.kernel<<<%c1_i32, %c1_i32>>> (%arg0 : index) = (%c1 : index) to (%c1 : index) step (%c1 : index) {
+    %load = fir.load %alloca : !fir.ref<f32>
+    "fir.end"() : () -> ()
+  }
+  return
+}
+
+// -----
+
+// Test fir.rebox live-in - should fail (box without data clause)
+func.func @test_fir_rebox(%arg0: !fir.box<!fir.array<?xf32>> {fir.bindc_name = "arr"}) {
+  %0 = fir.dummy_scope : !fir.dscope
+  %1 = fir.declare %arg0 dummy_scope %0 {uniq_name = "_QEarr"} : (!fir.box<!fir.array<?xf32>>, !fir.dscope) -> !fir.box<!fir.array<?xf32>>
+  // expected-note @below {{value}}
+  %2 = fir.rebox %1 : (!fir.box<!fir.array<?xf32>>) -> !fir.box<!fir.array<?xf32>>
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  acc.serial {
+    %c0 = arith.constant 0 : index
+    %3:3 = fir.box_dims %2, %c0 : (!fir.box<!fir.array<?xf32>>, index) -> (index, index, index)
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test fir.rebox with CUDA device attribute - should pass
+func.func @test_fir_rebox_cuda(%arg0: !fir.box<!fir.array<?xf32>> {cuf.data_attr = #cuf.cuda<device>, fir.bindc_name = "arr"}) {
+  %0 = fir.dummy_scope : !fir.dscope
+  %1 = fir.declare %arg0 dummy_scope %0 {data_attr = #cuf.cuda<device>, uniq_name = "_QEarr"} : (!fir.box<!fir.array<?xf32>>, !fir.dscope) -> !fir.box<!fir.array<?xf32>>
+  %2 = fir.rebox %1 : (!fir.box<!fir.array<?xf32>>) -> !fir.box<!fir.array<?xf32>>
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %c0 = arith.constant 0 : index
+    %3:3 = fir.box_dims %2, %c0 : (!fir.box<!fir.array<?xf32>>, index) -> (index, index, index)
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test fir.embox live-in - should fail
+func.func @test_fir_embox() {
+  %c10 = arith.constant 10 : index
+  %0 = fir.alloca !fir.array<10xf32>
+  %1 = fir.shape %c10 : (index) -> !fir.shape<1>
+  // expected-note @below {{value}}
+  %2 = fir.embox %0(%1) : (!fir.ref<!fir.array<10xf32>>, !fir.shape<1>) -> !fir.box<!fir.array<10xf32>>
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  acc.serial {
+    %c0 = arith.constant 0 : index
+    %3:3 = fir.box_dims %2, %c0 : (!fir.box<!fir.array<10xf32>>, index) -> (index, index, index)
+    acc.yield
+  }
+  return
+}
diff --git a/mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td b/mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td
index 94a4f8732fafa..37243ecf4e1ac 100644
--- a/mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td
@@ -363,4 +363,49 @@ def OffloadLiveInValueCanonicalization : Pass<"offload-livein-value-canonicaliza
   let dependentDialects = ["mlir::acc::OpenACCDialect"];
 }
 
+def OffloadTargetVerifier : Pass<"offload-target-verifier", "mlir::func::FuncOp"> {
+  let summary = "Verify values and symbols live into offload regions for legality";
+  let description = [{
+    This pass verifies that values and symbols used within OpenACC compute
+    constructs and other offload regions are legal for the target execution
+    model.
+
+    The pass performs two main checks:
+
+    1. **Live-in Value Verification**: Checks that all values that are live
+       into an offload region are valid for use in that region. This includes
+       checking that pointer-like and mappable types have appropriate data
+       clauses or device attributes.
+
+    2. **Symbol Use Verification**: Checks that all symbol references within
+       an offload region are valid for that region. This includes checking for
+       acc.routine declarations and acc.declare attributes.
+
+    The device_type option notes the target execution model:
+    - `none`, `nvidia`, `radeon`: Device execution (GPU offload)
+    - `host`, `multicore`: Host execution
+
+    When soft_check is enabled, the pass only emits debug messages for illegal
+    values/symbols instead of failing compilation. This is useful for
+    diagnostic purposes.
+  }];
+  let dependentDialects = ["mlir::acc::OpenACCDialect"];
+  let options = [
+    Option<"deviceType", "device-type", "mlir::acc::DeviceType",
+           "mlir::acc::DeviceType::None",
+           "Target device type for verification. Host/multicore uses host "
+           "region checking, all others use device region checking.",
+           [{::llvm::cl::values(
+              clEnumValN(mlir::acc::DeviceType::None, "none", "none"),
+              clEnumValN(mlir::acc::DeviceType::Host, "host", "host"),
+              clEnumValN(mlir::acc::DeviceType::Multicore, "multicore", "multicore"),
+              clEnumValN(mlir::acc::DeviceType::Nvidia, "nvidia", "nvidia"),
+              clEnumValN(mlir::acc::DeviceType::Radeon, "radeon", "radeon"))
+           }]>,
+    Option<"softCheck", "soft-check", "bool", "false",
+           "When true, illegal values are printed via LLVM_DEBUG instead of "
+           "failing compilation. Useful for diagnostic purposes.">
+  ];
+}
+
 #endif // MLIR_DIALECT_OPENACC_TRANSFORMS_PASSES
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/CMakeLists.txt b/mlir/lib/Dialect/OpenACC/Transforms/CMakeLists.txt
index 1e2f86964ac0d..20cfcccb9a42d 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/CMakeLists.txt
+++ b/mlir/lib/Dialect/OpenACC/Transforms/CMakeLists.txt
@@ -9,6 +9,7 @@ add_mlir_dialect_library(MLIROpenACCTransforms
   ACCSpecializeForHost.cpp
   LegalizeDataValues.cpp
   OffloadLiveInValueCanonicalization.cpp
+  OffloadTargetVerifier.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/OpenACC
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/OffloadTargetVerifier.cpp b/mlir/lib/Dialect/OpenACC/Transforms/OffloadTargetVerifier.cpp
new file mode 100644
index 0000000000000..93a91de29d385
--- /dev/null
+++ b/mlir/lib/Dialect/OpenACC/Transforms/OffloadTargetVerifier.cpp
@@ -0,0 +1,236 @@
+//===- OffloadTargetVerifier.cpp ------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This pass verifies that values and symbols used within offload regions are
+// legal for the target execution model.
+//
+// Overview:
+// ---------
+// Offload regions execute on a target device (e.g., GPU) where not all values
+// and symbols from the host context are accessible. This pass checks that
+// live-in values (values defined outside but used inside the region) and
+// symbol references are valid for device execution.
+//
+// The pass operates on any operation implementing `OffloadRegionOpInterface`,
+// which includes OpenACC compute constructs (`acc.parallel`, `acc.kernels`,
+// `acc.serial`) as well as GPU operations like `gpu.launch`.
+//
+// Verification:
+// -------------
+// For each offload region, the pass checks:
+//
+// 1. Live-in Values: Values flowing into the region must be valid for device
+//    use. This includes checking that data has been properly mapped via
+//    OpenACC data clauses (copyin, copyout, present, etc.) or is a scalar
+//    that can be passed by value.
+//
+// 2. Symbol References: Symbols referenced inside the region must be
+//    accessible on the device. This includes checking for proper `declare`
+//    attributes on globals or device-resident data attributes.
+//
+// Requirements:
+// -------------
+// 1. Target Region Identification: Operations representing offload regions
+//    must implement `acc::OffloadRegionOpInterface`.
+//
+// 2. OpenACCSupport Analysis: The pass relies on the `OpenACCSupport`
+//    analysis to determine value and symbol validity. This analysis provides
+//    dialect-specific hooks for checking legality through `isValidValueUse`
+//    and `isValidSymbolUse` methods. Custom dialect support can be registered
+//    by providing a derived `OpenACCSupport` analysis before running this
+//    pass.
+//
+// 3. Device Type: The `device_type` option specifies the target device.
+//    For `host` or `multicore` targets, verification of ACC compute
+//    constructs is not yet implemented.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Analysis/Liveness.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
+#include "mlir/IR/SymbolTable.h"
+#include "llvm/Support/Debug.h"
+
+namespace mlir {
+namespace acc {
+#define GEN_PASS_DEF_OFFLOADTARGETVERIFIER
+#include "mlir/Dialect/OpenACC/Transforms/Passes.h.inc"
+} // namespace acc
+} // namespace mlir
+
+#define DEBUG_TYPE "offload-target-verifier"
+
+using namespace mlir;
+
+namespace {
+
+class OffloadTargetVerifier
+    : public acc::impl::OffloadTargetVerifierBase<OffloadTargetVerifier> {
+public:
+  using OffloadTargetVerifierBase::OffloadTargetVerifierBase;
+
+  /// Returns true if the target device type corresponds to host execution.
+  bool isHostTarget() const {
+    return deviceType == acc::DeviceType::Host ||
+           deviceType == acc::DeviceType::Multicore;
+  }
+
+  /// Check live-in values for legality.
+  SmallVector<Value>
+  getIllegalLiveInValues(Region &region, Liveness &liveness,
+                         acc::OpenACCSupport &accSupport) const {
+    auto isInvalid = [&](Value val) -> bool {
+      return !accSupport.isValidValueUse(val, region);
+    };
+
+    SmallVector<Value> illegalValues(llvm::make_filter_range(
+        liveness.getLiveIn(&region.front()), isInvalid));
+
+    return illegalValues;
+  }
+
+  /// Check symbol uses for legality.
+  SmallVector<SymbolTable::SymbolUse>
+  getIllegalUsedSymbols(Region &region, acc::OpenACCSupport &accSupport) const {
+    auto symUses = SymbolTable::getSymbolUses(&region);
+
+    /...
[truncated]

``````````

</details>


https://github.com/llvm/llvm-project/pull/176467