mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-11-13 14:27:38 -06:00
Compare commits
| Author | SHA1 | Date | |
|---|---|---|---|
 AMIR
|
5193ef1d4f | ||
|
Max Krasnyansky
|
4df9b47976 | ||
|
thelittlefireman
|
b79696f1f5 | ||
|
Pascal
|
4f9dfa0e36 | ||
|
Max Krasnyansky
|
9daa0aa22b | ||
|
ynankani
|
fe14cad51b | ||
|
anujj
|
31953bcf74 | ||
|
Yonghye Kwon
|
ae93877c8f |
+1
-1
@@ -126,7 +126,7 @@ Legend:
|
||||
| TANH | ❌ | ✅ | ✅ | 🟡 | 🟡 | 🟡 | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
| TIMESTEP_EMBEDDING | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ | ❌ |
|
||||
| TOP_K | ❌ | ❌ | ✅ | ❌ | ❌ | ✅ | ✅ | ✅ | ✅ | 🟡 | ✅ | ❌ | ❌ |
|
||||
| TOP_K_MASK | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
|
||||
| TOP_K_MASK | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ |
|
||||
| TRI | ❌ | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
| TRUNC | ❌ | ❌ | ✅ | 🟡 | 🟡 | ❌ | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
| UPSCALE | ❌ | 🟡 | ✅ | ✅ | ❌ | ❌ | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
|
||||
@@ -19439,3 +19439,31 @@
|
||||
"Vulkan0","LIGHTNING_INDEXER","hsk=128,nh=64,kv=256,nb=512,ns=4,nm=1,type_K=q4_1","support","0","no","Vulkan"
|
||||
"Vulkan0","LIGHTNING_INDEXER","hsk=128,nh=64,kv=256,nb=512,ns=4,nm=1,type_K=q4_0","support","0","no","Vulkan"
|
||||
"Vulkan0","LIGHTNING_INDEXER","hsk=128,nh=64,kv=256,nb=512,ns=4,nm=1,type_K=iq4_nl","support","0","no","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=256,n_blk=4,n_pick=64,n_top=16,n_extra=3,score=1,nb=32,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=1001,n_blk=4,n_pick=70,n_top=16,n_extra=3,score=1,nb=7,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=256,n_blk=4,n_pick=64,n_top=64,n_extra=0,score=1,nb=33,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=512,n_blk=8,n_pick=16,n_top=32,n_extra=7,score=1,nb=5,nh=1,ns=4,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=4096,n_blk=4,n_pick=1024,n_top=512,n_extra=3,score=1,nb=1,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=4096,n_blk=4,n_pick=1024,n_top=512,n_extra=3,score=1,nb=64,nh=1,ns=2,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=256,n_blk=0,n_pick=256,n_top=16,n_extra=0,score=0,nb=32,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=1001,n_blk=0,n_pick=1001,n_top=64,n_extra=0,score=0,nb=7,nh=1,ns=3,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=4096,n_blk=0,n_pick=4096,n_top=2048,n_extra=0,score=0,nb=5,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=512,n_blk=0,n_pick=512,n_top=32,n_extra=5,score=1,nb=9,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=512,n_blk=16,n_pick=32,n_top=4,n_extra=0,score=1,nb=6,nh=4,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=1024,n_blk=8,n_pick=128,n_top=16,n_extra=8,score=0,nb=3,nh=2,ns=2,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=1001,n_blk=4,n_pick=70,n_top=16,n_extra=3,score=1,nb=7,nh=1,ns=1,real=0,bcast=1","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f16,n_kv=512,n_blk=8,n_pick=16,n_top=32,n_extra=7,score=1,nb=5,nh=1,ns=4,real=0,bcast=1","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=256,n_blk=4,n_pick=64,n_top=16,n_extra=3,score=1,nb=32,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=1001,n_blk=4,n_pick=70,n_top=16,n_extra=3,score=1,nb=7,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=256,n_blk=4,n_pick=64,n_top=64,n_extra=0,score=1,nb=33,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=512,n_blk=8,n_pick=16,n_top=32,n_extra=7,score=1,nb=5,nh=1,ns=4,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=4096,n_blk=4,n_pick=1024,n_top=512,n_extra=3,score=1,nb=1,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=4096,n_blk=4,n_pick=1024,n_top=512,n_extra=3,score=1,nb=64,nh=1,ns=2,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=256,n_blk=0,n_pick=256,n_top=16,n_extra=0,score=0,nb=32,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=1001,n_blk=0,n_pick=1001,n_top=64,n_extra=0,score=0,nb=7,nh=1,ns=3,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=4096,n_blk=0,n_pick=4096,n_top=2048,n_extra=0,score=0,nb=5,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=512,n_blk=0,n_pick=512,n_top=32,n_extra=5,score=1,nb=9,nh=1,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=512,n_blk=16,n_pick=32,n_top=4,n_extra=0,score=1,nb=6,nh=4,ns=1,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=1024,n_blk=8,n_pick=128,n_top=16,n_extra=8,score=0,nb=3,nh=2,ns=2,real=0,bcast=0","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=1001,n_blk=4,n_pick=70,n_top=16,n_extra=3,score=1,nb=7,nh=1,ns=1,real=0,bcast=1","support","1","yes","Vulkan"
|
||||
"Vulkan0","TOP_K_MASK","type=f32,n_kv=512,n_blk=8,n_pick=16,n_top=32,n_extra=7,score=1,nb=5,nh=1,ns=4,real=0,bcast=1","support","1","yes","Vulkan"
|
||||
|
||||
|
Can't render this file because it is too large.
|
@@ -2601,6 +2601,9 @@ static bool ggml_backend_cann_supports_op(ggml_backend_dev_t dev, const ggml_ten
|
||||
}
|
||||
case GGML_OP_SUM:
|
||||
return ggml_is_contiguous_rows(op->src[0]);
|
||||
case GGML_OP_RMS_NORM:
|
||||
// the cached gamma of ones does not track its type, so F32 and BF16 can not share it
|
||||
return op->src[0]->type != GGML_TYPE_BF16;
|
||||
case GGML_OP_L2_NORM:
|
||||
case GGML_OP_CROSS_ENTROPY_LOSS:
|
||||
case GGML_OP_DUP:
|
||||
@@ -2617,7 +2620,6 @@ static bool ggml_backend_cann_supports_op(ggml_backend_dev_t dev, const ggml_ten
|
||||
case GGML_OP_SUB:
|
||||
case GGML_OP_MUL:
|
||||
case GGML_OP_DIV:
|
||||
case GGML_OP_RMS_NORM:
|
||||
case GGML_OP_SQR:
|
||||
case GGML_OP_SQRT:
|
||||
case GGML_OP_CLAMP:
|
||||
|
||||
@@ -4144,6 +4144,50 @@ static void ggml_compute_forward_rms_norm_f32(
|
||||
}
|
||||
}
|
||||
|
||||
static void ggml_compute_forward_rms_norm_bf16(
|
||||
const ggml_compute_params * params,
|
||||
ggml_tensor * dst) {
|
||||
|
||||
const ggml_tensor * src0 = dst->src[0];
|
||||
|
||||
GGML_ASSERT(ggml_are_same_shape(src0, dst));
|
||||
GGML_ASSERT(src0->nb[0] == sizeof(ggml_bf16_t));
|
||||
|
||||
const int ith = params->ith;
|
||||
const int nth = params->nth;
|
||||
|
||||
GGML_TENSOR_UNARY_OP_LOCALS
|
||||
|
||||
float eps;
|
||||
memcpy(&eps, dst->op_params, sizeof(float));
|
||||
GGML_ASSERT(eps >= 0.0f);
|
||||
|
||||
for (int64_t i03 = 0; i03 < ne03; i03++) {
|
||||
for (int64_t i02 = 0; i02 < ne02; i02++) {
|
||||
for (int64_t i01 = ith; i01 < ne01; i01 += nth) {
|
||||
const ggml_bf16_t * x = (ggml_bf16_t *) ((char *) src0->data + i01*nb01 + i02*nb02 + i03*nb03);
|
||||
ggml_bf16_t * y = (ggml_bf16_t *) ((char *) dst->data + i01*nb1 + i02*nb2 + i03*nb3);
|
||||
|
||||
ggml_float sum = 0.0;
|
||||
for (int64_t i00 = 0; i00 < ne00; i00++) {
|
||||
const float v = GGML_BF16_TO_FP32(x[i00]);
|
||||
sum += (ggml_float)(v*v);
|
||||
}
|
||||
|
||||
const float mean = sum/ne00;
|
||||
const float scale = 1.0f/sqrtf(mean + eps);
|
||||
|
||||
// if you hit this, likely you got an inf somewhere earlier
|
||||
assert(scale > 0.0f);
|
||||
|
||||
for (int64_t i00 = 0; i00 < ne00; i00++) {
|
||||
y[i00] = GGML_FP32_TO_BF16(GGML_BF16_TO_FP32(x[i00])*scale);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
void ggml_compute_forward_rms_norm(
|
||||
const ggml_compute_params * params,
|
||||
ggml_tensor * dst) {
|
||||
@@ -4155,6 +4199,10 @@ void ggml_compute_forward_rms_norm(
|
||||
{
|
||||
ggml_compute_forward_rms_norm_f32<GGML_RMS_NORM_FUSE_OP_NONE>(params, dst);
|
||||
} break;
|
||||
case GGML_TYPE_BF16:
|
||||
{
|
||||
ggml_compute_forward_rms_norm_bf16(params, dst);
|
||||
} break;
|
||||
default:
|
||||
{
|
||||
GGML_ABORT("fatal error");
|
||||
|
||||
@@ -995,7 +995,8 @@ class tensor_traits_common : public tensor_traits_base {
|
||||
spacemit_kernels::rvv::forward_rms_norm_f32(params, op);
|
||||
return true;
|
||||
default:
|
||||
GGML_ABORT("fatal error");
|
||||
ggml_compute_forward_rms_norm(params, op);
|
||||
return true;
|
||||
}
|
||||
case GGML_OP_ADD:
|
||||
switch (op->src[0]->type) {
|
||||
|
||||
@@ -29,6 +29,7 @@ if (CUDAToolkit_FOUND)
|
||||
# 86 == RTX 3000, needs CUDA v11.1
|
||||
# 89 == RTX 4000, needs CUDA v11.8
|
||||
# 90 == Hopper H100/200, needs CUDA v11.8
|
||||
# 103 == Blackwell GB300, needs CUDA v12.9
|
||||
# 120 == Blackwell, needs CUDA v12.8, FP4 tensor cores
|
||||
#
|
||||
# XX-virtual == compile CUDA code as PTX, do JIT compilation to binary code on first run
|
||||
@@ -64,7 +65,7 @@ if (CUDAToolkit_FOUND)
|
||||
list(APPEND CMAKE_CUDA_ARCHITECTURES 120a-real)
|
||||
endif()
|
||||
if (CUDAToolkit_VERSION VERSION_GREATER_EQUAL "12.9")
|
||||
list(APPEND CMAKE_CUDA_ARCHITECTURES 121a-real)
|
||||
list(APPEND CMAKE_CUDA_ARCHITECTURES 103-real 121a-real)
|
||||
endif()
|
||||
endif()
|
||||
endif()
|
||||
|
||||
@@ -4165,22 +4165,25 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph,
|
||||
add = cgraph->nodes[node_idx+2];
|
||||
}
|
||||
|
||||
GGML_ASSERT(rms_norm->src[0]->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(rms_norm->type == GGML_TYPE_F32);
|
||||
|
||||
//rms norm only supports F32
|
||||
if (mul->src[0]->type != GGML_TYPE_F32 ||
|
||||
mul->src[1]->type != GGML_TYPE_F32 ||
|
||||
mul->type != GGML_TYPE_F32) {
|
||||
// the mul operand is a norm weight, which model files store as F32;
|
||||
// the add operand is an F32 bias or a residual with the type of x
|
||||
const ggml_type type = rms_norm->type;
|
||||
if ((type != GGML_TYPE_F32 && type != GGML_TYPE_BF16) || rms_norm->src[0]->type != type) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (add && (add->src[0]->type != GGML_TYPE_F32 ||
|
||||
add->src[1]->type != GGML_TYPE_F32 ||
|
||||
add->type != GGML_TYPE_F32) ) {
|
||||
const ggml_tensor * mul_src = mul->src[0] == rms_norm ? mul->src[1] : mul->src[0];
|
||||
if (mul_src->type != GGML_TYPE_F32 || mul->type != type) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (add) {
|
||||
const ggml_tensor * add_src = add->src[0] == mul ? add->src[1] : add->src[0];
|
||||
if ((add_src->type != GGML_TYPE_F32 && add_src->type != type) || add->type != type) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
//if rms norm is the B operand, then we don't handle broadcast
|
||||
if (rms_norm == mul->src[1] && !ggml_are_same_shape(mul->src[0], rms_norm)) {
|
||||
return false;
|
||||
@@ -4202,8 +4205,9 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph,
|
||||
const ggml_tensor * rms_norm = cgraph->nodes[node_idx];
|
||||
const ggml_tensor * scale = cgraph->nodes[node_idx+1];
|
||||
|
||||
GGML_ASSERT(rms_norm->src[0]->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(rms_norm->type == GGML_TYPE_F32);
|
||||
if (rms_norm->src[0]->type != GGML_TYPE_F32 || rms_norm->type != GGML_TYPE_F32) {
|
||||
return false;
|
||||
}
|
||||
|
||||
float bias;
|
||||
memcpy(&bias, (const float *) scale->op_params + 1, sizeof(float));
|
||||
@@ -6524,9 +6528,11 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
|
||||
case GGML_OP_SILU_BACK:
|
||||
return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32;
|
||||
case GGML_OP_NORM:
|
||||
case GGML_OP_RMS_NORM:
|
||||
case GGML_OP_L2_NORM:
|
||||
return ggml_is_contiguous_rows(op->src[0]);
|
||||
case GGML_OP_RMS_NORM:
|
||||
return (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_BF16) && op->type == op->src[0]->type &&
|
||||
ggml_is_contiguous_rows(op->src[0]);
|
||||
case GGML_OP_RMS_NORM_BACK:
|
||||
return ggml_is_contiguous(op->src[0]);
|
||||
case GGML_OP_NONE:
|
||||
|
||||
@@ -1,31 +1,36 @@
|
||||
static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_gcn(ggml_type type, int J, bool fallback) {
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
@@ -37,8 +42,11 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 512, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
@@ -50,43 +58,45 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
@@ -95,7 +105,10 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
|
||||
CASE(GGML_TYPE_Q2_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_K, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_K, 512, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_K, 512, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_K, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_K, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_K, 512, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_K, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, false);
|
||||
@@ -108,18 +121,20 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q3_K, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q3_K, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q3_K, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
@@ -131,6 +146,7 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_K, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
@@ -143,7 +159,11 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 56, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
@@ -151,126 +171,151 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 56, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
// CASE(GGML_TYPE_IQ1_S, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
// CASE(GGML_TYPE_IQ1_S, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
// CASE(GGML_TYPE_IQ2_XXS, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
// CASE(GGML_TYPE_IQ2_XXS, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
// CASE(GGML_TYPE_IQ2_XS, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
// CASE(GGML_TYPE_IQ2_XS, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
// CASE(GGML_TYPE_IQ2_S, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
// CASE(GGML_TYPE_IQ2_S, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
// CASE(GGML_TYPE_IQ3_XXS, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
// CASE(GGML_TYPE_IQ3_XXS, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
// CASE(GGML_TYPE_IQ3_S, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
// CASE(GGML_TYPE_IQ3_S, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_NVFP4, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
|
||||
+116
-64
@@ -1,4 +1,5 @@
|
||||
#include "norm.cuh"
|
||||
#include "convert.cuh"
|
||||
#include <cstdint>
|
||||
|
||||
template <int block_size>
|
||||
@@ -81,9 +82,35 @@ static __global__ void group_norm_f32(const float * x, float * dst, const int gr
|
||||
}
|
||||
}
|
||||
|
||||
template <int block_size, int vpt, bool do_vec, bool do_multiply = false, bool do_add = false, bool do_scale = false>
|
||||
static __global__ void rms_norm_f32(const float * x,
|
||||
float * dst,
|
||||
// 4 consecutive values of x as float4, x + 4*i is aligned to 4 values
|
||||
template <typename T>
|
||||
static __device__ __forceinline__ float4 rms_norm_load4(const T * x, const int i) {
|
||||
if constexpr (std::is_same_v<T, nv_bfloat16>) {
|
||||
__align__(8) nv_bfloat162 v[2];
|
||||
ggml_cuda_memcpy_1<sizeof(v)>(v, x + 4*i);
|
||||
const float2 lo = ggml_cuda_cast<float2>(v[0]);
|
||||
const float2 hi = ggml_cuda_cast<float2>(v[1]);
|
||||
return make_float4(lo.x, lo.y, hi.x, hi.y);
|
||||
} else {
|
||||
return ((const float4 *) x)[i];
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static __device__ __forceinline__ void rms_norm_store4(T * dst, const int i, const float4 r) {
|
||||
if constexpr (std::is_same_v<T, nv_bfloat16>) {
|
||||
const __align__(8) nv_bfloat16 v[4] = {
|
||||
ggml_cuda_cast<nv_bfloat16>(r.x), ggml_cuda_cast<nv_bfloat16>(r.y), ggml_cuda_cast<nv_bfloat16>(r.z), ggml_cuda_cast<nv_bfloat16>(r.w) };
|
||||
ggml_cuda_memcpy_1<sizeof(v)>(dst + 4*i, v);
|
||||
} else {
|
||||
((float4 *) dst)[i] = r;
|
||||
}
|
||||
}
|
||||
|
||||
// x and dst are float or nv_bfloat16, mul is float, add is float or the type of x
|
||||
template <typename T, int block_size, int vpt, bool do_vec, bool do_multiply = false, bool do_add = false, bool do_scale = false, typename T_add = float>
|
||||
static __global__ void rms_norm_f32(const T * x,
|
||||
T * dst,
|
||||
const int ncols,
|
||||
const int nchannels,
|
||||
const int nsamples,
|
||||
@@ -99,7 +126,7 @@ static __global__ void rms_norm_f32(const float * x,
|
||||
const uint3 mul_nrows_packed = make_uint3(0, 0, 0),
|
||||
const uint3 mul_nchannels_packed = make_uint3(0, 0, 0),
|
||||
const uint3 mul_nsamples_packed = make_uint3(0, 0, 0),
|
||||
const float * add = nullptr,
|
||||
const T_add * add = nullptr,
|
||||
const int64_t add_stride_row = 0,
|
||||
const int64_t add_stride_channel = 0,
|
||||
const int64_t add_stride_sample = 0,
|
||||
@@ -123,8 +150,8 @@ static __global__ void rms_norm_f32(const float * x,
|
||||
// grid.y and grid.z are clamped to the CUDA limit, iterate over the excess channels/samples
|
||||
for (int sample = blockIdx.z; sample < nsamples; sample += gridDim.z) {
|
||||
for (int channel = blockIdx.y; channel < nchannels; channel += gridDim.y) {
|
||||
const float * xc = x + sample*stride_sample + channel*stride_channel + row*stride_row;
|
||||
float * dstc = dst + ((sample*nchannels + channel)*nrows + row)*ncols;
|
||||
const T * xc = x + sample*stride_sample + channel*stride_channel + row*stride_row;
|
||||
T * dstc = dst + ((sample*nchannels + channel)*nrows + row)*ncols;
|
||||
|
||||
[[maybe_unused]] const float * mulc = nullptr;
|
||||
if constexpr (do_multiply) {
|
||||
@@ -134,7 +161,7 @@ static __global__ void rms_norm_f32(const float * x,
|
||||
mulc = mul + mul_sample * mul_stride_sample + mul_channel * mul_stride_channel + mul_row * mul_stride_row;
|
||||
}
|
||||
|
||||
[[maybe_unused]] const float * addc = nullptr;
|
||||
[[maybe_unused]] const T_add * addc = nullptr;
|
||||
if constexpr (do_add) {
|
||||
const int add_row = fastmodulo(row, add_nrows_packed);
|
||||
const int add_channel = fastmodulo(channel, add_nchannels_packed);
|
||||
@@ -144,10 +171,7 @@ static __global__ void rms_norm_f32(const float * x,
|
||||
|
||||
if constexpr (do_vec) {
|
||||
const int ncols4 = ncols / 4;
|
||||
const float4 * x4 = (const float4 *) xc;
|
||||
const float4 * mul4 = (const float4 *) mulc;
|
||||
const float4 * add4 = (const float4 *) addc;
|
||||
float4 * dst4 = (float4 *) dstc;
|
||||
|
||||
float4 xs[vpt];
|
||||
float tmp = 0.0f;
|
||||
@@ -156,7 +180,7 @@ static __global__ void rms_norm_f32(const float * x,
|
||||
for (int k = 0; k < vpt; k++) {
|
||||
const int i = tid + k*block_size;
|
||||
if (i < ncols4) {
|
||||
xs[k] = x4[i];
|
||||
xs[k] = rms_norm_load4(xc, i);
|
||||
tmp += xs[k].x*xs[k].x + xs[k].y*xs[k].y + xs[k].z*xs[k].z + xs[k].w*xs[k].w;
|
||||
}
|
||||
}
|
||||
@@ -176,7 +200,7 @@ static __global__ void rms_norm_f32(const float * x,
|
||||
r.w *= m.w;
|
||||
}
|
||||
if constexpr (do_add) {
|
||||
const float4 a = add4[i];
|
||||
const float4 a = rms_norm_load4(addc, i);
|
||||
r.x += a.x;
|
||||
r.y += a.y;
|
||||
r.z += a.z;
|
||||
@@ -188,14 +212,14 @@ static __global__ void rms_norm_f32(const float * x,
|
||||
r.z *= scale_out;
|
||||
r.w *= scale_out;
|
||||
}
|
||||
dst4[i] = r;
|
||||
rms_norm_store4(dstc, i, r);
|
||||
}
|
||||
}
|
||||
} else {
|
||||
float tmp = 0.0f; // partial sum for thread in warp
|
||||
|
||||
for (int col = tid; col < ncols; col += block_size) {
|
||||
const float xi = xc[col];
|
||||
const float xi = ggml_cuda_cast<float>(xc[col]);
|
||||
tmp += xi * xi;
|
||||
}
|
||||
|
||||
@@ -206,17 +230,18 @@ static __global__ void rms_norm_f32(const float * x,
|
||||
const float scale = rsqrtf(mean + eps);
|
||||
|
||||
for (int col = tid; col < ncols; col += block_size) {
|
||||
const float xi = ggml_cuda_cast<float>(xc[col]);
|
||||
if constexpr (do_multiply && do_add) {
|
||||
const int mul_col = fastmodulo(col, mul_ncols_packed);
|
||||
const int add_col = fastmodulo(col, add_ncols_packed);
|
||||
dstc[col] = scale * xc[col] * mulc[mul_col] + addc[add_col];
|
||||
dstc[col] = ggml_cuda_cast<T>(scale * xi * mulc[mul_col] + ggml_cuda_cast<float>(addc[add_col]));
|
||||
} else if constexpr (do_multiply) {
|
||||
const int mul_col = fastmodulo(col, mul_ncols_packed);
|
||||
dstc[col] = scale * xc[col] * mulc[mul_col];
|
||||
dstc[col] = ggml_cuda_cast<T>(scale * xi * mulc[mul_col]);
|
||||
} else if constexpr (do_scale) {
|
||||
dstc[col] = scale_out * (scale * xc[col]);
|
||||
dstc[col] = ggml_cuda_cast<T>(scale_out * (scale * xi));
|
||||
} else {
|
||||
dstc[col] = scale * xc[col];
|
||||
dstc[col] = ggml_cuda_cast<T>(scale * xi);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -387,7 +412,7 @@ static void group_norm_f32_cuda(
|
||||
do { \
|
||||
const dim3 block_dims(BLOCK_SIZE, 1, 1); \
|
||||
const ggml_cuda_kernel_launch_params launch_params{ blocks_num, block_dims, block_dims.x > WARP_SIZE ? 32 * sizeof(float) : 0, stream }; \
|
||||
ggml_cuda_kernel_launch(rms_norm_f32<BLOCK_SIZE, VPT, VEC, do_multiply, do_add, do_scale>, launch_params, \
|
||||
ggml_cuda_kernel_launch(rms_norm_f32<T, BLOCK_SIZE, VPT, VEC, do_multiply, do_add, do_scale, T_add>, launch_params, \
|
||||
x, dst, ncols, nchannels, nsamples, stride_row, stride_channel, stride_sample, eps, \
|
||||
mul, mul_stride_row, mul_stride_channel, mul_stride_sample, \
|
||||
mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed, \
|
||||
@@ -395,10 +420,10 @@ static void group_norm_f32_cuda(
|
||||
add_ncols_packed, add_nrows_packed, add_nchannels_packed, add_nsamples_packed, scale_out); \
|
||||
} while (0)
|
||||
|
||||
template <bool do_multiply, bool do_add, bool do_scale>
|
||||
template <bool do_multiply, bool do_add, bool do_scale, typename T_add = float, typename T>
|
||||
static void rms_norm_f32_dispatch(
|
||||
const dim3 blocks_num,
|
||||
const ggml_tensor * x_t, float * dst, const int ncols, const int nchannels, const int nsamples,
|
||||
const ggml_tensor * x_t, T * dst, const int ncols, const int nchannels, const int nsamples,
|
||||
const int64_t stride_row, const int64_t stride_channel, const int64_t stride_sample, const float eps,
|
||||
const ggml_tensor * mul_t, const int64_t mul_stride_row, const int64_t mul_stride_channel, const int64_t mul_stride_sample,
|
||||
const uint3 mul_ncols_packed, const uint3 mul_nrows_packed, const uint3 mul_nchannels_packed, const uint3 mul_nsamples_packed,
|
||||
@@ -408,20 +433,20 @@ static void rms_norm_f32_dispatch(
|
||||
static_assert(!do_add || do_multiply, "fusing add is not supported without multiplying");
|
||||
static_assert(!do_scale || !do_multiply, "fusing scale is not supported with multiplying");
|
||||
|
||||
const float * x = (const float *) x_t->data;
|
||||
const T * x = (const T *) x_t->data;
|
||||
const float * mul = mul_t ? (const float *) mul_t->data : nullptr;
|
||||
const float * add = add_t ? (const float *) add_t->data : nullptr;
|
||||
const T_add * add = add_t ? (const T_add *) add_t->data : nullptr;
|
||||
|
||||
// dst is written at row * ncols, i.e. always at a multiple of ncols, so only its base needs checking here
|
||||
bool vec = ncols % 4 == 0 && ncols > 0 && ((uintptr_t) dst & 15) == 0;
|
||||
vec = vec && ggml_cuda_is_aligned(x_t, 16);
|
||||
bool vec = ncols % 4 == 0 && ncols > 0 && ((uintptr_t) dst % (4*sizeof(T))) == 0;
|
||||
vec = vec && ggml_cuda_is_aligned(x_t, 4*sizeof(T));
|
||||
if constexpr (do_multiply) {
|
||||
vec = vec && mul_ncols_packed.z == (uint32_t) ncols
|
||||
&& ggml_cuda_is_aligned(mul_t, 16);
|
||||
}
|
||||
if constexpr (do_add) {
|
||||
vec = vec && add_ncols_packed.z == (uint32_t) ncols
|
||||
&& ggml_cuda_is_aligned(add_t, 16);
|
||||
&& ggml_cuda_is_aligned(add_t, 4*sizeof(T_add));
|
||||
}
|
||||
|
||||
const int n4 = ncols / 4;
|
||||
@@ -449,9 +474,9 @@ static void rms_norm_f32_dispatch(
|
||||
}
|
||||
#undef RMS_NORM_LAUNCH
|
||||
|
||||
template <bool do_scale = false>
|
||||
template <bool do_scale = false, typename T>
|
||||
static void rms_norm_f32_cuda(
|
||||
const ggml_tensor * x_t, float * dst, const int ncols, const int nrows, const int nchannels, const int nsamples,
|
||||
const ggml_tensor * x_t, T * dst, const int ncols, const int nrows, const int nchannels, const int nsamples,
|
||||
const int64_t stride_row, const int64_t stride_channel, const int64_t stride_sample, const float eps, cudaStream_t stream,
|
||||
const float scale_out = 1.0f) {
|
||||
const dim3 blocks_num(nrows, MIN(nchannels, UINT16_MAX), MIN(nsamples, UINT16_MAX));
|
||||
@@ -462,10 +487,11 @@ static void rms_norm_f32_cuda(
|
||||
stream, scale_out);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static void rms_norm_mul_f32_cuda(const ggml_tensor * x_t,
|
||||
const ggml_tensor * mul_t,
|
||||
const ggml_tensor * add_t,
|
||||
float * dst,
|
||||
T * dst,
|
||||
const int ncols,
|
||||
const int nrows,
|
||||
const int nchannels,
|
||||
@@ -511,12 +537,32 @@ static void rms_norm_mul_f32_cuda(const ggml_tensor * x_t,
|
||||
const uint3 add_nrows_packed = init_fastdiv_values(add_nrows);
|
||||
const uint3 add_nchannels_packed = init_fastdiv_values(add_nchannels);
|
||||
const uint3 add_nsamples_packed = init_fastdiv_values(add_nsamples);
|
||||
rms_norm_f32_dispatch</*do_multiply*/true, /*do_add*/true, /*do_scale*/false>(
|
||||
blocks_num, x_t, dst, ncols, nchannels, nsamples, stride_row, stride_channel, stride_sample, eps,
|
||||
mul_t, mul_stride_row, mul_stride_channel, mul_stride_sample,
|
||||
mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed,
|
||||
add_t, add_stride_row, add_stride_channel, add_stride_sample,
|
||||
add_ncols_packed, add_nrows_packed, add_nchannels_packed, add_nsamples_packed, stream);
|
||||
// the add operand is F32 or has the type of x (e.g. a BF16 residual)
|
||||
if (add_t->type == GGML_TYPE_F32) {
|
||||
rms_norm_f32_dispatch</*do_multiply*/true, /*do_add*/true, /*do_scale*/false, float>(
|
||||
blocks_num, x_t, dst, ncols, nchannels, nsamples, stride_row, stride_channel, stride_sample, eps,
|
||||
mul_t, mul_stride_row, mul_stride_channel, mul_stride_sample,
|
||||
mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed,
|
||||
add_t, add_stride_row, add_stride_channel, add_stride_sample,
|
||||
add_ncols_packed, add_nrows_packed, add_nchannels_packed, add_nsamples_packed, stream);
|
||||
} else {
|
||||
rms_norm_f32_dispatch</*do_multiply*/true, /*do_add*/true, /*do_scale*/false, T>(
|
||||
blocks_num, x_t, dst, ncols, nchannels, nsamples, stride_row, stride_channel, stride_sample, eps,
|
||||
mul_t, mul_stride_row, mul_stride_channel, mul_stride_sample,
|
||||
mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed,
|
||||
add_t, add_stride_row, add_stride_channel, add_stride_sample,
|
||||
add_ncols_packed, add_nrows_packed, add_nchannels_packed, add_nsamples_packed, stream);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// calls launch with the data of dst as float * or nv_bfloat16 *, the output types of the RMS_NORM kernels
|
||||
template <typename F>
|
||||
static void rms_norm_launch_typed(ggml_tensor * dst, const F & launch) {
|
||||
if (dst->type == GGML_TYPE_BF16) {
|
||||
launch((nv_bfloat16 *) dst->data);
|
||||
} else {
|
||||
launch((float *) dst->data);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -590,11 +636,10 @@ void ggml_cuda_op_group_norm(ggml_backend_cuda_context & ctx, ggml_tensor * dst)
|
||||
|
||||
void ggml_cuda_op_rms_norm(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
|
||||
const ggml_tensor * src0 = dst->src[0];
|
||||
float * dst_d = (float *) dst->data;
|
||||
cudaStream_t stream = ctx.stream();
|
||||
|
||||
GGML_ASSERT(src0->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT( dst->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_BF16);
|
||||
GGML_ASSERT( dst->type == src0->type);
|
||||
|
||||
GGML_TENSOR_UNARY_OP_LOCALS;
|
||||
|
||||
@@ -608,7 +653,9 @@ void ggml_cuda_op_rms_norm(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
|
||||
const int64_t s02 = nb02 / ts0;
|
||||
const int64_t s03 = nb03 / ts0;
|
||||
|
||||
rms_norm_f32_cuda(src0, dst_d, ne00, ne01, ne02, ne03, s01, s02, s03, eps, stream);
|
||||
rms_norm_launch_typed(dst, [&](auto * dst_d) {
|
||||
rms_norm_f32_cuda(src0, dst_d, ne00, ne01, ne02, ne03, s01, s02, s03, eps, stream);
|
||||
});
|
||||
}
|
||||
|
||||
void ggml_cuda_op_rms_norm_scale_fused(ggml_backend_cuda_context & ctx, ggml_tensor * dst, ggml_tensor * scale_tensor) {
|
||||
@@ -653,12 +700,12 @@ void ggml_cuda_op_rms_norm_fused(ggml_backend_cuda_context & ctx, ggml_tensor *
|
||||
GGML_ASSERT(false);
|
||||
}
|
||||
|
||||
float * dst_d = (float *) mul_tensor->data;
|
||||
cudaStream_t stream = ctx.stream();
|
||||
|
||||
GGML_ASSERT(rms_norm_src->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(dst->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(mul_tensor->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(rms_norm_src->type == GGML_TYPE_F32 || rms_norm_src->type == GGML_TYPE_BF16);
|
||||
GGML_ASSERT(dst->type == rms_norm_src->type);
|
||||
GGML_ASSERT(mul_tensor->type == rms_norm_src->type);
|
||||
GGML_ASSERT(mul_src->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(eps >= 0.0f);
|
||||
|
||||
const int64_t ne00 = rms_norm_src->ne[0];
|
||||
@@ -683,14 +730,16 @@ void ggml_cuda_op_rms_norm_fused(ggml_backend_cuda_context & ctx, ggml_tensor *
|
||||
const int mul_nchannels = mul_src->ne[2];
|
||||
const int mul_nsamples = mul_src->ne[3];
|
||||
|
||||
rms_norm_mul_f32_cuda(rms_norm_src, mul_src, nullptr, dst_d,
|
||||
ne00, ne01, ne02, ne03,
|
||||
/*s00*/ s01, s02, s03,
|
||||
/*mul_s00*/ mul_s01, mul_s02, mul_s03,
|
||||
mul_ncols, mul_nrows, mul_nchannels, mul_nsamples,
|
||||
/*add_s00*/ 0, 0, 0,
|
||||
0, 0, 0, 0,
|
||||
eps, stream);
|
||||
rms_norm_launch_typed(mul_tensor, [&](auto * dst_d) {
|
||||
rms_norm_mul_f32_cuda(rms_norm_src, mul_src, nullptr, dst_d,
|
||||
ne00, ne01, ne02, ne03,
|
||||
/*s00*/ s01, s02, s03,
|
||||
/*mul_s00*/ mul_s01, mul_s02, mul_s03,
|
||||
mul_ncols, mul_nrows, mul_nchannels, mul_nsamples,
|
||||
/*add_s00*/ 0, 0, 0,
|
||||
0, 0, 0, 0,
|
||||
eps, stream);
|
||||
});
|
||||
}
|
||||
|
||||
void ggml_cuda_op_rms_norm_fused_add(ggml_backend_cuda_context & ctx,
|
||||
@@ -722,13 +771,14 @@ void ggml_cuda_op_rms_norm_fused_add(ggml_backend_cuda_context & ctx,
|
||||
GGML_ASSERT(false);
|
||||
}
|
||||
|
||||
float * dst_d = (float *) add_tensor->data;
|
||||
cudaStream_t stream = ctx.stream();
|
||||
|
||||
GGML_ASSERT(rms_norm_src->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(dst->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(mul_tensor->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(add_tensor->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(rms_norm_src->type == GGML_TYPE_F32 || rms_norm_src->type == GGML_TYPE_BF16);
|
||||
GGML_ASSERT(dst->type == rms_norm_src->type);
|
||||
GGML_ASSERT(mul_tensor->type == rms_norm_src->type);
|
||||
GGML_ASSERT(add_tensor->type == rms_norm_src->type);
|
||||
GGML_ASSERT(mul_src->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(add_src->type == GGML_TYPE_F32 || add_src->type == rms_norm_src->type);
|
||||
GGML_ASSERT(eps >= 0.0f);
|
||||
|
||||
const int64_t ne00 = rms_norm_src->ne[0];
|
||||
@@ -764,14 +814,16 @@ void ggml_cuda_op_rms_norm_fused_add(ggml_backend_cuda_context & ctx,
|
||||
const int add_nchannels = add_src->ne[2];
|
||||
const int add_nsamples = add_src->ne[3];
|
||||
|
||||
rms_norm_mul_f32_cuda(rms_norm_src, mul_src,add_src,dst_d,
|
||||
ne00,ne01, ne02, ne03,
|
||||
/*s00*/ s01, s02, s03,
|
||||
/*mul_s00*/ mul_s01, mul_s02, mul_s03,
|
||||
mul_ncols, mul_nrows, mul_nchannels, mul_nsamples,
|
||||
/*add_s00*/ add_s01, add_s02, add_s03,
|
||||
add_ncols, add_nrows, add_nchannels, add_nsamples,
|
||||
eps, stream);
|
||||
rms_norm_launch_typed(add_tensor, [&](auto * dst_d) {
|
||||
rms_norm_mul_f32_cuda(rms_norm_src, mul_src, add_src, dst_d,
|
||||
ne00, ne01, ne02, ne03,
|
||||
/*s00*/ s01, s02, s03,
|
||||
/*mul_s00*/ mul_s01, mul_s02, mul_s03,
|
||||
mul_ncols, mul_nrows, mul_nchannels, mul_nsamples,
|
||||
/*add_s00*/ add_s01, add_s02, add_s03,
|
||||
add_ncols, add_nrows, add_nchannels, add_nsamples,
|
||||
eps, stream);
|
||||
});
|
||||
}
|
||||
|
||||
void ggml_cuda_op_rms_norm_back(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
|
||||
|
||||
@@ -316,6 +316,7 @@ struct ggml_hexagon_session {
|
||||
ggml_hexagon_opbatch* op_batch;
|
||||
ggml_hexagon_opqueue* op_queue;
|
||||
|
||||
std::unordered_set<ggml_hexagon_shared_buffer *> buffers;
|
||||
std::unordered_map<int, std::unique_ptr<ggml_hexagon_shared_buffer>> cloned_buffers;
|
||||
std::unordered_set<ggml_hexagon_session *> virt_peers;
|
||||
std::unordered_set<ggml_hexagon_session *> phys_peers;
|
||||
@@ -337,6 +338,7 @@ struct ggml_hexagon_session {
|
||||
ggml_hexagon_mdev_group mdev;
|
||||
ggml_backend_dev_t dev = nullptr;
|
||||
ggml_backend_hexagon_device_context * dev_ctx = nullptr;
|
||||
ggml_hexagon_shared_buffer * shm_buf = nullptr;
|
||||
ggml_hexagon_fence_buffer * fence_buf = nullptr;
|
||||
|
||||
ggml_hexagon_session(const ggml_hexagon_device_config & config, ggml_backend_dev_t dev = nullptr, uint32_t mdev_idx = 0, uint32_t mdev_count = 0) noexcept(false);
|
||||
@@ -373,8 +375,8 @@ struct ggml_hexagon_session {
|
||||
void flush_peers();
|
||||
void flush_pending(bool all = true);
|
||||
|
||||
ggml_hexagon_shared_buffer * mmap_tensor(const ggml_tensor * t);
|
||||
bool clone_buffer(const ggml_hexagon_shared_buffer*);
|
||||
void mmap_buffers();
|
||||
void clone_buffer(const ggml_hexagon_shared_buffer*);
|
||||
void release_buffer(const ggml_hexagon_shared_buffer*);
|
||||
void unclone_buffer(const ggml_hexagon_shared_buffer*);
|
||||
|
||||
@@ -434,7 +436,7 @@ struct ggml_hexagon_rpcmem_block {
|
||||
int fd = -1;
|
||||
size_t size = 0;
|
||||
|
||||
std::unordered_set<ggml_hexagon_session *> mapped_clones;
|
||||
std::unordered_set<ggml_hexagon_session *> clones;
|
||||
|
||||
ggml_hexagon_rpcmem_block(size_t size) {
|
||||
base = (uint8_t *) rpcmem_alloc2(RPCMEM_HEAP_ID_SYSTEM, RPCMEM_DEFAULT_FLAGS, size);
|
||||
@@ -468,16 +470,24 @@ struct ggml_hexagon_shared_buffer {
|
||||
size_t size() const { return mem ? mem->size : 0; }
|
||||
int fd() const { return mem ? mem->fd : -1; }
|
||||
|
||||
void mmap(bool extended = false) {
|
||||
void mmap() {
|
||||
if (!this->mem) return;
|
||||
if (this->mapped) return;
|
||||
|
||||
GGML_ASSERT(!this->pinned || !extended);
|
||||
GGML_ASSERT(!this->pinned || !this->extended);
|
||||
|
||||
this->extended = extended;
|
||||
fastrpc_map_flags flags = this->pinned ? FASTRPC_MAP_FD : (extended ? FASTRPC_MAP_FD_DELAYED_EXTENDED : FASTRPC_MAP_FD_DELAYED);
|
||||
fastrpc_map_flags flags = this->pinned ? FASTRPC_MAP_FD : (this->extended ? FASTRPC_MAP_FD_DELAYED_EXTENDED : FASTRPC_MAP_FD_DELAYED);
|
||||
|
||||
int err = fastrpc_mmap(sess->domain_id, fd(), (void *) base(), 0, size(), flags);
|
||||
if (err != 0) {
|
||||
if (this->extended) {
|
||||
flags = FASTRPC_MAP_FD_EXTENDED;
|
||||
}
|
||||
HEX_VERBOSE("ggml-hex: %s mmap failed (err 0x%08x), retrying in 100ms...\n", sess->c_name(), (unsigned) err);
|
||||
std::this_thread::sleep_for(std::chrono::milliseconds(100));
|
||||
err = fastrpc_mmap(sess->domain_id, fd(), (void *) base(), 0, size(), flags);
|
||||
}
|
||||
|
||||
if (err != 0) {
|
||||
GGML_LOG_ERROR("ggml-hex: %s buffer mapping failed : domain_id %d size %zu fd %d error 0x%08x\n", sess->c_name(),
|
||||
sess->domain_id, size(), fd(), (unsigned) err);
|
||||
@@ -550,10 +560,10 @@ struct ggml_hexagon_shared_buffer {
|
||||
|
||||
// Clone constructor for cross-session mapping
|
||||
ggml_hexagon_shared_buffer(ggml_hexagon_session * sess, const ggml_hexagon_shared_buffer & other) {
|
||||
this->sess = sess;
|
||||
this->mem = other.mem;
|
||||
this->mapped = false;
|
||||
this->pinned = other.pinned;
|
||||
this->sess = sess;
|
||||
this->mem = other.mem;
|
||||
this->mapped = false;
|
||||
this->pinned = false;
|
||||
this->extended = other.extended;
|
||||
}
|
||||
|
||||
@@ -563,19 +573,30 @@ struct ggml_hexagon_shared_buffer {
|
||||
};
|
||||
|
||||
struct ggml_hexagon_fence_buffer : public ggml_hexagon_shared_buffer {
|
||||
uint8_t * data = nullptr;
|
||||
uint32_t slot_count = 0;
|
||||
uint32_t slot_head = 0;
|
||||
std::vector<uint32_t> free_slots;
|
||||
ggml_backend_buffer backend_buffer{};
|
||||
|
||||
ggml_hexagon_fence_buffer(ggml_hexagon_session * sess, ggml_backend_buffer_type_t buft, size_t size)
|
||||
: ggml_hexagon_shared_buffer(sess, size, false /* pinned */),
|
||||
slot_count(size / HTP_FENCE_SLOT_SIZE),
|
||||
static size_t size_of(size_t n_slots) {
|
||||
return n_slots * HTP_FENCE_SLOT_SIZE;
|
||||
}
|
||||
|
||||
ggml_hexagon_fence_buffer(ggml_hexagon_shared_buffer * shm_buf, ggml_backend_buffer_type_t buft, size_t offset, size_t n_slots)
|
||||
: ggml_hexagon_shared_buffer(shm_buf->sess, *shm_buf),
|
||||
data(shm_buf->base() + offset),
|
||||
slot_count(n_slots),
|
||||
slot_head(0) {
|
||||
GGML_ASSERT((uintptr_t) data % HEX_L2_LINE_SIZE == 0);
|
||||
this->mapped = shm_buf->mapped;
|
||||
backend_buffer.buft = buft;
|
||||
backend_buffer.context = static_cast<ggml_hexagon_shared_buffer *>(this);
|
||||
backend_buffer.size = size;
|
||||
mmap(false);
|
||||
backend_buffer.size = size_of(n_slots);
|
||||
}
|
||||
|
||||
~ggml_hexagon_fence_buffer() {
|
||||
this->mapped = false;
|
||||
}
|
||||
|
||||
uint8_t * alloc_slot(uint32_t n_slots = 1) {
|
||||
@@ -583,11 +604,11 @@ struct ggml_hexagon_fence_buffer : public ggml_hexagon_shared_buffer {
|
||||
if (n_slots == 1 && !free_slots.empty()) {
|
||||
uint32_t slot = free_slots.back();
|
||||
free_slots.pop_back();
|
||||
ptr = base() + (size_t) slot * HTP_FENCE_SLOT_SIZE;
|
||||
ptr = data + (size_t) slot * HTP_FENCE_SLOT_SIZE;
|
||||
} else if (slot_head + n_slots <= slot_count) {
|
||||
uint32_t slot = slot_head;
|
||||
slot_head += n_slots;
|
||||
ptr = base() + (size_t) slot * HTP_FENCE_SLOT_SIZE;
|
||||
ptr = data + (size_t) slot * HTP_FENCE_SLOT_SIZE;
|
||||
}
|
||||
if (ptr) {
|
||||
memset(ptr, 0, (size_t) n_slots * HTP_FENCE_SLOT_SIZE);
|
||||
@@ -597,7 +618,7 @@ struct ggml_hexagon_fence_buffer : public ggml_hexagon_shared_buffer {
|
||||
|
||||
void free_slot(void * ptr, uint32_t n_slots = 1) {
|
||||
if (!ptr) return;
|
||||
uint32_t slot = ((uint8_t *) ptr - base()) / HTP_FENCE_SLOT_SIZE;
|
||||
uint32_t slot = ((uint8_t *) ptr - data) / HTP_FENCE_SLOT_SIZE;
|
||||
for (uint32_t i = 0; i < n_slots; i++) {
|
||||
free_slots.push_back(slot + i);
|
||||
}
|
||||
@@ -636,6 +657,7 @@ inline ggml_tensor ggml_hexagon_session::make_fence_tensor(const volatile void *
|
||||
|
||||
static void ggml_backend_hexagon_buffer_free_buffer(ggml_backend_buffer_t buffer) {
|
||||
auto sbuf = static_cast<ggml_hexagon_shared_buffer *>(buffer->context);
|
||||
sbuf->sess->buffers.erase(sbuf);
|
||||
sbuf->sess->unclone_buffer(sbuf);
|
||||
delete sbuf;
|
||||
}
|
||||
@@ -975,31 +997,44 @@ static bool ggml_backend_hexagon_host_buffer_type_is_host(ggml_backend_buffer_ty
|
||||
GGML_UNUSED(buft);
|
||||
}
|
||||
|
||||
struct ggml_backend_hexagon_alloc_buffer_n_plan_item {
|
||||
size_t size;
|
||||
int first;
|
||||
int last;
|
||||
};
|
||||
|
||||
using ggml_backend_hexagon_alloc_buffer_n_plan_t = std::vector<ggml_backend_hexagon_alloc_buffer_n_plan_item>;
|
||||
|
||||
static const char * ggml_hexagon_kv_layer_suffix(const struct ggml_tensor * t) {
|
||||
if (strncmp(t->name, "cache_", 6) != 0) {
|
||||
return NULL;
|
||||
}
|
||||
const char * p = strstr(t->name, "_l");
|
||||
if (!p || !isdigit((unsigned char)p[2])) {
|
||||
return NULL;
|
||||
}
|
||||
return p;
|
||||
}
|
||||
|
||||
struct ggml_backend_hexagon_alloc_unit {
|
||||
size_t size;
|
||||
int first;
|
||||
int last;
|
||||
};
|
||||
|
||||
struct ggml_backend_hexagon_alloc_buffer_n_plan_item {
|
||||
size_t size;
|
||||
std::vector<ggml_backend_hexagon_alloc_unit> units;
|
||||
};
|
||||
|
||||
using ggml_backend_hexagon_alloc_buffer_n_plan_t = std::vector<ggml_backend_hexagon_alloc_buffer_n_plan_item>;
|
||||
|
||||
static bool ggml_hexagon_tensor_layer(const struct ggml_tensor * t, int & kind, int & layer) {
|
||||
if (strncmp(t->name, "blk.", 4) == 0 && isdigit((unsigned char) t->name[4])) {
|
||||
char * end = NULL;
|
||||
long l = strtol(t->name + 4, &end, 10);
|
||||
if (end && *end == '.') {
|
||||
kind = 0;
|
||||
layer = (int) l;
|
||||
return true;
|
||||
}
|
||||
}
|
||||
if (strncmp(t->name, "cache_", 6) == 0) {
|
||||
const char * p = strstr(t->name, "_l");
|
||||
if (p && isdigit((unsigned char) p[2])) {
|
||||
char * end = NULL;
|
||||
long l = strtol(p + 2, &end, 10);
|
||||
if (end && *end == '\0') {
|
||||
kind = 1;
|
||||
layer = (int) l;
|
||||
return true;
|
||||
}
|
||||
}
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
static ggml_backend_hexagon_alloc_buffer_n_plan_t ggml_backend_hexagon_alloc_buffer_n_plan(
|
||||
ggml_backend_buffer_type_t buft, struct ggml_tensor ** tensors, int n_tensors) {
|
||||
ggml_backend_hexagon_alloc_buffer_n_plan_t plan;
|
||||
@@ -1012,61 +1047,72 @@ static ggml_backend_hexagon_alloc_buffer_n_plan_t ggml_backend_hexagon_alloc_buf
|
||||
int i = 0;
|
||||
while (i < n_tensors) {
|
||||
struct ggml_tensor * t = tensors[i];
|
||||
size_t unit_size = 0;
|
||||
int unit_first = i;
|
||||
int unit_last = i + 1;
|
||||
|
||||
if (t->data == NULL && t->view_src == NULL) {
|
||||
unit_size += GGML_PAD(ggml_backend_buft_get_alloc_size(buft, t), alignment);
|
||||
if (t->data != NULL || t->view_src != NULL) {
|
||||
i++;
|
||||
continue;
|
||||
}
|
||||
|
||||
const char * layer_suffix = ggml_hexagon_kv_layer_suffix(t);
|
||||
size_t unit_size = GGML_PAD(ggml_backend_buft_get_alloc_size(buft, t), alignment);
|
||||
int unit_first = i;
|
||||
int unit_last = i + 1;
|
||||
|
||||
while (unit_last < n_tensors) {
|
||||
int kind = -1, layer = -1;
|
||||
const bool has_layer = ggml_hexagon_tensor_layer(t, kind, layer);
|
||||
|
||||
while (has_layer && unit_last < n_tensors) {
|
||||
struct ggml_tensor * next = tensors[unit_last];
|
||||
|
||||
if (next->view_src != NULL) {
|
||||
if (next->data != NULL || next->view_src != NULL) {
|
||||
unit_last++;
|
||||
continue;
|
||||
}
|
||||
|
||||
if (layer_suffix != NULL) {
|
||||
const char * next_suffix = ggml_hexagon_kv_layer_suffix(next);
|
||||
if (next_suffix != NULL && strcmp(layer_suffix, next_suffix) == 0) {
|
||||
if (next->data == NULL) {
|
||||
unit_size += GGML_PAD(ggml_backend_buft_get_alloc_size(buft, next), alignment);
|
||||
}
|
||||
unit_last++;
|
||||
continue;
|
||||
}
|
||||
int next_kind = -1, next_layer = -1;
|
||||
if (!ggml_hexagon_tensor_layer(next, next_kind, next_layer) ||
|
||||
next_kind != kind || next_layer != layer) {
|
||||
break;
|
||||
}
|
||||
|
||||
break;
|
||||
size_t next_size = GGML_PAD(ggml_backend_buft_get_alloc_size(buft, next), alignment);
|
||||
if (unit_size + next_size > max_size) {
|
||||
break;
|
||||
}
|
||||
|
||||
unit_size += next_size;
|
||||
unit_last++;
|
||||
}
|
||||
|
||||
units.push_back({ unit_size, unit_first, unit_last });
|
||||
i = unit_last;
|
||||
}
|
||||
|
||||
size_t cur_buf_size = 0;
|
||||
int cur_buf_first = 0;
|
||||
|
||||
for (const auto & unit : units) {
|
||||
if (unit.size == 0) {
|
||||
continue;
|
||||
}
|
||||
|
||||
if (cur_buf_size > 0 && (cur_buf_size + unit.size) > max_size) {
|
||||
plan.push_back({ cur_buf_size, cur_buf_first, unit.first });
|
||||
cur_buf_size = 0;
|
||||
cur_buf_first = unit.first;
|
||||
bool placed = false;
|
||||
for (auto & item : plan) {
|
||||
if (item.size + unit.size <= max_size) {
|
||||
item.size += unit.size;
|
||||
item.units.push_back(unit);
|
||||
placed = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (!placed) {
|
||||
plan.push_back({ unit.size, { unit } });
|
||||
}
|
||||
|
||||
cur_buf_size += unit.size;
|
||||
}
|
||||
|
||||
if (cur_buf_size > 0) {
|
||||
plan.push_back({ cur_buf_size, cur_buf_first, n_tensors });
|
||||
std::stable_sort(plan.begin(), plan.end(), [](const auto & a, const auto & b) {
|
||||
return a.size > b.size;
|
||||
});
|
||||
|
||||
if (!plan.empty()) {
|
||||
for (const auto & unit : units) {
|
||||
if (unit.size == 0) {
|
||||
plan[0].units.push_back(unit);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return plan;
|
||||
@@ -1092,28 +1138,19 @@ static ggml_backend_buffer_t ggml_backend_hexagon_buffer_type_alloc_buffer_n(
|
||||
struct ggml_tallocr tallocr = ggml_tallocr_new(buffer);
|
||||
|
||||
struct ggml_tensor * t_failed = NULL;
|
||||
for (int j = item.first; j < item.last; j++) {
|
||||
struct ggml_tensor * t = tensors[j];
|
||||
if (t->data == NULL) {
|
||||
if (t->view_src == NULL) {
|
||||
for (const auto & unit : item.units) {
|
||||
for (int j = unit.first; j < unit.last; j++) {
|
||||
struct ggml_tensor * t = tensors[j];
|
||||
if (t->data == NULL && t->view_src == NULL) {
|
||||
if (ggml_tallocr_alloc(&tallocr, t) != GGML_STATUS_SUCCESS) {
|
||||
t_failed = t;
|
||||
break;
|
||||
}
|
||||
} else if (t->buffer == NULL) {
|
||||
if (ggml_backend_view_init(t) != GGML_STATUS_SUCCESS) {
|
||||
t_failed = t;
|
||||
break;
|
||||
}
|
||||
}
|
||||
} else {
|
||||
if (t->view_src != NULL && t->buffer == NULL) {
|
||||
if (ggml_backend_view_init(t) != GGML_STATUS_SUCCESS) {
|
||||
t_failed = t;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
if (t_failed != NULL) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (t_failed != NULL) {
|
||||
GGML_LOG_ERROR("%s: failed to initialize tensor %s\n", __func__, t_failed->name);
|
||||
@@ -1131,6 +1168,19 @@ static ggml_backend_buffer_t ggml_backend_hexagon_buffer_type_alloc_buffer_n(
|
||||
return NULL;
|
||||
}
|
||||
|
||||
for (int j = 0; j < n_tensors; j++) {
|
||||
struct ggml_tensor * t = tensors[j];
|
||||
if (t->view_src != NULL && t->buffer == NULL) {
|
||||
if (ggml_backend_view_init(t) != GGML_STATUS_SUCCESS) {
|
||||
GGML_LOG_ERROR("%s: failed to initialize tensor %s\n", __func__, t->name);
|
||||
for (ggml_backend_buffer_t b : buffers) {
|
||||
ggml_backend_buffer_free(b);
|
||||
}
|
||||
return NULL;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (buffers.size() == 1) {
|
||||
return buffers[0];
|
||||
}
|
||||
@@ -1200,7 +1250,6 @@ struct ggml_hexagon_opbatch {
|
||||
ggml_hexagon_session* sess;
|
||||
|
||||
std::vector<htp_opnode> ops; // htp_opnode of ops
|
||||
|
||||
std::vector<htp_buf_desc> h_bufs; // htp buffer descriptors
|
||||
std::vector<htp_tensor> h_tens; // htp tensor descriptors
|
||||
std::vector<htp_op_desc> h_ops; // htp op descriptors
|
||||
@@ -1529,45 +1578,45 @@ struct ggml_hexagon_registry {
|
||||
struct ggml_hexagon_opqueue {
|
||||
// Shared buffer for storing batches
|
||||
ggml_hexagon_shared_buffer *shm_buf;
|
||||
size_t shm_blk_size;
|
||||
uint8_t * data;
|
||||
size_t slot_size;
|
||||
size_t depth;
|
||||
|
||||
using opvec = std::vector<htp_opnode>;
|
||||
|
||||
std::vector<opvec> op_cache; // per batch op cache
|
||||
|
||||
ggml_hexagon_opqueue(ggml_hexagon_session *sess, size_t batch_size, size_t depth) : depth(depth) {
|
||||
static size_t size_of(size_t batch_size, size_t n_threads, size_t depth = 1) {
|
||||
size_t n_bufs = HTP_OP_MAX_BUFS;
|
||||
size_t n_ops = batch_size;
|
||||
size_t n_tensors = std::min<size_t>(n_ops * 4 + HTP_OP_MAX_INPUTS + HTP_OP_MAX_OUTPUTS, HTP_OP_MAX_TENSORS);
|
||||
|
||||
size_t tr_size = 0;
|
||||
if (opt_profile == 3) {
|
||||
tr_size = sess->n_threads * opt_optrace * sizeof(htp_trace_desc);
|
||||
tr_size = n_threads * opt_optrace * sizeof(htp_trace_desc);
|
||||
}
|
||||
|
||||
shm_blk_size = sizeof(htp_buf_desc) * n_bufs +
|
||||
sizeof(htp_tensor) * n_tensors +
|
||||
sizeof(htp_op_desc) * n_ops +
|
||||
sizeof(htp_prof_desc) * n_ops +
|
||||
tr_size;
|
||||
|
||||
shm_buf = new ggml_hexagon_shared_buffer(sess, shm_blk_size * depth, true /* pinned */);
|
||||
size_t slot_size = sizeof(htp_buf_desc) * n_bufs +
|
||||
sizeof(htp_tensor) * n_tensors +
|
||||
sizeof(htp_op_desc) * n_ops +
|
||||
sizeof(htp_prof_desc) * n_ops +
|
||||
tr_size;
|
||||
return slot_size * depth;
|
||||
}
|
||||
|
||||
ggml_hexagon_opqueue(ggml_hexagon_shared_buffer *shm_buf, size_t offset, size_t batch_size, size_t depth)
|
||||
: shm_buf(shm_buf),
|
||||
data(shm_buf->base() + offset),
|
||||
slot_size(size_of(batch_size, shm_buf->sess->n_threads)),
|
||||
depth(depth) {
|
||||
op_cache.resize(depth);
|
||||
|
||||
if (opt_verbose) {
|
||||
GGML_LOG_INFO("ggml-hex: %s allocated opqueue : batch-size %zu depth %zu shm-size %zu shm-block-size %zu\n",
|
||||
sess->c_name(), batch_size, depth, shm_buf->size(), shm_blk_size);
|
||||
GGML_LOG_INFO("ggml-hex: %s allocated opqueue : batch-size %zu depth %zu shm-size %zu slot-size %zu\n",
|
||||
shm_buf->c_name(), batch_size, depth, size_of(batch_size, shm_buf->sess->n_threads, depth), slot_size);
|
||||
}
|
||||
}
|
||||
|
||||
~ggml_hexagon_opqueue() {
|
||||
delete shm_buf;
|
||||
}
|
||||
|
||||
size_t shm_size() const { return shm_buf ? shm_buf->size() : 0; }
|
||||
|
||||
// push new batch
|
||||
bool push(htp_opbatch_req& req, dspqueue_buffer& dbuf, const ggml_hexagon_opbatch* op_batch, uint64_t seq) {
|
||||
static_assert(sizeof(htp_opbatch_req) % 8 == 0, "sizeof(htp_opbatch_req) must be multiple of 8");
|
||||
@@ -1603,13 +1652,13 @@ struct ggml_hexagon_opqueue {
|
||||
req.n_traces = 0;
|
||||
}
|
||||
|
||||
dbuf.ptr = shm_buf->base() + ((size_t) slot * shm_blk_size);
|
||||
dbuf.ptr = data + ((size_t) slot * slot_size);
|
||||
dbuf.fd = shm_buf->fd();
|
||||
dbuf.flags = DSPQUEUE_BUFFER_FLAG_FLUSH_SENDER | DSPQUEUE_BUFFER_FLAG_INVALIDATE_RECIPIENT;
|
||||
dbuf.offset = (uint8_t*) dbuf.ptr - (uint8_t*) shm_buf->base();
|
||||
dbuf.size = b_size + t_size + o_size + p_size + tr_size;
|
||||
|
||||
GGML_ASSERT(dbuf.size <= shm_blk_size);
|
||||
GGML_ASSERT(dbuf.size <= slot_size);
|
||||
|
||||
uint8_t * m_ptr = (uint8_t*) dbuf.ptr;
|
||||
uint8_t * b_ptr = m_ptr; m_ptr += b_size;
|
||||
@@ -1657,7 +1706,7 @@ struct ggml_hexagon_opqueue {
|
||||
}
|
||||
|
||||
const size_t m_size = b_size + t_size + o_size + p_size + tr_size;
|
||||
GGML_ASSERT(m_size <= shm_blk_size);
|
||||
GGML_ASSERT(m_size <= slot_size);
|
||||
|
||||
HEX_VERBOSE("ggml-hex: %s opqueue-pop batch #%llu : n-bufs %u n-tensors %u n-ops %u : m-size %zu b-size %zu t-size %zu o-size %zu\n",
|
||||
shm_buf->sess->c_name(), (unsigned long long) rsp.seq, rsp.n_bufs, rsp.n_tensors, rsp.n_ops,
|
||||
@@ -1775,6 +1824,7 @@ void ggml_hexagon_session::flush_batch(size_t min_ops) {
|
||||
if (op_batch->n_ops < min_ops) { return; }
|
||||
|
||||
op_batch->sort_buffers();
|
||||
this->mmap_buffers();
|
||||
|
||||
htp_opbatch_req req {};
|
||||
dspqueue_buffer dbuf{};
|
||||
@@ -1819,10 +1869,17 @@ void ggml_hexagon_session::flush_batch(size_t min_ops) {
|
||||
}
|
||||
|
||||
void ggml_hexagon_session::enqueue_op(const htp_opnode & node) {
|
||||
auto clone_tensor_buffer = [this](const ggml_tensor * t) {
|
||||
auto sbuf = this->mmap_tensor(t);
|
||||
if (!sbuf) return;
|
||||
if (sbuf->sess != this) {
|
||||
auto register_tensor = [&](const ggml_tensor * t) {
|
||||
if (!t) return;
|
||||
auto * sbuf = static_cast<ggml_hexagon_shared_buffer *>(t->buffer->context);
|
||||
if (ggml_backend_buffer_get_usage(t->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS) {
|
||||
sbuf->extended = opt_dma64;
|
||||
}
|
||||
if (sbuf->sess == this) {
|
||||
if (!sbuf->mapped) {
|
||||
this->buffers.insert(sbuf);
|
||||
}
|
||||
} else {
|
||||
this->clone_buffer(sbuf);
|
||||
}
|
||||
for (auto & sub : this->mdev.sessions) {
|
||||
@@ -1830,11 +1887,11 @@ void ggml_hexagon_session::enqueue_op(const htp_opnode & node) {
|
||||
}
|
||||
};
|
||||
|
||||
for (auto t : node.get_inputs()) {
|
||||
clone_tensor_buffer(t);
|
||||
for (const auto * src : node.get_inputs()) {
|
||||
register_tensor(src);
|
||||
}
|
||||
for (auto t : node.get_outputs()) {
|
||||
clone_tensor_buffer(t);
|
||||
for (const auto * dst : node.get_outputs()) {
|
||||
register_tensor(dst);
|
||||
}
|
||||
|
||||
const bool fits_tensors = op_batch->fit_tensors(node);
|
||||
@@ -1945,46 +2002,61 @@ void ggml_hexagon_session::enqueue_allreduce(
|
||||
this->enqueue_op(ar_node);
|
||||
}
|
||||
|
||||
ggml_hexagon_shared_buffer * ggml_hexagon_session::mmap_tensor(const ggml_tensor * t) {
|
||||
if (!t) return nullptr;
|
||||
void ggml_hexagon_session::mmap_buffers() {
|
||||
auto collect_session = [](ggml_hexagon_session * sess, std::unordered_set<ggml_hexagon_shared_buffer *> & unmapped) {
|
||||
for (auto * sbuf : sess->buffers) {
|
||||
if (!sbuf->mapped) {
|
||||
unmapped.insert(sbuf);
|
||||
}
|
||||
}
|
||||
for (auto & [fd, clone] : sess->cloned_buffers) {
|
||||
if (!clone->mapped) {
|
||||
unmapped.insert(clone.get());
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
auto sbuf = static_cast<ggml_hexagon_shared_buffer *>(t->buffer->context);
|
||||
if (!sbuf->mapped) {
|
||||
const bool is_weight = ggml_backend_buffer_get_usage(t->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS;
|
||||
const bool extended = opt_dma64 && is_weight;
|
||||
sbuf->mmap(extended);
|
||||
std::unordered_set<ggml_hexagon_shared_buffer *> unmapped;
|
||||
collect_session(this, unmapped);
|
||||
for (auto & sub : this->mdev.sessions) {
|
||||
collect_session(sub.get(), unmapped);
|
||||
}
|
||||
|
||||
if (!unmapped.empty()) {
|
||||
std::vector<ggml_hexagon_shared_buffer *> sorted(unmapped.begin(), unmapped.end());
|
||||
std::sort(sorted.begin(), sorted.end(), [](const auto * a, const auto * b) {
|
||||
if (a->size() != b->size()) {
|
||||
return a->size() > b->size();
|
||||
}
|
||||
if (a->fd() != b->fd()) {
|
||||
return a->fd() < b->fd();
|
||||
}
|
||||
return a->sess->domain_id < b->sess->domain_id;
|
||||
});
|
||||
for (auto * sbuf : sorted) {
|
||||
sbuf->mmap();
|
||||
}
|
||||
}
|
||||
return sbuf;
|
||||
}
|
||||
|
||||
bool ggml_hexagon_session::clone_buffer(const ggml_hexagon_shared_buffer * sbuf)
|
||||
void ggml_hexagon_session::clone_buffer(const ggml_hexagon_shared_buffer * sbuf)
|
||||
{
|
||||
GGML_ASSERT(sbuf && sbuf->mem);
|
||||
if (sbuf->sess == this) return true;
|
||||
if (sbuf->sess == this) return;
|
||||
|
||||
auto mem = sbuf->mem;
|
||||
int fd = mem->fd;
|
||||
|
||||
GGML_ASSERT(fd >= 0);
|
||||
|
||||
if (this->cloned_buffers.find(fd) != this->cloned_buffers.end()) return true;
|
||||
|
||||
GGML_ASSERT(sbuf->mapped);
|
||||
if (this->cloned_buffers.find(fd) != this->cloned_buffers.end()) return;
|
||||
|
||||
HEX_VERBOSE("ggml-hex: %s clone-buffer: %s base %p size %zu fd %d\n", this->name.c_str(),
|
||||
sbuf->c_name(), sbuf->base(), sbuf->size(), fd);
|
||||
|
||||
auto clone = std::make_unique<ggml_hexagon_shared_buffer>(this, *sbuf);
|
||||
try {
|
||||
clone->mmap(sbuf->extended);
|
||||
} catch (const std::exception & exc) {
|
||||
GGML_LOG_ERROR("ggml-hex: %s lazy mapping of buffer context failed: %s\n", this->c_name(), exc.what());
|
||||
return false;
|
||||
}
|
||||
|
||||
this->cloned_buffers[fd] = std::move(clone);
|
||||
mem->mapped_clones.insert(this);
|
||||
return true;
|
||||
mem->clones.insert(this);
|
||||
}
|
||||
|
||||
void ggml_hexagon_session::release_buffer(const ggml_hexagon_shared_buffer * sbuf) {
|
||||
@@ -2000,14 +2072,14 @@ void ggml_hexagon_session::release_buffer(const ggml_hexagon_shared_buffer * sbu
|
||||
auto clone = std::move(it->second);
|
||||
this->cloned_buffers.erase(it);
|
||||
}
|
||||
mem->mapped_clones.erase(this);
|
||||
mem->clones.erase(this);
|
||||
}
|
||||
|
||||
void ggml_hexagon_session::unclone_buffer(const ggml_hexagon_shared_buffer * sbuf) {
|
||||
GGML_ASSERT(sbuf && sbuf->mem);
|
||||
|
||||
auto mem = sbuf->mem;
|
||||
std::vector<ggml_hexagon_session *> sessions(mem->mapped_clones.begin(), mem->mapped_clones.end());
|
||||
std::vector<ggml_hexagon_session *> sessions(mem->clones.begin(), mem->clones.end());
|
||||
|
||||
for (auto * sess : sessions) {
|
||||
sess->release_buffer(sbuf);
|
||||
@@ -2276,9 +2348,14 @@ void ggml_hexagon_session::allocate(const ggml_hexagon_device_config & config) n
|
||||
}
|
||||
|
||||
// Allocate buffers and state for op batching
|
||||
this->op_queue = new ggml_hexagon_opqueue(this, opt_opbatch, opt_opqueue);
|
||||
const size_t n_fences = 512;
|
||||
const size_t opq_size = ggml_hexagon_opqueue::size_of(opt_opbatch, this->n_threads, opt_opqueue);
|
||||
const size_t fence_offset = GGML_PAD(opq_size, HEX_L2_LINE_SIZE);
|
||||
const size_t fence_size = ggml_hexagon_fence_buffer::size_of(n_fences);
|
||||
|
||||
this->fence_buf = new ggml_hexagon_fence_buffer(this, &dev_ctx->fence_buffer_type, 64 * 1024);
|
||||
this->shm_buf = new ggml_hexagon_shared_buffer(this, fence_offset + fence_size, true /* pinned */);
|
||||
this->op_queue = new ggml_hexagon_opqueue(this->shm_buf, 0, opt_opbatch, opt_opqueue);
|
||||
this->fence_buf = new ggml_hexagon_fence_buffer(this->shm_buf, &dev_ctx->fence_buffer_type, fence_offset, n_fences);
|
||||
if (this->mdev.count > 1) {
|
||||
this->mdev_fence_slot = this->alloc_fence(this->mdev.count);
|
||||
}
|
||||
@@ -2287,7 +2364,7 @@ void ggml_hexagon_session::allocate(const ggml_hexagon_device_config & config) n
|
||||
opt_vmem = ggml_hexagon_measure_max_vmem(this);
|
||||
GGML_LOG_INFO("ggml-hex: %s measured max vmem %zu\n", this->c_name(), opt_vmem);
|
||||
}
|
||||
const size_t shm_size = this->op_queue->shm_size();
|
||||
const size_t shm_size = this->shm_buf->size();
|
||||
this->max_vmem = (opt_vmem > shm_size) ? (opt_vmem - shm_size) : opt_vmem;
|
||||
|
||||
this->op_batch = new ggml_hexagon_opbatch(this, opt_opbatch, this->max_vmem);
|
||||
@@ -2328,8 +2405,6 @@ void ggml_hexagon_session::release() noexcept(true) {
|
||||
}
|
||||
}
|
||||
|
||||
delete this->op_batch;
|
||||
delete this->op_queue;
|
||||
for (auto & it : this->cpy_fence_slots) {
|
||||
free_fence((void *) it.second, 1);
|
||||
}
|
||||
@@ -2340,6 +2415,14 @@ void ggml_hexagon_session::release() noexcept(true) {
|
||||
delete this->fence_buf;
|
||||
this->fence_buf = nullptr;
|
||||
}
|
||||
|
||||
delete this->op_batch;
|
||||
delete this->op_queue;
|
||||
delete this->shm_buf;
|
||||
this->op_batch = nullptr;
|
||||
this->op_queue = nullptr;
|
||||
this->shm_buf = nullptr;
|
||||
|
||||
while (!this->cloned_buffers.empty()) {
|
||||
release_buffer(this->cloned_buffers.begin()->second.get());
|
||||
}
|
||||
@@ -2734,19 +2817,12 @@ static bool ggml_hexagon_cpy_tensor_async_phys(ggml_backend_t backend_src, ggml_
|
||||
auto sess_src = static_cast<ggml_hexagon_session *>(backend_src->context);
|
||||
auto sess_dst = static_cast<ggml_hexagon_session *>(backend_dst->context);
|
||||
|
||||
sess_src->mmap_tensor(src);
|
||||
auto sbuf_dst = sess_dst->mmap_tensor(dst);
|
||||
|
||||
if (!sess_src->clone_buffer(sbuf_dst)) { return false; }
|
||||
|
||||
const uint64_t src_key = ggml_hexagon_session_key(sess_src);
|
||||
auto & fence_slot = sess_dst->cpy_fence_slots[src_key];
|
||||
if (!fence_slot) {
|
||||
fence_slot = (volatile uint32_t *) sess_dst->alloc_fence(1);
|
||||
}
|
||||
|
||||
if (!sess_src->clone_buffer(sess_dst->fence_buf)) { return false; }
|
||||
|
||||
uint32_t fence_seq = sess_dst->next_fence_seq();
|
||||
|
||||
HEX_VERBOSE("ggml-hex: %s cpy-tensor-async %s -> %s size %zu : seq 0x%x\n",
|
||||
@@ -2766,11 +2842,6 @@ static bool ggml_hexagon_cpy_tensor_async_virt(ggml_backend_t backend_src, ggml_
|
||||
auto sess_src = static_cast<ggml_hexagon_session *>(backend_src->context);
|
||||
auto sess_dst = static_cast<ggml_hexagon_session *>(backend_dst->context);
|
||||
|
||||
auto sbuf_src = sess_src->mmap_tensor(src);
|
||||
sess_dst->mmap_tensor(dst);
|
||||
|
||||
if (!sess_dst->clone_buffer(sbuf_src)) { return false; }
|
||||
|
||||
HEX_VERBOSE("ggml-hex: %s cpy-tensor-async %s -> %s size %zu\n",
|
||||
sess_dst->name.c_str(), src->name, dst->name, ggml_nbytes(src));
|
||||
|
||||
@@ -2893,7 +2964,6 @@ static void ggml_backend_hexagon_event_wait(ggml_backend_t backend, ggml_backend
|
||||
return;
|
||||
}
|
||||
|
||||
sess->clone_buffer(hex_event->fence_sess->fence_buf);
|
||||
sess->add_peer(hex_event->sess);
|
||||
sess->enqueue_fence(&hex_event->fence_tensor, hex_event->seq, /* wait = */ true);
|
||||
}
|
||||
|
||||
@@ -82,6 +82,7 @@ struct htp_ops_context {
|
||||
struct htp_context {
|
||||
dspqueue_t dsp_queue;
|
||||
|
||||
struct htp_mmap shm;
|
||||
struct htp_mmap mmap[HTP_MAX_MMAPS];
|
||||
dma_queue_t dma[HTP_MAX_NTHREADS];
|
||||
struct htp_thread_trace trace[HTP_MAX_NTHREADS];
|
||||
|
||||
@@ -1089,7 +1089,12 @@ static int execute_op(struct htp_ops_context * octx) {
|
||||
}
|
||||
|
||||
static inline bool reuse_buf(struct htp_context *ctx, uint32_t *m_reuse, struct htp_buf_desc *b) {
|
||||
b->base = NULL;
|
||||
if (b->fd == ctx->shm.fd) {
|
||||
b->base = ctx->shm.base;
|
||||
return true;
|
||||
}
|
||||
|
||||
b->base = 0;
|
||||
|
||||
for (uint32_t i=0; i<HTP_MAX_MMAPS; i++) {
|
||||
struct htp_mmap *m = ctx->mmap + i;
|
||||
@@ -1203,7 +1208,7 @@ static void prep_op_bufs(struct htp_context *ctx, struct htp_buf_desc *bufs, uin
|
||||
|
||||
for (uint32_t i=0; i < n_bufs; i++) {
|
||||
struct htp_buf_desc *b = bufs + i;
|
||||
if (!(b->flags & HTP_BUF_EXTENDED)) {
|
||||
if (!(b->flags & HTP_BUF_EXTENDED) && b->fd != ctx->shm.fd) {
|
||||
b->base = 0;
|
||||
}
|
||||
if (!mmap_buf(ctx, b)) {
|
||||
@@ -1316,6 +1321,9 @@ static void process_opbatch(struct htp_context * ctx, const struct htp_opbatch_r
|
||||
n_bufs, n_tens, n_ops, req->n_traces, dbuf->size, b_size, t_size, o_size);
|
||||
|
||||
// Setup descriptor pointers
|
||||
ctx->shm.fd = dbuf->fd;
|
||||
ctx->shm.base = (uintptr_t) dbuf->ptr - dbuf->offset;
|
||||
|
||||
uint8_t * m_ptr = dbuf->ptr;
|
||||
struct htp_buf_desc* bufs = (struct htp_buf_desc*) m_ptr; m_ptr += b_size;
|
||||
struct htp_tensor* tens = (struct htp_tensor*) m_ptr; m_ptr += t_size;
|
||||
|
||||
@@ -1708,8 +1708,9 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
|
||||
case GGML_OP_ARGMAX:
|
||||
return has_simdgroup_reduction;
|
||||
case GGML_OP_NORM:
|
||||
case GGML_OP_RMS_NORM:
|
||||
return has_simdgroup_reduction && (ggml_is_contiguous_rows(op->src[0]));
|
||||
case GGML_OP_RMS_NORM:
|
||||
return has_simdgroup_reduction && op->src[0]->type == GGML_TYPE_F32 && (ggml_is_contiguous_rows(op->src[0]));
|
||||
case GGML_OP_ROPE:
|
||||
case GGML_OP_ROPE_BACK:
|
||||
return ggml_rope_get_truncate(op);
|
||||
|
||||
@@ -11966,7 +11966,7 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
|
||||
case GGML_OP_NORM:
|
||||
return true;
|
||||
case GGML_OP_RMS_NORM:
|
||||
return op->ne[0] % 4 == 0 && ggml_is_contiguous_rows(op->src[0]);
|
||||
return op->src[0]->type == GGML_TYPE_F32 && op->ne[0] % 4 == 0 && ggml_is_contiguous_rows(op->src[0]);
|
||||
case GGML_OP_L2_NORM:
|
||||
return ggml_is_contiguous_rows(op->src[0]);
|
||||
case GGML_OP_REPEAT:
|
||||
|
||||
@@ -1455,6 +1455,12 @@ static ggml_openvino_op_support is_op_supported_case(const ggml_tensor * op) {
|
||||
}
|
||||
break;
|
||||
}
|
||||
case GGML_OP_RMS_NORM: {
|
||||
if (op->type == GGML_TYPE_BF16) {
|
||||
return {false, "RMS_NORM with BF16 type is not supported"};
|
||||
}
|
||||
break;
|
||||
}
|
||||
case GGML_OP_ADD_ID: {
|
||||
// Keep support aligned with the CPU backend implementation, which only handles f32 inputs/output and i32 ids.
|
||||
if (op->type != GGML_TYPE_F32 || op->src[0]->type != GGML_TYPE_F32 || op->src[1]->type != GGML_TYPE_F32 ||
|
||||
|
||||
@@ -7598,8 +7598,9 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
|
||||
case GGML_OP_NORM:
|
||||
case GGML_OP_L2_NORM:
|
||||
case GGML_OP_GROUP_NORM:
|
||||
case GGML_OP_RMS_NORM:
|
||||
return true;
|
||||
case GGML_OP_RMS_NORM:
|
||||
return op->src[0]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32;
|
||||
case GGML_OP_RMS_NORM_BACK:
|
||||
return ggml_is_contiguous(op->src[0]);
|
||||
case GGML_OP_SCALE:
|
||||
|
||||
@@ -109,6 +109,7 @@ void ggml_vk_quantize_q8_1(ggml_backend_vk_context * ctx, vk_context& subctx, co
|
||||
void ggml_vk_dsv4_hc_comb(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * mixes, const ggml_tensor * scale, const ggml_tensor * base, ggml_tensor * dst);
|
||||
void ggml_vk_dsv4_hc_pre(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * x, const ggml_tensor * weights, ggml_tensor * dst);
|
||||
void ggml_vk_dsv4_hc_post(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * x, const ggml_tensor * residual, const ggml_tensor * post, const ggml_tensor * comb, ggml_tensor * dst, const ggml_tensor * gate_scale_in = nullptr);
|
||||
void ggml_vk_top_k_mask(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst);
|
||||
void ggml_vk_mul_mat(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx);
|
||||
bool ggml_vk_use_mul_mat_vec_id(const struct ggml_cgraph * cgraph, int node_idx);
|
||||
void ggml_vk_mul_mat_id(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx);
|
||||
|
||||
@@ -219,6 +219,29 @@ struct vk_op_dsv4_hc_post_push_constants {
|
||||
float gate_scale_out;
|
||||
};
|
||||
|
||||
struct vk_op_top_k_mask_push_constants {
|
||||
uint32_t n_kv;
|
||||
uint32_t n_top;
|
||||
uint32_t n_pick;
|
||||
uint32_t n_blk;
|
||||
uint32_t n_extra;
|
||||
uint32_t has_score;
|
||||
uint32_t ninf;
|
||||
|
||||
uint32_t ne1; uint32_t ne2;
|
||||
|
||||
uint32_t m1; uint32_t m2; uint32_t m3;
|
||||
uint32_t nbm1; uint32_t nbm2; uint32_t nbm3;
|
||||
uint32_t nbd1; uint32_t nbd2; uint32_t nbd3;
|
||||
|
||||
uint32_t m_offset;
|
||||
uint32_t k_offset;
|
||||
uint32_t s_offset;
|
||||
uint32_t b_offset;
|
||||
uint32_t e_offset;
|
||||
uint32_t d_offset;
|
||||
};
|
||||
|
||||
struct vk_op_count_experts_push_constants {
|
||||
uint32_t ne00;
|
||||
uint32_t ne01;
|
||||
@@ -972,6 +995,18 @@ template <> inline void init_pushconst_tensor_offsets(ggml_backend_vk_context *
|
||||
p.d_offset = get_misalign_bytes(ctx, dst) / ggml_type_size(dst->type);
|
||||
}
|
||||
|
||||
// the fifth source, extra, comes from dst
|
||||
template <> inline void init_pushconst_tensor_offsets(ggml_backend_vk_context * ctx, vk_op_top_k_mask_push_constants &p, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, const ggml_tensor * src3, ggml_tensor * dst) {
|
||||
const ggml_tensor * extra = dst->src[4];
|
||||
|
||||
p.m_offset = get_misalign_bytes(ctx, src0) / ggml_type_size(src0->type);
|
||||
p.k_offset = get_misalign_bytes(ctx, src1) / ggml_type_size(src1->type);
|
||||
p.s_offset = src2 ? get_misalign_bytes(ctx, src2) / ggml_type_size(src2->type) : 0;
|
||||
p.b_offset = src3 ? get_misalign_bytes(ctx, src3) / ggml_type_size(src3->type) : 0;
|
||||
p.e_offset = extra ? get_misalign_bytes(ctx, extra) / ggml_type_size(extra->type) : 0;
|
||||
p.d_offset = get_misalign_bytes(ctx, dst) / ggml_type_size(dst->type);
|
||||
}
|
||||
|
||||
template <typename T> size_t push_constant_size(const T &t) {
|
||||
static_assert(std::is_class<T>::value, "T must be a struct/class");
|
||||
GGML_UNUSED(t);
|
||||
|
||||
@@ -977,6 +977,8 @@ struct vk_device_struct {
|
||||
vk_pipeline pipeline_dsv4_hc_pre_gated_f32;
|
||||
vk_pipeline pipeline_dsv4_hc_post_f32;
|
||||
vk_pipeline pipeline_dsv4_hc_post_nocomb_f32;
|
||||
vk_pipeline pipeline_top_k_mask_f16;
|
||||
vk_pipeline pipeline_top_k_mask_f32;
|
||||
std::map<vk_solve_tri_pipeline_state, vk_pipeline> pipeline_solve_tri_f32;
|
||||
vk_pipeline pipeline_im2col_f32, pipeline_im2col_f32_f16;
|
||||
vk_pipeline pipeline_im2col_3d_f32, pipeline_im2col_3d_f32_f16;
|
||||
|
||||
@@ -3823,6 +3823,10 @@ void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
|
||||
ggml_vk_create_pipeline(device, device->pipeline_dsv4_hc_post_f32, "dsv4_hc_post_f32", dsv4_hc_post_f32_len, dsv4_hc_post_f32_data, "main", 5, sizeof(vk_op_dsv4_hc_post_push_constants), {256, 1, 1}, { 256, 1 }, 1);
|
||||
ggml_vk_create_pipeline(device, device->pipeline_dsv4_hc_post_nocomb_f32,"dsv4_hc_post_nocomb_f32",dsv4_hc_post_f32_len, dsv4_hc_post_f32_data, "main", 5, sizeof(vk_op_dsv4_hc_post_push_constants), {256, 1, 1}, { 256, 0 }, 1);
|
||||
|
||||
const uint32_t top_k_mask_wg = std::min({256u, device->properties.limits.maxComputeWorkGroupInvocations, device->properties.limits.maxComputeWorkGroupSize[0]});
|
||||
ggml_vk_create_pipeline(device, device->pipeline_top_k_mask_f16, "top_k_mask_f16", top_k_mask_f16_len, top_k_mask_f16_data, "main", 6, sizeof(vk_op_top_k_mask_push_constants), {1, 1, 1}, {top_k_mask_wg}, 1);
|
||||
ggml_vk_create_pipeline(device, device->pipeline_top_k_mask_f32, "top_k_mask_f32", top_k_mask_f32_len, top_k_mask_f32_data, "main", 6, sizeof(vk_op_top_k_mask_push_constants), {1, 1, 1}, {top_k_mask_wg}, 1);
|
||||
|
||||
for (auto &s : device->pipeline_solve_tri_f32) {
|
||||
const vk_solve_tri_pipeline_state &state = s.first;
|
||||
|
||||
@@ -7625,6 +7629,46 @@ void ggml_vk_dsv4_hc_post(ggml_backend_vk_context * ctx, vk_context& subctx, con
|
||||
ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { x_buf, r_buf, p_buf, c_buf, d_buf }, pc, { n_embd, n_tokens, 1 });
|
||||
}
|
||||
|
||||
void ggml_vk_top_k_mask(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst) {
|
||||
const ggml_tensor * mask = dst->src[0];
|
||||
const ggml_tensor * top_k = dst->src[1];
|
||||
const ggml_tensor * score = dst->src[2];
|
||||
const ggml_tensor * blocks = dst->src[3];
|
||||
const ggml_tensor * extra = dst->src[4];
|
||||
VK_LOG_DEBUG("ggml_vk_top_k_mask(" << mask << ", " << top_k << ", " << score << ", " << blocks << ", " << extra << ", " << dst << ")");
|
||||
|
||||
vk_pipeline pipeline = dst->type == GGML_TYPE_F16 ? ctx->device->pipeline_top_k_mask_f16 : ctx->device->pipeline_top_k_mask_f32;
|
||||
GGML_ASSERT(pipeline != nullptr);
|
||||
|
||||
ggml_pipeline_request_descriptor_sets(ctx, pipeline, 1);
|
||||
|
||||
const vk_subbuffer m_buf = ggml_vk_tensor_subbuffer(ctx, mask, true);
|
||||
const vk_subbuffer k_buf = ggml_vk_tensor_subbuffer(ctx, top_k, true);
|
||||
const vk_subbuffer s_buf = score ? ggml_vk_tensor_subbuffer(ctx, score, true) : k_buf;
|
||||
const vk_subbuffer b_buf = blocks ? ggml_vk_tensor_subbuffer(ctx, blocks, true) : k_buf;
|
||||
const vk_subbuffer e_buf = extra ? ggml_vk_tensor_subbuffer(ctx, extra, true) : k_buf;
|
||||
const vk_subbuffer d_buf = ggml_vk_tensor_subbuffer(ctx, dst, true);
|
||||
|
||||
vk_op_top_k_mask_push_constants pc = {
|
||||
(uint32_t)mask->ne[0],
|
||||
(uint32_t)top_k->ne[0],
|
||||
(uint32_t)(blocks ? blocks->ne[1] : score ? score->ne[0] : mask->ne[0]),
|
||||
blocks ? (uint32_t)blocks->ne[0] : 0,
|
||||
extra ? (uint32_t)extra->ne[0] : 0,
|
||||
score ? 1u : 0u,
|
||||
dst->type == GGML_TYPE_F16 ? 0xFC00u : 0xFF800000u,
|
||||
(uint32_t)dst->ne[1], (uint32_t)dst->ne[2],
|
||||
(uint32_t)mask->ne[1], (uint32_t)mask->ne[2], (uint32_t)mask->ne[3],
|
||||
ggml_vk_nb_elem(mask, 1), ggml_vk_nb_elem(mask, 2), ggml_vk_nb_elem(mask, 3),
|
||||
ggml_vk_nb_elem(dst, 1), ggml_vk_nb_elem(dst, 2), ggml_vk_nb_elem(dst, 3),
|
||||
0, 0, 0, 0, 0, 0,
|
||||
};
|
||||
init_pushconst_tensor_offsets(ctx, pc, mask, top_k, score, blocks, dst);
|
||||
|
||||
ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { m_buf, k_buf, s_buf, b_buf, e_buf, d_buf }, pc,
|
||||
{ (uint32_t)dst->ne[1], (uint32_t)dst->ne[2], (uint32_t)dst->ne[3] });
|
||||
}
|
||||
|
||||
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];
|
||||
@@ -13211,6 +13255,11 @@ bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgraph, in
|
||||
|
||||
break;
|
||||
|
||||
case GGML_OP_TOP_K_MASK:
|
||||
ggml_vk_top_k_mask(ctx, compute_ctx, node);
|
||||
|
||||
break;
|
||||
|
||||
case GGML_OP_GATED_DELTA_NET:
|
||||
ggml_vk_gated_delta_net(ctx, compute_ctx, node);
|
||||
|
||||
@@ -16264,11 +16313,11 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
|
||||
case GGML_OP_VIEW:
|
||||
case GGML_OP_PERMUTE:
|
||||
case GGML_OP_TRANSPOSE:
|
||||
case GGML_OP_RMS_NORM:
|
||||
return true;
|
||||
case GGML_OP_GROUP_NORM:
|
||||
return ggml_is_contiguous(op->src[0]);
|
||||
case GGML_OP_NORM:
|
||||
case GGML_OP_RMS_NORM:
|
||||
case GGML_OP_L2_NORM:
|
||||
return op->src[0]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32;
|
||||
case GGML_OP_ADD:
|
||||
@@ -16488,6 +16537,30 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
|
||||
}
|
||||
return true;
|
||||
}
|
||||
case GGML_OP_TOP_K_MASK:
|
||||
{
|
||||
const ggml_tensor * mask = op->src[0];
|
||||
if ((op->type != GGML_TYPE_F16 && op->type != GGML_TYPE_F32) || mask->type != op->type) {
|
||||
return false;
|
||||
}
|
||||
|
||||
// one workgroup per row
|
||||
const auto & max_wg = device->properties.limits.maxComputeWorkGroupCount;
|
||||
if (op->ne[1] > max_wg[0] || op->ne[2] > max_wg[1] || op->ne[3] > max_wg[2]) {
|
||||
return false;
|
||||
}
|
||||
|
||||
// the shader expects contiguous rows in mask and dst, and contiguous tensors in the other inputs
|
||||
if (!ggml_is_contiguous_rows(op) || !ggml_is_contiguous_rows(mask)) {
|
||||
return false;
|
||||
}
|
||||
for (int i = 1; i <= 4; ++i) {
|
||||
if (op->src[i] && !ggml_is_contiguous(op->src[i])) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
return true;
|
||||
}
|
||||
case GGML_OP_GATED_DELTA_NET:
|
||||
{
|
||||
const uint32_t S_v = op->src[2]->ne[0];
|
||||
|
||||
@@ -0,0 +1,95 @@
|
||||
#version 450
|
||||
|
||||
#extension GL_EXT_shader_16bit_storage : require
|
||||
#extension GL_KHR_memory_scope_semantics : enable
|
||||
|
||||
// One workgroup per row: fill the row with -inf, then copy the mask cells of the live picks and the extra cells.
|
||||
// The op only moves bits, so f16 needs no float16 arithmetic and the score check does not depend on float controls.
|
||||
|
||||
layout(constant_id = 0) const uint BLOCK_SIZE = 256;
|
||||
|
||||
layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
layout(push_constant) uniform parameter
|
||||
{
|
||||
uint n_kv;
|
||||
uint n_top;
|
||||
uint n_pick;
|
||||
uint n_blk; // 0: a pick is a cell
|
||||
uint n_extra;
|
||||
uint has_score;
|
||||
uint ninf; // bits of -inf: a 16-bit literal would need shaderInt16
|
||||
|
||||
uint ne1; uint ne2;
|
||||
|
||||
uint m1; uint m2; uint m3;
|
||||
uint nbm1; uint nbm2; uint nbm3;
|
||||
uint nbd1; uint nbd2; uint nbd3;
|
||||
|
||||
uint m_offset;
|
||||
uint k_offset;
|
||||
uint s_offset;
|
||||
uint b_offset;
|
||||
uint e_offset;
|
||||
uint d_offset;
|
||||
};
|
||||
|
||||
layout(binding = 0, std430) readonly buffer M { MASK_TYPE data_m[]; };
|
||||
layout(binding = 1, std430) readonly buffer K { int data_k[]; };
|
||||
layout(binding = 2, std430) readonly buffer S { uint data_s[]; };
|
||||
layout(binding = 3, std430) readonly buffer B { int data_b[]; };
|
||||
layout(binding = 4, std430) readonly buffer E { int data_e[]; };
|
||||
// coherent: the fill and the copy of a cell can come from different invocations
|
||||
layout(binding = 5, std430) coherent writeonly buffer D { MASK_TYPE data_d[]; };
|
||||
|
||||
void keep(const int c, const uint m_row, const uint d_row) {
|
||||
if (c >= 0 && uint(c) < n_kv) {
|
||||
data_d[d_row + uint(c)] = data_m[m_row + uint(c)];
|
||||
}
|
||||
}
|
||||
|
||||
// score > -inf
|
||||
bool is_live(const uint s) {
|
||||
return s != 0xFF800000u && (s & 0x7FFFFFFFu) <= 0x7F800000u;
|
||||
}
|
||||
|
||||
void main() {
|
||||
const uint tid = gl_LocalInvocationIndex;
|
||||
|
||||
const uint i1 = gl_WorkGroupID.x;
|
||||
const uint i2 = gl_WorkGroupID.y;
|
||||
const uint i3 = gl_WorkGroupID.z;
|
||||
|
||||
const uint row = i1 + ne1*(i2 + ne2*i3);
|
||||
const uint m_row = m_offset + (i1 % m1)*nbm1 + (i2 % m2)*nbm2 + (i3 % m3)*nbm3;
|
||||
const uint d_row = d_offset + i1*nbd1 + i2*nbd2 + i3*nbd3;
|
||||
|
||||
for (uint i = tid; i < n_kv; i += BLOCK_SIZE) {
|
||||
data_d[d_row + i] = MASK_TYPE(ninf);
|
||||
}
|
||||
|
||||
controlBarrier(gl_ScopeWorkgroup, gl_ScopeWorkgroup, gl_StorageSemanticsBuffer, gl_SemanticsAcquireRelease);
|
||||
|
||||
// one invocation per pick or extra cell; cells kept twice get the same value
|
||||
for (uint i = tid; i < n_top + n_extra; i += BLOCK_SIZE) {
|
||||
if (i >= n_top) {
|
||||
keep(data_e[e_offset + row*n_extra + (i - n_top)], m_row, d_row);
|
||||
continue;
|
||||
}
|
||||
|
||||
const int p = data_k[k_offset + row*n_top + i];
|
||||
if (p < 0 || uint(p) >= n_pick) {
|
||||
continue;
|
||||
}
|
||||
if (has_score != 0 && !is_live(data_s[s_offset + row*n_pick + uint(p)])) {
|
||||
continue;
|
||||
}
|
||||
if (n_blk == 0) {
|
||||
keep(p, m_row, d_row);
|
||||
continue;
|
||||
}
|
||||
for (uint k = 0; k < n_blk; ++k) {
|
||||
keep(data_b[b_offset + uint(p)*n_blk + k], m_row, d_row);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -1120,6 +1120,8 @@ void process_shaders() {
|
||||
string_to_spv("dsv4_hc_comb_f32", "dsv4_hc_comb.comp", {});
|
||||
string_to_spv("dsv4_hc_pre_f32", "dsv4_hc_pre.comp", {});
|
||||
string_to_spv("dsv4_hc_post_f32", "dsv4_hc_post.comp", {});
|
||||
string_to_spv("top_k_mask_f16", "top_k_mask.comp", {{"MASK_TYPE", "uint16_t"}});
|
||||
string_to_spv("top_k_mask_f32", "top_k_mask.comp", {{"MASK_TYPE", "uint"}});
|
||||
string_to_spv("cumsum_f32", "cumsum.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));
|
||||
string_to_spv("cumsum_multipass1_f32", "cumsum_multipass1.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));
|
||||
string_to_spv("cumsum_multipass2_f32", "cumsum_multipass2.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));
|
||||
|
||||
+7
-4
@@ -6,7 +6,7 @@ version = "0.0.0"
|
||||
dynamic = ["classifiers"]
|
||||
readme = "README.md"
|
||||
authors = [{name = "GGML", email = "ggml@ggml.ai"}]
|
||||
requires-python = '>=3.10,<3.15'
|
||||
requires-python = '>=3.10,<3.16'
|
||||
dependencies = [
|
||||
'numpy (>=2.2.6,<3.0.0); python_version < "3.14"',
|
||||
'numpy (>=2.3.0,<3.0.0); python_version >= "3.14"',
|
||||
@@ -40,9 +40,12 @@ packages = [
|
||||
|
||||
[tool.poetry.dependencies]
|
||||
torch = [
|
||||
{ version = "==2.11.0", source = "pypi", markers = "sys_platform == 'darwin'" },
|
||||
{ version = "==2.11.0+cpu", source = "pytorch", markers = "sys_platform == 'linux'" },
|
||||
{ version = "==2.11.0", source = "pypi", markers = "sys_platform == 'win32'" },
|
||||
{ version = "==2.11.0", source = "pypi", markers = "sys_platform == 'darwin' and python_version < '3.15'" },
|
||||
{ version = "==2.11.0+cpu", source = "pytorch", markers = "sys_platform == 'linux' and python_version < '3.15'" },
|
||||
{ version = "==2.11.0", source = "pypi", markers = "sys_platform == 'win32' and python_version < '3.15'" },
|
||||
{ version = "==2.15.0", source = "pypi", markers = "sys_platform == 'darwin' and python_version >= '3.15'" },
|
||||
{ version = "==2.15.0+cpu", source = "pytorch", markers = "sys_platform == 'linux' and python_version >= '3.15'" },
|
||||
{ version = "==2.15.0", source = "pypi", markers = "sys_platform == 'win32' and python_version >= '3.15'" },
|
||||
]
|
||||
|
||||
[tool.poetry.group.dev.dependencies]
|
||||
|
||||
@@ -1,3 +1,4 @@
|
||||
tabulate~=0.9.0
|
||||
GitPython~=3.1.43
|
||||
matplotlib~=3.10.0
|
||||
matplotlib~=3.10.0; python_version < "3.11"
|
||||
matplotlib~=3.11.2; python_version >= "3.11"
|
||||
|
||||
@@ -2,4 +2,5 @@
|
||||
--extra-index-url https://download.pytorch.org/whl/cpu
|
||||
|
||||
## Embedding Gemma requires PyTorch 2.6.0 or later, bumped to 2.11.0 for compatibility
|
||||
torch==2.11.0
|
||||
torch==2.11.0; python_version < "3.15"
|
||||
torch==2.15.0; python_version >= "3.15"
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
numpy~=2.2.6; python_version < "3.14"
|
||||
numpy~=2.4.6; python_version >= "3.14"
|
||||
numpy~=2.5.4; python_version >= "3.14"
|
||||
sentencepiece>=0.1.98,<0.3.0
|
||||
|
||||
transformers==4.57.6
|
||||
|
||||
@@ -1,4 +1,4 @@
|
||||
numpy~=2.2.6; python_version < "3.14"
|
||||
numpy~=2.4.6; python_version >= "3.14"
|
||||
numpy~=2.5.4; python_version >= "3.14"
|
||||
PySide6~=6.9.0
|
||||
gguf>=0.17.0
|
||||
|
||||
@@ -1,6 +1,7 @@
|
||||
datasets~=4.8.0
|
||||
matplotlib~=3.10.0
|
||||
matplotlib~=3.10.0; python_version < "3.11"
|
||||
matplotlib~=3.11.2; python_version >= "3.11"
|
||||
numpy~=2.2.6; python_version < "3.14"
|
||||
numpy~=2.4.6; python_version >= "3.14"
|
||||
numpy~=2.5.4; python_version >= "3.14"
|
||||
requests~=2.32.3
|
||||
tqdm~=4.67.1
|
||||
|
||||
@@ -1 +1 @@
|
||||
cffi~=1.16.0
|
||||
cffi~=2.1.1
|
||||
|
||||
@@ -1,9 +1,11 @@
|
||||
pytest~=8.3.3
|
||||
matplotlib~=3.10.0
|
||||
matplotlib~=3.10.0; python_version < "3.11"
|
||||
matplotlib~=3.11.2; python_version >= "3.11"
|
||||
numpy~=2.2.6; python_version < "3.14"
|
||||
numpy~=2.4.6; python_version >= "3.14"
|
||||
numpy~=2.5.4; python_version >= "3.14"
|
||||
openai~=2.14.0
|
||||
pandas~=2.3.3
|
||||
pandas~=2.3.3; python_version < "3.11"
|
||||
pandas~=3.0.6; python_version >= "3.11"
|
||||
prometheus-client~=0.20.0
|
||||
requests~=2.32.3
|
||||
wget~=3.2
|
||||
|
||||
@@ -68,8 +68,9 @@ log = logging.getLogger(__name__)
|
||||
|
||||
POLL_INTERVAL = 30
|
||||
JOB_TIMEOUT = 3600
|
||||
RUNNER_TIMEOUT = 300
|
||||
LOG_UPLOAD_TIMEOUT = 600
|
||||
CAPACITY_TIMEOUT = 1800
|
||||
CAPACITY_TIMEOUT = 300
|
||||
CAPACITY_POLL = 60
|
||||
MAX_CONCURRENT_JOBS = 5
|
||||
DEFAULT_RETRIES = 0
|
||||
@@ -82,6 +83,10 @@ class DeviceUnavailableError(Exception):
|
||||
"""Raised when the QDC device resource is not available (retryable)."""
|
||||
|
||||
|
||||
class RunnerUnavailableError(Exception):
|
||||
"""Raised when QDC runner or capacity is not available within timeout (non-retryable)."""
|
||||
|
||||
|
||||
_SCRIPTS_DIR = Path(__file__).parent
|
||||
_TESTS_DIR = _SCRIPTS_DIR / "tests"
|
||||
|
||||
@@ -253,8 +258,9 @@ DEVICE_PLATFORM: dict[str, Platform] = {
|
||||
# =============================================================================
|
||||
|
||||
|
||||
def wait_for_job(client, job_id: str, timeout: int) -> str:
|
||||
def wait_for_job(client, job_id: str, timeout: int, runner_timeout: int = RUNNER_TIMEOUT) -> str:
|
||||
elapsed = 0
|
||||
submitted_elapsed = 0
|
||||
last_state = None
|
||||
consecutive_errors = 0
|
||||
max_consecutive_errors = 5
|
||||
@@ -285,6 +291,26 @@ def wait_for_job(client, job_id: str, timeout: int) -> str:
|
||||
if raw != last_state:
|
||||
log.info("Job %s: %s", job_id, raw)
|
||||
last_state = raw
|
||||
|
||||
is_submitted = (status == JobState.SUBMITTED) or (
|
||||
isinstance(raw, str) and raw.lower() == "submitted"
|
||||
)
|
||||
if is_submitted:
|
||||
submitted_elapsed += POLL_INTERVAL
|
||||
if submitted_elapsed >= runner_timeout:
|
||||
try:
|
||||
qdc_api.abort_job(client, job_id)
|
||||
log.warning(
|
||||
"Aborted job %s: runner did not pick up job within %ds",
|
||||
job_id,
|
||||
runner_timeout,
|
||||
)
|
||||
except Exception as e:
|
||||
log.warning("Failed to abort job %s: %s", job_id, e)
|
||||
raise RunnerUnavailableError(
|
||||
f"Runner did not pick up job {job_id} within {runner_timeout}s"
|
||||
)
|
||||
|
||||
time.sleep(POLL_INTERVAL)
|
||||
elapsed += POLL_INTERVAL
|
||||
# Abort to free the QDC concurrency slot instead of leaking it.
|
||||
@@ -312,10 +338,10 @@ def wait_for_log_upload(client, job_id: str) -> None:
|
||||
log.warning("Timed out waiting for log upload after %ds", LOG_UPLOAD_TIMEOUT)
|
||||
|
||||
|
||||
def wait_for_capacity(client, max_jobs: int = MAX_CONCURRENT_JOBS) -> None:
|
||||
def wait_for_capacity(client, timeout: int = CAPACITY_TIMEOUT, max_jobs: int = MAX_CONCURRENT_JOBS) -> None:
|
||||
"""Block until the user's active (non-terminal) QDC job count is below max_jobs."""
|
||||
elapsed = 0
|
||||
while elapsed < CAPACITY_TIMEOUT:
|
||||
while elapsed < timeout:
|
||||
jobs_page = qdc_api.get_jobs_list(client, page_number=0, page_size=50)
|
||||
if jobs_page is None:
|
||||
log.warning(
|
||||
@@ -338,7 +364,7 @@ def wait_for_capacity(client, max_jobs: int = MAX_CONCURRENT_JOBS) -> None:
|
||||
time.sleep(CAPACITY_POLL)
|
||||
elapsed += CAPACITY_POLL
|
||||
raise TimeoutError(
|
||||
f"Capacity wait timed out after {CAPACITY_TIMEOUT}s"
|
||||
f"Capacity wait timed out after {timeout}s"
|
||||
)
|
||||
|
||||
|
||||
@@ -445,7 +471,11 @@ def fetch_logs_and_parse_tests(
|
||||
)
|
||||
|
||||
|
||||
def write_summary(result: JobResult, title: str = "QDC Test Results") -> None:
|
||||
def write_summary(
|
||||
result: JobResult,
|
||||
title: str = "QDC Test Results",
|
||||
note: str | None = None,
|
||||
) -> None:
|
||||
summary_path = os.environ.get("GITHUB_STEP_SUMMARY")
|
||||
if not summary_path:
|
||||
return
|
||||
@@ -454,8 +484,12 @@ def write_summary(result: JobResult, title: str = "QDC Test Results") -> None:
|
||||
|
||||
lines = [
|
||||
f"## {title}\n",
|
||||
f"Overall: {icon} {'PASSED' if result.passed else 'FAILED'}\n",
|
||||
]
|
||||
if note:
|
||||
lines.append(f"{note}\n")
|
||||
else:
|
||||
lines.append(f"Overall: {icon} {'PASSED' if result.passed else 'FAILED'}\n")
|
||||
|
||||
reportable = {n: ok for n, ok in result.tests.items() if "test_install" not in n}
|
||||
if reportable:
|
||||
lines += ["| Test | Result |", "| ---- | ------ |"]
|
||||
@@ -464,7 +498,7 @@ def write_summary(result: JobResult, title: str = "QDC Test Results") -> None:
|
||||
passed_n = sum(1 for v in reportable.values() if v)
|
||||
failed_n = sum(1 for v in reportable.values() if not v)
|
||||
lines += ["", f"**{passed_n} passed, {failed_n} failed**"]
|
||||
else:
|
||||
elif not note:
|
||||
lines.append("_No per-test data available._")
|
||||
|
||||
failed_names = [n for n, ok in reportable.items() if not ok]
|
||||
@@ -519,6 +553,8 @@ def parse_args() -> argparse.Namespace:
|
||||
help="Test suite to run (default: bench)")
|
||||
p.add_argument("--job-timeout", type=int, default=JOB_TIMEOUT, metavar="SECONDS",
|
||||
help=f"Max seconds to wait for job completion (default: {JOB_TIMEOUT})")
|
||||
p.add_argument("--runner-timeout", type=int, default=RUNNER_TIMEOUT, metavar="SECONDS",
|
||||
help=f"Max seconds to wait for runner to pick up job (default: {RUNNER_TIMEOUT})")
|
||||
p.add_argument("--retries", type=int, default=DEFAULT_RETRIES, metavar="N",
|
||||
help="Number of retries when device is unavailable (default: 0)")
|
||||
p.add_argument("--retry-delay", type=int, default=RETRY_DELAY, metavar="SECONDS",
|
||||
@@ -537,9 +573,11 @@ def _submit_and_run_job(client, args, spec, target_id, artifact_id) -> JobResult
|
||||
test failure).
|
||||
"""
|
||||
try:
|
||||
wait_for_capacity(client)
|
||||
wait_for_capacity(client, timeout=args.runner_timeout)
|
||||
except TimeoutError:
|
||||
raise DeviceUnavailableError("Capacity wait timed out — device busy")
|
||||
raise RunnerUnavailableError(
|
||||
f"Capacity wait timed out after {args.runner_timeout}s - device busy"
|
||||
)
|
||||
|
||||
job_name = spec.job_name_fmt.format(base="llama.cpp Hexagon tests")
|
||||
|
||||
@@ -559,11 +597,16 @@ def _submit_and_run_job(client, args, spec, target_id, artifact_id) -> JobResult
|
||||
job_parameters=[JobSubmissionParameter.WIFIENABLED],
|
||||
)
|
||||
if job_id is None:
|
||||
raise DeviceUnavailableError("Job submission failed — device may be unavailable")
|
||||
raise DeviceUnavailableError("Job submission failed - device may be unavailable")
|
||||
log.info("Job submitted: %s (device=%s)", job_id, args.device)
|
||||
|
||||
try:
|
||||
job_status = wait_for_job(client, job_id, timeout=args.job_timeout)
|
||||
job_status = wait_for_job(
|
||||
client,
|
||||
job_id,
|
||||
timeout=args.job_timeout,
|
||||
runner_timeout=args.runner_timeout,
|
||||
)
|
||||
except TimeoutError as e:
|
||||
raise DeviceUnavailableError(str(e))
|
||||
log.info("Job %s finished: %s", job_id, job_status)
|
||||
@@ -637,6 +680,14 @@ def main() -> int:
|
||||
try:
|
||||
result = _submit_and_run_job(client, args, spec, target_id, artifact_id)
|
||||
break
|
||||
except RunnerUnavailableError as e:
|
||||
log.warning("QDC runner unavailable for %s after %ds - job skipped: %s", args.device, args.runner_timeout, e)
|
||||
write_summary(
|
||||
JobResult(passed=True, tests={}),
|
||||
title=f"QDC Tests ({args.device})",
|
||||
note=f"SKIPPED: Runner did not pick up job within {args.runner_timeout}s. Skipped to avoid stalling CI.",
|
||||
)
|
||||
return 0
|
||||
except DeviceUnavailableError as e:
|
||||
if attempt < max_attempts:
|
||||
log.warning(
|
||||
|
||||
+68
-12
@@ -2975,6 +2975,7 @@ struct test_rms_norm_mul_rope : public test_case {
|
||||
int mode;
|
||||
const bool yarn;
|
||||
const bool truncate;
|
||||
const ggml_type type; // type of a, the weights are F32
|
||||
|
||||
std::string op_desc(ggml_tensor * t) override {
|
||||
GGML_UNUSED(t);
|
||||
@@ -2984,18 +2985,22 @@ struct test_rms_norm_mul_rope : public test_case {
|
||||
bool run_whole_graph() override { return true; }
|
||||
|
||||
std::string vars() override {
|
||||
return VARS_TO_STR11(ne, eps, multi_add, mul, rope, set_rows, broadcast, mode, set_rows_type, yarn, truncate);
|
||||
std::string s = VARS_TO_STR11(ne, eps, multi_add, mul, rope, set_rows, broadcast, mode, set_rows_type, yarn, truncate);
|
||||
if (type != GGML_TYPE_F32) {
|
||||
s += "," + VAR_TO_STR(type);
|
||||
}
|
||||
return s;
|
||||
}
|
||||
|
||||
test_rms_norm_mul_rope(std::array<int64_t, 4> ne, float eps = 1e-6f, bool multi_add = false,
|
||||
bool set_rows = false, bool broadcast = false, int mode = GGML_ROPE_TYPE_NORMAL,
|
||||
bool mul = true, bool rope = true, ggml_type set_rows_type = GGML_TYPE_F16,
|
||||
bool yarn = false, bool truncate = true)
|
||||
bool yarn = false, bool truncate = true, ggml_type type = GGML_TYPE_F32)
|
||||
: ne(ne), eps(eps), multi_add(multi_add), mul(mul), rope(rope), set_rows(set_rows), broadcast(broadcast),
|
||||
set_rows_type(set_rows_type), mode(mode), yarn(yarn), truncate(truncate) {}
|
||||
set_rows_type(set_rows_type), mode(mode), yarn(yarn), truncate(truncate), type(type) {}
|
||||
|
||||
ggml_tensor * build_graph(ggml_context * ctx) override {
|
||||
ggml_tensor * a = ggml_new_tensor_4d(ctx, GGML_TYPE_F32, ne[0], ne[1], ne[2], ne[3]);
|
||||
ggml_tensor * a = ggml_new_tensor_4d(ctx, type, ne[0], ne[1], ne[2], ne[3]);
|
||||
|
||||
ggml_tensor * b = nullptr;
|
||||
ggml_tensor * c = nullptr;
|
||||
@@ -3074,6 +3079,10 @@ struct test_rms_norm_mul_rope : public test_case {
|
||||
}
|
||||
|
||||
double max_nmse_err() override {
|
||||
if (type == GGML_TYPE_BF16) {
|
||||
// the fused kernel rounds to BF16 once, the CPU rounds after each op
|
||||
return 1e-4;
|
||||
}
|
||||
if (ne[0] == 8192) {
|
||||
return 5e-6;
|
||||
}
|
||||
@@ -4036,6 +4045,11 @@ struct test_rms_norm : public test_case {
|
||||
}
|
||||
}
|
||||
|
||||
double max_nmse_err() override {
|
||||
// the row scale can differ by a few F32 ulps between backends, so a few outputs round to the neighbouring BF16 value
|
||||
return type == GGML_TYPE_BF16 ? 1e-5 : test_case::max_nmse_err();
|
||||
}
|
||||
|
||||
float grad_eps() override {
|
||||
return 1.0f;
|
||||
}
|
||||
@@ -4090,6 +4104,8 @@ struct test_rms_norm_mul_add : public test_case {
|
||||
const bool post_mul;
|
||||
const bool alias_rms_input;
|
||||
const bool weight_broadcast;
|
||||
const ggml_type type_w; // type of b and d, GGML_TYPE_COUNT -> same as a
|
||||
const ggml_type type_add; // type of c, GGML_TYPE_COUNT -> same as b
|
||||
|
||||
std::string op_desc(ggml_tensor * t) override {
|
||||
GGML_UNUSED(t);
|
||||
@@ -4099,22 +4115,33 @@ struct test_rms_norm_mul_add : public test_case {
|
||||
bool run_whole_graph() override { return true; }
|
||||
|
||||
std::string vars() override {
|
||||
return VARS_TO_STR8(type, ne, eps, broadcast, multi_add, post_mul, alias_rms_input, weight_broadcast);
|
||||
std::string s = VARS_TO_STR8(type, ne, eps, broadcast, multi_add, post_mul, alias_rms_input, weight_broadcast);
|
||||
if (type_w != GGML_TYPE_COUNT) {
|
||||
s += "," + VAR_TO_STR(type_w);
|
||||
}
|
||||
if (type_add != GGML_TYPE_COUNT) {
|
||||
s += "," + VAR_TO_STR(type_add);
|
||||
}
|
||||
return s;
|
||||
}
|
||||
|
||||
test_rms_norm_mul_add(ggml_type type = GGML_TYPE_F32,
|
||||
std::array<int64_t, 4> ne = {64, 5, 4, 3},
|
||||
float eps = 1e-6f, bool broadcast = false, bool multi_add = false, bool post_mul = false,
|
||||
bool alias_rms_input = false, bool weight_broadcast = false)
|
||||
bool alias_rms_input = false, bool weight_broadcast = false, ggml_type type_w = GGML_TYPE_COUNT,
|
||||
ggml_type type_add = GGML_TYPE_COUNT)
|
||||
: type(type), ne(ne), eps(eps), broadcast(broadcast), multi_add(multi_add), post_mul(post_mul),
|
||||
alias_rms_input(alias_rms_input), weight_broadcast(weight_broadcast) {}
|
||||
alias_rms_input(alias_rms_input), weight_broadcast(weight_broadcast), type_w(type_w), type_add(type_add) {}
|
||||
|
||||
ggml_tensor * build_graph(ggml_context * ctx) override {
|
||||
std::array<int64_t, 4> broadcast_dims = {ne[0]*2, ne[1]*3, ne[2]*3, ne[3]*4};
|
||||
|
||||
const ggml_type tw = type_w == GGML_TYPE_COUNT ? type : type_w;
|
||||
const ggml_type ta = type_add == GGML_TYPE_COUNT ? tw : type_add;
|
||||
|
||||
ggml_tensor * a = ggml_new_tensor(ctx, type, 4, broadcast ? broadcast_dims.data() : ne.data());
|
||||
ggml_tensor * b = weight_broadcast ? ggml_new_tensor_1d(ctx, type, ne[0]) : ggml_new_tensor(ctx, type, 4, ne.data());
|
||||
ggml_tensor * c = ggml_new_tensor(ctx, type, 4, ne.data());
|
||||
ggml_tensor * b = weight_broadcast ? ggml_new_tensor_1d(ctx, tw, ne[0]) : ggml_new_tensor(ctx, tw, 4, ne.data());
|
||||
ggml_tensor * c = ggml_new_tensor(ctx, ta, 4, ne.data());
|
||||
ggml_tensor * d = nullptr;
|
||||
|
||||
ggml_set_param(a);
|
||||
@@ -4127,7 +4154,7 @@ struct test_rms_norm_mul_add : public test_case {
|
||||
// Use a, b and c early, so we don't end up with an OP_NONE between rms_norm and mul
|
||||
a = ggml_add(ctx, ggml_add(ctx, a, b), c);
|
||||
if (post_mul) {
|
||||
d = ggml_new_tensor_1d(ctx, type, 1);
|
||||
d = ggml_new_tensor_1d(ctx, tw, 1);
|
||||
ggml_set_param(d);
|
||||
ggml_set_name(d, "d");
|
||||
a = ggml_add(ctx, a, d);
|
||||
@@ -4151,6 +4178,11 @@ struct test_rms_norm_mul_add : public test_case {
|
||||
}
|
||||
}
|
||||
|
||||
double max_nmse_err() override {
|
||||
// fused ops round to BF16 once, the CPU rounds after each op
|
||||
return type == GGML_TYPE_BF16 ? 1e-4 : test_case::max_nmse_err();
|
||||
}
|
||||
|
||||
float grad_eps() override {
|
||||
return 1.0f;
|
||||
}
|
||||
@@ -11508,6 +11540,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
|
||||
for (bool v : { false, true }) {
|
||||
test_cases.emplace_back(new test_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, v, eps));
|
||||
test_cases.emplace_back(new test_rms_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, v, eps));
|
||||
test_cases.emplace_back(new test_rms_norm(GGML_TYPE_BF16, { n, 5, 4, 3 }, v, eps));
|
||||
}
|
||||
test_cases.emplace_back(new test_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, false, eps, true));
|
||||
test_cases.emplace_back(new test_norm_scale(GGML_TYPE_F32, { n, 5, 4, 3 }, eps, false, 1.5f));
|
||||
@@ -11526,12 +11559,14 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
|
||||
for (bool v : { false, true }) {
|
||||
test_cases.emplace_back(new test_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, v, eps));
|
||||
test_cases.emplace_back(new test_rms_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, v, eps));
|
||||
test_cases.emplace_back(new test_rms_norm(GGML_TYPE_BF16, { n, 5, 4, 3 }, v, eps));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// in-place tests
|
||||
test_cases.emplace_back(new test_rms_norm(GGML_TYPE_F32, {64, 5, 4, 3}, false, 1e-6f, true));
|
||||
test_cases.emplace_back(new test_rms_norm(GGML_TYPE_BF16, {64, 5, 4, 3}, false, 1e-6f, true));
|
||||
// shapes below exceed the CUDA gridDim.y/gridDim.z limit of 65535 (#27901)
|
||||
test_cases.emplace_back(new test_norm (GGML_TYPE_F32, {4, 1, 65536, 1}, false, 1e-6f));
|
||||
test_cases.emplace_back(new test_rms_norm (GGML_TYPE_F32, {4, 1, 65536, 1}, false, 1e-6f, false));
|
||||
@@ -11545,6 +11580,19 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_F32, {n, 3, 2, 1}, 1e-6f, false, false, false, false, true));
|
||||
}
|
||||
|
||||
// BF16 x with F32 weights, n > 1 because the NMSE of a single BF16 output is too noisy
|
||||
// the widths cover every CUDA block size, with and without 4-value loads
|
||||
for (uint32_t n : {127, 128, 512, 1024, 2048, 2816, 4096, 8192, 16385}) {
|
||||
test_cases.emplace_back(new test_rms_norm(GGML_TYPE_BF16, {n, 3, 2, 1}));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_rope({n, 3, 2, 1}, 1e-6f, false, false, true, GGML_ROPE_TYPE_NORMAL, true, false, GGML_TYPE_F16, false, true, GGML_TYPE_BF16));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_BF16, {n, 3, 2, 1}, 1e-6f, false, false, false, false, true, GGML_TYPE_F32));
|
||||
// BF16 residual added after the norm (post-norm models, e.g. Gemma)
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_BF16, {n, 3, 2, 1}, 1e-6f, false, false, false, false, true, GGML_TYPE_F32, GGML_TYPE_BF16));
|
||||
for (bool multi_add : {false, true}) {
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_BF16, {n, 1, 1, 1}, 1e-6f, false, multi_add, false, false, false, GGML_TYPE_F32));
|
||||
}
|
||||
}
|
||||
|
||||
for (ggml_type set_rows_type : { GGML_TYPE_F32, GGML_TYPE_F16 }) {
|
||||
test_cases.emplace_back(new test_rms_norm_mul_rope({ 256, 1, 1, 1 }, 1e-6f, false, true, false, GGML_ROPE_TYPE_NORMAL, false, false, set_rows_type));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_rope({ 128, 4, 3, 1 }, 1e-6f, false, true, false, GGML_ROPE_TYPE_NORMAL, false, false, set_rows_type));
|
||||
@@ -11559,6 +11607,9 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
|
||||
for (uint32_t n : { 64, 1025 }) {
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_F32, { n, 5, 4, 3 }, eps, false));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_F32, { n, 5, 4, 3 }, eps, true));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_BF16, { n, 5, 4, 3 }, eps, false, false, false, false, false, GGML_TYPE_F32));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_BF16, { n, 5, 4, 3 }, eps, true, false, false, false, false, GGML_TYPE_F32));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_BF16, { n, 5, 4, 3 }, eps, true, false, false, false, false, GGML_TYPE_F32, GGML_TYPE_BF16));
|
||||
test_cases.emplace_back(new test_norm_mul_add(GGML_TYPE_F32, { n, 5, 4, 3 }, eps, false));
|
||||
test_cases.emplace_back(new test_norm_mul_add(GGML_TYPE_F32, { n, 5, 4, 3 }, eps, true));
|
||||
test_cases.emplace_back(new test_add_rms_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, eps, false));
|
||||
@@ -11586,6 +11637,9 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_F32, { 256, 4, 3, 2 }, 1e-6f, false, false, true, false, true));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_F32, { 1536, 1, 1, 1 }, 1e-6f, false, false, false, true));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_F32, { 256, 4, 1, 1 }, 1e-6f, false, false, false, true));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_BF16, { 1536, 1, 1, 1 }, 1e-6f, false, false, true, false, false, GGML_TYPE_F32));
|
||||
// the ADD operand is the BF16 input itself, added in place
|
||||
test_cases.emplace_back(new test_rms_norm_mul_add(GGML_TYPE_BF16, { 1536, 1, 1, 1 }, 1e-6f, false, false, false, true, false, GGML_TYPE_F32));
|
||||
|
||||
test_cases.emplace_back(new test_rms_norm_mul_rope({128, 4, 7, 2}));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_rope({128, 4, 7, 2}, 1e-6f, false, true));
|
||||
@@ -13492,8 +13546,10 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
|
||||
|
||||
// RMS_NORM and fused variants at decode and prefill shapes
|
||||
for (int64_t n_tokens : {1, 512}) {
|
||||
test_cases.emplace_back(new test_rms_norm(GGML_TYPE_F32, {4096, n_tokens, 1, 1}));
|
||||
test_cases.emplace_back(new test_rms_norm(GGML_TYPE_F32, {128, 32*n_tokens, 1, 1}));
|
||||
for (ggml_type type : {GGML_TYPE_F32, GGML_TYPE_BF16}) {
|
||||
test_cases.emplace_back(new test_rms_norm(type, {4096, n_tokens, 1, 1}));
|
||||
test_cases.emplace_back(new test_rms_norm(type, {128, 32*n_tokens, 1, 1}));
|
||||
}
|
||||
test_cases.emplace_back(new test_rms_norm_mul_rope({4096, n_tokens, 1, 1}, 1e-6f, false, false, true, GGML_ROPE_TYPE_NORMAL, true, false));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_rope({128, 32, n_tokens, 1}, 1e-6f, false, false, true, GGML_ROPE_TYPE_NEOX));
|
||||
test_cases.emplace_back(new test_rms_norm_mul_rope({128, 8, n_tokens, 1}, 1e-6f, false, true, true, GGML_ROPE_TYPE_NEOX));
|
||||
|
||||
@@ -1,6 +1,8 @@
|
||||
-r ../../requirements/requirements-convert_legacy_llama.txt
|
||||
--extra-index-url https://download.pytorch.org/whl/cpu
|
||||
pillow~=11.3.0
|
||||
pillow~=12.3.0
|
||||
|
||||
torch==2.11.0 # check_requirements: ignore "=="
|
||||
torchvision==0.26.0; platform_machine != "s390x" # check_requirements: ignore "=="
|
||||
torch==2.11.0; python_version < "3.15" # check_requirements: ignore "=="
|
||||
torch==2.15.0; python_version >= "3.15" # check_requirements: ignore "=="
|
||||
torchvision==0.26.0; platform_machine != "s390x" and python_version < "3.15" # check_requirements: ignore "=="
|
||||
torchvision==0.29.2; platform_machine != "s390x" and python_version >= "3.15" # check_requirements: ignore "=="
|
||||
|
||||
@@ -2,7 +2,7 @@ pytest~=8.3.3
|
||||
pytest-xdist~=3.6
|
||||
filelock~=3.16
|
||||
numpy~=2.2.6; python_version < "3.14"
|
||||
numpy~=2.4.6; python_version >= "3.14"
|
||||
numpy~=2.5.4; python_version >= "3.14"
|
||||
openai~=2.14.0
|
||||
prometheus-client~=0.20.0
|
||||
requests~=2.32.3
|
||||
|
||||
@@ -1,6 +1,8 @@
|
||||
engine-strict=true
|
||||
ignore-scripts=true
|
||||
min-release-age=7
|
||||
# Allow root-declared remote tarball dependencies (llama-brand) on npm 12+
|
||||
allow-remote=root
|
||||
# The single-bundle production build needs more than the default heap on CI runners;
|
||||
# npm passes this to every script it runs, on every platform
|
||||
node-options=--max-old-space-size=6144
|
||||
|
||||
Reference in New Issue
Block a user