diff options
Diffstat (limited to 'source/slang/slang-emit-cuda.cpp')
| -rw-r--r-- | source/slang/slang-emit-cuda.cpp | 511 |
1 files changed, 152 insertions, 359 deletions
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp index 284652682..a151ab0e2 100644 --- a/source/slang/slang-emit-cuda.cpp +++ b/source/slang/slang-emit-cuda.cpp @@ -123,131 +123,6 @@ SlangResult CUDASourceEmitter::_calcCUDATextureTypeName(IRTextureTypeBase* texTy return SLANG_FAIL; } -SlangResult CUDASourceEmitter::calcScalarFuncName(HLSLIntrinsic::Op op, IRBasicType* type, StringBuilder& outBuilder) -{ - typedef HLSLIntrinsic::Op Op; - - UnownedStringSlice funcName; - - switch (op) - { - case Op::FRem: - { - if (type->getOp() == kIROp_FloatType || type->getOp() == kIROp_DoubleType) - { - funcName = HLSLIntrinsic::getInfo(op).funcName; - } - break; - } - default: break; - } - - if (funcName.getLength()) - { - outBuilder << funcName; - if (type->getOp() == kIROp_FloatType) - { - outBuilder << "f"; - } - return SLANG_OK; - } - - // Defer to the supers impl - return Super::calcScalarFuncName(op, type, outBuilder); -} - -void CUDASourceEmitter::emitSpecializedOperationDefinition(const HLSLIntrinsic* specOp) -{ - typedef HLSLIntrinsic::Op Op; - - if (auto vecType = as <IRVectorType>(specOp->returnType)) - { - // Converting to or from half vector types is implemented prelude as convert___half functions - // Get the from type -> if it's half we ignore - - if (specOp->op == Op::ConstructConvert) - { - auto signatureType = specOp->signatureType; - - // Need to have impl of convert_float, double, int, uint, in prelude - - const auto paramCount = signatureType->getParamCount(); - SLANG_UNUSED(paramCount); - - // We have 2 'params' and param 1 is the source type - SLANG_ASSERT(paramCount == 2); - IRType* paramType = signatureType->getParamType(1); - - auto vecParamType = as<IRVectorType>(paramType); - - if (auto baseType = as<IRBasicType>(vecParamType->getElementType())) - { - if (baseType->getBaseType() == BaseType::Half) - { - return; - } - } - } - - if (auto baseType = as<IRBasicType>(vecType->getElementType())) - { - if (baseType->getBaseType() == BaseType::Half) - { - switch (specOp->op) - { - case Op::Init: - - case Op::Add: - case Op::Mul: - case Op::Div: - case Op::Sub: - - case Op::Neg: - - case Op::ConstructFromScalar: - case Op::ConstructConvert: - - case Op::Leq: - case Op::Less: - case Op::Greater: - case Op::Geq: - case Op::Neq: - case Op::Eql: - { - return; - } - } - } - } - } - - switch (specOp->op) - { - case Op::Init: - { - // Special case handling - auto returnType = specOp->returnType; - - if (auto vecType = as <IRVectorType>(returnType)) - { - if (auto baseType = as<IRBasicType>(vecType->getElementType())) - { - if (baseType->getBaseType() == BaseType::Half) - { - // Defined already in cuda-prelude.h - return; - } - } - } - - break; - } - default: break; - } - - Super::emitSpecializedOperationDefinition(specOp); -} - SlangResult CUDASourceEmitter::calcTypeName(IRType* type, CodeGenTarget target, StringBuilder& out) { SLANG_UNUSED(target); @@ -322,25 +197,6 @@ SlangResult CUDASourceEmitter::calcTypeName(IRType* type, CodeGenTarget target, return Super::calcTypeName(type, target, out); } -const UnownedStringSlice* CUDASourceEmitter::getVectorElementNames(BaseType baseType, Index elemCount) -{ - static const UnownedStringSlice normal[] = { UnownedStringSlice::fromLiteral("x"), UnownedStringSlice::fromLiteral("y"), UnownedStringSlice::fromLiteral("z"), UnownedStringSlice::fromLiteral("w") }; - static const UnownedStringSlice half3[] = { UnownedStringSlice::fromLiteral("xy.x"), UnownedStringSlice::fromLiteral("xy.y"), UnownedStringSlice::fromLiteral("z") }; - static const UnownedStringSlice half4[] = { UnownedStringSlice::fromLiteral("xy.x"), UnownedStringSlice::fromLiteral("xy.y"), UnownedStringSlice::fromLiteral("zw.x"), UnownedStringSlice::fromLiteral("zw.y")}; - - if (baseType == BaseType::Half) - { - switch (elemCount) - { - default: break; - case 3: return half3; - case 4: return half4; - } - } - - return normal; -} - void CUDASourceEmitter::emitLayoutSemanticsImpl(IRInst* inst, char const* uniformSemanticSpelling) { Super::emitLayoutSemanticsImpl(inst, uniformSemanticSpelling); @@ -436,49 +292,6 @@ void CUDASourceEmitter::emitGlobalRTTISymbolPrefix() m_writer->emit("__constant__ "); } -void CUDASourceEmitter::emitCall(const HLSLIntrinsic* specOp, IRInst* inst, const IRUse* operands, int numOperands, const EmitOpInfo& inOuterPrec) -{ - switch (specOp->op) - { - case HLSLIntrinsic::Op::Init: - { - // For CUDA vector types we construct with make_ - - auto writer = m_writer; - - IRType* retType = specOp->returnType; - - if (IRVectorType* vecType = as<IRVectorType>(retType)) - { - if (numOperands == getIntVal(vecType->getElementCount())) - { - // Get the type name - writer->emit("make_"); - emitType(retType); - writer->emitChar('('); - - for (int i = 0; i < numOperands; ++i) - { - if (i > 0) - { - writer->emit(", "); - } - emitOperand(operands[i].get(), getInfo(EmitOp::General)); - } - - writer->emitChar(')'); - return; - } - } - // Just use the default - break; - } - default: break; - } - - return Super::emitCall(specOp, inst, operands, numOperands, inOuterPrec); -} - void CUDASourceEmitter::emitLoopControlDecorationImpl(IRLoopControlDecoration* decl) { if (decl->getMode() == kIRLoopControl_Unroll) @@ -487,59 +300,25 @@ void CUDASourceEmitter::emitLoopControlDecorationImpl(IRLoopControlDecoration* d } } -static bool _areEquivalent(IRType* a, IRType* b) -{ - if (a == b) - { - return true; - } - if (a->getOp() != b->getOp()) - { - return false; - } - - switch (a->getOp()) - { - case kIROp_VectorType: - { - IRVectorType* vecA = static_cast<IRVectorType*>(a); - IRVectorType* vecB = static_cast<IRVectorType*>(b); - return getIntVal(vecA->getElementCount()) == getIntVal(vecB->getElementCount()) && - _areEquivalent(vecA->getElementType(), vecB->getElementType()); - } - case kIROp_MatrixType: - { - IRMatrixType* matA = static_cast<IRMatrixType*>(a); - IRMatrixType* matB = static_cast<IRMatrixType*>(b); - return getIntVal(matA->getColumnCount()) == getIntVal(matB->getColumnCount()) && - getIntVal(matA->getRowCount()) == getIntVal(matB->getRowCount()) && - _areEquivalent(matA->getElementType(), matB->getElementType()); - } - default: - { - return as<IRBasicType>(a) != nullptr; - } - } -} - void CUDASourceEmitter::_emitInitializerListValue(IRType* dstType, IRInst* value) { // When constructing a matrix or vector from a single value this is handled by the default path switch (value->getOp()) { - case kIROp_MakeMatrix: case kIROp_MakeVector: + case kIROp_MakeMatrix: { IRType* type = value->getDataType(); // If the types are the same, we can can just break down and use - if (_areEquivalent(dstType, type)) + if (dstType == type) { if (auto vecType = as<IRVectorType>(type)) { if (UInt(getIntVal(vecType->getElementCount())) == value->getOperandCount()) { + emitType(type); _emitInitializerList(vecType->getElementType(), value->getOperands(), value->getOperandCount()); return; } @@ -551,20 +330,25 @@ void CUDASourceEmitter::_emitInitializerListValue(IRType* dstType, IRInst* value // TODO(JS): If num cols = 1, then it *doesn't* actually return a vector. // That could be argued is an error because we want swizzling or [] to work. - IRType* rowType = m_typeSet.addVectorType(matType->getElementType(), int(colCount)); - IRVectorType* rowVectorType = as<IRVectorType>(rowType); + IRBuilder builder(matType->getModule()); + builder.setInsertBefore(matType); const Index operandCount = Index(value->getOperandCount()); // Can init, with vectors. // For now special case if the rowVectorType is not actually a vector (when elementSize == 1) - if (operandCount == rowCount || rowVectorType == nullptr) + if (operandCount == rowCount) { - // We have to output vectors - - // Emit the braces for the Matrix struct, contains an row array. + // Emit the braces for the Matrix struct, and then each row vector in its own line. + emitType(matType); m_writer->emit("{\n"); m_writer->indent(); - _emitInitializerList(rowType, value->getOperands(), rowCount); + for (Index i = 0; i < rowCount; ++i) + { + if (i != 0) m_writer->emit(",\n"); + emitType(matType->getElementType()); + m_writer->emit(colCount); + _emitInitializerList(matType->getElementType(), value->getOperand(i)->getOperands(), colCount); + } m_writer->dedent(); m_writer->emit("\n}"); return; @@ -575,21 +359,18 @@ void CUDASourceEmitter::_emitInitializerListValue(IRType* dstType, IRInst* value IRType* elementType = matType->getElementType(); IRUse* operands = value->getOperands(); - // Emit the braces for the Matrix struct, and the array of rows - m_writer->emit("{\n"); - m_writer->indent(); + // Emit the braces for the Matrix struct, and the elements of each row in its own line. + emitType(matType); m_writer->emit("{\n"); m_writer->indent(); for (Index i = 0; i < rowCount; ++i) { - if (i != 0) m_writer->emit(", "); - _emitInitializerList(elementType, operands, colCount); + if (i != 0) m_writer->emit(",\n"); + _emitInitializerListContent(elementType, operands, colCount); operands += colCount; } m_writer->dedent(); m_writer->emit("\n}"); - m_writer->dedent(); - m_writer->emit("\n}"); return; } } @@ -603,116 +384,157 @@ void CUDASourceEmitter::_emitInitializerListValue(IRType* dstType, IRInst* value emitOperand(value, getInfo(EmitOp::General)); } -void CUDASourceEmitter::_emitInitializerList(IRType* elementType, IRUse* operands, Index operandCount) +void CUDASourceEmitter::_emitInitializerListContent(IRType* elementType, IRUse* operands, Index operandCount) { - m_writer->emit("{\n"); - m_writer->indent(); - for (Index i = 0; i < operandCount; ++i) { if (i != 0) m_writer->emit(", "); _emitInitializerListValue(elementType, operands[i].get()); } - - m_writer->dedent(); - m_writer->emit("\n}"); } -void CUDASourceEmitter::_emitGetHalfVectorElement(IRInst* base, Index index, Index vecSize, const EmitOpInfo& inOuterPrec) -{ - SLANG_ASSERT(index < vecSize); - - EmitOpInfo outerPrec = inOuterPrec; - - auto prec = getInfo(EmitOp::Postfix); - const bool needClose = maybeEmitParens(outerPrec, prec); - emitOperand(base, leftSide(outerPrec, prec)); +void CUDASourceEmitter::_emitInitializerList(IRType* elementType, IRUse* operands, Index operandCount) +{ + m_writer->emit("{\n"); + m_writer->indent(); - m_writer->emit("."); + _emitInitializerListContent(elementType, operands, operandCount); - switch (vecSize) - { - default: - { - char const* kComponents[] = { "x", "y", "z", "w" }; - m_writer->emit(kComponents[index]); - break; - } - case 3: - { - char const* kComponents[] = { "xy.x", "xy.y", "z"}; - m_writer->emit(kComponents[index]); - break; - } - case 4: - { - char const* kComponents[] = { "xy.x", "xy.y", "zw.x", "zw.y" }; - m_writer->emit(kComponents[index]); - break; - } - } + m_writer->dedent(); + m_writer->emit("\n}"); +} - maybeCloseParens(needClose); +void CUDASourceEmitter::emitIntrinsicCallExprImpl(IRCall* inst, IRTargetIntrinsicDecoration* targetIntrinsic, EmitOpInfo const& inOuterPrec) +{ + if (targetIntrinsic->getDefinition().startsWith("__half")) + m_extensionTracker->requireBaseType(BaseType::Half); + Super::emitIntrinsicCallExprImpl(inst, targetIntrinsic, inOuterPrec); } bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec) { switch(inst->getOp()) { - case kIROp_swizzle: + case kIROp_MakeVector: + case kIROp_MakeVectorFromScalar: { - // We need to special case for half types. - auto swizzleInst = static_cast<IRSwizzle*>(inst); - - IRInst* baseInst = swizzleInst->getBase(); - IRType* baseType = baseInst->getDataType(); - - // If we are swizzling from a built in type, - if (as<IRBasicType>(baseType)) + m_writer->emit("make_"); + emitType(inst->getDataType()); + m_writer->emit("("); + bool isFirst = true; + char xyzwNames[] = "xyzw"; + for (UInt i = 0; i < inst->getOperandCount(); i++) { - // Just use the default behavior + auto arg = inst->getOperand(i); + if (auto vectorType = as<IRVectorType>(arg->getDataType())) + { + for (int j = 0; j < cast<IRIntLit>(vectorType->getElementCount())->getValue(); j++) + { + if (isFirst) + isFirst = false; + else + m_writer->emit(", "); + auto outerPrec = getInfo(EmitOp::General); + auto prec = getInfo(EmitOp::Postfix); + emitOperand(arg, leftSide(outerPrec, prec)); + m_writer->emit("."); + m_writer->emitChar(xyzwNames[j]); + } + } + else + { + if (isFirst) + isFirst = false; + else + m_writer->emit(", "); + emitOperand(arg, getInfo(EmitOp::General)); + } } - else if (auto vecType = as<IRVectorType>(baseType)) + m_writer->emit(")"); + return true; + } + case kIROp_FloatCast: + case kIROp_CastIntToFloat: + case kIROp_IntCast: + case kIROp_CastFloatToInt: + { + if (auto dstVectorType = as<IRVectorType>(inst->getDataType())) { - if (auto basicType = as<IRBasicType>(vecType->getElementType())) + m_writer->emit("make_"); + emitType(inst->getDataType()); + m_writer->emit("("); + bool isFirst = true; + char xyzwNames[] = "xyzw"; + for (UInt i = 0; i < inst->getOperandCount(); i++) { - if (basicType->getBaseType() == BaseType::Half) + auto arg = inst->getOperand(i); + if (auto vectorType = as<IRVectorType>(arg->getDataType())) { - const Index vecElementCount = Index(getIntVal(vecType->getElementCount())); - - const Index elementCount = Index(swizzleInst->getElementCount()); - if (elementCount == 1) - { - const Index index = Index(getIntVal(swizzleInst->getElementIndex(0))); - _emitGetHalfVectorElement(baseInst, index, vecElementCount, inOuterPrec); - } - else + for (int j = 0; j < cast<IRIntLit>(vectorType->getElementCount())->getValue(); j++) { - auto outerPrec = getInfo(EmitOp::General); - - m_writer->emit("make___half"); - m_writer->emitInt64(elementCount); + if (isFirst) + isFirst = false; + else + m_writer->emit(", "); m_writer->emit("("); - - for (Index i = 0; i < elementCount; ++i) - { - if (i) - { - m_writer->emit(", "); - } - - const Index index = Index(getIntVal(swizzleInst->getElementIndex(i))); - _emitGetHalfVectorElement(baseInst, index, vecElementCount, outerPrec); - } - + emitType(dstVectorType->getElementType()); m_writer->emit(")"); + auto outerPrec = getInfo(EmitOp::General); + auto prec = getInfo(EmitOp::Postfix); + emitOperand(arg, leftSide(outerPrec, prec)); + m_writer->emit("."); + m_writer->emitChar(xyzwNames[j]); } - return true; + } + else + { + if (isFirst) + isFirst = false; + else + m_writer->emit(", "); + m_writer->emit("("); + emitType(dstVectorType->getElementType()); + m_writer->emit(")"); + emitOperand(arg, getInfo(EmitOp::General)); } } + m_writer->emit(")"); + return true; } - break; + else if (auto matrixType = as<IRMatrixType>(inst->getDataType())) + { + m_writer->emit("make"); + emitType(inst->getDataType()); + m_writer->emit("("); + for (UInt i = 0; i < inst->getOperandCount(); i++) + { + auto arg = inst->getOperand(i); + if (i > 0) + m_writer->emit(", "); + emitOperand(arg, getInfo(EmitOp::General)); + } + m_writer->emit(")"); + return true; + } + return false; + } + case kIROp_MakeMatrix: + case kIROp_MakeMatrixFromScalar: + case kIROp_MatrixReshape: + { + m_writer->emit("make"); + emitType(inst->getDataType()); + m_writer->emit("("); + for (UInt i = 0; i < inst->getOperandCount(); i++) + { + auto arg = inst->getOperand(i); + if (i > 0) + m_writer->emit(", "); + emitOperand(arg, getInfo(EmitOp::General)); + } + m_writer->emit(")"); + return true; } case kIROp_MakeArray: { @@ -722,13 +544,9 @@ bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu IRType* elementType = arrayType->getElementType(); // Emit braces for the FixedArray struct. - m_writer->emit("{\n"); - m_writer->indent(); _emitInitializerList(elementType, inst->getOperands(), Index(inst->getOperandCount())); - m_writer->dedent(); - m_writer->emit("\n}"); return true; } case kIROp_WaveMaskBallot: @@ -820,7 +638,19 @@ void CUDASourceEmitter::emitVectorTypeNameImpl(IRType* elementType, IRIntegerVal void CUDASourceEmitter::emitSimpleTypeImpl(IRType* type) { - m_writer->emit(_getTypeName(type)); + switch (type->getOp()) + { + case kIROp_VectorType: + { + auto vectorType = as<IRVectorType>(type); + m_writer->emit(getVectorPrefix(vectorType->getElementType()->getOp())); + m_writer->emit(as<IRIntLit>(vectorType->getElementCount())->getValue()); + break; + } + default: + m_writer->emit(_getTypeName(type)); + break; + } } void CUDASourceEmitter::emitRateQualifiersImpl(IRRate* rate) @@ -907,27 +737,6 @@ void CUDASourceEmitter::emitPreModuleImpl() // Emit generated types/functions writer->emit("\n"); - - { - List<IRType*> types; - m_typeSet.getTypes(IRTypeSet::Kind::Matrix, types); - - // Emit the type definitions - for (auto type : types) - { - emitTypeDefinition(type); - } - } - - { - List<const HLSLIntrinsic*> intrinsics; - m_intrinsicSet.getIntrinsics(intrinsics); - // Emit all the intrinsics that were used - for (auto intrinsic : intrinsics) - { - _maybeEmitSpecializedOperationDefinition(intrinsic); - } - } } @@ -951,22 +760,6 @@ bool CUDASourceEmitter::tryEmitGlobalParamImpl(IRGlobalParam* varDecl, IRType* v void CUDASourceEmitter::emitModuleImpl(IRModule* module, DiagnosticSink* sink) { - // Setup all built in types used in the module - m_typeSet.addAllBuiltinTypes(module); - // If any matrix types are used, then we need appropriate vector types too. - m_typeSet.addVectorForMatrixTypes(); - - // We need to add some vector intrinsics - used for calculating thread ids - { - IRType* type = m_typeSet.addVectorType(m_typeSet.getBuilder().getBasicType(BaseType::UInt), 3); - IRType* args[] = { type, type }; - - _addIntrinsic(HLSLIntrinsic::Op::Add, type, args, SLANG_COUNT_OF(args)); - _addIntrinsic(HLSLIntrinsic::Op::Mul, type, args, SLANG_COUNT_OF(args)); - } - - // TODO(JS): We may need to generate types (for example for matrices) - CLikeSourceEmitter::emitModuleImpl(module, sink); // Emit all witness table definitions. |