[REFACTOR][TARGET] Cleanup backend target registration (#19759)

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.

Author: tqchen

CMakeLists.txt

@@ -348,7 +348,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
     )

src/target/llvm/intrin_rule_nvptx.cc …d/cuda/codegen/llvm/intrin_rule_nvptx.ccsrc/target/llvm/intrin_rule_nvptx.cc renamed to src/backend/cuda/codegen/llvm/intrin_rule_nvptx.cc src/target/llvm/intrin_rule_nvptx.cc renamed to src/backend/cuda/codegen/llvm/intrin_rule_nvptx.cc

@@ -18,7 +18,7 @@
  */
 
 /*!
- * \file register.cc
+ * \file target_kind.cc
  * \brief CUDA compiler backend static registration.
  */
 #include <dlpack/dlpack.h>

src/backend/cuda/codegen/register.cc src/backend/cuda/codegen/target_kind.ccsrc/backend/cuda/codegen/register.cc renamed to src/backend/cuda/codegen/target_kind.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

src/backend/cuda/op/register.cc

@@ -18,7 +18,7 @@
  */
 
 /*!
- * \file register.cc
+ * \file target_kind.cc
  * \brief Hexagon compiler backend static registration.
  */
 #include <dlpack/dlpack.h>

src/backend/cuda/op/target_builtin.cc

@@ -18,7 +18,7 @@
  */
 
 /*!
- * \file register.cc
+ * \file target_kind.cc
  * \brief Metal compiler backend static registration.
  */
 #include <dlpack/dlpack.h>

src/backend/hexagon/codegen/register.cc …c/backend/hexagon/codegen/target_kind.ccsrc/backend/hexagon/codegen/register.cc renamed to src/backend/hexagon/codegen/target_kind.cc src/backend/hexagon/codegen/register.cc renamed to src/backend/hexagon/codegen/target_kind.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

src/backend/metal/codegen/register.cc src/backend/metal/codegen/target_kind.ccsrc/backend/metal/codegen/register.cc renamed to src/backend/metal/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>