| 1 | namespace dnnl { |
|---|---|
| 2 | namespace impl { |
| 3 | namespace gpu { |
| 4 | namespace ocl { |
| 5 | const char *gen9_gemm_compute_kernel = R"==(/******************************************************************************* )==" "\n" |
| 6 | R"==(* Copyright 2019-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_post_ops.h" )==" "\n" |
| 21 | R"==(#include "gpu/ocl/ocl_types.h" )==" "\n" |
| 22 | R"==(#define GRX 8 )==" "\n" |
| 23 | R"==(#if DT_F32 == 1 )==" "\n" |
| 24 | R"==(#if UNROLL_M <= 1 * GRX )==" "\n" |
| 25 | R"==(#define FLOATX float )==" "\n" |
| 26 | R"==(#define SIZEX 1 )==" "\n" |
| 27 | R"==(#elif UNROLL_M <= 2 * GRX )==" "\n" |
| 28 | R"==(#define FLOATX float2 )==" "\n" |
| 29 | R"==(#define SIZEX 2 )==" "\n" |
| 30 | R"==(#elif UNROLL_M <= 3 * GRX )==" "\n" |
| 31 | R"==(#define FLOATX float3 )==" "\n" |
| 32 | R"==(#define SIZEX 3 )==" "\n" |
| 33 | R"==(#else )==" "\n" |
| 34 | R"==(#define FLOATX float4 )==" "\n" |
| 35 | R"==(#define SIZEX 4 )==" "\n" |
| 36 | R"==(#endif )==" "\n" |
| 37 | R"==(#if UNROLL_N <= 1 * GRX )==" "\n" |
| 38 | R"==(#define FLOATY float )==" "\n" |
| 39 | R"==(#define SIZEY 1 )==" "\n" |
| 40 | R"==(#elif UNROLL_N <= 2 * GRX )==" "\n" |
| 41 | R"==(#define FLOATY float2 )==" "\n" |
| 42 | R"==(#define SIZEY 2 )==" "\n" |
| 43 | R"==(#elif UNROLL_N <= 3 * GRX )==" "\n" |
| 44 | R"==(#define FLOATY float3 )==" "\n" |
| 45 | R"==(#define SIZEY 3 )==" "\n" |
| 46 | R"==(#else )==" "\n" |
| 47 | R"==(#define FLOATY float4 )==" "\n" |
| 48 | R"==(#define SIZEY 4 )==" "\n" |
| 49 | R"==(#endif )==" "\n" |
| 50 | R"==(#define SHUFFLE(X, Y) intel_sub_group_shuffle(X, Y) )==" "\n" |
| 51 | R"==(#define SHUFFLE_DOWN(X, Y) intel_sub_group_shuffle_down(X, X, Y) )==" "\n" |
| 52 | R"==(#define SHUFFLE_UP(X, Y) intel_sub_group_shuffle_up(X, X, Y) )==" "\n" |
| 53 | R"==(#elif DT_F16 == 1 )==" "\n" |
| 54 | R"==(#if UNROLL_M <= 1 * GRX )==" "\n" |
| 55 | R"==(#define FLOATX half )==" "\n" |
| 56 | R"==(#define SIZEX 1 )==" "\n" |
| 57 | R"==(#elif UNROLL_M <= 2 * GRX )==" "\n" |
| 58 | R"==(#define FLOATX half2 )==" "\n" |
| 59 | R"==(#define SIZEX 2 )==" "\n" |
| 60 | R"==(#elif UNROLL_M <= 3 * GRX )==" "\n" |
| 61 | R"==(#define FLOATX half3 )==" "\n" |
| 62 | R"==(#define SIZEX 3 )==" "\n" |
| 63 | R"==(#elif UNROLL_M <= 4 * GRX )==" "\n" |
| 64 | R"==(#define FLOATX half4 )==" "\n" |
| 65 | R"==(#define SIZEX 4 )==" "\n" |
| 66 | R"==(#else )==" "\n" |
| 67 | R"==(#define FLOATX half8 )==" "\n" |
| 68 | R"==(#define SIZEX 8 )==" "\n" |
| 69 | R"==(#endif )==" "\n" |
| 70 | R"==(#if UNROLL_N <= 1 * GRX )==" "\n" |
| 71 | R"==(#define FLOATY half )==" "\n" |
| 72 | R"==(#define SIZEY 1 )==" "\n" |
| 73 | R"==(#elif UNROLL_N <= 2 * GRX )==" "\n" |
| 74 | R"==(#define FLOATY half2 )==" "\n" |
| 75 | R"==(#define SIZEY 2 )==" "\n" |
| 76 | R"==(#elif UNROLL_N <= 3 * GRX )==" "\n" |
| 77 | R"==(#define FLOATY half3 )==" "\n" |
| 78 | R"==(#define SIZEY 3 )==" "\n" |
| 79 | R"==(#elif UNROLL_N <= 4 * GRX )==" "\n" |
| 80 | R"==(#define FLOATY half4 )==" "\n" |
| 81 | R"==(#define SIZEY 4 )==" "\n" |
| 82 | R"==(#else )==" "\n" |
| 83 | R"==(#define FLOATY half8 )==" "\n" |
| 84 | R"==(#define SIZEY 8 )==" "\n" |
| 85 | R"==(#endif )==" "\n" |
| 86 | R"==(#if SIZEY == 2 )==" "\n" |
| 87 | R"==(#define SHUFFLE(X, Y) as_half2(intel_sub_group_shuffle(as_float(X), Y)) )==" "\n" |
| 88 | R"==(#elif SIZEY == 4 )==" "\n" |
| 89 | R"==(#define SHUFFLE(X, Y) as_half4(intel_sub_group_shuffle(as_float2(X), Y)) )==" "\n" |
| 90 | R"==(#else )==" "\n" |
| 91 | R"==(#define SHUFFLE(X, Y) as_half8(intel_sub_group_shuffle(as_float4(X), Y)) )==" "\n" |
| 92 | R"==(#endif )==" "\n" |
| 93 | R"==(#if SIZEX == 2 )==" "\n" |
| 94 | R"==(#define SHUFFLE_UP(X, Y) \ )==" "\n" |
| 95 | R"==(as_half2(intel_sub_group_shuffle_up(as_float(X), as_float(X), Y)) )==" "\n" |
| 96 | R"==(#define SHUFFLE_DOWN(X, Y) \ )==" "\n" |
| 97 | R"==(as_half2(intel_sub_group_shuffle_down(as_float(X), as_float(X), Y)) )==" "\n" |
| 98 | R"==(#elif SIZEX == 4 )==" "\n" |
| 99 | R"==(#define SHUFFLE_UP(X, Y) \ )==" "\n" |
| 100 | R"==(as_half4(intel_sub_group_shuffle_up(as_float2(X), as_float2(X), Y)) )==" "\n" |
| 101 | R"==(#define SHUFFLE_DOWN(X, Y) \ )==" "\n" |
| 102 | R"==(as_half4(intel_sub_group_shuffle_down(as_float2(X), as_float2(X), Y)) )==" "\n" |
| 103 | R"==(#else )==" "\n" |
| 104 | R"==(#define SHUFFLE_UP(X, Y) \ )==" "\n" |
| 105 | R"==(as_half8(intel_sub_group_shuffle_up(as_float4(X), as_float4(X), Y)) )==" "\n" |
| 106 | R"==(#define SHUFFLE_DOWN(X, Y) \ )==" "\n" |
| 107 | R"==(as_half8(intel_sub_group_shuffle_down(as_float4(X), as_float4(X), Y)) )==" "\n" |
| 108 | R"==(#endif )==" "\n" |
| 109 | R"==(#endif )==" "\n" |
| 110 | R"==(#define AS_FLOATX(X, Y) *((__global FLOATX *)(X + Y)) )==" "\n" |
| 111 | R"==(#define AS_FLOATY(X, Y) *((__global FLOATY *)(X + Y)) )==" "\n" |
| 112 | R"==(#if UNROLL_M <= 1 * GRX )==" "\n" |
| 113 | R"==(#define CALC_X(x, a, b, R0, R1, R2, R3) \ )==" "\n" |
| 114 | R"==(bb = SHUFFLE(b, x); \ )==" "\n" |
| 115 | R"==(R0##x = mad(a, bb, R0##x); )==" "\n" |
| 116 | R"==(#elif UNROLL_M <= 2 * GRX )==" "\n" |
| 117 | R"==(#define CALC_X(x, a, b, R0, R1, R2, R3) \ )==" "\n" |
| 118 | R"==(bb = SHUFFLE(b, x); \ )==" "\n" |
| 119 | R"==(R0##x = mad(a.s0, bb, R0##x); \ )==" "\n" |
| 120 | R"==(R1##x = mad(a.s1, bb, R1##x); )==" "\n" |
| 121 | R"==(#elif UNROLL_M <= 3 * GRX )==" "\n" |
| 122 | R"==(#define CALC_X(x, a, b, R0, R1, R2, R3) \ )==" "\n" |
| 123 | R"==(bb = SHUFFLE(b, x); \ )==" "\n" |
| 124 | R"==(R0##x = mad(a.s0, bb, R0##x); \ )==" "\n" |
| 125 | R"==(R1##x = mad(a.s1, bb, R1##x); \ )==" "\n" |
| 126 | R"==(R2##x = mad(a.s2, bb, R2##x); )==" "\n" |
| 127 | R"==(#else )==" "\n" |
| 128 | R"==(#define CALC_X(x, a, b, R0, R1, R2, R3) \ )==" "\n" |
| 129 | R"==(bb = SHUFFLE(b, x); \ )==" "\n" |
| 130 | R"==(R0##x = mad(a.s0, bb, R0##x); \ )==" "\n" |
| 131 | R"==(R1##x = mad(a.s1, bb, R1##x); \ )==" "\n" |
| 132 | R"==(R2##x = mad(a.s2, bb, R2##x); \ )==" "\n" |
| 133 | R"==(R3##x = mad(a.s3, bb, R3##x); )==" "\n" |
| 134 | R"==(#endif )==" "\n" |
| 135 | R"==(#define CALC(a, b, R0, R1, R2, R3) \ )==" "\n" |
| 136 | R"==(CALC_X(0, a, b, R0, R1, R2, R3); \ )==" "\n" |
| 137 | R"==(CALC_X(1, a, b, R0, R1, R2, R3); \ )==" "\n" |
| 138 | R"==(CALC_X(2, a, b, R0, R1, R2, R3); \ )==" "\n" |
| 139 | R"==(CALC_X(3, a, b, R0, R1, R2, R3); \ )==" "\n" |
| 140 | R"==(CALC_X(4, a, b, R0, R1, R2, R3); \ )==" "\n" |
| 141 | R"==(CALC_X(5, a, b, R0, R1, R2, R3); \ )==" "\n" |
| 142 | R"==(CALC_X(6, a, b, R0, R1, R2, R3); \ )==" "\n" |
| 143 | R"==(CALC_X(7, a, b, R0, R1, R2, R3); )==" "\n" |
| 144 | R"==(#define INIT_C(n) \ )==" "\n" |
| 145 | R"==(FLOATY cc##n##0 = DATA_ZERO, cc##n##1 = DATA_ZERO; \ )==" "\n" |
| 146 | R"==(FLOATY cc##n##2 = DATA_ZERO, cc##n##3 = DATA_ZERO; \ )==" "\n" |
| 147 | R"==(FLOATY cc##n##4 = DATA_ZERO, cc##n##5 = DATA_ZERO; \ )==" "\n" |
| 148 | R"==(FLOATY cc##n##6 = DATA_ZERO, cc##n##7 = DATA_ZERO; )==" "\n" |
| 149 | R"==(#if WITH_ELTWISE == 1 )==" "\n" |
| 150 | R"==(#define POST_OP(val) \ )==" "\n" |
| 151 | R"==(do { \ )==" "\n" |
| 152 | R"==(if (last_k_block) \ )==" "\n" |
| 153 | R"==(val = fwd_eltwise( \ )==" "\n" |
| 154 | R"==(val, eltwise_alpha, eltwise_beta, eltwise_scale); \ )==" "\n" |
| 155 | R"==(} while (0) )==" "\n" |
| 156 | R"==(#else )==" "\n" |
| 157 | R"==(#define POST_OP(val) )==" "\n" |
| 158 | R"==(#endif )==" "\n" |
| 159 | R"==(#ifdef BETA_ZERO )==" "\n" |
| 160 | R"==(#define UPDATE(c, acc) \ )==" "\n" |
| 161 | R"==(do { \ )==" "\n" |
| 162 | R"==(DATA_T val = acc; \ )==" "\n" |
| 163 | R"==(POST_OP(val); \ )==" "\n" |
| 164 | R"==(c = REF_TO_DST(val); \ )==" "\n" |
| 165 | R"==(} while (0) )==" "\n" |
| 166 | R"==(#else )==" "\n" |
| 167 | R"==(#define UPDATE(c, acc) \ )==" "\n" |
| 168 | R"==(do { \ )==" "\n" |
| 169 | R"==(DATA_T val = DST_TO_REF(c) + acc; \ )==" "\n" |
| 170 | R"==(POST_OP(val); \ )==" "\n" |
| 171 | R"==(c = REF_TO_DST(val); \ )==" "\n" |
| 172 | R"==(} while (0) )==" "\n" |
| 173 | R"==(#endif )==" "\n" |
| 174 | R"==(#if SIZEX == 1 )==" "\n" |
| 175 | R"==(#define UPDATE_YY(X, Y, R0, R1, R2, R3) \ )==" "\n" |
| 176 | R"==(if (n > (Y)) { \ )==" "\n" |
| 177 | R"==(if ((m > 0)) { UPDATE(c[offsetC + 0], R0); } \ )==" "\n" |
| 178 | R"==(offsetC += ldc; \ )==" "\n" |
| 179 | R"==(} )==" "\n" |
| 180 | R"==(#elif SIZEX == 2 )==" "\n" |
| 181 | R"==(#define UPDATE_YY(X, Y, R0, R1, R2, R3) \ )==" "\n" |
| 182 | R"==(if (n > (Y)) { \ )==" "\n" |
| 183 | R"==(if ((m > 0)) { UPDATE(c[offsetC + 0], R0); } \ )==" "\n" |
| 184 | R"==(if ((m > 1)) { UPDATE(c[offsetC + 1], R1); } \ )==" "\n" |
| 185 | R"==(offsetC += ldc; \ )==" "\n" |
| 186 | R"==(} )==" "\n" |
| 187 | R"==(#elif SIZEX == 3 )==" "\n" |
| 188 | R"==(#define UPDATE_YY(X, Y, R0, R1, R2, R3) \ )==" "\n" |
| 189 | R"==(if (n > (Y)) { \ )==" "\n" |
| 190 | R"==(if ((m > 0)) { UPDATE(c[offsetC + 0], R0); } \ )==" "\n" |
| 191 | R"==(if ((m > 1)) { UPDATE(c[offsetC + 1], R1); } \ )==" "\n" |
| 192 | R"==(if ((m > 2)) { UPDATE(c[offsetC + 2], R2); } \ )==" "\n" |
| 193 | R"==(offsetC += ldc; \ )==" "\n" |
| 194 | R"==(} )==" "\n" |
| 195 | R"==(#else )==" "\n" |
| 196 | R"==(#define UPDATE_YY(X, Y, R0, R1, R2, R3) \ )==" "\n" |
| 197 | R"==(if (n > (Y)) { \ )==" "\n" |
| 198 | R"==(if ((m > 0)) { UPDATE(c[offsetC + 0], R0); } \ )==" "\n" |
| 199 | R"==(if ((m > 1)) { UPDATE(c[offsetC + 1], R1); } \ )==" "\n" |
| 200 | R"==(if ((m > 2)) { UPDATE(c[offsetC + 2], R2); } \ )==" "\n" |
| 201 | R"==(if ((m > 3)) { UPDATE(c[offsetC + 3], R3); } \ )==" "\n" |
| 202 | R"==(offsetC += ldc; \ )==" "\n" |
| 203 | R"==(} )==" "\n" |
| 204 | R"==(#endif )==" "\n" |
| 205 | R"==(#if SIZEY == 1 )==" "\n" |
| 206 | R"==(#define UPDATE_Y(X, R0, R1, R2, R3) \ )==" "\n" |
| 207 | R"==(UPDATE_YY(X, X *SIZEY + 0, R0##X, R1##X, R2##X, R3##X); )==" "\n" |
| 208 | R"==(#elif SIZEY == 2 )==" "\n" |
| 209 | R"==(#define UPDATE_Y(X, R0, R1, R2, R3) \ )==" "\n" |
| 210 | R"==(UPDATE_YY(X, X *SIZEY + 0, R0##X.s0, R1##X.s0, R2##X.s0, R3##X.s0); \ )==" "\n" |
| 211 | R"==(UPDATE_YY(X, X *SIZEY + 1, R0##X.s1, R1##X.s1, R2##X.s1, R3##X.s1); )==" "\n" |
| 212 | R"==(#elif SIZEY == 3 )==" "\n" |
| 213 | R"==(#define UPDATE_Y(X, R0, R1, R2, R3) \ )==" "\n" |
| 214 | R"==(UPDATE_YY(X, X *SIZEY + 0, R0##X.s0, R1##X.s0, R2##X.s0, R3##X.s0); \ )==" "\n" |
| 215 | R"==(UPDATE_YY(X, X *SIZEY + 1, R0##X.s1, R1##X.s1, R2##X.s1, R3##X.s1); \ )==" "\n" |
| 216 | R"==(UPDATE_YY(X, X *SIZEY + 2, R0##X.s2, R1##X.s2, R2##X.s2, R3##X.s2); )==" "\n" |
| 217 | R"==(#else )==" "\n" |
| 218 | R"==(#define UPDATE_Y(X, R0, R1, R2, R3) \ )==" "\n" |
| 219 | R"==(UPDATE_YY(X, X *SIZEY + 0, R0##X.s0, R1##X.s0, R2##X.s0, R3##X.s0); \ )==" "\n" |
| 220 | R"==(UPDATE_YY(X, X *SIZEY + 1, R0##X.s1, R1##X.s1, R2##X.s1, R3##X.s1); \ )==" "\n" |
| 221 | R"==(UPDATE_YY(X, X *SIZEY + 2, R0##X.s2, R1##X.s2, R2##X.s2, R3##X.s2); \ )==" "\n" |
| 222 | R"==(UPDATE_YY(X, X *SIZEY + 3, R0##X.s3, R1##X.s3, R2##X.s3, R3##X.s3); )==" "\n" |
| 223 | R"==(#endif )==" "\n" |
| 224 | R"==(__attribute__((intel_reqd_sub_group_size(GRX))) __kernel void gen9_gemm_compute( )==" "\n" |
| 225 | R"==(long m, long n, long k, __global DATA_T *base, int offsetA, int offsetB, )==" "\n" |
| 226 | R"==(__global DST_DATA_T *c, long offsetC, long ldc, int last_k_block, )==" "\n" |
| 227 | R"==(float eltwise_alpha, float eltwise_beta, float eltwise_scale) { )==" "\n" |
| 228 | R"==(int idx, idy, lid; )==" "\n" |
| 229 | R"==(idx = get_group_id(0); )==" "\n" |
| 230 | R"==(idy = get_group_id(1) * get_enqueued_local_size(1) + get_local_id(1); )==" "\n" |
| 231 | R"==(lid = get_local_id(0); )==" "\n" |
| 232 | R"==(m -= UNROLL_M * idx; )==" "\n" |
| 233 | R"==(if (m > UNROLL_M) m = UNROLL_M; )==" "\n" |
| 234 | R"==(n -= UNROLL_N * idy; )==" "\n" |
| 235 | R"==(if (n > UNROLL_N) n = UNROLL_N; )==" "\n" |
| 236 | R"==(m -= UNROLL_M * lid / GRX; )==" "\n" |
| 237 | R"==(offsetA += UNROLL_M * k * idx + UNROLL_M * lid / GRX; )==" "\n" |
| 238 | R"==(offsetB += UNROLL_N * k * idy + UNROLL_N * lid / GRX; )==" "\n" |
| 239 | R"==(offsetC += UNROLL_M * idx + UNROLL_N * ldc * idy + UNROLL_M * lid / GRX; )==" "\n" |
| 240 | R"==(INIT_C(0); )==" "\n" |
| 241 | R"==(INIT_C(1); )==" "\n" |
| 242 | R"==(INIT_C(2); )==" "\n" |
| 243 | R"==(INIT_C(3); )==" "\n" |
| 244 | R"==(INIT_C(4); )==" "\n" |
| 245 | R"==(INIT_C(5); )==" "\n" |
| 246 | R"==(INIT_C(6); )==" "\n" |
| 247 | R"==(INIT_C(7); )==" "\n" |
| 248 | R"==(FLOATX blockA = AS_FLOATX(base, offsetA); )==" "\n" |
| 249 | R"==(offsetA += UNROLL_M; )==" "\n" |
| 250 | R"==(FLOATY blockB = AS_FLOATY(base, offsetB); )==" "\n" |
| 251 | R"==(offsetB += UNROLL_N; )==" "\n" |
| 252 | R"==(for (int l = k; l > 0; l--) { )==" "\n" |
| 253 | R"==(FLOATY bb; )==" "\n" |
| 254 | R"==(CALC(blockA, blockB, cc0, cc1, cc2, cc3); )==" "\n" |
| 255 | R"==(blockB = AS_FLOATY(base, offsetB); )==" "\n" |
| 256 | R"==(offsetB += UNROLL_N; )==" "\n" |
| 257 | R"==(blockA = AS_FLOATX(base, offsetA); )==" "\n" |
| 258 | R"==(offsetA += UNROLL_M; )==" "\n" |
| 259 | R"==(} )==" "\n" |
| 260 | R"==(UPDATE_Y(0, cc0, cc1, cc2, cc3); )==" "\n" |
| 261 | R"==(#if UNROLL_N >= 2 * SIZEY )==" "\n" |
| 262 | R"==(UPDATE_Y(1, cc0, cc1, cc2, cc3); )==" "\n" |
| 263 | R"==(#endif )==" "\n" |
| 264 | R"==(#if UNROLL_N >= 3 * SIZEY )==" "\n" |
| 265 | R"==(UPDATE_Y(2, cc0, cc1, cc2, cc3); )==" "\n" |
| 266 | R"==(#endif )==" "\n" |
| 267 | R"==(#if UNROLL_N >= 4 * SIZEY )==" "\n" |
| 268 | R"==(UPDATE_Y(3, cc0, cc1, cc2, cc3); )==" "\n" |
| 269 | R"==(#endif )==" "\n" |
| 270 | R"==(#if UNROLL_N >= 5 * SIZEY )==" "\n" |
| 271 | R"==(UPDATE_Y(4, cc0, cc1, cc2, cc3); )==" "\n" |
| 272 | R"==(#endif )==" "\n" |
| 273 | R"==(#if UNROLL_N >= 6 * SIZEY )==" "\n" |
| 274 | R"==(UPDATE_Y(5, cc0, cc1, cc2, cc3); )==" "\n" |
| 275 | R"==(#endif )==" "\n" |
| 276 | R"==(#if UNROLL_N >= 7 * SIZEY )==" "\n" |
| 277 | R"==(UPDATE_Y(6, cc0, cc1, cc2, cc3); )==" "\n" |
| 278 | R"==(#endif )==" "\n" |
| 279 | R"==(#if UNROLL_N >= 8 * SIZEY )==" "\n" |
| 280 | R"==(UPDATE_Y(7, cc0, cc1, cc2, cc3); )==" "\n" |
| 281 | R"==(#endif )==" "\n" |
| 282 | R"==(} )==" "\n" |
| 283 | R"==()=="; |
| 284 | } |
| 285 | } |
| 286 | } |
| 287 | } |
Generated on 2022-Nov-30 from project oneDNN revision 84bc949a1
Powered by 
Code Browser 2.1
Generator usage only permitted with license.