in src/sksl/codegen/SkSLMetalCodeGenerator.cpp [878:1326]
bool MetalCodeGenerator::writeIntrinsicCall(const FunctionCall& c, IntrinsicKind kind) {
const ExpressionArray& arguments = c.arguments();
switch (kind) {
case k_textureRead_IntrinsicKind: {
this->writeExpression(*arguments[0], Precedence::kExpression);
this->write(".read(");
this->writeExpression(*arguments[1], Precedence::kSequence);
this->write(")");
return true;
}
case k_textureWrite_IntrinsicKind: {
this->writeExpression(*arguments[0], Precedence::kExpression);
this->write(".write(");
this->writeExpression(*arguments[2], Precedence::kSequence);
this->write(", ");
this->writeExpression(*arguments[1], Precedence::kSequence);
this->write(")");
return true;
}
case k_textureWidth_IntrinsicKind: {
this->writeExpression(*arguments[0], Precedence::kExpression);
this->write(".get_width()");
return true;
}
case k_textureHeight_IntrinsicKind: {
this->writeExpression(*arguments[0], Precedence::kExpression);
this->write(".get_height()");
return true;
}
case k_mod_IntrinsicKind: {
// fmod(x, y) in metal calculates x - y * trunc(x / y) instead of x - y * floor(x / y)
std::string tmpX = this->getTempVariable(arguments[0]->type());
std::string tmpY = this->getTempVariable(arguments[1]->type());
this->write("(" + tmpX + " = ");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(", " + tmpY + " = ");
this->writeExpression(*arguments[1], Precedence::kSequence);
this->write(", " + tmpX + " - " + tmpY + " * floor(" + tmpX + " / " + tmpY + "))");
return true;
}
case k_pow_IntrinsicKind: {
// The Metal equivalent to GLSL's pow() is powr(). Metal's pow() allows negative base
// values, which is presumably more expensive to compute.
this->write("powr(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(", ");
this->writeExpression(*arguments[1], Precedence::kSequence);
this->write(")");
return true;
}
// GLSL declares scalar versions of most geometric intrinsics, but these don't exist in MSL
case k_distance_IntrinsicKind: {
if (arguments[0]->type().columns() == 1) {
this->write("abs(");
this->writeExpression(*arguments[0], Precedence::kAdditive);
this->write(" - ");
this->writeExpression(*arguments[1], Precedence::kAdditive);
this->write(")");
} else {
this->writeSimpleIntrinsic(c);
}
return true;
}
case k_dot_IntrinsicKind: {
if (arguments[0]->type().columns() == 1) {
this->write("(");
this->writeExpression(*arguments[0], Precedence::kMultiplicative);
this->write(" * ");
this->writeExpression(*arguments[1], Precedence::kMultiplicative);
this->write(")");
} else {
this->writeSimpleIntrinsic(c);
}
return true;
}
case k_faceforward_IntrinsicKind: {
if (arguments[0]->type().columns() == 1) {
// ((((Nref) * (I) < 0) ? 1 : -1) * (N))
this->write("((((");
this->writeExpression(*arguments[2], Precedence::kSequence);
this->write(") * (");
this->writeExpression(*arguments[1], Precedence::kSequence);
this->write(") < 0) ? 1 : -1) * (");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write("))");
} else {
this->writeSimpleIntrinsic(c);
}
return true;
}
case k_length_IntrinsicKind: {
this->write(arguments[0]->type().columns() == 1 ? "abs(" : "length(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(")");
return true;
}
case k_normalize_IntrinsicKind: {
this->write(arguments[0]->type().columns() == 1 ? "sign(" : "normalize(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(")");
return true;
}
case k_packUnorm2x16_IntrinsicKind: {
this->write("pack_float_to_unorm2x16(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(")");
return true;
}
case k_unpackUnorm2x16_IntrinsicKind: {
this->write("unpack_unorm2x16_to_float(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(")");
return true;
}
case k_packSnorm2x16_IntrinsicKind: {
this->write("pack_float_to_snorm2x16(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(")");
return true;
}
case k_unpackSnorm2x16_IntrinsicKind: {
this->write("unpack_snorm2x16_to_float(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(")");
return true;
}
case k_packUnorm4x8_IntrinsicKind: {
this->write("pack_float_to_unorm4x8(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(")");
return true;
}
case k_unpackUnorm4x8_IntrinsicKind: {
this->write("unpack_unorm4x8_to_float(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(")");
return true;
}
case k_packSnorm4x8_IntrinsicKind: {
this->write("pack_float_to_snorm4x8(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(")");
return true;
}
case k_unpackSnorm4x8_IntrinsicKind: {
this->write("unpack_snorm4x8_to_float(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(")");
return true;
}
case k_packHalf2x16_IntrinsicKind: {
this->write("as_type<uint>(half2(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write("))");
return true;
}
case k_unpackHalf2x16_IntrinsicKind: {
this->write("float2(as_type<half2>(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write("))");
return true;
}
case k_floatBitsToInt_IntrinsicKind:
case k_floatBitsToUint_IntrinsicKind:
case k_intBitsToFloat_IntrinsicKind:
case k_uintBitsToFloat_IntrinsicKind: {
this->write(this->getBitcastIntrinsic(c.type()));
this->write("(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(")");
return true;
}
case k_degrees_IntrinsicKind: {
this->write("((");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(") * 57.2957795)");
return true;
}
case k_radians_IntrinsicKind: {
this->write("((");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(") * 0.0174532925)");
return true;
}
case k_dFdx_IntrinsicKind: {
this->write("dfdx");
this->writeArgumentList(c.arguments());
return true;
}
case k_dFdy_IntrinsicKind: {
if (!fRTFlipName.empty()) {
this->write("(" + fRTFlipName + ".y * dfdy");
} else {
this->write("(dfdy");
}
this->writeArgumentList(c.arguments());
this->write(")");
return true;
}
case k_inverse_IntrinsicKind: {
this->write(this->getInversePolyfill(arguments));
this->writeArgumentList(c.arguments());
return true;
}
case k_inversesqrt_IntrinsicKind: {
this->write("rsqrt");
this->writeArgumentList(c.arguments());
return true;
}
case k_atan_IntrinsicKind: {
this->write(c.arguments().size() == 2 ? "atan2" : "atan");
this->writeArgumentList(c.arguments());
return true;
}
case k_reflect_IntrinsicKind: {
if (arguments[0]->type().columns() == 1) {
// We need to synthesize `I - 2 * N * I * N`.
std::string tmpI = this->getTempVariable(arguments[0]->type());
std::string tmpN = this->getTempVariable(arguments[1]->type());
// (_skTempI = ...
this->write("(" + tmpI + " = ");
this->writeExpression(*arguments[0], Precedence::kSequence);
// , _skTempN = ...
this->write(", " + tmpN + " = ");
this->writeExpression(*arguments[1], Precedence::kSequence);
// , _skTempI - 2 * _skTempN * _skTempI * _skTempN)
this->write(", " + tmpI + " - 2 * " + tmpN + " * " + tmpI + " * " + tmpN + ")");
} else {
this->writeSimpleIntrinsic(c);
}
return true;
}
case k_refract_IntrinsicKind: {
if (arguments[0]->type().columns() == 1) {
// Metal does implement refract for vectors; rather than reimplementing refract from
// scratch, we can replace the call with `refract(float2(I,0), float2(N,0), eta).x`.
this->write("(refract(float2(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(", 0), float2(");
this->writeExpression(*arguments[1], Precedence::kSequence);
this->write(", 0), ");
this->writeExpression(*arguments[2], Precedence::kSequence);
this->write(").x)");
} else {
this->writeSimpleIntrinsic(c);
}
return true;
}
case k_roundEven_IntrinsicKind: {
this->write("rint");
this->writeArgumentList(c.arguments());
return true;
}
case k_bitCount_IntrinsicKind: {
this->write("popcount(");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write(")");
return true;
}
case k_findLSB_IntrinsicKind: {
// Create a temp variable to store the expression, to avoid double-evaluating it.
std::string skTemp = this->getTempVariable(arguments[0]->type());
std::string exprType = this->typeName(arguments[0]->type());
// ctz returns numbits(type) on zero inputs; GLSL documents it as generating -1 instead.
// Use select to detect zero inputs and force a -1 result.
// (_skTemp1 = (.....), select(ctz(_skTemp1), int4(-1), _skTemp1 == int4(0)))
this->write("(");
this->write(skTemp);
this->write(" = (");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write("), select(ctz(");
this->write(skTemp);
this->write("), ");
this->write(exprType);
this->write("(-1), ");
this->write(skTemp);
this->write(" == ");
this->write(exprType);
this->write("(0)))");
return true;
}
case k_findMSB_IntrinsicKind: {
// Create a temp variable to store the expression, to avoid double-evaluating it.
std::string skTemp1 = this->getTempVariable(arguments[0]->type());
std::string exprType = this->typeName(arguments[0]->type());
// GLSL findMSB is actually quite different from Metal's clz:
// - For signed negative numbers, it returns the first zero bit, not the first one bit!
// - For an empty input (0/~0 depending on sign), findMSB gives -1; clz is numbits(type)
// (_skTemp1 = (.....),
this->write("(");
this->write(skTemp1);
this->write(" = (");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write("), ");
// Signed input types might be negative; we need another helper variable to negate the
// input (since we can only find one bits, not zero bits).
std::string skTemp2;
if (arguments[0]->type().isSigned()) {
// ... _skTemp2 = (select(_skTemp1, ~_skTemp1, _skTemp1 < 0)),
skTemp2 = this->getTempVariable(arguments[0]->type());
this->write(skTemp2);
this->write(" = (select(");
this->write(skTemp1);
this->write(", ~");
this->write(skTemp1);
this->write(", ");
this->write(skTemp1);
this->write(" < 0)), ");
} else {
skTemp2 = skTemp1;
}
// ... select(int4(clz(_skTemp2)), int4(-1), _skTemp2 == int4(0)))
this->write("select(");
this->write(this->typeName(c.type()));
this->write("(clz(");
this->write(skTemp2);
this->write(")), ");
this->write(this->typeName(c.type()));
this->write("(-1), ");
this->write(skTemp2);
this->write(" == ");
this->write(exprType);
this->write("(0)))");
return true;
}
case k_sign_IntrinsicKind: {
if (arguments[0]->type().componentType().isInteger()) {
// Create a temp variable to store the expression, to avoid double-evaluating it.
std::string skTemp = this->getTempVariable(arguments[0]->type());
std::string exprType = this->typeName(arguments[0]->type());
// (_skTemp = (.....),
this->write("(");
this->write(skTemp);
this->write(" = (");
this->writeExpression(*arguments[0], Precedence::kSequence);
this->write("), ");
// ... select(select(int4(0), int4(-1), _skTemp < 0), int4(1), _skTemp > 0))
this->write("select(select(");
this->write(exprType);
this->write("(0), ");
this->write(exprType);
this->write("(-1), ");
this->write(skTemp);
this->write(" < 0), ");
this->write(exprType);
this->write("(1), ");
this->write(skTemp);
this->write(" > 0))");
} else {
this->writeSimpleIntrinsic(c);
}
return true;
}
case k_matrixCompMult_IntrinsicKind: {
this->writeMatrixCompMult();
this->writeSimpleIntrinsic(c);
return true;
}
case k_outerProduct_IntrinsicKind: {
this->writeOuterProduct();
this->writeSimpleIntrinsic(c);
return true;
}
case k_mix_IntrinsicKind: {
SkASSERT(c.arguments().size() == 3);
if (arguments[2]->type().componentType().isBoolean()) {
// The Boolean forms of GLSL mix() use the select() intrinsic in Metal.
this->write("select");
this->writeArgumentList(c.arguments());
return true;
}
// The basic form of mix() is supported by Metal as-is.
this->writeSimpleIntrinsic(c);
return true;
}
case k_equal_IntrinsicKind:
case k_greaterThan_IntrinsicKind:
case k_greaterThanEqual_IntrinsicKind:
case k_lessThan_IntrinsicKind:
case k_lessThanEqual_IntrinsicKind:
case k_notEqual_IntrinsicKind: {
this->write("(");
this->writeExpression(*c.arguments()[0], Precedence::kRelational);
switch (kind) {
case k_equal_IntrinsicKind:
this->write(" == ");
break;
case k_notEqual_IntrinsicKind:
this->write(" != ");
break;
case k_lessThan_IntrinsicKind:
this->write(" < ");
break;
case k_lessThanEqual_IntrinsicKind:
this->write(" <= ");
break;
case k_greaterThan_IntrinsicKind:
this->write(" > ");
break;
case k_greaterThanEqual_IntrinsicKind:
this->write(" >= ");
break;
default:
SK_ABORT("unsupported comparison intrinsic kind");
}
this->writeExpression(*c.arguments()[1], Precedence::kRelational);
this->write(")");
return true;
}
case k_storageBarrier_IntrinsicKind:
this->write("threadgroup_barrier(mem_flags::mem_device)");
return true;
case k_workgroupBarrier_IntrinsicKind:
this->write("threadgroup_barrier(mem_flags::mem_threadgroup)");
return true;
case k_atomicAdd_IntrinsicKind:
this->write("atomic_fetch_add_explicit(&");
this->writeExpression(*c.arguments()[0], Precedence::kSequence);
this->write(", ");
this->writeExpression(*c.arguments()[1], Precedence::kSequence);
this->write(", memory_order_relaxed)");
return true;
case k_atomicLoad_IntrinsicKind:
this->write("atomic_load_explicit(&");
this->writeExpression(*c.arguments()[0], Precedence::kSequence);
this->write(", memory_order_relaxed)");
return true;
case k_atomicStore_IntrinsicKind:
this->write("atomic_store_explicit(&");
this->writeExpression(*c.arguments()[0], Precedence::kSequence);
this->write(", ");
this->writeExpression(*c.arguments()[1], Precedence::kSequence);
this->write(", memory_order_relaxed)");
return true;
default:
return false;
}
}