pub struct MatvecKind {
pub name: &'static str,
pub module_name: &'static str,
pub fn_name: &'static str,
pub src: &'static str,
}Expand description
One quantized weight format the CUDA matvec path can consume.
The __global__ entry point named by Self::fn_name must have
the signature every kernel in this directory shares:
void <fn_name>(const unsigned char* weights, const float* x,
float* out, int rows, int row_bytes,
int n_blocks_per_row)one threadblock per output row, 256 threads striding the row’s
blocks, a tree reduction into out[row]. launch_matvec supplies
exactly that geometry, so a kernel written to a different one
returns wrong numbers rather than failing.
Fields§
§name: &'static strGGUF quant name – the key ferrox-core looks a kind up by,
which is QuantKind::name().
module_name: &'static strNVRTC module cache key. Must be unique per kind.
fn_name: &'static strThe __global__ entry point inside that module.
src: &'static strThe complete translation unit defining it.
Trait Implementations§
Source§impl Clone for MatvecKind
impl Clone for MatvecKind
Source§fn clone(&self) -> MatvecKind
fn clone(&self) -> MatvecKind
Returns a duplicate of the value. Read more
1.0.0 (const: unstable) · Source§fn clone_from(&mut self, source: &Self)
fn clone_from(&mut self, source: &Self)
Performs copy-assignment from
source. Read moreimpl Copy for MatvecKind
Auto Trait Implementations§
impl Freeze for MatvecKind
impl RefUnwindSafe for MatvecKind
impl Send for MatvecKind
impl Sync for MatvecKind
impl Unpin for MatvecKind
impl UnsafeUnpin for MatvecKind
impl UnwindSafe for MatvecKind
Blanket Implementations§
Source§impl<T> BorrowMut<T> for Twhere
T: ?Sized,
impl<T> BorrowMut<T> for Twhere
T: ?Sized,
Source§fn borrow_mut(&mut self) -> &mut T
fn borrow_mut(&mut self) -> &mut T
Mutably borrows from an owned value. Read more