/*******************************************************************************
* 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/include/post_ops.h"
#include "gpu/intel/include/types.h"
void get_strides(int mask, long dim0, long dim1, long dim2, long *str0,
long *str1, long *str2) {
int ndims = dim0 > 1 ? 3 : 2 long dims[3] dims[0] = (ndims > 2 && mask & (1 << 0)) ? dim0 : 1 dims[1] = mask & (1 << (ndims - 1)) ? dim1 : 1 dims[2] = mask & (1 << (ndims - 2)) ? dim2 : 1
*str0 = dims[0] == 1 ? 0 : dims[1] * dims[2] *str1 = dims[1] == 1 ? 0 : 1 *str2 = dims[2] == 1 ? 0 : dims[1]}
__kernel void ref_gemm(__global A_DATA_T *a, __global B_DATA_T *b,
__global C_DATA_T *c, __global BIA_DATA_T *bias, long offset_a0,
long offset_b0, long offset_c0, long offset_bias0, int transa,
int transb, long MB, long M, long N, long K, long stride_a_mb,
long stride_b_mb, long stride_c, long lda, long ldb, long ldc,
float eltwise_alpha, float eltwise_beta, float eltwise_scale,
int bias_mask,
#if WITH_HOST_WEI_ZP
int ao_value,
#else
__global int *ao,
#endif
#if WITH_HOST_SRC_ZP
int bo_value,
#else
__global int *bo,
#endif
#if WITH_HOST_DST_ZP
int c0_value,
#else
__global int *c0,
#endif
int c0_mask, float beta) {
// See note in src/gpu/intel/gemm/primitive.hpp wrt args swap
#if WITH_WEI_ZPOINTS
#define ATTR_A0 ao[0]
#else
#define ATTR_A0 0
#endif
#if WITH_SRC_ZPOINTS
#define ATTR_B0 bo[0]
#else
#define ATTR_B0 0
#endif
#if WITH_HOST_SRC_ZP
int *bo = &bo_value#endif
#if WITH_HOST_WEI_ZP
int *ao = &ao_value#endif
#if WITH_HOST_DST_ZP
int *c0 = &c0_value#endif
long n = get_global_id(0) long m = get_global_id(1) long mb = get_global_id(2)
#if WITH_BIAS
bias += offset_bias0
long b_strides[3] get_strides(
bias_mask, MB, M, N, &b_strides[0], &b_strides[1], &b_strides[2])#endif
a += offset_a0 b += offset_b0 c += offset_c0
#if WITH_DST_ZPOINTS
long c0_strides[3] get_strides(
c0_mask, MB, M, N, &c0_strides[0], &c0_strides[1], &c0_strides[2])#endif
long stride_a_m = transa ? lda : 1 long stride_a_k = transa ? 1 : lda long stride_b_k = transb ? ldb : 1 long stride_b_n = transb ? 1 : ldb
ACC_DATA_T acc = 0 for (long k = 0 long off_a = mb * stride_a_mb + m * stride_a_m + k * stride_a_k long off_b = mb * stride_b_mb + k * stride_b_k + n * stride_b_n acc += TO_ACC(A_TO_REF(a[off_a]) - ATTR_A0)
* TO_ACC(B_TO_REF(b[off_b]) - ATTR_B0) }
long off_c = mb * stride_c + n * ldc + m#if WITH_BIAS || NON_DEFAULT_ATTRS
POST_OP_DATA_T temp = (POST_OP_DATA_T)acc#if WITH_BIAS
long off_bias = mb * b_strides[0] + m * b_strides[1] + n * b_strides[2] temp += BIA_TO_REF(bias[off_bias])#endif
#if WITH_POST_OP
#if WITH_SUM
temp += (POST_OP_DATA_T)(beta * C_TO_REF(c[off_c]))#endif
temp = fwd_eltwise(temp, eltwise_alpha, eltwise_beta, eltwise_scale)#endif
#if WITH_DST_ZPOINTS
long off_c0 = mb * c0_strides[0] + m * c0_strides[1] + n * c0_strides[2] temp += c0[off_c0]#endif
c[off_c] = TO_C(temp)#else
c[off_c] = TO_C(acc)#endif
}