875 {
876 const ExpressionArray& arguments = c.arguments();
877 switch (kind) {
878 case k_textureRead_IntrinsicKind: {
880 this->
write(
".read(");
883 return true;
884 }
885 case k_textureWrite_IntrinsicKind: {
887 this->
write(
".write(");
892 return true;
893 }
894 case k_textureWidth_IntrinsicKind: {
896 this->
write(
".get_width()");
897 return true;
898 }
899 case k_textureHeight_IntrinsicKind: {
901 this->
write(
".get_height()");
902 return true;
903 }
904 case k_mod_IntrinsicKind: {
905
908 this->
write(
"(" + tmpX +
" = ");
910 this->
write(
", " + tmpY +
" = ");
912 this->
write(
", " + tmpX +
" - " + tmpY +
" * floor(" + tmpX +
" / " + tmpY +
"))");
913 return true;
914 }
915
916 case k_distance_IntrinsicKind: {
917 if (arguments[0]->
type().columns() == 1) {
923 } else {
925 }
926 return true;
927 }
928 case k_dot_IntrinsicKind: {
929 if (arguments[0]->
type().columns() == 1) {
935 } else {
937 }
938 return true;
939 }
940 case k_faceforward_IntrinsicKind: {
941 if (arguments[0]->
type().columns() == 1) {
942
945 this->
write(
") * (");
947 this->
write(
") < 0) ? 1 : -1) * (");
950 } else {
952 }
953 return true;
954 }
955 case k_length_IntrinsicKind: {
956 this->
write(arguments[0]->
type().columns() == 1 ?
"abs(" :
"length(");
959 return true;
960 }
961 case k_normalize_IntrinsicKind: {
962 this->
write(arguments[0]->
type().columns() == 1 ?
"sign(" :
"normalize(");
965 return true;
966 }
967 case k_packUnorm2x16_IntrinsicKind: {
968 this->
write(
"pack_float_to_unorm2x16(");
971 return true;
972 }
973 case k_unpackUnorm2x16_IntrinsicKind: {
974 this->
write(
"unpack_unorm2x16_to_float(");
977 return true;
978 }
979 case k_packSnorm2x16_IntrinsicKind: {
980 this->
write(
"pack_float_to_snorm2x16(");
983 return true;
984 }
985 case k_unpackSnorm2x16_IntrinsicKind: {
986 this->
write(
"unpack_snorm2x16_to_float(");
989 return true;
990 }
991 case k_packUnorm4x8_IntrinsicKind: {
992 this->
write(
"pack_float_to_unorm4x8(");
995 return true;
996 }
997 case k_unpackUnorm4x8_IntrinsicKind: {
998 this->
write(
"unpack_unorm4x8_to_float(");
1001 return true;
1002 }
1003 case k_packSnorm4x8_IntrinsicKind: {
1004 this->
write(
"pack_float_to_snorm4x8(");
1007 return true;
1008 }
1009 case k_unpackSnorm4x8_IntrinsicKind: {
1010 this->
write(
"unpack_snorm4x8_to_float(");
1013 return true;
1014 }
1015 case k_packHalf2x16_IntrinsicKind: {
1016 this->
write(
"as_type<uint>(half2(");
1019 return true;
1020 }
1021 case k_unpackHalf2x16_IntrinsicKind: {
1022 this->
write(
"float2(as_type<half2>(");
1025 return true;
1026 }
1027 case k_floatBitsToInt_IntrinsicKind:
1028 case k_floatBitsToUint_IntrinsicKind:
1029 case k_intBitsToFloat_IntrinsicKind:
1030 case k_uintBitsToFloat_IntrinsicKind: {
1035 return true;
1036 }
1037 case k_degrees_IntrinsicKind: {
1040 this->
write(
") * 57.2957795)");
1041 return true;
1042 }
1043 case k_radians_IntrinsicKind: {
1046 this->
write(
") * 0.0174532925)");
1047 return true;
1048 }
1049 case k_dFdx_IntrinsicKind: {
1050 this->
write(
"dfdx");
1052 return true;
1053 }
1054 case k_dFdy_IntrinsicKind: {
1057 } else {
1058 this->
write(
"(dfdy");
1059 }
1062 return true;
1063 }
1064 case k_inverse_IntrinsicKind: {
1067 return true;
1068 }
1069 case k_inversesqrt_IntrinsicKind: {
1070 this->
write(
"rsqrt");
1072 return true;
1073 }
1074 case k_atan_IntrinsicKind: {
1075 this->
write(c.arguments().size() == 2 ?
"atan2" :
"atan");
1077 return true;
1078 }
1079 case k_reflect_IntrinsicKind: {
1080 if (arguments[0]->
type().columns() == 1) {
1081
1084
1085
1086 this->
write(
"(" + tmpI +
" = ");
1088
1089
1090 this->
write(
", " + tmpN +
" = ");
1092
1093
1094 this->
write(
", " + tmpI +
" - 2 * " + tmpN +
" * " + tmpI +
" * " + tmpN +
")");
1095 } else {
1097 }
1098 return true;
1099 }
1100 case k_refract_IntrinsicKind: {
1101 if (arguments[0]->
type().columns() == 1) {
1102
1103
1104 this->
write(
"(refract(float2(");
1106 this->
write(
", 0), float2(");
1108 this->
write(
", 0), ");
1110 this->
write(
").x)");
1111 } else {
1113 }
1114 return true;
1115 }
1116 case k_roundEven_IntrinsicKind: {
1117 this->
write(
"rint");
1119 return true;
1120 }
1121 case k_bitCount_IntrinsicKind: {
1122 this->
write(
"popcount(");
1125 return true;
1126 }
1127 case k_findLSB_IntrinsicKind: {
1128
1130 std::string exprType = this->
typeName(arguments[0]->
type());
1131
1132
1133
1134
1135
1137 this->
write(skTemp);
1138 this->
write(
" = (");
1140 this->
write(
"), select(ctz(");
1141 this->
write(skTemp);
1143 this->
write(exprType);
1144 this->
write(
"(-1), ");
1145 this->
write(skTemp);
1146 this->
write(
" == ");
1147 this->
write(exprType);
1148 this->
write(
"(0)))");
1149 return true;
1150 }
1151 case k_findMSB_IntrinsicKind: {
1152
1154 std::string exprType = this->
typeName(arguments[0]->
type());
1155
1156
1157
1158
1159
1160
1162 this->
write(skTemp1);
1163 this->
write(
" = (");
1166
1167
1168
1169 std::string skTemp2;
1170 if (arguments[0]->
type().isSigned()) {
1171
1173 this->
write(skTemp2);
1174 this->
write(
" = (select(");
1175 this->
write(skTemp1);
1177 this->
write(skTemp1);
1179 this->
write(skTemp1);
1180 this->
write(
" < 0)), ");
1181 } else {
1182 skTemp2 = skTemp1;
1183 }
1184
1185
1186 this->
write(
"select(");
1188 this->
write(
"(clz(");
1189 this->
write(skTemp2);
1190 this->
write(
")), ");
1192 this->
write(
"(-1), ");
1193 this->
write(skTemp2);
1194 this->
write(
" == ");
1195 this->
write(exprType);
1196 this->
write(
"(0)))");
1197 return true;
1198 }
1199 case k_sign_IntrinsicKind: {
1200 if (arguments[0]->
type().componentType().isInteger()) {
1201
1203 std::string exprType = this->
typeName(arguments[0]->
type());
1204
1205
1207 this->
write(skTemp);
1208 this->
write(
" = (");
1211
1212
1213 this->
write(
"select(select(");
1214 this->
write(exprType);
1215 this->
write(
"(0), ");
1216 this->
write(exprType);
1217 this->
write(
"(-1), ");
1218 this->
write(skTemp);
1219 this->
write(
" < 0), ");
1220 this->
write(exprType);
1221 this->
write(
"(1), ");
1222 this->
write(skTemp);
1223 this->
write(
" > 0))");
1224 } else {
1226 }
1227 return true;
1228 }
1229 case k_matrixCompMult_IntrinsicKind: {
1232 return true;
1233 }
1234 case k_outerProduct_IntrinsicKind: {
1237 return true;
1238 }
1239 case k_mix_IntrinsicKind: {
1240 SkASSERT(c.arguments().size() == 3);
1241 if (arguments[2]->
type().componentType().isBoolean()) {
1242
1243 this->
write(
"select");
1245 return true;
1246 }
1247
1249 return true;
1250 }
1251 case k_equal_IntrinsicKind:
1252 case k_greaterThan_IntrinsicKind:
1253 case k_greaterThanEqual_IntrinsicKind:
1254 case k_lessThan_IntrinsicKind:
1255 case k_lessThanEqual_IntrinsicKind:
1256 case k_notEqual_IntrinsicKind: {
1259 switch (kind) {
1260 case k_equal_IntrinsicKind:
1261 this->
write(
" == ");
1262 break;
1263 case k_notEqual_IntrinsicKind:
1264 this->
write(
" != ");
1265 break;
1266 case k_lessThan_IntrinsicKind:
1268 break;
1269 case k_lessThanEqual_IntrinsicKind:
1270 this->
write(
" <= ");
1271 break;
1272 case k_greaterThan_IntrinsicKind:
1274 break;
1275 case k_greaterThanEqual_IntrinsicKind:
1276 this->
write(
" >= ");
1277 break;
1278 default:
1279 SK_ABORT(
"unsupported comparison intrinsic kind");
1280 }
1283 return true;
1284 }
1285 case k_storageBarrier_IntrinsicKind:
1286 this->
write(
"threadgroup_barrier(mem_flags::mem_device)");
1287 return true;
1288 case k_workgroupBarrier_IntrinsicKind:
1289 this->
write(
"threadgroup_barrier(mem_flags::mem_threadgroup)");
1290 return true;
1291 case k_atomicAdd_IntrinsicKind:
1292 this->
write(
"atomic_fetch_add_explicit(&");
1296 this->
write(
", memory_order_relaxed)");
1297 return true;
1298 case k_atomicLoad_IntrinsicKind:
1299 this->
write(
"atomic_load_explicit(&");
1301 this->
write(
", memory_order_relaxed)");
1302 return true;
1303 case k_atomicStore_IntrinsicKind:
1304 this->
write(
"atomic_store_explicit(&");
1308 this->
write(
", memory_order_relaxed)");
1309 return true;
1310 default:
1311 return false;
1312 }
1313}
#define SK_ABORT(message,...)