diff --git a/mlir-tensorrt/compiler/lib/Conversion/StablehloToScf/StablehloToScf.cpp b/mlir-tensorrt/compiler/lib/Conversion/StablehloToScf/StablehloToScf.cpp index c24c7ea69..99eb495da 100644 --- a/mlir-tensorrt/compiler/lib/Conversion/StablehloToScf/StablehloToScf.cpp +++ b/mlir-tensorrt/compiler/lib/Conversion/StablehloToScf/StablehloToScf.cpp @@ -7,14 +7,14 @@ // https://github.com/openxla/xla/blob/main/LICENSE for the license // information. // -// Changes are copyright (c) 2024, NVIDIA CORPORATION. All rights reserved. +// Changes are copyright (c) 2024-2025, NVIDIA CORPORATION. All rights reserved. // //===----------------------------------------------------------------------===// /// /// Implementation of a pass to convert stablehlo control flow ops to scf ops. /// //===----------------------------------------------------------------------===// - +#include "mlir-tensorrt/Conversion/Passes.h" #include "mlir-tensorrt/Transforms/Transforms.h" #include "mlir/Analysis/SliceAnalysis.h" #include "mlir/Dialect/Arith/IR/Arith.h" @@ -538,4 +538,4 @@ struct StablehloToScfPass } } }; -} // namespace \ No newline at end of file +} // namespace diff --git a/mlir-tensorrt/compiler/lib/Dialect/Plan/Transforms/AssignMemorySpaces.cpp b/mlir-tensorrt/compiler/lib/Dialect/Plan/Transforms/AssignMemorySpaces.cpp index 394c1de6e..80cc709f8 100644 --- a/mlir-tensorrt/compiler/lib/Dialect/Plan/Transforms/AssignMemorySpaces.cpp +++ b/mlir-tensorrt/compiler/lib/Dialect/Plan/Transforms/AssignMemorySpaces.cpp @@ -297,7 +297,7 @@ convertFuncRegionTypes(RewriterBase &rewriter, FunctionOpInterface funcOp, /// change are RankedTensorTypes where the encoding has been updated. Therefore, /// we only insert `tensor.cast` operations to cast the values back to their /// original types. -struct LogicalResult convertFuncUsers(RewriterBase &rewriter, +static LogicalResult convertFuncUsers(RewriterBase &rewriter, FunctionOpInterface func, const SymbolUserMap &userMap) { OpBuilder::InsertionGuard g(rewriter); diff --git a/mlir-tensorrt/compiler/lib/Dialect/Plan/Transforms/PostClusteringValidation.cpp b/mlir-tensorrt/compiler/lib/Dialect/Plan/Transforms/PostClusteringValidation.cpp index cc488a9d1..27b52b776 100644 --- a/mlir-tensorrt/compiler/lib/Dialect/Plan/Transforms/PostClusteringValidation.cpp +++ b/mlir-tensorrt/compiler/lib/Dialect/Plan/Transforms/PostClusteringValidation.cpp @@ -1,6 +1,6 @@ //===- PostClusteringValidation.cpp ---------------------------------------===// // -// SPDX-FileCopyrightText: Copyright 2024 NVIDIA CORPORATION & AFFILIATES. +// SPDX-FileCopyrightText: Copyright 2024-2025 NVIDIA CORPORATION & AFFILIATES. // All rights reserved. // SPDX-License-Identifier: Apache-2.0 // @@ -27,6 +27,7 @@ #include "mlir-tensorrt-dialect/TensorRT/IR/TensorRTDialect.h" #include "mlir-tensorrt/Dialect/CUDA/IR/CUDADialect.h" #include "mlir-tensorrt/Dialect/Plan/IR/Plan.h" +#include "mlir-tensorrt/Dialect/Plan/Transforms/Passes.h" #include "mlir-tensorrt/Dialect/TensorRTRuntime/IR/TensorRTRuntime.h" #include "mlir-tensorrt/Transforms/Passes.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" diff --git a/mlir-tensorrt/compiler/lib/Dialect/StablehloExt/IR/StablehloTensorKindOpInterfaceImpl.cpp b/mlir-tensorrt/compiler/lib/Dialect/StablehloExt/IR/StablehloTensorKindOpInterfaceImpl.cpp index 69d2bec55..6130a8382 100644 --- a/mlir-tensorrt/compiler/lib/Dialect/StablehloExt/IR/StablehloTensorKindOpInterfaceImpl.cpp +++ b/mlir-tensorrt/compiler/lib/Dialect/StablehloExt/IR/StablehloTensorKindOpInterfaceImpl.cpp @@ -39,7 +39,7 @@ constexpr std::pair OpToHostParametersOffsetAndSize() { // Simple macro for creating the appearance of a table. #define CASE(OpType, start, size) \ if constexpr (std::is_same_v) \ - return {start, size} + return { start, size } CASE(stablehlo::DynamicIotaOp, 0, 1); // Note that `stablehlo.dynamic_slice` and `stablehlo.dynamic_update_slice` diff --git a/mlir-tensorrt/compiler/lib/Transforms/MemRefCastElimination/MemRefCastElimination.cpp b/mlir-tensorrt/compiler/lib/Transforms/MemRefCastElimination/MemRefCastElimination.cpp index be71b15de..c9fe00a25 100644 --- a/mlir-tensorrt/compiler/lib/Transforms/MemRefCastElimination/MemRefCastElimination.cpp +++ b/mlir-tensorrt/compiler/lib/Transforms/MemRefCastElimination/MemRefCastElimination.cpp @@ -1,6 +1,6 @@ //===- MemRefCastElimination.cpp ------------------------------------------===// // -// SPDX-FileCopyrightText: Copyright 2024 NVIDIA CORPORATION & AFFILIATES. +// SPDX-FileCopyrightText: Copyright 2024-2025 NVIDIA CORPORATION & AFFILIATES. // All rights reserved. // SPDX-License-Identifier: Apache-2.0 // @@ -41,7 +41,7 @@ using namespace mlir; /// "then" and "else" terminator operands are produced by "compatible" cast /// operations that can be moved to act on the result of the if. The types of /// the new `scf.if` operation are returned as well. -FailureOr, SmallVector>> +static FailureOr, SmallVector>> isCastEliminationCandidate(scf::IfOp op) { SmallVector resultIndices; SmallVector newResultTypes; @@ -139,4 +139,4 @@ class MemRefCastEliminationPass } } }; -} // namespace \ No newline at end of file +} // namespace diff --git a/mlir-tensorrt/compiler/lib/Transforms/UnrollForLoops/UnrollForLoops.cpp b/mlir-tensorrt/compiler/lib/Transforms/UnrollForLoops/UnrollForLoops.cpp index f95ae0b63..019e35745 100644 --- a/mlir-tensorrt/compiler/lib/Transforms/UnrollForLoops/UnrollForLoops.cpp +++ b/mlir-tensorrt/compiler/lib/Transforms/UnrollForLoops/UnrollForLoops.cpp @@ -87,9 +87,9 @@ static std::string printWithoutRegions(Operation *op) { /// Unrolls `op` if its trip count is static and less than `unrollThreshold`. /// Returns `success()` if the loop is unrolled or ignored, `failure()` if the /// transformation fails. -LogicalResult unrollForLoopWithStaticTripCount(IRRewriter &rewriter, - scf::ForOp op, - uint64_t unrollThreshold) { +static LogicalResult +unrollForLoopWithStaticTripCount(IRRewriter &rewriter, scf::ForOp op, + uint64_t unrollThreshold) { std::optional tripCount = getConstantTripCount(op); if (!tripCount) return success(); diff --git a/mlir-tensorrt/compiler/test/lib/Analysis/CMakeLists.txt b/mlir-tensorrt/compiler/test/lib/Analysis/CMakeLists.txt index 56c4fad70..b5b069dce 100644 --- a/mlir-tensorrt/compiler/test/lib/Analysis/CMakeLists.txt +++ b/mlir-tensorrt/compiler/test/lib/Analysis/CMakeLists.txt @@ -1,9 +1,7 @@ -add_mlir_tensorrt_test_library(MLIRTensorRTTestTensorKindAnalysis - TestTensorKindAnalysis.cpp +add_mlir_tensorrt_test_library(MLIRTensorRTCompilerAnalysisTestPasses TestBoundsAnalysis.cpp LINK_LIBS PUBLIC - MLIRTensorRTAnalysis MLIRTensorRTPlanAnalysis MLIR_LIBS PUBLIC diff --git a/mlir-tensorrt/compiler/test/lib/Analysis/TestBoundsAnalysis.cpp b/mlir-tensorrt/compiler/test/lib/Analysis/TestBoundsAnalysis.cpp index 1f07b668a..f9f636593 100644 --- a/mlir-tensorrt/compiler/test/lib/Analysis/TestBoundsAnalysis.cpp +++ b/mlir-tensorrt/compiler/test/lib/Analysis/TestBoundsAnalysis.cpp @@ -1,6 +1,6 @@ //===- TestBoundsAnalysis.cpp ---------------------------------------------===// // -// Copyright (c) 2024, NVIDIA CORPORATION. All rights reserved. +// Copyright (c) 2024-2025, NVIDIA CORPORATION. All rights reserved. // //===----------------------------------------------------------------------===// /// @@ -20,6 +20,10 @@ using namespace mlir; using namespace mlir::dataflow; +namespace mlir { +void registerTestBoundsAnalysisPass(); +} + /// Print out the lattice information for the given value `v`. template static void printLatticeInfo(llvm::raw_ostream &os, Value v, @@ -135,9 +139,7 @@ struct TestTensorValueBoundsAnalysisPass }; } // namespace -namespace mlir { -void registerTestBoundsAnalysisPass() { +void mlir::registerTestBoundsAnalysisPass() { PassRegistration(); PassRegistration(); } -} // namespace mlir diff --git a/mlir-tensorrt/compiler/tools/mlir-tensorrt-opt/mlir-tensorrt-opt.cpp b/mlir-tensorrt/compiler/tools/mlir-tensorrt-opt/mlir-tensorrt-opt.cpp index 7220cbadb..2ce15e3c4 100644 --- a/mlir-tensorrt/compiler/tools/mlir-tensorrt-opt/mlir-tensorrt-opt.cpp +++ b/mlir-tensorrt/compiler/tools/mlir-tensorrt-opt/mlir-tensorrt-opt.cpp @@ -31,8 +31,10 @@ using namespace llvm; #ifdef MLIR_TRT_ENABLE_TESTING namespace mlir { +namespace tensorrt { void registerTestTensorKindAnalysisPass(); void registerTestTensorRTShapeInferencePass(); +} // namespace tensorrt #ifdef MLIR_TRT_ENABLE_HLO void registerTestBoundsAnalysisPass(); @@ -40,8 +42,8 @@ void registerTestBoundsAnalysisPass(); } // namespace mlir static void registerTestPasses() { - mlir::registerTestTensorKindAnalysisPass(); - mlir::registerTestTensorRTShapeInferencePass(); + mlir::tensorrt::registerTestTensorKindAnalysisPass(); + mlir::tensorrt::registerTestTensorRTShapeInferencePass(); IF_MLIR_TRT_ENABLE_HLO({ mlir::registerTestBoundsAnalysisPass(); }); } #endif // MLIR_TRT_ENABLE_TESTING diff --git a/mlir-tensorrt/executor/lib/CAPI/Common/Common.cpp b/mlir-tensorrt/executor/lib/CAPI/Common/Common.cpp index 855f4d5fc..1cfdc3e04 100644 --- a/mlir-tensorrt/executor/lib/CAPI/Common/Common.cpp +++ b/mlir-tensorrt/executor/lib/CAPI/Common/Common.cpp @@ -1,6 +1,6 @@ -//===- Common.cpp -----0---------------------------------------------------===// +//===- Common.cpp ---------------------------------------------------------===// // -// SPDX-FileCopyrightText: Copyright 2024 NVIDIA CORPORATION & AFFILIATES. +// SPDX-FileCopyrightText: Copyright 2024-2025 NVIDIA CORPORATION & AFFILIATES. // All rights reserved. // SPDX-License-Identifier: Apache-2.0 // @@ -398,7 +398,7 @@ mtrtFunctionSignatureGetNumOutputArgs(MTRT_FunctionSignature signature, return mtrtStatusGetOk(); } -MTRT_Status getTypeHelper(TypeUnionView typeUnionView, MTRT_Type *type) { +static MTRT_Status getTypeHelper(TypeUnionView typeUnionView, MTRT_Type *type) { // Allocate the TypeUnion object, populate it by moving in the // concrete object, and release it to be owned by the CAPI object. auto typeUnion = std::make_unique(); @@ -466,8 +466,8 @@ mtrtFunctionSignatureGetShapeFuncName(MTRT_FunctionSignature signature, return mtrtStatusGetOk(); } -MTRT_Status getBoundsHelper(BoundsUnionView boundsUnionView, - MTRT_Bounds *bounds) { +static MTRT_Status getBoundsHelper(BoundsUnionView boundsUnionView, + MTRT_Bounds *bounds) { // Allocate the BoundsUnion object, populate it by moving in the // concrete object, and release it to be owned by the CAPI object. auto boundsUnion = std::make_unique(); diff --git a/mlir-tensorrt/executor/lib/Runtime/API/API.cpp b/mlir-tensorrt/executor/lib/Runtime/API/API.cpp index c4bbc7bee..c346be37d 100644 --- a/mlir-tensorrt/executor/lib/Runtime/API/API.cpp +++ b/mlir-tensorrt/executor/lib/Runtime/API/API.cpp @@ -1,6 +1,6 @@ //===- Executable.cpp ------ ----------------------------------------------===// // -// SPDX-FileCopyrightText: Copyright 2024 NVIDIA CORPORATION & AFFILIATES. +// SPDX-FileCopyrightText: Copyright 2024-2025 NVIDIA CORPORATION & AFFILIATES. // All rights reserved. // SPDX-License-Identifier: Apache-2.0 // @@ -549,7 +549,7 @@ struct CompletionToken // completes. userData is a heap-allocated pointer to // std::shared_ptr>. Implementation posts a small job to IO // thread pool and returns. -void cuda_event_host_callback(void *userData) { +static void cuda_event_host_callback(void *userData) { // // userData is pointer-to-heap-allocated Ref. Ref *tokenPtr = static_cast *>(userData); diff --git a/mlir-tensorrt/executor/lib/Runtime/Backend/C/CUDAModule.cpp b/mlir-tensorrt/executor/lib/Runtime/Backend/C/CUDAModule.cpp index 61e2554bb..3980c83b7 100644 --- a/mlir-tensorrt/executor/lib/Runtime/Backend/C/CUDAModule.cpp +++ b/mlir-tensorrt/executor/lib/Runtime/Backend/C/CUDAModule.cpp @@ -65,7 +65,7 @@ int32_t mtrt_cuda_set_active_device(int32_t device) { return device; } -int32_t mtrt_get_device(int32_t device) { return device; } +int32_t mtrt_cuda_get_device(int32_t device) { return device; } CUstream mtrt_cuda_stream_create() { CUstream stream; @@ -108,12 +108,6 @@ void mtrt_cuda_free(CUstream stream, void *ptr, int8_t isHostPinned, HANDLE_CUDART_ERROR(cudaFreeAsync(ptr, stream), ); } -CUfunction mtrt_cumodule_load_func(CUmodule module, const char *funcName) { - CUfunction func; - cuModuleGetFunction(&func, module, funcName); - return func; -} - static StatusOr getDeviceArch(int32_t deviceNumber) { CUdevice deviceID; RETURN_ERROR_WITH_MSG_IF_CUDADRV_ERROR(cuDeviceGet(&deviceID, deviceNumber), diff --git a/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/CUDA/CUDAModule.cpp b/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/CUDA/CUDAModule.cpp index e475f269e..37724f267 100644 --- a/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/CUDA/CUDAModule.cpp +++ b/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/CUDA/CUDAModule.cpp @@ -530,6 +530,11 @@ registerCudaMemoryManagementOps(sol::state_view &lua, }; } +#if defined(__GNUC__) || defined(__clang__) +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wmissing-prototypes" +#endif + namespace mtrt { void registerLuaCudaRuntimeExtension() { registerLuaRuntimeExtension( diff --git a/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/Core/CoreModule.cpp b/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/Core/CoreModule.cpp index 58a7fbd9c..eab328120 100644 --- a/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/Core/CoreModule.cpp +++ b/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/Core/CoreModule.cpp @@ -1,6 +1,6 @@ //===- CoreModule.cpp -----------------------------------------------------===// // -// SPDX-FileCopyrightText: Copyright 2024 NVIDIA CORPORATION & AFFILIATES. +// SPDX-FileCopyrightText: Copyright 2024-2025 NVIDIA CORPORATION & AFFILIATES. // All rights reserved. // SPDX-License-Identifier: Apache-2.0 // @@ -43,6 +43,10 @@ using namespace mtrt; using namespace mtrt; +namespace mtrt { +void registerLuaCoreRuntimeExtension(); +} + //===----------------------------------------------------------------------===// // Templated helpers //===----------------------------------------------------------------------===// diff --git a/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/CuBLAS/CuBLASModule.cpp b/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/CuBLAS/CuBLASModule.cpp index 2c0923e0c..aec0c61a3 100644 --- a/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/CuBLAS/CuBLASModule.cpp +++ b/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/CuBLAS/CuBLASModule.cpp @@ -578,6 +578,11 @@ static void registerExecutorCuBLASModuleLuaRuntimeMethods( }; } +#if defined(__GNUC__) || defined(__clang__) +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wmissing-prototypes" +#endif + namespace mtrt { void registerLuaCublasRuntimeExtension() { registerLuaRuntimeExtension( diff --git a/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/TensorRT/TensorRTModule.cpp b/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/TensorRT/TensorRTModule.cpp index b7c91acb8..6b4a8683d 100644 --- a/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/TensorRT/TensorRTModule.cpp +++ b/mlir-tensorrt/executor/lib/Runtime/Backend/Lua/Modules/TensorRT/TensorRTModule.cpp @@ -792,6 +792,11 @@ static void registerExecutorTensorRTModuleLuaRuntimeMethods( }; } +#if defined(__GNUC__) || defined(__clang__) +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wmissing-prototypes" +#endif + namespace mtrt { void registerLuaTensorRTRuntimeExtension() { registerLuaRuntimeExtension( diff --git a/mlir-tensorrt/executor/lib/Support/DeviceInfo.cpp b/mlir-tensorrt/executor/lib/Support/DeviceInfo.cpp index d2bb38d0b..2c8711619 100644 --- a/mlir-tensorrt/executor/lib/Support/DeviceInfo.cpp +++ b/mlir-tensorrt/executor/lib/Support/DeviceInfo.cpp @@ -35,7 +35,7 @@ static Status makeCudaStringError(cudaError_t errCode, } #endif // MLIR_TRT_ENABLE_CUDA -static StatusOr +[[maybe_unused]] static StatusOr getDeviceInformationFromHostImpl(int cudaDeviceOridinal) { #ifdef MLIR_TRT_ENABLE_CUDA cudaDeviceProp properties; diff --git a/mlir-tensorrt/executor/lib/Target/Lua/TranslateToRuntimeExecutable.cpp b/mlir-tensorrt/executor/lib/Target/Lua/TranslateToRuntimeExecutable.cpp index bbdefc90f..24f52fc1a 100644 --- a/mlir-tensorrt/executor/lib/Target/Lua/TranslateToRuntimeExecutable.cpp +++ b/mlir-tensorrt/executor/lib/Target/Lua/TranslateToRuntimeExecutable.cpp @@ -1,6 +1,6 @@ //===- TranslateToRuntimeExecutable.cpp -----------------------------------===// // -// SPDX-FileCopyrightText: Copyright 2024 NVIDIA CORPORATION & AFFILIATES. +// SPDX-FileCopyrightText: Copyright 2024-2025 NVIDIA CORPORATION & AFFILIATES. // All rights reserved. // SPDX-License-Identifier: Apache-2.0 // @@ -479,7 +479,7 @@ static mtrt::flat::FunctionSignatureT generateSignature() { return signature; } -LogicalResult +static LogicalResult translateBoundsIfPresent(FunctionOpInterface func, unsigned argIndex, mtrt::flat::FunctionSignatureT &signature, bool isInput) { diff --git a/mlir-tensorrt/executor/test/lib/BufferizationTestPass.cpp b/mlir-tensorrt/executor/test/lib/BufferizationTestPass.cpp index 9e8a9e50a..d11603b13 100644 --- a/mlir-tensorrt/executor/test/lib/BufferizationTestPass.cpp +++ b/mlir-tensorrt/executor/test/lib/BufferizationTestPass.cpp @@ -1,6 +1,6 @@ //===- BufferizationTestPass.cpp ------------------------------------------===// // -// Copyright (c) 2024, NVIDIA CORPORATION. All rights reserved. +// Copyright (c) 2024-2025, NVIDIA CORPORATION. All rights reserved. // //===----------------------------------------------------------------------===// /// @@ -22,6 +22,10 @@ using namespace mlir; using namespace mlir::executor; +namespace mlir::executor { +void registerTestExecutorBufferizePass(); +} + namespace { class ExecutorBufferizationTestPass : public PassWrapper(); PassPipelineRegistration<> executorBufferizationPipeline( @@ -71,4 +74,3 @@ void registerTestExecutorBufferizePass() { pm.addPass(createCanonicalizerPass()); }); } -} // namespace mlir::executor diff --git a/mlir-tensorrt/executor/test/lib/Transforms/Clustering/TestClustering.cpp b/mlir-tensorrt/executor/test/lib/Transforms/Clustering/TestClustering.cpp index 0cad5a81a..fc9c84ad7 100644 --- a/mlir-tensorrt/executor/test/lib/Transforms/Clustering/TestClustering.cpp +++ b/mlir-tensorrt/executor/test/lib/Transforms/Clustering/TestClustering.cpp @@ -1,6 +1,6 @@ //===- TestClustering.cpp ------------------------------------------------===// // -// Copyright (c) 2023, NVIDIA CORPORATION. All rights reserved. +// Copyright (c) 2023-2025, NVIDIA CORPORATION. All rights reserved. // //===----------------------------------------------------------------------===// /// @@ -23,8 +23,12 @@ using namespace mlir; -scf::ExecuteRegionOp createScfRegionOpFromCluster(const Cluster &cluster, - RewriterBase &rewriter) { +namespace mlir::executor { +void registerTestClusteringTransformPass(); +} + +static scf::ExecuteRegionOp +createScfRegionOpFromCluster(const Cluster &cluster, RewriterBase &rewriter) { return cast(mlir::createRegionOpFromCluster( cluster, rewriter, [](OpBuilder &b, Location loc, TypeRange types, Attribute target) { @@ -215,8 +219,6 @@ class TestClusteringPass }; } // namespace -namespace mlir::executor { -void registerTestClusteringTransformPass() { +void executor::registerTestClusteringTransformPass() { PassRegistration(); } -} // namespace mlir::executor \ No newline at end of file diff --git a/mlir-tensorrt/tensorrt/lib/Target/TensorRTEncodingOpInterface/NetworkEncoder.cpp b/mlir-tensorrt/tensorrt/lib/Target/TensorRTEncodingOpInterface/NetworkEncoder.cpp index 58afd004f..e2d5e6a74 100644 --- a/mlir-tensorrt/tensorrt/lib/Target/TensorRTEncodingOpInterface/NetworkEncoder.cpp +++ b/mlir-tensorrt/tensorrt/lib/Target/TensorRTEncodingOpInterface/NetworkEncoder.cpp @@ -376,7 +376,7 @@ nvinfer1::ILayer *NvInferNetworkEncoder::addDequantizeLayer( #endif } -nvinfer1::IFillLayer *populateFillLayerParameters( +static nvinfer1::IFillLayer *populateFillLayerParameters( nvinfer1::IFillLayer *layer, const nvinfer1::Dims &staticShape, nvinfer1::ITensor *dynamicShape, std::optional alpha, std::optional beta, nvinfer1::ITensor *dynamicAlpha, diff --git a/mlir-tensorrt/tensorrt/lib/TensorRT/IR/TensorRT.cpp b/mlir-tensorrt/tensorrt/lib/TensorRT/IR/TensorRT.cpp index c5c0782d8..ac934dbcb 100644 --- a/mlir-tensorrt/tensorrt/lib/TensorRT/IR/TensorRT.cpp +++ b/mlir-tensorrt/tensorrt/lib/TensorRT/IR/TensorRT.cpp @@ -2011,7 +2011,8 @@ void tensorrt::IfOp::build( /// https://docs.nvidia.com/deeplearning/tensorrt/operators/index.html#layers-flow-control-constructs /// However, it is missing some checks on convolution/activation/fill/unary ops /// and therefore may give false positives. -bool isOperationSupportedInControlFlowBranchRegion(TensorRTOpInterface op) { +static bool +isOperationSupportedInControlFlowBranchRegion(TensorRTOpInterface op) { return !isa( op); diff --git a/mlir-tensorrt/tensorrt/test/lib/Target/TestPlugins.cpp b/mlir-tensorrt/tensorrt/test/lib/Target/TestPlugins.cpp index 58bab67d6..7465fc02a 100644 --- a/mlir-tensorrt/tensorrt/test/lib/Target/TestPlugins.cpp +++ b/mlir-tensorrt/tensorrt/test/lib/Target/TestPlugins.cpp @@ -601,6 +601,11 @@ REGISTER_TENSORRT_PLUGIN(TestQuickPluginCreator); REGISTER_TENSORRT_PLUGIN(TestQuickShapePluginCreator); #endif // MLIR_TRT_COMPILE_TIME_TENSORRT_VERSION_GTE(10, 9, 0) +#if defined(__GNUC__) || defined(__clang__) +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wmissing-prototypes" +#endif + extern "C" { /// Provide an exported C-style function for creating the plugin creator. TENSORRTTESTPLUGINS_EXPORT nvinfer1::IPluginCreatorInterface * diff --git a/mlir-tensorrt/tensorrt/test/lib/Target/TestV2Plugins.cpp b/mlir-tensorrt/tensorrt/test/lib/Target/TestV2Plugins.cpp index 4413a5a71..488829063 100644 --- a/mlir-tensorrt/tensorrt/test/lib/Target/TestV2Plugins.cpp +++ b/mlir-tensorrt/tensorrt/test/lib/Target/TestV2Plugins.cpp @@ -327,6 +327,11 @@ class TestV2InferShapePluginCreator : public nvinfer1::IPluginCreator { REGISTER_TENSORRT_PLUGIN(TestV2Plugin1Creator); REGISTER_TENSORRT_PLUGIN(TestV2InferShapePluginCreator); +#if defined(__GNUC__) || defined(__clang__) +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wmissing-prototypes" +#endif + extern "C" { /// Provide an exported C-style function for creating the plugin creator. TENSORRTTESTPLUGINS_EXPORT diff --git a/mlir-tensorrt/tensorrt/test/lib/TensorRT/CMakeLists.txt b/mlir-tensorrt/tensorrt/test/lib/TensorRT/CMakeLists.txt index 5ff11d1ab..dbb0f57ae 100644 --- a/mlir-tensorrt/tensorrt/test/lib/TensorRT/CMakeLists.txt +++ b/mlir-tensorrt/tensorrt/test/lib/TensorRT/CMakeLists.txt @@ -1,5 +1,6 @@ add_mlir_tensorrt_dialect_test_library(MLIRTensorRTTestTypeInferencePass TestTypeInferencePass.cpp + TestTensorKindAnalysis.cpp MLIR_LIBS PUBLIC MLIRTensorDialect @@ -9,4 +10,5 @@ add_mlir_tensorrt_dialect_test_library(MLIRTensorRTTestTypeInferencePass LINK_LIBS PUBLIC MLIRTensorRTTensorRTDialect + MLIRTensorRTAnalysis ) diff --git a/mlir-tensorrt/compiler/test/lib/Analysis/TestTensorKindAnalysis.cpp b/mlir-tensorrt/tensorrt/test/lib/TensorRT/TestTensorKindAnalysis.cpp similarity index 81% rename from mlir-tensorrt/compiler/test/lib/Analysis/TestTensorKindAnalysis.cpp rename to mlir-tensorrt/tensorrt/test/lib/TensorRT/TestTensorKindAnalysis.cpp index 3b46cba9d..4c9d42e40 100644 --- a/mlir-tensorrt/compiler/test/lib/Analysis/TestTensorKindAnalysis.cpp +++ b/mlir-tensorrt/tensorrt/test/lib/TensorRT/TestTensorKindAnalysis.cpp @@ -1,10 +1,25 @@ //===- TestTensorKindAnalysis.cpp -----------------------------------------===// // -// Copyright (c) 2023-2024, NVIDIA CORPORATION. All rights reserved. +// SPDX-FileCopyrightText: Copyright 2023-2025 NVIDIA CORPORATION & AFFILIATES. +// All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. // //===----------------------------------------------------------------------===// /// -/// Add tensor kind analysis test pass. +/// This file contains the implementation of the `test-tensor-kind-analysis` +/// pass. /// //===----------------------------------------------------------------------===// #include "mlir-tensorrt-common/Interfaces/TensorKindOpInterface.h" @@ -21,6 +36,10 @@ using namespace mlir; +namespace mlir::tensorrt { +void registerTestTensorKindAnalysisPass(); +} + /// Print out the lattice information for the given value `v`. static void printLatticeInfo(llvm::raw_ostream &os, Value v, DataFlowSolver &solver) { @@ -111,8 +130,6 @@ struct TestTensorKindPass }; } // namespace -namespace mlir { -void registerTestTensorKindAnalysisPass() { +void tensorrt::registerTestTensorKindAnalysisPass() { PassRegistration(); } -} // namespace mlir diff --git a/mlir-tensorrt/tensorrt/test/lib/TensorRT/TestTypeInferencePass.cpp b/mlir-tensorrt/tensorrt/test/lib/TensorRT/TestTypeInferencePass.cpp index d88b839e7..89aa6a114 100644 --- a/mlir-tensorrt/tensorrt/test/lib/TensorRT/TestTypeInferencePass.cpp +++ b/mlir-tensorrt/tensorrt/test/lib/TensorRT/TestTypeInferencePass.cpp @@ -1,3 +1,27 @@ +//===- TestTypeInferencePass.cpp ------------------------------------------===// +// +// SPDX-FileCopyrightText: Copyright 2024-2025 NVIDIA CORPORATION & AFFILIATES. +// All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +//===----------------------------------------------------------------------===// +/// +/// This file contains the implementation of the `test-tensorrt-shape-inference` +/// pass. +/// +//===----------------------------------------------------------------------===// #include "mlir-tensorrt-dialect/TensorRT/IR/TensorRTDialect.h" #include "mlir/Dialect/Arith/IR/Arith.h" #include "mlir/Dialect/Arith/Utils/Utils.h" @@ -13,6 +37,10 @@ using namespace mlir; +namespace mlir::tensorrt { +void registerTestTensorRTShapeInferencePass(); +} + struct DimOfReifyRankedShapedTypeOpInterface : public OpRewritePattern { using OpRewritePattern::OpRewritePattern; @@ -76,8 +104,6 @@ class TestTypeInferencePass }; } // namespace -namespace mlir { -void registerTestTensorRTShapeInferencePass() { +void tensorrt::registerTestTensorRTShapeInferencePass() { PassRegistration(); } -} // namespace mlir diff --git a/mlir-tensorrt/tensorrt/tools/tensorrt-opt/CMakeLists.txt b/mlir-tensorrt/tensorrt/tools/tensorrt-opt/CMakeLists.txt index 9e0cf81d8..7d2ecafe4 100644 --- a/mlir-tensorrt/tensorrt/tools/tensorrt-opt/CMakeLists.txt +++ b/mlir-tensorrt/tensorrt/tools/tensorrt-opt/CMakeLists.txt @@ -15,7 +15,4 @@ mtrt_add_tool(tensorrt-opt tensorrt-opt.cpp ) target_link_libraries(tensorrt-opt PRIVATE ${MTRTD_TEST_LIBS} - # TODO: this library is actually in the `compiler/` sub-project; - # we need to refactor this library locaiton. - MLIRTensorRTTestTensorKindAnalysis ) diff --git a/mlir-tensorrt/tensorrt/tools/tensorrt-opt/tensorrt-opt.cpp b/mlir-tensorrt/tensorrt/tools/tensorrt-opt/tensorrt-opt.cpp index 43a7fc87f..486e80d5b 100644 --- a/mlir-tensorrt/tensorrt/tools/tensorrt-opt/tensorrt-opt.cpp +++ b/mlir-tensorrt/tensorrt/tools/tensorrt-opt/tensorrt-opt.cpp @@ -34,10 +34,10 @@ #include "mlir/Tools/mlir-opt/MlirOptMain.h" #include "mlir/Transforms/Passes.h" -namespace mlir { +namespace mlir::tensorrt { void registerTestTensorKindAnalysisPass(); void registerTestTensorRTShapeInferencePass(); -} // namespace mlir +} // namespace mlir::tensorrt int main(int argc, char **argv) { mlir::DialectRegistry registry; @@ -45,8 +45,8 @@ int main(int argc, char **argv) { mlir::tensor::TensorDialect, mlir::arith::ArithDialect, mlir::affine::AffineDialect, mlir::quant::QuantDialect, mlir::scf::SCFDialect>(); - mlir::registerTestTensorKindAnalysisPass(); - mlir::registerTestTensorRTShapeInferencePass(); + mlir::tensorrt::registerTestTensorKindAnalysisPass(); + mlir::tensorrt::registerTestTensorRTShapeInferencePass(); mlir::func::registerInlinerExtension(registry); mlir::tensorrt::registerTensorRTPasses(); mlir::tensorrt::registerTensorRTTranslationPasses();