Skip to content

Export C API from libinfiniops.so to eliminate vendor compiler requirement for downstream consumers #593

@bitzyz

Description

@bitzyz

Currently, downstream consumers (e.g., InfiniTensor) that call Operator::Call() must include device-specific headers, which forces them to replicate platform-specific build logic (vendor compiler detection, SDK paths, compile definitions) in their own CMake configuration.

Example: InfiniTensor's redundancy
InfiniTensor's adapter kernels (src/kernels/*.cc) call InfiniOps operators like:

#include "cuda/nvidia/add/kernel.h"  // transitively includes kernel.cuh with __global__
#include "cuda/moore/add/kernel.h"   // transitively includes MUSA-specific headers
#include "core/kernel.h"

infini::ops::Add::Call(handle, config, input, other, output);

Operator::Call() is a template function defined in operator.h. Its dispatch mechanism (DispatchFunc<ActiveDevices>) is resolved at compile time, requiring all device specializations to be visible at the call site:

// operator.h
template <typename... Args>
static auto Call(const Handle& handle, const Config& config, Args&&... args) {
    auto op = Make(config, args...);
    // Make() uses DispatchFunc<ActiveDevices<Key>>() which expands to:
    //   case kCpu:    return Operator<Add, kCpu>(...);     ← needs cpu/add/add.h
    //   case kNvidia: return Operator<Add, kNvidia>(...);  ← needs cuda/nvidia/add/kernel.h
    //   case kMoore:  return Operator<Add, kMoore>(...);   ← needs cuda/moore/add/kernel.h
    return (*op)(handle, args...);
}

Even though libinfiniops.so already contains compiled code for all device specializations, the template dispatch code is generated in the caller's translation unit, so the caller must include device-specific headers and use the vendor compiler.

Proposed Solution

Use C++ explicit template instantiation to compile Operator::Call() into libinfiniops.so, and extern template declarations to prevent downstream re-instantiation.

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Type

    No type
    No fields configured for issues without a type.

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions