From cc459ab279f95ad516499ee63cc3d2c5adc11b87 Mon Sep 17 00:00:00 2001 From: tqchen Date: Sun, 14 Jun 2026 01:53:01 +0000 Subject: [PATCH] [REFACTOR][TARGET] Cleanup backend target registration The backend reorg leaves target-kind and op registration behind generic register.cc shims, which makes the new layout harder to navigate and keeps a few obsolete target paths alive. This change keeps backend registration beside the backend definitions while preserving the existing static initialization behavior. Backend target-kind registration units now use semantic target_kind.cc names, target builtin registration is colocated with the CUDA, Metal, and Trainium builtin definitions, the NVPTX intrinsic rules live with CUDA LLVM codegen, and the unused target/opt placeholder is removed. --- CMakeLists.txt | 2 +- .../cuda/codegen}/llvm/intrin_rule_nvptx.cc | 0 .../codegen/{register.cc => target_kind.cc} | 2 +- src/backend/cuda/op/register.cc | 35 ------------------- src/backend/cuda/op/target_builtin.cc | 4 +++ .../codegen/{register.cc => target_kind.cc} | 2 +- .../codegen/{register.cc => target_kind.cc} | 2 +- src/backend/metal/op/register.cc | 35 ------------------- src/backend/metal/op/target_builtin.cc | 4 +++ .../codegen/{register.cc => target_kind.cc} | 2 +- .../codegen/{register.cc => target_kind.cc} | 2 +- .../codegen/{register.cc => target_kind.cc} | 2 +- src/backend/trn/op/register.cc | 35 ------------------- src/backend/trn/op/target_builtin.cc | 4 +++ .../codegen/{register.cc => target_kind.cc} | 2 +- .../codegen/{register.cc => target_kind.cc} | 2 +- src/target/opt/README | 1 - 17 files changed, 21 insertions(+), 115 deletions(-) rename src/{target => backend/cuda/codegen}/llvm/intrin_rule_nvptx.cc (100%) rename src/backend/cuda/codegen/{register.cc => target_kind.cc} (99%) delete mode 100644 src/backend/cuda/op/register.cc rename src/backend/hexagon/codegen/{register.cc => target_kind.cc} (98%) rename src/backend/metal/codegen/{register.cc => target_kind.cc} (98%) delete mode 100644 src/backend/metal/op/register.cc rename src/backend/opencl/codegen/{register.cc => target_kind.cc} (99%) rename src/backend/rocm/codegen/{register.cc => target_kind.cc} (99%) rename src/backend/trn/codegen/{register.cc => target_kind.cc} (98%) delete mode 100644 src/backend/trn/op/register.cc rename src/backend/vulkan/codegen/{register.cc => target_kind.cc} (99%) rename src/backend/webgpu/codegen/{register.cc => target_kind.cc} (99%) delete mode 100644 src/target/opt/README diff --git a/CMakeLists.txt b/CMakeLists.txt index ad99c4c6acba..bc1d32232a72 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -349,7 +349,7 @@ tvm_file_glob(GLOB CODEGEN_SRCS src/backend/trn/codegen/*.cc src/backend/trn/op/*.cc src/backend/trn/transform/*.cc - src/backend/vulkan/codegen/register.cc + src/backend/vulkan/codegen/target_kind.cc src/backend/vulkan/codegen/vulkan_fallback_module.cc src/backend/webgpu/codegen/*.cc ) diff --git a/src/target/llvm/intrin_rule_nvptx.cc b/src/backend/cuda/codegen/llvm/intrin_rule_nvptx.cc similarity index 100% rename from src/target/llvm/intrin_rule_nvptx.cc rename to src/backend/cuda/codegen/llvm/intrin_rule_nvptx.cc diff --git a/src/backend/cuda/codegen/register.cc b/src/backend/cuda/codegen/target_kind.cc similarity index 99% rename from src/backend/cuda/codegen/register.cc rename to src/backend/cuda/codegen/target_kind.cc index d1fa868183c2..e36131e74f81 100644 --- a/src/backend/cuda/codegen/register.cc +++ b/src/backend/cuda/codegen/target_kind.cc @@ -18,7 +18,7 @@ */ /*! - * \file register.cc + * \file target_kind.cc * \brief CUDA compiler backend static registration. */ #include diff --git a/src/backend/cuda/op/register.cc b/src/backend/cuda/op/register.cc deleted file mode 100644 index 51a7edbfcf4e..000000000000 --- a/src/backend/cuda/op/register.cc +++ /dev/null @@ -1,35 +0,0 @@ -/* - * Licensed to the Apache Software Foundation (ASF) under one - * or more contributor license agreements. See the NOTICE file - * distributed with this work for additional information - * regarding copyright ownership. The ASF licenses this file - * to you 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. - */ - -/*! - * \file register.cc - * \brief CUDA backend op static registration. - */ -#include -#include - -namespace tvm { -namespace tirx { -namespace builtin { -void RegisterCudaTargetBuiltins(); -} // namespace builtin -} // namespace tirx -} // namespace tvm - -TVM_FFI_STATIC_INIT_BLOCK() { tvm::tirx::builtin::RegisterCudaTargetBuiltins(); } diff --git a/src/backend/cuda/op/target_builtin.cc b/src/backend/cuda/op/target_builtin.cc index 9ebaaf590343..005fe5b32263 100644 --- a/src/backend/cuda/op/target_builtin.cc +++ b/src/backend/cuda/op/target_builtin.cc @@ -22,6 +22,8 @@ * * builtin intrinsic operators specific to CUDA target. */ +#include +#include #include #include #include @@ -597,6 +599,8 @@ void RegisterDeviceIntrinsicAliases() { #undef TIRX_DEFINE_BUILTIN_FUNC +TVM_FFI_STATIC_INIT_BLOCK() { RegisterCudaTargetBuiltins(); } + } // namespace builtin } // namespace tirx } // namespace tvm diff --git a/src/backend/hexagon/codegen/register.cc b/src/backend/hexagon/codegen/target_kind.cc similarity index 98% rename from src/backend/hexagon/codegen/register.cc rename to src/backend/hexagon/codegen/target_kind.cc index d6576bfee74c..c71d734f9f4d 100644 --- a/src/backend/hexagon/codegen/register.cc +++ b/src/backend/hexagon/codegen/target_kind.cc @@ -18,7 +18,7 @@ */ /*! - * \file register.cc + * \file target_kind.cc * \brief Hexagon compiler backend static registration. */ #include diff --git a/src/backend/metal/codegen/register.cc b/src/backend/metal/codegen/target_kind.cc similarity index 98% rename from src/backend/metal/codegen/register.cc rename to src/backend/metal/codegen/target_kind.cc index e90651775e57..3fda53fe708c 100644 --- a/src/backend/metal/codegen/register.cc +++ b/src/backend/metal/codegen/target_kind.cc @@ -18,7 +18,7 @@ */ /*! - * \file register.cc + * \file target_kind.cc * \brief Metal compiler backend static registration. */ #include diff --git a/src/backend/metal/op/register.cc b/src/backend/metal/op/register.cc deleted file mode 100644 index 6a12ee1282c8..000000000000 --- a/src/backend/metal/op/register.cc +++ /dev/null @@ -1,35 +0,0 @@ -/* - * Licensed to the Apache Software Foundation (ASF) under one - * or more contributor license agreements. See the NOTICE file - * distributed with this work for additional information - * regarding copyright ownership. The ASF licenses this file - * to you 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. - */ - -/*! - * \file register.cc - * \brief Metal backend op static registration. - */ -#include -#include - -namespace tvm { -namespace tirx { -namespace builtin { -void RegisterMetalTargetBuiltins(); -} // namespace builtin -} // namespace tirx -} // namespace tvm - -TVM_FFI_STATIC_INIT_BLOCK() { tvm::tirx::builtin::RegisterMetalTargetBuiltins(); } diff --git a/src/backend/metal/op/target_builtin.cc b/src/backend/metal/op/target_builtin.cc index f3bd50b7ab59..08b3bfd95a6c 100644 --- a/src/backend/metal/op/target_builtin.cc +++ b/src/backend/metal/op/target_builtin.cc @@ -22,6 +22,8 @@ * * builtin intrinsic operators specific to Metal target. */ +#include +#include #include #include @@ -57,6 +59,8 @@ TIRX_DEFINE_BUILTIN_FUNC(simdgroup_multiply_accumulate) #undef TIRX_DEFINE_BUILTIN_FUNC +TVM_FFI_STATIC_INIT_BLOCK() { RegisterMetalTargetBuiltins(); } + } // namespace builtin } // namespace tirx } // namespace tvm diff --git a/src/backend/opencl/codegen/register.cc b/src/backend/opencl/codegen/target_kind.cc similarity index 99% rename from src/backend/opencl/codegen/register.cc rename to src/backend/opencl/codegen/target_kind.cc index 1b64cefeb717..dfdc10e062da 100644 --- a/src/backend/opencl/codegen/register.cc +++ b/src/backend/opencl/codegen/target_kind.cc @@ -18,7 +18,7 @@ */ /*! - * \file register.cc + * \file target_kind.cc * \brief OpenCL compiler backend static registration. */ #include diff --git a/src/backend/rocm/codegen/register.cc b/src/backend/rocm/codegen/target_kind.cc similarity index 99% rename from src/backend/rocm/codegen/register.cc rename to src/backend/rocm/codegen/target_kind.cc index 4d0c8b738b2b..d460ae29951f 100644 --- a/src/backend/rocm/codegen/register.cc +++ b/src/backend/rocm/codegen/target_kind.cc @@ -18,7 +18,7 @@ */ /*! - * \file register.cc + * \file target_kind.cc * \brief ROCm compiler backend static registration. */ #include diff --git a/src/backend/trn/codegen/register.cc b/src/backend/trn/codegen/target_kind.cc similarity index 98% rename from src/backend/trn/codegen/register.cc rename to src/backend/trn/codegen/target_kind.cc index ca5e5876bbef..2c1e9676e663 100644 --- a/src/backend/trn/codegen/register.cc +++ b/src/backend/trn/codegen/target_kind.cc @@ -18,7 +18,7 @@ */ /*! - * \file register.cc + * \file target_kind.cc * \brief Trainium compiler backend static registration. */ #include diff --git a/src/backend/trn/op/register.cc b/src/backend/trn/op/register.cc deleted file mode 100644 index 365d7d6a9651..000000000000 --- a/src/backend/trn/op/register.cc +++ /dev/null @@ -1,35 +0,0 @@ -/* - * Licensed to the Apache Software Foundation (ASF) under one - * or more contributor license agreements. See the NOTICE file - * distributed with this work for additional information - * regarding copyright ownership. The ASF licenses this file - * to you 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. - */ - -/*! - * \file register.cc - * \brief Trainium backend op static registration. - */ -#include -#include - -namespace tvm { -namespace tirx { -namespace builtin { -void RegisterTRNTargetBuiltins(); -} // namespace builtin -} // namespace tirx -} // namespace tvm - -TVM_FFI_STATIC_INIT_BLOCK() { tvm::tirx::builtin::RegisterTRNTargetBuiltins(); } diff --git a/src/backend/trn/op/target_builtin.cc b/src/backend/trn/op/target_builtin.cc index baf7ce05bb17..c0d915bb2a36 100644 --- a/src/backend/trn/op/target_builtin.cc +++ b/src/backend/trn/op/target_builtin.cc @@ -23,6 +23,8 @@ * * builtin intrinsic operators specific to Trainium target. */ +#include +#include #include #include #include @@ -161,6 +163,8 @@ void RegisterNKIIntrinsicAliases() { #undef TIRX_DEFINE_BUILTIN_FUNC +TVM_FFI_STATIC_INIT_BLOCK() { RegisterTRNTargetBuiltins(); } + } // namespace builtin } // namespace tirx } // namespace tvm diff --git a/src/backend/vulkan/codegen/register.cc b/src/backend/vulkan/codegen/target_kind.cc similarity index 99% rename from src/backend/vulkan/codegen/register.cc rename to src/backend/vulkan/codegen/target_kind.cc index 92d0d115a84d..d956196ddd66 100644 --- a/src/backend/vulkan/codegen/register.cc +++ b/src/backend/vulkan/codegen/target_kind.cc @@ -18,7 +18,7 @@ */ /*! - * \file register.cc + * \file target_kind.cc * \brief Vulkan compiler backend static registration. */ #include diff --git a/src/backend/webgpu/codegen/register.cc b/src/backend/webgpu/codegen/target_kind.cc similarity index 99% rename from src/backend/webgpu/codegen/register.cc rename to src/backend/webgpu/codegen/target_kind.cc index 49e90245023c..0447bce49927 100644 --- a/src/backend/webgpu/codegen/register.cc +++ b/src/backend/webgpu/codegen/target_kind.cc @@ -18,7 +18,7 @@ */ /*! - * \file register.cc + * \file target_kind.cc * \brief WebGPU compiler backend static registration. */ #include diff --git a/src/target/opt/README b/src/target/opt/README deleted file mode 100644 index 535e8c03d114..000000000000 --- a/src/target/opt/README +++ /dev/null @@ -1 +0,0 @@ -This folder contains optional builds for codegen depending on compilation options.