Логика mul_mat_vec_q в llama.cpp [закрыто] ⇐ C++

[Anonymous]

Я пытаюсь заменить квантование смешанного назначения Gemm cuda kernel в llama.cpp на мою реализацию. Для этого я должен понять, что DataRement и логика расчета в ядре mul_mat_vec_q .
Я пытался понять код, читая его, но не удалось, для большого числа неизвестных переменных и сложных параллелей. Vdr, blocks_per_row_x, blocks_per_col_y, blocks_per_iter, kbx, kby ...) среднее и как эта функция работает на Gemm. < /p>
template
// tell the compiler to use as many registers as it wants, see nwarps definition below
__launch_bounds__(calc_nwarps(ncols_y, get_device_table_id())*ggml_cuda_get_physical_warp_size(), 1)
static __global__ void mul_mat_vec_q(
const void * __restrict__ vx, const void * __restrict__ vy, float * __restrict__ dst,
const int ncols_x, const int nrows_x, const int nrows_y, const int nrows_dst) {

constexpr int qk = ggml_cuda_type_traits::qk;
constexpr int qi = ggml_cuda_type_traits::qi;
constexpr int vdr = get_vdr_mmvq(type);
constexpr mmvq_parameter_table_id table_id = get_device_table_id();
constexpr int nwarps = calc_nwarps(ncols_y, table_id);
constexpr int rows_per_cuda_block = calc_rows_per_block(ncols_y, table_id);
constexpr int warp_size = ggml_cuda_get_physical_warp_size();

constexpr vec_dot_q_cuda_t vec_dot_q_cuda = get_vec_dot_q_cuda(type);

const int tid = warp_size*threadIdx.y + threadIdx.x;
const int row0 = rows_per_cuda_block*blockIdx.x;
const int blocks_per_row_x = ncols_x / qk;
const int blocks_per_col_y = nrows_y / QK8_1;
constexpr int blocks_per_iter = vdr * nwarps * warp_size / qi;

// partial sum for each thread
float tmp[ncols_y][rows_per_cuda_block] = {{0.0f}};

const block_q8_1 * y = (const block_q8_1 *) vy;

for (int kbx = tid / (qi/vdr); kbx < blocks_per_row_x; kbx += blocks_per_iter) {
const int kby = kbx * (qk/QK8_1); // y block index that aligns with kbx

// x block quant index when casting the quants to int
const int kqs = vdr * (tid % (qi/vdr));

#pragma unroll
for (int j = 0; j < ncols_y; ++j) {
#pragma unroll
for (int i = 0; i < rows_per_cuda_block; ++i) {
tmp[j] += vec_dot_q_cuda(vx, &y[j*blocks_per_col_y + kby], (row0 + i)*blocks_per_row_x + kbx, kqs);
}
}
}

__shared__ float tmp_shared[nwarps-1 > 0 ? nwarps-1 : 1][ncols_y][rows_per_cuda_block][warp_size];
if (threadIdx.y > 0) {
#pragma unroll
for (int j = 0; j < ncols_y; ++j) {
#pragma unroll
for (int i = 0; i < rows_per_cuda_block; ++i) {
tmp_shared[threadIdx.y-1][j][threadIdx.x] = tmp[j];
}
}
}
__syncthreads();
if (threadIdx.y > 0) {
return;
}

// sum up partial sums and write back result
#pragma unroll
for (int j = 0; j < ncols_y; ++j) {
#pragma unroll
for (int i = 0; i < rows_per_cuda_block; ++i) {
#pragma unroll
for (int l = 0; l < nwarps-1; ++l) {
tmp[j] += tmp_shared[l][j][threadIdx.x];
}
tmp[j] = warp_reduce_sum(tmp[j]);
}

if (threadIdx.x < rows_per_cuda_block && (rows_per_cuda_block == 1 || row0 + threadIdx.x < (unsigned)nrows_dst)) {
dst[j*nrows_dst + row0 + threadIdx.x] = tmp[j][threadIdx.x];
}
}

GGML_UNUSED(nrows_x);
}


Подробнее здесь: https://stackoverflow.com/questions/795 ... -llama-cpp