upstream-archive/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-10-02 10:57:33 -05:00
Code Issues Packages Projects Releases Wiki Activity

Fix odd-rows for convert/getrows, possible with FP8

Browse Source
This commit is contained in:
Oliver Simons
2026-10-01 18:49:36 +02:00
parent 7c53b9683f
commit f2ec3f1caa
3 changed files with 19 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/convert.cu
+8
View File
@@ -29,6 +29,14 @@ static __global__ void dequantize_block(const void * __restrict__ vx, dst_t * __
const int64_t iybs = i00 - i00%qk; // y block start index
const int64_t y_offset = qr == 1 ? 1 : qk/2;
if constexpr (qk == 1) {
if (i00 + 1 == ne00) {
const int64_t iy0 = (i0203*ne01 + i01)*ne00 + i00;
y[iy0] = ggml_cuda_cast<dst_t>(ggml_cuda_f8_e4m3_to_fp32(((const uint8_t *) vx)[ib]));
continue;
}
}
// dequantize
float2 v;
dequantize_kernel(vx, ib, iqs, v);
ggml/src/ggml-cuda/getrows.cu
+8 -1
View File
@@ -30,6 +30,13 @@ static __global__ void k_get_rows(
const int iybs = i00 - i00%qk; // dst block start index
const int y_offset = qr == 1 ? 1 : qk/2;
if constexpr (qk == 1) {
if (i00 + 1 == ne00) {
dst_row[i00] = ggml_cuda_cast<dst_t>(ggml_cuda_f8_e4m3_to_fp32(((const uint8_t *) src0_row)[i00]));
continue;
}
}
// dequantize
float2 v;
dequantize_kernel(src0_row, ib, iqs, v);
@@ -178,7 +185,7 @@ static void get_rows_cuda_q(
const size_t s12 = nb12 / sizeof(int32_t);
// const size_t s13 = nb13 / sizeof(int32_t);
GGML_ASSERT(ne00 % 2 == 0);
GGML_ASSERT(qk == 1 || ne00 % 2 == 0);
GGML_ASSERT(ne12 > 0);
GGML_ASSERT(ne11 <= std::numeric_limits<uint32_t>::max() / ne12);
tests/test-backend-ops.cpp
+3
View File
@@ -9376,6 +9376,9 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
}
test_cases.emplace_back(new test_get_rows(GGML_TYPE_F32, 1, 8, 2, 1, 1, false));
for (int n : {1, 3, 257}) {
test_cases.emplace_back(new test_get_rows(GGML_TYPE_F8_E4M3, n, 5, 3, 1, 1, false));
}
for (ggml_type type : all_types) {
for (int b : {1, 7}) {
for (bool v : {false, true}) {