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 <dlpack/dlpack.h>
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 <tvm/ffi/function.h>
-#include <tvm/runtime/base.h>
-
-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 <tvm/ffi/function.h>
+#include <tvm/runtime/base.h>
 #include <tvm/tirx/builtin.h>
 #include <tvm/tirx/op.h>
 #include <tvm/tirx/op_attr_types.h>
@@ -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 <dlpack/dlpack.h>
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 <dlpack/dlpack.h>
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 <tvm/ffi/function.h>
-#include <tvm/runtime/base.h>
-
-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 <tvm/ffi/function.h>
+#include <tvm/runtime/base.h>
 #include <tvm/tirx/op.h>
 #include <tvm/tirx/op_attr_types.h>
 
@@ -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 <dlpack/dlpack.h>
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 <dlpack/dlpack.h>
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 <dlpack/dlpack.h>
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 <tvm/ffi/function.h>
-#include <tvm/runtime/base.h>
-
-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 <tvm/ffi/function.h>
+#include <tvm/runtime/base.h>
 #include <tvm/tirx/builtin.h>
 #include <tvm/tirx/op.h>
 #include <tvm/tirx/op_attr_types.h>
@@ -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 <dlpack/dlpack.h>
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 <dlpack/dlpack.h>
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.