#include "cache.h"
#include "ops.h"
#include "core/registration.h"

#include <torch/library.h>

// Note: overwrite the external definition for sharing same name between
// libraries use different ISAs.
#define TORCH_EXTENSION_NAME _C
// ...

TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
  // vLLM custom ops
  // ...

  // Activation ops

  // Activation function used in SwiGLU.
  ops.def("silu_and_mul(Tensor! out, Tensor input) -> ()");
  ops.impl("silu_and_mul", torch::kCPU, &silu_and_mul);

  // Activation function used in GeGLU with `none` approximation.
  ops.def("gelu_and_mul(Tensor! out, Tensor input) -> ()");
  ops.impl("gelu_and_mul", torch::kCPU, &gelu_and_mul);
  // ...
}

REGISTER_EXTENSION(TORCH_EXTENSION_NAME)

#define SHOW(expr) std::cout << #expr " = " << (expr) << '\n'
SHOW(x + 1);   // 展开为：std::cout << "x + 1" " = " << (x + 1) << '\n';

#define DECLARE_INIT(ns) static void TORCH_LIBRARY_init_##ns(torch::Library&);
DECLARE_INIT(myops)   // 展开为：static void TORCH_LIBRARY_init_myops(torch::Library&);

#define LOG(fmt, ...) std::printf(fmt, __VA_ARGS__)

#define TORCH_CHECK_MSG(cond, type, ...)                   \
  (::c10::detail::torchCheckMsgImpl(                       \
      "Expected " #cond                                    \
      " to be true, but got false.  "                      \
      "(Could this error message be improved?  If so, "    \
      "please report an enhancement request to PyTorch.)", \
      ##__VA_ARGS__))

#define C10_CONCATENATE_IMPL(s1, s2) s1##s2
#define C10_CONCATENATE(s1, s2) C10_CONCATENATE_IMPL(s1, s2)

#define C10_STRINGIZE_IMPL(x) #x
#define C10_STRINGIZE(x) C10_STRINGIZE_IMPL(x)

#define STR1(x) #x
#define STR2(x) STR1(x)
STR1(__LINE__)   // "__LINE__"   ← 实参紧挨 #，不展开
STR2(__LINE__)   // "42"         ← 实参先展开成 42，再传给 STR1

#ifdef __COUNTER__
#define C10_UID __COUNTER__
#define C10_ANONYMOUS_VARIABLE(str) C10_CONCATENATE(str, __COUNTER__)
#else
#define C10_UID __LINE__
#define C10_ANONYMOUS_VARIABLE(str) C10_CONCATENATE(str, __LINE__)
#endif

#ifdef _WIN32
#define C10_HIDDEN
#if defined(C10_BUILD_SHARED_LIBS)
#define C10_EXPORT __declspec(dllexport)
#define C10_IMPORT __declspec(dllimport)
#else
#define C10_EXPORT
#define C10_IMPORT
#endif
#else // _WIN32
#if defined(__GNUC__)
#define C10_EXPORT __attribute__((__visibility__("default")))
#define C10_HIDDEN __attribute__((__visibility__("hidden")))
#else // defined(__GNUC__)
#define C10_EXPORT
#define C10_HIDDEN
#endif // defined(__GNUC__)
#define C10_IMPORT C10_EXPORT
#endif // _WIN32

// torch/headeronly/core/ScalarType.h
#define AT_FORALL_SCALAR_TYPES(_) \
  _(uint8_t, Byte)                \
  _(int8_t, Char)                 \
  _(int16_t, Short)               \
  _(int, Int)                     \
  _(int64_t, Long)                \
  _(float, Float)                 \
  _(double, Double)

enum class ScalarType : int8_t {
#define DEFINE_ST_ENUM_VAL_(_1, n) n,
  AT_FORALL_SCALAR_TYPES_WITH_COMPLEX_AND_QINTS(DEFINE_ST_ENUM_VAL_)
#undef DEFINE_ENUM_ST_ENUM_VAL_
      Undefined,
  NumOptions
};

inline const char* toString(ScalarType t) {
#define DEFINE_CASE(_, name) \
  case ScalarType::name:     \
    return #name;

  switch (t) {
    AT_FORALL_SCALAR_TYPES_WITH_COMPLEX_AND_QINTS(DEFINE_CASE)
    default:
      return "UNKNOWN_SCALAR";
  }
#undef DEFINE_CASE
}

// c10/util/Exception.h
#define C10_THROW_ERROR(err_type, msg) \
  throw ::c10::err_type(               \
      {__func__, __FILE__, static_cast<uint32_t>(__LINE__)}, msg)

#define TORCH_CHECK(cond, ...)                     \
  if (C10_UNLIKELY_OR_CONST(!(cond))) {            \
    ::c10::detail::torchCheckFail(                 \
        __func__,                                  \
        __FILE__,                                  \
        static_cast<uint32_t>(__LINE__),           \
        TORCH_CHECK_MSG(cond, "", ##__VA_ARGS__)); \
  }

#if defined(__GNUC__) || defined(__ICL) || defined(__clang__)
#define C10_LIKELY(expr) (__builtin_expect(static_cast<bool>(expr), 1))
#define C10_UNLIKELY(expr) (__builtin_expect(static_cast<bool>(expr), 0))
#else
#define C10_LIKELY(expr) (expr)
#define C10_UNLIKELY(expr) (expr)
#endif

#if defined(__CUDACC__)
#define C10_UNLIKELY_OR_CONST(e) e
#else
#define C10_UNLIKELY_OR_CONST(e) C10_UNLIKELY(e)
#endif

namespace c10::detail {
template <typename... Args>
auto torchCheckMsgImpl(const char* /*msg*/, const Args&... args) {
  return ::c10::str(args...);
}
inline C10_API const char* torchCheckMsgImpl(const char* msg) {
  return msg;
}
// If there is just 1 user-provided C-string argument, use it.
inline C10_API const char* torchCheckMsgImpl(
    const char* /*msg*/,
    const char* args) {
  return args;
}
} // namespace c10::detail

#define TORCH_CHECK_MSG(cond, type, ...)                   \
  (::c10::detail::torchCheckMsgImpl(                       \
      "Expected " #cond                                    \
      " to be true, but got false.  "                      \
      "(Could this error message be improved?  If so, "    \
      "please report an enhancement request to PyTorch.)", \
      ##__VA_ARGS__))

template <typename... Args>
inline auto str(const Args&... args) {
  return detail::_str_wrapper<
      typename detail::CanonicalizeStrTypes<Args>::type...>::call(args...);
}

// c10/util/Exception.cpp
void torchCheckFail(
    const char* func,
    const char* file,
    uint32_t line,
    const std::string& msg) {
  throw ::c10::Error({func, file, line}, msg);
}

#define TORCH_INTERNAL_ASSERT(cond, ...)                                         \
  if (C10_UNLIKELY_OR_CONST(!(cond))) {                                          \
    ::c10::detail::torchInternalAssertFail(                                      \
        __func__,                                                                \
        __FILE__,                                                                \
        static_cast<uint32_t>(__LINE__),                                         \
        #cond                                                                    \
        " INTERNAL ASSERT FAILED at " C10_STRINGIZE(__FILE__) ":" C10_STRINGIZE( \
            __LINE__) ", please report a bug to PyTorch. ",                      \
        c10::str(__VA_ARGS__));                                                  \
  }
#endif

#ifdef STRIP_ERROR_MESSAGES
#define TORCH_CHECK_MSG(cond, type, ...) \
  (#cond #type " CHECK FAILED at " C10_STRINGIZE(__FILE__))

#define TORCH_CHECK_LINALG(cond, ...) \
  TORCH_CHECK_WITH_MSG(LinAlgError, cond, "LINALG", __VA_ARGS__)
#define TORCH_CHECK_INDEX(cond, ...) \
  TORCH_CHECK_WITH_MSG(IndexError, cond, "INDEX", __VA_ARGS__)
#define TORCH_CHECK_VALUE(cond, ...) \
  TORCH_CHECK_WITH_MSG(ValueError, cond, "VALUE", __VA_ARGS__)
#define TORCH_CHECK_TYPE(cond, ...) \
  TORCH_CHECK_WITH_MSG(TypeError, cond, "TYPE", __VA_ARGS__)
#define TORCH_CHECK_NOT_IMPLEMENTED(cond, ...) \
  TORCH_CHECK_WITH_MSG(NotImplementedError, cond, "TYPE", __VA_ARGS__)

// 注册表：全局单例
struct Registry {
  static Registry& get() { static Registry r; return r; }
  std::map<std::string, Factory> table;
};

// 注册器：构造函数把一个条目写进注册表
struct Registrar {
  Registrar(const char* name, Factory f) { Registry::get().table[name] = f; }
};

// 使用：在每个实现文件里定义一个 static Registrar 对象
static Registrar reg_foo("foo", &make_foo);   // 加载时自动执行

template <class SrcType, class ObjectPtrType, class... Args>
class Registry {
 public:
  typedef std::function<ObjectPtrType(Args...)> Creator;

  void Register(
      const SrcType& key,
      Creator creator,
      const RegistryPriority priority = REGISTRY_DEFAULT) {
    std::lock_guard<std::mutex> lock(register_mutex_);
    // ...
    registry_[key] = creator;
    priority_[key] = priority;
  }

  ObjectPtrType Create(const SrcType& key, Args... args) {
    auto it = registry_.find(key);
    if (it == registry_.end()) {
      // Returns nullptr if the key is not registered.
      return nullptr;
    }
    return it->second(args...);
  }
  // ...
 private:
  std::unordered_map<SrcType, Creator> registry_;
  std::unordered_map<SrcType, RegistryPriority> priority_;
  // ...
  std::mutex register_mutex_;
};

template <class SrcType, class ObjectPtrType, class... Args>
class Registerer {
 public:
  explicit Registerer(
      const SrcType& key,
      Registry<SrcType, ObjectPtrType, Args...>* registry,
      typename Registry<SrcType, ObjectPtrType, Args...>::Creator creator,
      const std::string& help_msg = "") {
    registry->Register(key, creator, help_msg);
  }
  // ...
};

#define C10_REGISTER_TYPED_CLASS(RegistryName, key, ...)                    \
  static Registerer##RegistryName C10_ANONYMOUS_VARIABLE(g_##RegistryName)( \
      key,                                                                  \
      RegistryName(),                                                       \
      Registerer##RegistryName::DefaultCreator<__VA_ARGS__>,                \
      ::c10::demangle_type<__VA_ARGS__>());

#define C10_REGISTER_GUARD_IMPL(DevType, DeviceGuardImpl)              \
  static ::c10::impl::DeviceGuardImplRegistrar C10_ANONYMOUS_VARIABLE( \
      g_##DeviceType)(::c10::DeviceType::DevType, new DeviceGuardImpl());

#define REGISTER_ARCH_DISPATCH(name, arch, fn) \
  template <> name##_DECLARE_DISPATCH_type::FnPtr TORCH_API DispatchStub<name##_DECLARE_DISPATCH_type::FnPtr, struct name##_DECLARE_DISPATCH_type>::arch = fn;

#define REGISTER_CUDA_DISPATCH(name, fn) \
  static RegisterCUDADispatch<struct name##_DECLARE_DISPATCH_type> name ## __register(name, fn);

#if defined(__CUDACC__)
#define REGISTER_DISPATCH(name, fn) REGISTER_CUDA_DISPATCH(name, fn)
#elif defined(__HIPCC__)
#define REGISTER_DISPATCH(name, fn) REGISTER_CUDA_DISPATCH(name, fn)
#elif defined(__OBJC__) && defined(USE_MPS)
#define REGISTER_DISPATCH(name, fn) REGISTER_MPS_DISPATCH(name, fn)
#elif defined(CPU_CAPABILITY)
// ...
#define REGISTER_DISPATCH(name, fn) REGISTER_ARCH_DISPATCH(name, CPU_CAPABILITY, fn)
#endif

#define TORCH_LIBRARY(ns, m)                                                   \
  static void TORCH_LIBRARY_init_##ns(torch::Library&);                        \
  static const torch::detail::TorchLibraryInit TORCH_LIBRARY_static_init_##ns( \
      torch::Library::DEF,                                                     \
      &TORCH_LIBRARY_init_##ns,                                                \
      C10_STRINGIZE(ns),                                                                     \
      std::nullopt,                                                            \
      __FILE__,                                                                \
      __LINE__);                                                               \
  void TORCH_LIBRARY_init_##ns(torch::Library& m)

TORCH_LIBRARY(myops, m) {
  m.def("scale_shift(Tensor x, float alpha, float beta) -> Tensor");
}

static void TORCH_LIBRARY_init_myops(torch::Library&);
static const torch::detail::TorchLibraryInit TORCH_LIBRARY_static_init_myops(
    torch::Library::DEF,
    &TORCH_LIBRARY_init_myops,
    "myops",
    std::nullopt,
    "ext.cpp",
    7);
void TORCH_LIBRARY_init_myops(torch::Library& m) {
  m.def("scale_shift(Tensor x, float alpha, float beta) -> Tensor");
}

namespace torch::detail {
class TorchLibraryInit final {
 private:
  using InitFn = void(Library&);
  Library lib_;

 public:
  TorchLibraryInit(
      Library::Kind kind,
      InitFn* fn,
      const char* ns,
      std::optional<c10::DispatchKey> k,
      const char* file,
      uint32_t line)
      : lib_(kind, ns, k, file, line) {
    fn(lib_);
  }
};
} // namespace torch::detail

Library::Library(Kind kind, std::string ns, std::optional<c10::DispatchKey> k, const char* file, uint32_t line)
  : kind_(kind)
  , ns_(ns == "_" ? std::nullopt : std::make_optional(std::move(ns)))
  , dispatch_key_(k.value_or(CatchAll) == CatchAll ? std::optional<c10::DispatchKey>() : k)
  , file_(file)
  , line_(line)
  {
    switch (kind_) {
      case DEF:
        // Only DEFs require library uniqueness; fragments
        // don't register a library
        registrars_.emplace_back(
          c10::Dispatcher::singleton().registerLibrary(
            // NOLINTNEXTLINE(bugprone-unchecked-optional-access)
            ns_.value(), debugString(file_, line_)
          )
        );
        [[fallthrough]];
      case FRAGMENT:
        TORCH_CHECK(
          ns_.has_value(),
          toString(kind_), ": cannot define ", toString(kind_), " with the wildcard namespace _ "
          "(every ", toString(kind_), " defines operators for a distinct namespace!) "
          "Did you mean to use TORCH_LIBRARY_IMPL instead?  "
          ERROR_CONTEXT
        );
        TORCH_INTERNAL_ASSERT(!dispatch_key_.has_value(), ERROR_CONTEXT);
        break;
      case IMPL:
        // Nothing to do, everything is OK
        break;
    }
  }

#define ERROR_CONTEXT "(Error occurred while processing ", toString(kind_), " block at ", file_, ":", line_, ")"

      registrars_.emplace_back(
        c10::Dispatcher::singleton().registerDef(
          std::move(schema),
          debugString(file_, line_),
          tags
        )
      );

Library& Library::_impl(const char* name_str, CppFunction&& f, _RegisterOrVerify rv) & {
  at::OperatorName name = _parseNameForLib(name_str);
  // ...
  auto dispatch_key = f.dispatch_key_.has_value() ? f.dispatch_key_ : dispatch_key_;
  switch (rv) {
    case _RegisterOrVerify::REGISTER:
      registrars_.emplace_back(
        c10::Dispatcher::singleton().registerImpl(
          std::move(name),
          dispatch_key,
          std::move(f.func_),
          f.cpp_signature_,
          std::move(f.schema_),
          debugString(std::move(f.debug_), file_, line_)
        )
      );
      break;
    // ...
  }
  return *this;
}

#define TORCH_LIBRARY_IMPL(ns, k, m) _TORCH_LIBRARY_IMPL(ns, k, m, C10_UID)

#define _TORCH_LIBRARY_IMPL(ns, k, m, uid)                                \
  static void C10_CONCATENATE(                                            \
      TORCH_LIBRARY_IMPL_init_##ns##_##k##_, uid)(torch::Library&);       \
  static const torch::detail::TorchLibraryInit C10_CONCATENATE(           \
      TORCH_LIBRARY_IMPL_static_init_##ns##_##k##_, uid)(                 \
      torch::Library::IMPL,                                               \
      &C10_CONCATENATE(TORCH_LIBRARY_IMPL_init_##ns##_##k##_, uid),       \
      C10_STRINGIZE(ns),                                                                \
      std::make_optional(c10::DispatchKey::k),                            \
      __FILE__,                                                           \
      __LINE__);                                                          \
  void C10_CONCATENATE(                                                   \
      TORCH_LIBRARY_IMPL_init_##ns##_##k##_, uid)(torch::Library & m)

static void TORCH_LIBRARY_IMPL_init_myops_CPU_0(torch::Library&);
static const torch::detail::TorchLibraryInit TORCH_LIBRARY_IMPL_static_init_myops_CPU_0(
    torch::Library::IMPL,
    &TORCH_LIBRARY_IMPL_init_myops_CPU_0,
    "myops",
    std::make_optional(c10::DispatchKey::CPU),
    "ext.cpp",
    23);
void TORCH_LIBRARY_IMPL_init_myops_CPU_0(torch::Library& m) {
  m.impl("scale_shift", &scale_shift_cpu);
}

// NB: The EXACT NAMING of the initializer functions (e.g.,
// TORCH_LIBRARY_init_aten) matters for the code analyzer;
// see the regexes at tools/code_analyzer/run_analyzer.sh

// registry.cpp
std::map<std::string, Factory> g_registry;          // 动态初始化：要跑 std::map 的构造函数

// foo.cpp
static Registrar reg_foo("foo", &make_foo);        // 构造函数里写 g_registry

std::map<std::string, Factory>& registry() {
  static std::map<std::string, Factory> r;   // 第一次调用时构造
  return r;
}

#define C10_DEFINE_TYPED_REGISTRY(                                         \
    RegistryName, SrcType, ObjectType, PtrType, ...)                       \
  C10_EXPORT ::c10::Registry<SrcType, PtrType<ObjectType>, ##__VA_ARGS__>* \
  RegistryName() {                                                         \
    static ::c10::Registry<SrcType, PtrType<ObjectType>, ##__VA_ARGS__>*   \
        registry = new ::c10::                                             \
            Registry<SrcType, PtrType<ObjectType>, ##__VA_ARGS__>();       \
    return registry;                                                       \
  }

  static Dispatcher& realSingleton();

  C10_ALWAYS_INLINE static Dispatcher& singleton() {
#if !defined C10_MOBILE
    // Implemented inline so that steady-state code needn't incur
    // function-call overhead. We can't just inline `realSingleton`
    // because the function-local static would get duplicated across
    // all DSOs that include & use this header, leading to multiple
    // singleton instances.
    static Dispatcher& s = realSingleton();
    return s;
#else
    // ...
    return realSingleton();
#endif
  }

// aten/src/ATen/core/dispatch/Dispatcher.cpp
C10_EXPORT Dispatcher& Dispatcher::realSingleton() {
  static Dispatcher _singleton;
  return _singleton;
}

// This one is being used by libc10.so
#ifdef C10_BUILD_MAIN_LIB
#define C10_API C10_EXPORT
#else
#define C10_API C10_IMPORT
#endif

// This one is being used by libtorch.so
#ifdef CAFFE2_BUILD_MAIN_LIB
#define TORCH_API C10_EXPORT
#else
#define TORCH_API C10_IMPORT
#endif

// libtorch_cuda.so (where torch_cuda_cu and torch_cuda_cpp are a part of the
// same api)
#ifdef TORCH_CUDA_BUILD_MAIN_LIB
#define TORCH_CUDA_CPP_API C10_EXPORT
#define TORCH_CUDA_CU_API C10_EXPORT
#else
#define TORCH_CUDA_CPP_API C10_IMPORT
#define TORCH_CUDA_CU_API C10_IMPORT
#endif

#pragma once

#include <Python.h>

#define _CONCAT(A, B) A##B
#define CONCAT(A, B) _CONCAT(A, B)

#define _STRINGIFY(A) #A
#define STRINGIFY(A) _STRINGIFY(A)

// A version of the TORCH_LIBRARY macro that expands the NAME, i.e. so NAME
// could be a macro instead of a literal token.
#define TORCH_LIBRARY_EXPAND(NAME, MODULE) TORCH_LIBRARY(NAME, MODULE)

// A version of the TORCH_LIBRARY_IMPL macro that expands the NAME, i.e. so NAME
// could be a macro instead of a literal token.
#define TORCH_LIBRARY_IMPL_EXPAND(NAME, DEVICE, MODULE) \
  TORCH_LIBRARY_IMPL(NAME, DEVICE, MODULE)

// REGISTER_EXTENSION allows the shared library to be loaded and initialized
// via python's import statement.
#define REGISTER_EXTENSION(NAME)                                               \
  PyMODINIT_FUNC CONCAT(PyInit_, NAME)() {                                     \
    static struct PyModuleDef module = {PyModuleDef_HEAD_INIT,                 \
                                        STRINGIFY(NAME), nullptr, 0, nullptr}; \
    return PyModule_Create(&module);                                           \
  }

// csrc/libtorch_stable/torch_bindings.cpp
STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
  // ...
  ops.def("permute_cols(Tensor A, Tensor perm) -> Tensor");
  // ...
}

STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
  // ...
  ops.impl("per_token_group_fp8_quant", TORCH_BOX(&per_token_group_quant_fp8));
  // ...
}

/// C10_NOINLINE - Functions whose declaration is annotated with this will not
/// be inlined.
#ifdef __GNUC__
#define C10_NOINLINE __attribute__((noinline))
#elif _MSC_VER
#define C10_NOINLINE __declspec(noinline)
#else
#define C10_NOINLINE
#endif

#if defined(_MSC_VER)
#define C10_ALWAYS_INLINE __forceinline
#elif __has_attribute(always_inline) || defined(__GNUC__)
#define C10_ALWAYS_INLINE __attribute__((__always_inline__)) inline
#else
#define C10_ALWAYS_INLINE inline
#endif

#ifndef HAS_DEMANGLE
#if defined(__ANDROID__) || defined(_WIN32) || defined(__EMSCRIPTEN__)
#define HAS_DEMANGLE 0
#elif defined(__APPLE__) && \
    (TARGET_IPHONE_SIMULATOR || TARGET_OS_SIMULATOR || TARGET_OS_IPHONE)
#define HAS_DEMANGLE 0
#else
#define HAS_DEMANGLE 1
#endif
#endif // HAS_DEMANGLE

#if defined(__CUDACC__) || defined(__HIPCC__)
// Designates functions callable from the host (CPU) and the device (GPU)
#define C10_HOST_DEVICE __host__ __device__
#define C10_DEVICE __device__
#define C10_HOST __host__
// ...
#else
#define C10_HOST_DEVICE
#define C10_HOST
#define C10_DEVICE
#endif

#if __CUDA_ARCH__ == 750
constexpr uint32_t CUDA_MAX_THREADS_PER_SM = 1024;
#elif __CUDA_ARCH__ == 860 || __CUDA_ARCH__ == 870 || __CUDA_ARCH__ == 890 || \
    __CUDA_ARCH__ == 1200
constexpr uint32_t CUDA_MAX_THREADS_PER_SM = 1536;
#else
constexpr uint32_t CUDA_MAX_THREADS_PER_SM = 2048;
#endif

inline __device__ float2 bf1622float2(const __nv_bfloat162 val) {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 800
  assert(false);
#else
  return __bfloat1622float2(val);
#endif
  __builtin_unreachable();  // Suppress missing return statement warning
}

// aten/src/ATen/native/SummaryOps.cpp
Tensor
_bincount_cpu(const Tensor& self, const std::optional<Tensor>& weights_opt, int64_t minlength) {
  // See [Note: hacky wrapper removal for optional tensor]
  c10::MaybeOwned<Tensor> weights_maybe_owned = at::borrow_from_optional_tensor(weights_opt);
  const Tensor& weights = *weights_maybe_owned;

  return AT_DISPATCH_INTEGRAL_TYPES(self.scalar_type(), "bincount_cpu", [&] {
    // ...
  });
}

// ${generated_comment}
#define TORCH_ASSERT_ONLY_METHOD_OPERATORS
#include <torch/library.h>

namespace at {
TORCH_LIBRARY(aten, m) {
  ${aten_schema_registrations};
  // Distributed Ops
  // Implementations located in torch/csrc/jit/runtime/register_distributed_ops.cpp
  m.def("get_gradients(int context_id) -> Dict(Tensor, Tensor)");
}
${schema_registrations}
}  // namespace at

m.def("bincount(Tensor self, Tensor? weights=None, SymInt minlength=0) -> Tensor", tags_N);

// aten::bincount(Tensor self, Tensor? weights=None, SymInt minlength=0) -> Tensor
inline at::Tensor bincount(const at::Tensor & self, const ::std::optional<at::Tensor> & weights={}, int64_t minlength=0) {
    return at::_ops::bincount::call(self, weights, minlength);
}
// 以及一个 bincount_symint(..., c10::SymInt minlength=0) 重载

struct TORCH_API bincount {
  using schema = at::Tensor (const at::Tensor &, const ::std::optional<at::Tensor> &, c10::SymInt);
  using ptr_schema = schema*;
  // See Note [static constexpr char* members for windows NVCC]
  static constexpr const char* name = "aten::bincount";
  static constexpr const char* overload_name = "";
  static constexpr const char* schema_str = "bincount(Tensor self, Tensor? weights=None, SymInt minlength=0) -> Tensor";
  static at::Tensor call(const at::Tensor & self, const ::std::optional<at::Tensor> & weights, c10::SymInt minlength);
  static at::Tensor redispatch(c10::DispatchKeySet dispatchKeySet, const at::Tensor & self, const ::std::optional<at::Tensor> & weights, c10::SymInt minlength);
};

// aten::bincount(Tensor self, Tensor? weights=None, SymInt minlength=0) -> Tensor
static C10_NOINLINE c10::TypedOperatorHandle<bincount::schema> create_bincount_typed_handle() {
  return c10::Dispatcher::singleton()
      .findSchemaOrThrow(bincount::name, bincount::overload_name)
      .typed<bincount::schema>();
}

// aten::bincount(Tensor self, Tensor? weights=None, SymInt minlength=0) -> Tensor
at::Tensor bincount::call(const at::Tensor & self, const ::std::optional<at::Tensor> & weights, c10::SymInt minlength) {

    static auto op = create_bincount_typed_handle();
    return op.call(self, weights, minlength);
}

namespace {

at::Tensor wrapper_CPU__bincount(const at::Tensor & self, const ::std::optional<at::Tensor> & weights, c10::SymInt minlength) {
    // No device check

  // DeviceGuard omitted
  return at::native::_bincount_cpu(self, weights, minlength.guard_int(__FILE__, __LINE__));
}

} // anonymous namespace

TORCH_LIBRARY_IMPL(aten, CPU, m) {
    m.impl("bincount",
TORCH_FN(wrapper_CPU__bincount));
    // ... 同一分片里其他几百个 CPU kernel
}

#define TORCH_FN_TYPE(func)                                           \
  ::c10::CompileTimeFunctionPointer<                                  \
      std::remove_pointer_t<std::remove_reference_t<decltype(func)>>, \
      func>
#define TORCH_FN(func) TORCH_FN_TYPE(func)()

// NB: TORCH_LIBRARY_IMPL must be in an anonymous namespace to avoid
// ambiguity with conflicting identifiers that may have been defined in
// at namespace already.
namespace {

${dispatch_anonymous_definitions}

${static_init_dispatch_registrations}

} // anonymous namespace

TORCH_API at::Tensor _bincount_cpu(const at::Tensor & self, const ::std::optional<at::Tensor> & weights={}, int64_t minlength=0);

Library& Library::_def(c10::FunctionSchema&& schema, c10::OperatorName* out_name, const std::vector<at::Tag>& tags, _RegisterOrVerify rv) & {
  TORCH_CHECK(kind_ == DEF || kind_ == FRAGMENT,
    DEF_PRELUDE,
    "Cannot define an operator inside of a ", toString(kind_), " block.  "
    "All def()s should be placed in the (unique) TORCH_LIBRARY block for their namespace.  ",
    ERROR_CONTEXT
  );
  TORCH_INTERNAL_ASSERT(ns_.has_value(), ERROR_CONTEXT);
  TORCH_INTERNAL_ASSERT(!dispatch_key_.has_value(), ERROR_CONTEXT);
  auto ns_opt = schema.getNamespace();
  if (ns_opt.has_value()) {
    // Note [Redundancy in registration code is OK]
    // ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
    // In an earlier version of this code, I made it an error to explicitly
    // specify the namespace, even when the namespaces match.  I've decided
    // to relax this constraint because sometimes we code generate registrations
    // and you cannot conveniently tell what the enclosing context will be;
    // ...
    TORCH_CHECK(*ns_opt == *ns_,
      "Explicitly provided namespace (", *ns_opt, ") in schema string "
      "does not match namespace of enclosing ", toString(kind_), " block (", *ns_, ").  "
      "Move this definition to the (unique) TORCH_LIBRARY block corresponding to this namespace "
      "(and consider deleting the namespace from your schema string.)  ",
      ERROR_CONTEXT
    );
  } else {
    bool b = schema.setNamespaceIfNotSet(ns_->c_str());
    TORCH_INTERNAL_ASSERT(b, ERROR_CONTEXT);
  }
  // ...
      registrars_.emplace_back(
        c10::Dispatcher::singleton().registerDef(
          std::move(schema),
          debugString(file_, line_),
          tags
        )
      );
  // ...
  return *this;
}

#define DEF_PRELUDE "def(\"", schema.operator_name(), "\"): "

// Note: overwrite the external definition for sharing same name between
// libraries use different ISAs.
#define TORCH_EXTENSION_NAME _C

TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {

  ops.def("silu_and_mul(Tensor! out, Tensor input) -> ()");
  ops.impl("silu_and_mul", torch::kCPU, &silu_and_mul);

REGISTER_EXTENSION(TORCH_EXTENSION_NAME)

// minic10/macros/Macros.h
#pragma once

#include <cstdint>
#include <exception>
#include <sstream>
#include <string>

// ---------------------------------------------------------------------------
// 1. 可见性：对照 torch/headeronly/macros/Export.h 的 C10_EXPORT / C10_API
// ---------------------------------------------------------------------------
#ifdef _WIN32
#define MINI_EXPORT __declspec(dllexport)
#define MINI_IMPORT __declspec(dllimport)
#define MINI_HIDDEN
#else
#if defined(__GNUC__)  // gcc 和 clang 都定义 __GNUC__
#define MINI_EXPORT __attribute__((__visibility__("default")))
#define MINI_HIDDEN __attribute__((__visibility__("hidden")))
#else
#define MINI_EXPORT
#define MINI_HIDDEN
#endif
#define MINI_IMPORT MINI_EXPORT
#endif

// 编译 libminic10 本体时由构建系统定义 MINIC10_BUILD_MAIN_LIB（对应 C10_BUILD_MAIN_LIB）
#ifdef MINIC10_BUILD_MAIN_LIB
#define MINI_API MINI_EXPORT
#else
#define MINI_API MINI_IMPORT
#endif

// ---------------------------------------------------------------------------
// 2. 预处理器工具：字符串化、拼接、唯一标识符（对照 C10_STRINGIZE / C10_CONCATENATE / C10_UID）
// ---------------------------------------------------------------------------
#define MINI_STRINGIZE_IMPL(x) #x
#define MINI_STRINGIZE(x) MINI_STRINGIZE_IMPL(x)

#define MINI_CONCATENATE_IMPL(a, b) a##b
#define MINI_CONCATENATE(a, b) MINI_CONCATENATE_IMPL(a, b)

#ifdef __COUNTER__
#define MINI_UID __COUNTER__
#else
#define MINI_UID __LINE__
#endif
#define MINI_ANONYMOUS_VARIABLE(str) MINI_CONCATENATE(str, MINI_UID)

// ---------------------------------------------------------------------------
// 3. 分支预测提示（对照 C10_LIKELY / C10_UNLIKELY）
// ---------------------------------------------------------------------------
#if defined(__GNUC__) || defined(__clang__)
#define MINI_LIKELY(expr) (__builtin_expect(static_cast<bool>(expr), 1))
#define MINI_UNLIKELY(expr) (__builtin_expect(static_cast<bool>(expr), 0))
#else
#define MINI_LIKELY(expr) (expr)
#define MINI_UNLIKELY(expr) (expr)
#endif

// ---------------------------------------------------------------------------
// 4. MINI_CHECK：对照 c10/util/Exception.h 的 TORCH_CHECK
// ---------------------------------------------------------------------------
namespace minic10 {

struct SourceLocation {
  const char* function;
  const char* file;
  uint32_t line;
};

// MINI_API 不只导出函数：也让 typeinfo/vtable 具有默认可见性，
// 否则跨 .so 边界 catch (const Error&) 会失败（RTTI 按类型信息地址比较）。
class MINI_API Error : public std::exception {
  std::string msg_;

 public:
  Error(SourceLocation loc, const std::string& msg) {
    std::ostringstream ss;
    ss << msg << "\nException raised from " << loc.function << " at " << loc.file << ":"
       << loc.line;
    msg_ = ss.str();
  }
  const char* what() const noexcept override { return msg_.c_str(); }
};

namespace detail {

// 把任意多个可以 << 的参数拼成一个 std::string（对照 c10::str）
template <typename... Args>
inline std::string str(const Args&... args) {
  std::ostringstream ss;
  (ss << ... << args);  // C++17 折叠表达式
  return ss.str();
}

// 三个重载，对照 c10::detail::torchCheckMsgImpl：
//   没有用户消息       -> 用默认消息（const char*，零开销）
//   只有一个字符串字面量 -> 直接用它（不构造 ostringstream）
//   其他               -> 拼接
inline const char* checkMsg(const char* default_msg) { return default_msg; }
inline const char* checkMsg(const char* /*default_msg*/, const char* user_msg) { return user_msg; }
template <typename... Args>
inline std::string checkMsg(const char* /*default_msg*/, const Args&... args) {
  return str(args...);
}

[[noreturn]] inline void checkFail(const char* func, const char* file, uint32_t line,
                                   const std::string& msg) {
  throw Error({func, file, line}, msg);
}
[[noreturn]] inline void checkFail(const char* func, const char* file, uint32_t line,
                                   const char* msg) {
  throw Error({func, file, line}, std::string(msg));
}

}  // namespace detail
}  // namespace minic10

// 用法：MINI_CHECK(cond);  MINI_CHECK(cond, "expected ", x, " got ", y);
// 消息参数只在 cond 为假时才求值 —— 因为它们在 if 的花括号里。
#define MINI_CHECK(cond, ...)                                                           \
  if (MINI_UNLIKELY(!(cond))) {                                                         \
    ::minic10::detail::checkFail(                                                       \
        __func__, __FILE__, static_cast<uint32_t>(__LINE__),                            \
        ::minic10::detail::checkMsg("Expected " #cond " to be true, but got false.",    \
                                    ##__VA_ARGS__));                                    \
  }

#define MINI_INTERNAL_ASSERT(cond, ...)                                                 \
  if (MINI_UNLIKELY(!(cond))) {                                                         \
    ::minic10::detail::checkFail(                                                       \
        __func__, __FILE__, static_cast<uint32_t>(__LINE__),                            \
        ::minic10::detail::str(#cond " INTERNAL ASSERT FAILED at " __FILE__ ":"         \
                               MINI_STRINGIZE(__LINE__) ". ", ##__VA_ARGS__));          \
  }

// check_demo.cpp
#include <cstdio>
#include "minic10/macros/Macros.h"
int calls = 0;
int expensive() { ++calls; return 42; }
void f(int x) {
  MINI_CHECK(x > 0, "x must be positive, got ", x, " (expensive=", expensive(), ")");
  MINI_CHECK(x < 100, "single literal message");
  MINI_CHECK(x != 7);
}
int main() {
  f(1); std::printf("ok, expensive() called %d times\n", calls);
  for (int v : {-3, 200, 7}) {
    try { f(v); } catch (const minic10::Error& e) { std::printf("--\n%s\n", e.what()); }
  }
  std::printf("expensive() called %d times\n", calls);
}

if ((__builtin_expect(static_cast<bool>(!(x > 0)), 0))) {
  ::minic10::detail::checkFail(
      __func__, "check_demo.cpp", static_cast<uint32_t>(6),
      ::minic10::detail::checkMsg("Expected " "x > 0" " to be true, but got false.",
                                  "x must be positive, got ", x, " (expensive=", expensive(), ")"));
};

// minic10/dispatch/Dispatcher.h（节选）
class MINI_API Dispatcher final {
 public:
  // 对照 c10::Dispatcher::realSingleton / singleton：
  // 真正的 static 对象在 Dispatcher.cpp 里，只有一份；
  // 头文件里的 singleton() 只缓存一个引用，避免函数调用开销。
  static Dispatcher& realSingleton();
  static Dispatcher& singleton() {
    static Dispatcher& s = realSingleton();
    return s;
  }

  void registerLibrary(const std::string& ns, std::string debug);
  OperatorEntry& registerDef(const std::string& qualname, std::string debug);
  void registerImpl(const std::string& qualname, DispatchKey key, KernelFunction fn, std::string debug);

  OperatorHandle findOpOrThrow(const std::string& qualname) const;
  std::vector<std::string> registeredOps() const;

 private:
  Dispatcher() = default;
  OperatorEntry& findOrRegisterName_(const std::string& qualname);

  // unique_ptr 保证 OperatorEntry 地址稳定，OperatorHandle 可以长期持有指针
  std::unordered_map<std::string, std::unique_ptr<OperatorEntry>> operators_;
  std::unordered_map<std::string, std::string> libraries_;   // ns -> debug
};

// minic10/dispatch/Dispatcher.cpp
#include "minic10/dispatch/Dispatcher.h"

namespace minic10 {

// 函数内静态对象：第一次调用时构造，与翻译单元的静态初始化顺序无关。
Dispatcher& Dispatcher::realSingleton() {
  static Dispatcher _singleton;
  return _singleton;
}

void Dispatcher::registerLibrary(const std::string& ns, std::string debug) {
  auto found = libraries_.find(ns);
  MINI_CHECK(found == libraries_.end(),
             "Only a single MINI_LIBRARY can be used to register the namespace ", ns,
             "; please put all of your definitions in a single MINI_LIBRARY block. "
             "If you were trying to specify implementations, consider using MINI_LIBRARY_IMPL "
             "(which allows for multiple definitions). "
             "Previous registration: ", found->second, "; latest registration: ", debug);
  libraries_.emplace(ns, std::move(debug));
}

OperatorEntry& Dispatcher::findOrRegisterName_(const std::string& qualname) {
  auto it = operators_.find(qualname);
  if (it == operators_.end()) {
    auto entry = std::make_unique<OperatorEntry>();
    entry->name = qualname;
    it = operators_.emplace(qualname, std::move(entry)).first;
  }
  return *it->second;
}

OperatorEntry& Dispatcher::registerDef(const std::string& qualname, std::string debug) {
  OperatorEntry& op = findOrRegisterName_(qualname);
  MINI_CHECK(!op.has_schema, "Tried to register an operator (", qualname,
             ") with the same name and overload name multiple times. ",
             "Previous: ", op.def_debug, "; latest: ", debug);
  op.has_schema = true;
  op.def_debug = std::move(debug);
  return op;
}

void Dispatcher::registerImpl(const std::string& qualname, DispatchKey key, KernelFunction fn,
                              std::string debug) {
  // 注意：impl 可以先于 def 到达（静态初始化顺序不确定），所以这里也 find-or-register
  OperatorEntry& op = findOrRegisterName_(qualname);
  auto idx = static_cast<size_t>(key);
  MINI_CHECK(!op.kernels[idx].isValid(), "Overriding a previously registered kernel for ", qualname,
             " at DispatchKey ", toString(key), ". Previous: ", op.kernel_debug[idx],
             "; latest: ", debug);
  op.kernels[idx] = fn;
  op.kernel_debug[idx] = std::move(debug);
}

OperatorHandle Dispatcher::findOpOrThrow(const std::string& qualname) const {
  auto it = operators_.find(qualname);
  MINI_CHECK(it != operators_.end() && it->second->has_schema,
             "Could not find schema for ", qualname,
             ". Is the library that registers it linked in (and not dropped as an unreferenced "
             "object file of a static library)?");
  return OperatorHandle(it->second.get());
}

// ...

}  // namespace minic10

struct OperatorEntry {
  std::string name;           // "minic10::add"
  bool has_schema = false;    // def() 过了没有；impl 可能先于 def 到达
  std::string def_debug;      // "registered at add.cpp:12"
  std::array<KernelFunction, static_cast<size_t>(DispatchKey::NumKeys)> kernels{};
  std::array<std::string, static_cast<size_t>(DispatchKey::NumKeys)> kernel_debug{};

  const KernelFunction& lookup(DispatchKey k) const { return kernels[static_cast<size_t>(k)]; }
};

// minic10/library.h：对照 torch/library.h
#pragma once
#include <cstdint>
#include <optional>
#include <string>
#include "minic10/core/DispatchKey.h"
#include "minic10/dispatch/Dispatcher.h"
#include "minic10/macros/Macros.h"

namespace minic10 {

// 对照 torch::Library：一个宏块对应一个 Library 对象，
// 它知道自己的命名空间、DispatchKey（IMPL 才有）和出生地（文件、行号）。
class Library final {
 public:
  enum Kind { DEF, IMPL, FRAGMENT };

  Library(Kind kind, std::string ns, std::optional<DispatchKey> key, const char* file, uint32_t line)
      : kind_(kind), ns_(std::move(ns)), key_(key), file_(file), line_(line) {
    if (kind_ == DEF) {
      Dispatcher::singleton().registerLibrary(ns_, debugString());
    }
    MINI_CHECK(kind_ == IMPL || !key_.has_value(), "MINI_LIBRARY blocks cannot have a DispatchKey ",
               errorContext());
  }
  Library(const Library&) = delete;
  Library& operator=(const Library&) = delete;

  // def("add")：登记 schema "ns::add"。mini-c10 没有 schema 解析器，只记名字。
  Library& def(const char* name) {
    MINI_CHECK(kind_ == DEF || kind_ == FRAGMENT, "def(\"", name,
               "\"): Cannot define an operator inside of a MINI_LIBRARY_IMPL block. ", errorContext());
    Dispatcher::singleton().registerDef(qualify(name), debugString());
    return *this;
  }

  // impl("add", &add_cpu)：把函数指针登记到本块的 DispatchKey 下
  template <class Return, class... Args>
  Library& impl(const char* name, Return (*fn)(Args...)) {
    MINI_CHECK(kind_ == IMPL && key_.has_value(), "impl(\"", name,
               "\"): must be used inside a MINI_LIBRARY_IMPL block with a DispatchKey. ", errorContext());
    Dispatcher::singleton().registerImpl(qualify(name), *key_, KernelFunction::makeFromUnboxed(fn),
                                         debugString());
    return *this;
  }

 private:
  std::string qualify(const char* name) const { return ns_ + "::" + name; }
  std::string debugString() const { return detail::str("registered at ", file_, ":", line_); }
  std::string errorContext() const {
    return detail::str("(Error occurred while processing ", kindName(), " block at ", file_, ":", line_, ")");
  }
  const char* kindName() const {
    switch (kind_) {
      case DEF: return "MINI_LIBRARY";
      case IMPL: return "MINI_LIBRARY_IMPL";
      case FRAGMENT: return "MINI_LIBRARY_FRAGMENT";
    }
    return "(unknown)";
  }

  Kind kind_;
  std::string ns_;
  std::optional<DispatchKey> key_;
  const char* file_;
  uint32_t line_;
};

namespace detail {
// 对照 torch::detail::TorchLibraryInit：一个静态对象，构造函数里
// 先构造 Library，再调用用户写的函数体。它的唯一作用就是"在加载时被构造"。
class LibraryInit final {
  using InitFn = void(Library&);
  Library lib_;
 public:
  LibraryInit(Library::Kind kind, InitFn* fn, const char* ns, std::optional<DispatchKey> key,
              const char* file, uint32_t line)
      : lib_(kind, ns, key, file, line) {
    fn(lib_);
  }
};
}  // namespace detail

}  // namespace minic10

// MINI_LIBRARY(minic10, m) { m.def("add"); }
// 展开为：一个 static 函数的声明、一个 static const LibraryInit 对象、以及该函数的定义头。
// 用户写的 { ... } 就是函数体。
#define MINI_LIBRARY(ns, m)                                                       \
  static void MINI_LIBRARY_init_##ns(::minic10::Library&);                        \
  static const ::minic10::detail::LibraryInit MINI_LIBRARY_static_init_##ns(      \
      ::minic10::Library::DEF, &MINI_LIBRARY_init_##ns, MINI_STRINGIZE(ns),       \
      std::nullopt, __FILE__, __LINE__);                                          \
  void MINI_LIBRARY_init_##ns(::minic10::Library& m)

// 同一命名空间可以有多个 FRAGMENT，所以名字里要拼一个唯一 id（__COUNTER__）
#define MINI_LIBRARY_FRAGMENT(ns, m) _MINI_LIBRARY_FRAGMENT(ns, m, MINI_UID)
#define _MINI_LIBRARY_FRAGMENT(ns, m, uid)                                                    \
  static void MINI_CONCATENATE(MINI_LIBRARY_FRAGMENT_init_##ns##_, uid)(::minic10::Library&); \
  static const ::minic10::detail::LibraryInit MINI_CONCATENATE(                               \
      MINI_LIBRARY_FRAGMENT_static_init_##ns##_, uid)(                                        \
      ::minic10::Library::FRAGMENT,                                                           \
      &MINI_CONCATENATE(MINI_LIBRARY_FRAGMENT_init_##ns##_, uid), MINI_STRINGIZE(ns),         \
      std::nullopt, __FILE__, __LINE__);                                                      \
  void MINI_CONCATENATE(MINI_LIBRARY_FRAGMENT_init_##ns##_, uid)(::minic10::Library& m)

// MINI_LIBRARY_IMPL(minic10, CPU, m) { m.impl("add", &add_cpu); }
#define MINI_LIBRARY_IMPL(ns, k, m) _MINI_LIBRARY_IMPL(ns, k, m, MINI_UID)
#define _MINI_LIBRARY_IMPL(ns, k, m, uid)                                                       \
  static void MINI_CONCATENATE(MINI_LIBRARY_IMPL_init_##ns##_##k##_, uid)(::minic10::Library&); \
  static const ::minic10::detail::LibraryInit MINI_CONCATENATE(                                 \
      MINI_LIBRARY_IMPL_static_init_##ns##_##k##_, uid)(                                        \
      ::minic10::Library::IMPL,                                                                 \
      &MINI_CONCATENATE(MINI_LIBRARY_IMPL_init_##ns##_##k##_, uid), MINI_STRINGIZE(ns),         \
      std::make_optional(::minic10::DispatchKey::k), __FILE__, __LINE__);                       \
  void MINI_CONCATENATE(MINI_LIBRARY_IMPL_init_##ns##_##k##_, uid)(::minic10::Library& m)

// minic10/ops/RegisterSchema.cpp：对照 torchgen 生成的 RegisterSchema.cpp
// 一个命名空间只能有一个 MINI_LIBRARY 块，所有 def 放在这里。
#include "minic10/library.h"

MINI_LIBRARY(minic10, m) {
  m.def("add");
  m.def("mul");
}

static void MINI_LIBRARY_init_minic10(::minic10::Library&);
static const ::minic10::detail::LibraryInit MINI_LIBRARY_static_init_minic10(
    ::minic10::Library::DEF, &MINI_LIBRARY_init_minic10, "minic10",
    std::nullopt, "minic10/ops/RegisterSchema.cpp", 5);
void MINI_LIBRARY_init_minic10(::minic10::Library& m) {
  m.def("add");
  m.def("mul");
}

// minic10/ops/add.cpp：第 3 篇的 CPU kernel、第 4 篇的 Meta kernel，本篇改成自注册。
// 这个文件里没有任何符号被别的翻译单元引用。
#include "minic10/core/Dispatch.h"
#include "minic10/core/Tensor.h"
#include "minic10/library.h"
#include "minic10/macros/Macros.h"

#include <string>

namespace minic10 {
namespace {   // 内部链接：kernel 函数不导出（对照 RegisterCPU.cpp 的匿名命名空间）

std::string shapeStr(const Tensor& t) {
  std::string s = "[";
  for (size_t i = 0; i < t.sizes().size(); ++i) s += (i ? ", " : "") + std::to_string(t.sizes()[i]);
  return s + "]";
}

Tensor add_cpu(const Tensor& a, const Tensor& b) {
  MINI_CHECK(a.sizes() == b.sizes(), "add: shape mismatch, ", shapeStr(a), " vs ", shapeStr(b));
  MINI_CHECK(a.dtype() == b.dtype(), "add: dtype mismatch, ", toString(a.dtype()), " vs ", toString(b.dtype()));
  Tensor out = empty(a.sizes(), a.dtype());
  MINI_DISPATCH_FLOATING_TYPES(a.dtype(), "add_cpu", [&] {
    const scalar_t* pa = a.data_ptr<scalar_t>();
    const scalar_t* pb = b.data_ptr<scalar_t>();
    scalar_t* po = out.data_ptr<scalar_t>();
    for (int64_t i = 0; i < a.numel(); ++i) po[i] = pa[i] + pb[i];
  });
  return out;
}

Tensor add_meta(const Tensor& a, const Tensor& b) {
  MINI_CHECK(a.sizes() == b.sizes(), "add: shape mismatch");
  return empty_meta(a.sizes(), a.dtype());   // 只算形状，不碰数据
}

}  // namespace

// 静态注册：两个 static 对象在 libminic10 加载时构造，把函数指针写进 Dispatcher 的表
MINI_LIBRARY_IMPL(minic10, CPU, m) {
  m.impl("add", &add_cpu);
}

MINI_LIBRARY_IMPL(minic10, Meta, m) {
  m.impl("add", &add_meta);
}

}  // namespace minic10

static void MINI_LIBRARY_IMPL_init_minic10_CPU_0(::minic10::Library&);
static const ::minic10::detail::LibraryInit MINI_LIBRARY_IMPL_static_init_minic10_CPU_0(
    ::minic10::Library::IMPL, &MINI_LIBRARY_IMPL_init_minic10_CPU_0, "minic10",
    std::make_optional(::minic10::DispatchKey::CPU), "minic10/ops/add.cpp", 40);
void MINI_LIBRARY_IMPL_init_minic10_CPU_0(::minic10::Library& m) {
  m.impl("add", &add_cpu);
}

// minic10/ops/ops.h：公开的 C++ API。对照生成的 ATen/Functions.h：
// 这里不直接调用 kernel，而是经 Dispatcher 查表 —— 所以 ops/add.cpp 不需要被任何人引用。
#pragma once
#include "minic10/core/Tensor.h"
#include "minic10/dispatch/Dispatcher.h"

namespace minic10 {

inline Tensor add(const Tensor& a, const Tensor& b) {
  // 对照 Operators_*.cpp 里的 static auto op = create_add_Tensor_typed_handle();
  static const OperatorHandle op = Dispatcher::singleton().findOpOrThrow("minic10::add");
  return Dispatcher::singleton().call<Tensor, const Tensor&, const Tensor&>(op, a, b);
}

inline Tensor mul(const Tensor& a, const Tensor& b) {
  static const OperatorHandle op = Dispatcher::singleton().findOpOrThrow("minic10::mul");
  return Dispatcher::singleton().call<Tensor, const Tensor&, const Tensor&>(op, a, b);
}

}  // namespace minic10

// examples/main.cpp
#include <cstdio>
#include "minic10/core/Tensor.h"
#include "minic10/dispatch/Dispatcher.h"
#include "minic10/ops/ops.h"

using namespace minic10;

int main() {
  std::printf("registered ops:");
  for (const auto& n : Dispatcher::singleton().registeredOps()) std::printf(" %s", n.c_str());
  std::printf("\n");

  Tensor a = empty({2, 3}, ScalarType::Float);
  Tensor b = empty({2, 3}, ScalarType::Float);
  for (int64_t i = 0; i < a.numel(); ++i) { a.data_ptr<float>()[i] = float(i); b.data_ptr<float>()[i] = 10.f; }

  try {
    Tensor c = add(a, b);
    Tensor d = mul(a, b);
    std::printf("add[5]=%g mul[5]=%g\n", c.data_ptr<float>()[5], d.data_ptr<float>()[5]);

    Tensor am = empty_meta({4, 4}, ScalarType::Float);
    Tensor cm = add(am, am);
    std::printf("meta add: key=%s numel=%lld data=%p\n", toString(cm.key()), (long long)cm.numel(), cm.data_ptr<float>());

    Tensor bad = empty({3, 2}, ScalarType::Float);
    add(a, bad);
  } catch (const Error& e) {
    std::printf("caught minic10::Error:\n%s\n", e.what());
  }
  return 0;
}