| #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]; |
| } |
| } |
|
|
