[Mlir-commits] [mlir] 8dfec25 - [mlir][acc] Add OffloadTargetVerifier pass (#176467)

[llvmlistbot at llvm.org]

Author: Razvan Lupusoru
Date: 2026-01-20T17:17:08Z
New Revision: 8dfec259742648c12cca4470eedf583c4e21801f

URL: https://github.com/llvm/llvm-project/commit/8dfec259742648c12cca4470eedf583c4e21801f
DIFF: https://github.com/llvm/llvm-project/commit/8dfec259742648c12cca4470eedf583c4e21801f.diff

LOG: [mlir][acc] Add OffloadTargetVerifier pass (#176467)

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

Added: 
    flang/test/Transforms/OpenACC/offload-target-verifier.fir
    mlir/lib/Dialect/OpenACC/Transforms/OffloadTargetVerifier.cpp
    mlir/test/Dialect/OpenACC/offload-target-verifier.mlir

Modified: 
    flang/include/flang/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.h
    flang/lib/Optimizer/OpenACC/Analysis/FIROpenACCSupportAnalysis.cpp
    mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td
    mlir/lib/Dialect/OpenACC/Transforms/CMakeLists.txt
    mlir/lib/Dialect/OpenACC/Utils/OpenACCUtils.cpp

Removed: 
    


################################################################################
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..91a7c7d6489a4
--- /dev/null
+++ b/mlir/lib/Dialect/OpenACC/Transforms/OffloadTargetVerifier.cpp
@@ -0,0 +1,234 @@
+//===- 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);
+
+    // When there are no symbols used in the region, there are no illegal ones.
+    if (!symUses.has_value())
+      return {};
+
+    auto isInvalidSymbol = [&](const SymbolTable::SymbolUse &symUse) -> bool {
+      Operation *definingOp = nullptr;
+      return !accSupport.isValidSymbolUse(symUse.getUser(),
+                                          symUse.getSymbolRef(), &definingOp);
+    };
+
+    auto invalidSyms =
+        llvm::make_filter_range(symUses.value(), isInvalidSymbol);
+    SmallVector<SymbolTable::SymbolUse> invalidSymsList(invalidSyms);
+    return invalidSymsList;
+  }
+
+  /// Check if the region has illegal live-in values.
+  bool hasIllegalLiveInValues(Operation *regionOp,
+                              acc::OpenACCSupport &accSupport) const {
+    if (regionOp->getNumRegions() == 0)
+      return false;
+
+    Liveness liveness(regionOp);
+    SmallVector<Value> invalidValues =
+        getIllegalLiveInValues(regionOp->getRegion(0), liveness, accSupport);
+
+    bool hasIllegalValues = !invalidValues.empty();
+
+    if (hasIllegalValues) {
+      if (softCheck) {
+        // Emit warnings for each illegal value.
+        auto diag = regionOp->emitWarning("offload target verifier: ")
+                    << invalidValues.size() << " illegal live-in value(s)";
+        for (auto [idx, invalidValue] : llvm::enumerate(invalidValues)) {
+          diag.attachNote(invalidValue.getLoc()) << "value: " << invalidValue;
+        }
+      } else {
+        accSupport.emitNYI(regionOp->getLoc(),
+                           "offload target verifier failed due to " +
+                               Twine(invalidValues.size()) +
+                               " illegal live-in value(s)");
+      }
+    }
+
+    return hasIllegalValues;
+  }
+
+  /// Check if the region has illegal symbol uses.
+  bool hasIllegalSymbolUses(Operation *regionOp,
+                            acc::OpenACCSupport &accSupport) const {
+    if (regionOp->getNumRegions() == 0)
+      return false;
+
+    SmallVector<SymbolTable::SymbolUse> invalidSyms =
+        getIllegalUsedSymbols(regionOp->getRegion(0), accSupport);
+
+    bool hasIllegalSymbols = !invalidSyms.empty();
+
+    if (hasIllegalSymbols) {
+      auto getSymName = [&](SymbolTable::SymbolUse symUse) -> std::string {
+        return symUse.getSymbolRef().getLeafReference().str();
+      };
+      std::string invalidString =
+          llvm::join(llvm::map_range(invalidSyms, getSymName), ", ");
+
+      // Emit only warnings when softCheck is enabled.
+      if (softCheck)
+        regionOp->emitWarning("offload target verifier: illegal symbol(s): ")
+            << invalidString;
+      else
+        accSupport.emitNYI(regionOp->getLoc(),
+                           "offload target verifier failed due to illegal "
+                           "symbol(s): " +
+                               invalidString);
+    }
+
+    return hasIllegalSymbols;
+  }
+
+  void runOnOperation() override {
+    LLVM_DEBUG(llvm::dbgs() << "Enter OffloadTargetVerifier()\n");
+    func::FuncOp func = getOperation();
+
+    // Try to get cached parent analysis first, fall back to local analysis.
+    auto cachedAnalysis =
+        getCachedParentAnalysis<acc::OpenACCSupport>(func->getParentOp());
+    acc::OpenACCSupport &accSupport = cachedAnalysis
+                                          ? cachedAnalysis->get()
+                                          : getAnalysis<acc::OpenACCSupport>();
+
+    bool hasErrors = false;
+
+    func.walk([&](Operation *op) {
+      // Only process offload region operations.
+      if (!isa<acc::OffloadRegionOpInterface>(op))
+        return WalkResult::advance();
+
+      // TODO: Host/multicore verification for ACC compute constructs is not yet
+      // implemented.
+      if (isHostTarget() && isa<ACC_COMPUTE_CONSTRUCT_OPS>(op)) {
+        accSupport.emitNYI(op->getLoc(),
+                           "host/multicore verification for ACC compute "
+                           "constructs");
+        return WalkResult::advance();
+      }
+
+      // Check for illegal live-in values.
+      bool hasIllegalValues = hasIllegalLiveInValues(op, accSupport);
+      if (hasIllegalValues)
+        hasErrors = true;
+
+      // Check for illegal symbol uses.
+      bool hasIllegalSyms = hasIllegalSymbolUses(op, accSupport);
+      if (hasIllegalSyms)
+        hasErrors = true;
+
+      if (!hasIllegalValues && !hasIllegalSyms && softCheck)
+        op->emitRemark("offload target verifier: passed validity check");
+
+      return WalkResult::advance();
+    });
+
+    if (hasErrors && !softCheck)
+      signalPassFailure();
+
+    LLVM_DEBUG(llvm::dbgs() << "Exit OffloadTargetVerifier()\n");
+  }
+};
+
+} // namespace

diff  --git a/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtils.cpp b/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtils.cpp
index bd3dda48d44b4..26243b2ae84be 100644
--- a/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtils.cpp
+++ b/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtils.cpp
@@ -9,6 +9,7 @@
 #include "mlir/Dialect/OpenACC/OpenACCUtils.h"
 
 #include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/Dominance.h"
 #include "mlir/IR/SymbolTable.h"
 #include "mlir/Interfaces/FunctionInterfaces.h"
@@ -249,6 +250,12 @@ bool mlir::acc::isDeviceValue(mlir::Value val) {
 }
 
 bool mlir::acc::isValidValueUse(mlir::Value val, mlir::Region &region) {
+  // Types that can be passed by value are legal.
+  Type type = val.getType();
+  if (type.isIntOrIndexOrFloat() || isa<mlir::ComplexType>(type) ||
+      llvm::isa<mlir::VectorType>(type))
+    return true;
+
   // If this is produced by an ACC data entry operation, it is valid.
   if (isa_and_nonnull<ACC_DATA_ENTRY_OPS>(val.getDefiningOp()))
     return true;

diff  --git a/mlir/test/Dialect/OpenACC/offload-target-verifier.mlir b/mlir/test/Dialect/OpenACC/offload-target-verifier.mlir
new file mode 100644
index 0000000000000..d4380de8e2a31
--- /dev/null
+++ b/mlir/test/Dialect/OpenACC/offload-target-verifier.mlir
@@ -0,0 +1,231 @@
+// RUN: mlir-opt %s --pass-pipeline="builtin.module(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 = memref.alloca() : memref<i32>
+  %livein = memref.load %alloca[] : memref<i32>
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %accalloca = memref.alloca() : memref<i32>
+    memref.store %livein, %accalloca[] : memref<i32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test memref live-in without data clause - should fail
+func.func @test_memref_f32() {
+  // expected-note @below {{value}}
+  %livein = memref.alloca() : memref<f32>
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  acc.serial {
+    %load = memref.load %livein[] : memref<f32>
+    %accalloca = memref.alloca() : memref<f32>
+    memref.store %load, %accalloca[] : memref<f32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test memref with copyin data clause - should pass
+func.func @test_memref_f32_copyin() {
+  %alloca = memref.alloca() : memref<f32>
+  %livein = acc.copyin varPtr(%alloca : memref<f32>) -> memref<f32>
+  // expected-remark @below {{passed validity check}}
+  acc.serial dataOperands(%livein : memref<f32>) {
+    %load = memref.load %livein[] : memref<f32>
+    %accalloca = memref.alloca() : memref<f32>
+    memref.store %load, %accalloca[] : memref<f32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test memref with private clause - should pass (privatized values are not live-in)
+acc.private.recipe @privatization_memref_f32 : memref<f32> init {
+^bb0(%arg0: memref<f32>):
+  %0 = memref.alloca() : memref<f32>
+  acc.yield %0 : memref<f32>
+}
+
+func.func @test_memref_f32_private() {
+  %livein = memref.alloca() : memref<f32>
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %private = acc.private varPtr(%livein : memref<f32>) recipe(@privatization_memref_f32) -> memref<f32>
+    %load = memref.load %private[] : memref<f32>
+    %accalloca = memref.alloca() : memref<f32>
+    memref.store %load, %accalloca[] : memref<f32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test llvm.ptr live-in without data clause - should fail
+func.func @test_llvmptr_f64() {
+  %c1 = arith.constant 1 : i64
+  // expected-note @below {{value}}
+  %alloca = llvm.alloca %c1 x f64 : (i64) -> !llvm.ptr
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  acc.serial {
+    %c1_inner = arith.constant 1 : i64
+    %load = llvm.load %alloca : !llvm.ptr -> f64
+    %accalloca = llvm.alloca %c1_inner x f64 : (i64) -> !llvm.ptr
+    llvm.store %load, %accalloca : f64, !llvm.ptr
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test global symbol without declare attribute - should fail
+memref.global @global_array : memref<10xf32> = uninitialized
+
+func.func @test_global_symbol_no_declare() {
+  // expected-warning @below {{illegal symbol(s): global_array}}
+  acc.serial {
+    %livein = memref.get_global @global_array : memref<10xf32>
+    %c0 = arith.constant 0 : index
+    %loaded = memref.load %livein[%c0] : memref<10xf32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test memref with GPU address space (device data) - should pass
+func.func @test_memref_gpu_address_space() {
+  %alloca = memref.alloca() : memref<f32, #gpu.address_space<global>>
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %load = memref.load %alloca[] : memref<f32, #gpu.address_space<global>>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test global symbol with acc.declare attribute - should pass
+memref.global @global_array_declared : memref<10xf32> = dense<0.0> {acc.declare = #acc.declare<dataClause = acc_create>}
+
+func.func @test_global_symbol_with_declare() {
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %livein = memref.get_global @global_array_declared : memref<10xf32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test gpu.launch region (another OffloadRegionOpInterface)
+func.func @test_gpu_launch() {
+  %c1 = arith.constant 1 : index
+  // expected-note @below {{value}}
+  %alloca = memref.alloca() : memref<f32>
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  gpu.launch blocks(%bx, %by, %bz) in (%grid_x = %c1, %grid_y = %c1, %grid_z = %c1)
+             threads(%tx, %ty, %tz) in (%block_x = %c1, %block_y = %c1, %block_z = %c1) {
+    %load = memref.load %alloca[] : memref<f32>
+    gpu.terminator
+  }
+  return
+}
+
+// -----
+
+// Test acc.parallel region
+func.func @test_acc_parallel() {
+  // expected-note @below {{value}}
+  %alloca = memref.alloca() : memref<f32>
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  acc.parallel {
+    %load = memref.load %alloca[] : memref<f32>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test acc.kernels region
+func.func @test_acc_kernels() {
+  // expected-note @below {{value}}
+  %alloca = memref.alloca() : memref<f32>
+  // expected-warning @below {{1 illegal live-in value(s)}}
+  acc.kernels {
+    %load = memref.load %alloca[] : memref<f32>
+    acc.terminator
+  }
+  return
+}
+
+// -----
+
+// Test device global (memref.global with GPU address space) - should pass
+memref.global @device_global : memref<10xf32, #gpu.address_space<global>> = uninitialized
+
+func.func @test_device_global() {
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %livein = memref.get_global @device_global : memref<10xf32, #gpu.address_space<global>>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test complex scalar (complex types can be passed by value)
+func.func @test_complex_scalar() {
+  %alloca = memref.alloca() : memref<complex<f32>>
+  %livein = memref.load %alloca[] : memref<complex<f32>>
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %accalloca = memref.alloca() : memref<complex<f32>>
+    memref.store %livein, %accalloca[] : memref<complex<f32>>
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test index type scalar
+func.func @test_index_scalar() {
+  %c10 = arith.constant 10 : index
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %c1 = arith.constant 1 : index
+    %sum = arith.addi %c10, %c1 : index
+    acc.yield
+  }
+  return
+}
+
+// -----
+
+// Test f64 scalar
+func.func @test_f64_scalar() {
+  %alloca = memref.alloca() : memref<f64>
+  %livein = memref.load %alloca[] : memref<f64>
+  // expected-remark @below {{passed validity check}}
+  acc.serial {
+    %accalloca = memref.alloca() : memref<f64>
+    memref.store %livein, %accalloca[] : memref<f64>
+    acc.yield
+  }
+  return
+}