mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-09-04 02:37:27 +02:00
sycl : add Kronecker product FWHT support for sizes 384, 640, 768, 1280 (#28016)
This commit is contained in:
Jingxin (Philip) Li
@@ -1,6 +1,50 @@
|
|||||||
#include "fwht.hpp"
|
#include "fwht.hpp"
|
||||||
|
|
||||||
#include <cmath>
|
#include <cmath>
|
||||||
|
#define P 1.0f
|
||||||
|
#define N -1.0f
|
||||||
|
|
||||||
|
// constant Hadamard matrix via Paley I construction
|
||||||
|
static constexpr float H12[12][12] = {
|
||||||
|
{ P, P, P, P, P, P, P, P, P, P, P, P },
|
||||||
|
{ P, N, P, N, P, P, P, N, N, N, P, N },
|
||||||
|
{ P, N, N, P, N, P, P, P, N, N, N, P },
|
||||||
|
{ P, P, N, N, P, N, P, P, P, N, N, N },
|
||||||
|
{ P, N, P, N, N, P, N, P, P, P, N, N },
|
||||||
|
{ P, N, N, P, N, N, P, N, P, P, P, N },
|
||||||
|
{ P, N, N, N, P, N, N, P, N, P, P, P },
|
||||||
|
{ P, P, N, N, N, P, N, N, P, N, P, P },
|
||||||
|
{ P, P, P, N, N, N, P, N, N, P, N, P },
|
||||||
|
{ P, P, P, P, N, N, N, P, N, N, P, N },
|
||||||
|
{ P, N, P, P, P, N, N, N, P, N, N, P },
|
||||||
|
{ P, P, N, P, P, P, N, N, N, P, N, N }
|
||||||
|
};
|
||||||
|
|
||||||
|
static constexpr float H20[20][20] = {
|
||||||
|
{ P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P },
|
||||||
|
{ P, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N },
|
||||||
|
{ P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P },
|
||||||
|
{ P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P },
|
||||||
|
{ P, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N },
|
||||||
|
{ P, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N },
|
||||||
|
{ P, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N },
|
||||||
|
{ P, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N },
|
||||||
|
{ P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P },
|
||||||
|
{ P, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N },
|
||||||
|
{ P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P },
|
||||||
|
{ P, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N },
|
||||||
|
{ P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P },
|
||||||
|
{ P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P },
|
||||||
|
{ P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P },
|
||||||
|
{ P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P },
|
||||||
|
{ P, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N },
|
||||||
|
{ P, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N },
|
||||||
|
{ P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P },
|
||||||
|
{ P, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N }
|
||||||
|
};
|
||||||
|
|
||||||
|
#undef P
|
||||||
|
#undef N
|
||||||
|
|
||||||
template <int N>
|
template <int N>
|
||||||
static void fwht_kernel(const float * __restrict__ src, float * __restrict__ dst, const int64_t n_rows,
|
static void fwht_kernel(const float * __restrict__ src, float * __restrict__ dst, const int64_t n_rows,
|
||||||
@@ -80,6 +124,122 @@ static void launch_fwht(const float * src, float * dst, const int64_t n_rows, co
|
|||||||
});
|
});
|
||||||
}
|
}
|
||||||
|
|
||||||
|
template <int N, int m>
|
||||||
|
static void kronecker_kernel(const float * __restrict__ src,
|
||||||
|
float * __restrict__ dst,
|
||||||
|
const int64_t n_rows,
|
||||||
|
const float scale,
|
||||||
|
const sycl::nd_item<2> & item) {
|
||||||
|
static_assert(m == 12 || m == 20, "block size has to be 12 or 20.");
|
||||||
|
|
||||||
|
const sycl::sub_group sg = item.get_sub_group();
|
||||||
|
|
||||||
|
const int64_t r = item.get_global_id(0);
|
||||||
|
if (r >= n_rows) {
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
src += r * N;
|
||||||
|
dst += r * N;
|
||||||
|
|
||||||
|
constexpr int blocks_per_group = N / m;
|
||||||
|
constexpr int el_w = blocks_per_group / WARP_SIZE;
|
||||||
|
static_assert(el_w >= 1 && blocks_per_group % WARP_SIZE == 0, "blocks_per_group must be a multiple of WARP_SIZE");
|
||||||
|
float reg[el_w * m];
|
||||||
|
const int lane = sg.get_local_linear_id();
|
||||||
|
|
||||||
|
#pragma unroll
|
||||||
|
for (int i = 0; i < el_w; ++i) {
|
||||||
|
const int b_idx = i * WARP_SIZE + lane;
|
||||||
|
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < m; ++j) {
|
||||||
|
reg[i * m + j] = src[b_idx * m + j] * scale;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
#pragma unroll
|
||||||
|
for (int b = 0; b < el_w; ++b) {
|
||||||
|
float z[m] = { 0.0f };
|
||||||
|
|
||||||
|
#pragma unroll
|
||||||
|
for (int i = 0; i < m; ++i) {
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < m; ++j) {
|
||||||
|
const float h = (m == 12 ? H12[j][i] : H20[j][i]);
|
||||||
|
z[i] += reg[b * m + j] * h;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
#pragma unroll
|
||||||
|
for (int i = 0; i < m; ++i) {
|
||||||
|
reg[b * m + i] = z[i];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
#pragma unroll
|
||||||
|
for (int h = 1; h < WARP_SIZE; h *= 2) {
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < el_w; ++j) {
|
||||||
|
#pragma unroll
|
||||||
|
for (int k = 0; k < m; ++k) {
|
||||||
|
const float val = reg[j * m + k];
|
||||||
|
const float val2 = dpct::permute_sub_group_by_xor(sg, val, h, WARP_SIZE);
|
||||||
|
|
||||||
|
reg[j * m + k] = (lane & h) == 0 ? val + val2 : val2 - val;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
#pragma unroll
|
||||||
|
for (int h = WARP_SIZE; h < blocks_per_group; h *= 2) {
|
||||||
|
const int step = h / WARP_SIZE;
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < el_w; j += 2 * step) {
|
||||||
|
#pragma unroll
|
||||||
|
for (int s = 0; s < step; ++s) {
|
||||||
|
#pragma unroll
|
||||||
|
for (int k = 0; k < m; ++k) {
|
||||||
|
const float x = reg[(j + s) * m + k];
|
||||||
|
const float y = reg[(j + s + step) * m + k];
|
||||||
|
|
||||||
|
reg[(j + s) * m + k] = x + y;
|
||||||
|
reg[(j + s + step) * m + k] = x - y;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
#pragma unroll
|
||||||
|
for (int i = 0; i < el_w; ++i) {
|
||||||
|
const int b_idx = i * WARP_SIZE + lane;
|
||||||
|
#pragma unroll
|
||||||
|
for (int k = 0; k < m; ++k) {
|
||||||
|
dst[b_idx * m + k] = reg[i * m + k];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
template <int N, int m>
|
||||||
|
static void launch_kronecker(const float * src,
|
||||||
|
float * dst,
|
||||||
|
const int64_t n_rows,
|
||||||
|
const float scale,
|
||||||
|
dpct::queue_ptr stream) {
|
||||||
|
constexpr int rows_per_block = 4;
|
||||||
|
|
||||||
|
const int64_t num_blocks = (n_rows + rows_per_block - 1) / rows_per_block;
|
||||||
|
|
||||||
|
// dim 1 is the fastest-varying, so a sub-group is exactly one row's WARP_SIZE lanes.
|
||||||
|
const sycl::range<2> global(num_blocks * rows_per_block, WARP_SIZE);
|
||||||
|
const sycl::range<2> local(rows_per_block, WARP_SIZE);
|
||||||
|
|
||||||
|
stream->parallel_for(sycl::nd_range<2>(global, local),
|
||||||
|
[=](sycl::nd_item<2> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||||
|
kronecker_kernel<N, m>(src, dst, n_rows, scale, item);
|
||||||
|
});
|
||||||
|
}
|
||||||
|
|
||||||
bool ggml_sycl_op_fwht(ggml_backend_sycl_context & ctx, const ggml_tensor * src, ggml_tensor * dst) {
|
bool ggml_sycl_op_fwht(ggml_backend_sycl_context & ctx, const ggml_tensor * src, ggml_tensor * dst) {
|
||||||
if (src->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
|
if (src->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
|
||||||
return false;
|
return false;
|
||||||
@@ -113,6 +273,18 @@ bool ggml_sycl_op_fwht(ggml_backend_sycl_context & ctx, const ggml_tensor * src,
|
|||||||
case 512:
|
case 512:
|
||||||
launch_fwht<512>(src_d, dst_d, rows, scale, stream);
|
launch_fwht<512>(src_d, dst_d, rows, scale, stream);
|
||||||
return true;
|
return true;
|
||||||
|
case 384:
|
||||||
|
launch_kronecker<384, 12>(src_d, dst_d, rows, scale, stream);
|
||||||
|
return true;
|
||||||
|
case 768:
|
||||||
|
launch_kronecker<768, 12>(src_d, dst_d, rows, scale, stream);
|
||||||
|
return true;
|
||||||
|
case 640:
|
||||||
|
launch_kronecker<640, 20>(src_d, dst_d, rows, scale, stream);
|
||||||
|
return true;
|
||||||
|
case 1280:
|
||||||
|
launch_kronecker<1280, 20>(src_d, dst_d, rows, scale, stream);
|
||||||
|
return true;
|
||||||
default:
|
default:
|
||||||
return false;
|
return false;
|
||||||
}
|
}
|
||||||
|
|||||||
+111
-13
@@ -4660,6 +4660,51 @@ struct test_mul_mat : public test_case {
|
|||||||
}
|
}
|
||||||
};
|
};
|
||||||
|
|
||||||
|
#define P 1.0f
|
||||||
|
#define N -1.0f
|
||||||
|
|
||||||
|
// constant Hadamard matrix via Paley I construction
|
||||||
|
static constexpr float H12[12][12] = {
|
||||||
|
{ P, P, P, P, P, P, P, P, P, P, P, P },
|
||||||
|
{ P, N, P, N, P, P, P, N, N, N, P, N },
|
||||||
|
{ P, N, N, P, N, P, P, P, N, N, N, P },
|
||||||
|
{ P, P, N, N, P, N, P, P, P, N, N, N },
|
||||||
|
{ P, N, P, N, N, P, N, P, P, P, N, N },
|
||||||
|
{ P, N, N, P, N, N, P, N, P, P, P, N },
|
||||||
|
{ P, N, N, N, P, N, N, P, N, P, P, P },
|
||||||
|
{ P, P, N, N, N, P, N, N, P, N, P, P },
|
||||||
|
{ P, P, P, N, N, N, P, N, N, P, N, P },
|
||||||
|
{ P, P, P, P, N, N, N, P, N, N, P, N },
|
||||||
|
{ P, N, P, P, P, N, N, N, P, N, N, P },
|
||||||
|
{ P, P, N, P, P, P, N, N, N, P, N, N }
|
||||||
|
};
|
||||||
|
|
||||||
|
static constexpr float H20[20][20] = {
|
||||||
|
{ P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P },
|
||||||
|
{ P, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N },
|
||||||
|
{ P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P },
|
||||||
|
{ P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P },
|
||||||
|
{ P, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N },
|
||||||
|
{ P, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N },
|
||||||
|
{ P, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N },
|
||||||
|
{ P, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N },
|
||||||
|
{ P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P },
|
||||||
|
{ P, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N },
|
||||||
|
{ P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P },
|
||||||
|
{ P, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N },
|
||||||
|
{ P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P },
|
||||||
|
{ P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P },
|
||||||
|
{ P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P },
|
||||||
|
{ P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P },
|
||||||
|
{ P, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N },
|
||||||
|
{ P, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N },
|
||||||
|
{ P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P },
|
||||||
|
{ P, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N }
|
||||||
|
};
|
||||||
|
|
||||||
|
#undef P
|
||||||
|
#undef N
|
||||||
|
|
||||||
// GGML_HINT_SRC0_IS_HADAMARD
|
// GGML_HINT_SRC0_IS_HADAMARD
|
||||||
struct test_mul_mat_hadamard : public test_mul_mat {
|
struct test_mul_mat_hadamard : public test_mul_mat {
|
||||||
test_mul_mat_hadamard(ggml_type type_a = GGML_TYPE_F32, ggml_type type_b = GGML_TYPE_F32,
|
test_mul_mat_hadamard(ggml_type type_a = GGML_TYPE_F32, ggml_type type_b = GGML_TYPE_F32,
|
||||||
@@ -4684,20 +4729,57 @@ struct test_mul_mat_hadamard : public test_mul_mat {
|
|||||||
void initialize_tensors(ggml_context * ctx) override {
|
void initialize_tensors(ggml_context * ctx) override {
|
||||||
for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
|
for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
|
||||||
if (strcmp(t->name, "a") == 0) {
|
if (strcmp(t->name, "a") == 0) {
|
||||||
const int64_t n_cols = t->ne[0];
|
const int64_t n_cols = t->ne[0];
|
||||||
const int64_t n_rows = ggml_nrows(t);
|
const int64_t n_rows = ggml_nrows(t);
|
||||||
std::vector<float> data(n_cols * n_rows);
|
std::vector<float> data(n_cols * n_rows);
|
||||||
float scale = 1.0f / sqrtf((float)n_cols);
|
float scale = 1.0f / sqrtf((float) n_cols);
|
||||||
for (int64_t r = 0; r < n_rows; r++) {
|
|
||||||
float * row_data = data.data() + r * n_cols;
|
auto is_pow2 = [](const int64_t a) {
|
||||||
for (int64_t i = 0; i < n_cols; i++) {
|
return (a > 0) && ((a & (a - 1)) == 0);
|
||||||
int pop = 0;
|
};
|
||||||
int64_t val = r & i;
|
|
||||||
while (val) {
|
const bool is_kronecker =
|
||||||
pop += (val & 1);
|
((n_cols % 12 == 0) && is_pow2(n_cols / 12)) || ((n_cols % 20 == 0) && is_pow2(n_cols / 20));
|
||||||
val >>= 1;
|
|
||||||
|
if (is_kronecker) {
|
||||||
|
const int64_t B = (n_cols % 12 == 0 && is_pow2(n_cols / 12)) ? 12 : 20;
|
||||||
|
const int64_t M = n_cols / B;
|
||||||
|
for (int64_t r = 0; r < n_rows; r++) {
|
||||||
|
float * row_data = data.data() + r * n_cols;
|
||||||
|
const int64_t r_mod = r % n_cols;
|
||||||
|
const int64_t r_b = r_mod / B;
|
||||||
|
const int64_t r_m = r_mod % B;
|
||||||
|
|
||||||
|
for (int64_t i = 0; i < n_cols; i++) {
|
||||||
|
const int64_t c_b = i / B;
|
||||||
|
const int64_t c_m = i % B;
|
||||||
|
|
||||||
|
int pop = 0;
|
||||||
|
int64_t val = r_b & c_b;
|
||||||
|
while (val) {
|
||||||
|
pop += (val & 1);
|
||||||
|
val >>= 1;
|
||||||
|
}
|
||||||
|
const float sign_m = (pop % 2 == 0) ? 1.0f : -1.0f;
|
||||||
|
const float sign_b = (B == 12) ? H12[c_m][r_m] : H20[c_m][r_m];
|
||||||
|
|
||||||
|
row_data[i] = scale * sign_b * sign_m;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
else if (is_pow2(n_cols)) {
|
||||||
|
for (int64_t r = 0; r < n_rows; r++) {
|
||||||
|
float * row_data = data.data() + r * n_cols;
|
||||||
|
for (int64_t i = 0; i < n_cols; i++) {
|
||||||
|
int pop_cnt = 0;
|
||||||
|
int64_t val = r & i;
|
||||||
|
while (val) {
|
||||||
|
pop_cnt += (val & 1);
|
||||||
|
val >>= 1;
|
||||||
|
}
|
||||||
|
row_data[i] = (pop_cnt % 2 == 0) ? scale : -scale;
|
||||||
}
|
}
|
||||||
row_data[i] = (pop % 2 == 0) ? scale : -scale;
|
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
ggml_backend_tensor_set(t, data.data(), 0, data.size() * sizeof(float));
|
ggml_backend_tensor_set(t, data.data(), 0, data.size() * sizeof(float));
|
||||||
@@ -9298,6 +9380,14 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
|
|||||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 256, 512, 256)); // many rows
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 256, 512, 256)); // many rows
|
||||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 32, 1, 32)); // too small (N<64)
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 32, 1, 32)); // too small (N<64)
|
||||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1024, 1, 1024)); // too big (N>512)
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1024, 1, 1024)); // too big (N>512)
|
||||||
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 1, 384)); // m=12 (N=384)
|
||||||
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 32, 384)); // m=12 (batch)
|
||||||
|
test_cases.emplace_back(
|
||||||
|
new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 4, 384, { 2, 3 })); // m=12 (multi-dim)
|
||||||
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 768, 1, 768)); // m=12 (N=768)
|
||||||
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 1, 640)); // m=20 (N=640)
|
||||||
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 32, 640)); // m=20 (batch)
|
||||||
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1280, 1, 1280)); // m=20 (N=1280)
|
||||||
|
|
||||||
#if 0
|
#if 0
|
||||||
// > 4GB A matrix. Too slow to be enabled by default.
|
// > 4GB A matrix. Too slow to be enabled by default.
|
||||||
@@ -10500,7 +10590,15 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
|
|||||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 128, 2048, 128));
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 128, 2048, 128));
|
||||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 256, 2048, 256));
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 256, 2048, 256));
|
||||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 512, 2048, 512));
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 512, 2048, 512));
|
||||||
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 1, 384)); // m=12 (N=384)
|
||||||
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 32, 384)); // m=12 (batch)
|
||||||
|
test_cases.emplace_back(
|
||||||
|
new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 4, 384, { 2, 3 })); // m=12 (multi-dim)
|
||||||
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 768, 1, 768)); // m=12 (N=768)
|
||||||
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 1, 640)); // m=20 (N=640)
|
||||||
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 32, 640)); // m=20 (batch)
|
||||||
|
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1280, 1, 1280)); // m=20 (N=1280)
|
||||||
|
|
||||||
test_cases.emplace_back(new test_solve_tri(GGML_TYPE_F32, { 64, 64, 4, 4 }, { 32, 64, 4, 4 }));
|
test_cases.emplace_back(new test_solve_tri(GGML_TYPE_F32, { 64, 64, 4, 4 }, { 32, 64, 4, 4 }));
|
||||||
test_cases.emplace_back(new test_solve_tri(GGML_TYPE_F32, { 128, 128, 4, 2 }, { 32, 128, 4, 2 }));
|
test_cases.emplace_back(new test_solve_tri(GGML_TYPE_F32, { 128, 128, 4, 2 }, { 32, 128, 4, 2 }));
|
||||||
// qwen3next with CHUNK_SIZE 64
|
// qwen3next with CHUNK_SIZE 64
|
||||||
|
|||||||
Reference in New Issue
Block a user