879 {
880 const ExpressionArray& arguments = c.arguments();
881 switch (kind) {
882 case k_textureRead_IntrinsicKind: {
884 this->
write(
".read(");
887 return true;
888 }
889 case k_textureWrite_IntrinsicKind: {
891 this->
write(
".write(");
896 return true;
897 }
898 case k_textureWidth_IntrinsicKind: {
900 this->
write(
".get_width()");
901 return true;
902 }
903 case k_textureHeight_IntrinsicKind: {
905 this->
write(
".get_height()");
906 return true;
907 }
908 case k_mod_IntrinsicKind: {
909
912 this->
write(
"(" + tmpX +
" = ");
914 this->
write(
", " + tmpY +
" = ");
916 this->
write(
", " + tmpX +
" - " + tmpY +
" * floor(" + tmpX +
" / " + tmpY +
"))");
917 return true;
918 }
919
920 case k_distance_IntrinsicKind: {
921 if (arguments[0]->
type().columns() == 1) {
927 } else {
929 }
930 return true;
931 }
932 case k_dot_IntrinsicKind: {
933 if (arguments[0]->
type().columns() == 1) {
939 } else {
941 }
942 return true;
943 }
944 case k_faceforward_IntrinsicKind: {
945 if (arguments[0]->
type().columns() == 1) {
946
949 this->
write(
") * (");
951 this->
write(
") < 0) ? 1 : -1) * (");
954 } else {
956 }
957 return true;
958 }
959 case k_length_IntrinsicKind: {
960 this->
write(arguments[0]->
type().columns() == 1 ?
"abs(" :
"length(");
963 return true;
964 }
965 case k_normalize_IntrinsicKind: {
966 this->
write(arguments[0]->
type().columns() == 1 ?
"sign(" :
"normalize(");
969 return true;
970 }
971 case k_packUnorm2x16_IntrinsicKind: {
972 this->
write(
"pack_float_to_unorm2x16(");
975 return true;
976 }
977 case k_unpackUnorm2x16_IntrinsicKind: {
978 this->
write(
"unpack_unorm2x16_to_float(");
981 return true;
982 }
983 case k_packSnorm2x16_IntrinsicKind: {
984 this->
write(
"pack_float_to_snorm2x16(");
987 return true;
988 }
989 case k_unpackSnorm2x16_IntrinsicKind: {
990 this->
write(
"unpack_snorm2x16_to_float(");
993 return true;
994 }
995 case k_packUnorm4x8_IntrinsicKind: {
996 this->
write(
"pack_float_to_unorm4x8(");
999 return true;
1000 }
1001 case k_unpackUnorm4x8_IntrinsicKind: {
1002 this->
write(
"unpack_unorm4x8_to_float(");
1005 return true;
1006 }
1007 case k_packSnorm4x8_IntrinsicKind: {
1008 this->
write(
"pack_float_to_snorm4x8(");
1011 return true;
1012 }
1013 case k_unpackSnorm4x8_IntrinsicKind: {
1014 this->
write(
"unpack_snorm4x8_to_float(");
1017 return true;
1018 }
1019 case k_packHalf2x16_IntrinsicKind: {
1020 this->
write(
"as_type<uint>(half2(");
1023 return true;
1024 }
1025 case k_unpackHalf2x16_IntrinsicKind: {
1026 this->
write(
"float2(as_type<half2>(");
1029 return true;
1030 }
1031 case k_floatBitsToInt_IntrinsicKind:
1032 case k_floatBitsToUint_IntrinsicKind:
1033 case k_intBitsToFloat_IntrinsicKind:
1034 case k_uintBitsToFloat_IntrinsicKind: {
1039 return true;
1040 }
1041 case k_degrees_IntrinsicKind: {
1044 this->
write(
") * 57.2957795)");
1045 return true;
1046 }
1047 case k_radians_IntrinsicKind: {
1050 this->
write(
") * 0.0174532925)");
1051 return true;
1052 }
1053 case k_dFdx_IntrinsicKind: {
1054 this->
write(
"dfdx");
1056 return true;
1057 }
1058 case k_dFdy_IntrinsicKind: {
1061 } else {
1062 this->
write(
"(dfdy");
1063 }
1066 return true;
1067 }
1068 case k_inverse_IntrinsicKind: {
1071 return true;
1072 }
1073 case k_inversesqrt_IntrinsicKind: {
1074 this->
write(
"rsqrt");
1076 return true;
1077 }
1078 case k_atan_IntrinsicKind: {
1079 this->
write(c.arguments().size() == 2 ?
"atan2" :
"atan");
1081 return true;
1082 }
1083 case k_reflect_IntrinsicKind: {
1084 if (arguments[0]->
type().columns() == 1) {
1085
1088
1089
1090 this->
write(
"(" + tmpI +
" = ");
1092
1093
1094 this->
write(
", " + tmpN +
" = ");
1096
1097
1098 this->
write(
", " + tmpI +
" - 2 * " + tmpN +
" * " + tmpI +
" * " + tmpN +
")");
1099 } else {
1101 }
1102 return true;
1103 }
1104 case k_refract_IntrinsicKind: {
1105 if (arguments[0]->
type().columns() == 1) {
1106
1107
1108 this->
write(
"(refract(float2(");
1110 this->
write(
", 0), float2(");
1112 this->
write(
", 0), ");
1114 this->
write(
").x)");
1115 } else {
1117 }
1118 return true;
1119 }
1120 case k_roundEven_IntrinsicKind: {
1121 this->
write(
"rint");
1123 return true;
1124 }
1125 case k_bitCount_IntrinsicKind: {
1126 this->
write(
"popcount(");
1129 return true;
1130 }
1131 case k_findLSB_IntrinsicKind: {
1132
1134 std::string exprType = this->
typeName(arguments[0]->
type());
1135
1136
1137
1138
1139
1141 this->
write(skTemp);
1142 this->
write(
" = (");
1144 this->
write(
"), select(ctz(");
1145 this->
write(skTemp);
1147 this->
write(exprType);
1148 this->
write(
"(-1), ");
1149 this->
write(skTemp);
1150 this->
write(
" == ");
1151 this->
write(exprType);
1152 this->
write(
"(0)))");
1153 return true;
1154 }
1155 case k_findMSB_IntrinsicKind: {
1156
1158 std::string exprType = this->
typeName(arguments[0]->
type());
1159
1160
1161
1162
1163
1164
1166 this->
write(skTemp1);
1167 this->
write(
" = (");
1170
1171
1172
1173 std::string skTemp2;
1174 if (arguments[0]->
type().isSigned()) {
1175
1177 this->
write(skTemp2);
1178 this->
write(
" = (select(");
1179 this->
write(skTemp1);
1181 this->
write(skTemp1);
1183 this->
write(skTemp1);
1184 this->
write(
" < 0)), ");
1185 } else {
1186 skTemp2 = skTemp1;
1187 }
1188
1189
1190 this->
write(
"select(");
1192 this->
write(
"(clz(");
1193 this->
write(skTemp2);
1194 this->
write(
")), ");
1196 this->
write(
"(-1), ");
1197 this->
write(skTemp2);
1198 this->
write(
" == ");
1199 this->
write(exprType);
1200 this->
write(
"(0)))");
1201 return true;
1202 }
1203 case k_sign_IntrinsicKind: {
1204 if (arguments[0]->
type().componentType().isInteger()) {
1205
1207 std::string exprType = this->
typeName(arguments[0]->
type());
1208
1209
1211 this->
write(skTemp);
1212 this->
write(
" = (");
1215
1216
1217 this->
write(
"select(select(");
1218 this->
write(exprType);
1219 this->
write(
"(0), ");
1220 this->
write(exprType);
1221 this->
write(
"(-1), ");
1222 this->
write(skTemp);
1223 this->
write(
" < 0), ");
1224 this->
write(exprType);
1225 this->
write(
"(1), ");
1226 this->
write(skTemp);
1227 this->
write(
" > 0))");
1228 } else {
1230 }
1231 return true;
1232 }
1233 case k_matrixCompMult_IntrinsicKind: {
1236 return true;
1237 }
1238 case k_outerProduct_IntrinsicKind: {
1241 return true;
1242 }
1243 case k_mix_IntrinsicKind: {
1244 SkASSERT(c.arguments().size() == 3);
1245 if (arguments[2]->
type().componentType().isBoolean()) {
1246
1247 this->
write(
"select");
1249 return true;
1250 }
1251
1253 return true;
1254 }
1255 case k_equal_IntrinsicKind:
1256 case k_greaterThan_IntrinsicKind:
1257 case k_greaterThanEqual_IntrinsicKind:
1258 case k_lessThan_IntrinsicKind:
1259 case k_lessThanEqual_IntrinsicKind:
1260 case k_notEqual_IntrinsicKind: {
1263 switch (kind) {
1264 case k_equal_IntrinsicKind:
1265 this->
write(
" == ");
1266 break;
1267 case k_notEqual_IntrinsicKind:
1268 this->
write(
" != ");
1269 break;
1270 case k_lessThan_IntrinsicKind:
1272 break;
1273 case k_lessThanEqual_IntrinsicKind:
1274 this->
write(
" <= ");
1275 break;
1276 case k_greaterThan_IntrinsicKind:
1278 break;
1279 case k_greaterThanEqual_IntrinsicKind:
1280 this->
write(
" >= ");
1281 break;
1282 default:
1283 SK_ABORT(
"unsupported comparison intrinsic kind");
1284 }
1287 return true;
1288 }
1289 case k_storageBarrier_IntrinsicKind:
1290 this->
write(
"threadgroup_barrier(mem_flags::mem_device)");
1291 return true;
1292 case k_workgroupBarrier_IntrinsicKind:
1293 this->
write(
"threadgroup_barrier(mem_flags::mem_threadgroup)");
1294 return true;
1295 case k_atomicAdd_IntrinsicKind:
1296 this->
write(
"atomic_fetch_add_explicit(&");
1300 this->
write(
", memory_order_relaxed)");
1301 return true;
1302 case k_atomicLoad_IntrinsicKind:
1303 this->
write(
"atomic_load_explicit(&");
1305 this->
write(
", memory_order_relaxed)");
1306 return true;
1307 case k_atomicStore_IntrinsicKind:
1308 this->
write(
"atomic_store_explicit(&");
1312 this->
write(
", memory_order_relaxed)");
1313 return true;
1314 case k_loadFloatBuffer_IntrinsicKind: {
1318 c.position(),
1319 c.arguments()[0]->clone());
1320
1322 return true;
1323 }
1324 default:
1325 return false;
1326 }
1327}
#define SK_ABORT(message,...)
static std::unique_ptr< Expression > LoadFloatBuffer(const Context &context, const SkSL::ShaderCaps &shaderCaps, Position position, std::unique_ptr< Expression > idx)