Skip to content

Commit 3e841d8

Browse files
lhezshawngu-quicquic-sszot
authored and
mglambda
committed
opencl: fix for small models (ggml-org#11950)
* opencl: fix small shape gemv, remove unused extensions * opencl: fix `transpose_16`, `dump_tensor`, enforce subgroup size * opencl: fix for token length < 4 * opencl: use wave size of 64 for all Adreno GPUs --------- Co-authored-by: Shawn Gu <[email protected]> Co-authored-by: Skyler Szot <[email protected]>
1 parent 3e1a083 commit 3e841d8

File tree

6 files changed

+67
-59
lines changed

6 files changed

+67
-59
lines changed

‎ggml/src/ggml-opencl/ggml-opencl.cpp‎

Lines changed: 25 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -444,19 +444,8 @@ static ggml_backend_opencl_context * ggml_cl2_init(ggml_backend_dev_t dev) {
444444
backend_ctx->gpu_family = GPU_FAMILY::ADRENO;
445445
backend_ctx->adreno_gen = get_adreno_gpu_gen(default_device->name);
446446

447-
// Default wave size is 128, A8x uses 64.
448-
if (backend_ctx->adreno_gen == ADRENO_GPU_GEN::A8X) {
449-
backend_ctx->adreno_wave_size = 64;
450-
} else if (backend_ctx->adreno_gen == ADRENO_GPU_GEN::A7X ||
451-
backend_ctx->adreno_gen == ADRENO_GPU_GEN::X1E) {
452-
backend_ctx->adreno_wave_size = 128;
453-
} else {
454-
backend_ctx->adreno_wave_size = 128;
455-
GGML_LOG_WARN("ggml_opencl: Unsupported Adreno GPU: %s, "
456-
"using wave size %d, "
457-
"may not work as expected\n",
458-
backend_ctx->device_name.c_str(), backend_ctx->adreno_wave_size);
459-
}
447+
// Use wave size of 64 for all Adreno GPUs.
448+
backend_ctx->adreno_wave_size = 64;
460449
} else if (strstr(default_device->name, "Intel")) {
461450
backend_ctx->gpu_family = GPU_FAMILY::INTEL;
462451
} else {
@@ -1376,6 +1365,11 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
13761365
int M = tensor->ne[1]; // ne01
13771366
int K = tensor->ne[0]; // ne00
13781367

1368+
//For matrix-vector multiplication kernel, we assume K is a multiple of 32
1369+
GGML_ASSERT(K % 32 == 0);
1370+
//For transpose kernels, we assume K is a multiple of 4 (satisfied by prior assert), and M is a multiple of 4
1371+
GGML_ASSERT(M % 4 == 0);
1372+
13791373
// transpose is out of place, so we need to allocate transposed buffers
13801374
// <----------------------------------------------------------------------------------> //
13811375
// use sub_buffer of max buffer size instead
@@ -1416,36 +1410,36 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
14161410
cl_mem qT_d_image1D;
14171411
cl_mem dT_d_image1D;
14181412

1419-
cl_image_format img_fmt_1d = { CL_RGBA, CL_FLOAT };
1413+
cl_image_format img_fmt_1d = { CL_RGBA, CL_HALF_FLOAT };
14201414
cl_image_desc img_desc_1d;
14211415

14221416
memset(&img_desc_1d, 0, sizeof(img_desc_1d));
14231417
img_desc_1d.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
1424-
img_desc_1d.image_width = M * K / 8 / 4;
1418+
img_desc_1d.image_width = M * K / 4 / 4;
14251419
img_desc_1d.buffer = extra->q;
14261420
q_d_image1D = clCreateImage(context, 0, &img_fmt_1d, &img_desc_1d, NULL, &err);
14271421
CL_CHECK(err);
14281422

1429-
img_fmt_1d = { CL_RGBA, CL_FLOAT };
1423+
img_fmt_1d = { CL_RGBA, CL_HALF_FLOAT };
14301424
memset(&img_desc_1d, 0, sizeof(img_desc_1d));
14311425
img_desc_1d.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
1432-
img_desc_1d.image_width = M * K / 8 / 4;
1426+
img_desc_1d.image_width = M * K / 4 / 4;
14331427
img_desc_1d.buffer = qT_d;
14341428
qT_d_image1D = clCreateImage(context, 0, &img_fmt_1d, &img_desc_1d, NULL, &err);
14351429
CL_CHECK(err);
14361430

1437-
img_fmt_1d = { CL_RGBA, CL_FLOAT };
1431+
img_fmt_1d = { CL_RGBA, CL_HALF_FLOAT };
14381432
memset(&img_desc_1d, 0, sizeof(img_desc_1d));
14391433
img_desc_1d.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
1440-
img_desc_1d.image_width = M * K / 32 / 4 / 2;
1434+
img_desc_1d.image_width = M * K / 32 / 4;
14411435
img_desc_1d.buffer = extra->d;
14421436
d_d_image1D = clCreateImage(context, 0, &img_fmt_1d, &img_desc_1d, NULL, &err);
14431437
CL_CHECK(err);
14441438

1445-
img_fmt_1d = { CL_RGBA, CL_FLOAT };
1439+
img_fmt_1d = { CL_RGBA, CL_HALF_FLOAT };
14461440
memset(&img_desc_1d, 0, sizeof(img_desc_1d));
14471441
img_desc_1d.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
1448-
img_desc_1d.image_width = M * K / 32 / 4 / 2;
1442+
img_desc_1d.image_width = M * K / 32 / 4;
14491443
img_desc_1d.buffer = dT_d;
14501444
dT_d_image1D = clCreateImage(context, 0, &img_fmt_1d, &img_desc_1d, NULL, &err);
14511445
CL_CHECK(err);
@@ -1454,8 +1448,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
14541448
// set up and call the transpose kernels
14551449
// <----------------------------------------------------------------------------------> //
14561450
// weights
1457-
int height_q = M / 8;
1458-
int width_q = K / 8 / 4;
1451+
int height_q = M / 4;
1452+
int width_q = K / 4 / 4;
14591453
kernel = backend_ctx->kernel_transpose_16;
14601454

14611455
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &q_d_image1D));
@@ -1469,8 +1463,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
14691463
CL_CHECK(clWaitForEvents(1, &evt));
14701464

14711465
// scales
1472-
int height_s = M / 8;
1473-
int width_s = K / 32 / 8;
1466+
int height_s = M / 4;
1467+
int width_s = K / 32 / 4;
14741468

14751469
kernel = backend_ctx->kernel_transpose_16;
14761470
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_d_image1D));
@@ -1864,7 +1858,6 @@ static void dump_tensor(ggml_backend_t backend, const struct ggml_tensor * tenso
18641858
void * buf_d;
18651859
#endif
18661860

1867-
#ifdef GGML_USE_OPENCL
18681861
// Make sure everything is done.
18691862
CL_CHECK(clFinish(queue));
18701863

@@ -1900,7 +1893,6 @@ static void dump_tensor(ggml_backend_t backend, const struct ggml_tensor * tenso
19001893
extra->offset, ggml_nbytes(tensor), buf, 0, NULL, NULL));
19011894
CL_CHECK(clFinish(queue));
19021895
#endif // GGML_OPENCL_SOA_Q
1903-
#endif // GGML_USE_OPENCL
19041896

19051897
// Open file and dump.
19061898
char fname[512];
@@ -2865,6 +2857,9 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
28652857
CL_CHECK(status);
28662858

28672859
int height_B = N/4;
2860+
if (height_B == 0) {
2861+
height_B = 1;
2862+
}
28682863
int width_B = K/4;
28692864
int padded_height_B = (N + padding)/4;
28702865

@@ -3013,11 +3008,12 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
30133008
}
30143009

30153010
if (N == 1) {
3016-
local_work_size[0] = backend_ctx->adreno_wave_size; // localsize
3011+
size_t wavesize = backend_ctx->adreno_wave_size;
3012+
local_work_size[0] = wavesize; // localsize
30173013
local_work_size[1] = 4; // reduce factor
30183014
local_work_size[2] = 1;
30193015

3020-
global_work_size[0] = M / 2;
3016+
global_work_size[0] = (((M / 2) + wavesize - 1) / wavesize) * wavesize;
30213017
global_work_size[1] = 4; // reduce factor
30223018
global_work_size[2] = 1;
30233019
}

‎ggml/src/ggml-opencl/kernels/ggml-opencl.cl‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1797,6 +1797,9 @@ kernel void kernel_mul_mat_f16_f16(
17971797
//------------------------------------------------------------------------------
17981798
// mul_mat_f16_f32_1row
17991799
//------------------------------------------------------------------------------
1800+
#ifdef ADRENO_GPU
1801+
REQD_SUBGROUP_SIZE_64
1802+
#endif
18001803
kernel void kernel_mul_mat_f16_f32_1row(
18011804
global char * src0,
18021805
ulong offset0,

‎ggml/src/ggml-opencl/kernels/ggml-opencl_gemv_noshuffle.cl‎

Lines changed: 8 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,11 @@
11
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
22
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
3-
#pragma OPENCL EXTENSION cl_qcom_subgroup_uniform_load: enable
4-
#pragma OPENCL EXTENSION cl_qcom_subgroup_constant_load: enable
5-
#pragma OPENCL EXTENSION cl_qcom_extra_vector_types : enable
3+
4+
#ifdef cl_qcom_reqd_sub_group_size
65
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
6+
#define ADRENO_GPU 1
7+
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
8+
#endif
79

810
// assume
911
#define QK4_0 32
@@ -186,8 +188,9 @@
186188
total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
187189
total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
188190

189-
190-
__attribute__((qcom_reqd_sub_group_size("full")))
191+
#ifdef ADRENO_GPU
192+
REQD_SUBGROUP_SIZE_64
193+
#endif
191194
__kernel void kernel_gemv_noshuffle(
192195
__read_only image1d_buffer_t src0_q, // quantized A
193196
global half2 * src0_d, // A scales

‎ggml/src/ggml-opencl/kernels/ggml-opencl_gemv_noshuffle_general.cl‎

Lines changed: 8 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,11 @@
11
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
22
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
3-
#pragma OPENCL EXTENSION cl_qcom_subgroup_uniform_load: enable
4-
#pragma OPENCL EXTENSION cl_qcom_subgroup_constant_load: enable
5-
#pragma OPENCL EXTENSION cl_qcom_extra_vector_types : enable
3+
4+
#ifdef cl_qcom_reqd_sub_group_size
65
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
6+
#define ADRENO_GPU 1
7+
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
8+
#endif
79

810
// assume
911
#define QK4_0 32
@@ -186,8 +188,9 @@
186188
total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
187189
total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
188190

189-
190-
__attribute__((qcom_reqd_sub_group_size("full")))
191+
#ifdef ADRENO_GPU
192+
REQD_SUBGROUP_SIZE_64
193+
#endif
191194
__kernel void kernel_gemv_noshuffle(
192195
__read_only image1d_buffer_t src0_q, // quantized A
193196
global half2 * src0_d, // A scales

‎ggml/src/ggml-opencl/kernels/ggml-opencl_mul_mat_Ab_Bi_8x4.cl‎

Lines changed: 10 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,16 @@
77
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
88
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
99

10-
__attribute__((qcom_reqd_sub_group_size("full")))
10+
#ifdef cl_qcom_reqd_sub_group_size
11+
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
12+
#define ADRENO_GPU 1
13+
#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
14+
#endif
15+
16+
#ifdef ADRENO_GPU
17+
REQD_SUBGROUP_SIZE_128
18+
#endif
19+
1120
kernel void kernel_mul_mat_Ab_Bi_8x4(
1221
global const ushort * src0_q, // quantized A
1322
global const half * src0_d, // A scales

‎ggml/src/ggml-opencl/kernels/ggml-opencl_transpose_16.cl‎

Lines changed: 13 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,6 @@
1-
// 16-bit transpose, loading/storing an 8x8 tile of elements
1+
// 16-bit transpose, loading/storing a 4x4 tile of elements
2+
3+
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
24

35
kernel void kernel_transpose_16(
46
__read_only image1d_buffer_t input,
@@ -9,24 +11,16 @@ kernel void kernel_transpose_16(
911

1012
const int i = get_global_id(0);
1113
const int j = get_global_id(1);
12-
const int i_3 = i<<3;
13-
const int j_3 = j<<3;
14+
const int i_2 = i<<2;
15+
const int j_2 = j<<2;
1416

15-
ushort8 temp0 = as_ushort8(read_imagef(input, (j_3+0)*cols+i));
16-
ushort8 temp1 = as_ushort8(read_imagef(input, (j_3+1)*cols+i));
17-
ushort8 temp2 = as_ushort8(read_imagef(input, (j_3+2)*cols+i));
18-
ushort8 temp3 = as_ushort8(read_imagef(input, (j_3+3)*cols+i));
19-
ushort8 temp4 = as_ushort8(read_imagef(input, (j_3+4)*cols+i));
20-
ushort8 temp5 = as_ushort8(read_imagef(input, (j_3+5)*cols+i));
21-
ushort8 temp6 = as_ushort8(read_imagef(input, (j_3+6)*cols+i));
22-
ushort8 temp7 = as_ushort8(read_imagef(input, (j_3+7)*cols+i));
17+
half4 temp0 = read_imageh(input, (j_2+0)*cols+i);
18+
half4 temp1 = read_imageh(input, (j_2+1)*cols+i);
19+
half4 temp2 = read_imageh(input, (j_2+2)*cols+i);
20+
half4 temp3 = read_imageh(input, (j_2+3)*cols+i);
2321

24-
write_imagef(output, (i_3+0)*rows+j, as_float4((ushort8)(temp0.s0, temp1.s0, temp2.s0, temp3.s0, temp4.s0, temp5.s0, temp6.s0, temp7.s0)));
25-
write_imagef(output, (i_3+1)*rows+j, as_float4((ushort8)(temp0.s1, temp1.s1, temp2.s1, temp3.s1, temp4.s1, temp5.s1, temp6.s1, temp7.s1)));
26-
write_imagef(output, (i_3+2)*rows+j, as_float4((ushort8)(temp0.s2, temp1.s2, temp2.s2, temp3.s2, temp4.s2, temp5.s2, temp6.s2, temp7.s2)));
27-
write_imagef(output, (i_3+3)*rows+j, as_float4((ushort8)(temp0.s3, temp1.s3, temp2.s3, temp3.s3, temp4.s3, temp5.s3, temp6.s3, temp7.s3)));
28-
write_imagef(output, (i_3+4)*rows+j, as_float4((ushort8)(temp0.s4, temp1.s4, temp2.s4, temp3.s4, temp4.s4, temp5.s4, temp6.s4, temp7.s4)));
29-
write_imagef(output, (i_3+5)*rows+j, as_float4((ushort8)(temp0.s5, temp1.s5, temp2.s5, temp3.s5, temp4.s5, temp5.s5, temp6.s5, temp7.s5)));
30-
write_imagef(output, (i_3+6)*rows+j, as_float4((ushort8)(temp0.s6, temp1.s6, temp2.s6, temp3.s6, temp4.s6, temp5.s6, temp6.s6, temp7.s6)));
31-
write_imagef(output, (i_3+7)*rows+j, as_float4((ushort8)(temp0.s7, temp1.s7, temp2.s7, temp3.s7, temp4.s7, temp5.s7, temp6.s7, temp7.s7)));
22+
write_imageh(output, (i_2+0)*rows+j, (half4)(temp0.s0, temp1.s0, temp2.s0, temp3.s0));
23+
write_imageh(output, (i_2+1)*rows+j, (half4)(temp0.s1, temp1.s1, temp2.s1, temp3.s1));
24+
write_imageh(output, (i_2+2)*rows+j, (half4)(temp0.s2, temp1.s2, temp2.s2, temp3.s2));
25+
write_imageh(output, (i_2+3)*rows+j, (half4)(temp0.s3, temp1.s3, temp2.s3, temp3.s3));
3226
}

0 commit comments

Comments
 (0)