bool
TargetNVC0::isModSupported(const Instruction *insn, int s, Modifier mod) const
{
if (!isFloatType(insn->dType)) {
switch (insn->op) {
case OP_ABS:
case OP_NEG:
case OP_CVT:
case OP_CEIL:
case OP_FLOOR:
case OP_TRUNC:
case OP_AND:
case OP_OR:
case OP_XOR:
break;
case OP_ADD:
if (mod.abs())
return false;
if (insn->src(s ? 0 : 1).mod.neg())
return false;
break;
case OP_SUB:
if (s == 0)
return insn->src(1).mod.neg() ? false : true;
break;
default:
return false;
}
}
if (s > 3)
return false;
return (mod & Modifier(opInfo[insn->op].srcMods[s])) == mod;
}
PyType PyType::lookupType(const std::string &typeNameIn, ULONG64 module)
{
if (debuggingTypeEnabled())
DebugPrint() << "lookup type '" << typeNameIn << "'";
std::string typeName = typeNameIn;
trimBack(typeName);
trimFront(typeName);
if (isPointerType(typeName) || isArrayType(typeName))
return PyType(0, 0, typeName);
if (typeName.find("enum ") == 0)
typeName.erase(0, 5);
if (endsWith(typeName, " const"))
typeName.erase(typeName.length() - 6);
if (typeName == "__int64" || typeName == "unsigned __int64")
typeName.erase(typeName.find("__"), 2);
const static std::regex typeNameRE("^[a-zA-Z_][a-zA-Z0-9_]*!?[a-zA-Z0-9_<>:, \\*\\&\\[\\]]*$");
if (!std::regex_match(typeName, typeNameRE))
return PyType();
CIDebugSymbols *symbols = ExtensionCommandContext::instance()->symbols();
ULONG typeId;
HRESULT result = S_FALSE;
if (module != 0 && !isIntegralType(typeName) && !isFloatType(typeName))
result = symbols->GetTypeId(module, typeName.c_str(), &typeId);
if (FAILED(result) || result == S_FALSE)
result = symbols->GetSymbolTypeId(typeName.c_str(), &typeId, &module);
if (FAILED(result))
return createUnresolvedType(typeName);
return PyType(module, typeId, typeName);
}
//--------------------------------------------------------------
bool ofxMuiNumberData::addValue(float val, ofxMuiRange _bounds, ofxMuiRange _range) {
// prepare iterators
// start inserting
historyLength.push_back(_defaults->dataHistoryLength);
deque<float> initQ(_defaults->dataHistoryLength, val); // create a history
values.push_back(initQ); // add the queue
bounds.push_back(_bounds);
ranges.push_back(_range);
displayPrecision.push_back(_defaults->displayPrecision);
if(isBoolType() || isIntType()) {
incrementValues.push_back(_defaults->incrementInt);
} else if(isFloatType()) {
incrementValues.push_back(_defaults->incrementFloat);
}
bool boundsChanged = false;
bool rangeChanged = false;
bool valueChanged = false;
checkRange(boundsChanged, rangeChanged, getNumValues()-1);
constrainValue(valueChanged, getNumValues()-1);
}
std::string PyType::name(bool withModule) const
{
if (m_name.empty()) {
auto symbols = ExtensionCommandContext::instance()->symbols();
ULONG size = 0;
symbols->GetTypeName(m_module, m_typeId, NULL, 0, &size);
if (size == 0)
return std::string();
std::string typeName(size - 1, '\0');
if (FAILED(symbols->GetTypeName(m_module, m_typeId, &typeName[0], size, &size)))
return std::string();
m_name = typeName;
}
if (withModule && !isIntegralType(m_name) && !isFloatType(m_name)) {
std::string completeName = module();
if (!completeName.empty())
completeName.append("!");
completeName.append(m_name);
return completeName;
}
return m_name;
}
// gets the ith item of a number list
// 0.0 if no such type or out of bounds
double Tag::getListItemAsFloat(int32_t i) const {
if (isIntType(payload->tagList.type))
return payload->tagList.values->at(i)->tagInt;
if (isFloatType(payload->tagList.type))
return payload->tagList.values->at(i)->tagFloat;
return 0.0;
}
// get value as string
// might contain '\n' if list or compound
std::string Tag::asString() const {
if (isIntType(type)) return std::to_string(payload->tagInt);
else if (isFloatType(type)) return std::to_string(payload->tagFloat);
else if (type == tagTypeString) return *payload->tagString;
else if (isListType(type) || type == tagTypeCompound) {
std::string str = std::to_string(getListSize()) + " entries\n{\n";
for (int32_t i = 0; i < getListSize(); i++) {
// limit output to 10-15 lines
if (isListType(type) && i >= 10 && getListSize() > 15) {
str += " ... and " + std::to_string(getListSize()-10) + " more\n";
break;
}
Tag * tag = getListItemAsTag(i);
if (tag != NULL) {
std::string content = " " + tag->toString();
// indent content
for (size_t pos = content.find_first_of('\n'); pos != std::string::npos; pos = content.find_first_of('\n', pos+1)) {
content = content.replace(pos, 1, "\n ");
}
str += content + "\n";
}
else str += " ERROR\n";
}
str += "}";
return str;
}
return "";
}
bool
TargetNV50::isOpSupported(operation op, DataType ty) const
{
if (ty == TYPE_F64 && chipset < 0xa0)
return false;
switch (op) {
case OP_PRERET:
return chipset >= 0xa0;
case OP_TXG:
return chipset >= 0xa3 && chipset != 0xaa && chipset != 0xac;
case OP_POW:
case OP_SQRT:
case OP_DIV:
case OP_MOD:
case OP_SET_AND:
case OP_SET_OR:
case OP_SET_XOR:
case OP_SLCT:
case OP_SELP:
case OP_POPCNT:
case OP_INSBF:
case OP_EXTBF:
case OP_EXIT: // want exit modifier instead (on NOP if required)
case OP_MEMBAR:
return false;
case OP_SAD:
return ty == TYPE_S32;
case OP_SET:
return !isFloatType(ty);
default:
return true;
}
}
// gets the ith item of a list as string
// "" if out of bounds
// may contain '\n' if list or compound
std::string Tag::getListItemAsString(int32_t i) const {
if (isIntType(payload->tagList.type))
return std::to_string(payload->tagList.values->at(i)->tagInt);
if (isFloatType(payload->tagList.type))
return std::to_string(payload->tagList.values->at(i)->tagFloat);
if (payload->tagList.type == tagTypeString)
return *payload->tagList.values->at(i)->tagString;
// no primitive type, use Tag::asString()
Tag * tag = getListItemAsTag(i);
if (tag) {
std::string str = tag->asString();
delete tag;
return str;
}
return "";
}
int PyType::code() const
{
if (!m_resolved)
return TypeCodeUnresolvable;
if (m_tag < 0) {
// try to parse typeName
const std::string &typeName = name();
if (typeName.empty())
return TypeCodeUnresolvable;
if (isPointerType(typeName))
return TypeCodePointer;
if (isArrayType(typeName))
return TypeCodeArray;
if (typeName.find("<function>") == 0)
return TypeCodeFunction;
if (isIntegralType(typeName))
return TypeCodeIntegral;
if (isFloatType(typeName))
return TypeCodeFloat;
IDebugSymbolGroup2 *sg = 0;
if (FAILED(ExtensionCommandContext::instance()->symbols()->CreateSymbolGroup2(&sg)))
return TypeCodeStruct;
if (knownType(name(), 0) != KT_Unknown)
return TypeCodeStruct;
const std::string helperValueName = SymbolGroupValue::pointedToSymbolName(0, name(true));
ULONG index = DEBUG_ANY_ID;
if (SUCCEEDED(sg->AddSymbol(helperValueName.c_str(), &index)))
m_tag = PyValue(index, sg).tag();
sg->Release();
}
switch (m_tag) {
case SymTagUDT: return TypeCodeStruct;
case SymTagEnum: return TypeCodeEnum;
case SymTagTypedef: return TypeCodeTypedef;
case SymTagFunctionType: return TypeCodeFunction;
case SymTagPointerType: return TypeCodePointer;
case SymTagArrayType: return TypeCodeArray;
case SymTagBaseType: return isIntegralType(name()) ? TypeCodeIntegral : TypeCodeFloat;
default: break;
}
return TypeCodeStruct;
}
//--------------------------------------------------------------
bool ofxMuiNumberData::insertValue(int index, float val, ofxMuiRange _bounds, ofxMuiRange _range) {
if(isValidInsertionIndex(index)) {
// prepare iterators
valuesIter = values.begin() + index;
historyLengthIter = historyLength.begin() + index;
boundsIter = bounds.begin() + index;
rangesIter = ranges.begin() + index;
displayPrecisionIter = displayPrecision.begin() + index;
incrementValuesIter = incrementValues.begin() + index;
// start inserting
historyLength.insert(historyLengthIter, _defaults->dataHistoryLength);
deque<float> initQ(_defaults->dataHistoryLength, val); // create a history
values.push_back(initQ); // add the queue
values.insert(valuesIter, initQ);
bounds.insert(boundsIter, _bounds);
ranges.insert(rangesIter, _range);
// if it's global, take it from the other
displayPrecision.insert(displayPrecisionIter, _defaults->displayPrecision);
if(isBoolType() || isIntType()) {
incrementValues.insert(incrementValuesIter, _defaults->incrementInt);
} else if(isFloatType()) {
incrementValues.insert(incrementValuesIter, _defaults->incrementFloat);
}
bool boundsChanged = false;
bool rangeChanged = false;
bool valueChanged = false;
checkRange(boundsChanged, rangeChanged, index);
constrainValue(valueChanged,index);
} else {
return false;
}
}
bool
TargetGM107::isOpSupported(operation op, DataType ty) const
{
switch (op) {
case OP_SAD:
case OP_POW:
case OP_DIV:
case OP_MOD:
return false;
case OP_SQRT:
if (ty == TYPE_F64)
return false;
return chipset >= NVISA_GM200_CHIPSET;
case OP_XMAD:
if (isFloatType(ty))
return false;
break;
default:
break;
}
return true;
}
// long: rrr, arr, rcr, acr, rrc, arc, gcr, grr
// short: rr, ar, rc, gr
// immd: ri, gi
bool
TargetNV50::insnCanLoad(const Instruction *i, int s,
const Instruction *ld) const
{
DataFile sf = ld->src(0).getFile();
if (sf == FILE_IMMEDIATE && (i->predSrc >= 0 || i->flagsDef >= 0))
return false;
if (s >= opInfo[i->op].srcNr)
return false;
if (!(opInfo[i->op].srcFiles[s] & (1 << (int)sf)))
return false;
if (s == 2 && i->src(1).getFile() != FILE_GPR)
return false;
// NOTE: don't rely on flagsDef
for (int d = 0; i->defExists(d); ++d)
if (i->def(d).getFile() == FILE_FLAGS)
return false;
unsigned mode = 0;
for (int z = 0; z < Target::operationSrcNr[i->op]; ++z) {
DataFile zf = (z == s) ? sf : i->src(z).getFile();
switch (zf) {
case FILE_GPR:
break;
case FILE_MEMORY_SHARED:
case FILE_SHADER_INPUT:
mode |= 1 << (z * 2);
break;
case FILE_MEMORY_CONST:
mode |= 2 << (z * 2);
break;
case FILE_IMMEDIATE:
mode |= 3 << (z * 2);
default:
break;
}
}
switch (mode) {
case 0x00:
case 0x01:
case 0x03:
case 0x08:
case 0x09:
case 0x0c:
case 0x20:
case 0x21:
break;
case 0x0d:
if (ld->bb->getProgram()->getType() != Program::TYPE_GEOMETRY)
return false;
default:
return false;
}
uint8_t ldSize;
if ((i->op == OP_MUL || i->op == OP_MAD) && !isFloatType(i->dType)) {
// 32-bit MUL will be split into 16-bit MULs
if (ld->src(0).isIndirect(0))
return false;
if (sf == FILE_IMMEDIATE)
return false;
ldSize = 2;
} else {
ldSize = typeSizeof(ld->dType);
}
if (sf == FILE_IMMEDIATE)
return true;
// Check if memory access is encodable:
if (ldSize < 4 && sf == FILE_SHADER_INPUT) // no < 4-byte aligned a[] access
return false;
if (ld->getSrc(0)->reg.data.offset > (int32_t)(127 * ldSize))
return false;
if (ld->src(0).isIndirect(0)) {
for (int z = 0; i->srcExists(z); ++z)
if (i->src(z).isIndirect(0))
return false;
// s[] access only possible in CP, $aX always applies
if (sf == FILE_MEMORY_SHARED)
return true;
if (!ld->bb) // can't check type ...
return false;
Program::Type pt = ld->bb->getProgram()->getType();
// $aX applies to c[] only in VP, FP, GP if p[] is not accessed
if (pt == Program::TYPE_COMPUTE)
return false;
if (pt == Program::TYPE_GEOMETRY) {
if (sf == FILE_MEMORY_CONST)
return i->src(s).getFile() != FILE_SHADER_INPUT;
return sf == FILE_SHADER_INPUT;
}
return sf == FILE_MEMORY_CONST;
}
return true;
}
llvm::Value* genFloatPrimitiveMethodCall(Function& function, const SEM::Type* type, const String& methodName, const SEM::FunctionType functionType,
llvm::ArrayRef<SEM::Value> templateArgs, PendingResultArray args, llvm::Value* const hintResultValue) {
auto& module = function.module();
auto& builder = function.getBuilder();
const auto& typeName = type->getObjectType()->name().first();
const auto methodID = module.context().getMethodID(CanonicalizeMethodName(methodName));
const auto methodOwner = methodID.isConstructor() ? nullptr : args[0].resolveWithoutBind(function);
if (methodName == "__move_to") {
const auto moveToPtr = args[1].resolve(function);
const auto moveToPosition = args[2].resolve(function);
const auto destPtr = builder.CreateInBoundsGEP(moveToPtr, moveToPosition);
const auto castedDestPtr = builder.CreatePointerCast(destPtr, genPointerType(module, type));
genMoveStore(function, methodOwner, castedDestPtr, type);
return ConstantGenerator(module).getVoidUndef();
} else if (methodName == "create") {
return ConstantGenerator(module).getPrimitiveFloat(typeName, 0.0);
} else if (methodName == "__setdead" || methodName == "__set_dead") {
// Do nothing.
return ConstantGenerator(module).getVoidUndef();
} else if (methodName == "__islive" || methodName == "__is_live") {
return ConstantGenerator(module).getI1(true);
} else if (methodName.starts_with("implicit_cast_") || methodName.starts_with("cast_")) {
const auto argType = functionType.parameterTypes().front();
const auto operand = args[0].resolve(function);
const auto selfType = genType(module, type);
if (isFloatType(module, argType)) {
if (methodName.starts_with("implicit_cast_")) {
return builder.CreateFPExt(operand, selfType);
} else {
return builder.CreateFPTrunc(operand, selfType);
}
} else if (isUnsignedIntegerType(module, argType)) {
return builder.CreateUIToFP(operand, selfType);
} else if (isSignedIntegerType(module, argType)) {
return builder.CreateSIToFP(operand, selfType);
} else {
llvm_unreachable("Unknown float cast source type.");
}
} else if (isUnaryOp(methodName)) {
const auto zero = ConstantGenerator(module).getPrimitiveFloat(typeName, 0.0);
if (methodName == "implicit_cast" || methodName == "cast") {
return callCastMethod(function, methodOwner, type, methodName, templateArgs.front().typeRefType(), hintResultValue);
} else if (methodName == "implicit_copy" || methodName == "copy" || methodName == "plus") {
return methodOwner;
} else if (methodName == "minus") {
return builder.CreateFNeg(methodOwner);
} else if (methodName == "isZero") {
return builder.CreateFCmpOEQ(methodOwner, zero);
} else if (methodName == "isPositive") {
return builder.CreateFCmpOGT(methodOwner, zero);
} else if (methodName == "isNegative") {
return builder.CreateFCmpOLT(methodOwner, zero);
} else if (methodName == "abs") {
// Generates: (value < 0) ? -value : value.
const auto lessThanZero = builder.CreateFCmpOLT(methodOwner, zero);
return builder.CreateSelect(lessThanZero, builder.CreateFNeg(methodOwner), methodOwner);
} else if (methodName == "sqrt") {
llvm::Type* const intrinsicTypes[] = { methodOwner->getType() };
const auto sqrtIntrinsic = llvm::Intrinsic::getDeclaration(module.getLLVMModulePtr(), llvm::Intrinsic::sqrt, intrinsicTypes);
llvm::Value* const sqrtArgs[] = { methodOwner };
return builder.CreateCall(sqrtIntrinsic, sqrtArgs);
} else {
llvm_unreachable("Unknown primitive unary op.");
}
} else if (isBinaryOp(methodName)) {
const auto operand = args[1].resolveWithoutBind(function);
if (methodName == "add") {
return builder.CreateFAdd(methodOwner, operand);
} else if (methodName == "subtract") {
return builder.CreateFSub(methodOwner, operand);
} else if (methodName == "multiply") {
return builder.CreateFMul(methodOwner, operand);
} else if (methodName == "divide") {
return builder.CreateFDiv(methodOwner, operand);
} else if (methodName == "modulo") {
return builder.CreateFRem(methodOwner, operand);
} else if (methodName == "equal") {
return builder.CreateFCmpOEQ(methodOwner, operand);
} else if (methodName == "not_equal") {
return builder.CreateFCmpONE(methodOwner, operand);
} else if (methodName == "less_than") {
return builder.CreateFCmpOLT(methodOwner, operand);
} else if (methodName == "less_than_or_equal") {
return builder.CreateFCmpOLE(methodOwner, operand);
} else if (methodName == "greater_than") {
return builder.CreateFCmpOGT(methodOwner, operand);
} else if (methodName == "greater_than_or_equal") {
return builder.CreateFCmpOGE(methodOwner, operand);
} else if (methodName == "compare") {
const auto isLessThan = builder.CreateFCmpOLT(methodOwner, operand);
const auto isGreaterThan = builder.CreateFCmpOGT(methodOwner, operand);
const auto minusOne = ConstantGenerator(module).getI8(-1);
const auto zero = ConstantGenerator(module).getI8(0);
const auto plusOne = ConstantGenerator(module).getI8(1);
return builder.CreateSelect(isLessThan, minusOne,
builder.CreateSelect(isGreaterThan, plusOne, zero));
} else {
llvm_unreachable("Unknown primitive binary op.");
}
} else {
printf("%s\n", methodName.c_str());
llvm_unreachable("Unknown primitive method.");
}
}
// Return true when an instruction requires to set up a barrier because it
// doesn't operate at a fixed latency. Variable latency instructions are memory
// operations, double precision operations, special function unit operations
// and other low throughput instructions.
bool
TargetGM107::isBarrierRequired(const Instruction *insn) const
{
const OpClass cl = getOpClass(insn->op);
if (insn->dType == TYPE_F64 || insn->sType == TYPE_F64)
return true;
switch (cl) {
case OPCLASS_ATOMIC:
case OPCLASS_LOAD:
case OPCLASS_STORE:
case OPCLASS_SURFACE:
case OPCLASS_TEXTURE:
return true;
case OPCLASS_SFU:
switch (insn->op) {
case OP_COS:
case OP_EX2:
case OP_LG2:
case OP_LINTERP:
case OP_PINTERP:
case OP_RCP:
case OP_RSQ:
case OP_SIN:
case OP_SQRT:
return true;
default:
break;
}
break;
case OPCLASS_BITFIELD:
switch (insn->op) {
case OP_BFIND:
case OP_POPCNT:
return true;
default:
break;
}
break;
case OPCLASS_CONTROL:
switch (insn->op) {
case OP_EMIT:
case OP_RESTART:
return true;
default:
break;
}
break;
case OPCLASS_OTHER:
switch (insn->op) {
case OP_AFETCH:
case OP_PFETCH:
case OP_PIXLD:
case OP_SHFL:
return true;
case OP_RDSV:
return !isCS2RSV(insn->getSrc(0)->reg.data.sv.sv);
default:
break;
}
break;
case OPCLASS_ARITH:
if ((insn->op == OP_MUL || insn->op == OP_MAD) &&
!isFloatType(insn->dType))
return true;
break;
case OPCLASS_CONVERT:
if (insn->def(0).getFile() != FILE_PREDICATE &&
insn->src(0).getFile() != FILE_PREDICATE)
return true;
break;
default:
break;
}
return false;
}
bool isIntOrFloatType(const QByteArray &type)
{
return isIntType(type) || isFloatType(type);
}
bool OpenCLParser::convert(std::string fileNameIN, std::string fileNameOUT) {
if ( access( fileNameIN.c_str(), F_OK ) == -1 ) {
LOG(ERROR) << "kernel source file = '" << fileNameIN.c_str() << "' doesn't exist";
return false;
}
if ( access( fileNameIN.c_str(), R_OK ) == -1 ) {
LOG(ERROR) << "kernel source file = '" << fileNameIN.c_str() << "' isn't readable";
return false;
}
if ( access( fileNameOUT.c_str(), F_OK ) == 0 ) {
struct stat statIN;
if (stat(fileNameIN.c_str(), &statIN) == -1) {
perror(fileNameIN.c_str());
return false;
}
struct stat statOUT;
if (stat(fileNameOUT.c_str(), &statOUT) == -1) {
perror(fileNameOUT.c_str());
return false;
}
if ( statOUT.st_mtime > statIN.st_mtime ) {
DLOG(INFO) << "kernel source file = '" << fileNameOUT.c_str() << "' up-to-date";
return true;
}
}
std::ifstream file;
file.open(fileNameIN.c_str(), std::ifstream::in );
if ( ! file.is_open() ) {
LOG(ERROR) << "failed to open file = '" << fileNameIN.c_str() << "' for reading";
return false;
}
std::string line;
std::string kernel_buffer;
std::string kernel_name;
std::string kernel_type;
std::string kernel_name_typed;
std::string kernel_line_typed;
std::string kernel_modified;
std::string type_replace;
std::string stdOpenCL;
stdOpenCL += "// This file was auto-generated from file '" + fileNameIN + "' to conform to standard OpenCL\n";
bool recording = false;
while (std::getline(file, line)) {
if ( isAttributeLine(line) ) {
if ( recording ) {
recording = false;
}
kernel_name_typed = getTypedKernelName(line);
kernel_line_typed = "__kernel void " + kernel_name_typed + getTypedKernelLine(line) + " {";
if ( isFloatType(kernel_name_typed) ) {
type_replace = "float";
}
if ( isDoubleType(kernel_name_typed) ) {
type_replace = "double";
}
kernel_modified = kernel_line_typed + "\n" + kernel_buffer;
boost::regex re;
re = boost::regex("\\sT\\s", boost::regex::perl);
kernel_modified = boost::regex_replace(kernel_modified, re, " "+type_replace+" ");
re = boost::regex("\\sT\\*\\s", boost::regex::perl);
kernel_modified = boost::regex_replace(kernel_modified, re, " "+type_replace+"* ");
stdOpenCL += kernel_modified;
continue;
}
if ( isTemplateKernelLine(line) ) {
kernel_name = getKernelName(line);
kernel_type = getKernelType(line);
DLOG(INFO)<<"found template kernel '"<<kernel_name<<"' with type '"<<kernel_type<<"'";
if ( recording == false ) {
recording = true;
} else {
LOG(ERROR) << "error parsing kernel source file = '" << fileNameIN.c_str() << "'";
return false;
}
continue;
}
if ( recording ) {
kernel_buffer += line + "\n";
} else {
kernel_buffer = "";
stdOpenCL += line + "\n";
}
}
std::ofstream out(fileNameOUT.c_str());
out << stdOpenCL;
out.close();
DLOG(INFO) << "convert AMD OpenCL '"<<fileNameIN.c_str()<<"' to standard OpenCL '"<<fileNameOUT.c_str()<<"'";
return true;
}
