Something went wrong. Try again.
A single-model (Bonsai-27B) inference engine for Intel Arc GPUs
Something went wrong. Try again.
C++
at extra-debug
1234567891011121314151617181920212223242526272829303132333435363738394041424344454647484950#include "qxmx_device.h"
#include <cstring>#include <numeric>#include <string>
namespace qx {
bool device_init(device* d, size_t budget_bytes, char* err, size_t errlen) { try { sycl::device dev = sycl::device(sycl::gpu_selector_v); d->q = sycl::queue(dev, sycl::property::queue::in_order{}); d->budget = budget_bytes; d->global_mem = dev.get_info<sycl::info::device::global_mem_size>(); d->max_slm = dev.get_info<sycl::info::device::local_mem_size>(); return true; } catch (const std::exception& e) { std::snprintf(err, errlen, "device init: %s", e.what()); return false; }}
std::uint64_t device::checksum(const void* p, size_t nbytes) { if (nbytes == 0) return 0; constexpr size_t LWS = 64; constexpr size_t CHUNK = 4096; /* bytes per work-group */ size_t ngroups = (nbytes + CHUNK - 1) / CHUNK; std::uint64_t* partials = chk_shared<std::uint64_t>(q, ngroups, "device.checksum.partials"); q.parallel_for( sycl::nd_range<1>(ngroups * LWS, sycl::range<1>(LWS)), [=](sycl::nd_item<1> it) { size_t g = it.get_group(0); size_t lid = it.get_local_id(0); size_t lws = it.get_local_range(0); size_t base = g * CHUNK; size_t end = base + CHUNK; if (end > nbytes) end = nbytes; std::uint64_t local = 0; const uint8_t* bp = (const uint8_t*)p; for (size_t i = base + lid; i < end; i += lws) local += bp[i]; std::uint64_t wg = sycl::reduce_over_group(it.get_group(), local, std::plus<>()); if (lid == 0) partials[g] = wg; }).wait(); std::uint64_t sum = 0; for (size_t i = 0; i < ngroups; i++) sum += partials[i]; sycl::free(partials, q); return sum;}
} // namespace qx