Commit 29dbbad2b for llama.cpp
commit 29dbbad2b1e498d395d3e05fcc9885440f78cedd
Author: Georgi Gerganov <ggerganov@gmail.com>
Date: Thu Nov 12 20:58:24 2026 +0200
ggml : prune (#29801)
* ggml : prune deprecated API
* cont : bump RPC version
* cont : minor
* ggml : remove outdated documentation from ggml.h
The 172-line comment block at the top of the header has not been updated in a
substantive way since the backend v2 split (Nov 2023) and now documents APIs
that no longer exist:
- `ggml_set_f32()`, `ggml_get_f32_1d()`, `ggml_set_f32_1d()` were removed
- `ggml_conv_1d_1s()`, `ggml_conv_1d_2s()` ops no longer exist
- `ggml_set_param()` no longer takes a context argument
- `ggml_graph_compute_with_ctx()` moved to the CPU backend header
- the optimization API now lives in `ggml-opt.h`
The block also ended with 5 empty `TODO` sections.
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* ggml : remove DIAG_MASK_INF and DIAG_MASK_ZERO (#31946)
* rpc : bump RPC protocol major version
The tensor op enum is sent over the wire, so removing GGML_OP_ADD1,
GGML_OP_DIAG_MASK_INF and GGML_OP_DIAG_MASK_ZERO shifts the value of every
op that follows them and peers built before and after this change would
misinterpret each other's ops.
The handshake compares the major version for exact equality and ignores the
patch version, so bumping the major version is what rejects incompatible
peers instead of silently misbehaving.
Assisted-by: dsh:llama.cpp/ggml-org/DeepSeek-V4.1-Flash-1017-GGUF (Mac Studio, M5 Ultra, 256GB)
---------
Co-authored-by: losewayy <1134285612@qq.com>
diff --git a/docs/ops.md b/docs/ops.md
index c8204465b..3c2bcd239 100644
--- a/docs/ops.md
+++ b/docs/ops.md
@@ -40,7 +40,6 @@ Legend:
| CROSS_ENTROPY_LOSS_BACK | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
| CUMSUM | ❌ | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
| DIAG | ❌ | ❌ | ✅ | ✅ | 🟡 | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
-| DIAG_MASK_INF | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | ❌ | 🟡 | ✅ | ✅ | ❌ | ❌ | ❌ |
| DIV | ❌ | ✅ | ✅ | ✅ | ❌ | 🟡 | 🟡 | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
| DSV4_HC_COMB | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
| DSV4_HC_POST | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
diff --git a/docs/ops/BLAS.csv b/docs/ops/BLAS.csv
index 3af72452a..eabdd6f81 100644
--- a/docs/ops/BLAS.csv
+++ b/docs/ops/BLAS.csv
@@ -8741,9 +8741,6 @@
"BLAS","ROUND","type=f32,ne=[1024,1024,1,1]","support","0","no","BLAS"
"BLAS","TRUNC","type=f32,ne=[7,1,5,3]","support","0","no","BLAS"
"BLAS","TRUNC","type=f32,ne=[1024,1024,1,1]","support","0","no","BLAS"
-"BLAS","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","0","no","BLAS"
-"BLAS","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","0","no","BLAS"
-"BLAS","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","0","no","BLAS"
"BLAS","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","0","no","BLAS"
"BLAS","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","0","no","BLAS"
"BLAS","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","0","no","BLAS"
diff --git a/docs/ops/CANN.csv b/docs/ops/CANN.csv
index 8434461a5..de317ad15 100644
--- a/docs/ops/CANN.csv
+++ b/docs/ops/CANN.csv
@@ -8764,9 +8764,6 @@
"CANN0","ROUND","type=f32,ne=[1024,1024,1,1]","support","0","no","CANN"
"CANN0","TRUNC","type=f32,ne=[7,1,5,3]","support","0","no","CANN"
"CANN0","TRUNC","type=f32,ne=[1024,1024,1,1]","support","0","no","CANN"
-"CANN0","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","1","yes","CANN"
-"CANN0","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","1","yes","CANN"
-"CANN0","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","1","yes","CANN"
"CANN0","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","CANN"
"CANN0","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","CANN"
"CANN0","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","CANN"
diff --git a/docs/ops/CPU.csv b/docs/ops/CPU.csv
index 5d491fc4b..49803f227 100644
--- a/docs/ops/CPU.csv
+++ b/docs/ops/CPU.csv
@@ -12962,9 +12962,6 @@
"CPU","ROUND","type=f32,ne=[1024,1024,1,1]","support","1","yes","CPU"
"CPU","TRUNC","type=f32,ne=[7,1,5,3]","support","1","yes","CPU"
"CPU","TRUNC","type=f32,ne=[1024,1024,1,1]","support","1","yes","CPU"
-"CPU","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","1","yes","CPU"
-"CPU","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","1","yes","CPU"
-"CPU","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","1","yes","CPU"
"CPU","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","CPU"
"CPU","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","CPU"
"CPU","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","CPU"
diff --git a/docs/ops/CUDA.csv b/docs/ops/CUDA.csv
index 7499a4927..2b8c8b7a8 100644
--- a/docs/ops/CUDA.csv
+++ b/docs/ops/CUDA.csv
@@ -8725,9 +8725,6 @@
"CUDA0","ROUND","type=f32,ne=[1024,1024,1,1]","support","1","yes","CUDA"
"CUDA0","TRUNC","type=f32,ne=[7,1,5,3]","support","1","yes","CUDA"
"CUDA0","TRUNC","type=f32,ne=[1024,1024,1,1]","support","1","yes","CUDA"
-"CUDA0","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","1","yes","CUDA"
-"CUDA0","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","1","yes","CUDA"
-"CUDA0","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","1","yes","CUDA"
"CUDA0","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","CUDA"
"CUDA0","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","CUDA"
"CUDA0","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","CUDA"
diff --git a/docs/ops/ET.csv b/docs/ops/ET.csv
index 91774c58e..56a8f9705 100644
--- a/docs/ops/ET.csv
+++ b/docs/ops/ET.csv
@@ -9637,9 +9637,6 @@
"ET","ROUND","type=f32,ne=[1024,1024,1,1]","support","1","yes","ET"
"ET","TRUNC","type=f32,ne=[7,1,5,3]","support","0","no","ET"
"ET","TRUNC","type=f32,ne=[1024,1024,1,1]","support","1","yes","ET"
-"ET","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","0","no","ET"
-"ET","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","0","no","ET"
-"ET","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","0","no","ET"
"ET","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","ET"
"ET","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","ET"
"ET","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","ET"
diff --git a/docs/ops/Hexagon.csv b/docs/ops/Hexagon.csv
index c9608b715..160c41499 100644
--- a/docs/ops/Hexagon.csv
+++ b/docs/ops/Hexagon.csv
@@ -14564,9 +14564,6 @@
"HTP0","ROUND","type=f32,ne=[1024,1024,1,1]","support","0","no","HTP"
"HTP0","TRUNC","type=f32,ne=[7,1,5,3]","support","0","no","HTP"
"HTP0","TRUNC","type=f32,ne=[1024,1024,1,1]","support","0","no","HTP"
-"HTP0","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","0","no","HTP"
-"HTP0","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","0","no","HTP"
-"HTP0","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","0","no","HTP"
"HTP0","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","HTP"
"HTP0","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","HTP"
"HTP0","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","HTP"
diff --git a/docs/ops/Metal.csv b/docs/ops/Metal.csv
index ec34bb66b..8d38300a0 100644
--- a/docs/ops/Metal.csv
+++ b/docs/ops/Metal.csv
@@ -9628,9 +9628,6 @@
"MTL0","ROUND","type=f32,ne=[1024,1024,1,1]","support","1","yes","MTL"
"MTL0","TRUNC","type=f32,ne=[7,1,5,3]","support","1","yes","MTL"
"MTL0","TRUNC","type=f32,ne=[1024,1024,1,1]","support","1","yes","MTL"
-"MTL0","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","0","no","MTL"
-"MTL0","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","0","no","MTL"
-"MTL0","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","0","no","MTL"
"MTL0","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","MTL"
"MTL0","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","MTL"
"MTL0","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","MTL"
diff --git a/docs/ops/OpenCL.csv b/docs/ops/OpenCL.csv
index 4a6e90456..8b595e4b9 100644
--- a/docs/ops/OpenCL.csv
+++ b/docs/ops/OpenCL.csv
@@ -14887,9 +14887,6 @@
"OpenCL","ROUND","type=f32,ne=[1024,1024,1,1]","support","1","yes","OpenCL"
"OpenCL","TRUNC","type=f32,ne=[7,1,5,3]","support","1","yes","OpenCL"
"OpenCL","TRUNC","type=f32,ne=[1024,1024,1,1]","support","1","yes","OpenCL"
-"OpenCL","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","1","yes","OpenCL"
-"OpenCL","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","1","yes","OpenCL"
-"OpenCL","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","0","no","OpenCL"
"OpenCL","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","OpenCL"
"OpenCL","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","OpenCL"
"OpenCL","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","OpenCL"
diff --git a/docs/ops/SYCL.csv b/docs/ops/SYCL.csv
index 3ef1ca1d7..e01a536a7 100644
--- a/docs/ops/SYCL.csv
+++ b/docs/ops/SYCL.csv
@@ -14534,9 +14534,6 @@
"SYCL0","ROUND","type=f32,ne=[1024,1024,1,1]","support","1","yes","SYCL"
"SYCL0","TRUNC","type=f32,ne=[7,1,5,3]","support","1","yes","SYCL"
"SYCL0","TRUNC","type=f32,ne=[1024,1024,1,1]","support","1","yes","SYCL"
-"SYCL0","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","1","yes","SYCL"
-"SYCL0","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","1","yes","SYCL"
-"SYCL0","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","1","yes","SYCL"
"SYCL0","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","SYCL"
"SYCL0","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","SYCL"
"SYCL0","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","SYCL"
diff --git a/docs/ops/Vulkan.csv b/docs/ops/Vulkan.csv
index 13e67eb79..65ac8bf33 100644
--- a/docs/ops/Vulkan.csv
+++ b/docs/ops/Vulkan.csv
@@ -12313,9 +12313,6 @@
"Vulkan0","ROUND","type=f32,ne=[1024,1024,1,1]","support","1","yes","Vulkan"
"Vulkan0","TRUNC","type=f32,ne=[7,1,5,3]","support","1","yes","Vulkan"
"Vulkan0","TRUNC","type=f32,ne=[1024,1024,1,1]","support","1","yes","Vulkan"
-"Vulkan0","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","1","yes","Vulkan"
-"Vulkan0","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","1","yes","Vulkan"
-"Vulkan0","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","1","yes","Vulkan"
"Vulkan0","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","Vulkan"
"Vulkan0","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","Vulkan"
"Vulkan0","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","Vulkan"
diff --git a/docs/ops/WebGPU.csv b/docs/ops/WebGPU.csv
index 1da1a0f04..3b874fd68 100644
--- a/docs/ops/WebGPU.csv
+++ b/docs/ops/WebGPU.csv
@@ -10560,9 +10560,6 @@
"WebGPU: WebGPU","ROUND","type=f32,ne=[1024,1024,1,1]","support","1","yes","WebGPU"
"WebGPU: WebGPU","TRUNC","type=f32,ne=[7,1,5,3]","support","1","yes","WebGPU"
"WebGPU: WebGPU","TRUNC","type=f32,ne=[1024,1024,1,1]","support","1","yes","WebGPU"
-"WebGPU: WebGPU","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","0","no","WebGPU"
-"WebGPU: WebGPU","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","0","no","WebGPU"
-"WebGPU: WebGPU","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","0","no","WebGPU"
"WebGPU: WebGPU","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","WebGPU"
"WebGPU: WebGPU","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","WebGPU"
"WebGPU: WebGPU","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","1","yes","WebGPU"
diff --git a/docs/ops/ZenDNN.csv b/docs/ops/ZenDNN.csv
index 19d579fb8..b309f1a0d 100644
--- a/docs/ops/ZenDNN.csv
+++ b/docs/ops/ZenDNN.csv
@@ -9319,9 +9319,6 @@
"ZenDNN","ROUND","type=f32,ne=[1024,1024,1,1]","support","0","no","ZenDNN"
"ZenDNN","TRUNC","type=f32,ne=[7,1,5,3]","support","0","no","ZenDNN"
"ZenDNN","TRUNC","type=f32,ne=[1024,1024,1,1]","support","0","no","ZenDNN"
-"ZenDNN","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","0","no","ZenDNN"
-"ZenDNN","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","0","no","ZenDNN"
-"ZenDNN","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","0","no","ZenDNN"
"ZenDNN","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","0","no","ZenDNN"
"ZenDNN","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","0","no","ZenDNN"
"ZenDNN","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","0","no","ZenDNN"
diff --git a/docs/ops/zDNN.csv b/docs/ops/zDNN.csv
index 005a1cb4d..a44ef3ebb 100644
--- a/docs/ops/zDNN.csv
+++ b/docs/ops/zDNN.csv
@@ -14550,9 +14550,6 @@
"zDNN","ROUND","type=f32,ne=[1024,1024,1,1]","support","0","no","zDNN"
"zDNN","TRUNC","type=f32,ne=[7,1,5,3]","support","0","no","zDNN"
"zDNN","TRUNC","type=f32,ne=[1024,1024,1,1]","support","0","no","zDNN"
-"zDNN","DIAG_MASK_INF","type=f32,ne=[10,10,1,1],n_past=5","support","0","no","zDNN"
-"zDNN","DIAG_MASK_INF","type=f32,ne=[10,10,3,1],n_past=5","support","0","no","zDNN"
-"zDNN","DIAG_MASK_INF","type=f32,ne=[10,10,3,2],n_past=5","support","0","no","zDNN"
"zDNN","SOFT_MAX","type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","0","no","zDNN"
"zDNN","SOFT_MAX","type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","0","no","zDNN"
"zDNN","SOFT_MAX","type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0","support","0","no","zDNN"
diff --git a/ggml/include/ggml-rpc.h b/ggml/include/ggml-rpc.h
index f32ddeee2..da1d4e1f1 100644
--- a/ggml/include/ggml-rpc.h
+++ b/ggml/include/ggml-rpc.h
@@ -6,26 +6,31 @@
extern "C" {
#endif
-#define RPC_PROTO_MAJOR_VERSION 9
+#define RPC_PROTO_MAJOR_VERSION 10
#define RPC_PROTO_MINOR_VERSION 0
#define RPC_PROTO_PATCH_VERSION 0
#ifdef __cplusplus
-static_assert(GGML_OP_COUNT == 102, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION");
+static_assert(GGML_OP_COUNT == 99, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION");
#endif
#define GGML_RPC_MAX_SERVERS 16
// backend API
GGML_BACKEND_API ggml_backend_t ggml_backend_rpc_init(const char * endpoint, uint32_t device);
+
GGML_BACKEND_API bool ggml_backend_is_rpc(ggml_backend_t backend);
GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_rpc_buffer_type(const char * endpoint, uint32_t device);
GGML_BACKEND_API void ggml_backend_rpc_get_device_memory(const char * endpoint, uint32_t device, size_t * free, size_t * total);
-GGML_BACKEND_API void ggml_backend_rpc_start_server(const char * endpoint, const char * cache_dir,
- size_t n_threads, size_t n_devices, ggml_backend_dev_t * devices);
+GGML_BACKEND_API void ggml_backend_rpc_start_server(
+ const char * endpoint,
+ const char * cache_dir,
+ size_t n_threads,
+ size_t n_devices,
+ ggml_backend_dev_t * devices);
GGML_BACKEND_API ggml_backend_reg_t ggml_backend_rpc_reg(void);
GGML_BACKEND_API ggml_backend_reg_t ggml_backend_rpc_add_server(const char * endpoint);
diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h
index 973aea25b..2341fa6c1 100644
--- a/ggml/include/ggml.h
+++ b/ggml/include/ggml.h
@@ -1,178 +1,5 @@
#pragma once
-//
-// GGML Tensor Library
-//
-// This documentation is still a work in progress.
-// If you wish some specific topics to be covered, feel free to drop a comment:
-//
-// https://github.com/ggml-org/whisper.cpp/issues/40
-//
-// ## Overview
-//
-// This library implements:
-//
-// - a set of tensor operations
-// - automatic differentiation
-// - basic optimization algorithms
-//
-// The aim of this library is to provide a minimalistic approach for various machine learning tasks. This includes,
-// but is not limited to, the following:
-//
-// - linear regression
-// - support vector machines
-// - neural networks
-//
-// The library allows the user to define a certain function using the available tensor operations. This function
-// definition is represented internally via a computation graph. Each tensor operation in the function definition
-// corresponds to a node in the graph. Having the computation graph defined, the user can choose to compute the
-// function's value and/or its gradient with respect to the input variables. Optionally, the function can be optimized
-// using one of the available optimization algorithms.
-//
-// For example, here we define the function: f(x) = a*x^2 + b
-//
-// {
-// struct ggml_init_params params = {
-// .mem_size = 16*1024*1024,
-// .mem_buffer = NULL,
-// };
-//
-// // memory allocation happens here
-// struct ggml_context * ctx = ggml_init(params);
-//
-// struct ggml_tensor * x = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 1);
-//
-// ggml_set_param(ctx, x); // x is an input variable
-//
-// struct ggml_tensor * a = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 1);
-// struct ggml_tensor * b = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 1);
-// struct ggml_tensor * x2 = ggml_mul(ctx, x, x);
-// struct ggml_tensor * f = ggml_add(ctx, ggml_mul(ctx, a, x2), b);
-//
-// ...
-// }
-//
-// Notice that the function definition above does not involve any actual computation. The computation is performed only
-// when the user explicitly requests it. For example, to compute the function's value at x = 2.0:
-//
-// {
-// ...
-//
-// struct ggml_cgraph * gf = ggml_new_graph(ctx);
-// ggml_build_forward_expand(gf, f);
-//
-// // set the input variable and parameter values
-// ggml_set_f32(x, 2.0f);
-// ggml_set_f32(a, 3.0f);
-// ggml_set_f32(b, 4.0f);
-//
-// ggml_graph_compute_with_ctx(ctx, &gf, n_threads);
-//
-// printf("f = %f\n", ggml_get_f32_1d(f, 0));
-//
-// ...
-// }
-//
-// The actual computation is performed in the ggml_graph_compute() function.
-//
-// The ggml_new_tensor_...() functions create new tensors. They are allocated in the memory buffer provided to the
-// ggml_init() function. You have to be careful not to exceed the memory buffer size. Therefore, you have to know
-// in advance how much memory you need for your computation. Alternatively, you can allocate a large enough memory
-// and after defining the computation graph, call the ggml_used_mem() function to find out how much memory was
-// actually needed.
-//
-// The ggml_set_param() function marks a tensor as an input variable. This is used by the automatic
-// differentiation and optimization algorithms.
-//
-// The described approach allows to define the function graph once and then compute its forward or backward graphs
-// multiple times. All computations will use the same memory buffer allocated in the ggml_init() function. This way
-// the user can avoid the memory allocation overhead at runtime.
-//
-// The library supports multi-dimensional tensors - up to 4 dimensions. The FP16 and FP32 data types are first class
-// citizens, but in theory the library can be extended to support FP8 and integer data types.
-//
-// Each tensor operation produces a new tensor. Initially the library was envisioned to support only the use of unary
-// and binary operations. Most of the available operations fall into one of these two categories. With time, it became
-// clear that the library needs to support more complex operations. The way to support these operations is not clear
-// yet, but a few examples are demonstrated in the following operations:
-//
-// - ggml_permute()
-// - ggml_conv_1d_1s()
-// - ggml_conv_1d_2s()
-//
-// For each tensor operator, the library implements a forward and backward computation function. The forward function
-// computes the output tensor value given the input tensor values. The backward function computes the adjoint of the
-// input tensors given the adjoint of the output tensor. For a detailed explanation of what this means, take a
-// calculus class, or watch the following video:
-//
-// What is Automatic Differentiation?
-// https://www.youtube.com/watch?v=wG_nF1awSSY
-//
-//
-// ## Tensor data (struct ggml_tensor)
-//
-// The tensors are stored in memory via the ggml_tensor struct. The structure provides information about the size of
-// the tensor, the data type, and the memory buffer where the tensor data is stored. Additionally, it contains
-// pointers to the "source" tensors - i.e. the tensors that were used to compute the current tensor. For example:
-//
-// {
-// struct ggml_tensor * c = ggml_add(ctx, a, b);
-//
-// assert(c->src[0] == a);
-// assert(c->src[1] == b);
-// }
-//
-// The multi-dimensional tensors are stored in row-major order. The ggml_tensor struct contains fields for the
-// number of elements in each dimension ("ne") as well as the number of bytes ("nb", a.k.a. stride). This allows
-// to store tensors that are not contiguous in memory, which is useful for operations such as transposition and
-// permutation. All tensor operations have to take the stride into account and not assume that the tensor is
-// contiguous in memory.
-//
-// The data of the tensor is accessed via the "data" pointer. For example:
-//
-// {
-// const int nx = 2;
-// const int ny = 3;
-//
-// struct ggml_tensor * a = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, nx, ny);
-//
-// for (int y = 0; y < ny; y++) {
-// for (int x = 0; x < nx; x++) {
-// *(float *) ((char *) a->data + y*a->nb[1] + x*a->nb[0]) = x + y;
-// }
-// }
-//
-// ...
-// }
-//
-// Alternatively, there are helper functions, such as ggml_get_f32_1d() and ggml_set_f32_1d() that can be used.
-//
-// ## The matrix multiplication operator (ggml_mul_mat)
-//
-// TODO
-//
-//
-// ## Multi-threading
-//
-// TODO
-//
-//
-// ## Overview of ggml.c
-//
-// TODO
-//
-//
-// ## SIMD optimizations
-//
-// TODO
-//
-//
-// ## Debugging ggml
-//
-// TODO
-//
-//
-
#ifdef GGML_SHARED
# if defined(_WIN32) && !defined(__MINGW32__)
# ifdef GGML_BUILD
@@ -445,7 +272,6 @@ extern "C" {
// the precision parameters are stored as ggml_tensor.op_params to the respective ops
enum ggml_prec {
GGML_PREC_UNDEFINED = 0,
- GGML_PREC_DEFAULT = 0, // note: deprecated, use GGML_PREC_UNDEFINED
GGML_PREC_F32 = 10,
GGML_PREC_BF16 = 15,
GGML_PREC_F16 = 20,
@@ -499,7 +325,6 @@ extern "C" {
GGML_OP_DUP,
GGML_OP_ADD,
GGML_OP_ADD_ID,
- GGML_OP_ADD1,
GGML_OP_ACC,
GGML_OP_SUB,
GGML_OP_MUL,
@@ -541,8 +366,6 @@ extern "C" {
GGML_OP_GET_ROWS_BACK,
GGML_OP_SET_ROWS,
GGML_OP_DIAG,
- GGML_OP_DIAG_MASK_INF,
- GGML_OP_DIAG_MASK_ZERO,
GGML_OP_SOFT_MAX,
GGML_OP_SOFT_MAX_BACK,
GGML_OP_ROPE,
@@ -765,10 +588,6 @@ extern "C" {
GGML_API size_t ggml_type_size(enum ggml_type type); // size in bytes for all elements in a block
GGML_API size_t ggml_row_size (enum ggml_type type, int64_t ne); // size in bytes for all elements in a row
- GGML_DEPRECATED(
- GGML_API double ggml_type_sizef(enum ggml_type type), // ggml_type_size()/ggml_blck_size() as float
- "use ggml_row_size() instead");
-
GGML_API const char * ggml_type_name(enum ggml_type type);
GGML_API const char * ggml_op_name (enum ggml_op op);
GGML_API const char * ggml_op_symbol(enum ggml_op op);
@@ -937,18 +756,6 @@ extern "C" {
struct ggml_tensor * b,
struct ggml_tensor * ids);
- GGML_DEPRECATED(GGML_API struct ggml_tensor * ggml_add1(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- struct ggml_tensor * b),
- "use ggml_add instead");
-
- GGML_DEPRECATED(GGML_API struct ggml_tensor * ggml_add1_inplace(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- struct ggml_tensor * b),
- "use ggml_add_inplace instead");
-
// dst = a
// view(dst, nb1, nb2, nb3, offset) += b
// return dst
@@ -1500,13 +1307,6 @@ extern "C" {
struct ggml_tensor * a,
struct ggml_tensor * b);
- // change the precision of a matrix multiplication
- // set to GGML_PREC_F32 for higher precision (useful for phi-2)
- GGML_DEPRECATED(GGML_API void ggml_mul_mat_set_prec(
- struct ggml_tensor * a,
- enum ggml_prec prec),
- "use ggml_prec_set_acc() instead");
-
// change the hint of a matrix multiplication
GGML_API void ggml_mul_mat_set_hint(
struct ggml_tensor * a,
@@ -1778,30 +1578,6 @@ extern "C" {
struct ggml_context * ctx,
struct ggml_tensor * a);
- // set elements above the diagonal to -INF
- GGML_API struct ggml_tensor * ggml_diag_mask_inf(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- int n_past);
-
- // in-place, returns view(a)
- GGML_API struct ggml_tensor * ggml_diag_mask_inf_inplace(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- int n_past);
-
- // set elements above the diagonal to 0
- GGML_API struct ggml_tensor * ggml_diag_mask_zero(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- int n_past);
-
- // in-place, returns view(a)
- GGML_API struct ggml_tensor * ggml_diag_mask_zero_inplace(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- int n_past);
-
GGML_API struct ggml_tensor * ggml_clamp(
struct ggml_context * ctx,
struct ggml_tensor * a,
@@ -2005,36 +1781,6 @@ extern "C" {
float beta_fast,
float beta_slow);
- GGML_DEPRECATED(GGML_API struct ggml_tensor * ggml_rope_custom(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- struct ggml_tensor * b,
- int n_dims,
- int mode,
- int n_ctx_orig,
- float freq_base,
- float freq_scale,
- float ext_factor,
- float attn_factor,
- float beta_fast,
- float beta_slow),
- "use ggml_rope_ext instead");
-
- GGML_DEPRECATED(GGML_API struct ggml_tensor * ggml_rope_custom_inplace(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- struct ggml_tensor * b,
- int n_dims,
- int mode,
- int n_ctx_orig,
- float freq_base,
- float freq_scale,
- float ext_factor,
- float attn_factor,
- float beta_fast,
- float beta_slow),
- "use ggml_rope_ext_inplace instead");
-
// rotary position embedding backward, i.e compute dx from dy
// a - dy
GGML_API struct ggml_tensor * ggml_rope_ext_back(
@@ -2340,7 +2086,7 @@ extern "C" {
GGML_SCALE_MODE_BILINEAR = 1,
GGML_SCALE_MODE_BICUBIC = 2,
- GGML_SCALE_MODE_COUNT
+ GGML_SCALE_MODE_COUNT = 3
};
enum ggml_scale_flag {
@@ -2356,18 +2102,6 @@ extern "C" {
int scale_factor,
enum ggml_scale_mode mode);
- // interpolate
- // interpolate scale to specified dimensions
- GGML_DEPRECATED(GGML_API struct ggml_tensor * ggml_upscale_ext(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- int ne0,
- int ne1,
- int ne2,
- int ne3,
- enum ggml_scale_mode mode),
- "use ggml_interpolate instead");
-
// Up- or downsamples the input to the specified size.
// 2D scale modes (eg. bilinear) are applied to the first two dimensions.
GGML_API struct ggml_tensor * ggml_interpolate(
@@ -2536,11 +2270,6 @@ extern "C" {
float max_bias,
float logit_softcap);
- GGML_DEPRECATED(GGML_API void ggml_flash_attn_ext_set_prec(
- struct ggml_tensor * a,
- enum ggml_prec prec),
- "use ggml_prec_set_acc() instead");
-
GGML_API enum ggml_prec ggml_flash_attn_ext_get_prec(
const struct ggml_tensor * a);
diff --git a/ggml/src/ggml-alloc.c b/ggml/src/ggml-alloc.c
index b0c067b49..65d2bea27 100644
--- a/ggml/src/ggml-alloc.c
+++ b/ggml/src/ggml-alloc.c
@@ -23,11 +23,8 @@ bool ggml_op_can_inplace(enum ggml_op op) {
switch (op) {
case GGML_OP_FILL:
case GGML_OP_SCALE:
- case GGML_OP_DIAG_MASK_ZERO:
- case GGML_OP_DIAG_MASK_INF:
case GGML_OP_ADD:
case GGML_OP_ADD_ID:
- case GGML_OP_ADD1:
case GGML_OP_SUB:
case GGML_OP_MUL:
case GGML_OP_DIV:
diff --git a/ggml/src/ggml-backend-meta.cpp b/ggml/src/ggml-backend-meta.cpp
index 991398316..595a5824a 100644
--- a/ggml/src/ggml-backend-meta.cpp
+++ b/ggml/src/ggml-backend-meta.cpp
@@ -939,7 +939,6 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
case GGML_OP_ADD_ID: {
split_state = handle_bin_bcast(src_ss);
} break;
- case GGML_OP_ADD1:
case GGML_OP_ACC: {
split_state = handle_generic(src_ss, /*scalar_only =*/ true);
} break;
@@ -1020,9 +1019,7 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
case GGML_OP_SET_ROWS: {
split_state = handle_set_rows(src_ss);
} break;
- case GGML_OP_DIAG:
- case GGML_OP_DIAG_MASK_INF:
- case GGML_OP_DIAG_MASK_ZERO: {
+ case GGML_OP_DIAG: {
split_state = handle_generic(src_ss, /*scalar_only =*/ true);
} break;
case GGML_OP_SOFT_MAX:
diff --git a/ggml/src/ggml-cann/aclnn_ops.cpp b/ggml/src/ggml-cann/aclnn_ops.cpp
index 66e011ef0..e602e4bf3 100644
--- a/ggml/src/ggml-cann/aclnn_ops.cpp
+++ b/ggml/src/ggml-cann/aclnn_ops.cpp
@@ -1347,31 +1347,6 @@ void ggml_cann_rms_norm(ggml_backend_cann_context & ctx, ggml_tensor * dst) {
GGML_CANN_CALL_ACLNN_OP(ctx, RmsNorm, acl_src.get(), acl_gamma.get(), eps, acl_dst.get(), acl_rstd.get());
}
-// TODO: performace is low.
-void ggml_cann_diag_mask(ggml_backend_cann_context & ctx, ggml_tensor * dst, float value) {
- ggml_tensor * src = dst->src[0];
-
- acl_tensor_ptr acl_src = ggml_cann_create_tensor(src);
- acl_tensor_ptr acl_dst = ggml_cann_create_tensor(dst);
-
- const int n_past = ((int32_t *) dst->op_params)[0];
-
- ggml_cann_pool_alloc one_tensor_allocator(ctx.pool(), ggml_nbytes(src));
- void * buffer = one_tensor_allocator.get();
-
- acl_tensor_ptr mask_tensor = ggml_cann_create_tensor(buffer, ggml_cann_type_mapping(src->type),
- ggml_type_size(src->type), src->ne, src->nb, GGML_MAX_DIMS);
-
- aclnn_fill_scalar(ctx, value, mask_tensor.get());
-
- float alphaValue = 1.0f;
- acl_scalar_ptr alpha = ggml_cann_create_scalar(&alphaValue, aclDataType::ACL_FLOAT);
-
- GGML_CANN_CALL_ACLNN_OP(ctx, InplaceTriu, mask_tensor.get(), n_past + 1);
- GGML_CANN_CALL_ACLNN_OP(ctx, Tril, acl_src.get(), n_past + 1, acl_dst.get());
- GGML_CANN_CALL_ACLNN_OP(ctx, InplaceAdd, acl_dst.get(), mask_tensor.get(), alpha.get());
-}
-
/**
* @brief Permutes the dimensions of a tensor according to a specified order.
*
diff --git a/ggml/src/ggml-cann/aclnn_ops.h b/ggml/src/ggml-cann/aclnn_ops.h
index 678f4d654..f470042cf 100644
--- a/ggml/src/ggml-cann/aclnn_ops.h
+++ b/ggml/src/ggml-cann/aclnn_ops.h
@@ -466,21 +466,6 @@ void ggml_cann_dup(ggml_backend_cann_context & ctx, ggml_tensor * dst);
*/
void ggml_cann_rms_norm(ggml_backend_cann_context & ctx, ggml_tensor * dst);
-/**
- * @brief Applies a diagonal mask to the tensor with a specified value.
- *
- * @details This function creates a mask tensor filled with ones, then applies
- * an upper triangular and lower triangular operation to it based on
- * the number of past elements specified. Afterward, it adds the masked
- * tensor to the destination tensor in-place.
- *
- * @param ctx The backend CANN context used for operations.
- * @param dst The destination tensor where the result will be stored. dst->op is
- * `GGML_OP_DIAG_MASK`
- * @param value The value to use for masking.
- */
-void ggml_cann_diag_mask(ggml_backend_cann_context & ctx, ggml_tensor * dst, float value);
-
/**
* @brief Performs an image-to-column transformation on the input tensor.
*
diff --git a/ggml/src/ggml-cann/ggml-cann.cpp b/ggml/src/ggml-cann/ggml-cann.cpp
index b49aee73d..0bfb96009 100644
--- a/ggml/src/ggml-cann/ggml-cann.cpp
+++ b/ggml/src/ggml-cann/ggml-cann.cpp
@@ -1793,7 +1793,6 @@ static bool ggml_cann_compute_forward(ggml_backend_cann_context & ctx, struct gg
ggml_cann_dup(ctx, dst);
break;
case GGML_OP_ADD:
- case GGML_OP_ADD1:
ggml_cann_binary_op<aclnn_add>(ctx, dst);
break;
case GGML_OP_SUB:
@@ -1958,9 +1957,6 @@ static bool ggml_cann_compute_forward(ggml_backend_cann_context & ctx, struct gg
case GGML_OP_PERMUTE:
case GGML_OP_TRANSPOSE:
break;
- case GGML_OP_DIAG_MASK_INF:
- ggml_cann_diag_mask(ctx, dst, -INFINITY);
- break;
case GGML_OP_SOFT_MAX:
ggml_cann_softmax(ctx, dst);
break;
@@ -2618,7 +2614,6 @@ static bool ggml_backend_cann_supports_op(ggml_backend_dev_t dev, const ggml_ten
case GGML_OP_TRANSPOSE:
case GGML_OP_NORM:
case GGML_OP_ADD:
- case GGML_OP_ADD1:
case GGML_OP_SUB:
case GGML_OP_MUL:
case GGML_OP_DIV:
@@ -2626,7 +2621,6 @@ static bool ggml_backend_cann_supports_op(ggml_backend_dev_t dev, const ggml_ten
case GGML_OP_SQR:
case GGML_OP_SQRT:
case GGML_OP_CLAMP:
- case GGML_OP_DIAG_MASK_INF:
case GGML_OP_SUM_ROWS:
case GGML_OP_ARGSORT:
case GGML_OP_ACC:
diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c
index 40e93ccab..a76af8487 100644
--- a/ggml/src/ggml-cpu/ggml-cpu.c
+++ b/ggml/src/ggml-cpu/ggml-cpu.c
@@ -1782,10 +1782,6 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm
{
ggml_compute_forward_add_id(params, tensor);
} break;
- case GGML_OP_ADD1:
- {
- ggml_compute_forward_add1(params, tensor);
- } break;
case GGML_OP_ACC:
{
ggml_compute_forward_acc(params, tensor);
@@ -1926,14 +1922,6 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm
{
ggml_compute_forward_diag(params, tensor);
} break;
- case GGML_OP_DIAG_MASK_INF:
- {
- ggml_compute_forward_diag_mask_inf(params, tensor);
- } break;
- case GGML_OP_DIAG_MASK_ZERO:
- {
- ggml_compute_forward_diag_mask_zero(params, tensor);
- } break;
case GGML_OP_SOFT_MAX:
{
ggml_compute_forward_soft_max(params, tensor);
@@ -2285,7 +2273,6 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) {
case GGML_OP_CONT:
case GGML_OP_ADD:
case GGML_OP_ADD_ID:
- case GGML_OP_ADD1:
case GGML_OP_ACC:
case GGML_OP_CUMSUM:
case GGML_OP_TRI:
@@ -2407,8 +2394,6 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) {
{
n_tasks = 1;
} break;
- case GGML_OP_DIAG_MASK_ZERO:
- case GGML_OP_DIAG_MASK_INF:
case GGML_OP_SOFT_MAX_BACK:
case GGML_OP_ROPE:
case GGML_OP_ROPE_BACK:
@@ -2889,7 +2874,6 @@ struct ggml_cplan ggml_graph_plan(
} break;
case GGML_OP_ADD:
case GGML_OP_ADD_ID:
- case GGML_OP_ADD1:
{
if (ggml_is_quantized(node->src[0]->type) || node->src[0]->type == GGML_TYPE_F8_E4M3) {
cur = ggml_type_size(GGML_TYPE_F32) * node->src[0]->ne[0] * n_tasks;
diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
index c987a087c..eed5d69e3 100644
--- a/ggml/src/ggml-cpu/ops.cpp
+++ b/ggml/src/ggml-cpu/ops.cpp
@@ -5704,94 +5704,6 @@ void ggml_compute_forward_diag(
}
}
-// ggml_compute_forward_diag_mask_inf
-
-static void ggml_compute_forward_diag_mask_f32(
- const ggml_compute_params * params,
- ggml_tensor * dst,
- const float value) {
-
- const ggml_tensor * src0 = dst->src[0];
-
- const int ith = params->ith;
- const int nth = params->nth;
-
- const int n_past = ((int32_t *) dst->op_params)[0];
- const bool inplace = src0->data == dst->data;
-
- GGML_ASSERT(n_past >= 0);
-
- if (!inplace) {
- if (ith == 0) {
- // memcpy needs to be synchronized across threads to avoid race conditions.
- // => do it in INIT phase
- GGML_ASSERT(ggml_nelements(dst) == ggml_nelements(src0));
- GGML_ASSERT(ggml_is_contiguous(dst) && ggml_is_contiguous(src0));
- memcpy(
- ((char *) dst->data),
- ((char *) src0->data),
- ggml_nbytes(dst));
- }
- ggml_barrier(params->threadpool);
- }
-
- // TODO: handle transposed/permuted matrices
-
- const int n = ggml_nrows(src0);
- const int nc = src0->ne[0];
- const int nr = src0->ne[1];
- const int nz = n/nr;
-
- GGML_ASSERT( dst->nb[0] == sizeof(float));
- GGML_ASSERT(src0->nb[0] == sizeof(float));
-
- for (int k = 0; k < nz; k++) {
- for (int j = ith; j < nr; j += nth) {
- for (int i = n_past; i < nc; i++) {
- if (i > n_past + j) {
- *(float *)((char *) dst->data + k*dst->nb[2] + j*dst->nb[1] + i*dst->nb[0]) = value;
- }
- }
- }
- }
-}
-
-void ggml_compute_forward_diag_mask_inf(
- const ggml_compute_params * params,
- ggml_tensor * dst) {
-
- const ggml_tensor * src0 = dst->src[0];
-
- switch (src0->type) {
- case GGML_TYPE_F32:
- {
- ggml_compute_forward_diag_mask_f32(params, dst, -INFINITY);
- } break;
- default:
- {
- GGML_ABORT("fatal error");
- }
- }
-}
-
-void ggml_compute_forward_diag_mask_zero(
- const ggml_compute_params * params,
- ggml_tensor * dst) {
-
- const ggml_tensor * src0 = dst->src[0];
-
- switch (src0->type) {
- case GGML_TYPE_F32:
- {
- ggml_compute_forward_diag_mask_f32(params, dst, 0);
- } break;
- default:
- {
- GGML_ABORT("fatal error");
- }
- }
-}
-
// ggml_compute_forward_soft_max
static void ggml_compute_forward_soft_max_f32(
diff --git a/ggml/src/ggml-cpu/ops.h b/ggml/src/ggml-cpu/ops.h
index 00a2ea7dc..aac527a64 100644
--- a/ggml/src/ggml-cpu/ops.h
+++ b/ggml/src/ggml-cpu/ops.h
@@ -56,8 +56,6 @@ void ggml_compute_forward_get_rows(const struct ggml_compute_params * params, st
void ggml_compute_forward_get_rows_back(const struct ggml_compute_params * params, struct ggml_tensor * dst);
void ggml_compute_forward_set_rows(const struct ggml_compute_params * params, struct ggml_tensor * dst);
void ggml_compute_forward_diag(const struct ggml_compute_params * params, struct ggml_tensor * dst);
-void ggml_compute_forward_diag_mask_inf(const struct ggml_compute_params * params, struct ggml_tensor * dst);
-void ggml_compute_forward_diag_mask_zero(const struct ggml_compute_params * params, struct ggml_tensor * dst);
void ggml_compute_forward_soft_max(const struct ggml_compute_params * params, struct ggml_tensor * dst);
void ggml_compute_forward_soft_max_ext_back(const struct ggml_compute_params * params, struct ggml_tensor * dst);
void ggml_compute_forward_rope(const struct ggml_compute_params * params, struct ggml_tensor * dst);
diff --git a/ggml/src/ggml-cuda/diagmask.cu b/ggml/src/ggml-cuda/diagmask.cu
deleted file mode 100644
index 4b713ba22..000000000
--- a/ggml/src/ggml-cuda/diagmask.cu
+++ /dev/null
@@ -1,40 +0,0 @@
-#include "diagmask.cuh"
-
-static __global__ void diag_mask_inf_f32(const float * x, float * dst, const int ncols, const int rows_per_channel, const int n_past) {
- const int col = blockDim.y*blockIdx.y + threadIdx.y;
- const int row = blockDim.x*blockIdx.x + threadIdx.x;
-
- if (col >= ncols) {
- return;
- }
-
- const int i = row*ncols + col;
- //dst[i] = col > (n_past + row % rows_per_channel) ? -INFINITY : x[i];
- //dst[i] = x[i] - (col > n_past + row % rows_per_channel) * INT_MAX; // equivalent within rounding error but slightly faster on GPU
- dst[i] = x[i] - (col > n_past + row % rows_per_channel) * FLT_MAX;
-}
-
-static void diag_mask_inf_f32_cuda(const float * x, float * dst, const int ncols_x, const int nrows_x, const int rows_per_channel, const int n_past, cudaStream_t stream) {
- const dim3 block_dims(1, CUDA_DIAG_MASK_INF_BLOCK_SIZE, 1);
- const int block_num_x = (ncols_x + CUDA_DIAG_MASK_INF_BLOCK_SIZE - 1) / CUDA_DIAG_MASK_INF_BLOCK_SIZE;
- const dim3 block_nums(nrows_x, block_num_x, 1);
- diag_mask_inf_f32<<<block_nums, block_dims, 0, stream>>>(x, dst, ncols_x, rows_per_channel, n_past);
-}
-
-void ggml_cuda_op_diag_mask_inf(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
- const ggml_tensor * src0 = dst->src[0];
- const float * src0_d = (const float *)src0->data;
- 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);
-
- const int64_t ne00 = src0->ne[0];
- const int64_t ne01 = src0->ne[1];
- const int nrows0 = ggml_nrows(src0);
-
- const int n_past = ((int32_t *) dst->op_params)[0];
-
- diag_mask_inf_f32_cuda(src0_d, dst_d, ne00, nrows0, ne01, n_past, stream);
-}
diff --git a/ggml/src/ggml-cuda/diagmask.cuh b/ggml/src/ggml-cuda/diagmask.cuh
deleted file mode 100644
index 6cdbef17e..000000000
--- a/ggml/src/ggml-cuda/diagmask.cuh
+++ /dev/null
@@ -1,5 +0,0 @@
-#include "common.cuh"
-
-#define CUDA_DIAG_MASK_INF_BLOCK_SIZE 32
-
-void ggml_cuda_op_diag_mask_inf(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index dd384d91f..fed14ddfc 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -25,7 +25,6 @@
#include "ggml-cuda/cpy.cuh"
#include "ggml-cuda/cross-entropy-loss.cuh"
#include "ggml-cuda/cumsum.cuh"
-#include "ggml-cuda/diagmask.cuh"
#include "ggml-cuda/diag.cuh"
#include "ggml-cuda/fattn.cuh"
#include "ggml-cuda/fp8.cuh"
@@ -2742,7 +2741,6 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg
ggml_cuda_dup(ctx, dst);
break;
case GGML_OP_ADD:
- case GGML_OP_ADD1: // TODO: more efficient implementation
ggml_cuda_op_add(ctx, dst);
break;
case GGML_OP_ADD_ID:
@@ -2937,9 +2935,6 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg
case GGML_OP_DIAG:
ggml_cuda_op_diag(ctx, dst);
break;
- case GGML_OP_DIAG_MASK_INF:
- ggml_cuda_op_diag_mask_inf(ctx, dst);
- break;
case GGML_OP_SOFT_MAX:
ggml_cuda_op_soft_max(ctx, dst);
break;
@@ -6540,7 +6535,6 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_OP_PERMUTE:
case GGML_OP_TRANSPOSE:
case GGML_OP_ADD_ID:
- case GGML_OP_ADD1:
case GGML_OP_SQR:
case GGML_OP_SQRT:
case GGML_OP_SIN:
@@ -6587,8 +6581,6 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
}
case GGML_OP_CONT:
return true;
- case GGML_OP_DIAG_MASK_INF:
- return true;
case GGML_OP_SOFT_MAX:
return true;
case GGML_OP_SOFT_MAX_BACK: {
diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt
index 0570fb835..cdf4a5d54 100644
--- a/ggml/src/ggml-opencl/CMakeLists.txt
+++ b/ggml/src/ggml-opencl/CMakeLists.txt
@@ -76,7 +76,6 @@ set(GGML_OPENCL_KERNELS
clamp
cpy
cvt
- diag_mask_inf
diag
div
dsv4_hc
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 91675b6d4..7fa1a8e35 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -798,7 +798,6 @@ struct ggml_backend_opencl_context {
cl_program program_moe_add_id_glu;
cl_program program_clamp;
cl_program program_cvt;
- cl_program program_diag_mask_inf;
cl_program program_gelu;
cl_program program_gemv_noshuffle_general;
cl_program program_gemv_noshuffle;
@@ -886,7 +885,6 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_rms_norm_mul_add = nullptr; // fused rms_norm(x)*w + b (residual)
cl_kernel kernel_l2_norm_f32;
cl_kernel kernel_group_norm, kernel_group_norm_mul_add;
- cl_kernel kernel_diag_mask_inf, kernel_diag_mask_inf_8;
cl_kernel kernel_diag_f32;
cl_kernel kernel_soft_max, kernel_soft_max_4;
cl_kernel kernel_soft_max_f16, kernel_soft_max_4_f16;
@@ -2291,23 +2289,6 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
GGML_LOG_CONT(".");
}
- // diag_mask_inf
- {
-#ifdef GGML_OPENCL_EMBED_KERNELS
- const std::string kernel_src {
- #include "diag_mask_inf.cl.h"
- };
-#else
- const std::string kernel_src = read_file("diag_mask_inf.cl");
-#endif
- backend_ctx->program_diag_mask_inf =
- build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);
-
- CL_CHECK((backend_ctx->kernel_diag_mask_inf_8 = clCreateKernel(backend_ctx->program_diag_mask_inf, "kernel_diag_mask_inf_8", &err), err));
- CL_CHECK((backend_ctx->kernel_diag_mask_inf = clCreateKernel(backend_ctx->program_diag_mask_inf, "kernel_diag_mask_inf", &err), err));
- GGML_LOG_CONT(".");
- }
-
// diag
{
#ifdef GGML_OPENCL_EMBED_KERNELS
@@ -12214,8 +12195,6 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
return true;
case GGML_OP_DIAG:
return true;
- case GGML_OP_DIAG_MASK_INF:
- return op->ne[3] == 1;
case GGML_OP_ROPE: {
const int mode = ((const int32_t *) op->op_params)[2];
if (!ggml_rope_get_truncate(op)) {
@@ -37948,68 +37927,6 @@ static void ggml_cl_set(ggml_backend_t backend, const ggml_tensor * src0, const
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
}
-static void ggml_cl_diag_mask_inf(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
- GGML_ASSERT(src0);
- GGML_ASSERT(src0->extra);
- GGML_ASSERT(dst);
- GGML_ASSERT(dst->extra);
-
- UNUSED(src1);
-
- int n_past = ((int32_t *)(dst->op_params))[0];
-
- const int ne00 = src0 ? src0->ne[0] : 0;
- const int ne01 = src0 ? src0->ne[1] : 0;
- const int ne02 = src0 ? src0->ne[2] : 0;
-
- ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
-
- ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *)src0->extra;
- ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra;
-
- cl_ulong offset0 = extra0->offset + src0->view_offs;
- cl_ulong offsetd = extrad->offset + dst->view_offs;
-
- cl_kernel kernel;
-
- if (ne00%8 == 0) {
- kernel = backend_ctx->kernel_diag_mask_inf_8;
-
- CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->data_device));
- CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0));
- CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extrad->data_device));
- CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetd));
- CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &ne00));
- CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), &ne01));
- CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &n_past));
-
- size_t global_work_size[] = {(size_t)ne00*ne01*ne02/8, 1, 1};
- size_t local_work_size[] = {64, 1, 1};
-
- backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
- } else {
- kernel = backend_ctx->kernel_diag_mask_inf;
-
- CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->data_device));
- CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0));
- CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extrad->data_device));
- CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetd));
- CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &ne00));
- CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), &ne01));
- CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &n_past));
-
- size_t global_work_size[] = {(size_t)ne00, (size_t)ne01, (size_t)ne02};
- size_t local_work_size[] = {64, 1, 1};
-
- size_t * local_work_size_ptr = local_work_size;
- if (ne00 % 64 != 0 && !backend_ctx->non_uniform_workgroups) {
- local_work_size_ptr = nullptr; // Let driver choose the work-group sizes.
- }
-
- backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size_ptr, dst);
- }
-}
-
static void ggml_cl_diag(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
GGML_ASSERT(src0);
GGML_ASSERT(src0->extra);
@@ -39935,12 +39852,6 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor
}
func = ggml_cl_diag;
break;
- case GGML_OP_DIAG_MASK_INF:
- if (!any_on_device) {
- return false;
- }
- func = ggml_cl_diag_mask_inf;
- break;
case GGML_OP_SOFT_MAX:
if (!any_on_device) {
return false;
diff --git a/ggml/src/ggml-opencl/kernels/diag_mask_inf.cl b/ggml/src/ggml-opencl/kernels/diag_mask_inf.cl
deleted file mode 100644
index 36eff0439..000000000
--- a/ggml/src/ggml-opencl/kernels/diag_mask_inf.cl
+++ /dev/null
@@ -1,58 +0,0 @@
-#pragma OPENCL EXTENSION cl_khr_fp16 : enable
-
-//------------------------------------------------------------------------------
-// diag_mask_inf kernels
-//------------------------------------------------------------------------------
-kernel void kernel_diag_mask_inf(
- global float * src0,
- ulong offset0,
- global float * dst,
- ulong offsetd,
- int ne00,
- int ne01,
- int n_past
-) {
- src0 = (global float*)((global char*)src0 + offset0);
- dst = (global float*)((global char*)dst + offsetd);
-
- int i02 = get_global_id(2);
- int i01 = get_global_id(1);
- int i00 = get_global_id(0);
-
- if (i00 > n_past + i01) {
- dst[i02*ne01*ne00 + i01*ne00 + i00] = -INFINITY;
- } else {
- dst[i02*ne01*ne00 + i01*ne00 + i00] = src0[i02*ne01*ne00 + i01*ne00 + i00];
- }
-}
-
-kernel void kernel_diag_mask_inf_8(
- global float4 * src0,
- ulong offset0,
- global float4 * dst,
- ulong offsetd,
- int ne00,
- int ne01,
- int n_past
-) {
- src0 = (global float4*)((global char*)src0 + offset0);
- dst = (global float4*)((global char*)dst + offsetd);
-
- int i = 2*get_global_id(0);
-
- dst[i+0] = src0[i+0];
- dst[i+1] = src0[i+1];
- int i4 = 4*i;
- int i02 = i4/(ne00*ne01); i4 -= i02*ne00*ne01;
- int i01 = i4/(ne00); i4 -= i01*ne00;
- int i00 = i4;
- for (int k = 3; k >= 0; --k) {
- if (i00 + 4 + k <= n_past + i01) {
- break;
- }
- (&dst[i+1])[k] = -INFINITY;
- if (i00 + k > n_past + i01) {
- (&dst[i])[k] = -INFINITY;
- }
- }
-}
diff --git a/ggml/src/ggml-openvino/openvino/op_table.cpp b/ggml/src/ggml-openvino/openvino/op_table.cpp
index 23b3a2cff..857fbf175 100644
--- a/ggml/src/ggml-openvino/openvino/op_table.cpp
+++ b/ggml/src/ggml-openvino/openvino/op_table.cpp
@@ -33,7 +33,6 @@ std::unordered_map<std::string, CreatorFunction> get_supported_ops() {
using namespace ov::op;
return {
{"GGML_OP_ADD", op::translate_add },
- {"GGML_OP_ADD1", op::translate_1to1_match_2_inputs<v1::Add> },
{"GGML_OP_ADD_ID", op::translate_add_id },
{"GGML_OP_CONCAT", op::translate_concat },
{"GGML_OP_CONT", op::translate_cont },
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index cb9fa4742..fdd526425 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -2517,23 +2517,6 @@ k_argsort_f32_i32(const float *x, int *dst, const int ncols, int ncols_pad,
}
}
-static void diag_mask_inf_f32(const float * x, float * dst, const int ncols, const int rows_per_channel, const int n_past,
- const sycl::nd_item<3> &item_ct1) {
- const int col = item_ct1.get_local_range(1) * item_ct1.get_group(1) +
- item_ct1.get_local_id(1);
- const int row = item_ct1.get_local_range(2) * item_ct1.get_group(2) +
- item_ct1.get_local_id(2);
-
- if (col >= ncols) {
- return;
- }
-
- const int i = row*ncols + col;
- //dst[i] = col > (n_past + row % rows_per_channel) ? -INFINITY : x[i];
- //dst[i] = x[i] - (col > n_past + row % rows_per_channel) * INT_MAX; // equivalent within rounding error but slightly faster on GPU
- dst[i] = x[i] - (col > n_past + row % rows_per_channel) * FLT_MAX;
-}
-
static void scale_f32(const float * x, float * dst, const float scale, const float bias, const int k,
const sycl::nd_item<3> &item_ct1) {
const int i = item_ct1.get_local_range(2) * item_ct1.get_group(2) +
@@ -3139,20 +3122,6 @@ static void argmax_f32_i32_sycl(const float *x, int *dst, const int ncols,
});
});
}
-static void diag_mask_inf_f32_sycl(const float *x, float *dst,
- const int ncols_x, const int nrows_x,
- const int rows_per_channel, const int n_past,
- queue_ptr stream) {
- const sycl::range<3> block_dims(1, SYCL_DIAG_MASK_INF_BLOCK_SIZE, 1);
- const int block_num_x = (ncols_x + SYCL_DIAG_MASK_INF_BLOCK_SIZE - 1) / SYCL_DIAG_MASK_INF_BLOCK_SIZE;
- const sycl::range<3> block_nums(1, block_num_x, nrows_x);
- stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](sycl::nd_item<3> item_ct1) {
- diag_mask_inf_f32(x, dst, ncols_x,
- rows_per_channel, n_past,
- item_ct1);
- });
-}
static dpct::err0 ggml_sycl_cpy_tensor_2d(void *dst,
const struct ggml_tensor *src,
@@ -3536,25 +3505,6 @@ inline void ggml_sycl_op_argmax(ggml_backend_sycl_context & ctx, ggml_tensor * d
argmax_f32_i32_sycl(src0_dd, dst_dd, ncols, nrows, main_stream);
}
-inline void ggml_sycl_op_diag_mask_inf(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
- GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32);
- GGML_ASSERT( dst->type == GGML_TYPE_F32);
- dpct::queue_ptr main_stream = ctx.stream();
- SYCL_CHECK(ggml_sycl_set_device(ctx.device));
- const float * src0_dd = static_cast<const float *>(dst->src[0]->data);
- float * dst_dd = static_cast<float *>(dst->data);
-
- const int64_t ne00 = dst->src[0]->ne[0];
- const int64_t ne01 = dst->src[0]->ne[1];
- const int nrows0 = ggml_nrows(dst->src[0]);
- // the kernel indexes in 32-bit
- GGML_ASSERT(ne00 * nrows0 <= INT32_MAX);
-
- const int n_past = ((int32_t *) dst->op_params)[0];
-
- diag_mask_inf_f32_sycl(src0_dd, dst_dd, ne00, nrows0, ne01, n_past, main_stream);
-}
-
static void tri_f32_sycl(
const float * src,
float * dst,
@@ -6045,11 +5995,6 @@ static void ggml_sycl_scale(ggml_backend_sycl_context & ctx, ggml_tensor * dst)
ggml_sycl_op_scale(ctx, dst);
}
-static void ggml_sycl_diag_mask_inf(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
- scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
- ggml_sycl_op_diag_mask_inf(ctx, dst);
-}
-
static void ggml_sycl_pool2d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
ggml_sycl_op_pool2d(ctx, dst);
@@ -6193,7 +6138,6 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg
ggml_sycl_dup(ctx, dst);
break;
case GGML_OP_ADD:
- case GGML_OP_ADD1: // TODO: more efficient implementation
ggml_sycl_add(ctx, dst);
break;
case GGML_OP_ADD_ID:
@@ -6396,9 +6340,6 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg
case GGML_OP_TRI:
ggml_sycl_op_tri(ctx, dst);
break;
- case GGML_OP_DIAG_MASK_INF:
- ggml_sycl_diag_mask_inf(ctx, dst);
- break;
case GGML_OP_SOFT_MAX:
ggml_sycl_op_soft_max(ctx, dst);
break;
@@ -7639,7 +7580,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
case GGML_OP_PERMUTE:
case GGML_OP_TRANSPOSE:
case GGML_OP_ADD:
- case GGML_OP_ADD1:
case GGML_OP_ADD_ID:
case GGML_OP_SUB:
case GGML_OP_COUNT_EQUAL:
@@ -7674,8 +7614,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
op->type == GGML_TYPE_F32 &&
ggml_is_contiguous(src0);
}
- case GGML_OP_DIAG_MASK_INF:
- return true;
case GGML_OP_SOFT_MAX:
return true;
case GGML_OP_SOFT_MAX_BACK: {
diff --git a/ggml/src/ggml-sycl/presets.hpp b/ggml/src/ggml-sycl/presets.hpp
index eeb803ca1..e9fa6c8a6 100644
--- a/ggml/src/ggml-sycl/presets.hpp
+++ b/ggml/src/ggml-sycl/presets.hpp
@@ -35,7 +35,6 @@
#define SYCL_SCALE_BLOCK_SIZE 256
#define SYCL_ROPE_BLOCK_SIZE 256
#define SYCL_ALIBI_BLOCK_SIZE 32
-#define SYCL_DIAG_MASK_INF_BLOCK_SIZE 32
#define SYCL_QUANTIZE_BLOCK_SIZE 256
#define SYCL_DEQUANTIZE_BLOCK_SIZE 256
#define SYCL_GET_ROWS_BLOCK_SIZE 256
diff --git a/ggml/src/ggml-vulkan/ggml-vulkan-common.h b/ggml/src/ggml-vulkan/ggml-vulkan-common.h
index 2c39a525f..4e05fb818 100644
--- a/ggml/src/ggml-vulkan/ggml-vulkan-common.h
+++ b/ggml/src/ggml-vulkan/ggml-vulkan-common.h
@@ -172,7 +172,6 @@ void ggml_vk_l2_norm(ggml_backend_vk_context * ctx, vk_context& subctx, const gg
void ggml_vk_unary(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
void ggml_vk_xielu(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
void ggml_vk_glu(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
-void ggml_vk_diag_mask_inf(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst);
void ggml_vk_soft_max(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, ggml_tensor * dst);
void ggml_vk_soft_max_back(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
void ggml_vk_topk_moe(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_cgraph * cgraph, int node_idx);
diff --git a/ggml/src/ggml-vulkan/ggml-vulkan-debug.cpp b/ggml/src/ggml-vulkan/ggml-vulkan-debug.cpp
index 792093ccf..2c9c09c06 100644
--- a/ggml/src/ggml-vulkan/ggml-vulkan-debug.cpp
+++ b/ggml/src/ggml-vulkan/ggml-vulkan-debug.cpp
@@ -1035,8 +1035,6 @@ void ggml_vk_check_results_0(ggml_backend_vk_context * ctx, ggml_cgraph * cgraph
} else if (tensor->op == GGML_OP_SCALE) {
const float * params = (const float *)tensor->op_params;
tensor_clone = ggml_scale_bias(ggml_ctx, src_clone[0], params[0], params[1]);
- } else if (tensor->op == GGML_OP_ADD1) {
- tensor_clone = ggml_add1(ggml_ctx, src_clone[0], src_clone[1]);
} else if (tensor->op == GGML_OP_ARANGE) {
const float start = ggml_get_op_params_f32(tensor, 0);
const float stop = ggml_get_op_params_f32(tensor, 1);
@@ -1101,8 +1099,6 @@ void ggml_vk_check_results_0(ggml_backend_vk_context * ctx, ggml_cgraph * cgraph
}
} else if (tensor->op == GGML_OP_SOFT_MAX_BACK) {
tensor_clone = ggml_soft_max_ext_back(ggml_ctx, src_clone[0], src_clone[1], ((float *)tensor->op_params)[0], ((float *)tensor->op_params)[1]);
- } else if (tensor->op == GGML_OP_DIAG_MASK_INF) {
- tensor_clone = ggml_diag_mask_inf(ggml_ctx, src_clone[0], tensor->op_params[0]);
} else if (tensor->op == GGML_OP_ROPE || tensor->op == GGML_OP_ROPE_BACK) {
const int n_dims = ((int32_t *) tensor->op_params)[1];
const int mode = ((int32_t *) tensor->op_params)[2];
diff --git a/ggml/src/ggml-vulkan/ggml-vulkan-push-constants.h b/ggml/src/ggml-vulkan/ggml-vulkan-push-constants.h
index e78f87b1b..1ef00477b 100644
--- a/ggml/src/ggml-vulkan/ggml-vulkan-push-constants.h
+++ b/ggml/src/ggml-vulkan/ggml-vulkan-push-constants.h
@@ -462,12 +462,6 @@ struct vk_op_add_id_push_constants {
uint32_t s21;
};
-struct vk_op_diag_mask_push_constants {
- uint32_t ncols;
- uint32_t rows_per_channel;
- int32_t n_past;
-};
-
struct vk_op_rope_push_constants {
uint32_t rope_mode;
uint32_t nrows;
diff --git a/ggml/src/ggml-vulkan/ggml-vulkan-types.h b/ggml/src/ggml-vulkan/ggml-vulkan-types.h
index 02ccb8ba0..81fd5b8a8 100644
--- a/ggml/src/ggml-vulkan/ggml-vulkan-types.h
+++ b/ggml/src/ggml-vulkan/ggml-vulkan-types.h
@@ -943,7 +943,6 @@ struct vk_device_struct {
vk_pipeline pipeline_leaky_relu[2];
vk_pipeline pipeline_silu_back_f32;
- vk_pipeline pipeline_diag_mask_inf_f32;
vk_pipeline pipeline_soft_max_f32, pipeline_soft_max_f32_f16;
vk_pipeline pipeline_soft_max_f32_wg512, pipeline_soft_max_f32_f16_wg512;
vk_pipeline pipeline_soft_max_back_f32;
diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
index 041497d6a..22641f5fb 100644
--- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp
+++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
@@ -3676,10 +3676,6 @@ void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
CREATE_UNARY_MUL(softplus, 3)
#undef CREATE_UNARY_MUL
- ggml_vk_create_pipeline(device, device->pipeline_add1_f16_f16, "add1_f16_f16", add1_f16_f16_len, add1_f16_f16_data, "main", 3, sizeof(vk_op_binary_push_constants), {512, 1, 1}, {}, 1);
- ggml_vk_create_pipeline(device, device->pipeline_add1_f16_f32, "add1_f16_f32", add1_f16_f32_len, add1_f16_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {512, 1, 1}, {}, 1);
- ggml_vk_create_pipeline(device, device->pipeline_add1_f32_f32, "add1_f32_f32", add1_f32_f32_len, add1_f32_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {512, 1, 1}, {}, 1);
-
ggml_vk_create_pipeline(device, device->pipeline_arange_f32, "arange_f32", arange_f32_len, arange_f32_data, "main", 1, sizeof(vk_op_push_constants), {512, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_fill_f32, "fill_f32", fill_f32_len, fill_f32_data, "main", 1, sizeof(vk_op_push_constants), {512, 1, 1}, {}, 1);
@@ -3700,8 +3696,6 @@ void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
ggml_vk_create_pipeline(device, device->pipeline_silu_back_f32, "silu_back_f32", silu_back_f32_len, silu_back_f32_data, "main", 3, sizeof(vk_op_push_constants), {512, 1, 1}, {}, 1);
- ggml_vk_create_pipeline(device, device->pipeline_diag_mask_inf_f32, "diag_mask_inf_f32", diag_mask_inf_f32_len, diag_mask_inf_f32_data, "main", 2, sizeof(vk_op_diag_mask_push_constants), {1, 512, 1}, {}, 1, true);
-
ggml_vk_create_pipeline(device, device->pipeline_soft_max_f32, "soft_max_f32", soft_max_f32_len, soft_max_f32_data, "main", 4, sizeof(vk_op_soft_max_push_constants), {1, 1, 1}, { device->subgroup_size }, 1);
ggml_vk_create_pipeline(device, device->pipeline_soft_max_f32_wg512, "soft_max_f32_wg512", soft_max_f32_len, soft_max_f32_data, "main", 4, sizeof(vk_op_soft_max_push_constants), {1, 1, 1}, { 512 }, 1);
ggml_vk_create_pipeline(device, device->pipeline_soft_max_f32_f16, "soft_max_f32_f16", soft_max_f32_f16_len, soft_max_f32_f16_data, "main", 4, sizeof(vk_op_soft_max_push_constants), {1, 1, 1}, { device->subgroup_size }, 1);
@@ -9477,11 +9471,6 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const
break;
}
return nullptr;
- case GGML_OP_DIAG_MASK_INF:
- if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
- return ctx->device->pipeline_diag_mask_inf_f32;
- }
- return nullptr;
case GGML_OP_SOFT_MAX:
GGML_ASSERT(!src1 || src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16);
GGML_ASSERT(!src2 || src2->type == GGML_TYPE_F32);
@@ -9860,17 +9849,6 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const
return pipeline;
}
return nullptr;
- case GGML_OP_ADD1:
- if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) {
- return ctx->device->pipeline_add1_f16_f16;
- }
- if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F16) {
- return ctx->device->pipeline_add1_f16_f32;
- }
- if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
- return ctx->device->pipeline_add1_f32_f32;
- }
- return nullptr;
case GGML_OP_ARANGE:
if (dst->type == GGML_TYPE_F32) {
return ctx->device->pipeline_arange_f32;
@@ -10009,9 +9987,6 @@ static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, co
elements = { groups_x, groups_y, uint32_t(CEIL_DIV(uint64_t(total_workgroups), groups_xy)) };
GGML_ASSERT(elements[2] <= ctx->device->properties.limits.maxComputeWorkGroupCount[2]);
} break;
- case GGML_OP_DIAG_MASK_INF:
- elements = { (uint32_t)ggml_nrows(src0), (uint32_t)ne00, 1 };
- break;
case GGML_OP_ROPE:
case GGML_OP_ROPE_BACK:
{
@@ -10145,7 +10120,6 @@ static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, co
case GGML_OP_SUB:
case GGML_OP_DIV:
case GGML_OP_MUL:
- case GGML_OP_ADD1:
case GGML_OP_OUT_PROD:
case GGML_OP_ARANGE:
case GGML_OP_FILL:
@@ -11040,21 +11014,6 @@ void ggml_vk_sqrt(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_
ggml_vk_op_f32(ctx, subctx, src0, nullptr, nullptr, nullptr, dst, GGML_OP_SQRT, vk_op_unary_push_constants_init(src0, dst));
}
-void ggml_vk_add1(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
- const uint32_t src0_type_size = ggml_type_size(src0->type);
- const uint32_t src1_type_size = ggml_type_size(src1->type);
- const uint32_t dst_type_size = ggml_type_size(dst->type);
-
- ggml_vk_op_f32<vk_op_binary_push_constants>(ctx, subctx, src0, src1, nullptr, nullptr, dst, GGML_OP_ADD1, {
- (uint32_t)ggml_nelements(src0),
- (uint32_t)src0->ne[0], (uint32_t)src0->ne[1], (uint32_t)src0->ne[2],(uint32_t)src0->ne[3], (uint32_t)src0->nb[0] / src0_type_size, (uint32_t)src0->nb[1] / src0_type_size, (uint32_t)src0->nb[2] / src0_type_size, (uint32_t)src0->nb[3] / src0_type_size,
- (uint32_t)src1->ne[0], (uint32_t)src1->ne[1], (uint32_t)src1->ne[2],(uint32_t)src1->ne[3], (uint32_t)src1->nb[0] / src1_type_size, (uint32_t)src1->nb[1] / src1_type_size, (uint32_t)src1->nb[2] / src1_type_size, (uint32_t)src1->nb[3] / src1_type_size,
- (uint32_t) dst->ne[0], (uint32_t) dst->ne[1], (uint32_t) dst->ne[2],(uint32_t) dst->ne[3], (uint32_t) dst->nb[0] / dst_type_size, (uint32_t) dst->nb[1] / dst_type_size, (uint32_t) dst->nb[2] / dst_type_size, (uint32_t) dst->nb[3] / dst_type_size,
- 0,
- 0.0f, 0.0f, 0,
- });
-}
-
void ggml_vk_arange(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst) {
VK_LOG_DEBUG("ggml_vk_arange(dst=" << dst << ", ne=" << ggml_nelements(dst) << ")");
@@ -11569,11 +11528,6 @@ void ggml_vk_glu(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_t
});
}
-void ggml_vk_diag_mask_inf(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) {
- int32_t * op_params = (int32_t *)dst->op_params;
- ggml_vk_op_f32<vk_op_diag_mask_push_constants>(ctx, subctx, src0, nullptr, nullptr, nullptr, dst, GGML_OP_DIAG_MASK_INF, { (uint32_t)src0->ne[0], (uint32_t)src0->ne[1], op_params[0] });
-}
-
void ggml_vk_soft_max(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, ggml_tensor * dst) {
float * op_params = (float *)dst->op_params;
@@ -12945,10 +12899,6 @@ bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgraph, in
case GGML_OP_UPSCALE:
ggml_vk_upscale(ctx, compute_ctx, src0, node);
- break;
- case GGML_OP_ADD1:
- ggml_vk_add1(ctx, compute_ctx, src0, src1, node);
-
break;
case GGML_OP_ARANGE:
ggml_vk_arange(ctx, compute_ctx, node);
@@ -13099,10 +13049,6 @@ bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgraph, in
default:
return false;
}
- break;
- case GGML_OP_DIAG_MASK_INF:
- ggml_vk_diag_mask_inf(ctx, compute_ctx, src0, node);
-
break;
case GGML_OP_SOFT_MAX:
if (ctx->fused_topk_moe_mode != TOPK_MOE_COUNT) {
@@ -16369,10 +16315,6 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
case GGML_OP_CONCAT: {
return ggml_vk_concat_supported(op->src[0], op->src[1], op);
}
- case GGML_OP_ADD1:
- return (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32)
- || (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F32)
- || (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F16);
case GGML_OP_ARANGE:
return op->type == GGML_TYPE_F32;
case GGML_OP_FILL:
@@ -16383,8 +16325,6 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
case GGML_OP_PAD_REFLECT_1D:
case GGML_OP_ROLL:
return op->src[0]->type == GGML_TYPE_F32;
- case GGML_OP_DIAG_MASK_INF:
- return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32;
case GGML_OP_SOFT_MAX:
return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32
&& (!op->src[1] || (op->src[1]->type == GGML_TYPE_F32 || op->src[1]->type == GGML_TYPE_F16));
diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/diag_mask_inf.comp b/ggml/src/ggml-vulkan/vulkan-shaders/diag_mask_inf.comp
deleted file mode 100644
index 9cef8a8ec..000000000
--- a/ggml/src/ggml-vulkan/vulkan-shaders/diag_mask_inf.comp
+++ /dev/null
@@ -1,34 +0,0 @@
-#version 450
-
-#extension GL_EXT_shader_16bit_storage : require
-#extension GL_EXT_control_flow_attributes : enable
-
-layout (push_constant) uniform parameter
-{
- uint ncols;
- uint rows_per_channel;
- uint n_past;
-} p;
-
-#include "types.glsl"
-
-layout(local_size_x = 1, local_size_y = 512, local_size_z = 1) in;
-
-layout (binding = 0) readonly buffer X {A_TYPE data_a[];};
-layout (binding = 1) writeonly buffer D {D_TYPE data_d[];};
-
-void main() {
- const uint col = gl_GlobalInvocationID.y;
- const uint row = gl_GlobalInvocationID.x;
-
- if (col >= p.ncols) {
- return;
- }
-
- const uint i = row*p.ncols + col;
- if (col > p.n_past + row % p.rows_per_channel) {
- data_d[i] = D_TYPE(uintBitsToFloat(0xFF800000));
- } else {
- data_d[i] = D_TYPE(data_a[i]);
- }
-}
diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp
index f9ae46ef8..7ccccfa1d 100644
--- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp
+++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp
@@ -1075,8 +1075,6 @@ void process_shaders() {
string_to_spv("silu_back_f32", "silu_back.comp", {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}});
- string_to_spv("diag_mask_inf_f32", "diag_mask_inf.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float"}});
-
string_to_spv("soft_max_f32", "soft_max.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}}));
string_to_spv("soft_max_f32_f16", "soft_max.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"B_TYPE", "float16_t"}, {"D_TYPE", "float"}}));
string_to_spv("soft_max_back_f32", "soft_max_back.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}}));
diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c
index c6a080d22..1fca258a5 100644
--- a/ggml/src/ggml.c
+++ b/ggml/src/ggml.c
@@ -714,13 +714,13 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = {
.from_float_ref = (ggml_from_float_t) quantize_row_q4_1_ref,
},
[4] = { // GGML_TYPE_Q4_2
- .type_name = "DEPRECATED",
+ .type_name = "REMOVED",
.blck_size = 0,
.type_size = 0,
.is_quantized = false,
},
[5] = { // GGML_TYPE_Q4_3
- .type_name = "DEPRECATED",
+ .type_name = "REMOVED",
.blck_size = 0,
.type_size = 0,
.is_quantized = false,
@@ -1009,7 +1009,6 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = {
"DUP",
"ADD",
"ADD_ID",
- "ADD1",
"ACC",
"SUB",
"MUL",
@@ -1051,8 +1050,6 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = {
"GET_ROWS_BACK",
"SET_ROWS",
"DIAG",
- "DIAG_MASK_INF",
- "DIAG_MASK_ZERO",
"SOFT_MAX",
"SOFT_MAX_BACK",
"ROPE",
@@ -1117,7 +1114,7 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = {
"GLU",
};
-static_assert(GGML_OP_COUNT == 102, "GGML_OP_COUNT != 102");
+static_assert(GGML_OP_COUNT == 99, "GGML_OP_COUNT != 99");
static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = {
"none",
@@ -1125,7 +1122,6 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = {
"x",
"x+y",
"x[i]+y",
- "x+y",
"view(x,nb,offset)+=y->x",
"x-y",
"x*y",
@@ -1167,8 +1163,6 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = {
"get_rows_back(x)",
"set_rows(x)",
"diag(x)",
- "diag_mask_inf(x)",
- "diag_mask_zero(x)",
"soft_max(x)",
"soft_max_back(x)",
"rope(x)",
@@ -1233,7 +1227,7 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = {
"glu(x)",
};
-static_assert(GGML_OP_COUNT == 102, "GGML_OP_COUNT != 102");
+static_assert(GGML_OP_COUNT == 99, "GGML_OP_COUNT != 99");
static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2");
@@ -1360,12 +1354,6 @@ size_t ggml_row_size(enum ggml_type type, int64_t ne) {
return ggml_type_size(type)*ne/ggml_blck_size(type);
}
-double ggml_type_sizef(enum ggml_type type) {
- assert(type >= 0);
- assert(type < GGML_TYPE_COUNT);
- return ((double)(type_traits[type].type_size))/type_traits[type].blck_size;
-}
-
const char * ggml_type_name(enum ggml_type type) {
assert(type >= 0);
assert(type < GGML_TYPE_COUNT);
@@ -2172,39 +2160,6 @@ struct ggml_tensor * ggml_add_id(
return result;
}
-// ggml_add1
-
-static struct ggml_tensor * ggml_add1_impl(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- struct ggml_tensor * b,
- bool inplace) {
- GGML_ASSERT(ggml_is_scalar(b));
- GGML_ASSERT(ggml_is_padded_1d(a));
-
- struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
-
- result->op = GGML_OP_ADD1;
- result->src[0] = a;
- result->src[1] = b;
-
- return result;
-}
-
-struct ggml_tensor * ggml_add1(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- struct ggml_tensor * b) {
- return ggml_add1_impl(ctx, a, b, false);
-}
-
-struct ggml_tensor * ggml_add1_inplace(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- struct ggml_tensor * b) {
- return ggml_add1_impl(ctx, a, b, true);
-}
-
// ggml_acc
static struct ggml_tensor * ggml_acc_impl(
@@ -3419,16 +3374,6 @@ struct ggml_tensor * ggml_mul_mat(
return result;
}
-void ggml_mul_mat_set_prec(
- struct ggml_tensor * a,
- enum ggml_prec prec) {
- GGML_ASSERT(a->op == GGML_OP_MUL_MAT);
-
- const int32_t prec_i32 = (int32_t) prec;
-
- ggml_set_op_params_i32(a, 0, prec_i32);
-}
-
void ggml_mul_mat_set_hint(
struct ggml_tensor * a,
enum ggml_op_hint hint) {
@@ -4102,70 +4047,6 @@ struct ggml_tensor * ggml_diag(
return result;
}
-// ggml_diag_mask_inf
-
-static struct ggml_tensor * ggml_diag_mask_inf_impl(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- int n_past,
- bool inplace) {
- struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
-
- int32_t params[] = { n_past };
- ggml_set_op_params(result, params, sizeof(params));
-
- result->op = GGML_OP_DIAG_MASK_INF;
- result->src[0] = a;
-
- return result;
-}
-
-struct ggml_tensor * ggml_diag_mask_inf(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- int n_past) {
- return ggml_diag_mask_inf_impl(ctx, a, n_past, false);
-}
-
-struct ggml_tensor * ggml_diag_mask_inf_inplace(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- int n_past) {
- return ggml_diag_mask_inf_impl(ctx, a, n_past, true);
-}
-
-// ggml_diag_mask_zero
-
-static struct ggml_tensor * ggml_diag_mask_zero_impl(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- int n_past,
- bool inplace) {
- struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
-
- int32_t params[] = { n_past };
- ggml_set_op_params(result, params, sizeof(params));
-
- result->op = GGML_OP_DIAG_MASK_ZERO;
- result->src[0] = a;
-
- return result;
-}
-
-struct ggml_tensor * ggml_diag_mask_zero(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- int n_past) {
- return ggml_diag_mask_zero_impl(ctx, a, n_past, false);
-}
-
-struct ggml_tensor * ggml_diag_mask_zero_inplace(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- int n_past) {
- return ggml_diag_mask_zero_impl(ctx, a, n_past, true);
-}
-
// ggml_clamp
static struct ggml_tensor * ggml_clamp_impl(
@@ -4487,44 +4368,6 @@ struct ggml_tensor * ggml_rope_ext_inplace(
);
}
-struct ggml_tensor * ggml_rope_custom(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- struct ggml_tensor * b,
- int n_dims,
- int mode,
- int n_ctx_orig,
- float freq_base,
- float freq_scale,
- float ext_factor,
- float attn_factor,
- float beta_fast,
- float beta_slow) {
- return ggml_rope_impl(
- ctx, a, b, NULL, n_dims, NULL, mode, n_ctx_orig, freq_base, freq_scale,
- ext_factor, attn_factor, beta_fast, beta_slow, false
- );
-}
-
-struct ggml_tensor * ggml_rope_custom_inplace(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- struct ggml_tensor * b,
- int n_dims,
- int mode,
- int n_ctx_orig,
- float freq_base,
- float freq_scale,
- float ext_factor,
- float attn_factor,
- float beta_fast,
- float beta_slow) {
- return ggml_rope_impl(
- ctx, a, b, NULL, n_dims, NULL, mode, n_ctx_orig, freq_base, freq_scale,
- ext_factor, attn_factor, beta_fast, beta_slow, true
- );
-}
-
// Apparently solving `n_rot = 2pi * x * base^((2 * max_pos_emb) / n_dims)` for x, we get
// `corr_dim(n_rot) = n_dims * log(max_pos_emb / (n_rot * 2pi)) / (2 * log(base))`
static float ggml_rope_yarn_corr_dim(int n_dims, int n_ctx_orig, float n_rot, float base) {
@@ -5264,17 +5107,6 @@ struct ggml_tensor * ggml_upscale(
return ggml_interpolate_impl(ctx, a, a->ne[0] * scale_factor, a->ne[1] * scale_factor, a->ne[2], a->ne[3], mode);
}
-struct ggml_tensor * ggml_upscale_ext(
- struct ggml_context * ctx,
- struct ggml_tensor * a,
- int ne0,
- int ne1,
- int ne2,
- int ne3,
- enum ggml_scale_mode mode) {
- return ggml_interpolate_impl(ctx, a, ne0, ne1, ne2, ne3, mode);
-}
-
struct ggml_tensor * ggml_interpolate(
struct ggml_context * ctx,
struct ggml_tensor * a,
@@ -5654,16 +5486,6 @@ struct ggml_tensor * ggml_flash_attn_ext_rows(
}
-void ggml_flash_attn_ext_set_prec(
- struct ggml_tensor * a,
- enum ggml_prec prec) {
- GGML_ASSERT(a->op == GGML_OP_FLASH_ATTN_EXT);
-
- const int32_t prec_i32 = (int32_t) prec;
-
- ggml_set_op_params_i32(a, 3, prec_i32); // scale is on first pos, max_bias on second
-}
-
enum ggml_prec ggml_flash_attn_ext_get_prec(
const struct ggml_tensor * a) {
GGML_ASSERT(a->op == GGML_OP_FLASH_ATTN_EXT);
@@ -6873,22 +6695,6 @@ static void ggml_acc_or_set(
ggml_build_forward_expand(cgraph, cgraph->grads[isrc]);
}
-static void ggml_add1_or_set(
- struct ggml_context * ctx,
- struct ggml_cgraph * cgraph,
- size_t isrc,
- struct ggml_tensor * tensor) {
- struct ggml_tensor * src = cgraph->visited_hash_set.keys[isrc];
- GGML_ASSERT(src);
- if (cgraph->grads[isrc]) {
- cgraph->grads[isrc] = ggml_add1_impl(ctx, cgraph->grads[isrc], tensor, cgraph->grad_accs[isrc]);
- } else {
- cgraph->grads[isrc] = ggml_repeat(ctx, tensor, src);
- }
- ggml_format_name(cgraph->grads[isrc], "grad for %s", src->name);
- ggml_build_forward_expand(cgraph, cgraph->grads[isrc]);
-}
-
static void ggml_sub_or_set(
struct ggml_context * ctx,
struct ggml_cgraph * cgraph,
@@ -6943,14 +6749,6 @@ static void ggml_compute_backward(
ggml_add_or_set(ctx, cgraph, isrc1, tmp);
}
} break;
- case GGML_OP_ADD1: {
- if (src0_needs_grads) {
- ggml_add_or_set(ctx, cgraph, isrc0, grad);
- }
- if (src1_needs_grads) {
- ggml_add_or_set(ctx, cgraph, isrc1, ggml_mean(ctx, grad)); // TODO: should probably be sum instead of mean
- }
- } break;
case GGML_OP_ACC: {
if (src0_needs_grads) {
ggml_add_or_set(ctx, cgraph, isrc0, grad);
@@ -7023,7 +6821,7 @@ static void ggml_compute_backward(
} break;
case GGML_OP_SUM: {
if (src0_needs_grads) {
- ggml_add1_or_set(ctx, cgraph, isrc0, grad);
+ ggml_add_or_set(ctx, cgraph, isrc0, ggml_repeat(ctx, grad, src0));
}
} break;
case GGML_OP_SUM_ROWS: {
@@ -7033,7 +6831,7 @@ static void ggml_compute_backward(
} break;
case GGML_OP_MEAN: {
if (src0_needs_grads) {
- ggml_add1_or_set(ctx, cgraph, isrc0, ggml_scale_impl(ctx, grad, 1.0f/src0->ne[0], 0.0, false));
+ ggml_add_or_set(ctx, cgraph, isrc0, ggml_repeat(ctx, ggml_scale_impl(ctx, grad, 1.0f/src0->ne[0], 0.0, false), src0));
}
} break;
case GGML_OP_REPEAT: {
@@ -7223,20 +7021,6 @@ static void ggml_compute_backward(
// noop
}
} break;
- case GGML_OP_DIAG_MASK_INF: {
- if (src0_needs_grads) {
- /* ggml_diag_mask_inf_impl() shouldn't be here */
- /* ref: https://github.com/ggml-org/llama.cpp/pull/4203#discussion_r1412377992 */
- const int n_past = ((const int32_t *) tensor->op_params)[0];
- ggml_add_or_set(ctx, cgraph, isrc0, ggml_diag_mask_zero_impl(ctx, grad, n_past, false));
- }
- } break;
- case GGML_OP_DIAG_MASK_ZERO: {
- if (src0_needs_grads) {
- const int n_past = ((const int32_t *) tensor->op_params)[0];
- ggml_add_or_set(ctx, cgraph, isrc0, ggml_diag_mask_zero_impl(ctx, grad, n_past, false));
- }
- } break;
case GGML_OP_TRI: {
if (src0_needs_grads) {
ggml_add_or_set(ctx, cgraph, isrc0, ggml_tri(ctx, grad, (enum ggml_tri_type) ggml_get_op_params_i32(tensor, 0)));
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 4d7fa1188..3abdff5fd 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -6495,33 +6495,6 @@ struct test_trunc : public test_case {
}
};
-// GGML_OP_DIAG_MASK_INF
-struct test_diag_mask_inf : public test_case {
- const ggml_type type;
- const std::array<int64_t, 4> ne;
- const int n_past;
-
- std::string vars() override {
- return VARS_TO_STR3(type, ne, n_past);
- }
-
- test_diag_mask_inf(ggml_type type = GGML_TYPE_F32,
- std::array<int64_t, 4> ne = {10, 10, 3, 2},
- int n_past = 5)
- : type(type), ne(ne), n_past(n_past) {}
-
- ggml_tensor * build_graph(ggml_context * ctx) override {
- ggml_tensor * a = ggml_new_tensor(ctx, type, 4, ne.data());
- ggml_set_param(a);
- ggml_set_name(a, "a");
-
- ggml_tensor * out = ggml_diag_mask_inf(ctx, a, n_past);
- ggml_set_name(out, "out");
-
- return out;
- }
-};
-
// GGML_OP_SOFT_MAX
struct test_soft_max : public test_case {
const ggml_type type;
@@ -12418,10 +12391,6 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_trunc (type, {1024, 1024, 1, 1}));
}
- test_cases.emplace_back(new test_diag_mask_inf(GGML_TYPE_F32, {10, 10, 1, 1}, 5));
- test_cases.emplace_back(new test_diag_mask_inf(GGML_TYPE_F32, {10, 10, 3, 1}, 5));
- test_cases.emplace_back(new test_diag_mask_inf(GGML_TYPE_F32, {10, 10, 3, 2}, 5));
-
#if 0
std::uniform_int_distribution<> dist_ne1(1, 50);
int exponent = 1;
@@ -12962,7 +12931,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
for (int kv : { 113, 512, 1024, }) {
if (nr2 != 1 && kv != 512) continue;
for (int nb : { 1, 3, 32, 75, }) {
- for (ggml_prec prec : {GGML_PREC_F32, GGML_PREC_BF16, GGML_PREC_DEFAULT}) {
+ for (ggml_prec prec : {GGML_PREC_F32, GGML_PREC_BF16, GGML_PREC_UNDEFINED}) {
if (hsk != 128 && prec != GGML_PREC_F32) continue;
for (ggml_type type_KV : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_BF16, GGML_TYPE_Q8_0, GGML_TYPE_Q5_1, GGML_TYPE_Q5_0, GGML_TYPE_Q4_1, GGML_TYPE_Q4_0, GGML_TYPE_IQ4_NL}) {
if (type_KV != GGML_TYPE_F16 && hsk != 64 && hsk != 72) continue;