From 1679f8d84fd0d148f80e794df37586193c41c0ce Mon Sep 17 00:00:00 2001 From: leejet Date: Wed, 12 Aug 2026 00:10:04 +0800 Subject: [PATCH 1/2] ggml : add native INT8 convrot support --- ggml/include/ggml-rpc.h | 4 +- ggml/include/ggml.h | 21 + ggml/src/ggml-backend-meta.cpp | 6 + ggml/src/ggml-cpu/ggml-cpu.c | 163 +++++++ ggml/src/ggml-cpu/ggml-cpu.cpp | 50 ++ ggml/src/ggml-cpu/ops.cpp | 142 ++++++ ggml/src/ggml-cpu/ops.h | 2 + ggml/src/ggml-cuda/fwht.cu | 70 +++ ggml/src/ggml-cuda/fwht.cuh | 1 + ggml/src/ggml-cuda/ggml-cuda.cu | 451 ++++++++++++++++++ ggml/src/ggml-vulkan/ggml-vulkan.cpp | 181 ++++++- .../vulkan-shaders/mul_mat_i8_tensorwise.comp | 58 +++ .../mul_mat_i8_tensorwise_cm1.comp | 87 ++++ .../vulkan-shaders/quantize_i8_convrot.comp | 102 ++++ .../vulkan-shaders/vulkan-shaders-gen.cpp | 8 + ggml/src/ggml.c | 99 +++- tests/CMakeLists.txt | 1 + tests/test-backend-ops.cpp | 69 +++ tests/test-int8-convrot.cpp | 303 ++++++++++++ 19 files changed, 1812 insertions(+), 6 deletions(-) create mode 100644 ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_i8_tensorwise.comp create mode 100644 ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_i8_tensorwise_cm1.comp create mode 100644 ggml/src/ggml-vulkan/vulkan-shaders/quantize_i8_convrot.comp create mode 100644 tests/test-int8-convrot.cpp diff --git a/ggml/include/ggml-rpc.h b/ggml/include/ggml-rpc.h index 276aea00ea1..f07729d43bb 100644 --- a/ggml/include/ggml-rpc.h +++ b/ggml/include/ggml-rpc.h @@ -8,10 +8,10 @@ extern "C" { #define RPC_PROTO_MAJOR_VERSION 5 #define RPC_PROTO_MINOR_VERSION 0 -#define RPC_PROTO_PATCH_VERSION 0 +#define RPC_PROTO_PATCH_VERSION 1 #ifdef __cplusplus -static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION"); +static_assert(GGML_OP_COUNT == 103, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION"); #endif #define GGML_RPC_MAX_SERVERS 16 diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index 5cb49d0ee48..b0afa547408 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -590,6 +590,9 @@ extern "C" { GGML_OP_GLU, + GGML_OP_REGULAR_HADAMARD, + GGML_OP_QUANTIZE_I8_CONVROT, + GGML_OP_COUNT, }; @@ -1430,6 +1433,14 @@ extern "C" { struct ggml_tensor * a, struct ggml_tensor * b); + GGML_API struct ggml_tensor * ggml_mul_mat_i8_tensorwise( + struct ggml_context * ctx, + struct ggml_tensor * weight, + struct ggml_tensor * input, + struct ggml_tensor * weight_scale, + struct ggml_tensor * bias, + int convrot_group_size); + // change the precision of a matrix multiplication // set to GGML_PREC_F32 for higher precision (useful for phi-2) GGML_API void ggml_mul_mat_set_prec( @@ -1441,6 +1452,16 @@ extern "C" { struct ggml_tensor * a, enum ggml_op_hint hint); + GGML_API struct ggml_tensor * ggml_regular_hadamard( + struct ggml_context * ctx, + struct ggml_tensor * a, + int group_size); + + GGML_API struct ggml_tensor * ggml_quantize_i8_convrot( + struct ggml_context * ctx, + struct ggml_tensor * a, + int group_size); + // indirect matrix multiplication GGML_API struct ggml_tensor * ggml_mul_mat_id( struct ggml_context * ctx, diff --git a/ggml/src/ggml-backend-meta.cpp b/ggml/src/ggml-backend-meta.cpp index 7654ea1f30f..ca20b22a825 100644 --- a/ggml/src/ggml-backend-meta.cpp +++ b/ggml/src/ggml-backend-meta.cpp @@ -1009,6 +1009,12 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state( case GGML_OP_GLU: { split_state = handle_generic(src_ss, /*scalar_only =*/ false); } break; + case GGML_OP_REGULAR_HADAMARD: { + split_state = handle_per_row(src_ss); + } break; + case GGML_OP_QUANTIZE_I8_CONVROT: { + split_state = handle_generic(src_ss, /*scalar_only =*/ true); + } break; default: { GGML_ABORT("ggml op not implemented: %s", ggml_op_name(tensor->op)); split_state = {GGML_BACKEND_SPLIT_AXIS_UNKNOWN, {0}, {1}, 1}; diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 7918845cca0..1a25c2cdc18 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -1161,6 +1161,147 @@ void ggml_set_f32_nd(const struct ggml_tensor * tensor, int i0, int i1, int i2, // ggml_compute_forward_mul_mat +static void ggml_regular_hadamard_group_f32(float * values, int group_size) { + const float scale = 1.0f / sqrtf((float) group_size); + for (int i = 0; i < group_size; ++i) { + values[i] *= scale; + } + + for (int stride = 1; stride < group_size; stride *= 4) { + for (int base = 0; base < group_size; base += 4 * stride) { + for (int j = 0; j < stride; ++j) { + const int i0 = base + j; + const int i1 = i0 + stride; + const int i2 = i1 + stride; + const int i3 = i2 + stride; + const float a = values[i0]; + const float b = values[i1]; + const float c = values[i2]; + const float d = values[i3]; + values[i0] = a + b + c - d; + values[i1] = a + b - c + d; + values[i2] = a - b + c + d; + values[i3] = -a + b + c + d; + } + } + } +} + +static void ggml_compute_forward_mul_mat_i8_f32( + const struct ggml_compute_params * params, + struct ggml_tensor * dst) { + const struct ggml_tensor * src0 = dst->src[0]; + const struct ggml_tensor * src1 = dst->src[1]; + + GGML_ASSERT(src0->type == GGML_TYPE_I8); + GGML_ASSERT(src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_I8); + GGML_ASSERT(dst->type == GGML_TYPE_F32); + GGML_ASSERT(src0->ne[2] == 1 && src0->ne[3] == 1); + GGML_ASSERT(ggml_is_contiguous(src0)); + GGML_ASSERT(ggml_is_contiguous(src1)); + GGML_ASSERT(ggml_is_contiguous(dst)); + + const int64_t k = src0->ne[0]; + const int64_t n = src0->ne[1]; + const int64_t rows = ggml_nrows(dst); + const bool prequantized = src1->type == GGML_TYPE_I8; + const int64_t rows_padded = GGML_PAD(rows, 4); + const int64_t scale_rows = (rows * (int64_t) sizeof(float) + k - 1) / k; + const int convrot_group_size = ggml_get_op_params_i32(dst, 2); + const size_t qbytes = prequantized ? 0 : (size_t) ggml_nelements(src1) * sizeof(int8_t); + const size_t scale_offset = GGML_PAD(qbytes, sizeof(float)); + + GGML_ASSERT(src1->ne[0] == k); + GGML_ASSERT(!prequantized || (src1->op == GGML_OP_QUANTIZE_I8_CONVROT && src1->ne[1] == rows_padded + scale_rows)); + GGML_ASSERT(convrot_group_size == 0 || (convrot_group_size <= 256 && k % convrot_group_size == 0)); + GGML_ASSERT(prequantized || params->wsize >= scale_offset + (size_t) rows * sizeof(float)); + + int hadamard_size = convrot_group_size; + while (hadamard_size > 1 && hadamard_size % 4 == 0) { + hadamard_size /= 4; + } + GGML_ASSERT(hadamard_size == 0 || hadamard_size == 1); + + int8_t * qdata = prequantized ? (int8_t *) src1->data : (int8_t *) params->wdata; + float * scales = prequantized ? NULL : (float *) ((char *) params->wdata + scale_offset); + const float * packed_scales = prequantized ? (const float *) ((const int8_t *) src1->data + k * rows_padded) : NULL; + const float * src1_data = prequantized ? NULL : (const float *) src1->data; + + for (int64_t row = params->ith; !prequantized && row < rows; row += params->nth) { + const float * src_row = src1_data + row * k; + int8_t * qrow = qdata + row * k; + float amax = 0.0f; + if (convrot_group_size == 0) { + for (int64_t i = 0; i < k; ++i) { + amax = MAX(amax, fabsf(src_row[i])); + } + } else { + float values[256]; + for (int64_t group = 0; group < k; group += convrot_group_size) { + memcpy(values, src_row + group, (size_t) convrot_group_size * sizeof(float)); + ggml_regular_hadamard_group_f32(values, convrot_group_size); + for (int i = 0; i < convrot_group_size; ++i) { + amax = MAX(amax, fabsf(values[i])); + } + } + } + + const float row_scale = amax / 127.0f; + scales[row] = row_scale; + if (row_scale == 0.0f) { + memset(qrow, 0, (size_t) k); + continue; + } + + const float inv_scale = 1.0f / row_scale; + if (convrot_group_size == 0) { + for (int64_t i = 0; i < k; ++i) { + int value = (int) lrintf(src_row[i] * inv_scale); + value = MAX(-127, MIN(127, value)); + qrow[i] = (int8_t) value; + } + } else { + float values[256]; + for (int64_t group = 0; group < k; group += convrot_group_size) { + memcpy(values, src_row + group, (size_t) convrot_group_size * sizeof(float)); + ggml_regular_hadamard_group_f32(values, convrot_group_size); + for (int i = 0; i < convrot_group_size; ++i) { + int value = (int) lrintf(values[i] * inv_scale); + value = MAX(-127, MIN(127, value)); + qrow[group + i] = (int8_t) value; + } + } + } + } + + if (!prequantized) { + ggml_barrier(params->threadpool); + } + + float * dst_data = (float *) dst->data; + const int64_t output_elements = n * rows; + for (int64_t index = params->ith; index < output_elements; index += params->nth) { + const int64_t output_row = index / n; + const int64_t i01 = index - output_row * n; + const int8_t * weight_row = (const int8_t *) ((const char *) src0->data + i01 * src0->nb[1]); + const int8_t * activation_row = prequantized ? (const int8_t *) ((const char *) src1->data + output_row * src1->nb[1]) : qdata + output_row * k; + + int32_t sum = 0; + for (int64_t i = 0; i < k; ++i) { + sum += (int32_t) weight_row[i] * (int32_t) activation_row[i]; + } + const float activation_scale = prequantized ? packed_scales[output_row] : scales[output_row]; + float value = (float) sum * activation_scale; + if (dst->src[2] != NULL) { + value *= ((const float *) dst->src[2]->data)[i01]; + } + if (dst->src[3] != NULL) { + value += ((const float *) dst->src[3]->data)[i01]; + } + dst_data[index] = value; + } +} + static void ggml_compute_forward_mul_mat_one_chunk( const struct ggml_compute_params * params, struct ggml_tensor * dst, @@ -1258,6 +1399,13 @@ void ggml_compute_forward_mul_mat( const struct ggml_tensor * src0 = dst->src[0]; const struct ggml_tensor * src1 = dst->src[1]; + if (src0->type == GGML_TYPE_I8 && + (src1->type == GGML_TYPE_F32 || + (src1->type == GGML_TYPE_I8 && src1->op == GGML_OP_QUANTIZE_I8_CONVROT))) { + ggml_compute_forward_mul_mat_i8_f32(params, dst); + return; + } + const int32_t hint = ggml_get_op_params_i32(dst, 1); if (hint == GGML_HINT_SRC0_IS_HADAMARD && !params->use_ref) { ggml_compute_forward_fwht(params, dst); @@ -2032,6 +2180,14 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm { ggml_compute_forward_glu(params, tensor); } break; + case GGML_OP_REGULAR_HADAMARD: + { + ggml_compute_forward_regular_hadamard(params, tensor); + } break; + case GGML_OP_QUANTIZE_I8_CONVROT: + { + ggml_compute_forward_quantize_i8_convrot(params, tensor); + } break; case GGML_OP_GET_REL_POS: { ggml_compute_forward_get_rel_pos(params, tensor); @@ -2330,6 +2486,8 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) { case GGML_OP_MUL_MAT: case GGML_OP_MUL_MAT_ID: case GGML_OP_OUT_PROD: + case GGML_OP_REGULAR_HADAMARD: + case GGML_OP_QUANTIZE_I8_CONVROT: { n_tasks = n_threads; } break; @@ -2847,6 +3005,11 @@ struct ggml_cplan ggml_graph_plan( } break; case GGML_OP_MUL_MAT: { + if (node->src[0]->type == GGML_TYPE_I8 && node->src[1]->type == GGML_TYPE_F32) { + const size_t qbytes = (size_t) ggml_nelements(node->src[1]) * sizeof(int8_t); + cur = GGML_PAD(qbytes, sizeof(float)) + (size_t) ggml_nrows(node->src[1]) * sizeof(float); + break; + } const enum ggml_type vec_dot_type = type_traits_cpu[node->src[0]->type].vec_dot_type; if (node->src[1]->type != vec_dot_type) { diff --git a/ggml/src/ggml-cpu/ggml-cpu.cpp b/ggml/src/ggml-cpu/ggml-cpu.cpp index c0c9aa3cf09..68835eb2f65 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.cpp +++ b/ggml/src/ggml-cpu/ggml-cpu.cpp @@ -451,7 +451,57 @@ static bool ggml_backend_cpu_device_supports_op(ggml_backend_dev_t dev, const st op->type != GGML_TYPE_IQ1_S && op->type != GGML_TYPE_IQ1_M; // missing type_traits.from_float case GGML_OP_MUL_MAT: + if (src0->type == GGML_TYPE_I8) { + const ggml_tensor * weight_scale = op->src[2]; + const ggml_tensor * bias = op->src[3]; + const int convrot_group_size = ggml_get_op_params_i32(op, 2); + const ggml_tensor * packed_src = src1->src[0]; + const int64_t packed_rows = packed_src != nullptr ? GGML_PAD(ggml_nrows(packed_src), 4) : 0; + const int64_t packed_scale_rows = packed_src != nullptr ? (ggml_nrows(packed_src) * (int64_t) sizeof(float) + src0->ne[0] - 1) / src0->ne[0] : 0; + const bool packed_input = src1->type == GGML_TYPE_I8 && + src1->op == GGML_OP_QUANTIZE_I8_CONVROT && packed_src != nullptr && + src1->ne[0] == src0->ne[0] && src1->ne[1] == packed_rows + packed_scale_rows; + + int hadamard_size = convrot_group_size; + while (hadamard_size > 1 && hadamard_size % 4 == 0) { + hadamard_size /= 4; + } + + return (src1->type == GGML_TYPE_F32 || packed_input) && + op->type == GGML_TYPE_F32 && src0->ne[2] == 1 && src0->ne[3] == 1 && + ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && ggml_is_contiguous(op) && + (weight_scale == nullptr || + (weight_scale->type == GGML_TYPE_F32 && ggml_is_contiguous(weight_scale) && ggml_nelements(weight_scale) == src0->ne[1])) && + (bias == nullptr || + (bias->type == GGML_TYPE_F32 && ggml_is_contiguous(bias) && ggml_nelements(bias) == src0->ne[1])) && + (convrot_group_size == 0 || + (convrot_group_size <= 256 && src0->ne[0] % convrot_group_size == 0 && hadamard_size == 1)); + } return src1->type == GGML_TYPE_F32 || src1->type == ggml_get_type_traits_cpu(src0->type)->vec_dot_type; + case GGML_OP_QUANTIZE_I8_CONVROT: { + int group_size = ggml_get_op_params_i32(op, 0); + if (src0->type != GGML_TYPE_F32 || op->type != GGML_TYPE_I8 || + !ggml_is_contiguous(src0) || !ggml_is_contiguous(op) || + group_size <= 0 || group_size > 256 || src0->ne[0] % group_size != 0) { + return false; + } + while (group_size > 1 && group_size % 4 == 0) { + group_size /= 4; + } + return group_size == 1; + } + case GGML_OP_REGULAR_HADAMARD: { + int group_size = ggml_get_op_params_i32(op, 0); + if (src0->type != GGML_TYPE_F32 || op->type != GGML_TYPE_F32 || + !ggml_is_contiguous(src0) || !ggml_is_contiguous(op) || + group_size <= 0 || src0->ne[0] % group_size != 0) { + return false; + } + while (group_size > 1 && group_size % 4 == 0) { + group_size /= 4; + } + return group_size == 1; + } case GGML_OP_SOFT_MAX_BACK: { if (op->src[0]->type != GGML_TYPE_F32 || op->src[1]->type != GGML_TYPE_F32) { return false; diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 42ec809ce52..c5d11f67725 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -11919,6 +11919,148 @@ void ggml_compute_forward_fwht(const ggml_compute_params * params, ggml_tensor * } } +void ggml_compute_forward_regular_hadamard(const ggml_compute_params * params, ggml_tensor * dst) { + const ggml_tensor * src = dst->src[0]; + + GGML_ASSERT(src->type == GGML_TYPE_F32); + GGML_ASSERT(dst->type == GGML_TYPE_F32); + GGML_ASSERT(ggml_is_contiguous(src)); + GGML_ASSERT(ggml_is_contiguous(dst)); + + const int group_size = ggml_get_op_params_i32(dst, 0); + GGML_ASSERT(group_size > 0 && src->ne[0] % group_size == 0); + + const int64_t group_count = ggml_nelements(src) / group_size; + const int64_t groups_per_thread = (group_count + params->nth - 1) / params->nth; + const int64_t start_group = params->ith * groups_per_thread; + const int64_t end_group = MIN(start_group + groups_per_thread, group_count); + const float scale = 1.0f / sqrtf((float) group_size); + + const float * src_data = (const float *) src->data; + float * dst_data = (float *) dst->data; + + for (int64_t group = start_group; group < end_group; ++group) { + const float * src_group = src_data + group * group_size; + float * dst_group = dst_data + group * group_size; + + for (int i = 0; i < group_size; ++i) { + dst_group[i] = src_group[i] * scale; + } + + for (int stride = 1; stride < group_size; stride *= 4) { + for (int base = 0; base < group_size; base += 4 * stride) { + for (int j = 0; j < stride; ++j) { + const int i0 = base + j; + const int i1 = i0 + stride; + const int i2 = i1 + stride; + const int i3 = i2 + stride; + const float a = dst_group[i0]; + const float b = dst_group[i1]; + const float c = dst_group[i2]; + const float d = dst_group[i3]; + dst_group[i0] = a + b + c - d; + dst_group[i1] = a + b - c + d; + dst_group[i2] = a - b + c + d; + dst_group[i3] = -a + b + c + d; + } + } + } + } +} + +void ggml_compute_forward_quantize_i8_convrot(const ggml_compute_params * params, ggml_tensor * dst) { + const ggml_tensor * src = dst->src[0]; + + GGML_ASSERT(src->type == GGML_TYPE_F32); + GGML_ASSERT(dst->type == GGML_TYPE_I8); + GGML_ASSERT(ggml_is_contiguous(src)); + GGML_ASSERT(ggml_is_contiguous(dst)); + + const int group_size = ggml_get_op_params_i32(dst, 0); + const int64_t k = src->ne[0]; + const int64_t rows = ggml_nrows(src); + const int64_t rows_padded = GGML_PAD(rows, 4); + const int64_t scale_rows = (rows * (int64_t) sizeof(float) + k - 1) / k; + const float transform_scale = 1.0f / sqrtf((float) group_size); + GGML_ASSERT(group_size > 0 && group_size <= 256 && k % group_size == 0); + GGML_ASSERT(dst->ne[0] == k && dst->ne[1] == rows_padded + scale_rows); + + const float * src_data = (const float *) src->data; + char * dst_data = (char *) dst->data; + float * dst_scales = (float *) (dst_data + k * rows_padded); + + for (int64_t row = params->ith; row < rows; row += params->nth) { + const float * src_row = src_data + row * k; + int8_t * dst_row = (int8_t *) (dst_data + row * dst->nb[1]); + float amax = 0.0f; + float values[256]; + + for (int64_t group = 0; group < k; group += group_size) { + for (int i = 0; i < group_size; ++i) { + values[i] = src_row[group + i] * transform_scale; + } + for (int stride = 1; stride < group_size; stride *= 4) { + for (int base = 0; base < group_size; base += 4 * stride) { + for (int j = 0; j < stride; ++j) { + const int i0 = base + j; + const int i1 = i0 + stride; + const int i2 = i1 + stride; + const int i3 = i2 + stride; + const float a = values[i0]; + const float b = values[i1]; + const float c = values[i2]; + const float d = values[i3]; + values[i0] = a + b + c - d; + values[i1] = a + b - c + d; + values[i2] = a - b + c + d; + values[i3] = -a + b + c + d; + } + } + } + for (int i = 0; i < group_size; ++i) { + amax = MAX(amax, fabsf(values[i])); + } + } + + const float row_scale = amax / 127.0f; + const float inv_scale = row_scale == 0.0f ? 0.0f : 1.0f / row_scale; + dst_scales[row] = row_scale; + + for (int64_t group = 0; group < k; group += group_size) { + for (int i = 0; i < group_size; ++i) { + values[i] = src_row[group + i] * transform_scale; + } + for (int stride = 1; stride < group_size; stride *= 4) { + for (int base = 0; base < group_size; base += 4 * stride) { + for (int j = 0; j < stride; ++j) { + const int i0 = base + j; + const int i1 = i0 + stride; + const int i2 = i1 + stride; + const int i3 = i2 + stride; + const float a = values[i0]; + const float b = values[i1]; + const float c = values[i2]; + const float d = values[i3]; + values[i0] = a + b + c - d; + values[i1] = a + b - c + d; + values[i2] = a - b + c + d; + values[i3] = -a + b + c + d; + } + } + } + for (int i = 0; i < group_size; ++i) { + int value = row_scale == 0.0f ? 0 : (int) lrintf(values[i] * inv_scale); + value = MAX(-127, MIN(127, value)); + dst_row[group + i] = (int8_t) value; + } + } + } + + if (params->ith == 0 && rows_padded > rows) { + memset(dst_data + rows * dst->nb[1], 0, (size_t) (rows_padded - rows) * dst->nb[1]); + } +} + // ggml_compute_forward_lightning_indexer void ggml_compute_forward_lightning_indexer( diff --git a/ggml/src/ggml-cpu/ops.h b/ggml/src/ggml-cpu/ops.h index 4c1642a6760..66bcf913d6c 100644 --- a/ggml/src/ggml-cpu/ops.h +++ b/ggml/src/ggml-cpu/ops.h @@ -118,6 +118,8 @@ void ggml_compute_forward_cross_entropy_loss_back(const struct ggml_compute_para void ggml_compute_forward_opt_step_adamw(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_mul_mat(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_fwht(const struct ggml_compute_params * params, struct ggml_tensor * dst); +void ggml_compute_forward_regular_hadamard(const struct ggml_compute_params * params, struct ggml_tensor * dst); +void ggml_compute_forward_quantize_i8_convrot(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_opt_step_sgd(const struct ggml_compute_params * params, struct ggml_tensor * dst); #ifdef __cplusplus } diff --git a/ggml/src/ggml-cuda/fwht.cu b/ggml/src/ggml-cuda/fwht.cu index 184dc254c72..7a21ca0fc40 100644 --- a/ggml/src/ggml-cuda/fwht.cu +++ b/ggml/src/ggml-cuda/fwht.cu @@ -58,6 +58,44 @@ __global__ void fwht_cuda(const float * src, float * dst, const int64_t n_rows, } } +template +__global__ void regular_hadamard_cuda(const float * src, float * dst, const int64_t n_groups, const float scale) { + const int64_t group = blockIdx.x; + const int tid = threadIdx.x; + if (group >= n_groups) { + return; + } + + __shared__ float values[N]; + src += group * N; + dst += group * N; + + values[tid] = src[tid] * scale; + __syncthreads(); + +#pragma unroll + for (int stride = 1; stride < N; stride *= 4) { + if (tid < N / 4) { + const int base = (tid / stride) * 4 * stride + tid % stride; + const int i0 = base; + const int i1 = i0 + stride; + const int i2 = i1 + stride; + const int i3 = i2 + stride; + const float a = values[i0]; + const float b = values[i1]; + const float c = values[i2]; + const float d = values[i3]; + values[i0] = a + b + c - d; + values[i1] = a + b - c + d; + values[i2] = a - b + c + d; + values[i3] = -a + b + c + d; + } + __syncthreads(); + } + + dst[tid] = values[tid]; +} + bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst) { GGML_ASSERT(ggml_are_same_shape(src, dst)); if (!ggml_is_contiguous(src) || !ggml_is_contiguous(dst)) { @@ -99,3 +137,35 @@ bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src, return false; } } + +bool ggml_cuda_op_regular_hadamard(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst, int group_size) { + GGML_ASSERT(ggml_are_same_shape(src, dst)); + if (src->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32 || + !ggml_is_contiguous(src) || !ggml_is_contiguous(dst) || + group_size <= 0 || ggml_nelements(src) % group_size != 0) { + return false; + } + + const int64_t n_groups = ggml_nelements(src) / group_size; + const float scale = 1.0f / sqrtf((float) group_size); + const float * src_d = (const float *) src->data; + float * dst_d = (float *) dst->data; + cudaStream_t stream = ctx.stream(); + + switch (group_size) { + case 4: + regular_hadamard_cuda<4><<>>(src_d, dst_d, n_groups, scale); + return true; + case 16: + regular_hadamard_cuda<16><<>>(src_d, dst_d, n_groups, scale); + return true; + case 64: + regular_hadamard_cuda<64><<>>(src_d, dst_d, n_groups, scale); + return true; + case 256: + regular_hadamard_cuda<256><<>>(src_d, dst_d, n_groups, scale); + return true; + default: + return false; + } +} diff --git a/ggml/src/ggml-cuda/fwht.cuh b/ggml/src/ggml-cuda/fwht.cuh index cf3df94cafa..d4d28fad670 100644 --- a/ggml/src/ggml-cuda/fwht.cuh +++ b/ggml/src/ggml-cuda/fwht.cuh @@ -2,3 +2,4 @@ // Returns whether the Fast Walsh-Hadamard transform could be used. bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst); +bool ggml_cuda_op_regular_hadamard(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst, int group_size); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 1d4f4dfbd21..aa15717dbe3 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1616,6 +1616,387 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const } } +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) +__global__ void quantize_rowwise_i8_cuda(const float * src, int8_t * dst, float * scales, int64_t k, int64_t rows) { + const int64_t row = blockIdx.x; + const int tid = threadIdx.x; + if (row >= rows) { + return; + } + + __shared__ float maxima[256]; + const float * src_row = src + row * k; + int8_t * dst_row = dst + row * k; + float local_max = 0.0f; + for (int64_t i = tid; i < k; i += blockDim.x) { + local_max = fmaxf(local_max, fabsf(src_row[i])); + } + maxima[tid] = local_max; + __syncthreads(); + + for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) { + if (tid < stride) { + maxima[tid] = fmaxf(maxima[tid], maxima[tid + stride]); + } + __syncthreads(); + } + + const float scale = maxima[0] / 127.0f; + if (tid == 0) { + scales[row] = scale; + } + if (scale == 0.0f) { + for (int64_t i = tid; i < k; i += blockDim.x) { + dst_row[i] = 0; + } + return; + } + + const float inv_scale = 1.0f / scale; + for (int64_t i = tid; i < k; i += blockDim.x) { + int value = __float2int_rn(src_row[i] * inv_scale); + value = max(-127, min(127, value)); + dst_row[i] = (int8_t) value; + } +} + +__global__ void quantize_rowwise_i8_convrot_cuda( + const float * src, int8_t * dst, float * scales, int64_t k, int64_t rows, int64_t dst_stride) { + constexpr int group_size = 256; + constexpr int group_threads = group_size / 4; + constexpr int groups_per_wave = 1024 / group_threads; + + const int64_t row = blockIdx.x; + const int tid = threadIdx.x; + if (row >= rows) { + return; + } + + extern __shared__ float values[]; + __shared__ float maxima[32]; + const float * src_row = src + row * k; + int8_t * dst_row = dst + row * dst_stride; + + for (int64_t i = tid; i < k; i += blockDim.x) { + values[i] = src_row[i] * (1.0f / 16.0f); + } + __syncthreads(); + + const int group_lane = tid % group_threads; + const int wave_group = tid / group_threads; + const int64_t groups = k / group_size; + for (int64_t group_base = 0; group_base < groups; group_base += groups_per_wave) { + const int64_t group = group_base + wave_group; +#pragma unroll + for (int stride = 1; stride < group_size; stride *= 4) { + if (group < groups) { + float * group_values = values + group * group_size; + const int base = (group_lane / stride) * 4 * stride + group_lane % stride; + const int i0 = base; + const int i1 = i0 + stride; + const int i2 = i1 + stride; + const int i3 = i2 + stride; + const float a = group_values[i0]; + const float b = group_values[i1]; + const float c = group_values[i2]; + const float d = group_values[i3]; + group_values[i0] = a + b + c - d; + group_values[i1] = a + b - c + d; + group_values[i2] = a - b + c + d; + group_values[i3] = -a + b + c + d; + } + __syncthreads(); + } + } + + float local_max = 0.0f; + for (int64_t i = tid; i < k; i += blockDim.x) { + local_max = fmaxf(local_max, fabsf(values[i])); + } + for (int offset = 16; offset > 0; offset >>= 1) { + local_max = fmaxf(local_max, __shfl_down_sync(0xffffffff, local_max, offset)); + } + const int lane = tid & 31; + const int warp = tid >> 5; + if (lane == 0) { + maxima[warp] = local_max; + } + __syncthreads(); + + if (warp == 0) { + local_max = lane < blockDim.x / 32 ? maxima[lane] : 0.0f; + for (int offset = 16; offset > 0; offset >>= 1) { + local_max = fmaxf(local_max, __shfl_down_sync(0xffffffff, local_max, offset)); + } + if (lane == 0) { + maxima[0] = local_max / 127.0f; + if (scales != nullptr) { + scales[row] = maxima[0]; + } else { + *((float *) (dst_row + k)) = maxima[0]; + } + } + } + __syncthreads(); + + const float row_scale = maxima[0]; + const float inv_scale = row_scale == 0.0f ? 0.0f : 1.0f / row_scale; + for (int64_t i = tid; i < k; i += blockDim.x) { + int value = row_scale == 0.0f ? 0 : __float2int_rn(values[i] * inv_scale); + value = max(-127, min(127, value)); + dst_row[i] = (int8_t) value; + } +} + +__global__ void convrot_group_amax_i8_cuda( + const float * src, float * partial_max, int64_t k, int64_t groups, int64_t total_groups) { + const int64_t index = blockIdx.x; + const int tid = threadIdx.x; + if (index >= total_groups) { + return; + } + + __shared__ float values[256]; + __shared__ float maxima[2]; + const int64_t row = index / groups; + const int64_t group = index - row * groups; + const float * input = src + row * k + group * 256; + values[tid] = input[tid] * (1.0f / 16.0f); + values[tid + 64] = input[tid + 64] * (1.0f / 16.0f); + values[tid + 128] = input[tid + 128] * (1.0f / 16.0f); + values[tid + 192] = input[tid + 192] * (1.0f / 16.0f); + __syncthreads(); + +#pragma unroll + for (int stride = 1; stride < 256; stride *= 4) { + const int base = (tid / stride) * 4 * stride + tid % stride; + const int i0 = base; + const int i1 = i0 + stride; + const int i2 = i1 + stride; + const int i3 = i2 + stride; + const float a = values[i0]; + const float b = values[i1]; + const float c = values[i2]; + const float d = values[i3]; + values[i0] = a + b + c - d; + values[i1] = a + b - c + d; + values[i2] = a - b + c + d; + values[i3] = -a + b + c + d; + __syncthreads(); + } + + float value = fmaxf(fabsf(values[tid]), fabsf(values[tid + 64])); + value = fmaxf(value, fabsf(values[tid + 128])); + value = fmaxf(value, fabsf(values[tid + 192])); + for (int offset = 16; offset > 0; offset >>= 1) { + value = fmaxf(value, __shfl_down_sync(0xffffffff, value, offset)); + } + if ((tid & 31) == 0) { + maxima[tid >> 5] = value; + } + __syncthreads(); + if (tid == 0) { + partial_max[index] = fmaxf(maxima[0], maxima[1]); + } +} + +__global__ void reduce_convrot_row_amax_i8_cuda(const float * partial_max, float * scales, int64_t groups, int64_t rows) { + const int64_t row = blockIdx.x; + const int tid = threadIdx.x; + if (row >= rows) { + return; + } + + __shared__ float maxima[256]; + float value = 0.0f; + for (int64_t group = tid; group < groups; group += blockDim.x) { + value = fmaxf(value, partial_max[row * groups + group]); + } + maxima[tid] = value; + __syncthreads(); + for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) { + if (tid < stride) { + maxima[tid] = fmaxf(maxima[tid], maxima[tid + stride]); + } + __syncthreads(); + } + if (tid == 0) { + scales[row] = maxima[0] / 127.0f; + } +} + +__global__ void convrot_group_quantize_i8_cuda( + const float * src, int8_t * dst, const float * scales, int64_t k, int64_t groups, int64_t total_groups, int64_t dst_stride) { + const int64_t index = blockIdx.x; + const int tid = threadIdx.x; + if (index >= total_groups) { + return; + } + + __shared__ float values[256]; + const int64_t row = index / groups; + const int64_t group = index - row * groups; + const float * input = src + row * k + group * 256; + int8_t * output = dst + row * dst_stride + group * 256; + values[tid] = input[tid] * (1.0f / 16.0f); + values[tid + 64] = input[tid + 64] * (1.0f / 16.0f); + values[tid + 128] = input[tid + 128] * (1.0f / 16.0f); + values[tid + 192] = input[tid + 192] * (1.0f / 16.0f); + __syncthreads(); + +#pragma unroll + for (int stride = 1; stride < 256; stride *= 4) { + const int base = (tid / stride) * 4 * stride + tid % stride; + const int i0 = base; + const int i1 = i0 + stride; + const int i2 = i1 + stride; + const int i3 = i2 + stride; + const float a = values[i0]; + const float b = values[i1]; + const float c = values[i2]; + const float d = values[i3]; + values[i0] = a + b + c - d; + values[i1] = a + b - c + d; + values[i2] = a - b + c + d; + values[i3] = -a + b + c + d; + __syncthreads(); + } + + const float row_scale = scales[row]; + const float inv_scale = row_scale == 0.0f ? 0.0f : 1.0f / row_scale; + const int indices[4] = { tid, tid + 64, tid + 128, tid + 192 }; +#pragma unroll + for (int i = 0; i < 4; ++i) { + int value = row_scale == 0.0f ? 0 : __float2int_rn(values[indices[i]] * inv_scale); + value = max(-127, min(127, value)); + output[indices[i]] = (int8_t) value; + } + if (group == 0 && tid == 0 && dst_stride != k) { + *((float *) (dst + row * dst_stride + k)) = row_scale; + } +} + +__global__ void dequantize_i32_rows_cuda( + const int32_t * src, const float * scales, const float * weight_scales, const float * bias, float * dst, int64_t n, int64_t rows) { + const int64_t output = (int64_t) blockIdx.x * blockDim.x + threadIdx.x; + const int64_t row = blockIdx.y; + if (output < n && row < rows) { + const int64_t index = row * n + output; + const float activation_scale = scales[row]; + float value = (float) src[index] * activation_scale; + if (weight_scales != nullptr) { + value *= weight_scales[output]; + } + if (bias != nullptr) { + value += bias[output]; + } + dst[index] = value; + } +} + +static void ggml_cuda_quantize_i8_convrot( + ggml_backend_cuda_context & ctx, const float * src, int8_t * dst, float * scales, int64_t k, int64_t rows, int64_t rows_padded, int64_t dst_stride) { + cudaStream_t stream = ctx.stream(); + if (rows_padded > rows) { + CUDA_CHECK(cudaMemsetAsync(dst + dst_stride * rows, 0, (size_t) dst_stride * (rows_padded - rows), stream)); + } + + const size_t nbytes_shared = (size_t) k * sizeof(float); + const size_t smpbo = ggml_cuda_info().devices[ggml_cuda_get_device()].smpbo; + if (nbytes_shared + 128 <= smpbo) { + CUDA_SET_SHARED_MEMORY_LIMIT(quantize_rowwise_i8_convrot_cuda, smpbo - 128); + quantize_rowwise_i8_convrot_cuda<<>>(src, dst, scales, k, rows, dst_stride); + return; + } + + const int64_t groups = k / 256; + const int64_t total_groups = rows * groups; + ggml_cuda_pool_alloc partial_max(ctx.pool(), total_groups); + ggml_cuda_pool_alloc packed_scales(ctx.pool()); + float * scales_d = scales; + if (scales_d == nullptr) { + scales_d = packed_scales.alloc(rows); + } + convrot_group_amax_i8_cuda<<>>(src, partial_max.get(), k, groups, total_groups); + reduce_convrot_row_amax_i8_cuda<<>>(partial_max.get(), scales_d, groups, rows); + convrot_group_quantize_i8_cuda<<>>(src, dst, scales_d, k, groups, total_groups, dst_stride); +} + +static bool ggml_cuda_op_quantize_i8_convrot( + ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst, int group_size) { + const int64_t rows = ggml_nrows(src); + const int64_t rows_padded = GGML_PAD(rows, 4); + const int64_t scale_rows = (rows * (int64_t) sizeof(float) + src->ne[0] - 1) / src->ne[0]; + if (src->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_I8 || + !ggml_is_contiguous(src) || !ggml_is_contiguous(dst) || + group_size != 256 || src->ne[0] % group_size != 0 || + dst->ne[0] != src->ne[0] || dst->ne[1] != rows_padded + scale_rows) { + return false; + } + + int8_t * dst_data = (int8_t *) dst->data; + ggml_cuda_quantize_i8_convrot(ctx, (const float *) src->data, dst_data, (float *) (dst_data + src->ne[0] * rows_padded), src->ne[0], rows, rows_padded, src->ne[0]); + return true; +} + +static void ggml_cuda_mul_mat_i8( + ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + GGML_ASSERT(src0->type == GGML_TYPE_I8); + GGML_ASSERT(src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_I8); + GGML_ASSERT(dst->type == GGML_TYPE_F32); + GGML_ASSERT(src0->ne[2] == 1 && src0->ne[3] == 1); + GGML_ASSERT(ggml_is_contiguous(src0)); + GGML_ASSERT(ggml_is_contiguous(src1)); + GGML_ASSERT(ggml_is_contiguous(dst)); + + const int64_t k = src0->ne[0]; + const int64_t n = src0->ne[1]; + const int64_t rows = ggml_nrows(dst); + const bool prequantized = src1->type == GGML_TYPE_I8; + const int64_t rows_padded = GGML_PAD(rows, 4); + const int64_t scale_rows = (rows * (int64_t) sizeof(float) + k - 1) / k; + const int64_t qstride = k; + const int convrot_group_size = ggml_get_op_params_i32(dst, 2); + GGML_ASSERT(src1->ne[0] == k); + GGML_ASSERT(!prequantized || (src1->op == GGML_OP_QUANTIZE_I8_CONVROT && src1->ne[1] == rows_padded + scale_rows)); + GGML_ASSERT(convrot_group_size == 0 || convrot_group_size == 256); + + ggml_cuda_pool_alloc qdata(ctx.pool()); + ggml_cuda_pool_alloc scales(ctx.pool()); + cudaStream_t stream = ctx.stream(); + int8_t * qdata_d = prequantized ? (int8_t *) src1->data : qdata.alloc((size_t) k * rows_padded); + float * scales_d = prequantized ? (float *) ((int8_t *) src1->data + k * rows_padded) : scales.alloc(rows); + + if (prequantized) { + GGML_ASSERT(src1->op == GGML_OP_QUANTIZE_I8_CONVROT); + } else if (convrot_group_size == 0) { + if (rows_padded > rows) { + CUDA_CHECK(cudaMemsetAsync(qdata_d + k * rows, 0, (size_t) k * (rows_padded - rows), stream)); + } + quantize_rowwise_i8_cuda<<>>((const float *) src1->data, qdata_d, scales_d, k, rows); + } else { + ggml_cuda_quantize_i8_convrot(ctx, (const float *) src1->data, qdata_d, scales_d, k, rows, rows_padded, k); + } + + const float * weight_scales = dst->src[2] != nullptr ? (const float *) dst->src[2]->data : nullptr; + const float * bias = dst->src[3] != nullptr ? (const float *) dst->src[3]->data : nullptr; + ggml_cuda_pool_alloc accum(ctx.pool(), (size_t) n * rows_padded); + + const int32_t alpha = 1; + const int32_t beta = 0; + CUBLAS_CHECK(cublasSetStream(ctx.cublas_handle(), stream)); + CUBLAS_CHECK(cublasGemmEx(ctx.cublas_handle(), CUBLAS_OP_T, CUBLAS_OP_N, + (int) n, (int) rows_padded, (int) k, + &alpha, src0->data, CUDA_R_8I, (int) k, + qdata_d, CUDA_R_8I, (int) qstride, + &beta, accum.get(), CUDA_R_32I, (int) n, + CUBLAS_COMPUTE_32I, CUBLAS_GEMM_DEFAULT_TENSOR_OP)); + + const dim3 dequant_grid((n + 255) / 256, rows, 1); + dequantize_i32_rows_cuda<<>>(accum.get(), scales_d, weight_scales, bias, (float *) dst->data, n, rows); +} +#endif + static void ggml_cuda_mul_mat_cublas(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { ggml_type compute_type = src0->type; if (ggml_is_quantized(compute_type)) { @@ -1810,6 +2191,15 @@ static bool ggml_cuda_should_fuse_mul_mat_vec_q(const ggml_tensor * tensor) { } static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + if (src0->type == GGML_TYPE_I8) { +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + ggml_cuda_mul_mat_i8(ctx, src0, src1, dst); +#else + GGML_ABORT("INT8 tensorwise matmul is only implemented for CUDA"); +#endif + return; + } + GGML_TENSOR_BINARY_OP_LOCALS const int32_t hint = ggml_get_op_params_i32(dst, 1); @@ -2217,6 +2607,16 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg case GGML_OP_MUL_MAT_ID: ggml_cuda_mul_mat_id(ctx, dst); break; + case GGML_OP_REGULAR_HADAMARD: + GGML_ASSERT(ggml_cuda_op_regular_hadamard(ctx, dst->src[0], dst, ggml_get_op_params_i32(dst, 0))); + break; + case GGML_OP_QUANTIZE_I8_CONVROT: +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + GGML_ASSERT(ggml_cuda_op_quantize_i8_convrot(ctx, dst->src[0], dst, ggml_get_op_params_i32(dst, 0))); +#else + GGML_ABORT("INT8 convrot quantization is only implemented for CUDA"); +#endif + break; case GGML_OP_OUT_PROD: ggml_cuda_out_prod(ctx, dst); break; @@ -4885,6 +5285,33 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g { struct ggml_tensor * a = op->src[0]; struct ggml_tensor * b = op->src[1]; + if (a->type == GGML_TYPE_I8) { +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + const ggml_tensor * weight_scale = op->src[2]; + const ggml_tensor * bias = op->src[3]; + const int convrot_group_size = ggml_get_op_params_i32(op, 2); + const ggml_tensor * packed_src = b->src[0]; + const int64_t packed_rows = packed_src != nullptr ? GGML_PAD(ggml_nrows(packed_src), 4) : 0; + const int64_t packed_scale_rows = packed_src != nullptr ? (ggml_nrows(packed_src) * (int64_t) sizeof(float) + a->ne[0] - 1) / a->ne[0] : 0; + const bool packed_input = b->type == GGML_TYPE_I8 && + b->op == GGML_OP_QUANTIZE_I8_CONVROT && packed_src != nullptr && + b->ne[0] == a->ne[0] && b->ne[1] == packed_rows + packed_scale_rows; + return op->op == GGML_OP_MUL_MAT && + (b->type == GGML_TYPE_F32 || packed_input) && + op->type == GGML_TYPE_F32 && + a->ne[0] % 4 == 0 && a->ne[1] % 4 == 0 && + a->ne[2] == 1 && a->ne[3] == 1 && + ggml_is_contiguous(a) && ggml_is_contiguous(b) && ggml_is_contiguous(op) && + (weight_scale == nullptr || + (weight_scale->type == GGML_TYPE_F32 && ggml_is_contiguous(weight_scale) && ggml_nelements(weight_scale) == a->ne[1])) && + (bias == nullptr || + (bias->type == GGML_TYPE_F32 && ggml_is_contiguous(bias) && ggml_nelements(bias) == a->ne[1])) && + (convrot_group_size == 0 || (convrot_group_size == 256 && a->ne[0] % 256 == 0)) && + turing_mma_available(ggml_cuda_info().devices[dev_ctx->device].cc); +#else + return false; +#endif + } if (a->nb[0] != ggml_element_size(a) || b->nb[0] != ggml_element_size(b)) { return false; // TODO this could in principle be implemented though currently there is no use case. } @@ -5260,6 +5687,30 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_OP_DIAG: case GGML_OP_SOLVE_TRI: return true; + case GGML_OP_REGULAR_HADAMARD: { + const ggml_tensor * src = op->src[0]; + const int group_size = ggml_get_op_params_i32(op, 0); + return src != nullptr && src->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 && + ggml_is_contiguous(src) && ggml_is_contiguous(op) && group_size > 0 && + src->ne[0] % group_size == 0 && + (group_size == 4 || group_size == 16 || group_size == 64 || group_size == 256); + } + case GGML_OP_QUANTIZE_I8_CONVROT: +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + { + const ggml_tensor * src = op->src[0]; + const int group_size = ggml_get_op_params_i32(op, 0); + const int64_t rows = src != nullptr ? ggml_nrows(src) : 0; + const int64_t rows_padded = GGML_PAD(rows, 4); + const int64_t scale_rows = src != nullptr ? (rows * (int64_t) sizeof(float) + src->ne[0] - 1) / src->ne[0] : 0; + return src != nullptr && src->type == GGML_TYPE_F32 && op->type == GGML_TYPE_I8 && + ggml_is_contiguous(src) && ggml_is_contiguous(op) && group_size == 256 && + src->ne[0] % group_size == 0 && op->ne[0] == src->ne[0] && + op->ne[1] == rows_padded + scale_rows; + } +#else + return false; +#endif case GGML_OP_LIGHTNING_INDEXER: return ggml_cuda_lightning_indexer_supported(dev_ctx->device, op); diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 45fa97f8129..f66f42bca4f 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -911,6 +911,9 @@ struct vk_device_struct { vk_pipeline pipeline_matmul_split_k_reduce; vk_pipeline pipeline_quantize_q8_1_x4; + vk_pipeline pipeline_quantize_i8_convrot; + vk_pipeline pipeline_mul_mat_i8_tensorwise; + vk_pipeline pipeline_mul_mat_i8_tensorwise_cm1; vk_pipeline pipeline_dequant[GGML_TYPE_COUNT]; vk_pipeline pipeline_dequant_mul_mat_vec_f32_f32[DMMV_WG_SIZE_COUNT][GGML_TYPE_COUNT][mul_mat_vec_max_cols]; @@ -2002,6 +2005,19 @@ struct vk_quantize_q8_1_push_constants { uint32_t num_blocks; }; +struct vk_quantize_i8_convrot_push_constants { + uint32_t k; + uint32_t rows; +}; + +struct vk_mul_mat_i8_tensorwise_push_constants { + uint32_t n; + uint32_t rows; + uint32_t k; + uint32_t has_bias; + uint32_t row_offset; +}; + struct vk_op_flash_attn_split_k_reduce_push_constants { uint32_t D; uint32_t ne1; @@ -5455,6 +5471,18 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_quantize_q8_1_x4, "quantize_q8_1_x4", quantize_q8_1_x4_len, quantize_q8_1_x4_data, "main", 2, sizeof(vk_quantize_q8_1_push_constants), {32 * device->subgroup_size / 8, 1, 1}, { device->subgroup_size }, 1); } +#if defined(GGML_VULKAN_INTEGER_DOT_GLSLC_SUPPORT) + if (device->integer_dot_product) { + ggml_vk_create_pipeline(device, device->pipeline_quantize_i8_convrot, "quantize_i8_convrot", quantize_i8_convrot_len, quantize_i8_convrot_data, "main", 3, sizeof(vk_quantize_i8_convrot_push_constants), {1, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_mul_mat_i8_tensorwise, "mul_mat_i8_tensorwise", mul_mat_i8_tensorwise_len, mul_mat_i8_tensorwise_data, "main", 6, sizeof(vk_mul_mat_i8_tensorwise_push_constants), {16, 16, 1}, {}, 1); +#if defined(GGML_VULKAN_COOPMAT_GLSLC_SUPPORT) + if (device->coopmat_int_support) { + ggml_vk_create_pipeline(device, device->pipeline_mul_mat_i8_tensorwise_cm1, "mul_mat_i8_tensorwise_cm1", mul_mat_i8_tensorwise_cm1_len, mul_mat_i8_tensorwise_cm1_data, "main", 6, sizeof(vk_mul_mat_i8_tensorwise_push_constants), {4 * device->coopmat_int_n, 4 * device->coopmat_int_m, 1}, {device->subgroup_size, device->coopmat_int_m, device->coopmat_int_n, device->coopmat_int_k}, 1, true, true, device->subgroup_size); + } +#endif + } +#endif + for (uint32_t i = 0; i < p021_max_gqa_ratio; ++i) { if (device->subgroup_arithmetic && device->subgroup_require_full_support) { ggml_vk_create_pipeline2(device, device->pipeline_mul_mat_vec_p021_f16_f32[i], "mul_mat_vec_p021_f16_f32"+std::to_string(i+1), mul_mat_vec_p021_f16_f32_subgroup_add_len, mul_mat_vec_p021_f16_f32_subgroup_add_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_p021_push_constants), {1, 1, 1}, {device->subgroup_size, i + 1}, 1, true, true); @@ -9954,6 +9982,142 @@ static void ggml_vk_fwht(ggml_backend_vk_context * ctx, vk_context& subctx, cons ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { src_buf, dst_buf }, pc, { workgroups_x, 1, 1 }); } +static bool ggml_vk_can_use_quantize_i8_convrot(const vk_device & device, const ggml_tensor * dst) { +#if defined(GGML_VULKAN_INTEGER_DOT_GLSLC_SUPPORT) + const ggml_tensor * src = dst->src[0]; + if (!device->integer_dot_product || src == nullptr || + src->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_I8 || + !ggml_is_contiguous(src) || !ggml_is_contiguous(dst)) { + return false; + } + + const int group_size = ggml_get_op_params_i32(dst, 0); + const int64_t rows = ggml_nrows(src); + const int64_t rows_padded = GGML_PAD(rows, 4); + const int64_t scale_rows = (rows * (int64_t) sizeof(float) + src->ne[0] - 1) / src->ne[0]; + return group_size == 256 && src->ne[0] % 256 == 0 && + src->ne[0] <= std::numeric_limits::max() && + rows <= std::numeric_limits::max() && + dst->ne[0] == src->ne[0] && dst->ne[1] == rows_padded + scale_rows; +#else + GGML_UNUSED(device); + GGML_UNUSED(dst); + return false; +#endif +} + +static bool ggml_vk_can_use_mul_mat_i8_tensorwise(const vk_device & device, const ggml_tensor * dst) { +#if defined(GGML_VULKAN_INTEGER_DOT_GLSLC_SUPPORT) + const ggml_tensor * weight = dst->src[0]; + const ggml_tensor * input = dst->src[1]; + const ggml_tensor * weight_scale = dst->src[2]; + const ggml_tensor * bias = dst->src[3]; + const ggml_tensor * logical_input = input != nullptr ? input->src[0] : nullptr; + if (!device->integer_dot_product || weight == nullptr || input == nullptr || + weight->type != GGML_TYPE_I8 || input->type != GGML_TYPE_I8 || + input->op != GGML_OP_QUANTIZE_I8_CONVROT || logical_input == nullptr || + logical_input->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32 || + weight_scale == nullptr || weight->ne[2] != 1 || weight->ne[3] != 1 || + !ggml_is_contiguous(weight) || !ggml_is_contiguous(input) || + !ggml_is_contiguous(logical_input) || !ggml_is_contiguous(weight_scale) || + !ggml_is_contiguous(dst)) { + return false; + } + + const int convrot_group_size = ggml_get_op_params_i32(dst, 2); + const int64_t rows = ggml_nrows(logical_input); + const int64_t rows_padded = GGML_PAD(rows, 4); + const int64_t scale_rows = (rows * (int64_t) sizeof(float) + weight->ne[0] - 1) / weight->ne[0]; + const bool valid_bias = bias == nullptr || + (bias->type == GGML_TYPE_F32 && ggml_is_contiguous(bias) && ggml_nelements(bias) == weight->ne[1]); + return convrot_group_size == 256 && ggml_get_op_params_i32(input, 0) == 256 && + weight->ne[0] % 256 == 0 && input->ne[0] == weight->ne[0] && + input->ne[1] == rows_padded + scale_rows && dst->ne[0] == weight->ne[1] && + ggml_nrows(dst) == rows && weight_scale->type == GGML_TYPE_F32 && + ggml_nelements(weight_scale) == weight->ne[1] && valid_bias && + weight->ne[0] / 4 <= std::numeric_limits::max() && + weight->ne[1] <= std::numeric_limits::max() && + rows <= std::numeric_limits::max(); +#else + GGML_UNUSED(device); + GGML_UNUSED(dst); + return false; +#endif +} + +static void ggml_vk_quantize_i8_convrot(ggml_backend_vk_context * ctx, vk_context & subctx, const ggml_tensor * src, ggml_tensor * dst) { + GGML_ASSERT(ggml_vk_can_use_quantize_i8_convrot(ctx->device, dst)); + vk_pipeline pipeline = ctx->device->pipeline_quantize_i8_convrot; + ggml_pipeline_request_descriptor_sets(ctx, pipeline, 1); + + const uint32_t rows = (uint32_t) ggml_nrows(src); + const uint64_t rows_padded = GGML_PAD((uint64_t) rows, 4); + const uint64_t scale_offset = (uint64_t) src->ne[0] * rows_padded; + const vk_subbuffer src_buf = ggml_vk_tensor_subbuffer(ctx, src); + const vk_subbuffer dst_buf = ggml_vk_tensor_subbuffer(ctx, dst); + const vk_subbuffer quantized_buf = { dst_buf.buffer, dst_buf.offset, scale_offset }; + const vk_subbuffer scales_buf = { + dst_buf.buffer, + dst_buf.offset + scale_offset, + (uint64_t) rows * sizeof(float), + }; + const vk_quantize_i8_convrot_push_constants pc = { + (uint32_t) src->ne[0], + rows, + }; + const uint32_t workgroups = std::min(rows, ctx->device->properties.limits.maxComputeWorkGroupCount[0]); + ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { src_buf, quantized_buf, scales_buf }, pc, { workgroups, 1, 1 }); +} + +static void ggml_vk_mul_mat_i8_tensorwise(ggml_backend_vk_context * ctx, vk_context & subctx, const ggml_tensor * weight, const ggml_tensor * input, ggml_tensor * dst) { + GGML_ASSERT(ggml_vk_can_use_mul_mat_i8_tensorwise(ctx->device, dst)); + const uint32_t rows = (uint32_t) ggml_nrows(dst); + uint32_t coop_rows = 0; +#if defined(GGML_VULKAN_INTEGER_DOT_GLSLC_SUPPORT) && defined(GGML_VULKAN_COOPMAT_GLSLC_SUPPORT) + if (ctx->device->coopmat_int_support && weight->ne[0] % ctx->device->coopmat_int_k == 0 && weight->ne[1] % (4 * ctx->device->coopmat_int_n) == 0) { + const uint32_t tile_rows = 4 * ctx->device->coopmat_int_m; + coop_rows = rows / tile_rows * tile_rows; + } +#endif + + const ggml_tensor * weight_scale = dst->src[2]; + const ggml_tensor * bias = dst->src[3]; + const uint64_t rows_padded = GGML_PAD((uint64_t) rows, 4); + const uint64_t scale_offset = (uint64_t) weight->ne[0] * rows_padded; + const vk_subbuffer weight_buf = ggml_vk_tensor_subbuffer(ctx, weight); + const vk_subbuffer input_buf = ggml_vk_tensor_subbuffer(ctx, input); + const vk_subbuffer quantized_buf = { input_buf.buffer, input_buf.offset, scale_offset }; + const vk_subbuffer activation_scales_buf = { + input_buf.buffer, + input_buf.offset + scale_offset, + (uint64_t) rows * sizeof(float), + }; + const vk_subbuffer weight_scales_buf = ggml_vk_tensor_subbuffer(ctx, weight_scale); + const vk_subbuffer bias_buf = bias != nullptr ? ggml_vk_tensor_subbuffer(ctx, bias) : weight_scales_buf; + const vk_subbuffer dst_buf = ggml_vk_tensor_subbuffer(ctx, dst); + const auto dispatch = [&](vk_pipeline pipeline, uint32_t row_offset, uint32_t dispatch_rows) { + if (dispatch_rows == 0) { + return; + } + ggml_pipeline_request_descriptor_sets(ctx, pipeline, 1); + const vk_mul_mat_i8_tensorwise_push_constants pc = { + (uint32_t) weight->ne[1], + rows, + (uint32_t) weight->ne[0], + bias != nullptr ? 1u : 0u, + row_offset, + }; + ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, + { weight_buf, quantized_buf, weight_scales_buf, activation_scales_buf, bias_buf, dst_buf }, + pc, { (uint32_t) weight->ne[1], dispatch_rows, 1 }); + }; + + if (coop_rows > 0) { + dispatch(ctx->device->pipeline_mul_mat_i8_tensorwise_cm1, 0, coop_rows); + } + dispatch(ctx->device->pipeline_mul_mat_i8_tensorwise, coop_rows, rows - coop_rows); +} + static void ggml_vk_mul_mat(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx) { ggml_tensor * dst = cgraph->nodes[node_idx]; ggml_tensor * src0 = dst->src[0]; @@ -15593,9 +15757,17 @@ static bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgr case GGML_OP_LEAKY_RELU: ggml_vk_leaky_relu(ctx, compute_ctx, src0, node); + break; + case GGML_OP_QUANTIZE_I8_CONVROT: + ggml_vk_quantize_i8_convrot(ctx, compute_ctx, src0, node); + break; case GGML_OP_MUL_MAT: - ggml_vk_mul_mat(ctx, compute_ctx, cgraph, node_idx); + if (src0->type == GGML_TYPE_I8) { + ggml_vk_mul_mat_i8_tensorwise(ctx, compute_ctx, src0, src1, node); + } else { + ggml_vk_mul_mat(ctx, compute_ctx, cgraph, node_idx); + } break; case GGML_OP_MUL_MAT_ID: @@ -16430,7 +16602,7 @@ static bool ggml_vk_can_fuse(const ggml_backend_vk_context * ctx, const struct g const ggml_tensor *mul = cgraph->nodes[node_idx]; const ggml_tensor *add = cgraph->nodes[node_idx + 1]; - if (!mm_add_ok(mul, add)) { + if (mul->src[0]->type == GGML_TYPE_I8 || !mm_add_ok(mul, add)) { return false; } if (ops.size() == 3) { @@ -17977,10 +18149,15 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm default: return false; } + case GGML_OP_QUANTIZE_I8_CONVROT: + return ggml_vk_can_use_quantize_i8_convrot(device, op); case GGML_OP_MUL_MAT: case GGML_OP_MUL_MAT_ID: { ggml_type src0_type = op->src[0]->type; + if (src0_type == GGML_TYPE_I8) { + return op->op == GGML_OP_MUL_MAT && ggml_vk_can_use_mul_mat_i8_tensorwise(device, op); + } if (op->op == GGML_OP_MUL_MAT_ID) { if (!device->mul_mat_id_s[src0_type] && !device->mul_mat_id_m[src0_type] && !device->mul_mat_id_l[src0_type]) { // If there's not enough shared memory for row_ids and the result tile, fallback to CPU diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_i8_tensorwise.comp b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_i8_tensorwise.comp new file mode 100644 index 00000000000..08b45c23571 --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_i8_tensorwise.comp @@ -0,0 +1,58 @@ +#version 450 + +#extension GL_EXT_control_flow_attributes : require +#extension GL_EXT_integer_dot_product : require + +layout(local_size_x = 16, local_size_y = 16, local_size_z = 1) in; + +layout(push_constant) uniform parameter +{ + uint n; + uint rows; + uint k; + uint has_bias; + uint row_offset; +} p; + +layout(binding = 0, std430) readonly buffer W { int data_w[]; }; +layout(binding = 1, std430) readonly buffer A { int data_a[]; }; +layout(binding = 2, std430) readonly buffer WS { float data_ws[]; }; +layout(binding = 3, std430) readonly buffer AS { float data_as[]; }; +layout(binding = 4, std430) readonly buffer B { float data_b[]; }; +layout(binding = 5, std430) writeonly buffer D { float data_d[]; }; + +shared int weights[16 * 8]; +shared int activations[16 * 8]; + +void main() { + const uint output_index = gl_WorkGroupID.x * 16 + gl_LocalInvocationID.x; + const uint row = p.row_offset + gl_WorkGroupID.y * 16 + gl_LocalInvocationID.y; + const uint lx = gl_LocalInvocationID.x; + const uint ly = gl_LocalInvocationID.y; + const uint k_packed = p.k / 4; + int sum = 0; + + for (uint k_base = 0; k_base < k_packed; k_base += 8) { + if (ly < 8) { + weights[lx * 8 + ly] = output_index < p.n ? data_w[output_index * k_packed + k_base + ly] : 0; + } + if (lx < 8) { + activations[ly * 8 + lx] = row < p.rows ? data_a[row * k_packed + k_base + lx] : 0; + } + barrier(); + + [[unroll]] + for (uint i = 0; i < 8; ++i) { + sum += dotPacked4x8EXT(weights[lx * 8 + i], activations[ly * 8 + i]); + } + barrier(); + } + + if (output_index < p.n && row < p.rows) { + float value = float(sum) * data_as[row] * data_ws[output_index]; + if (p.has_bias != 0) { + value += data_b[output_index]; + } + data_d[row * p.n + output_index] = value; + } +} diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_i8_tensorwise_cm1.comp b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_i8_tensorwise_cm1.comp new file mode 100644 index 00000000000..35cea6bf879 --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_i8_tensorwise_cm1.comp @@ -0,0 +1,87 @@ +#version 460 + +#extension GL_EXT_control_flow_attributes : require +#extension GL_EXT_shader_8bit_storage : require +#extension GL_EXT_shader_explicit_arithmetic_types_int8 : require +#extension GL_EXT_shader_explicit_arithmetic_types_int32 : require +#extension GL_KHR_cooperative_matrix : require +#extension GL_KHR_memory_scope_semantics : require +#extension GL_KHR_shader_subgroup_basic : require + +layout(constant_id = 0) const uint BLOCK_SIZE = 32; +layout(constant_id = 1) const uint CM_M = 16; +layout(constant_id = 2) const uint CM_N = 16; +layout(constant_id = 3) const uint CM_K = 16; +const uint TILE_M = 4; +const uint TILE_N = 4; + +layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in; + +layout(push_constant) uniform parameter +{ + uint n; + uint rows; + uint k; + uint has_bias; + uint row_offset; +} p; + +layout(binding = 0, std430) readonly buffer W { int8_t data_w[]; }; +layout(binding = 1, std430) readonly buffer A { int8_t data_a[]; }; +layout(binding = 2, std430) readonly buffer WS { float data_ws[]; }; +layout(binding = 3, std430) readonly buffer AS { float data_as[]; }; +layout(binding = 4, std430) readonly buffer B { float data_b[]; }; +layout(binding = 5, std430) writeonly buffer D { float data_d[]; }; + +shared int32_t accumulators[1024]; + +void main() { + const uint output_base = gl_WorkGroupID.x * CM_N * TILE_N; + const uint row_base = gl_WorkGroupID.y * CM_M * TILE_M; + coopmat sums[TILE_M * TILE_N]; + [[unroll]] + for (uint i = 0; i < TILE_M * TILE_N; ++i) { + sums[i] = coopmat(0); + } + + for (uint k_base = 0; k_base < p.k; k_base += CM_K) { + coopmat activations[TILE_M]; + coopmat weights[TILE_N]; + [[unroll]] + for (uint tile_m = 0; tile_m < TILE_M; ++tile_m) { + coopMatLoad(activations[tile_m], data_a, (row_base + tile_m * CM_M) * p.k + k_base, p.k, gl_CooperativeMatrixLayoutRowMajor); + } + [[unroll]] + for (uint tile_n = 0; tile_n < TILE_N; ++tile_n) { + coopMatLoad(weights[tile_n], data_w, (output_base + tile_n * CM_N) * p.k + k_base, p.k, gl_CooperativeMatrixLayoutColumnMajor); + } + [[unroll]] + for (uint tile_m = 0; tile_m < TILE_M; ++tile_m) { + [[unroll]] + for (uint tile_n = 0; tile_n < TILE_N; ++tile_n) { + const uint index = tile_m * TILE_N + tile_n; + sums[index] = coopMatMulAdd(activations[tile_m], weights[tile_n], sums[index]); + } + } + } + + [[unroll]] + for (uint tile_m = 0; tile_m < TILE_M; ++tile_m) { + [[unroll]] + for (uint tile_n = 0; tile_n < TILE_N; ++tile_n) { + coopMatStore(sums[tile_m * TILE_N + tile_n], accumulators, 0, CM_N, gl_CooperativeMatrixLayoutRowMajor); + controlBarrier(gl_ScopeSubgroup, gl_ScopeSubgroup, gl_StorageSemanticsShared, gl_SemanticsAcquireRelease); + + for (uint index = gl_SubgroupInvocationID; index < CM_M * CM_N; index += gl_SubgroupSize) { + const uint row = row_base + tile_m * CM_M + index / CM_N; + const uint output_index = output_base + tile_n * CM_N + index % CM_N; + float value = float(accumulators[index]) * data_as[row] * data_ws[output_index]; + if (p.has_bias != 0) { + value += data_b[output_index]; + } + data_d[row * p.n + output_index] = value; + } + controlBarrier(gl_ScopeSubgroup, gl_ScopeSubgroup, gl_StorageSemanticsShared, gl_SemanticsAcquireRelease); + } + } +} diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/quantize_i8_convrot.comp b/ggml/src/ggml-vulkan/vulkan-shaders/quantize_i8_convrot.comp new file mode 100644 index 00000000000..1c69d9dde2e --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/quantize_i8_convrot.comp @@ -0,0 +1,102 @@ +#version 450 + +#extension GL_EXT_control_flow_attributes : require + +layout(local_size_x = 64, local_size_y = 1, local_size_z = 1) in; + +layout(push_constant) uniform parameter +{ + uint k; + uint rows; +} p; + +layout(binding = 0, std430) readonly buffer A { float data_a[]; }; +layout(binding = 1, std430) writeonly buffer D { uint data_d[]; }; +layout(binding = 2, std430) writeonly buffer S { float data_s[]; }; + +shared float values[256]; +shared float maxima[64]; + +void transform_group(const uint row, const uint group) { + const uint tid = gl_LocalInvocationID.x; + const uint base_offset = row * p.k + group * 256; + + [[unroll]] + for (uint i = 0; i < 4; ++i) { + const uint index = tid + i * 64; + values[index] = data_a[base_offset + index] * (1.0 / 16.0); + } + barrier(); + + [[unroll]] + for (uint stride = 1; stride < 256; stride *= 4) { + const uint base = (tid / stride) * 4 * stride + tid % stride; + const uint i0 = base; + const uint i1 = i0 + stride; + const uint i2 = i1 + stride; + const uint i3 = i2 + stride; + const float a = values[i0]; + const float b = values[i1]; + const float c = values[i2]; + const float d = values[i3]; + values[i0] = a + b + c - d; + values[i1] = a + b - c + d; + values[i2] = a - b + c + d; + values[i3] = -a + b + c + d; + barrier(); + } +} + +uint quantize_value(const float value, const float inv_scale) { + const int quantized = clamp(int(roundEven(value * inv_scale)), -127, 127); + return uint(quantized) & 0xffu; +} + +void quantize_row(const uint row) { + const uint tid = gl_LocalInvocationID.x; + const uint groups = p.k / 256; + float local_max = 0.0; + + for (uint group = 0; group < groups; ++group) { + transform_group(row, group); + [[unroll]] + for (uint i = 0; i < 4; ++i) { + local_max = max(local_max, abs(values[tid + i * 64])); + } + barrier(); + } + + maxima[tid] = local_max; + barrier(); + [[unroll]] + for (uint stride = 32; stride > 0; stride >>= 1) { + if (tid < stride) { + maxima[tid] = max(maxima[tid], maxima[tid + stride]); + } + barrier(); + } + + if (tid == 0) { + data_s[row] = maxima[0] / 127.0; + } + barrier(); + + const float row_scale = maxima[0] / 127.0; + const float inv_scale = row_scale == 0.0 ? 0.0 : 1.0 / row_scale; + for (uint group = 0; group < groups; ++group) { + transform_group(row, group); + const uint index = tid * 4; + const uint packed = quantize_value(values[index], inv_scale) + | (quantize_value(values[index + 1], inv_scale) << 8) + | (quantize_value(values[index + 2], inv_scale) << 16) + | (quantize_value(values[index + 3], inv_scale) << 24); + data_d[(row * p.k + group * 256) / 4 + tid] = packed; + barrier(); + } +} + +void main() { + for (uint row = gl_WorkGroupID.x; row < p.rows; row += gl_NumWorkGroups.x) { + quantize_row(row); + } +} diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp index c93d6eecee1..a0ae1c44cc1 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp @@ -876,6 +876,14 @@ void process_shaders() { string_to_spv("quantize_q8_1_x4", "quantize_q8_1.comp", {{"QBLOCK_X4", "1"}}); string_to_spv("quantize_q8_1_x4_subgroup", "quantize_q8_1.comp", {{"QBLOCK_X4", "1"}, {"USE_SUBGROUPS", "1"}}); +#if defined(GGML_VULKAN_INTEGER_DOT_GLSLC_SUPPORT) + string_to_spv("quantize_i8_convrot", "quantize_i8_convrot.comp", {}); + string_to_spv("mul_mat_i8_tensorwise", "mul_mat_i8_tensorwise.comp", {}); +#if defined(GGML_VULKAN_COOPMAT_GLSLC_SUPPORT) + string_to_spv("mul_mat_i8_tensorwise", "mul_mat_i8_tensorwise_cm1.comp", {}, true, true); +#endif +#endif + string_to_spv("mul_f32", "mul.comp", {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}}); string_to_spv("div_f32", "div.comp", {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}}); diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index da7f3a5f2e3..b9fe5116574 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -1098,9 +1098,12 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = { "OPT_STEP_SGD", "GLU", + + "REGULAR_HADAMARD", + "QUANTIZE_I8_CONVROT", }; -static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT != 101"); +static_assert(GGML_OP_COUNT == 103, "GGML_OP_COUNT != 103"); static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "none", @@ -1213,9 +1216,12 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "sgd(x)", "glu(x)", + + "regular_hadamard(x)", + "quantize_i8_convrot(x)", }; -static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT != 101"); +static_assert(GGML_OP_COUNT == 103, "GGML_OP_COUNT != 103"); static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2"); @@ -3292,6 +3298,51 @@ struct ggml_tensor * ggml_mul_mat( return result; } +struct ggml_tensor * ggml_mul_mat_i8_tensorwise( + struct ggml_context * ctx, + struct ggml_tensor * weight, + struct ggml_tensor * input, + struct ggml_tensor * weight_scale, + struct ggml_tensor * bias, + int convrot_group_size) { + GGML_ASSERT(weight->type == GGML_TYPE_I8); + GGML_ASSERT(input->type == GGML_TYPE_F32 || input->type == GGML_TYPE_I8); + GGML_ASSERT(weight_scale != NULL && weight_scale->type == GGML_TYPE_F32); + GGML_ASSERT(ggml_is_contiguous(weight_scale)); + GGML_ASSERT(ggml_nelements(weight_scale) == weight->ne[1]); + GGML_ASSERT(bias == NULL || (bias->type == GGML_TYPE_F32 && ggml_is_contiguous(bias))); + GGML_ASSERT(bias == NULL || ggml_nelements(bias) == weight->ne[1]); + GGML_ASSERT(convrot_group_size == 0 || weight->ne[0] % convrot_group_size == 0); + + int n = convrot_group_size; + while (n > 1 && n % 4 == 0) { + n /= 4; + } + GGML_ASSERT(n == 0 || n == 1); + + const struct ggml_tensor * logical_input = input; + if (input->type == GGML_TYPE_I8) { + GGML_ASSERT(input->op == GGML_OP_QUANTIZE_I8_CONVROT); + GGML_ASSERT(input->src[0] != NULL); + logical_input = input->src[0]; + const int64_t rows_padded = GGML_PAD(ggml_nrows(logical_input), 4); + const int64_t scale_rows = (ggml_nrows(logical_input) * (int64_t) sizeof(float) + weight->ne[0] - 1) / weight->ne[0]; + GGML_ASSERT(input->ne[0] == weight->ne[0]); + GGML_ASSERT(input->ne[1] == rows_padded + scale_rows); + } + GGML_ASSERT(ggml_can_mul_mat(weight, logical_input)); + + const int64_t ne[4] = { weight->ne[1], logical_input->ne[1], logical_input->ne[2], logical_input->ne[3] }; + struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne); + result->op = GGML_OP_MUL_MAT; + result->src[0] = weight; + result->src[1] = input; + result->src[2] = weight_scale; + result->src[3] = bias; + ggml_set_op_params_i32(result, 2, convrot_group_size); + return result; +} + void ggml_mul_mat_set_prec( struct ggml_tensor * a, enum ggml_prec prec) { @@ -3312,6 +3363,50 @@ void ggml_mul_mat_set_hint( ggml_set_op_params_i32(a, 1, hint_i32); } +struct ggml_tensor * ggml_regular_hadamard( + struct ggml_context * ctx, + struct ggml_tensor * a, + int group_size) { + GGML_ASSERT(a->type == GGML_TYPE_F32); + GGML_ASSERT(group_size > 0 && a->ne[0] % group_size == 0); + + int n = group_size; + while (n > 1 && n % 4 == 0) { + n /= 4; + } + GGML_ASSERT(n == 1); + + struct ggml_tensor * result = ggml_dup_tensor(ctx, a); + result->op = GGML_OP_REGULAR_HADAMARD; + result->src[0] = a; + ggml_set_op_params_i32(result, 0, group_size); + return result; +} + +struct ggml_tensor * ggml_quantize_i8_convrot( + struct ggml_context * ctx, + struct ggml_tensor * a, + int group_size) { + GGML_ASSERT(a->type == GGML_TYPE_F32); + GGML_ASSERT(ggml_is_contiguous(a)); + GGML_ASSERT(group_size > 0 && a->ne[0] % group_size == 0); + + int n = group_size; + while (n > 1 && n % 4 == 0) { + n /= 4; + } + GGML_ASSERT(n == 1); + + const int64_t rows_padded = GGML_PAD(ggml_nrows(a), 4); + const int64_t scale_rows = (ggml_nrows(a) * (int64_t) sizeof(float) + a->ne[0] - 1) / a->ne[0]; + const int64_t ne[2] = { a->ne[0], rows_padded + scale_rows }; + struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_I8, 2, ne); + result->op = GGML_OP_QUANTIZE_I8_CONVROT; + result->src[0] = a; + ggml_set_op_params_i32(result, 0, group_size); + return result; +} + // ggml_mul_mat_id /* diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 419e1eba4c2..b2fe6b26899 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -288,6 +288,7 @@ endif() if (NOT GGML_BACKEND_DL) # these tests use the backends directly and cannot be built with dynamic loading llama_build_and_test(test-barrier.cpp) + llama_build_and_test(test-int8-convrot.cpp) llama_build_and_test(test-quantize-fns.cpp) llama_build_and_test(test-quantize-perf.cpp) llama_build_and_test(test-rope.cpp) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index a4c22156c63..70c655a01b6 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -4407,6 +4407,69 @@ struct test_mul_mat_hadamard : public test_mul_mat { } }; +struct test_mul_mat_i8_convrot : public test_case { + const int64_t k; + const int64_t n; + const int64_t rows; + const bool use_bias; + + std::string vars() override { + return VARS_TO_STR4(k, n, rows, use_bias); + } + + test_mul_mat_i8_convrot(int64_t k, int64_t n, int64_t rows, bool use_bias) + : k(k), n(n), rows(rows), use_bias(use_bias) {} + + ggml_tensor * build_graph(ggml_context * ctx) override { + ggml_tensor * input = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, k, rows); + ggml_tensor * weight = ggml_new_tensor_2d(ctx, GGML_TYPE_I8, k, n); + ggml_tensor * weight_scale = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n); + ggml_tensor * bias = use_bias ? ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n) : nullptr; + ggml_set_name(input, "input"); + ggml_set_name(weight, "weight"); + ggml_set_name(weight_scale, "weight_scale"); + if (bias != nullptr) { + ggml_set_name(bias, "bias"); + } + + ggml_tensor * quantized = ggml_quantize_i8_convrot(ctx, input, 256); + ggml_set_name(quantized, "quantized"); + ggml_tensor * out = ggml_mul_mat_i8_tensorwise(ctx, weight, quantized, weight_scale, bias, 256); + ggml_set_name(out, "out"); + return out; + } + + void initialize_tensors(ggml_context * ctx) override { + for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != nullptr; t = ggml_get_next_tensor(ctx, t)) { + if (strcmp(t->name, "weight") == 0) { + std::vector data(ggml_nelements(t)); + for (size_t i = 0; i < data.size(); ++i) { + data[i] = (int8_t) ((i * 17 + i / (size_t) k * 11) % 61 - 30); + } + ggml_backend_tensor_set(t, data.data(), 0, data.size()); + } else if (strcmp(t->name, "weight_scale") == 0) { + init_tensor_uniform(t, 0.01f, 0.25f); + } else if (strcmp(t->name, "input") == 0 || strcmp(t->name, "bias") == 0) { + init_tensor_uniform(t, -2.0f, 2.0f); + } + } + } + + std::string op_desc(ggml_tensor * t) override { + GGML_UNUSED(t); + return "MUL_MAT_I8_CONVROT"; + } + + uint64_t op_flops(ggml_tensor * t) override { + GGML_UNUSED(t); + return 2 * k * n * rows; + } + + bool run_whole_graph() override { + return true; + } +}; + static void init_mul_mat_id_tensors(ggml_context * ctx, int n_mats) { std::random_device rd; std::default_random_engine rng(rd()); @@ -8847,6 +8910,10 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 32, 1, 32)); // too small (N<64) test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1024, 1, 1024)); // too big (N>512) + test_cases.emplace_back(new test_mul_mat_i8_convrot(256, 4, 1, true)); + test_cases.emplace_back(new test_mul_mat_i8_convrot(512, 128, 65, true)); + test_cases.emplace_back(new test_mul_mat_i8_convrot(512, 128, 65, false)); + #if 0 // > 4GB A matrix. Too slow to be enabled by default. test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F16, 900000, 3, 2592, {1, 1}, {1, 1})); @@ -9910,6 +9977,8 @@ static std::vector> make_test_cases_perf() { test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 256, 2048, 256)); test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 512, 2048, 512)); + test_cases.emplace_back(new test_mul_mat_i8_convrot(512, 128, 65, true)); + test_cases.emplace_back(new test_solve_tri(GGML_TYPE_F32, { 64, 64, 4, 4 }, { 32, 64, 4, 4 })); test_cases.emplace_back(new test_solve_tri(GGML_TYPE_F32, { 128, 128, 4, 2 }, { 32, 128, 4, 2 })); // qwen3next with CHUNK_SIZE 64 diff --git a/tests/test-int8-convrot.cpp b/tests/test-int8-convrot.cpp new file mode 100644 index 00000000000..cfdfc33f684 --- /dev/null +++ b/tests/test-int8-convrot.cpp @@ -0,0 +1,303 @@ +#include "ggml-backend.h" +#include "ggml-cpu.h" +#include "ggml.h" + +#include +#include +#include +#include +#include +#include + +static bool nearly_equal(float actual, float expected, float tolerance = 1e-5f) { + return std::fabs(actual - expected) <= tolerance * std::max(1.0f, std::fabs(expected)); +} + +static float regular_hadamard_reference(const std::vector & input, int output) { + static constexpr int h4[4][4] = { + { 1, 1, 1, -1}, + { 1, 1, -1, 1}, + { 1, -1, 1, 1}, + {-1, 1, 1, 1}, + }; + + float sum = 0.0f; + for (int column = 0; column < (int) input.size(); ++column) { + int row_digit = output; + int column_digit = column; + int sign = 1; + while (row_digit != 0 || column_digit != 0) { + sign *= h4[row_digit % 4][column_digit % 4]; + row_digit /= 4; + column_digit /= 4; + } + sum += sign * input[column]; + } + return sum / std::sqrt((float) input.size()); +} + +int main(int argc, char ** argv) { + const bool use_cuda = argc == 2 && std::strcmp(argv[1], "--cuda") == 0; + const bool use_vulkan = argc == 2 && std::strcmp(argv[1], "--vulkan") == 0; + if (argc > 2 || (argc == 2 && !use_cuda && !use_vulkan)) { + std::fprintf(stderr, "usage: %s [--cuda|--vulkan]\n", argv[0]); + return 1; + } + + ggml_init_params params = { + 1024 * 1024, + nullptr, + true, + }; + ggml_context * ctx = ggml_init(params); + ggml_tensor * x = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, 4, 2); + ggml_tensor * w = ggml_new_tensor_2d(ctx, GGML_TYPE_I8, 4, 3); + ggml_tensor * rot = ggml_regular_hadamard(ctx, x, 4); + ggml_tensor * out = ggml_mul_mat(ctx, w, rot); + ggml_tensor * x256 = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 256); + ggml_tensor * r256 = ggml_regular_hadamard(ctx, x256, 256); + ggml_tensor * w256 = ggml_new_tensor_2d(ctx, GGML_TYPE_I8, 256, 4); + ggml_tensor * ws256 = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 4); + ggml_tensor * b256 = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 4); + ggml_tensor * q256 = ggml_quantize_i8_convrot(ctx, x256, 256); + ggml_tensor * fused256 = ggml_mul_mat_i8_tensorwise(ctx, w256, q256, ws256, b256, 256); + const int large_k = 512; + const int large_n = 128; + const int large_rows = 65; + ggml_tensor * x_large = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, large_k, large_rows); + ggml_tensor * w_large = ggml_new_tensor_2d(ctx, GGML_TYPE_I8, large_k, large_n); + ggml_tensor * ws_large = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, large_n); + ggml_tensor * b_large = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, large_n); + ggml_tensor * q_large = ggml_quantize_i8_convrot(ctx, x_large, 256); + ggml_tensor * fused_large = ggml_mul_mat_i8_tensorwise(ctx, w_large, q_large, ws_large, b_large, 256); + ggml_tensor * fused_large_no_bias = ggml_mul_mat_i8_tensorwise(ctx, w_large, q_large, ws_large, nullptr, 256); + ggml_tensor * fused_large_direct = ggml_mul_mat_i8_tensorwise(ctx, w_large, x_large, ws_large, b_large, 256); + + ggml_cgraph * graph = ggml_new_graph(ctx); + if (!use_vulkan) { + ggml_build_forward_expand(graph, out); + ggml_build_forward_expand(graph, r256); + ggml_build_forward_expand(graph, fused_large_direct); + } + ggml_build_forward_expand(graph, fused256); + ggml_build_forward_expand(graph, fused_large); + ggml_build_forward_expand(graph, fused_large_no_bias); + + ggml_backend_t backend = nullptr; + if (use_cuda || use_vulkan) { + const char * backend_name = use_cuda ? "CUDA" : "Vulkan"; + ggml_backend_load_all(); + for (size_t i = 0; i < ggml_backend_dev_count(); ++i) { + ggml_backend_dev_t device = ggml_backend_dev_get(i); + if (ggml_backend_dev_type(device) == GGML_BACKEND_DEVICE_TYPE_GPU && std::strstr(ggml_backend_dev_name(device), backend_name) != nullptr) { + backend = ggml_backend_dev_init(device, nullptr); + break; + } + } + if (backend == nullptr) { + std::fprintf(stderr, "%s backend not found\n", backend_name); + return 1; + } + } else { + backend = ggml_backend_cpu_init(); + ggml_backend_cpu_set_n_threads(backend, 2); + } + if (!ggml_backend_supports_op(backend, q_large) || + !ggml_backend_supports_op(backend, fused_large) || + !ggml_backend_supports_op(backend, fused_large_no_bias) || + (!use_vulkan && !ggml_backend_supports_op(backend, fused_large_direct))) { + std::fprintf(stderr, "backend does not report packed INT8 convrot matmul support\n"); + return 1; + } + ggml_backend_buffer_t buffer = ggml_backend_alloc_ctx_tensors(ctx, backend); + + const std::vector x_data = { + 1.0f, 2.0f, 3.0f, 4.0f, + -1.0f, 0.5f, 2.0f, -3.0f, + }; + const std::vector w_data = { + 1, 2, 3, 4, + -2, 1, 0, 2, + 127, -127, 64, -64, + }; + std::vector x256_data(256); + std::vector w256_data(256 * 4); + for (int i = 0; i < 256; ++i) { + x256_data[i] = (float) ((i * 17) % 29 - 14) / 7.0f; + for (int output = 0; output < 4; ++output) { + w256_data[output * 256 + i] = (int8_t) ((i * (output + 3) + output * 11) % 31 - 15); + } + } + const std::vector ws256_data = { 0.125f, 0.25f, 0.5f, 1.5f }; + const std::vector b256_data = { -1.0f, 0.5f, 2.0f, -3.0f }; + std::vector x_large_data(large_k * large_rows); + std::vector w_large_data(large_k * large_n); + std::vector ws_large_data(large_n); + std::vector b_large_data(large_n); + for (int row = 0; row < large_rows; ++row) { + for (int i = 0; i < large_k; ++i) { + x_large_data[row * large_k + i] = (float) ((row * 13 + i * 17) % 43 - 21) / 9.0f; + } + } + for (int output = 0; output < large_n; ++output) { + const int pattern = output % 7; + for (int i = 0; i < large_k; ++i) { + w_large_data[output * large_k + i] = (int8_t) ((i * (pattern + 3) + pattern * 11) % 61 - 30); + } + ws_large_data[output] = (float) (output % 11 + 1) / 100.0f; + b_large_data[output] = (float) (output % 9 - 4) / 7.0f; + } + ggml_backend_tensor_set(x, x_data.data(), 0, ggml_nbytes(x)); + ggml_backend_tensor_set(w, w_data.data(), 0, ggml_nbytes(w)); + ggml_backend_tensor_set(x256, x256_data.data(), 0, ggml_nbytes(x256)); + ggml_backend_tensor_set(w256, w256_data.data(), 0, ggml_nbytes(w256)); + ggml_backend_tensor_set(ws256, ws256_data.data(), 0, ggml_nbytes(ws256)); + ggml_backend_tensor_set(b256, b256_data.data(), 0, ggml_nbytes(b256)); + ggml_backend_tensor_set(x_large, x_large_data.data(), 0, ggml_nbytes(x_large)); + ggml_backend_tensor_set(w_large, w_large_data.data(), 0, ggml_nbytes(w_large)); + ggml_backend_tensor_set(ws_large, ws_large_data.data(), 0, ggml_nbytes(ws_large)); + ggml_backend_tensor_set(b_large, b_large_data.data(), 0, ggml_nbytes(b_large)); + if (ggml_backend_graph_compute(backend, graph) != GGML_STATUS_SUCCESS) { + std::fprintf(stderr, "graph compute failed\n"); + return 1; + } + + std::vector rot_data(ggml_nelements(rot)); + std::vector out_data(ggml_nelements(out)); + if (!use_vulkan) { + ggml_backend_tensor_get(rot, rot_data.data(), 0, ggml_nbytes(rot)); + ggml_backend_tensor_get(out, out_data.data(), 0, ggml_nbytes(out)); + } + + const std::vector expected_rot = { + 1.0f, 2.0f, 3.0f, 4.0f, + 2.25f, -2.75f, -1.25f, 0.25f, + }; + if (!use_vulkan) { + for (size_t i = 0; i < expected_rot.size(); ++i) { + if (!nearly_equal(rot_data[i], expected_rot[i])) { + std::fprintf(stderr, "regular Hadamard mismatch at %zu: %.8f != %.8f\n", i, rot_data[i], expected_rot[i]); + return 1; + } + } + } + + std::vector r256_data(ggml_nelements(r256)); + if (!use_vulkan) { + ggml_backend_tensor_get(r256, r256_data.data(), 0, ggml_nbytes(r256)); + } + for (int i = 0; i < 256; ++i) { + const float expected = regular_hadamard_reference(x256_data, i); + if (!use_vulkan && !nearly_equal(r256_data[i], expected)) { + std::fprintf(stderr, "regular Hadamard-256 mismatch at %d: %.8f != %.8f\n", i, r256_data[i], expected); + return 1; + } + r256_data[i] = expected; + } + + float fused_amax = 0.0f; + for (float value : r256_data) { + fused_amax = std::max(fused_amax, std::fabs(value)); + } + const float fused_scale = fused_amax / 127.0f; + std::vector fused_quantized(256); + for (int i = 0; i < 256; ++i) { + int value = (int) std::lrint(r256_data[i] / fused_scale); + value = std::max(-127, std::min(127, value)); + fused_quantized[i] = (int8_t) value; + } + std::vector expected_fused(4); + for (int output = 0; output < 4; ++output) { + int32_t sum = 0; + for (int i = 0; i < 256; ++i) { + sum += (int32_t) w256_data[output * 256 + i] * (int32_t) fused_quantized[i]; + } + expected_fused[output] = (float) sum * fused_scale * ws256_data[output] + b256_data[output]; + } + std::vector fused_data(4); + ggml_backend_tensor_get(fused256, fused_data.data(), 0, ggml_nbytes(fused256)); + for (size_t i = 0; i < expected_fused.size(); ++i) { + if (!nearly_equal(fused_data[i], expected_fused[i])) { + std::fprintf(stderr, "fused INT8 convrot mismatch at %zu: %.8f != %.8f\n", i, fused_data[i], expected_fused[i]); + return 1; + } + } + + std::vector q_large_data(ggml_nbytes(q_large)); + std::vector fused_large_data(ggml_nelements(fused_large)); + std::vector fused_large_no_bias_data(ggml_nelements(fused_large_no_bias)); + std::vector fused_large_direct_data(ggml_nelements(fused_large_direct)); + ggml_backend_tensor_get(q_large, q_large_data.data(), 0, ggml_nbytes(q_large)); + ggml_backend_tensor_get(fused_large, fused_large_data.data(), 0, ggml_nbytes(fused_large)); + ggml_backend_tensor_get(fused_large_no_bias, fused_large_no_bias_data.data(), 0, ggml_nbytes(fused_large_no_bias)); + if (!use_vulkan) { + ggml_backend_tensor_get(fused_large_direct, fused_large_direct_data.data(), 0, ggml_nbytes(fused_large_direct)); + } + const int large_rows_padded = (large_rows + 3) & ~3; + std::vector large_activation_scales(large_rows); + std::memcpy(large_activation_scales.data(), q_large_data.data() + large_k * large_rows_padded, large_rows * sizeof(float)); + for (int row = 0; row < large_rows; ++row) { + int32_t sums[7] = {}; + for (int pattern = 0; pattern < 7; ++pattern) { + for (int i = 0; i < large_k; ++i) { + sums[pattern] += (int32_t) q_large_data[row * large_k + i] * (int32_t) w_large_data[pattern * large_k + i]; + } + } + for (int output = 0; output < large_n; ++output) { + const int32_t sum = sums[output % 7]; + const float expected_no_bias = (float) sum * large_activation_scales[row] * ws_large_data[output]; + const float expected = expected_no_bias + b_large_data[output]; + const size_t index = (size_t) row * large_n + output; + if (!nearly_equal(fused_large_data[index], expected)) { + std::fprintf(stderr, "large fused INT8 convrot mismatch at %zu: %.8f != %.8f\n", index, fused_large_data[index], expected); + return 1; + } + if (!nearly_equal(fused_large_no_bias_data[index], expected_no_bias)) { + std::fprintf(stderr, "large no-bias fused INT8 convrot mismatch at %zu: %.8f != %.8f\n", index, fused_large_no_bias_data[index], expected_no_bias); + return 1; + } + if (!use_vulkan && !nearly_equal(fused_large_direct_data[index], expected)) { + std::fprintf(stderr, "large direct fused INT8 convrot mismatch at %zu: %.8f != %.8f\n", index, fused_large_direct_data[index], expected); + return 1; + } + } + } + + std::vector expected_out(6); + for (int row = 0; row < 2; ++row) { + const float * activation = expected_rot.data() + row * 4; + float amax = 0.0f; + for (int i = 0; i < 4; ++i) { + amax = std::max(amax, std::fabs(activation[i])); + } + const float scale = amax / 127.0f; + int8_t quantized[4]; + for (int i = 0; i < 4; ++i) { + int value = (int) std::lrint(activation[i] / scale); + value = std::max(-127, std::min(127, value)); + quantized[i] = (int8_t) value; + } + for (int output = 0; output < 3; ++output) { + int32_t sum = 0; + for (int i = 0; i < 4; ++i) { + sum += (int32_t) w_data[output * 4 + i] * (int32_t) quantized[i]; + } + expected_out[row * 3 + output] = (float) sum * scale; + } + } + + if (!use_vulkan) { + for (size_t i = 0; i < expected_out.size(); ++i) { + if (!nearly_equal(out_data[i], expected_out[i])) { + std::fprintf(stderr, "INT8 matmul mismatch at %zu: %.8f != %.8f\n", i, out_data[i], expected_out[i]); + return 1; + } + } + } + + ggml_backend_buffer_free(buffer); + ggml_backend_free(backend); + ggml_free(ctx); + std::printf("INT8 convrot %s test passed\n", use_cuda ? "CUDA" : (use_vulkan ? "Vulkan" : "CPU")); + return 0; +} From 1a41cef5feaca23384790fcf743820bd416b4fac Mon Sep 17 00:00:00 2001 From: leejet Date: Thu, 13 Aug 2026 00:22:26 +0800 Subject: [PATCH 2/2] ggml : remove standalone regular Hadamard op --- ggml/include/ggml-rpc.h | 2 +- ggml/include/ggml.h | 7 +-- ggml/src/ggml-backend-meta.cpp | 3 -- ggml/src/ggml-cpu/ggml-cpu.c | 5 -- ggml/src/ggml-cpu/ggml-cpu.cpp | 12 ----- ggml/src/ggml-cpu/ops.cpp | 49 -------------------- ggml/src/ggml-cpu/ops.h | 1 - ggml/src/ggml-cuda/fwht.cu | 70 ---------------------------- ggml/src/ggml-cuda/fwht.cuh | 1 - ggml/src/ggml-cuda/ggml-cuda.cu | 11 ----- ggml/src/ggml.c | 28 ++--------- tests/test-int8-convrot.cpp | 82 +-------------------------------- 12 files changed, 7 insertions(+), 264 deletions(-) diff --git a/ggml/include/ggml-rpc.h b/ggml/include/ggml-rpc.h index f07729d43bb..a88eab44011 100644 --- a/ggml/include/ggml-rpc.h +++ b/ggml/include/ggml-rpc.h @@ -11,7 +11,7 @@ extern "C" { #define RPC_PROTO_PATCH_VERSION 1 #ifdef __cplusplus -static_assert(GGML_OP_COUNT == 103, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION"); +static_assert(GGML_OP_COUNT == 102, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION"); #endif #define GGML_RPC_MAX_SERVERS 16 diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index b0afa547408..70fcc21e063 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -590,7 +590,6 @@ extern "C" { GGML_OP_GLU, - GGML_OP_REGULAR_HADAMARD, GGML_OP_QUANTIZE_I8_CONVROT, GGML_OP_COUNT, @@ -1452,11 +1451,7 @@ extern "C" { struct ggml_tensor * a, enum ggml_op_hint hint); - GGML_API struct ggml_tensor * ggml_regular_hadamard( - struct ggml_context * ctx, - struct ggml_tensor * a, - int group_size); - + // Packs row-wise I8 activations followed by one F32 scale per row. GGML_API struct ggml_tensor * ggml_quantize_i8_convrot( struct ggml_context * ctx, struct ggml_tensor * a, diff --git a/ggml/src/ggml-backend-meta.cpp b/ggml/src/ggml-backend-meta.cpp index ca20b22a825..b0f3b1d3812 100644 --- a/ggml/src/ggml-backend-meta.cpp +++ b/ggml/src/ggml-backend-meta.cpp @@ -1009,9 +1009,6 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state( case GGML_OP_GLU: { split_state = handle_generic(src_ss, /*scalar_only =*/ false); } break; - case GGML_OP_REGULAR_HADAMARD: { - split_state = handle_per_row(src_ss); - } break; case GGML_OP_QUANTIZE_I8_CONVROT: { split_state = handle_generic(src_ss, /*scalar_only =*/ true); } break; diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 1a25c2cdc18..cd598af6e47 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -2180,10 +2180,6 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm { ggml_compute_forward_glu(params, tensor); } break; - case GGML_OP_REGULAR_HADAMARD: - { - ggml_compute_forward_regular_hadamard(params, tensor); - } break; case GGML_OP_QUANTIZE_I8_CONVROT: { ggml_compute_forward_quantize_i8_convrot(params, tensor); @@ -2486,7 +2482,6 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) { case GGML_OP_MUL_MAT: case GGML_OP_MUL_MAT_ID: case GGML_OP_OUT_PROD: - case GGML_OP_REGULAR_HADAMARD: case GGML_OP_QUANTIZE_I8_CONVROT: { n_tasks = n_threads; diff --git a/ggml/src/ggml-cpu/ggml-cpu.cpp b/ggml/src/ggml-cpu/ggml-cpu.cpp index 68835eb2f65..aec18532be8 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.cpp +++ b/ggml/src/ggml-cpu/ggml-cpu.cpp @@ -490,18 +490,6 @@ static bool ggml_backend_cpu_device_supports_op(ggml_backend_dev_t dev, const st } return group_size == 1; } - case GGML_OP_REGULAR_HADAMARD: { - int group_size = ggml_get_op_params_i32(op, 0); - if (src0->type != GGML_TYPE_F32 || op->type != GGML_TYPE_F32 || - !ggml_is_contiguous(src0) || !ggml_is_contiguous(op) || - group_size <= 0 || src0->ne[0] % group_size != 0) { - return false; - } - while (group_size > 1 && group_size % 4 == 0) { - group_size /= 4; - } - return group_size == 1; - } case GGML_OP_SOFT_MAX_BACK: { if (op->src[0]->type != GGML_TYPE_F32 || op->src[1]->type != GGML_TYPE_F32) { return false; diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index c5d11f67725..fbbab2b47d8 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -11919,55 +11919,6 @@ void ggml_compute_forward_fwht(const ggml_compute_params * params, ggml_tensor * } } -void ggml_compute_forward_regular_hadamard(const ggml_compute_params * params, ggml_tensor * dst) { - const ggml_tensor * src = dst->src[0]; - - GGML_ASSERT(src->type == GGML_TYPE_F32); - GGML_ASSERT(dst->type == GGML_TYPE_F32); - GGML_ASSERT(ggml_is_contiguous(src)); - GGML_ASSERT(ggml_is_contiguous(dst)); - - const int group_size = ggml_get_op_params_i32(dst, 0); - GGML_ASSERT(group_size > 0 && src->ne[0] % group_size == 0); - - const int64_t group_count = ggml_nelements(src) / group_size; - const int64_t groups_per_thread = (group_count + params->nth - 1) / params->nth; - const int64_t start_group = params->ith * groups_per_thread; - const int64_t end_group = MIN(start_group + groups_per_thread, group_count); - const float scale = 1.0f / sqrtf((float) group_size); - - const float * src_data = (const float *) src->data; - float * dst_data = (float *) dst->data; - - for (int64_t group = start_group; group < end_group; ++group) { - const float * src_group = src_data + group * group_size; - float * dst_group = dst_data + group * group_size; - - for (int i = 0; i < group_size; ++i) { - dst_group[i] = src_group[i] * scale; - } - - for (int stride = 1; stride < group_size; stride *= 4) { - for (int base = 0; base < group_size; base += 4 * stride) { - for (int j = 0; j < stride; ++j) { - const int i0 = base + j; - const int i1 = i0 + stride; - const int i2 = i1 + stride; - const int i3 = i2 + stride; - const float a = dst_group[i0]; - const float b = dst_group[i1]; - const float c = dst_group[i2]; - const float d = dst_group[i3]; - dst_group[i0] = a + b + c - d; - dst_group[i1] = a + b - c + d; - dst_group[i2] = a - b + c + d; - dst_group[i3] = -a + b + c + d; - } - } - } - } -} - void ggml_compute_forward_quantize_i8_convrot(const ggml_compute_params * params, ggml_tensor * dst) { const ggml_tensor * src = dst->src[0]; diff --git a/ggml/src/ggml-cpu/ops.h b/ggml/src/ggml-cpu/ops.h index 66bcf913d6c..3b57916798f 100644 --- a/ggml/src/ggml-cpu/ops.h +++ b/ggml/src/ggml-cpu/ops.h @@ -118,7 +118,6 @@ void ggml_compute_forward_cross_entropy_loss_back(const struct ggml_compute_para void ggml_compute_forward_opt_step_adamw(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_mul_mat(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_fwht(const struct ggml_compute_params * params, struct ggml_tensor * dst); -void ggml_compute_forward_regular_hadamard(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_quantize_i8_convrot(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_opt_step_sgd(const struct ggml_compute_params * params, struct ggml_tensor * dst); #ifdef __cplusplus diff --git a/ggml/src/ggml-cuda/fwht.cu b/ggml/src/ggml-cuda/fwht.cu index 7a21ca0fc40..184dc254c72 100644 --- a/ggml/src/ggml-cuda/fwht.cu +++ b/ggml/src/ggml-cuda/fwht.cu @@ -58,44 +58,6 @@ __global__ void fwht_cuda(const float * src, float * dst, const int64_t n_rows, } } -template -__global__ void regular_hadamard_cuda(const float * src, float * dst, const int64_t n_groups, const float scale) { - const int64_t group = blockIdx.x; - const int tid = threadIdx.x; - if (group >= n_groups) { - return; - } - - __shared__ float values[N]; - src += group * N; - dst += group * N; - - values[tid] = src[tid] * scale; - __syncthreads(); - -#pragma unroll - for (int stride = 1; stride < N; stride *= 4) { - if (tid < N / 4) { - const int base = (tid / stride) * 4 * stride + tid % stride; - const int i0 = base; - const int i1 = i0 + stride; - const int i2 = i1 + stride; - const int i3 = i2 + stride; - const float a = values[i0]; - const float b = values[i1]; - const float c = values[i2]; - const float d = values[i3]; - values[i0] = a + b + c - d; - values[i1] = a + b - c + d; - values[i2] = a - b + c + d; - values[i3] = -a + b + c + d; - } - __syncthreads(); - } - - dst[tid] = values[tid]; -} - bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst) { GGML_ASSERT(ggml_are_same_shape(src, dst)); if (!ggml_is_contiguous(src) || !ggml_is_contiguous(dst)) { @@ -137,35 +99,3 @@ bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src, return false; } } - -bool ggml_cuda_op_regular_hadamard(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst, int group_size) { - GGML_ASSERT(ggml_are_same_shape(src, dst)); - if (src->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32 || - !ggml_is_contiguous(src) || !ggml_is_contiguous(dst) || - group_size <= 0 || ggml_nelements(src) % group_size != 0) { - return false; - } - - const int64_t n_groups = ggml_nelements(src) / group_size; - const float scale = 1.0f / sqrtf((float) group_size); - const float * src_d = (const float *) src->data; - float * dst_d = (float *) dst->data; - cudaStream_t stream = ctx.stream(); - - switch (group_size) { - case 4: - regular_hadamard_cuda<4><<>>(src_d, dst_d, n_groups, scale); - return true; - case 16: - regular_hadamard_cuda<16><<>>(src_d, dst_d, n_groups, scale); - return true; - case 64: - regular_hadamard_cuda<64><<>>(src_d, dst_d, n_groups, scale); - return true; - case 256: - regular_hadamard_cuda<256><<>>(src_d, dst_d, n_groups, scale); - return true; - default: - return false; - } -} diff --git a/ggml/src/ggml-cuda/fwht.cuh b/ggml/src/ggml-cuda/fwht.cuh index d4d28fad670..cf3df94cafa 100644 --- a/ggml/src/ggml-cuda/fwht.cuh +++ b/ggml/src/ggml-cuda/fwht.cuh @@ -2,4 +2,3 @@ // Returns whether the Fast Walsh-Hadamard transform could be used. bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst); -bool ggml_cuda_op_regular_hadamard(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst, int group_size); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index aa15717dbe3..68cb8be2c8a 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -2607,9 +2607,6 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg case GGML_OP_MUL_MAT_ID: ggml_cuda_mul_mat_id(ctx, dst); break; - case GGML_OP_REGULAR_HADAMARD: - GGML_ASSERT(ggml_cuda_op_regular_hadamard(ctx, dst->src[0], dst, ggml_get_op_params_i32(dst, 0))); - break; case GGML_OP_QUANTIZE_I8_CONVROT: #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) GGML_ASSERT(ggml_cuda_op_quantize_i8_convrot(ctx, dst->src[0], dst, ggml_get_op_params_i32(dst, 0))); @@ -5687,14 +5684,6 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_OP_DIAG: case GGML_OP_SOLVE_TRI: return true; - case GGML_OP_REGULAR_HADAMARD: { - const ggml_tensor * src = op->src[0]; - const int group_size = ggml_get_op_params_i32(op, 0); - return src != nullptr && src->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 && - ggml_is_contiguous(src) && ggml_is_contiguous(op) && group_size > 0 && - src->ne[0] % group_size == 0 && - (group_size == 4 || group_size == 16 || group_size == 64 || group_size == 256); - } case GGML_OP_QUANTIZE_I8_CONVROT: #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) { diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index b9fe5116574..e8e603361d2 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -1099,11 +1099,10 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = { "GLU", - "REGULAR_HADAMARD", "QUANTIZE_I8_CONVROT", }; -static_assert(GGML_OP_COUNT == 103, "GGML_OP_COUNT != 103"); +static_assert(GGML_OP_COUNT == 102, "GGML_OP_COUNT != 102"); static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "none", @@ -1217,11 +1216,10 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "glu(x)", - "regular_hadamard(x)", "quantize_i8_convrot(x)", }; -static_assert(GGML_OP_COUNT == 103, "GGML_OP_COUNT != 103"); +static_assert(GGML_OP_COUNT == 102, "GGML_OP_COUNT != 102"); static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2"); @@ -3298,7 +3296,7 @@ struct ggml_tensor * ggml_mul_mat( return result; } -struct ggml_tensor * ggml_mul_mat_i8_tensorwise( +GGML_API struct ggml_tensor * ggml_mul_mat_i8_tensorwise( struct ggml_context * ctx, struct ggml_tensor * weight, struct ggml_tensor * input, @@ -3363,26 +3361,6 @@ void ggml_mul_mat_set_hint( ggml_set_op_params_i32(a, 1, hint_i32); } -struct ggml_tensor * ggml_regular_hadamard( - struct ggml_context * ctx, - struct ggml_tensor * a, - int group_size) { - GGML_ASSERT(a->type == GGML_TYPE_F32); - GGML_ASSERT(group_size > 0 && a->ne[0] % group_size == 0); - - int n = group_size; - while (n > 1 && n % 4 == 0) { - n /= 4; - } - GGML_ASSERT(n == 1); - - struct ggml_tensor * result = ggml_dup_tensor(ctx, a); - result->op = GGML_OP_REGULAR_HADAMARD; - result->src[0] = a; - ggml_set_op_params_i32(result, 0, group_size); - return result; -} - struct ggml_tensor * ggml_quantize_i8_convrot( struct ggml_context * ctx, struct ggml_tensor * a, diff --git a/tests/test-int8-convrot.cpp b/tests/test-int8-convrot.cpp index cfdfc33f684..d70c4d61261 100644 --- a/tests/test-int8-convrot.cpp +++ b/tests/test-int8-convrot.cpp @@ -50,12 +50,7 @@ int main(int argc, char ** argv) { true, }; ggml_context * ctx = ggml_init(params); - ggml_tensor * x = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, 4, 2); - ggml_tensor * w = ggml_new_tensor_2d(ctx, GGML_TYPE_I8, 4, 3); - ggml_tensor * rot = ggml_regular_hadamard(ctx, x, 4); - ggml_tensor * out = ggml_mul_mat(ctx, w, rot); ggml_tensor * x256 = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 256); - ggml_tensor * r256 = ggml_regular_hadamard(ctx, x256, 256); ggml_tensor * w256 = ggml_new_tensor_2d(ctx, GGML_TYPE_I8, 256, 4); ggml_tensor * ws256 = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 4); ggml_tensor * b256 = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 4); @@ -75,8 +70,6 @@ int main(int argc, char ** argv) { ggml_cgraph * graph = ggml_new_graph(ctx); if (!use_vulkan) { - ggml_build_forward_expand(graph, out); - ggml_build_forward_expand(graph, r256); ggml_build_forward_expand(graph, fused_large_direct); } ggml_build_forward_expand(graph, fused256); @@ -111,15 +104,6 @@ int main(int argc, char ** argv) { } ggml_backend_buffer_t buffer = ggml_backend_alloc_ctx_tensors(ctx, backend); - const std::vector x_data = { - 1.0f, 2.0f, 3.0f, 4.0f, - -1.0f, 0.5f, 2.0f, -3.0f, - }; - const std::vector w_data = { - 1, 2, 3, 4, - -2, 1, 0, 2, - 127, -127, 64, -64, - }; std::vector x256_data(256); std::vector w256_data(256 * 4); for (int i = 0; i < 256; ++i) { @@ -147,8 +131,6 @@ int main(int argc, char ** argv) { ws_large_data[output] = (float) (output % 11 + 1) / 100.0f; b_large_data[output] = (float) (output % 9 - 4) / 7.0f; } - ggml_backend_tensor_set(x, x_data.data(), 0, ggml_nbytes(x)); - ggml_backend_tensor_set(w, w_data.data(), 0, ggml_nbytes(w)); ggml_backend_tensor_set(x256, x256_data.data(), 0, ggml_nbytes(x256)); ggml_backend_tensor_set(w256, w256_data.data(), 0, ggml_nbytes(w256)); ggml_backend_tensor_set(ws256, ws256_data.data(), 0, ggml_nbytes(ws256)); @@ -162,37 +144,9 @@ int main(int argc, char ** argv) { return 1; } - std::vector rot_data(ggml_nelements(rot)); - std::vector out_data(ggml_nelements(out)); - if (!use_vulkan) { - ggml_backend_tensor_get(rot, rot_data.data(), 0, ggml_nbytes(rot)); - ggml_backend_tensor_get(out, out_data.data(), 0, ggml_nbytes(out)); - } - - const std::vector expected_rot = { - 1.0f, 2.0f, 3.0f, 4.0f, - 2.25f, -2.75f, -1.25f, 0.25f, - }; - if (!use_vulkan) { - for (size_t i = 0; i < expected_rot.size(); ++i) { - if (!nearly_equal(rot_data[i], expected_rot[i])) { - std::fprintf(stderr, "regular Hadamard mismatch at %zu: %.8f != %.8f\n", i, rot_data[i], expected_rot[i]); - return 1; - } - } - } - - std::vector r256_data(ggml_nelements(r256)); - if (!use_vulkan) { - ggml_backend_tensor_get(r256, r256_data.data(), 0, ggml_nbytes(r256)); - } + std::vector r256_data(256); for (int i = 0; i < 256; ++i) { - const float expected = regular_hadamard_reference(x256_data, i); - if (!use_vulkan && !nearly_equal(r256_data[i], expected)) { - std::fprintf(stderr, "regular Hadamard-256 mismatch at %d: %.8f != %.8f\n", i, r256_data[i], expected); - return 1; - } - r256_data[i] = expected; + r256_data[i] = regular_hadamard_reference(x256_data, i); } float fused_amax = 0.0f; @@ -263,38 +217,6 @@ int main(int argc, char ** argv) { } } - std::vector expected_out(6); - for (int row = 0; row < 2; ++row) { - const float * activation = expected_rot.data() + row * 4; - float amax = 0.0f; - for (int i = 0; i < 4; ++i) { - amax = std::max(amax, std::fabs(activation[i])); - } - const float scale = amax / 127.0f; - int8_t quantized[4]; - for (int i = 0; i < 4; ++i) { - int value = (int) std::lrint(activation[i] / scale); - value = std::max(-127, std::min(127, value)); - quantized[i] = (int8_t) value; - } - for (int output = 0; output < 3; ++output) { - int32_t sum = 0; - for (int i = 0; i < 4; ++i) { - sum += (int32_t) w_data[output * 4 + i] * (int32_t) quantized[i]; - } - expected_out[row * 3 + output] = (float) sum * scale; - } - } - - if (!use_vulkan) { - for (size_t i = 0; i < expected_out.size(); ++i) { - if (!nearly_equal(out_data[i], expected_out[i])) { - std::fprintf(stderr, "INT8 matmul mismatch at %zu: %.8f != %.8f\n", i, out_data[i], expected_out[i]); - return 1; - } - } - } - ggml_backend_buffer_free(buffer); ggml_backend_free(backend); ggml_free(ctx);