Hongqiang Wang
1548a240e3
* opencl: add extended elementwise unary ops (sgn, step, elu, hardswish, hardsigmoid, floor, ceil, round, trunc) Adds nine GGML_UNARY_OP_* elementwise ops that were falling back to CPU on the OpenCL backend, following the same variant shape as the existing ABS op: f32, f32_4 (vec4), f16, f16_4 (vec4), and stride-addressed f32_nc / f16_nc for non-contiguous inputs. New kernels/unary_ext.cl (macro-generated), a shared ggml_cl_unary_ext dispatch helper mirroring ggml_cl_abs, the supports_op cases, and the compute-forward cases. Values are computed in float (the f16 variants read/write half and convert), so the conditional ops (step, elu) match the CPU reference; the vec4 forms use select() for the branch. Validated with test-backend-ops on Adreno 840 and 850 (E17): all nine ops pass every case including the vec4 and non-contiguous variants (8/8 or 14/14). * opencl: dispatch a contiguous f32 copy over the whole device kernel_cpy_f32_f32 maps one workgroup to each (i01,i02,i03) row and strides the row across that workgroup's lanes, and the host launches ne01*MIN(64,ne00) work items. A tensor with few long rows therefore runs on a single workgroup. The mamba2 and gated-delta-net recurrent state cache is one row of 524288 floats, copied once per layer per graph, and lands on 64 work items. When both sides are contiguous the copy is a linear move, so dispatch it over the whole device: one work item per float4. Gated on ggml_is_contiguous for both tensors and equal element counts, so copies already spread over many rows keep the existing path. The kernel is created optionally, so a driver that rejects it falls back rather than aborting. vload4/vstore4 rather than a float4 cast: they require only the scalar type's alignment, and these buffers carry an arbitrary 4-byte view offset. CPY, DUP and CONT are 217/217 on Adreno 840 and 740 with the path enabled and disabled. GGML_OPENCL_CPY_FLAT=0 forces the old kernel. * opencl: support all easy-copy types in CONCAT CONCAT was F32-only. Extend it to every "easy-copy" type -- any non-quantized type with a block size of 1 and an element size of 1, 2, 4 or 8 bytes, i.e. f16/bf16/i8/i16/i32/i64 as well as f32. The kernels are keyed by element SIZE rather than by type, which is what CUDA already does for the same op: one kernel per byte width (b1/b2/b4/b8) plus the packed b4 fast path, instead of one per ggml type. supports_op gates on the same property, so a new type of a supported width is picked up with no further work. Validated with test-backend-ops on Adreno 840 / A8X and X2-90 / X2E.
314 lines
8.6 KiB
Common Lisp
314 lines
8.6 KiB
Common Lisp
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
|
|
|
//------------------------------------------------------------------------------
|
|
// cpy
|
|
//------------------------------------------------------------------------------
|
|
|
|
kernel void kernel_cpy_f16_f16(
|
|
global half * src0,
|
|
ulong offset0,
|
|
global half * dst,
|
|
ulong offsetd,
|
|
int ne00,
|
|
int ne01,
|
|
int ne02,
|
|
int ne03,
|
|
ulong nb00,
|
|
ulong nb01,
|
|
ulong nb02,
|
|
ulong nb03,
|
|
int ne0,
|
|
int ne1,
|
|
int ne2,
|
|
int ne3,
|
|
ulong nb0,
|
|
ulong nb1,
|
|
ulong nb2,
|
|
ulong nb3
|
|
) {
|
|
src0 = (global half*)((global char*)src0 + offset0);
|
|
dst = (global half*)((global char*)dst + offsetd);
|
|
|
|
int i03 = get_group_id(2);
|
|
int i02 = get_group_id(1);
|
|
int i01 = get_group_id(0);
|
|
|
|
int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
|
|
|
|
int i3 = n / (ne2*ne1*ne0);
|
|
int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
|
|
int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
|
|
int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
|
|
|
|
global half * dst_data = (global half *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
|
|
|
|
for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
|
|
global const half * src = (global half *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
|
|
dst_data[i00] = src[0];
|
|
}
|
|
}
|
|
|
|
kernel void kernel_cpy_f16_f32(
|
|
global half * src0,
|
|
ulong offset0,
|
|
global float * dst,
|
|
ulong offsetd,
|
|
int ne00,
|
|
int ne01,
|
|
int ne02,
|
|
int ne03,
|
|
ulong nb00,
|
|
ulong nb01,
|
|
ulong nb02,
|
|
ulong nb03,
|
|
int ne0,
|
|
int ne1,
|
|
int ne2,
|
|
int ne3,
|
|
ulong nb0,
|
|
ulong nb1,
|
|
ulong nb2,
|
|
ulong nb3
|
|
) {
|
|
|
|
src0 = (global half*)((global char*)src0 + offset0);
|
|
dst = (global float*)((global char*)dst + offsetd);
|
|
|
|
int i03 = get_group_id(2);
|
|
int i02 = get_group_id(1);
|
|
int i01 = get_group_id(0);
|
|
|
|
int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
|
|
|
|
int i3 = n / (ne2*ne1*ne0);
|
|
int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
|
|
int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
|
|
int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
|
|
|
|
global float * dst_data = (global float *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
|
|
|
|
for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
|
|
global half * src = (global half *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
|
|
dst_data[i00] = src[0];
|
|
}
|
|
}
|
|
|
|
kernel void kernel_cpy_f32_f16(
|
|
global float * src0,
|
|
ulong offset0,
|
|
global half * dst,
|
|
ulong offsetd,
|
|
int ne00,
|
|
int ne01,
|
|
int ne02,
|
|
int ne03,
|
|
ulong nb00,
|
|
ulong nb01,
|
|
ulong nb02,
|
|
ulong nb03,
|
|
int ne0,
|
|
int ne1,
|
|
int ne2,
|
|
int ne3,
|
|
ulong nb0,
|
|
ulong nb1,
|
|
ulong nb2,
|
|
ulong nb3
|
|
) {
|
|
src0 = (global float*)((global char*)src0 + offset0);
|
|
dst = (global half*)((global char*)dst + offsetd);
|
|
|
|
int i03 = get_group_id(2);
|
|
int i02 = get_group_id(1);
|
|
int i01 = get_group_id(0);
|
|
|
|
int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
|
|
|
|
int i3 = n / (ne2*ne1*ne0);
|
|
int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
|
|
int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
|
|
int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
|
|
|
|
global half * dst_data = (global half *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
|
|
|
|
for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
|
|
global const float * src = (global float *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
|
|
|
|
dst_data[i00] = src[0];
|
|
}
|
|
}
|
|
|
|
kernel void kernel_cpy_f32_f32(
|
|
global float * src0,
|
|
ulong offset0,
|
|
global float * dst,
|
|
ulong offsetd,
|
|
int ne00,
|
|
int ne01,
|
|
int ne02,
|
|
int ne03,
|
|
ulong nb00,
|
|
ulong nb01,
|
|
ulong nb02,
|
|
ulong nb03,
|
|
int ne0,
|
|
int ne1,
|
|
int ne2,
|
|
int ne3,
|
|
ulong nb0,
|
|
ulong nb1,
|
|
ulong nb2,
|
|
ulong nb3
|
|
) {
|
|
src0 = (global float*)((global char*)src0 + offset0);
|
|
dst = (global float*)((global char*)dst + offsetd);
|
|
|
|
int i03 = get_group_id(2);
|
|
int i02 = get_group_id(1);
|
|
int i01 = get_group_id(0);
|
|
|
|
int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
|
|
|
|
int i3 = n / (ne2*ne1*ne0);
|
|
int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
|
|
int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
|
|
int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
|
|
|
|
global float * dst_data = (global float *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
|
|
|
|
for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
|
|
global const float * src = (global float *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
|
|
|
|
dst_data[i00] = src[0];
|
|
}
|
|
}
|
|
|
|
kernel void kernel_cpy_f32_f32_pack(
|
|
global float * src0,
|
|
ulong offset0,
|
|
global float * dst,
|
|
ulong offsetd,
|
|
int ne00,
|
|
int ne01,
|
|
int ne02,
|
|
int ne03,
|
|
ulong nb00,
|
|
ulong nb01,
|
|
ulong nb02,
|
|
ulong nb03,
|
|
int ne0,
|
|
int ne1,
|
|
int ne2,
|
|
int ne3,
|
|
ulong nb0,
|
|
ulong nb1,
|
|
ulong nb2,
|
|
ulong nb3
|
|
) {
|
|
src0 = (global float*)((global char*)src0 + offset0);
|
|
dst = (global float*)((global char*)dst + offsetd);
|
|
|
|
int lsz = get_local_size(0);
|
|
int tpr = min(ne00, lsz); // threads per row
|
|
int rpw = lsz / tpr; // rows per workgroup
|
|
int lid = get_local_id(0);
|
|
int row = get_group_id(0)*rpw + lid / tpr;
|
|
int lane = lid - (lid / tpr) * tpr;
|
|
|
|
int nrows = ne01*ne02*ne03;
|
|
if (row >= nrows) {
|
|
return;
|
|
}
|
|
|
|
int i01 = row % ne01;
|
|
int t = row / ne01;
|
|
int i02 = t % ne02;
|
|
int i03 = t / ne02;
|
|
|
|
// linear index of the first element of this row, unflattened over dst dims
|
|
long n = (long)row * ne00;
|
|
int i3 = (int)(n / ((long)ne2*ne1*ne0));
|
|
long rm = n - (long)i3*ne2*ne1*ne0;
|
|
int i2 = (int)(rm / ((long)ne1*ne0));
|
|
rm -= (long)i2*ne1*ne0;
|
|
int i1 = (int)(rm / ne0);
|
|
int i0 = (int)(rm - (long)i1*ne0);
|
|
|
|
global float * dst_data = (global float *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
|
|
|
|
for (int i00 = lane; i00 < ne00; i00 += tpr) {
|
|
global const float * src = (global float *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
|
|
dst_data[i00] = src[0];
|
|
}
|
|
}
|
|
|
|
kernel void kernel_cpy_i32_i32(
|
|
global int * src0,
|
|
ulong offset0,
|
|
global int * dst,
|
|
ulong offsetd,
|
|
int ne00,
|
|
int ne01,
|
|
int ne02,
|
|
int ne03,
|
|
ulong nb00,
|
|
ulong nb01,
|
|
ulong nb02,
|
|
ulong nb03,
|
|
int ne0,
|
|
int ne1,
|
|
int ne2,
|
|
int ne3,
|
|
ulong nb0,
|
|
ulong nb1,
|
|
ulong nb2,
|
|
ulong nb3
|
|
) {
|
|
src0 = (global int*)((global char*)src0 + offset0);
|
|
dst = (global int*)((global char*)dst + offsetd);
|
|
|
|
int i03 = get_group_id(2);
|
|
int i02 = get_group_id(1);
|
|
int i01 = get_group_id(0);
|
|
|
|
int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
|
|
|
|
int i3 = n / (ne2*ne1*ne0);
|
|
int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
|
|
int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
|
|
int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
|
|
|
|
global int * dst_data = (global int *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
|
|
|
|
for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
|
|
global const int * src = (global int *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
|
|
|
|
dst_data[i00] = src[0];
|
|
}
|
|
}
|
|
|
|
// Contiguous f32 copy, one work item per float4 over the whole tensor. The kernels above map
|
|
// one workgroup to each row, which leaves a tensor with few long rows on a single compute unit.
|
|
// vload4/vstore4 rather than a float4 cast: these buffers carry an arbitrary 4-byte view offset.
|
|
kernel void kernel_cpy_f32_f32_flat(
|
|
global float * src0,
|
|
ulong offset0,
|
|
global float * dst,
|
|
ulong offsetd,
|
|
ulong ne,
|
|
ulong n4
|
|
) {
|
|
src0 = (global float*)((global char*)src0 + offset0);
|
|
dst = (global float*)((global char*)dst + offsetd);
|
|
|
|
const ulong i = get_global_id(0);
|
|
|
|
if (i < n4) {
|
|
vstore4(vload4(i, src0), i, dst);
|
|
} else if (i == n4) {
|
|
for (ulong t = n4 * 4; t < ne; ++t) {
|
|
dst[t] = src0[t];
|
|
}
|
|
}
|
|
}
|