| 1 | namespace dnnl { |
|---|---|
| 2 | namespace impl { |
| 3 | namespace gpu { |
| 4 | namespace ocl { |
| 5 | const char *xe_lp_conv_nhwc_fwd_first_x8s8x_kernel = R"==(/******************************************************************************* )==" "\n" |
| 6 | R"==(* Copyright 2020-2022 Intel Corporation )==" "\n" |
| 7 | R"==(* )==" "\n" |
| 8 | R"==(* Licensed under the Apache License, Version 2.0 (the "License"); )==" "\n" |
| 9 | R"==(* you may not use this file except in compliance with the License. )==" "\n" |
| 10 | R"==(* You may obtain a copy of the License at )==" "\n" |
| 11 | R"==(* )==" "\n" |
| 12 | R"==(* http: )==" "\n" |
| 13 | R"==(* )==" "\n" |
| 14 | R"==(* Unless required by applicable law or agreed to in writing, software )==" "\n" |
| 15 | R"==(* distributed under the License is distributed on an "AS IS" BASIS, )==" "\n" |
| 16 | R"==(* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. )==" "\n" |
| 17 | R"==(* See the License for the specific language governing permissions and )==" "\n" |
| 18 | R"==(* limitations under the License. )==" "\n" |
| 19 | R"==(*******************************************************************************/ )==" "\n" |
| 20 | R"==(#include "gpu/ocl/ocl_math_utils.h" )==" "\n" |
| 21 | R"==(#include "gpu/ocl/ocl_post_ops.h" )==" "\n" |
| 22 | R"==(#include "gpu/ocl/ocl_scales.h" )==" "\n" |
| 23 | R"==(#include "gpu/ocl/ocl_types.h" )==" "\n" |
| 24 | R"==(#define KDHW_SIZE (KH * KW * KD) )==" "\n" |
| 25 | R"==(#define CX1_OFFSET (sizeof(C00) / sizeof(int)) )==" "\n" |
| 26 | R"==(#if (PW % 4 == 0) && (SRC_SLM_SIZE % 4 == 0) )==" "\n" |
| 27 | R"==(#define WRITE_SLM_BLOCK(p, data) block_write(p, data) )==" "\n" |
| 28 | R"==(#define WRITE_SLM_BLOCK_SHORT(p, data) block_write_us(p, data) )==" "\n" |
| 29 | R"==(#else )==" "\n" |
| 30 | R"==(#define WRITE_SLM_BLOCK(p, data) block_write_emu(p, data) )==" "\n" |
| 31 | R"==(#define WRITE_SLM_BLOCK_SHORT(p, data) block_write_us_emu(p, data) )==" "\n" |
| 32 | R"==(#endif )==" "\n" |
| 33 | R"==(#define GET_INT_BLOCK(src_slm, slm_index, src_global) \ )==" "\n" |
| 34 | R"==(uchar4 res = 0; \ )==" "\n" |
| 35 | R"==(for (int j = 0; j < IC; j++) { \ )==" "\n" |
| 36 | R"==(((uchar *)&res)[j] = src_global[j]; \ )==" "\n" |
| 37 | R"==(} \ )==" "\n" |
| 38 | R"==(src_slm[slm_index] = as_int(res); \ )==" "\n" |
| 39 | R"==(src_global += IC; )==" "\n" |
| 40 | R"==(#define BLOCK_READ_SRC(data, idx) \ )==" "\n" |
| 41 | R"==(data = intel_sub_group_block_read8((__global uint *)&src[idx]); )==" "\n" |
| 42 | R"==(#define BLOCK_READ_WHT(data, idx) \ )==" "\n" |
| 43 | R"==(data = as_int(intel_sub_group_block_read((__global uint *)&wei[idx])); )==" "\n" |
| 44 | R"==(#define BLOCK_READ_WHT8(data, idx) \ )==" "\n" |
| 45 | R"==(data = as_int8(intel_sub_group_block_read8((__global uint *)&wei[idx])); )==" "\n" |
| 46 | R"==(#define BLOCK_READ_BIA(data, idx) \ )==" "\n" |
| 47 | R"==(data = as_float4(intel_sub_group_block_read4((__global uint *)&bias[idx])); )==" "\n" |
| 48 | R"==(#if SCALES_PER_OC )==" "\n" |
| 49 | R"==(#define SCALE_VEC4 scales.s01230123 )==" "\n" |
| 50 | R"==(#define SCALE scales )==" "\n" |
| 51 | R"==(#elif SCALES_COMMON )==" "\n" |
| 52 | R"==(#define SCALE_VEC4 runtime_scales[0] )==" "\n" |
| 53 | R"==(#define SCALE runtime_scales[0] )==" "\n" |
| 54 | R"==(#else )==" "\n" |
| 55 | R"==(#define SCALE_VEC4 1 )==" "\n" |
| 56 | R"==(#define SCALE 1 )==" "\n" |
| 57 | R"==(#endif )==" "\n" |
| 58 | R"==(#define OC_PADD8 ((OC % 8) ? (OC / 8 + 1) * 8 : OC) )==" "\n" |
| 59 | R"==(#if DST_DT_S8 || DST_DT_U8 )==" "\n" |
| 60 | R"==(#define OC_BLOCK_READ_BOUND 4 )==" "\n" |
| 61 | R"==(#define OC_BLOCK_WRITE_BOUND 16 )==" "\n" |
| 62 | R"==(#else )==" "\n" |
| 63 | R"==(#define OC_BLOCK_READ_BOUND 1 )==" "\n" |
| 64 | R"==(#define OC_BLOCK_WRITE_BOUND 4 )==" "\n" |
| 65 | R"==(#endif )==" "\n" |
| 66 | R"==(inline DST_DATA4_T read_oc_block4(const __global DST_DATA_T *dst, int off) { )==" "\n" |
| 67 | R"==(const int local_id = get_sub_group_local_id(); )==" "\n" |
| 68 | R"==(#if OC % OC_BLOCK != 0 )==" "\n" |
| 69 | R"==(int tail = OC - off; )==" "\n" |
| 70 | R"==(if (tail < OC_BLOCK) { )==" "\n" |
| 71 | R"==(return (DST_DATA4_T)( )==" "\n" |
| 72 | R"==(local_id < tail - 8 * 0 ? dst[0 * 8 + local_id] : 0, )==" "\n" |
| 73 | R"==(local_id < tail - 8 * 1 ? dst[1 * 8 + local_id] : 0, )==" "\n" |
| 74 | R"==(local_id < tail - 8 * 2 ? dst[2 * 8 + local_id] : 0, )==" "\n" |
| 75 | R"==(local_id < tail - 8 * 3 ? dst[3 * 8 + local_id] : 0); )==" "\n" |
| 76 | R"==(} )==" "\n" |
| 77 | R"==(#endif )==" "\n" |
| 78 | R"==(#if OC % OC_BLOCK_READ_BOUND != 0 )==" "\n" |
| 79 | R"==(return (DST_DATA4_T)(dst[0 * 8 + local_id], dst[1 * 8 + local_id], )==" "\n" |
| 80 | R"==(dst[2 * 8 + local_id], dst[3 * 8 + local_id]); )==" "\n" |
| 81 | R"==(#else )==" "\n" |
| 82 | R"==(return BLOCK_READ_DST4(dst); )==" "\n" |
| 83 | R"==(#endif )==" "\n" |
| 84 | R"==(} )==" "\n" |
| 85 | R"==(inline void write_oc_block4( )==" "\n" |
| 86 | R"==(__global DST_DATA_T *dst, int off, DST_DATA4_T value) { )==" "\n" |
| 87 | R"==(const int local_id = get_sub_group_local_id(); )==" "\n" |
| 88 | R"==(#if OC % OC_BLOCK != 0 )==" "\n" |
| 89 | R"==(int tail = OC - off; )==" "\n" |
| 90 | R"==(if (tail < OC_BLOCK) { )==" "\n" |
| 91 | R"==(if (local_id < tail) dst[0 * 8 + local_id] = value.s0; )==" "\n" |
| 92 | R"==(if (local_id < tail - 8 * 1) dst[1 * 8 + local_id] = value.s1; )==" "\n" |
| 93 | R"==(if (local_id < tail - 8 * 2) dst[2 * 8 + local_id] = value.s2; )==" "\n" |
| 94 | R"==(if (local_id < tail - 8 * 3) dst[3 * 8 + local_id] = value.s3; )==" "\n" |
| 95 | R"==(return; )==" "\n" |
| 96 | R"==(} )==" "\n" |
| 97 | R"==(#endif )==" "\n" |
| 98 | R"==(#if OC % OC_BLOCK_WRITE_BOUND != 0 )==" "\n" |
| 99 | R"==(dst[0 * 8 + local_id] = value.s0; )==" "\n" |
| 100 | R"==(dst[1 * 8 + local_id] = value.s1; )==" "\n" |
| 101 | R"==(dst[2 * 8 + local_id] = value.s2; )==" "\n" |
| 102 | R"==(dst[3 * 8 + local_id] = value.s3; )==" "\n" |
| 103 | R"==(return; )==" "\n" |
| 104 | R"==(#else )==" "\n" |
| 105 | R"==(BLOCK_WRITE_DST4(dst, value); )==" "\n" |
| 106 | R"==(return; )==" "\n" |
| 107 | R"==(#endif )==" "\n" |
| 108 | R"==(} )==" "\n" |
| 109 | R"==(__attribute__((intel_reqd_sub_group_size(SUB_GROUP_SIZE))) )==" "\n" |
| 110 | R"==(__attribute__((reqd_work_group_size(LWS_0, LWS_1, LWS_2))) __kernel void )==" "\n" |
| 111 | R"==(conv_nhwc_fwd_first_x8s8x(const __global uchar *src, const __global char *wei, )==" "\n" |
| 112 | R"==(const __global float *bias, __global DST_DATA_T *dst POST_OP_ARGS, )==" "\n" |
| 113 | R"==(const __global float *runtime_scales, )==" "\n" |
| 114 | R"==(const __global int *src_compensation, const __global int *src_zpoints, )==" "\n" |
| 115 | R"==(const __global int *dst_compensation) { )==" "\n" |
| 116 | R"==(const int group_oc = get_group_id(0) * OC_GROUP; )==" "\n" |
| 117 | R"==(const int group_mb = get_group_id(2) * MB_GROUP; )==" "\n" |
| 118 | R"==(const int group_sp = get_group_id(1) * SP_GROUP; )==" "\n" |
| 119 | R"==(const int sub_group_id = get_sub_group_id(); )==" "\n" |
| 120 | R"==(const int sub_local_id = get_sub_group_local_id(); )==" "\n" |
| 121 | R"==(const int oc = (sub_group_id % OC_GROUP); )==" "\n" |
| 122 | R"==(const int sp = (sub_group_id / OC_GROUP); )==" "\n" |
| 123 | R"==(const int g = (group_oc + oc) / OC_NCHUNK; )==" "\n" |
| 124 | R"==(const int group_ic = IC_NCHUNK * g; )==" "\n" |
| 125 | R"==(const int god = group_sp / (OW_PADDED * OH); )==" "\n" |
| 126 | R"==(const int gohw = group_sp % (OW_PADDED * OH); )==" "\n" |
| 127 | R"==(const int goh = gohw / OW_PADDED; )==" "\n" |
| 128 | R"==(const int gow = OW_BLOCK * (gohw % OW_PADDED); )==" "\n" |
| 129 | R"==(const int gid = god * SD; )==" "\n" |
| 130 | R"==(const int gih = goh * SH; )==" "\n" |
| 131 | R"==(const int giw = gow * SW; )==" "\n" |
| 132 | R"==(const int local_ow = OW_BLOCK * sp; )==" "\n" |
| 133 | R"==(const int local_iw = local_ow * SW; )==" "\n" |
| 134 | R"==(const int od = god; )==" "\n" |
| 135 | R"==(const int ow = gow + local_ow; )==" "\n" |
| 136 | R"==(const int oh = goh; )==" "\n" |
| 137 | R"==(const int id = gid - PD; )==" "\n" |
| 138 | R"==(const int iw = giw + local_iw - PW; )==" "\n" |
| 139 | R"==(const int ih = gih - PH; )==" "\n" |
| 140 | R"==(__local uint S_slice[SRC_SLM_SIZE * KH * KD]; )==" "\n" |
| 141 | R"==(__local uint *S_part = S_slice + (sp * SW * OW_BLOCK + PW); )==" "\n" |
| 142 | R"==(__local SRC_MMAD_DATA_T *S_work = S_slice + (sp * SW * OW_BLOCK); )==" "\n" |
| 143 | R"==(#if DST_NHWC )==" "\n" |
| 144 | R"==(dst += group_mb * OD * OH * OW * G * OC; )==" "\n" |
| 145 | R"==(dst += (OW * OH * od + OW * oh + ow) * G * OC; )==" "\n" |
| 146 | R"==(dst += OC_BLOCK * (group_oc + oc); )==" "\n" |
| 147 | R"==(#else )==" "\n" |
| 148 | R"==(dst += OC_BLOCK * OD * OH * OW * MB_BLOCK * (group_oc + oc); )==" "\n" |
| 149 | R"==(dst += OC_BLOCK * OD * OH * OW * OC_NCHUNK * G * MB_BLOCK )==" "\n" |
| 150 | R"==(* (group_mb / MB_BLOCK); )==" "\n" |
| 151 | R"==(dst += OC_BLOCK * (group_mb % MB_BLOCK); )==" "\n" |
| 152 | R"==(dst += OC_BLOCK * MB_BLOCK * (OW * OH * od + OW * oh + ow); )==" "\n" |
| 153 | R"==(#endif )==" "\n" |
| 154 | R"==(src += (group_mb >= MB ? MB - 1 : group_mb) * ID * IH * IW * G * IC; )==" "\n" |
| 155 | R"==(src += (IW * IH * id + IW * ih + iw + PW) * G * IC; )==" "\n" |
| 156 | R"==(wei += 4 * KDHW_SIZE * OC_BLOCK * (group_oc + oc); )==" "\n" |
| 157 | R"==(/* WORK WITH SLM */ )==" "\n" |
| 158 | R"==(const bool left_tail = iw < 0; )==" "\n" |
| 159 | R"==(const bool left_nozero_tail = sub_group_id == 0 && iw >= 0; )==" "\n" |
| 160 | R"==(const bool right_tail = (iw + PW + SLM_TAIL >= IW) && (iw + PW < IW); )==" "\n" |
| 161 | R"==(const bool empty = (iw + PW >= IW); )==" "\n" |
| 162 | R"==(const bool right_nozero_tail )==" "\n" |
| 163 | R"==(= sp == (LWS_1 - 1) && (iw + PW + SLM_TAIL < IW); )==" "\n" |
| 164 | R"==(barrier(CLK_LOCAL_MEM_FENCE); )==" "\n" |
| 165 | R"==(/* KD */ )==" "\n" |
| 166 | R"==(#if KD > 1 )==" "\n" |
| 167 | R"==(for (int kd = 0; kd < KD; kd++) { )==" "\n" |
| 168 | R"==(if (kd * (1 + DD) + id < 0 || kd * (1 + DD) + id >= ID) { )==" "\n" |
| 169 | R"==(S_part += SRC_SLM_SIZE * KH; )==" "\n" |
| 170 | R"==(src += IC * IW * IH * (1 + DD); )==" "\n" |
| 171 | R"==(continue; )==" "\n" |
| 172 | R"==(} )==" "\n" |
| 173 | R"==(#endif )==" "\n" |
| 174 | R"==(/* KH */ )==" "\n" |
| 175 | R"==(#if KH > 1 )==" "\n" |
| 176 | R"==(for (int kh = 0; kh < KH; kh++) { )==" "\n" |
| 177 | R"==(if (kh * (1 + DH) + ih < 0 || kh * (1 + DH) + ih >= IH) { )==" "\n" |
| 178 | R"==(S_part += SRC_SLM_SIZE; )==" "\n" |
| 179 | R"==(src += IC * IW * (1 + DH); )==" "\n" |
| 180 | R"==(continue; )==" "\n" |
| 181 | R"==(} )==" "\n" |
| 182 | R"==(#endif )==" "\n" |
| 183 | R"==(/* KW */ )==" "\n" |
| 184 | R"==(/* left tail */ )==" "\n" |
| 185 | R"==(#if PW > 0 )==" "\n" |
| 186 | R"==(if (left_tail) { )==" "\n" |
| 187 | R"==(for (int i = -PW; i < 0; i++) { )==" "\n" |
| 188 | R"==(S_part[i] = 0; )==" "\n" |
| 189 | R"==(} )==" "\n" |
| 190 | R"==(} )==" "\n" |
| 191 | R"==(#endif )==" "\n" |
| 192 | R"==(/* right tail */ )==" "\n" |
| 193 | R"==(#if ZERO_TAIL > 0 )==" "\n" |
| 194 | R"==(if (right_tail) { )==" "\n" |
| 195 | R"==(for (int i = SLM_TAIL; )==" "\n" |
| 196 | R"==(i < SW * OW_BLOCK + (KW - 1) * (1 + DW) - PW; i++) { )==" "\n" |
| 197 | R"==(S_part[i] = 0; )==" "\n" |
| 198 | R"==(} )==" "\n" |
| 199 | R"==(} )==" "\n" |
| 200 | R"==(#if SLM_NCHUNK < OW_NCHUNK )==" "\n" |
| 201 | R"==(if (empty) { )==" "\n" |
| 202 | R"==(for (int i = 0; i < SW * OW_BLOCK + (KW - 1) * (1 + DW) - PW; )==" "\n" |
| 203 | R"==(i++) { )==" "\n" |
| 204 | R"==(WRITE_SLM_BLOCK(S_part + i * 8, 0); )==" "\n" |
| 205 | R"==(} )==" "\n" |
| 206 | R"==(} )==" "\n" |
| 207 | R"==(#endif )==" "\n" |
| 208 | R"==(#endif )==" "\n" |
| 209 | R"==(#if SLM_NCHUNK < OW_NCHUNK )==" "\n" |
| 210 | R"==(if (iw + PW < IW) { )==" "\n" |
| 211 | R"==(#endif )==" "\n" |
| 212 | R"==(#if OW_NCHUNK > LWS_1 )==" "\n" |
| 213 | R"==(/* Copy tails in case of multigroups */ )==" "\n" |
| 214 | R"==(if (ow < OW) { )==" "\n" |
| 215 | R"==(#if PW > 0 )==" "\n" |
| 216 | R"==(if (left_nozero_tail) { )==" "\n" |
| 217 | R"==(__global uchar *s = (__global uchar *)src; )==" "\n" |
| 218 | R"==(for (int i = -PW; i < 0; i++) { )==" "\n" |
| 219 | R"==(GET_INT_BLOCK(S_part, i, s); )==" "\n" |
| 220 | R"==(} )==" "\n" |
| 221 | R"==(} )==" "\n" |
| 222 | R"==(#endif )==" "\n" |
| 223 | R"==(if (right_nozero_tail) { )==" "\n" |
| 224 | R"==(__global uchar *s = (__global uchar *)src; )==" "\n" |
| 225 | R"==(for (int i = SW * OW_BLOCK; )==" "\n" |
| 226 | R"==(i < SW * OW_BLOCK + (KW - 1) * (1 + DW) - PW; )==" "\n" |
| 227 | R"==(i++) { )==" "\n" |
| 228 | R"==(GET_INT_BLOCK(S_part, i, s); )==" "\n" |
| 229 | R"==(} )==" "\n" |
| 230 | R"==(} )==" "\n" |
| 231 | R"==(#endif )==" "\n" |
| 232 | R"==(#if SLM_TAIL != OW_BLOCK * SW )==" "\n" |
| 233 | R"==(/* Copy last block to SLM */ )==" "\n" |
| 234 | R"==(if (right_tail) { )==" "\n" |
| 235 | R"==(__global uchar *s = (__global uchar *)src; )==" "\n" |
| 236 | R"==(__attribute__((opencl_unroll_hint)) )==" "\n" |
| 237 | R"==(for (int i = 0; i < SLM_TAIL; i++) { )==" "\n" |
| 238 | R"==(GET_INT_BLOCK(S_part, i, s); )==" "\n" |
| 239 | R"==(} )==" "\n" |
| 240 | R"==(} else { )==" "\n" |
| 241 | R"==(#endif )==" "\n" |
| 242 | R"==(#if (SW * OW_BLOCK) % 8 == 0 )==" "\n" |
| 243 | R"==(/* Copy block to SLM */ )==" "\n" |
| 244 | R"==(#if IC == 4 )==" "\n" |
| 245 | R"==(__attribute__((opencl_unroll_hint)) )==" "\n" |
| 246 | R"==(for (int i = 0; i < SW * OW_BLOCK; i += 8) { )==" "\n" |
| 247 | R"==(WRITE_SLM_BLOCK(S_part + i, )==" "\n" |
| 248 | R"==(intel_sub_group_block_read((const __global )==" "\n" |
| 249 | R"==(uint *)(&src[i * IC]))); )==" "\n" |
| 250 | R"==(} )==" "\n" |
| 251 | R"==(#else )==" "\n" |
| 252 | R"==(__global uchar *s = (__global uchar *)src; )==" "\n" |
| 253 | R"==(__attribute__((opencl_unroll_hint)) )==" "\n" |
| 254 | R"==(for (int i = 0; i < SW * OW_BLOCK; i += 8) { )==" "\n" |
| 255 | R"==(const int local_id = get_sub_group_local_id(); )==" "\n" |
| 256 | R"==(uchar4 res = 0; )==" "\n" |
| 257 | R"==(for (int j = 0; j < IC; j++) { )==" "\n" |
| 258 | R"==(res[j] = s[local_id * IC + j]; )==" "\n" |
| 259 | R"==(} )==" "\n" |
| 260 | R"==(s += IC * 8; )==" "\n" |
| 261 | R"==(WRITE_SLM_BLOCK(S_part + i, as_uint(res)); )==" "\n" |
| 262 | R"==(} )==" "\n" |
| 263 | R"==(#endif )==" "\n" |
| 264 | R"==(#else )==" "\n" |
| 265 | R"==(#if IC == 4 )==" "\n" |
| 266 | R"==(__attribute__((opencl_unroll_hint)) )==" "\n" |
| 267 | R"==(for (int i = 0; i < SW * OW_BLOCK; i += 4) { )==" "\n" |
| 268 | R"==(WRITE_SLM_BLOCK_SHORT(S_part + i, )==" "\n" |
| 269 | R"==(intel_sub_group_block_read_us( )==" "\n" |
| 270 | R"==((const __global ushort *)(&src[i * IC]))); )==" "\n" |
| 271 | R"==(} )==" "\n" |
| 272 | R"==(#else )==" "\n" |
| 273 | R"==(__global uchar *s = (__global uchar *)src; )==" "\n" |
| 274 | R"==(__attribute__((opencl_unroll_hint)) )==" "\n" |
| 275 | R"==(for (int i = 0; i < SW * OW_BLOCK; i += 4) { )==" "\n" |
| 276 | R"==(const int local_id = get_sub_group_local_id(); )==" "\n" |
| 277 | R"==(uchar2 res = 0; )==" "\n" |
| 278 | R"==(#if IC == 3 )==" "\n" |
| 279 | R"==(res[0] = local_id % 2 ? s[local_id * IC / 2 + 1] )==" "\n" |
| 280 | R"==(: s[0 + local_id * IC / 2]; )==" "\n" |
| 281 | R"==(res[1] = local_id % 2 ?: s[1 + local_id * IC / 2]; )==" "\n" |
| 282 | R"==(#elif IC == 2 )==" "\n" |
| 283 | R"==(res[0] = local_id % 2 ?: s[0 + local_id * IC / 2]; )==" "\n" |
| 284 | R"==(res[1] = local_id % 2 ?: s[1 + local_id * IC / 2]; )==" "\n" |
| 285 | R"==(#else )==" "\n" |
| 286 | R"==(res[0] = local_id % 2 ?: s[0 + local_id * IC / 2]; )==" "\n" |
| 287 | R"==(#endif )==" "\n" |
| 288 | R"==(s += IC * 4; )==" "\n" |
| 289 | R"==(WRITE_SLM_BLOCK_SHORT(S_part + i, as_ushort(res)); )==" "\n" |
| 290 | R"==(} )==" "\n" |
| 291 | R"==(#endif )==" "\n" |
| 292 | R"==(#endif )==" "\n" |
| 293 | R"==(#if SLM_TAIL != OW_BLOCK * SW )==" "\n" |
| 294 | R"==(} )==" "\n" |
| 295 | R"==(#endif )==" "\n" |
| 296 | R"==(#if OW_NCHUNK > LWS_1 )==" "\n" |
| 297 | R"==(} )==" "\n" |
| 298 | R"==(#endif )==" "\n" |
| 299 | R"==(#if SLM_NCHUNK < OW_NCHUNK )==" "\n" |
| 300 | R"==(} )==" "\n" |
| 301 | R"==(#endif )==" "\n" |
| 302 | R"==(#if KH > 1 )==" "\n" |
| 303 | R"==(S_part += SRC_SLM_SIZE; )==" "\n" |
| 304 | R"==(src += IC * IW * (1 + DH); )==" "\n" |
| 305 | R"==(} )==" "\n" |
| 306 | R"==(S_part -= SRC_SLM_SIZE * KH; )==" "\n" |
| 307 | R"==(src -= IC * KH * IW * (1 + DH); )==" "\n" |
| 308 | R"==(#endif )==" "\n" |
| 309 | R"==(#if KD > 1 )==" "\n" |
| 310 | R"==(S_part += SRC_SLM_SIZE * KH; )==" "\n" |
| 311 | R"==(src += IC * IW * IH * (1 + DD); )==" "\n" |
| 312 | R"==(} )==" "\n" |
| 313 | R"==(#endif )==" "\n" |
| 314 | R"==(barrier(CLK_LOCAL_MEM_FENCE); )==" "\n" |
| 315 | R"==(SRC_MMAD_DATA8_T S; )==" "\n" |
| 316 | R"==(int8 W0 = 0, W1 = 0, W2 = 0, W3 = 0; )==" "\n" |
| 317 | R"==(int W00 = 0, W10 = 0, W20 = 0, W30 = 0; )==" "\n" |
| 318 | R"==(int8 C00 = 0; )==" "\n" |
| 319 | R"==(int8 C10 = 0; )==" "\n" |
| 320 | R"==(int8 C20 = 0; )==" "\n" |
| 321 | R"==(int8 C30 = 0; )==" "\n" |
| 322 | R"==(#if OW_BLOCK == 12 )==" "\n" |
| 323 | R"==(SRC_MMAD_DATA4_T SS; )==" "\n" |
| 324 | R"==(int4 C01 = 0; )==" "\n" |
| 325 | R"==(int4 C11 = 0; )==" "\n" |
| 326 | R"==(int4 C21 = 0; )==" "\n" |
| 327 | R"==(int4 C31 = 0; )==" "\n" |
| 328 | R"==(#endif )==" "\n" |
| 329 | R"==(#if OW_BLOCK == 16 )==" "\n" |
| 330 | R"==(int8 C01 = 0; )==" "\n" |
| 331 | R"==(int8 C11 = 0; )==" "\n" |
| 332 | R"==(int8 C21 = 0; )==" "\n" |
| 333 | R"==(int8 C31 = 0; )==" "\n" |
| 334 | R"==(#endif )==" "\n" |
| 335 | R"==(for (int i = 0; i < KDHW_SIZE - KDHW_SIZE % 8; i += 8) { )==" "\n" |
| 336 | R"==(const int ihw = (i + sub_local_id) % (KW * KH); )==" "\n" |
| 337 | R"==(const int filter_iw = (ihw % KW) * (1 + DW); )==" "\n" |
| 338 | R"==(const int filter_ih = ihw / KW; )==" "\n" |
| 339 | R"==(const int filter_id = (i + sub_local_id) / (KH * KW); )==" "\n" |
| 340 | R"==(const int filter = (filter_ih * (1 + DH) + ih >= 0) )==" "\n" |
| 341 | R"==(&& (filter_ih * (1 + DH) + ih < IH) )==" "\n" |
| 342 | R"==(&& (filter_id * (1 + DD) + id >= 0 )==" "\n" |
| 343 | R"==(&& filter_id * (1 + DD) + id < ID); )==" "\n" |
| 344 | R"==(BLOCK_READ_WHT8(W0, 0); )==" "\n" |
| 345 | R"==(#if OC_PADD8 * 4 > OC_BLOCK )==" "\n" |
| 346 | R"==(BLOCK_READ_WHT8(W1, KDHW_SIZE * OC_BLOCK); )==" "\n" |
| 347 | R"==(#endif )==" "\n" |
| 348 | R"==(#if OC_PADD8 * 4 > OC_BLOCK * 2 )==" "\n" |
| 349 | R"==(BLOCK_READ_WHT8(W2, 2 * KDHW_SIZE * OC_BLOCK); )==" "\n" |
| 350 | R"==(#endif )==" "\n" |
| 351 | R"==(#if OC_PADD8 * 4 > OC_BLOCK * 3 )==" "\n" |
| 352 | R"==(BLOCK_READ_WHT8(W3, 3 * KDHW_SIZE * OC_BLOCK); )==" "\n" |
| 353 | R"==(#endif )==" "\n" |
| 354 | R"==(if (filter) { )==" "\n" |
| 355 | R"==(S.s0 = S_work[SW * 0 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 356 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 357 | R"==(S.s1 = S_work[SW * 1 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 358 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 359 | R"==(S.s2 = S_work[SW * 2 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 360 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 361 | R"==(S.s3 = S_work[SW * 3 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 362 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 363 | R"==(S.s4 = S_work[SW * 4 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 364 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 365 | R"==(S.s5 = S_work[SW * 5 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 366 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 367 | R"==(S.s6 = S_work[SW * 6 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 368 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 369 | R"==(S.s7 = S_work[SW * 7 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 370 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 371 | R"==(#if OW_BLOCK == 12 )==" "\n" |
| 372 | R"==(SS.s0 = S_work[SW * 8 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 373 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 374 | R"==(SS.s1 = S_work[SW * 9 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 375 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 376 | R"==(SS.s2 = S_work[SW * 10 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 377 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 378 | R"==(SS.s3 = S_work[SW * 11 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 379 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 380 | R"==(#endif )==" "\n" |
| 381 | R"==(} else { )==" "\n" |
| 382 | R"==(S = 0; )==" "\n" |
| 383 | R"==(#if OW_BLOCK == 12 )==" "\n" |
| 384 | R"==(SS = 0; )==" "\n" |
| 385 | R"==(#endif )==" "\n" |
| 386 | R"==(} )==" "\n" |
| 387 | R"==(C00 = mmad8x8(S, W0, C00); )==" "\n" |
| 388 | R"==(C10 = mmad8x8(S, W1, C10); )==" "\n" |
| 389 | R"==(C20 = mmad8x8(S, W2, C20); )==" "\n" |
| 390 | R"==(C30 = mmad8x8(S, W3, C30); )==" "\n" |
| 391 | R"==(#if OW_BLOCK == 12 )==" "\n" |
| 392 | R"==(C01 = mmad8x4(SS, W0, C01); )==" "\n" |
| 393 | R"==(C11 = mmad8x4(SS, W1, C11); )==" "\n" |
| 394 | R"==(C21 = mmad8x4(SS, W2, C21); )==" "\n" |
| 395 | R"==(C31 = mmad8x4(SS, W3, C31); )==" "\n" |
| 396 | R"==(#endif )==" "\n" |
| 397 | R"==(#if OW_BLOCK == 16 )==" "\n" |
| 398 | R"==(if (filter) { )==" "\n" |
| 399 | R"==(S.s0 = S_work[SW * 8 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 400 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 401 | R"==(S.s1 = S_work[SW * 9 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 402 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 403 | R"==(S.s2 = S_work[SW * 10 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 404 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 405 | R"==(S.s3 = S_work[SW * 11 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 406 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 407 | R"==(S.s4 = S_work[SW * 12 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 408 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 409 | R"==(S.s5 = S_work[SW * 13 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 410 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 411 | R"==(S.s6 = S_work[SW * 14 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 412 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 413 | R"==(S.s7 = S_work[SW * 15 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 414 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 415 | R"==(} else { )==" "\n" |
| 416 | R"==(S = 0; )==" "\n" |
| 417 | R"==(} )==" "\n" |
| 418 | R"==(C01 = mmad8x8(S, W0, C01); )==" "\n" |
| 419 | R"==(C11 = mmad8x8(S, W1, C11); )==" "\n" |
| 420 | R"==(C21 = mmad8x8(S, W2, C21); )==" "\n" |
| 421 | R"==(C31 = mmad8x8(S, W3, C31); )==" "\n" |
| 422 | R"==(#endif )==" "\n" |
| 423 | R"==(wei += OC_BLOCK * 8; )==" "\n" |
| 424 | R"==(} )==" "\n" |
| 425 | R"==(for (int i = KDHW_SIZE - KDHW_SIZE % 8; i < KDHW_SIZE; i++) { )==" "\n" |
| 426 | R"==(const int ihw = (i) % (KW * KH); )==" "\n" |
| 427 | R"==(const int filter_iw = (ihw % KW) * (1 + DW); )==" "\n" |
| 428 | R"==(const int filter_ih = ihw / KW; )==" "\n" |
| 429 | R"==(const int filter_id = (i) / (KH * KW); )==" "\n" |
| 430 | R"==(const int filter = (filter_ih * (1 + DH) + ih >= 0) )==" "\n" |
| 431 | R"==(&& (filter_ih * (1 + DH) + ih < IH) )==" "\n" |
| 432 | R"==(&& (filter_id * (1 + DD) + id >= 0 )==" "\n" |
| 433 | R"==(&& filter_id * (1 + DD) + id < ID); )==" "\n" |
| 434 | R"==(if (filter) { )==" "\n" |
| 435 | R"==(BLOCK_READ_WHT(W00, 0); )==" "\n" |
| 436 | R"==(#if OC_PADD8 * 4 > OC_BLOCK )==" "\n" |
| 437 | R"==(BLOCK_READ_WHT(W10, KDHW_SIZE * OC_BLOCK); )==" "\n" |
| 438 | R"==(#endif )==" "\n" |
| 439 | R"==(#if OC_PADD8 * 4 > OC_BLOCK * 2 )==" "\n" |
| 440 | R"==(BLOCK_READ_WHT(W20, 2 * KDHW_SIZE * OC_BLOCK); )==" "\n" |
| 441 | R"==(#endif )==" "\n" |
| 442 | R"==(#if OC_PADD8 * 4 > OC_BLOCK * 3 )==" "\n" |
| 443 | R"==(BLOCK_READ_WHT(W30, 3 * KDHW_SIZE * OC_BLOCK); )==" "\n" |
| 444 | R"==(#endif )==" "\n" |
| 445 | R"==(S.s0 = S_work[SW * 0 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 446 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 447 | R"==(S.s1 = S_work[SW * 1 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 448 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 449 | R"==(S.s2 = S_work[SW * 2 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 450 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 451 | R"==(S.s3 = S_work[SW * 3 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 452 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 453 | R"==(S.s4 = S_work[SW * 4 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 454 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 455 | R"==(S.s5 = S_work[SW * 5 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 456 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 457 | R"==(S.s6 = S_work[SW * 6 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 458 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 459 | R"==(S.s7 = S_work[SW * 7 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 460 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 461 | R"==(#if OW_BLOCK == 12 )==" "\n" |
| 462 | R"==(SS.s0 = S_work[SW * 8 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 463 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 464 | R"==(SS.s1 = S_work[SW * 9 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 465 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 466 | R"==(SS.s2 = S_work[SW * 10 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 467 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 468 | R"==(SS.s3 = S_work[SW * 11 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 469 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 470 | R"==(#endif )==" "\n" |
| 471 | R"==(C00.s0 = idot4(AS_SRC_DATA4_T(S.s0), as_char4(W00), C00.s0); )==" "\n" |
| 472 | R"==(C00.s1 = idot4(AS_SRC_DATA4_T(S.s1), as_char4(W00), C00.s1); )==" "\n" |
| 473 | R"==(C00.s2 = idot4(AS_SRC_DATA4_T(S.s2), as_char4(W00), C00.s2); )==" "\n" |
| 474 | R"==(C00.s3 = idot4(AS_SRC_DATA4_T(S.s3), as_char4(W00), C00.s3); )==" "\n" |
| 475 | R"==(C00.s4 = idot4(AS_SRC_DATA4_T(S.s4), as_char4(W00), C00.s4); )==" "\n" |
| 476 | R"==(C00.s5 = idot4(AS_SRC_DATA4_T(S.s5), as_char4(W00), C00.s5); )==" "\n" |
| 477 | R"==(C00.s6 = idot4(AS_SRC_DATA4_T(S.s6), as_char4(W00), C00.s6); )==" "\n" |
| 478 | R"==(C00.s7 = idot4(AS_SRC_DATA4_T(S.s7), as_char4(W00), C00.s7); )==" "\n" |
| 479 | R"==(C10.s0 = idot4(AS_SRC_DATA4_T(S.s0), as_char4(W10), C10.s0); )==" "\n" |
| 480 | R"==(C10.s1 = idot4(AS_SRC_DATA4_T(S.s1), as_char4(W10), C10.s1); )==" "\n" |
| 481 | R"==(C10.s2 = idot4(AS_SRC_DATA4_T(S.s2), as_char4(W10), C10.s2); )==" "\n" |
| 482 | R"==(C10.s3 = idot4(AS_SRC_DATA4_T(S.s3), as_char4(W10), C10.s3); )==" "\n" |
| 483 | R"==(C10.s4 = idot4(AS_SRC_DATA4_T(S.s4), as_char4(W10), C10.s4); )==" "\n" |
| 484 | R"==(C10.s5 = idot4(AS_SRC_DATA4_T(S.s5), as_char4(W10), C10.s5); )==" "\n" |
| 485 | R"==(C10.s6 = idot4(AS_SRC_DATA4_T(S.s6), as_char4(W10), C10.s6); )==" "\n" |
| 486 | R"==(C10.s7 = idot4(AS_SRC_DATA4_T(S.s7), as_char4(W10), C10.s7); )==" "\n" |
| 487 | R"==(C20.s0 = idot4(AS_SRC_DATA4_T(S.s0), as_char4(W20), C20.s0); )==" "\n" |
| 488 | R"==(C20.s1 = idot4(AS_SRC_DATA4_T(S.s1), as_char4(W20), C20.s1); )==" "\n" |
| 489 | R"==(C20.s2 = idot4(AS_SRC_DATA4_T(S.s2), as_char4(W20), C20.s2); )==" "\n" |
| 490 | R"==(C20.s3 = idot4(AS_SRC_DATA4_T(S.s3), as_char4(W20), C20.s3); )==" "\n" |
| 491 | R"==(C20.s4 = idot4(AS_SRC_DATA4_T(S.s4), as_char4(W20), C20.s4); )==" "\n" |
| 492 | R"==(C20.s5 = idot4(AS_SRC_DATA4_T(S.s5), as_char4(W20), C20.s5); )==" "\n" |
| 493 | R"==(C20.s6 = idot4(AS_SRC_DATA4_T(S.s6), as_char4(W20), C20.s6); )==" "\n" |
| 494 | R"==(C20.s7 = idot4(AS_SRC_DATA4_T(S.s7), as_char4(W20), C20.s7); )==" "\n" |
| 495 | R"==(C30.s0 = idot4(AS_SRC_DATA4_T(S.s0), as_char4(W30), C30.s0); )==" "\n" |
| 496 | R"==(C30.s1 = idot4(AS_SRC_DATA4_T(S.s1), as_char4(W30), C30.s1); )==" "\n" |
| 497 | R"==(C30.s2 = idot4(AS_SRC_DATA4_T(S.s2), as_char4(W30), C30.s2); )==" "\n" |
| 498 | R"==(C30.s3 = idot4(AS_SRC_DATA4_T(S.s3), as_char4(W30), C30.s3); )==" "\n" |
| 499 | R"==(C30.s4 = idot4(AS_SRC_DATA4_T(S.s4), as_char4(W30), C30.s4); )==" "\n" |
| 500 | R"==(C30.s5 = idot4(AS_SRC_DATA4_T(S.s5), as_char4(W30), C30.s5); )==" "\n" |
| 501 | R"==(C30.s6 = idot4(AS_SRC_DATA4_T(S.s6), as_char4(W30), C30.s6); )==" "\n" |
| 502 | R"==(C30.s7 = idot4(AS_SRC_DATA4_T(S.s7), as_char4(W30), C30.s7); )==" "\n" |
| 503 | R"==(#if OW_BLOCK == 12 )==" "\n" |
| 504 | R"==(C01.s0 = idot4(AS_SRC_DATA4_T(SS.s0), as_char4(W00), C01.s0); )==" "\n" |
| 505 | R"==(C01.s1 = idot4(AS_SRC_DATA4_T(SS.s1), as_char4(W00), C01.s1); )==" "\n" |
| 506 | R"==(C01.s2 = idot4(AS_SRC_DATA4_T(SS.s2), as_char4(W00), C01.s2); )==" "\n" |
| 507 | R"==(C01.s3 = idot4(AS_SRC_DATA4_T(SS.s3), as_char4(W00), C01.s3); )==" "\n" |
| 508 | R"==(C11.s0 = idot4(AS_SRC_DATA4_T(SS.s0), as_char4(W10), C11.s0); )==" "\n" |
| 509 | R"==(C11.s1 = idot4(AS_SRC_DATA4_T(SS.s1), as_char4(W10), C11.s1); )==" "\n" |
| 510 | R"==(C11.s2 = idot4(AS_SRC_DATA4_T(SS.s2), as_char4(W10), C11.s2); )==" "\n" |
| 511 | R"==(C11.s3 = idot4(AS_SRC_DATA4_T(SS.s3), as_char4(W10), C11.s3); )==" "\n" |
| 512 | R"==(C21.s0 = idot4(AS_SRC_DATA4_T(SS.s0), as_char4(W20), C21.s0); )==" "\n" |
| 513 | R"==(C21.s1 = idot4(AS_SRC_DATA4_T(SS.s1), as_char4(W20), C21.s1); )==" "\n" |
| 514 | R"==(C21.s2 = idot4(AS_SRC_DATA4_T(SS.s2), as_char4(W20), C21.s2); )==" "\n" |
| 515 | R"==(C21.s3 = idot4(AS_SRC_DATA4_T(SS.s3), as_char4(W20), C21.s3); )==" "\n" |
| 516 | R"==(C31.s0 = idot4(AS_SRC_DATA4_T(SS.s0), as_char4(W30), C31.s0); )==" "\n" |
| 517 | R"==(C31.s1 = idot4(AS_SRC_DATA4_T(SS.s1), as_char4(W30), C31.s1); )==" "\n" |
| 518 | R"==(C31.s2 = idot4(AS_SRC_DATA4_T(SS.s2), as_char4(W30), C31.s2); )==" "\n" |
| 519 | R"==(C31.s3 = idot4(AS_SRC_DATA4_T(SS.s3), as_char4(W30), C31.s3); )==" "\n" |
| 520 | R"==(#endif )==" "\n" |
| 521 | R"==(#if OW_BLOCK == 16 )==" "\n" |
| 522 | R"==(S.s0 = S_work[SW * 8 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 523 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 524 | R"==(S.s1 = S_work[SW * 9 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 525 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 526 | R"==(S.s2 = S_work[SW * 10 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 527 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 528 | R"==(S.s3 = S_work[SW * 11 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 529 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 530 | R"==(S.s4 = S_work[SW * 12 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 531 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 532 | R"==(S.s5 = S_work[SW * 13 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 533 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 534 | R"==(S.s6 = S_work[SW * 14 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 535 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 536 | R"==(S.s7 = S_work[SW * 15 + SRC_SLM_SIZE * KH * filter_id )==" "\n" |
| 537 | R"==(+ SRC_SLM_SIZE * filter_ih + filter_iw]; )==" "\n" |
| 538 | R"==(C01.s0 = idot4(AS_SRC_DATA4_T(S.s0), as_char4(W00), C01.s0); )==" "\n" |
| 539 | R"==(C01.s1 = idot4(AS_SRC_DATA4_T(S.s1), as_char4(W00), C01.s1); )==" "\n" |
| 540 | R"==(C01.s2 = idot4(AS_SRC_DATA4_T(S.s2), as_char4(W00), C01.s2); )==" "\n" |
| 541 | R"==(C01.s3 = idot4(AS_SRC_DATA4_T(S.s3), as_char4(W00), C01.s3); )==" "\n" |
| 542 | R"==(C01.s4 = idot4(AS_SRC_DATA4_T(S.s4), as_char4(W00), C01.s4); )==" "\n" |
| 543 | R"==(C01.s5 = idot4(AS_SRC_DATA4_T(S.s5), as_char4(W00), C01.s5); )==" "\n" |
| 544 | R"==(C01.s6 = idot4(AS_SRC_DATA4_T(S.s6), as_char4(W00), C01.s6); )==" "\n" |
| 545 | R"==(C01.s7 = idot4(AS_SRC_DATA4_T(S.s7), as_char4(W00), C01.s7); )==" "\n" |
| 546 | R"==(C11.s0 = idot4(AS_SRC_DATA4_T(S.s0), as_char4(W10), C11.s0); )==" "\n" |
| 547 | R"==(C11.s1 = idot4(AS_SRC_DATA4_T(S.s1), as_char4(W10), C11.s1); )==" "\n" |
| 548 | R"==(C11.s2 = idot4(AS_SRC_DATA4_T(S.s2), as_char4(W10), C11.s2); )==" "\n" |
| 549 | R"==(C11.s3 = idot4(AS_SRC_DATA4_T(S.s3), as_char4(W10), C11.s3); )==" "\n" |
| 550 | R"==(C11.s4 = idot4(AS_SRC_DATA4_T(S.s4), as_char4(W10), C11.s4); )==" "\n" |
| 551 | R"==(C11.s5 = idot4(AS_SRC_DATA4_T(S.s5), as_char4(W10), C11.s5); )==" "\n" |
| 552 | R"==(C11.s6 = idot4(AS_SRC_DATA4_T(S.s6), as_char4(W10), C11.s6); )==" "\n" |
| 553 | R"==(C11.s7 = idot4(AS_SRC_DATA4_T(S.s7), as_char4(W10), C11.s7); )==" "\n" |
| 554 | R"==(C21.s0 = idot4(AS_SRC_DATA4_T(S.s0), as_char4(W20), C21.s0); )==" "\n" |
| 555 | R"==(C21.s1 = idot4(AS_SRC_DATA4_T(S.s1), as_char4(W20), C21.s1); )==" "\n" |
| 556 | R"==(C21.s2 = idot4(AS_SRC_DATA4_T(S.s2), as_char4(W20), C21.s2); )==" "\n" |
| 557 | R"==(C21.s3 = idot4(AS_SRC_DATA4_T(S.s3), as_char4(W20), C21.s3); )==" "\n" |
| 558 | R"==(C21.s4 = idot4(AS_SRC_DATA4_T(S.s4), as_char4(W20), C21.s4); )==" "\n" |
| 559 | R"==(C21.s5 = idot4(AS_SRC_DATA4_T(S.s5), as_char4(W20), C21.s5); )==" "\n" |
| 560 | R"==(C21.s6 = idot4(AS_SRC_DATA4_T(S.s6), as_char4(W20), C21.s6); )==" "\n" |
| 561 | R"==(C21.s7 = idot4(AS_SRC_DATA4_T(S.s7), as_char4(W20), C21.s7); )==" "\n" |
| 562 | R"==(C31.s0 = idot4(AS_SRC_DATA4_T(S.s0), as_char4(W30), C31.s0); )==" "\n" |
| 563 | R"==(C31.s1 = idot4(AS_SRC_DATA4_T(S.s1), as_char4(W30), C31.s1); )==" "\n" |
| 564 | R"==(C31.s2 = idot4(AS_SRC_DATA4_T(S.s2), as_char4(W30), C31.s2); )==" "\n" |
| 565 | R"==(C31.s3 = idot4(AS_SRC_DATA4_T(S.s3), as_char4(W30), C31.s3); )==" "\n" |
| 566 | R"==(C31.s4 = idot4(AS_SRC_DATA4_T(S.s4), as_char4(W30), C31.s4); )==" "\n" |
| 567 | R"==(C31.s5 = idot4(AS_SRC_DATA4_T(S.s5), as_char4(W30), C31.s5); )==" "\n" |
| 568 | R"==(C31.s6 = idot4(AS_SRC_DATA4_T(S.s6), as_char4(W30), C31.s6); )==" "\n" |
| 569 | R"==(C31.s7 = idot4(AS_SRC_DATA4_T(S.s7), as_char4(W30), C31.s7); )==" "\n" |
| 570 | R"==(#endif )==" "\n" |
| 571 | R"==(} )==" "\n" |
| 572 | R"==(wei += OC_BLOCK; )==" "\n" |
| 573 | R"==(} )==" "\n" |
| 574 | R"==(DST_DATA16_T R1, R2, R3, R4; )==" "\n" |
| 575 | R"==(#if SCALES_PER_OC )==" "\n" |
| 576 | R"==(float4 scales; )==" "\n" |
| 577 | R"==(block_read_scales( )==" "\n" |
| 578 | R"==(&scales, (group_oc + oc) * OC_BLOCK, sub_local_id, runtime_scales); )==" "\n" |
| 579 | R"==(#endif )==" "\n" |
| 580 | R"==(#if WITH_BIAS )==" "\n" |
| 581 | R"==(float4 bia; )==" "\n" |
| 582 | R"==(BLOCK_READ_BIA(bia, (group_oc + oc) * OC_BLOCK); )==" "\n" |
| 583 | R"==(#define QUANTIZE_ADD_BIAS() tmp = SCALE * fma(tmp, (float4)1, bia); )==" "\n" |
| 584 | R"==(#define QUANTIZE_ADD_BIAS_4() \ )==" "\n" |
| 585 | R"==(tmp0 = ((float8)SCALE_VEC4) * fma(tmp0, (float8)1, bia.s01230123); \ )==" "\n" |
| 586 | R"==(tmp1 = ((float8)SCALE_VEC4) * fma(tmp1, (float8)1, bia.s01230123); )==" "\n" |
| 587 | R"==(#else )==" "\n" |
| 588 | R"==(#define QUANTIZE_ADD_BIAS() tmp *= SCALE; )==" "\n" |
| 589 | R"==(#define QUANTIZE_ADD_BIAS_4() \ )==" "\n" |
| 590 | R"==(tmp0 *= SCALE_VEC4; \ )==" "\n" |
| 591 | R"==(tmp1 *= SCALE_VEC4; )==" "\n" |
| 592 | R"==(#endif )==" "\n" |
| 593 | R"==(#if WITH_POST_OP )==" "\n" |
| 594 | R"==(#define APPLY_POST_OPS_COMMON(nelems, accum, sum) \ )==" "\n" |
| 595 | R"==(for (int didx = 0; didx < nelems; ++didx) { \ )==" "\n" |
| 596 | R"==(float tmp_i = accum[didx]; \ )==" "\n" |
| 597 | R"==(SUM_DATA_T d_i = sum[didx]; \ )==" "\n" |
| 598 | R"==(const int po_mb = group_mb % MB; \ )==" "\n" |
| 599 | R"==(const int po_oc \ )==" "\n" |
| 600 | R"==(= (oc * OC_BLOCK + ((didx * SUB_GROUP_SIZE) % OC_BLOCK) \ )==" "\n" |
| 601 | R"==(+ sub_local_id) \ )==" "\n" |
| 602 | R"==(% (OC * G); \ )==" "\n" |
| 603 | R"==(APPLY_POST_OPS_SERIAL_BINARY_2D( \ )==" "\n" |
| 604 | R"==(tmp_i, float, d_i, SUM_DATA_T, po_mb, 1, po_oc, 1); \ )==" "\n" |
| 605 | R"==(accum[didx] = tmp_i; \ )==" "\n" |
| 606 | R"==(} )==" "\n" |
| 607 | R"==(#if DST_NHWC )==" "\n" |
| 608 | R"==(#define DO_POST_OP() \ )==" "\n" |
| 609 | R"==(do { \ )==" "\n" |
| 610 | R"==(SUM_DATA4_T d; \ )==" "\n" |
| 611 | R"==(if (WITH_SUM) \ )==" "\n" |
| 612 | R"==(d = AS_SUM_DATA4_T( \ )==" "\n" |
| 613 | R"==(read_oc_block4(dst, (group_oc + oc) * OC_BLOCK)); \ )==" "\n" |
| 614 | R"==(APPLY_POST_OPS_COMMON(4, tmp, d); \ )==" "\n" |
| 615 | R"==(} while (0) )==" "\n" |
| 616 | R"==(#define DO_POST_OP_4() \ )==" "\n" |
| 617 | R"==({ \ )==" "\n" |
| 618 | R"==(SUM_DATA16_T d; \ )==" "\n" |
| 619 | R"==(if (WITH_SUM) \ )==" "\n" |
| 620 | R"==(d = AS_SUM_DATA16_T( \ )==" "\n" |
| 621 | R"==((DST_DATA16_T)(read_oc_block4(dst + G * OC * 0, \ )==" "\n" |
| 622 | R"==((group_oc + oc) * OC_BLOCK), \ )==" "\n" |
| 623 | R"==(read_oc_block4(dst + G * OC * 1, \ )==" "\n" |
| 624 | R"==((group_oc + oc) * OC_BLOCK), \ )==" "\n" |
| 625 | R"==(read_oc_block4(dst + G * OC * 2, \ )==" "\n" |
| 626 | R"==((group_oc + oc) * OC_BLOCK), \ )==" "\n" |
| 627 | R"==(read_oc_block4(dst + G * OC * 3, \ )==" "\n" |
| 628 | R"==((group_oc + oc) * OC_BLOCK))); \ )==" "\n" |
| 629 | R"==(float16 tmp_x16 = (float16)(tmp0, tmp1); \ )==" "\n" |
| 630 | R"==(APPLY_POST_OPS_COMMON(16, tmp_x16, d); \ )==" "\n" |
| 631 | R"==(tmp0 = tmp_x16.s01234567; \ )==" "\n" |
| 632 | R"==(tmp1 = tmp_x16.s89abcdef; \ )==" "\n" |
| 633 | R"==(} )==" "\n" |
| 634 | R"==(#else )==" "\n" |
| 635 | R"==(#define DO_POST_OP() \ )==" "\n" |
| 636 | R"==({ \ )==" "\n" |
| 637 | R"==(SUM_DATA4_T d; \ )==" "\n" |
| 638 | R"==(if (WITH_SUM) d = AS_SUM_DATA4_T(BLOCK_READ_DST4(dst)); \ )==" "\n" |
| 639 | R"==(APPLY_POST_OPS_COMMON(4, tmp, d); \ )==" "\n" |
| 640 | R"==(} )==" "\n" |
| 641 | R"==(#define DO_POST_OP_4() \ )==" "\n" |
| 642 | R"==({ \ )==" "\n" |
| 643 | R"==(SUM_DATA16_T d; \ )==" "\n" |
| 644 | R"==(if (WITH_SUM) d = AS_SUM_DATA16_T(BLOCK_READ_DST16(dst)); \ )==" "\n" |
| 645 | R"==(float16 tmp_x16 = (float16)(tmp0, tmp1); \ )==" "\n" |
| 646 | R"==(APPLY_POST_OPS_COMMON(16, tmp_x16, d); \ )==" "\n" |
| 647 | R"==(tmp0 = tmp_x16.s01234567; \ )==" "\n" |
| 648 | R"==(tmp1 = tmp_x16.s89abcdef; \ )==" "\n" |
| 649 | R"==(} )==" "\n" |
| 650 | R"==(#endif )==" "\n" |
| 651 | R"==(#else )==" "\n" |
| 652 | R"==(#define DO_POST_OP() ; )==" "\n" |
| 653 | R"==(#define DO_POST_OP_4() ; )==" "\n" |
| 654 | R"==(#endif )==" "\n" |
| 655 | R"==(#define PACK(C0, C1, C2, C3, idx) \ )==" "\n" |
| 656 | R"==(do { \ )==" "\n" |
| 657 | R"==(tmp.s0 = ((int *)&C0)[idx]; \ )==" "\n" |
| 658 | R"==(tmp.s1 = ((int *)&C1)[idx]; \ )==" "\n" |
| 659 | R"==(tmp.s2 = ((int *)&C2)[idx]; \ )==" "\n" |
| 660 | R"==(tmp.s3 = ((int *)&C3)[idx]; \ )==" "\n" |
| 661 | R"==(} while (0) )==" "\n" |
| 662 | R"==(#define PACK_4(C0, C1, C2, C3, idx) \ )==" "\n" |
| 663 | R"==(do { \ )==" "\n" |
| 664 | R"==(tmp0.s0 = ((int *)&C0)[idx]; \ )==" "\n" |
| 665 | R"==(tmp0.s1 = ((int *)&C1)[idx]; \ )==" "\n" |
| 666 | R"==(tmp0.s2 = ((int *)&C2)[idx]; \ )==" "\n" |
| 667 | R"==(tmp0.s3 = ((int *)&C3)[idx]; \ )==" "\n" |
| 668 | R"==(\ )==" "\n" |
| 669 | R"==(tmp0.s4 = ((int *)&C0)[idx + 1]; \ )==" "\n" |
| 670 | R"==(tmp0.s5 = ((int *)&C1)[idx + 1]; \ )==" "\n" |
| 671 | R"==(tmp0.s6 = ((int *)&C2)[idx + 1]; \ )==" "\n" |
| 672 | R"==(tmp0.s7 = ((int *)&C3)[idx + 1]; \ )==" "\n" |
| 673 | R"==(\ )==" "\n" |
| 674 | R"==(tmp1.s0 = ((int *)&C0)[idx + 2]; \ )==" "\n" |
| 675 | R"==(tmp1.s1 = ((int *)&C1)[idx + 2]; \ )==" "\n" |
| 676 | R"==(tmp1.s2 = ((int *)&C2)[idx + 2]; \ )==" "\n" |
| 677 | R"==(tmp1.s3 = ((int *)&C3)[idx + 2]; \ )==" "\n" |
| 678 | R"==(\ )==" "\n" |
| 679 | R"==(tmp1.s4 = ((int *)&C0)[idx + 3]; \ )==" "\n" |
| 680 | R"==(tmp1.s5 = ((int *)&C1)[idx + 3]; \ )==" "\n" |
| 681 | R"==(tmp1.s6 = ((int *)&C2)[idx + 3]; \ )==" "\n" |
| 682 | R"==(tmp1.s7 = ((int *)&C3)[idx + 3]; \ )==" "\n" |
| 683 | R"==(} while (0) )==" "\n" |
| 684 | R"==(#define CONVERT_PACK() \ )==" "\n" |
| 685 | R"==(do { \ )==" "\n" |
| 686 | R"==(tmp_cvt = CONVERT_DST_DATA4_T(tmp); \ )==" "\n" |
| 687 | R"==(} while (0) )==" "\n" |
| 688 | R"==(#define CONVERT_PACK_4() \ )==" "\n" |
| 689 | R"==(do { \ )==" "\n" |
| 690 | R"==(R.s01234567 = CONVERT_DST_DATA8_T(tmp0); \ )==" "\n" |
| 691 | R"==(R.s89abcdef = CONVERT_DST_DATA8_T(tmp1); \ )==" "\n" |
| 692 | R"==(} while (0) )==" "\n" |
| 693 | R"==(#if DST_NHWC )==" "\n" |
| 694 | R"==(#define STORE_DST(C0, C1, C2, C3, i) \ )==" "\n" |
| 695 | R"==(do { \ )==" "\n" |
| 696 | R"==(PACK(C0, C1, C2, C3, i); \ )==" "\n" |
| 697 | R"==(QUANTIZE_ADD_BIAS(); \ )==" "\n" |
| 698 | R"==(DO_POST_OP(); \ )==" "\n" |
| 699 | R"==(CONVERT_PACK(); \ )==" "\n" |
| 700 | R"==(write_oc_block4(dst, (group_oc + oc) * OC_BLOCK, tmp_cvt); \ )==" "\n" |
| 701 | R"==(dst += OC; \ )==" "\n" |
| 702 | R"==(} while (0) )==" "\n" |
| 703 | R"==(#define STORE_DST_4(C0, C1, C2, C3, i) \ )==" "\n" |
| 704 | R"==(do { \ )==" "\n" |
| 705 | R"==(PACK_4(C0, C1, C2, C3, i); \ )==" "\n" |
| 706 | R"==(QUANTIZE_ADD_BIAS_4(); \ )==" "\n" |
| 707 | R"==(DO_POST_OP_4(); \ )==" "\n" |
| 708 | R"==(CONVERT_PACK_4(); \ )==" "\n" |
| 709 | R"==(write_oc_block4(dst, (group_oc + oc) * OC_BLOCK, R.s0123); \ )==" "\n" |
| 710 | R"==(write_oc_block4(dst + OC * 1, (group_oc + oc) * OC_BLOCK, R.s4567); \ )==" "\n" |
| 711 | R"==(write_oc_block4(dst + OC * 2, (group_oc + oc) * OC_BLOCK, R.s89ab); \ )==" "\n" |
| 712 | R"==(write_oc_block4(dst + OC * 3, (group_oc + oc) * OC_BLOCK, R.scdef); \ )==" "\n" |
| 713 | R"==(dst += 4 * OC; \ )==" "\n" |
| 714 | R"==(} while (0) )==" "\n" |
| 715 | R"==(#else )==" "\n" |
| 716 | R"==(#define STORE_DST(C0, C1, C2, C3, i) \ )==" "\n" |
| 717 | R"==(do { \ )==" "\n" |
| 718 | R"==(PACK(C0, C1, C2, C3, i); \ )==" "\n" |
| 719 | R"==(if (MB % MB_BLOCK == 0 || group_mb < MB) { \ )==" "\n" |
| 720 | R"==(QUANTIZE_ADD_BIAS(); \ )==" "\n" |
| 721 | R"==(DO_POST_OP(); \ )==" "\n" |
| 722 | R"==(CONVERT_PACK(); \ )==" "\n" |
| 723 | R"==(} else { \ )==" "\n" |
| 724 | R"==(tmp_cvt = 0; \ )==" "\n" |
| 725 | R"==(} \ )==" "\n" |
| 726 | R"==(BLOCK_WRITE_DST4(dst, tmp_cvt); \ )==" "\n" |
| 727 | R"==(dst += OC_BLOCK * MB_BLOCK; \ )==" "\n" |
| 728 | R"==(} while (0) )==" "\n" |
| 729 | R"==(#define STORE_DST_4(C0, C1, C2, C3, i) \ )==" "\n" |
| 730 | R"==(do { \ )==" "\n" |
| 731 | R"==(PACK_4(C0, C1, C2, C3, i); \ )==" "\n" |
| 732 | R"==(QUANTIZE_ADD_BIAS_4(); \ )==" "\n" |
| 733 | R"==(DO_POST_OP_4(); \ )==" "\n" |
| 734 | R"==(CONVERT_PACK_4(); \ )==" "\n" |
| 735 | R"==(BLOCK_WRITE_DST16(dst, R); \ )==" "\n" |
| 736 | R"==(dst += 4 * OC_BLOCK; \ )==" "\n" |
| 737 | R"==(} while (0) )==" "\n" |
| 738 | R"==(#endif )==" "\n" |
| 739 | R"==(if (ow < OW) { )==" "\n" |
| 740 | R"==(float4 tmp; )==" "\n" |
| 741 | R"==(DST_DATA4_T tmp_cvt; )==" "\n" |
| 742 | R"==(float8 tmp0, tmp1; )==" "\n" |
| 743 | R"==(DST_DATA16_T R; )==" "\n" |
| 744 | R"==(#if OW_TAIL )==" "\n" |
| 745 | R"==(if (ow + OW_BLOCK < OW) { )==" "\n" |
| 746 | R"==(#endif )==" "\n" |
| 747 | R"==(#if !DST_NHWC && MB_BLOCK == 32 )==" "\n" |
| 748 | R"==(STORE_DST(C00, C10, C20, C30, 0); )==" "\n" |
| 749 | R"==(STORE_DST(C00, C10, C20, C30, 1); )==" "\n" |
| 750 | R"==(STORE_DST(C00, C10, C20, C30, 2); )==" "\n" |
| 751 | R"==(STORE_DST(C00, C10, C20, C30, 3); )==" "\n" |
| 752 | R"==(STORE_DST(C00, C10, C20, C30, 4); )==" "\n" |
| 753 | R"==(STORE_DST(C00, C10, C20, C30, 5); )==" "\n" |
| 754 | R"==(STORE_DST(C00, C10, C20, C30, 6); )==" "\n" |
| 755 | R"==(STORE_DST(C00, C10, C20, C30, 7); )==" "\n" |
| 756 | R"==(#if OW_BLOCK >= 12 )==" "\n" |
| 757 | R"==(STORE_DST(C01, C11, C21, C31, 0); )==" "\n" |
| 758 | R"==(STORE_DST(C01, C11, C21, C31, 1); )==" "\n" |
| 759 | R"==(STORE_DST(C01, C11, C21, C31, 2); )==" "\n" |
| 760 | R"==(STORE_DST(C01, C11, C21, C31, 3); )==" "\n" |
| 761 | R"==(#endif )==" "\n" |
| 762 | R"==(#if OW_BLOCK == 16 )==" "\n" |
| 763 | R"==(STORE_DST(C01, C11, C21, C31, 4); )==" "\n" |
| 764 | R"==(STORE_DST(C01, C11, C21, C31, 5); )==" "\n" |
| 765 | R"==(STORE_DST(C01, C11, C21, C31, 6); )==" "\n" |
| 766 | R"==(STORE_DST(C01, C11, C21, C31, 7); )==" "\n" |
| 767 | R"==(#endif )==" "\n" |
| 768 | R"==(#else )==" "\n" |
| 769 | R"==(STORE_DST_4(C00, C10, C20, C30, 0); )==" "\n" |
| 770 | R"==(STORE_DST_4(C00, C10, C20, C30, 4); )==" "\n" |
| 771 | R"==(#if OW_BLOCK >= 12 )==" "\n" |
| 772 | R"==(STORE_DST_4(C01, C11, C21, C31, 0); )==" "\n" |
| 773 | R"==(#endif )==" "\n" |
| 774 | R"==(#if OW_BLOCK >= 16 )==" "\n" |
| 775 | R"==(STORE_DST_4(C01, C11, C21, C31, 4); )==" "\n" |
| 776 | R"==(#endif )==" "\n" |
| 777 | R"==(#endif )==" "\n" |
| 778 | R"==(#if OW_TAIL )==" "\n" |
| 779 | R"==(} else { )==" "\n" |
| 780 | R"==(#if OW_TAIL < 4 )==" "\n" |
| 781 | R"==(for (int i = 0; i < OW_TAIL; i++) { )==" "\n" |
| 782 | R"==(STORE_DST(C00, C10, C20, C30, i); )==" "\n" |
| 783 | R"==(} )==" "\n" |
| 784 | R"==(#else )==" "\n" |
| 785 | R"==(#if !DST_NHWC && MB_BLOCK == 32 )==" "\n" |
| 786 | R"==(STORE_DST(C00, C10, C20, C30, 0); )==" "\n" |
| 787 | R"==(STORE_DST(C00, C10, C20, C30, 1); )==" "\n" |
| 788 | R"==(STORE_DST(C00, C10, C20, C30, 2); )==" "\n" |
| 789 | R"==(STORE_DST(C00, C10, C20, C30, 3); )==" "\n" |
| 790 | R"==(#else )==" "\n" |
| 791 | R"==(STORE_DST_4(C00, C10, C20, C30, 0); )==" "\n" |
| 792 | R"==(#endif )==" "\n" |
| 793 | R"==(#endif )==" "\n" |
| 794 | R"==(#if OW_TAIL > 4 )==" "\n" |
| 795 | R"==(#if OW_TAIL < 8 )==" "\n" |
| 796 | R"==(for (int i = 4; i < OW_TAIL; i++) { )==" "\n" |
| 797 | R"==(STORE_DST(C00, C10, C20, C30, i); )==" "\n" |
| 798 | R"==(} )==" "\n" |
| 799 | R"==(#else )==" "\n" |
| 800 | R"==(#if !DST_NHWC && MB_BLOCK == 32 )==" "\n" |
| 801 | R"==(STORE_DST(C00, C10, C20, C30, 4); )==" "\n" |
| 802 | R"==(STORE_DST(C00, C10, C20, C30, 5); )==" "\n" |
| 803 | R"==(STORE_DST(C00, C10, C20, C30, 6); )==" "\n" |
| 804 | R"==(STORE_DST(C00, C10, C20, C30, 7); )==" "\n" |
| 805 | R"==(#else )==" "\n" |
| 806 | R"==(STORE_DST_4(C00, C10, C20, C30, 4); )==" "\n" |
| 807 | R"==(#endif )==" "\n" |
| 808 | R"==(#endif )==" "\n" |
| 809 | R"==(#if OW_TAIL > 8 )==" "\n" |
| 810 | R"==(#if OW_TAIL < 12 )==" "\n" |
| 811 | R"==(for (int i = 8; i < OW_TAIL; i++) { )==" "\n" |
| 812 | R"==(STORE_DST(C01, C11, C21, C31, i - CX1_OFFSET); )==" "\n" |
| 813 | R"==(} )==" "\n" |
| 814 | R"==(#else )==" "\n" |
| 815 | R"==(#if !DST_NHWC && MB_BLOCK == 32 )==" "\n" |
| 816 | R"==(STORE_DST(C01, C11, C21, C31, 0); )==" "\n" |
| 817 | R"==(STORE_DST(C01, C11, C21, C31, 1); )==" "\n" |
| 818 | R"==(STORE_DST(C01, C11, C21, C31, 2); )==" "\n" |
| 819 | R"==(STORE_DST(C01, C11, C21, C31, 3); )==" "\n" |
| 820 | R"==(#else )==" "\n" |
| 821 | R"==(STORE_DST_4(C01, C11, C21, C31, 0); )==" "\n" |
| 822 | R"==(#endif )==" "\n" |
| 823 | R"==(#endif )==" "\n" |
| 824 | R"==(#if OW_TAIL > 12 )==" "\n" |
| 825 | R"==(#if OW_TAIL < 16 )==" "\n" |
| 826 | R"==(for (int i = 12; i < OW_TAIL; i++) { )==" "\n" |
| 827 | R"==(STORE_DST(C01, C11, C21, C31, i - CX1_OFFSET); )==" "\n" |
| 828 | R"==(} )==" "\n" |
| 829 | R"==(#else )==" "\n" |
| 830 | R"==(#if !DST_NHWC && MB_BLOCK == 32 )==" "\n" |
| 831 | R"==(STORE_DST(C01, C11, C21, C31, 4); )==" "\n" |
| 832 | R"==(STORE_DST(C01, C11, C21, C31, 5); )==" "\n" |
| 833 | R"==(STORE_DST(C01, C11, C21, C31, 6); )==" "\n" |
| 834 | R"==(STORE_DST(C01, C11, C21, C31, 7); )==" "\n" |
| 835 | R"==(#else )==" "\n" |
| 836 | R"==(STORE_DST_4(C01, C11, C21, C31, 4); )==" "\n" |
| 837 | R"==(#endif )==" "\n" |
| 838 | R"==(#endif )==" "\n" |
| 839 | R"==(#endif )==" "\n" |
| 840 | R"==(#endif )==" "\n" |
| 841 | R"==(#endif )==" "\n" |
| 842 | R"==(} )==" "\n" |
| 843 | R"==(#endif )==" "\n" |
| 844 | R"==(} )==" "\n" |
| 845 | R"==(} )==" "\n" |
| 846 | R"==()=="; |
| 847 | } |
| 848 | } |
| 849 | } |
| 850 | } |
Generated on 2022-Nov-30 from project oneDNN revision 84bc949a1
Powered by 
Code Browser 2.1
Generator usage only permitted with license.