Function gemv_4bit_inference

csrc/xpu_ops.cpp:34–53

template <typename T, int BITS>
void gemv_4bit_inference(
    int m, int n, int k, T* A, unsigned char* B, float* absmax, float* datatype, T* out, int lda, int ldb, int ldc,
    int blocksize, sycl::queue* stream
) {

    auto& queue = *stream;

    const size_t GROUP_SIZE = 128; // workgroup_size
    const size_t SUBG_SIZE = 32;   // subgroup_size
    const size_t NUM_PER_THREAD = GROUP_SIZE / SUBG_SIZE;
    size_t workgroup_num = (n + NUM_PER_THREAD - 1) / NUM_PER_THREAD;

    kgemv_4bit_inference<T, GROUP_SIZE, NUM_PER_THREAD, SUBG_SIZE, BITS> kfn(
        m, n, k, A, B, absmax, datatype, out, lda, ldb, ldc, blocksize
    );

    sycl_comp_kernel_submit<decltype(kfn), 1, SUBG_SIZE>(
        sycl::nd_range<1>(sycl::range<1>(GROUP_SIZE * workgroup_num), sycl::range<1>(GROUP_SIZE)), queue, kfn
    );
}

//==============================================================
//                   TEMPLATE DEFINITIONS