Skip to content

Commit f5aeac5

Browse files
wangshuai09wangshuai09
committed
cann: add quantize_fp16_q4_0
1 parent 5af1609 commit f5aeac5

File tree

6 files changed

+271
-12
lines changed

6 files changed

+271
-12
lines changed

‎ggml/src/ggml-cann.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -627,7 +627,6 @@ GGML_CALL static void* ggml_backend_cann_buffer_get_base(
627627
GGML_CALL static void ggml_backend_cann_transform_q4_0(ggml_tensor* tensor,
628628
const void* src,
629629
void* dst) {
630-
GGML_ASSERT(tensor->op == GGML_OP_NONE);
631630

632631
int64_t n_elems = ggml_nelements(tensor);
633632
int64_t groups = n_elems / QK4_0;
@@ -679,7 +678,6 @@ GGML_CALL static void ggml_backend_cann_transform_q4_0(ggml_tensor* tensor,
679678
*/
680679
GGML_CALL static void ggml_backend_cann_transform_back_q4_0(
681680
const ggml_tensor* tensor, void* src, void* dst) {
682-
GGML_ASSERT(tensor->op == GGML_OP_NONE);
683681

684682
int64_t n_elems = ggml_nelements(tensor);
685683
int64_t groups = n_elems / QK4_0;
@@ -1694,6 +1692,7 @@ GGML_CALL static bool ggml_backend_cann_supports_op(ggml_backend_t backend,
16941692
case GGML_TYPE_F32:
16951693
case GGML_TYPE_F16:
16961694
case GGML_TYPE_Q8_0:
1695+
case GGML_TYPE_Q4_0:
16971696
return true;
16981697
default:
16991698
return false;

‎ggml/src/ggml-cann/aclnn_ops.cpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -910,6 +910,13 @@ void ggml_cann_dup(ggml_backend_cann_context& ctx, ggml_tensor* dst) {
910910
((ggml_tensor*)dst->extra)->ne);
911911
return;
912912
}
913+
if (dst->type == GGML_TYPE_Q4_0) {
914+
aclrtlaunch_ascendc_quantize_f16_q4_0(
915+
2, ctx.stream(), src->data, dst->data,
916+
((ggml_tensor*)src->extra)->ne, ((ggml_tensor*)src->extra)->nb,
917+
((ggml_tensor*)dst->extra)->ne);
918+
return;
919+
}
913920
if (dst->type == GGML_TYPE_F16) {
914921
if (ggml_are_same_shape(src, dst)) {
915922
cann_copy(ctx, acl_src, acl_dst);

‎ggml/src/ggml-cann/kernels/CMakeLists.txt

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,7 @@ file(GLOB SRC_FILES
99
get_row_q8_0.cpp
1010
quantize_f32_q8_0.cpp
1111
quantize_f16_q8_0.cpp
12+
quantize_f16_q4_0.cpp
1213
dup.cpp
1314
)
1415

@@ -29,4 +30,4 @@ ascendc_library(ascendc_kernels STATIC
2930
${SRC_FILES}
3031
)
3132

32-
#ascendc_compile_definitions(ascendc_kernels PRIVATE -DASCENDC_DUMP)
33+
# ascendc_compile_definitions(ascendc_kernels PRIVATE -DASCENDC_DUMP)

‎ggml/src/ggml-cann/kernels/ascendc_kernels.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,7 @@
88

99
#include "aclrtlaunch_ascendc_quantize_f32_q8_0.h"
1010
#include "aclrtlaunch_ascendc_quantize_f16_q8_0.h"
11+
#include "aclrtlaunch_ascendc_quantize_f16_q4_0.h"
1112

1213
#include "aclrtlaunch_ascendc_dup_by_rows_fp16.h"
1314
#include "aclrtlaunch_ascendc_dup_by_rows_fp32.h"

‎ggml/src/ggml-cann/kernels/quantize_f16_q4_0.cpp

Lines changed: 231 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,231 @@
1+
#include "kernel_operator.h"
2+
3+
using namespace AscendC;
4+
5+
#define BUFFER_NUM 2
6+
#define Group_Size 32
7+
8+
class QUANTIZE_F16_Q4_0 {
9+
public:
10+
__aicore__ inline QUANTIZE_F16_Q4_0() {}
11+
__aicore__ inline void init(GM_ADDR input, GM_ADDR output,
12+
int64_t *input_ne_ub, size_t *input_nb_ub,
13+
int64_t *output_ne_ub) {
14+
int64_t op_block_num = GetBlockNum();
15+
int64_t op_block_idx = GetBlockIdx();
16+
17+
for (int i = 0; i < 4; i++) {
18+
input_ne[i] = input_ne_ub[i];
19+
input_stride[i] = input_nb_ub[i] / input_nb_ub[0];
20+
output_ne[i] = output_ne_ub[i];
21+
}
22+
23+
output_stride[0] = 1;
24+
for (int i = 1; i < 4; i++) {
25+
output_stride[i] = output_stride[i - 1] * output_ne[i - 1];
26+
}
27+
28+
// scale saved one by one:. [group1_scale, group2_scale, ...]
29+
scale_ne = input_ne;
30+
scale_stride[0] = 1;
31+
scale_stride[1] = input_ne[0] / Group_Size;
32+
for (int i = 2; i < 4; i++) {
33+
scale_stride[i] = scale_stride[i - 1] * scale_ne[i - 1];
34+
}
35+
36+
// split input tensor by rows.
37+
uint64_t nr = input_ne[1] * input_ne[2] * input_ne[3];
38+
dr = nr / op_block_num;
39+
40+
uint64_t tails = nr % op_block_num;
41+
if (op_block_idx < tails) {
42+
dr += 1;
43+
ir = dr * op_block_idx;
44+
} else {
45+
ir = dr * op_block_idx + tails;
46+
}
47+
48+
group_size_in_row = scale_stride[1];
49+
int64_t scale_offset = output_ne[0] * output_ne[1] * output_ne[2] *
50+
output_ne[3] * sizeof(uint8_t) / 2;
51+
52+
input_gm.SetGlobalBuffer((__gm__ half *)input);
53+
output_gm.SetGlobalBuffer((__gm__ int4b_t *)output);
54+
scale_gm.SetGlobalBuffer((__gm__ half *)(output + scale_offset + ir *
55+
group_size_in_row *
56+
sizeof(half)));
57+
58+
pipe.InitBuffer(input_queue, BUFFER_NUM, Group_Size * sizeof(half));
59+
pipe.InitBuffer(output_queue, BUFFER_NUM, Group_Size * sizeof(int4b_t));
60+
pipe.InitBuffer(work_queue, 1, 32);
61+
pipe.InitBuffer(max_queue, 1, 32);
62+
pipe.InitBuffer(min_queue, 1, 32);
63+
pipe.InitBuffer(scale_queue, 1, 32);
64+
pipe.InitBuffer(int8_queue, 1, 32);
65+
pipe.InitBuffer(cast_queue , 1 , Group_Size * sizeof(float));
66+
}
67+
68+
__aicore__ inline void copy_in(uint32_t offset) {
69+
LocalTensor<half> input_local = input_queue.AllocTensor<half>();
70+
DataCopy(input_local, input_gm[offset], Group_Size);
71+
input_queue.EnQue(input_local);
72+
}
73+
74+
__aicore__ inline void copy_out(uint32_t offset) {
75+
LocalTensor<int4b_t> output_local = output_queue.DeQue<int4b_t>();
76+
DataCopy(output_gm[offset], output_local, Group_Size);
77+
output_queue.FreeTensor(output_local);
78+
}
79+
80+
__aicore__ inline half calculate_group(int64_t row, int64_t group) {
81+
const int64_t i3 = row / (input_ne[1] * input_ne[2]);
82+
const int64_t i2 = (row - i3 * input_ne[1] * input_ne[2]) / input_ne[1];
83+
const int64_t i1 =
84+
row - i3 * input_ne[1] * input_ne[2] - i2 * input_ne[1];
85+
86+
const int64_t input_offset = i1 * input_stride[1] +
87+
i2 * input_stride[2] +
88+
i3 * input_stride[3] + Group_Size * group;
89+
90+
const int64_t output_offset = i1 * output_stride[1] +
91+
i2 * output_stride[2] +
92+
i3 * output_stride[3] + Group_Size * group;
93+
94+
PRINTF("output offset %d \n", output_offset);
95+
PRINTF("group %d \n", group);
96+
PRINTF("i1 %d, i2 %d, i3 %d, output_stride1 %d\n", i1, i2, i3, output_stride[1]);
97+
98+
copy_in(input_offset);
99+
LocalTensor<half> input_local = input_queue.DeQue<half>();
100+
LocalTensor<int4b_t> output_local = output_queue.AllocTensor<int4b_t>();
101+
LocalTensor<float> work_local = work_queue.AllocTensor<float>();
102+
LocalTensor<float> max_local = max_queue.AllocTensor<float>();
103+
LocalTensor<float> min_local = min_queue.AllocTensor<float>();
104+
LocalTensor<float> cast_local = cast_queue.AllocTensor<float>();
105+
LocalTensor<int8_t> int8_local = int8_queue.AllocTensor<int8_t>();
106+
107+
// TODO: OPTIMIZE
108+
Cast(cast_local, input_local, RoundMode::CAST_NONE, Group_Size);
109+
ReduceMax(max_local, cast_local, work_local, Group_Size);
110+
ReduceMin(min_local, cast_local, work_local, Group_Size);
111+
const float max_value = max_local.GetValue(0);
112+
const float min_value = min_local.GetValue(0);
113+
float d = max_value;
114+
if (min_value < 0 && (-1 * min_value) > max_value) {
115+
d = min_value;
116+
}
117+
PRINTF("d %f \n", d);
118+
pipe_barrier(PIPE_ALL);
119+
d = d / (-8);
120+
if (d != 0) {
121+
Muls(cast_local, cast_local, 1.0f / d, Group_Size);
122+
}
123+
124+
//
125+
Cast(input_local, cast_local, RoundMode::CAST_ROUND, Group_Size);
126+
Cast(output_local, input_local, RoundMode::CAST_ROUND, Group_Size);
127+
128+
output_queue.EnQue(output_local);
129+
130+
//
131+
PRINTF("output: ");
132+
for(int i =0; i<32; i++) {
133+
PRINTF("%f, ", cast_local.GetValue(i));
134+
}
135+
PRINTF("\n");
136+
copy_out(output_offset);
137+
138+
input_queue.FreeTensor(input_local);
139+
work_queue.FreeTensor(work_local);
140+
max_queue.FreeTensor(max_local);
141+
min_queue.FreeTensor(min_local);
142+
int8_queue.FreeTensor(int8_local);
143+
cast_queue.FreeTensor(cast_local);
144+
return (half)d;
145+
}
146+
147+
__aicore__ inline void calculate() {
148+
LocalTensor<half> scale_local = scale_queue.AllocTensor<half>();
149+
uint32_t scale_local_offset = 0;
150+
uint32_t scale_global_offset = 0;
151+
for (int64_t i = ir; i < ir + dr; i++) {
152+
for (int64_t j = 0; j < group_size_in_row; j++) {
153+
half scale = calculate_group(i, j);
154+
scale_local.SetValue(scale_local_offset++, scale);
155+
if (scale_local_offset == 16) {
156+
scale_local_offset = 0;
157+
// TODO: OPTIMIZE ME
158+
pipe_barrier(PIPE_ALL);
159+
DataCopy(scale_gm[scale_global_offset], scale_local, 16);
160+
pipe_barrier(PIPE_ALL);
161+
scale_global_offset += 16;
162+
}
163+
}
164+
}
165+
166+
if (scale_local_offset != 0) {
167+
pipe_barrier(PIPE_ALL);
168+
DataCopyExtParams dataCopyParams;
169+
dataCopyParams.blockCount = 1;
170+
dataCopyParams.blockLen = scale_local_offset * sizeof(half);
171+
DataCopyPad(scale_gm[scale_global_offset], scale_local,
172+
dataCopyParams);
173+
pipe_barrier(PIPE_ALL);
174+
}
175+
}
176+
177+
private:
178+
int64_t input_ne[4];
179+
size_t input_stride[4];
180+
181+
int64_t *scale_ne;
182+
size_t scale_stride[4];
183+
184+
int64_t output_ne[4];
185+
size_t output_stride[4];
186+
187+
int64_t group_size_in_row;
188+
189+
int64_t ir;
190+
int64_t dr;
191+
192+
TPipe pipe;
193+
GlobalTensor<half> input_gm;
194+
GlobalTensor<half> scale_gm;
195+
GlobalTensor<int4b_t> output_gm;
196+
TQue<QuePosition::VECIN, BUFFER_NUM> input_queue;
197+
TQue<QuePosition::VECOUT, BUFFER_NUM> output_queue;
198+
TQue<QuePosition::VECIN, 1> work_queue;
199+
TQue<QuePosition::VECOUT, 1> max_queue;
200+
TQue<QuePosition::VECOUT, 1> min_queue;
201+
TQue<QuePosition::VECOUT, 1> scale_queue;
202+
TQue<QuePosition::VECOUT, 1> cast_queue;
203+
TQue<QuePosition::VECOUT, 1> int8_queue;
204+
TQue<QuePosition::VECOUT, 1> const_15_queue;
205+
206+
};
207+
208+
template <typename T>
209+
__aicore__ inline void copy_to_ub(GM_ADDR gm, T *ub, size_t size) {
210+
auto gm_ptr = (__gm__ uint8_t *)gm;
211+
auto ub_ptr = (uint8_t *)(ub);
212+
for (int32_t i = 0; i < size; ++i, ++ub_ptr, ++gm_ptr) {
213+
*ub_ptr = *gm_ptr;
214+
}
215+
}
216+
217+
extern "C" __global__ __aicore__ void ascendc_quantize_f16_q4_0(
218+
GM_ADDR input_gm, GM_ADDR output_gm, GM_ADDR input_ne_gm,
219+
GM_ADDR input_nb_gm, GM_ADDR output_ne_gm) {
220+
int64_t input_ne_ub[4];
221+
size_t input_nb_ub[4];
222+
int64_t output_ne_ub[4];
223+
224+
copy_to_ub(input_ne_gm, input_ne_ub, 32);
225+
copy_to_ub(input_nb_gm, input_nb_ub, 32);
226+
copy_to_ub(output_ne_gm, output_ne_ub, 32);
227+
228+
QUANTIZE_F16_Q4_0 op;
229+
op.init(input_gm, output_gm, input_ne_ub, input_nb_ub, output_ne_ub);
230+
op.calculate();
231+
}

‎tests/test-backend-ops.cpp

Lines changed: 29 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -482,6 +482,24 @@ struct test_case {
482482
std::vector<float> f1 = tensor_to_float(t1);
483483
std::vector<float> f2 = tensor_to_float(t2);
484484

485+
printf("f1:\n ");
486+
for (int i=0; i<128; i++) {
487+
printf("%f,", f1[i]);
488+
if ((i+1)%32==0) {
489+
printf("\n");
490+
}
491+
}
492+
printf("\n");
493+
printf("f2: \n");
494+
for (int i=0; i<128; i++) {
495+
printf("%f,", f2[i]);
496+
if ((i+1)%32==0) {
497+
printf("\n");
498+
}
499+
}
500+
printf("\n");
501+
// printf("%f, %f \n", f1[0], f2[0]);
502+
485503
for (size_t i = 0; i < f1.size(); i++) {
486504
// check for nans
487505
if (std::isnan(f1[i]) || std::isnan(f2[i])) {
@@ -2170,17 +2188,19 @@ static bool test_backend(ggml_backend_t backend, test_mode mode, const char * op
21702188
test_cases.emplace_back(new test_dup(GGML_TYPE_I16, {10, 8, 3, 1}, {0, 2, 1, 3}));
21712189
test_cases.emplace_back(new test_dup(GGML_TYPE_I16, {10, 8, 3, 1}, {1, 2, 0, 3}));
21722190

2173-
for (ggml_type type_src : {GGML_TYPE_F16, GGML_TYPE_F32}) {
2174-
for (ggml_type type_dst : all_types) {
2175-
test_cases.emplace_back(new test_cpy(type_src, type_dst, {256, 4, 4, 4}));
2176-
test_cases.emplace_back(new test_cpy(type_src, type_dst, {256, 2, 3, 4}, {0, 2, 1, 3})); // cpy by rows
2177-
}
2178-
}
2179-
for (ggml_type type_src : {GGML_TYPE_F16, GGML_TYPE_F32}) {
2180-
for (ggml_type type_dst : {GGML_TYPE_F16, GGML_TYPE_F32}) {
2181-
test_cases.emplace_back(new test_cpy(type_src, type_dst, {256, 2, 3, 4}, {1, 0, 2, 3})); // cpy not-contiguous
2191+
for (ggml_type type_src : {GGML_TYPE_F16}) {
2192+
for (ggml_type type_dst : {GGML_TYPE_Q4_0}) {
2193+
// test_cases.emplace_back(new test_cpy(type_src, type_dst, {256, 4, 4, 4}));
2194+
// test_cases.emplace_back(new test_cpy(type_src, type_dst, {256, 1, 1, 1}));
2195+
test_cases.emplace_back(new test_cpy(type_src, type_dst, {32, 4, 1, 1}));
2196+
// test_cases.emplace_back(new test_cpy(type_src, type_dst, {256, 2, 3, 4}, {0, 2, 1, 3})); // cpy by rows
21822197
}
21832198
}
2199+
// for (ggml_type type_src : {GGML_TYPE_F16, GGML_TYPE_F32}) {
2200+
// for (ggml_type type_dst : {GGML_TYPE_F16, GGML_TYPE_F32}) {
2201+
// test_cases.emplace_back(new test_cpy(type_src, type_dst, {256, 2, 3, 4}, {1, 0, 2, 3})); // cpy not-contiguous
2202+
// }
2203+
// }
21842204

21852205
test_cases.emplace_back(new test_cont());
21862206

0 commit comments

Comments
 (0)