onednn-src 0.1.13

/*******************************************************************************
* Copyright 2019 Intel Corporation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
*     http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*******************************************************************************/

#include "gpu/intel/binary/common.h"

#if IS_TENSOR_OP && IS_DENSE && IS_SAME_MD && !WITH_BINARY_POST_OP
KERNEL_ATTR
__kernel void simple_binary(__global DATA_T *src0, __global DATA_T *src1,
#if IS_TERNARY
        __global char *src2,
#endif
        __global DST_DATA_T *dst POST_OP_ARGS, __global float *src0_scale,
        __global float *src1_scale) {
    off_t off = GWS_GET_IDX();

    float tmp_src0 = SRC0_TO_FLOAT(src0[off]);
    float tmp_src1 = SRC1_TO_FLOAT(src1[off]);
#if IS_TERNARY
    char tmp_src2 = src2[off];
#endif
    float d = 0;

#if WITH_SRC0_SCALE
    tmp_src0 = tmp_src0 * (*src0_scale);
#endif
#if WITH_SRC1_SCALE
    tmp_src1 = tmp_src1 * (*src1_scale);
#endif
#if IS_TERNARY
    d = ternary_op(BINARY_ALG, tmp_src0, tmp_src1, tmp_src2);
#else
    d = binary_op(BINARY_ALG, tmp_src0, tmp_src1);
#endif

    float dst_data;
#if WITH_SUM
    dst_data = CONVERT_FLOAT_T(dst[off]);
#endif

    APPLY_POST_OPS_SERIAL(d, dst_data, 0, 0, 0, 0, 0, 0);
    dst[off] = TO_DST(d);
}
#else
KERNEL_ATTR
__kernel void simple_binary(__global SRC0_DATA_T *src0,
        __global SRC1_DATA_T *src1,
#if IS_TERNARY
        __global char *src2,
#endif
        __global DST_DATA_T *dst POST_OP_ARGS, __global float *src0_scale,
        __global float *src1_scale) {

    // since gws = no. of total elems in A, id will be the logical offset
    off_t dims0[6] = {0};
    dims0[0] = GWS_GET_D0();
    dims0[1] = GWS_GET_D1();
    dims0[2] = GWS_GET_D2();
    dims0[3] = GWS_GET_D3();
    dims0[4] = GWS_GET_D4();
    dims0[5] = GWS_GET_D5();
    int d1_block = GWS_GET_D1_BLOCK();
    off_t dims0_po[6]
            = {dims0[0], dims0[1], dims0[2], dims0[3], dims0[4], dims0[5]};
    off_t d1_init = GWS_GET_D1();

    off_t dst_off = DST_OFF(
            dims0[0], dims0[1], dims0[2], dims0[3], dims0[4], dims0[5]);
#if TENSOR_OP
    off_t src0_off = SRC0_OFF(
            dims0[0], dims0[1], dims0[2], dims0[3], dims0[4], dims0[5]);
    off_t src1_off = SRC1_OFF(
            dims0[0], dims0[1], dims0[2], dims0[3], dims0[4], dims0[5]);

#if IS_TERNARY
    off_t src2_off = SRC2_OFF(
            dims0[0], dims0[1], dims0[2], dims0[3], dims0[4], dims0[5]);
#endif

#else
    off_t src0_off
            = SRC0_OFF(dims0[0] * !SRC0_BCAST_DIM0, dims0[1] * !SRC0_BCAST_DIM1,
                    dims0[2] * !SRC0_BCAST_DIM2, dims0[3] * !SRC0_BCAST_DIM3,
                    dims0[4] * !SRC0_BCAST_DIM4, dims0[5] * !SRC0_BCAST_DIM5);
    off_t src1_off
            = SRC1_OFF(dims0[0] * !SRC1_BCAST_DIM0, dims0[1] * !SRC1_BCAST_DIM1,
                    dims0[2] * !SRC1_BCAST_DIM2, dims0[3] * !SRC1_BCAST_DIM3,
                    dims0[4] * !SRC1_BCAST_DIM4, dims0[5] * !SRC1_BCAST_DIM5);
#if IS_TERNARY
    off_t src2_off
            = SRC2_OFF(dims0[0] * !SRC2_BCAST_DIM0, dims0[1] * !SRC2_BCAST_DIM1,
                    dims0[2] * !SRC2_BCAST_DIM2, dims0[3] * !SRC2_BCAST_DIM3,
                    dims0[4] * !SRC2_BCAST_DIM4, dims0[5] * !SRC2_BCAST_DIM5);
#endif

#endif

    // SRC1_D1 = IC for SRC1, using the dispatch ap
    int block_size = d1_block;

    if (dims0[0] >= DST_D0) {
        for (int ic = 0; ic < block_size; ++ic) {
            dst[dst_off] = TO_DST(0.0f);
            dst_off++;
        }

        return;
    }

    if (d1_init + block_size <= DST_D1) {
        for (int ic = 0; ic < block_size; ++ic) {
            float tmp_src0 = SRC0_TO_FLOAT(src0[src0_off]);
            float tmp_src1 = SRC1_TO_FLOAT(src1[src1_off]);
#if IS_TERNARY
            char tmp_src2 = src2[src2_off];
#endif
            float d = 0;

#if WITH_SRC0_SCALE
            tmp_src0 = tmp_src0 * (*src0_scale);
#endif
#if WITH_SRC1_SCALE
            tmp_src1 = tmp_src1 * (*src1_scale);
#endif

#if IS_TERNARY
            d = ternary_op(BINARY_ALG, tmp_src0, tmp_src1, tmp_src2);
#else
            d = binary_op(BINARY_ALG, tmp_src0, tmp_src1);
#endif

            float dst_data;
#if WITH_SUM
            dst_data = CONVERT_FLOAT_T(dst[dst_off]);
#endif
            APPLY_POST_OPS_SERIAL(d, dst_data, dims0_po[0], dims0_po[1],
                    dims0_po[2], dims0_po[3], dims0_po[4], dims0_po[5]);

            dst[dst_off] = TO_DST(d);

#if USE_UNROLL_16B || SRC0_UNROLL_16B
            src0_off++;
            dst_off++;
            ++dims0_po[1];
            if (USE_UNROLL_16B && (SRC1_D1 > 1)) {
                src1_off++;
            } else if (SRC0_UNROLL_16B && (SRC1_D1 > 1)) {
                src1_off += SRC1_S1_0; // Equilvalent stride in plain format
            }
#endif
        }
    } else {
        for (int ic = 0; ic < DST_D1 - d1_init; ic++) {
            float tmp_src0 = SRC0_TO_FLOAT(src0[src0_off]);
            float tmp_src1 = SRC1_TO_FLOAT(src1[src1_off]);
#if IS_TERNARY
            char tmp_src2 = src2[src2_off];
#endif
            float d = 0;

#if WITH_SRC0_SCALE
            tmp_src0 = tmp_src0 * (*src0_scale);
#endif
#if WITH_SRC1_SCALE
            tmp_src1 = tmp_src1 * (*src1_scale);
#endif

#if IS_TERNARY
            d = ternary_op(BINARY_ALG, tmp_src0, tmp_src1, tmp_src2);
#else
            d = binary_op(BINARY_ALG, tmp_src0, tmp_src1);
#endif

            float dst_data;
#if WITH_SUM
            dst_data = CONVERT_FLOAT_T(dst[dst_off]);
#endif
            APPLY_POST_OPS_SERIAL(d, dst_data, dims0_po[0], dims0_po[1],
                    dims0_po[2], dims0_po[3], dims0_po[4], dims0_po[5]);

            dst[dst_off] = TO_DST(d);

#if USE_UNROLL_16B || SRC0_UNROLL_16B
            src0_off++;
            dst_off++;
            ++dims0_po[1];
            if (USE_UNROLL_16B && (SRC1_D1 > 1)) {
                src1_off++;
            } else if (SRC0_UNROLL_16B && (SRC1_D1 > 1)) {
                src1_off += SRC1_S1_0; // Equilvalent stride in plain format
            }
#endif
        }
#if DST_D1 != DST_PD1
        for (int ic = 0; ic < min(DST_PD1 - DST_D1, block_size); ic++) {
            dst[dst_off] = TO_DST(0.0f);
            dst_off++;
        }
#endif
    }
}
#endif