AuroraMiddleware/HLSLcc
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

bring latest changes from upstream ( 1155504498e0 )

Browse Source 
- Bunch of fixes for issues found during SRP Lightweight and HD RP development
- Fix indentation issues from generated Metal output
- Metal Tessellation fixes
- Improve ES2 support

- Merge "fix issue with negation of (bool->int) result" #22 (https://github.com/Unity-Technologies/HLSLcc/pull/22)
- Merge "Updated to support basic parsing of resources for Shader Model 5.1" #27 (https://github.com/Unity-Technologies/HLSLcc/pull/27)
- Partial merge of "bugfix request" #37 (https://github.com/Unity-Technologies/HLSLcc/pull/37)
This commit is contained in:
Antti Tapaninen 2018-12-22 16:20:26 -08:00
parent cb9809a23c
commit 5a261a7fe8
20 changed files with 668 additions and 327 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

73
.editorconfig Normal file
View File

@ -0,0 +1,73 @@
# see http://editorconfig.org/ for docs on this file
root = true
[*]
# help with sharing files across os's (i.e. network share or through local vm)
end_of_line = lf
#charset temporarily disabled due to bug in VS2017 changing to UTF-8 with BOM (https://favro.com/card/c564ede4ed3337f7b17986b6/Uni-17877)
#charset = utf-8
trim_trailing_whitespace = true
insert_final_newline = true
# formattable file extensions (keep in sync with format.ini from unity-meta repo)
#
# Note: We need to split the formattable files configs into shorter duplicate entries (logically grouped)
# due to known issue in VS editorconfig extension where there is a limit of 51 characters (empirically determined).
# see: https://github.com/editorconfig/editorconfig-visualstudio/issues/21
#
## uncrustify
[*.{c,h,cpp,hpp,m,mm,cc,cs}]
indent_style = space
indent_size = 4
## generic formatter (shaders)
[*.{cg,cginc,glslinc,hlsl,shader,y,ypp,yy}]
indent_style = space
indent_size = 4
## generic formatter (misc)
[*.{asm,s,S,pch,pchmm,java,sh,uss}]
indent_style = space
indent_size = 4
## perltidy
[*.{pl,pm,t,it}]
indent_style = space
indent_size = 4
## unity special
[*.{bindings,mem.xml}]
indent_style = space
indent_size = 4
# other filetypes we want to overwrite default configuration to preserve the standard
[{Makefile,makefile}]
# TAB characters are part of the Makefile format
indent_style = tab
[*.{md,markdown}]
# trailing whitespace is significant in markdown (bad choice, bad!)
trim_trailing_whitespace = false
# keep these and the VS stuff below in sync with .hgeol's CRLF extensions
[*.{vcproj,bat,cmd,xaml,tt,t4,ttinclude}]
end_of_line = crlf
# this VS-specific stuff is based on experiments to see how VS will modify a file after it has been manually edited.
# the settings are meant to closely match what VS does to minimize unnecessary diffs. this duplicates some settings in *
# but let's be explicit here to be safe (in case someone wants to copy-paste this out to another .editorconfig).
[*.{vcxproj,vcxproj.filters,csproj,props,targets}]
indent_style = space
indent_size = 2
end_of_line = crlf
charset = utf-8-bom
trim_trailing_whitespace = true
insert_final_newline = false
[*.{sln,sln.template}]
indent_style = tab
indent_size = 4
end_of_line = crlf
charset = utf-8
trim_trailing_whitespace = true
insert_final_newline = false

2
include/ShaderInfo.h
View File

@ -239,6 +239,8 @@ struct ResourceBinding
uint32_t ui32BindPoint;
uint32_t ui32BindCount;
uint32_t ui32Flags;
uint32_t ui32Space;
uint32_t ui32RangeID;
REFLECT_RESOURCE_DIMENSION eDimension;
RESOURCE_RETURN_TYPE ui32ReturnType;
uint32_t ui32NumSamples;

4
include/hlslcc.h
View File

@ -425,6 +425,7 @@ public:
}
MemberDefinitions m_SharedFunctionMembers;
std::vector<std::string> m_SharedDependencies;
BindingSlotAllocator m_SharedTextureSlots, m_SharedSamplerSlots;
BindingSlotAllocator m_SharedBufferSlots;
@ -439,6 +440,7 @@ public:
m_NextSpecID = kArraySizeConstantID + 1;
m_SpecConstantMap.clear();
m_SharedFunctionMembers.clear();
m_SharedDependencies.clear();
}
// Retrieve or allocate a layout slot for Vulkan specialization constant
@ -481,7 +483,7 @@ public:
virtual bool OnConstantBuffer(const std::string &name, size_t bufferSize, size_t memberCount) { return true; }
// Returns false if this constant var is not needed for this shader. See above.
virtual bool OnConstant(const std::string &name, int bindIndex, SHADER_VARIABLE_TYPE cType, int rows, int cols, bool isMatrix, int arraySize) { return true; }
virtual bool OnConstant(const std::string &name, int bindIndex, SHADER_VARIABLE_TYPE cType, int rows, int cols, bool isMatrix, int arraySize, bool isUsed) { return true; }
virtual void OnConstantBufferBinding(const std::string &name, int bindIndex) {}
virtual void OnTextureBinding(const std::string &name, int bindIndex, int samplerIndex, bool multisampled, HLSLCC_TEX_DIMENSION dim, bool isUAV) {}

5
src/HLSLCrossCompilerContext.cpp
View File

@ -293,3 +293,8 @@ bool HLSLCrossCompilerContext::IsVulkan() const
{
return (flags & HLSLCC_FLAG_VULKAN_BINDINGS) != 0;
}
bool HLSLCrossCompilerContext::IsSwitch() const
{
return (flags & HLSLCC_FLAG_NVN_TARGET) != 0;
}

36
src/HLSLccToolkit.cpp
View File

@ -494,4 +494,40 @@ namespace HLSLcc
if (ePos < 0 && pointPos < 0 && !fpcheck(f))
bcatcstr(b, ".0");
}
bstring GetEarlyMain(HLSLCrossCompilerContext *psContext)
{
bstring *oldString = psContext->currentGLSLString;
bstring *str = &psContext->psShader->asPhases[psContext->currentPhase].earlyMain;
int indent = psContext->indent;
if (psContext->psShader->eTargetLanguage == LANG_METAL && !psContext->indent)
++psContext->indent;
psContext->currentGLSLString = str;
psContext->AddIndentation();
psContext->currentGLSLString = oldString;
psContext->indent = indent;
return *str;
}
bstring GetPostShaderCode(HLSLCrossCompilerContext *psContext)
{
bstring *oldString = psContext->currentGLSLString;
bstring *str = &psContext->psShader->asPhases[psContext->currentPhase].postShaderCode;
int indent = psContext->indent;
if (psContext->psShader->eTargetLanguage == LANG_METAL && !psContext->indent)
++psContext->indent;
psContext->psShader->asPhases[psContext->currentPhase].hasPostShaderCode = 1;
psContext->currentGLSLString = str;
psContext->AddIndentation();
psContext->currentGLSLString = oldString;
psContext->indent = indent;
return *str;
}
}

22
src/Operand.cpp
View File

@ -113,8 +113,6 @@ uint32_t Operand::GetNumSwizzleElements(uint32_t _ui32CompMask /* = OPERAND_4_CO
switch (eType)
{
case OPERAND_TYPE_INPUT_THREAD_ID_IN_GROUP_FLATTENED:
return 1; // TODO: does mask make any sense here?
case OPERAND_TYPE_INPUT_THREAD_ID_IN_GROUP:
case OPERAND_TYPE_INPUT_THREAD_ID:
case OPERAND_TYPE_INPUT_THREAD_GROUP_ID:
@ -368,6 +366,8 @@ SHADER_VARIABLE_TYPE Operand::GetDataType(HLSLCrossCompilerContext* psContext, S
break;
}
case OPERAND_TYPE_INPUT:
case OPERAND_TYPE_INPUT_PATCH_CONSTANT:
case OPERAND_TYPE_INPUT_CONTROL_POINT:
{
const uint32_t ui32Register = aui32ArraySizes[iIndexDims - 1];
int regSpace = GetRegisterSpace(psContext);
@ -398,7 +398,7 @@ SHADER_VARIABLE_TYPE Operand::GetDataType(HLSLCrossCompilerContext* psContext, S
case NAME_RENDER_TARGET_ARRAY_INDEX:
case NAME_VIEWPORT_ARRAY_INDEX:
case NAME_SAMPLE_INDEX:
return SVT_INT;
return (psContext->psShader->eTargetLanguage == LANG_METAL) ? SVT_UINT : SVT_INT;
case NAME_IS_FRONT_FACE:
return SVT_UINT;
@ -416,22 +416,14 @@ SHADER_VARIABLE_TYPE Operand::GetDataType(HLSLCrossCompilerContext* psContext, S
if (psIn->eSystemValueType == NAME_IS_FRONT_FACE)
return SVT_UINT;
if (eSpecialName == NAME_PRIMITIVE_ID || eSpecialName == NAME_VERTEX_ID)
{
return SVT_INT;
}
//UINT in DX, INT in GL.
if (psIn->eSystemValueType == NAME_INSTANCE_ID ||
psIn->eSystemValueType == NAME_PRIMITIVE_ID ||
if (psIn->eSystemValueType == NAME_PRIMITIVE_ID ||
psIn->eSystemValueType == NAME_VERTEX_ID ||
psIn->eSystemValueType == NAME_INSTANCE_ID ||
psIn->eSystemValueType == NAME_RENDER_TARGET_ARRAY_INDEX ||
psIn->eSystemValueType == NAME_VIEWPORT_ARRAY_INDEX ||
psIn->eSystemValueType == NAME_SAMPLE_INDEX
)
{
return SVT_INT;
}
psIn->eSystemValueType == NAME_SAMPLE_INDEX)
return (psContext->psShader->eTargetLanguage == LANG_METAL) ? SVT_UINT : SVT_INT;
if (psIn->eMinPrec != MIN_PRECISION_DEFAULT)
{

9
src/decode.cpp
View File

@ -311,15 +311,8 @@ uint32_t DecodeOperand(const uint32_t *pui32Tokens, Operand* psOperand)
}
}
if (psOperand->eType == OPERAND_TYPE_OUTPUT_DEPTH_GREATER_EQUAL ||
psOperand->eType == OPERAND_TYPE_OUTPUT_DEPTH_LESS_EQUAL)
{
psOperand->ui32RegisterNumber = -1;
psOperand->ui32CompMask = -1;
}
// Used only for Metal
if (psOperand->eType == OPERAND_TYPE_OUTPUT_DEPTH)
if (psOperand->eType == OPERAND_TYPE_OUTPUT_DEPTH_GREATER_EQUAL || psOperand->eType == OPERAND_TYPE_OUTPUT_DEPTH_LESS_EQUAL || psOperand->eType == OPERAND_TYPE_OUTPUT_DEPTH)
{
psOperand->ui32RegisterNumber = 0;
psOperand->ui32CompMask = 1;

4
src/internal_includes/HLSLCrossCompilerContext.h
View File

@ -32,6 +32,10 @@ public:
// Returns true if VULKAN_BINDINGS flag is set
bool IsVulkan() const;
// Helper functions for checking flags
// Returns true if HLSLCC_FLAG_NVN_TARGET flag is set
bool IsSwitch() const;
Shader* psShader;
GLSLCrossDependencyData* psDependencies;
const char *inputPrefix; // Prefix for shader inputs

3
src/internal_includes/HLSLccToolkit.h
View File

@ -61,6 +61,9 @@ namespace HLSLcc
// Helper function to print floats with full precision
void PrintFloat(bstring b, float f);
bstring GetEarlyMain(HLSLCrossCompilerContext *psContext);
bstring GetPostShaderCode(HLSLCrossCompilerContext *psContext);
// Flags for ForeachOperand
// Process suboperands
#define FEO_FLAG_SUBOPERAND 1

1
src/internal_includes/toGLSL.h
View File

@ -122,6 +122,7 @@ private:
// A <function name, body text> map of extra helper functions we'll need.
FunctionDefinitions m_FunctionDefinitions;
std::vector<std::string> m_FunctionDefinitionsOrder;
std::set<uint32_t> m_DeclaredRenderTarget;
int m_NumDeclaredWhileTrueLoops;

4
src/internal_includes/toMetal.h
View File

@ -57,8 +57,8 @@ private:
void DeclareConstantBuffer(const ConstantBuffer *psCBuf, uint32_t ui32BindingPoint);
void DeclareStructType(const std::string &name, const std::vector<ShaderVar> &contents, bool withinCB = false, uint32_t cumulativeOffset = 0, bool stripUnused = false);
void DeclareStructType(const std::string &name, const std::vector<ShaderVarType> &contents, bool withinCB = false, uint32_t cumulativeOffset = 0);
void DeclareStructVariable(const std::string &parentName, const ShaderVar &var, bool withinCB = false, uint32_t cumulativeOffset = 0);
void DeclareStructVariable(const std::string &parentName, const ShaderVarType &var, bool withinCB = false, uint32_t cumulativeOffset = 0);
void DeclareStructVariable(const std::string &parentName, const ShaderVar &var, bool withinCB = false, uint32_t cumulativeOffset = 0, bool isUsed = true);
void DeclareStructVariable(const std::string &parentName, const ShaderVarType &var, bool withinCB = false, uint32_t cumulativeOffset = 0, bool isUsed = true);
void DeclareBufferVariable(const Declaration *psDecl, bool isRaw, bool isUAV);
void DeclareResource(const Declaration *psDecl);

17
src/reflect.cpp
View File

@ -192,7 +192,7 @@ static void ReadPatchConstantSignatures(const uint32_t* pui32Tokens,
}
}
static const uint32_t* ReadResourceBinding(const uint32_t* pui32FirstResourceToken, const uint32_t* pui32Tokens, ResourceBinding* psBinding, uint32_t decodeFlags)
static const uint32_t* ReadResourceBinding(ShaderInfo* psShaderInfo, const uint32_t* pui32FirstResourceToken, const uint32_t* pui32Tokens, ResourceBinding* psBinding, uint32_t decodeFlags)
{
uint32_t ui32NameOffset = *pui32Tokens++;
@ -202,10 +202,17 @@ static const uint32_t* ReadResourceBinding(const uint32_t* pui32FirstResourceTok
psBinding->eType = (ResourceType) * pui32Tokens++;
psBinding->ui32ReturnType = (RESOURCE_RETURN_TYPE)*pui32Tokens++;
psBinding->eDimension = (REFLECT_RESOURCE_DIMENSION)*pui32Tokens++;
psBinding->ui32NumSamples = *pui32Tokens++;
psBinding->ui32NumSamples = *pui32Tokens++; // fxc generates 2^32 - 1 for non MS images
psBinding->ui32BindPoint = *pui32Tokens++;
psBinding->ui32BindCount = *pui32Tokens++;
psBinding->ui32Flags = *pui32Tokens++;
if (((psShaderInfo->ui32MajorVersion >= 5) && (psShaderInfo->ui32MinorVersion >= 1)) ||
(psShaderInfo->ui32MajorVersion > 5))
{
psBinding->ui32Space = *pui32Tokens++;
psBinding->ui32RangeID = *pui32Tokens++;
}
psBinding->ePrecision = REFLECT_RESOURCE_PRECISION_UNKNOWN;
if (decodeFlags & HLSLCC_FLAG_SAMPLER_PRECISION_ENCODED_IN_NAME)
@ -337,9 +344,9 @@ static const uint32_t* ReadConstantBuffer(ShaderInfo* psShaderInfo,
if (psShaderInfo->ui32MajorVersion >= 5)
{
/*uint32_t StartTexture = * */ pui32VarToken++;
/*uint32_t TextureSize = * */pui32VarToken++;
/*uint32_t TextureSize = * */ pui32VarToken++;
/*uint32_t StartSampler = * */ pui32VarToken++;
/*uint32_t SamplerSize = * */pui32VarToken++;
/*uint32_t SamplerSize = * */ pui32VarToken++;
}
psVar->haveDefaultValue = 0;
@ -407,7 +414,7 @@ static void ReadResources(const uint32_t* pui32Tokens,//in
for (i = 0; i < ui32NumResourceBindings; ++i)
{
pui32ResourceBindings = ReadResourceBinding(pui32FirstToken, pui32ResourceBindings, psResBindings + i, decodeFlags);
pui32ResourceBindings = ReadResourceBinding(psShaderInfo, pui32FirstToken, pui32ResourceBindings, psResBindings + i, decodeFlags);
ASSERT(psResBindings[i].ui32BindPoint < MAX_RESOURCE_BINDINGS);
}

97
src/toGLSL.cpp
View File

@ -308,11 +308,20 @@ static void AddVersionDependentCode(HLSLCrossCompilerContext* psContext)
if (psContext->psShader->eShaderType == PIXEL_SHADER &&
(psContext->psShader->eTargetLanguage == LANG_ES_100 || psContext->psShader->eTargetLanguage == LANG_ES_300 || psContext->psShader->eTargetLanguage == LANG_ES_310 || (psContext->flags & HLSLCC_FLAG_NVN_TARGET)))
{
if ((psContext->flags & HLSLCC_FLAG_VULKAN_BINDINGS) || (psContext->flags & HLSLCC_FLAG_NVN_TARGET))
bcatcstr(glsl, "precision highp float;\n");
else if (psContext->psShader->eTargetLanguage == LANG_ES_100)
if (psContext->psShader->eTargetLanguage == LANG_ES_100)
{
// gles 2.0 shaders can have mediump as default if the GPU doesn't have highp support
bcatcstr(glsl, "#ifdef GL_FRAGMENT_PRECISION_HIGH\nprecision highp float;\n#else\nprecision mediump float;\n#endif\n");
bcatcstr(glsl,
"#ifdef GL_FRAGMENT_PRECISION_HIGH\n"
" precision highp float;\n"
"#else\n"
" precision mediump float;\n"
"#endif\n");
}
else
{
bcatcstr(glsl, "precision highp float;\n");
}
// Define default int precision to highp to avoid issues on platforms that actually implement mediump
bcatcstr(glsl, "precision highp int;\n");
@ -495,19 +504,29 @@ static void DoHullShaderPassthrough(HLSLCrossCompilerContext *psContext)
outputName = oss.str();
}
const char * prec = HavePrecisionQualifiers(psContext) ? "highp " : "";
const char * prec = "";
if (HavePrecisionQualifiers(psContext))
{
if (psSig->eMinPrec != MIN_PRECISION_DEFAULT)
prec = "mediump ";
else
prec = "highp ";
}
int inLoc = psContext->psDependencies->GetVaryingLocation(inputName, HULL_SHADER, true);
int outLoc = psContext->psDependencies->GetVaryingLocation(outputName, HULL_SHADER, false);
psContext->AddIndentation();
if (ui32NumComponents > 1) // TODO Precision
bformata(glsl, "in %s%s%d %s%s%d[];\n", prec, Type, ui32NumComponents, psContext->inputPrefix, psSig->semanticName.c_str(), psSig->ui32SemanticIndex);
if (ui32NumComponents > 1)
bformata(glsl, "layout(location = %d) in %s%s%d %s%s%d[];\n", inLoc, prec, Type, ui32NumComponents, psContext->inputPrefix, psSig->semanticName.c_str(), psSig->ui32SemanticIndex);
else
bformata(glsl, "in %s%s %s%s%d[];\n", prec, Type, psContext->inputPrefix, psSig->semanticName.c_str(), psSig->ui32SemanticIndex);
bformata(glsl, "layout(location = %d) in %s%s %s%s%d[];\n", inLoc, prec, Type, psContext->inputPrefix, psSig->semanticName.c_str(), psSig->ui32SemanticIndex);
psContext->AddIndentation();
if (ui32NumComponents > 1) // TODO Precision
bformata(glsl, "out %s%s%d %s%s%d[];\n", prec, Type, ui32NumComponents, psContext->outputPrefix, psSig->semanticName.c_str(), psSig->ui32SemanticIndex);
if (ui32NumComponents > 1)
bformata(glsl, "layout(location = %d) out %s%s%d %s%s%d[];\n", inLoc, prec, Type, ui32NumComponents, psContext->outputPrefix, psSig->semanticName.c_str(), psSig->ui32SemanticIndex);
else
bformata(glsl, "out %s%s %s%s%d[];\n", prec, Type, psContext->outputPrefix, psSig->semanticName.c_str(), psSig->ui32SemanticIndex);
bformata(glsl, "layout(location = %d) out %s%s %s%s%d[];\n", inLoc, prec, Type, psContext->outputPrefix, psSig->semanticName.c_str(), psSig->ui32SemanticIndex);
}
psContext->AddIndentation();
@ -624,7 +643,7 @@ bool ToGLSL::Translate()
psContext->DoDataTypeAnalysis(&phase);
phase.ResolveUAVProperties();
psShader->ResolveStructuredBufferBindingSlots(&phase);
if (!psContext->IsVulkan())
if (!psContext->IsVulkan() && !psContext->IsSwitch())
phase.PruneConstArrays();
}
@ -796,7 +815,26 @@ bool ToGLSL::Translate()
bcatcstr(glsl, "}\n");
// Print out extra functions we generated, in reverse order for potential dependencies
std::for_each(m_FunctionDefinitions.rbegin(), m_FunctionDefinitions.rend(), [&extensions](const FunctionDefinitions::value_type &p)
{
bcatcstr(extensions, p.second.c_str());
bcatcstr(extensions, "\n");
});
// Concat extensions and glsl for the final shader code.
if (m_NeedUnityInstancingArraySizeDecl)
{
if (psContext->flags & HLSLCC_FLAG_VULKAN_BINDINGS)
{
bformata(extensions, "layout(constant_id = %d) const int %s = 2;\n", kArraySizeConstantID, UNITY_RUNTIME_INSTANCING_ARRAY_SIZE_MACRO);
}
else
{
bcatcstr(extensions, "#ifndef " UNITY_RUNTIME_INSTANCING_ARRAY_SIZE_MACRO "\n\t#define " UNITY_RUNTIME_INSTANCING_ARRAY_SIZE_MACRO " 2\n#endif\n");
}
}
bconcat(extensions, glsl);
bdestroy(glsl);
psContext->glsl = extensions;
@ -858,6 +896,7 @@ bool ToGLSL::Translate()
}
}
bstring generatedFunctionsKeyword = bfromcstr("\n// Generated functions\n\n");
bstring beforeMain = NULL;
bstring beforeMainKeyword = NULL;
@ -878,6 +917,9 @@ bool ToGLSL::Translate()
DeclareSpecializationConstants(psShader->asPhases[0]);
}
// Search and replace string, for injecting generated functions that need to be after default precision declarations
bconcat(glsl, generatedFunctionsKeyword);
// Search and replace string, for injecting stuff from translation that need to be after normal declarations and before main
if (!HaveDynamicIndexing(psContext))
{
@ -910,12 +952,19 @@ bool ToGLSL::Translate()
bcatcstr(glsl, "}\n");
// Print out extra functions we generated, in reverse order for potential dependencies
std::for_each(m_FunctionDefinitions.rbegin(), m_FunctionDefinitions.rend(), [&extensions](const FunctionDefinitions::value_type &p)
// Print out extra functions we generated in generation order to satisfy dependencies
{
bcatcstr(extensions, p.second.c_str());
bcatcstr(extensions, "\n");
});
bstring generatedFunctions = bfromcstr("");
for (std::vector<std::string>::const_iterator funcNameIter = m_FunctionDefinitionsOrder.begin(); funcNameIter != m_FunctionDefinitionsOrder.end(); ++funcNameIter)
{
const FunctionDefinitions::const_iterator definition = m_FunctionDefinitions.find(*funcNameIter);
ASSERT(definition != m_FunctionDefinitions.end());
bcatcstr(generatedFunctions, definition->second.c_str());
bcatcstr(generatedFunctions, "\n");
}
bfindreplace(glsl, generatedFunctionsKeyword, generatedFunctions, 0);
bdestroy(generatedFunctions);
}
// Concat extensions and glsl for the final shader code.
if (m_NeedUnityInstancingArraySizeDecl)
@ -959,6 +1008,7 @@ bool ToGLSL::DeclareExtraFunction(const std::string &name, bstring body)
if (m_FunctionDefinitions.find(name) != m_FunctionDefinitions.end())
return true;
m_FunctionDefinitions.insert(std::make_pair(name, (const char *)body->data));
m_FunctionDefinitionsOrder.push_back(name);
return false;
}
@ -1010,14 +1060,14 @@ void ToGLSL::UseExtraFunctionDependency(const std::string &name)
{
UseExtraFunctionDependency("op_modi");
bformata(code, "int op_and(int a, int b) { int result = 0; int n = 1; for (int i = 0; i < BITWISE_BIT_COUNT; i++) { if ((op_modi(a, 2) == 1) && (op_modi(b, 2) == 1)) { result += n; } a = a / 2; b = b / 2; n = n * 2; if (!(a > 0 && b > 0)) { break; } } return result; }\n");
bformata(code, "int op_and(int a, int b) { int result = 0; int n = 1; for (int i = 0; i < BITWISE_BIT_COUNT; i++) { if ((op_modi(a, 2) != 0) && (op_modi(b, 2) != 0)) { result += n; } a = a / 2; b = b / 2; n = n * 2; if (!(a > 0 && b > 0)) { break; } } return result; }\n");
PrintComponentWrapper2(code, "op_and", "ivec2", "ivec3", "ivec4");
}
else if (name == "op_or")
{
UseExtraFunctionDependency("op_modi");
bformata(code, "int op_or(int a, int b) { int result = 0; int n = 1; for (int i = 0; i < BITWISE_BIT_COUNT; i++) { if ((op_modi(a, 2) == 1) || (op_modi(b, 2) == 1)) { result += n; } a = a / 2; b = b / 2; n = n * 2; if (!(a > 0 || b > 0)) { break; } } return result; }\n");
bformata(code, "int op_or(int a, int b) { int result = 0; int n = 1; for (int i = 0; i < BITWISE_BIT_COUNT; i++) { if ((op_modi(a, 2) != 0) || (op_modi(b, 2) != 0)) { result += n; } a = a / 2; b = b / 2; n = n * 2; if (!(a > 0 || b > 0)) { break; } } return result; }\n");
PrintComponentWrapper2(code, "op_or", "ivec2", "ivec3", "ivec4");
}
else if (name == "op_xor")
@ -1042,6 +1092,15 @@ void ToGLSL::UseExtraFunctionDependency(const std::string &name)
bformata(code, "int op_not(int value) { return -value - 1; }\n");
PrintComponentWrapper1(code, "op_not", "ivec2", "ivec3", "ivec4");
}
else if (name == "int_bitfieldInsert")
{
// Can't use the name 'bitfieldInsert' because Adreno fails with "can't redefine/overload built-in functions!"
bcatcstr(code,
"int int_bitfieldInsert(int base, int insert, int offset, int bits) {\n"
" uint mask = ~(uint(0xffffffff) << uint(bits)) << uint(offset);\n"
" return int((uint(base) & ~mask) | ((uint(insert) << uint(offset)) & mask));\n"
"}\n");
}
else
{
match = false;

138
src/toGLSLDeclaration.cpp
View File

@ -196,7 +196,7 @@ static void DeclareInput(
{
if (regSpace == 0)
{
if (psSig->semanticName == "POS" && psSig->ui32SemanticIndex == 0)
if ((psSig->semanticName == "POS" || psSig->semanticName == "SV_Position") && psSig->ui32SemanticIndex == 0)
return;
}
}
@ -250,8 +250,8 @@ static void DeclareInput(
if (psShader->eShaderType == VERTEX_SHADER && !(psContext->flags & HLSLCC_FLAG_DISABLE_EXPLICIT_LOCATIONS))
addLocation = true;
// Add intra-shader locations if requested in flags
if (psShader->eShaderType != VERTEX_SHADER && (psContext->flags & HLSLCC_FLAG_SEPARABLE_SHADER_OBJECTS))
// Add intra-shader locations if supported
if (psShader->eShaderType != VERTEX_SHADER)
addLocation = true;
if (addLocation)
@ -339,14 +339,14 @@ bool ToGLSL::RenderTargetDeclared(uint32_t input)
void ToGLSL::AddBuiltinInput(const Declaration* psDecl, const char* builtinName)
{
Shader* psShader = psContext->psShader; const Operand* op = &psDecl->asOperands[0];
const int regSpace = op->GetRegisterSpace(psContext); ASSERT(regSpace == 0);
const uint32_t ui32Reg = op->ui32RegisterNumber, ui32CompMask = op->ui32CompMask;
Shader* psShader = psContext->psShader;
const Operand* psOperand = &psDecl->asOperands[0];
const int regSpace = psOperand->GetRegisterSpace(psContext);
ASSERT(regSpace == 0);
// we need to at least mark if they are scalars or not (as we might need to use vector ctor)
if (op->GetNumInputElements(psContext) == 1)
psShader->abScalarInput[regSpace][ui32Reg] |= (int)ui32CompMask;
if (psOperand->GetNumInputElements(psContext) == 1)
psShader->abScalarInput[regSpace][psOperand->ui32RegisterNumber] |= (int)psOperand->ui32CompMask;
}
void ToGLSL::AddBuiltinOutput(const Declaration* psDecl, int arrayElements, const char* builtinName)
@ -471,9 +471,6 @@ void ToGLSL::HandleOutputRedirect(const Declaration *psDecl, const char *Precisi
psContext->AddIndentation();
bformata(glsl, "%s vec4 phase%d_Output%d_%d;\n", Precision, psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
psPhase->hasPostShaderCode = 1;
psContext->currentGLSLString = &psPhase->postShaderCode;
while (comp < 4)
{
int numComps = 0;
@ -496,39 +493,36 @@ void ToGLSL::HandleOutputRedirect(const Declaration *psDecl, const char *Precisi
mask = psSig->ui32Mask;
((Operand *)psOperand)->ui32CompMask = 1 << comp;
psContext->AddIndentation();
TranslateOperand(psOperand, TO_FLAG_NAME_ONLY);
bcatcstr(psPhase->postShaderCode, " = ");
bstring str = GetPostShaderCode(psContext);
TranslateOperand(str, psOperand, TO_FLAG_NAME_ONLY);
bcatcstr(str, " = ");
if (psSig->eComponentType == INOUT_COMPONENT_SINT32)
{
bformata(psPhase->postShaderCode, HaveBitEncodingOps(psContext->psShader->eTargetLanguage) ? "floatBitsToInt(" : "int(");
bformata(str, HaveBitEncodingOps(psContext->psShader->eTargetLanguage) ? "floatBitsToInt(" : "int(");
hasCast = 1;
}
else if (psSig->eComponentType == INOUT_COMPONENT_UINT32)
{
bformata(psPhase->postShaderCode, HaveBitEncodingOps(psContext->psShader->eTargetLanguage) ? "floatBitsToUint(" : "int(");
bformata(str, HaveBitEncodingOps(psContext->psShader->eTargetLanguage) ? "floatBitsToUint(" : "int(");
hasCast = 1;
}
bformata(psPhase->postShaderCode, "phase%d_Output%d_%d.", psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
bformata(str, "phase%d_Output%d_%d.", psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
// Print out mask
for (i = 0; i < 4; i++)
{
if ((mask & (1 << i)) == 0)
continue;
bformata(psPhase->postShaderCode, "%c", "xyzw"[i]);
bformata(str, "%c", "xyzw"[i]);
}
if (hasCast)
bcatcstr(psPhase->postShaderCode, ")");
bcatcstr(str, ")");
comp += numComps;
bcatcstr(psPhase->postShaderCode, ";\n");
bcatcstr(str, ";\n");
}
psContext->currentGLSLString = &psContext->glsl;
((Operand *)psOperand)->ui32CompMask = origMask;
if (regSpace == 0)
psShader->asPhases[psContext->currentPhase].acOutputNeedsRedirect[psOperand->ui32RegisterNumber] = 0xfe;
@ -648,9 +642,7 @@ void ToGLSL::AddUserOutput(const Declaration* psDecl)
{
if (psShader->eTargetLanguage == LANG_ES_100 && !psContext->EnableExtension("GL_EXT_frag_depth"))
{
bcatcstr(psContext->extensions, "#ifdef GL_EXT_frag_depth\n");
bcatcstr(psContext->extensions, "#define gl_FragDepth gl_FragDepthEXT\n");
bcatcstr(psContext->extensions, "#endif\n");
}
break;
}
@ -759,7 +751,7 @@ void ToGLSL::AddUserOutput(const Declaration* psDecl)
}
}
if (HaveInOutLocationQualifier(psContext->psShader->eTargetLanguage) && (psContext->flags & HLSLCC_FLAG_SEPARABLE_SHADER_OBJECTS))
if (HaveInOutLocationQualifier(psContext->psShader->eTargetLanguage))
{
bformata(glsl, "layout(location = %d) ", psContext->psDependencies->GetVaryingLocation(std::string(OutputName), psShader->eShaderType, false));
}
@ -891,6 +883,34 @@ void ToGLSL::DeclareUBOConstants(const uint32_t ui32BindingPoint, const Constant
bformata(glsl, "#endif\n#undef UNITY_UNIFORM\n");
}
bool DeclareRWStructuredBufferTemplateTypeAsInteger(HLSLCrossCompilerContext* psContext, const Operand* psOperand)
{
// with cases like: RWStructuredBuffer<int4> myBuffer; /*...*/ AtomicMin(myBuffer[0].x , myInt);
// if we translate RWStructuredBuffer template type to uint, incorrect version of the function might be called ( AtomicMin(uint..) instead of AtomicMin(int..) )
// we try to avoid this case by using integer type in those cases
if (psContext && psOperand)
{
const bool isVulkan = (psContext->flags & HLSLCC_FLAG_VULKAN_BINDINGS) != 0;
if (!isVulkan)
{
if (psContext->psShader && HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
{
uint32_t ui32BindingPoint = psOperand->ui32RegisterNumber;
const ResourceBinding* psBinding = NULL;
psContext->psShader->sInfo.GetResourceFromBindingPoint(RGROUP_UAV, ui32BindingPoint, &psBinding);
if (psBinding)
{
const ConstantBuffer* psBuffer = NULL;
psContext->psShader->sInfo.GetConstantBufferFromBindingPoint(RGROUP_UAV, psBinding->ui32BindPoint, &psBuffer);
if (psBuffer && psBuffer->asVars.size() == 1 && psBuffer->asVars[0].sType.Type == SVT_INT /*&& psContext->IsSwitch()*/)
return true;
}
}
}
}
return false;
}
static void DeclareBufferVariable(HLSLCrossCompilerContext* psContext, uint32_t ui32BindingPoint,
const Operand* psOperand, const uint32_t ui32GloballyCoherentAccess,
const uint32_t isRaw, const uint32_t isUAV, const uint32_t hasEmbeddedCounter, const uint32_t stride, bstring glsl)
@ -910,7 +930,12 @@ static void DeclareBufferVariable(HLSLCrossCompilerContext* psContext, uint32_t
// Declare the struct type for structured buffers
if (!isRaw)
bformata(glsl, " struct %s_type {\n\tuint[%d] value;\n};\n\n", BufName.c_str(), stride / 4);
{
const char* typeStr = "uint";
if (isUAV && DeclareRWStructuredBufferTemplateTypeAsInteger(psContext, psOperand))
typeStr = "int";
bformata(glsl, " struct %s_type {\n\t%s[%d] value;\n};\n\n", BufName.c_str(), typeStr, stride / 4);
}
if (isVulkan)
{
@ -1513,18 +1538,17 @@ void ToGLSL::HandleInputRedirect(const Declaration *psDecl, const char *Precisio
else
bformata(glsl, "%s vec4 phase%d_Input%d_%d;\n", Precision, psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
psContext->currentGLSLString = &psPhase->earlyMain;
psContext->indent++;
// Do a conditional loop. In normal cases needsLooping == 0 so this is only run once.
do
{
int comp = 0;
psContext->AddIndentation();
bstring str = GetEarlyMain(psContext);
if (needsLooping)
bformata(psPhase->earlyMain, "phase%d_Input%d_%d[%d] = vec4(", psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber, i);
bformata(str, "phase%d_Input%d_%d[%d] = vec4(", psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber, i);
else
bformata(psPhase->earlyMain, "phase%d_Input%d_%d = vec4(", psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
bformata(str, "phase%d_Input%d_%d = vec4(", psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
while (comp < 4)
{
@ -1541,12 +1565,12 @@ void ToGLSL::HandleInputRedirect(const Declaration *psDecl, const char *Precisio
numComps = GetNumberBitsSet(psSig->ui32Mask);
if (psSig->eComponentType == INOUT_COMPONENT_SINT32)
{
bformata(psPhase->earlyMain, HaveBitEncodingOps(psContext->psShader->eTargetLanguage) ? "intBitsToFloat(" : "float(");
bformata(str, HaveBitEncodingOps(psContext->psShader->eTargetLanguage) ? "intBitsToFloat(" : "float(");
hasCast = 1;
}
else if (psSig->eComponentType == INOUT_COMPONENT_UINT32)
{
bformata(psPhase->earlyMain, HaveBitEncodingOps(psContext->psShader->eTargetLanguage) ? "uintBitsToFloat(" : "float(");
bformata(str, HaveBitEncodingOps(psContext->psShader->eTargetLanguage) ? "uintBitsToFloat(" : "float(");
hasCast = 1;
}
@ -1557,7 +1581,7 @@ void ToGLSL::HandleInputRedirect(const Declaration *psDecl, const char *Precisio
// And the component mask
psOperand->ui32CompMask = 1 << comp;
TranslateOperand(psOperand, TO_FLAG_NAME_ONLY);
TranslateOperand(str, psOperand, TO_FLAG_NAME_ONLY);
// Restore the original array size value and mask
psOperand->ui32CompMask = origMask;
@ -1565,23 +1589,22 @@ void ToGLSL::HandleInputRedirect(const Declaration *psDecl, const char *Precisio
psOperand->aui32ArraySizes[0] = origArraySize;
if (hasCast)
bcatcstr(psPhase->earlyMain, ")");
bcatcstr(str, ")");
comp += numComps;
}
else // no signature found -> fill with zero
{
bcatcstr(psPhase->earlyMain, "0");
bcatcstr(str, "0");
comp++;
}
if (comp < 4)
bcatcstr(psPhase->earlyMain, ", ");
bcatcstr(str, ", ");
}
bcatcstr(psPhase->earlyMain, ");\n");
bcatcstr(str, ");\n");
}
while ((--i) >= 0);
psContext->currentGLSLString = &psContext->glsl;
psContext->indent--;
if (regSpace == 0)
@ -1648,13 +1671,16 @@ void ToGLSL::TranslateDeclaration(const Declaration* psDecl)
*/
if (HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
AddBuiltinInput(psDecl, "(gl_FrontFacing ? 0xffffffffu : uint(0))"); // Hi Adreno.
AddBuiltinInput(psDecl, "(gl_FrontFacing ? 0xffffffffu : uint(0))"); // Old ES3.0 Adrenos treat 0u as const int
else
AddBuiltinInput(psDecl, "(gl_FrontFacing ? int(1) : int(0))");
AddBuiltinInput(psDecl, "(gl_FrontFacing ? 1 : 0)");
break;
}
case NAME_SAMPLE_INDEX:
{
// Using gl_SampleID requires either GL_OES_sample_variables or #version 320 es
if (IsESLanguage(psContext->psShader->eTargetLanguage))
psContext->RequireExtension("GL_OES_sample_variables");
AddBuiltinInput(psDecl, "gl_SampleID");
break;
}
@ -1693,7 +1719,10 @@ void ToGLSL::TranslateDeclaration(const Declaration* psDecl)
AddBuiltinOutput(psDecl, 0, "gl_Layer");
if (psShader->eShaderType == VERTEX_SHADER)
{
psContext->RequireExtension("GL_AMD_vertex_shader_layer");
if (psContext->IsVulkan())
psContext->RequireExtension("GL_ARB_shader_viewport_layer_array");
else
psContext->RequireExtension("GL_AMD_vertex_shader_layer");
}
break;
@ -1955,6 +1984,7 @@ void ToGLSL::TranslateDeclaration(const Declaration* psDecl)
case NAME_POSITION:
{
AddBuiltinInput(psDecl, "gl_FragCoord");
bcatcstr(GetEarlyMain(psContext), "vec4 hlslcc_FragCoord = vec4(gl_FragCoord.xyz, 1.0/gl_FragCoord.w);\n");
break;
}
default:
@ -2669,6 +2699,22 @@ void ToGLSL::TranslateDeclaration(const Declaration* psDecl)
});
bcatcstr(glsl, ");\n");
}
else if (psContext->IsSwitch())
{
bstring glsl = *psContext->currentGLSLString;
bformata(glsl, "const vec4 ImmCB_%d[] = vec4[%d] (\n", psContext->currentPhase, psDecl->asImmediateConstBuffer.size());
bool isFirst = true;
std::for_each(psDecl->asImmediateConstBuffer.begin(), psDecl->asImmediateConstBuffer.end(), [&](const ICBVec4 &data)
{
if (!isFirst)
{
bcatcstr(glsl, ",\n");
}
isFirst = false;
bformata(glsl, "vec4(uintBitsToFloat(uint(0x%Xu)), uintBitsToFloat(uint(0x%Xu)), uintBitsToFloat(uint(0x%Xu)), uintBitsToFloat(uint(0x%Xu)))", data.a, data.b, data.c, data.d);
});
bcatcstr(glsl, ");\n");
}
else
{
// TODO: This is only ever accessed as a float currently. Do trickery if we ever see ints accessed from an array.
@ -3343,7 +3389,7 @@ bool ToGLSL::TranslateSystemValue(const Operand *psOperand, const ShaderInfo::In
{
case NAME_POSITION:
if (psContext->psShader->eShaderType == PIXEL_SHADER)
result = "gl_FragCoord";
result = "hlslcc_FragCoord";
else
result = "gl_Position";
return true;
@ -3385,9 +3431,9 @@ bool ToGLSL::TranslateSystemValue(const Operand *psOperand, const ShaderInfo::In
return true;
case NAME_IS_FRONT_FACE:
if (HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
result = "(gl_FrontFacing ? 0xffffffffu : uint(0))";
result = "(gl_FrontFacing ? 0xffffffffu : uint(0))"; // Old ES3.0 Adrenos treat 0u as const int
else
result = "(gl_FrontFacing ? int(1) : int(0))";
result = "(gl_FrontFacing ? 1 : 0)";
if (pui32IgnoreSwizzle)
*pui32IgnoreSwizzle = 1;
return true;

151
src/toGLSLInstruction.cpp
View File

@ -18,6 +18,7 @@ using namespace HLSLcc;
const char* GetSamplerType(HLSLCrossCompilerContext* psContext,
const RESOURCE_DIMENSION eDimension,
const uint32_t ui32RegisterNumber);
bool DeclareRWStructuredBufferTemplateTypeAsInteger(HLSLCrossCompilerContext* psContext, const Operand* psOperand);
// This function prints out the destination name, possible destination writemask, assignment operator
// and any possible conversions needed based on the eSrcType+ui32SrcElementCount (type and size of data expected to be coming in)
@ -209,7 +210,7 @@ void ToGLSL::AddComparison(Instruction* psInst, ComparisonType eType,
if (HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
bcatcstr(glsl, " * 0xFFFFFFFFu");
else
bcatcstr(glsl, " * 0xFFFFFFFF");
bcatcstr(glsl, " * -1"); // GLSL ES 2 spec: high precision ints are guaranteed to have a range of at least (-2^16, 2^16)
}
}
@ -258,10 +259,9 @@ void ToGLSL::AddComparison(Instruction* psInst, ComparisonType eType,
bcatcstr(glsl, "1.0 : 0.0");
else
{
if (HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
bcatcstr(glsl, "0xFFFFFFFFu : uint(0u)"); // Adreno can't handle 0u.
else
bcatcstr(glsl, "0xFFFFFFFF : int(0)");
// Old ES3.0 Adrenos treat 0u as const int.
// GLSL ES 2 spec: high precision ints are guaranteed to have a range of at least (-2^16, 2^16)
bcatcstr(glsl, HaveUnsignedTypes(psContext->psShader->eTargetLanguage) ? ") ? 0xFFFFFFFFu : uint(0)" : ") ? -1 : 0");
}
AddAssignPrologue(needsParenthesis, true);
bcatcstr(glsl, "; }\n");
@ -288,15 +288,12 @@ void ToGLSL::AddComparison(Instruction* psInst, ComparisonType eType,
if (!isBoolDest)
{
if (floatResult)
{
bcatcstr(glsl, ") ? 1.0 : 0.0");
}
else
{
if (HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
bcatcstr(glsl, ") ? 0xFFFFFFFFu : uint(0u)"); // Adreno can't handle 0u.
else
bcatcstr(glsl, ") ? 0xFFFFFFFF : int(0)");
// Old ES3.0 Adrenos treat 0u as const int.
// GLSL ES 2 spec: high precision ints are guaranteed to have a range of at least (-2^16, 2^16)
bcatcstr(glsl, HaveUnsignedTypes(psContext->psShader->eTargetLanguage) ? ") ? 0xFFFFFFFFu : uint(0)" : ") ? -1 : 0");
}
}
AddAssignPrologue(needsParenthesis);
@ -369,10 +366,7 @@ void ToGLSL::AddMOVCBinaryOp(const Operand *pDest, const Operand *src0, Operand
else
{
if (s0Type == SVT_UINT || s0Type == SVT_UINT16)
if (HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
bcatcstr(glsl, " != uint(0u)) ? "); // Adreno doesn't understand 0u.
else
bcatcstr(glsl, " != int(0)) ? ");
bcatcstr(glsl, HaveUnsignedTypes(psContext->psShader->eTargetLanguage) ? " != uint(0)) ? " : " != 0) ? "); // Old ES3.0 Adrenos treat 0u as const int.
else if (s0Type == SVT_BOOL)
bcatcstr(glsl, ") ? ");
else
@ -396,13 +390,18 @@ void ToGLSL::AddMOVCBinaryOp(const Operand *pDest, const Operand *src0, Operand
{
// TODO: We can actually do this in one op using mix().
int srcElem = -1;
SHADER_VARIABLE_TYPE dstType = pDest->GetDataType(psContext);
SHADER_VARIABLE_TYPE s0Type = src0->GetDataType(psContext);
// Use an extra temp if dest is also one of the sources. Without this some swizzle combinations
// might alter the source before all components are handled.
const char* tempName = "hlslcc_movcTemp";
bool dstIsSrc1 = (pDest->eType == src1->eType) && (pDest->ui32RegisterNumber == src1->ui32RegisterNumber);
bool dstIsSrc2 = (pDest->eType == src2->eType) && (pDest->ui32RegisterNumber == src2->ui32RegisterNumber);
const std::string tempName = "hlslcc_movcTemp";
bool dstIsSrc1 = (pDest->eType == src1->eType)
&& (dstType == src1->GetDataType(psContext))
&& (pDest->ui32RegisterNumber == src1->ui32RegisterNumber);
bool dstIsSrc2 = (pDest->eType == src2->eType)
&& (dstType == src2->GetDataType(psContext))
&& (pDest->ui32RegisterNumber == src2->ui32RegisterNumber);
if (dstIsSrc1 || dstIsSrc2)
{
@ -413,9 +412,14 @@ void ToGLSL::AddMOVCBinaryOp(const Operand *pDest, const Operand *src0, Operand
int numComponents = (pDest->eType == OPERAND_TYPE_TEMP) ?
psContext->psShader->GetTempComponentCount(eDestType, pDest->ui32RegisterNumber) :
pDest->iNumComponents;
bformata(glsl, "%s %s = ", HLSLcc::GetConstructorForType(psContext, eDestType, numComponents), tempName);
TranslateOperand(glsl, pDest, TO_FLAG_NAME_ONLY);
bcatcstr(glsl, ";\n");
const char* constructorStr = HLSLcc::GetConstructorForType(psContext, eDestType, numComponents, false);
bformata(glsl, "%s %s = ", constructorStr, tempName.c_str());
TranslateOperand(pDest, TO_FLAG_NAME_ONLY);
bformata(glsl, ";\n");
// Override OPERAND_TYPE_TEMP name temporarily
const_cast<Operand *>(pDest)->specialName.assign(tempName);
}
for (destElem = 0; destElem < 4; ++destElem)
@ -448,29 +452,20 @@ void ToGLSL::AddMOVCBinaryOp(const Operand *pDest, const Operand *src0, Operand
}
}
if (!dstIsSrc1)
TranslateOperand(src1, SVTTypeToFlag(eDestType), 1 << srcElem);
else
{
bformata(glsl, "%s", tempName);
TranslateOperandSwizzleWithMask(glsl, psContext, src1, 1 << srcElem, 0);
}
TranslateOperand(src1, SVTTypeToFlag(eDestType), 1 << srcElem);
bcatcstr(glsl, " : ");
if (!dstIsSrc2)
TranslateOperand(src2, SVTTypeToFlag(eDestType), 1 << srcElem);
else
{
bformata(glsl, "%s", tempName);
TranslateOperandSwizzleWithMask(glsl, psContext, src2, 1 << srcElem, 0);
}
TranslateOperand(src2, SVTTypeToFlag(eDestType), 1 << srcElem);
AddAssignPrologue(numParenthesis);
}
if (dstIsSrc1 || dstIsSrc2)
{
const_cast<Operand *>(pDest)->specialName.clear();
psContext->AddIndentation();
TranslateOperand(glsl, pDest, TO_FLAG_NAME_ONLY);
bformata(glsl, " = %s;\n", tempName.c_str());
--psContext->indent;
psContext->AddIndentation();
bcatcstr(glsl, "}\n");
@ -530,12 +525,9 @@ void ToGLSL::CallBinaryOp(const char* name, Instruction* psInst,
AddAssignToDest(&psInst->asOperands[dest], eDataType, dstSwizCount, &needsParenthesis);
// Horrible Adreno bug workaround:
// All pre-ES3.1 Adreno GLES3.0 drivers fail in cases like this:
// vec4 a.xyz = b.xyz + c.yzw;
// Attempt to detect this and fall back to component-wise binary op.
if ((psContext->psShader->eTargetLanguage == LANG_ES_300) &&
((src0AccessCount > 1 && !(src0AccessMask & OPERAND_4_COMPONENT_MASK_X)) || (src1AccessCount > 1 && !(src1AccessMask & OPERAND_4_COMPONENT_MASK_X))))
// Adreno 3xx fails on binary ops that operate on vectors
bool opComponentWiseOnAdreno = (!strcmp("&", name) || !strcmp("|", name) || !strcmp("^", name) || !strcmp(">>", name) || !strcmp("<<", name));
if (psContext->psShader->eTargetLanguage == LANG_ES_300 && opComponentWiseOnAdreno)
{
uint32_t i;
int firstPrinted = 0;
@ -566,6 +558,7 @@ void ToGLSL::CallBinaryOp(const char* name, Instruction* psInst,
bformata(glsl, " %s ", name);
TranslateOperand(&psInst->asOperands[src1], ui32Flags, destMask);
}
AddAssignPrologue(needsParenthesis, isEmbedded);
}
@ -997,10 +990,7 @@ void ToGLSL::GetResInfoData(Instruction* psInst, int index, int destElem)
bcatcstr(glsl, "(");
if (dim < (index + 1))
{
if (HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
bcatcstr(glsl, eResInfoReturnType == RESINFO_INSTRUCTION_RETURN_UINT ? "uint(0u)" : "0.0");
else
bcatcstr(glsl, eResInfoReturnType == RESINFO_INSTRUCTION_RETURN_UINT ? "int(0)" : "0.0");
bcatcstr(glsl, eResInfoReturnType == RESINFO_INSTRUCTION_RETURN_UINT ? (HaveUnsignedTypes(psContext->psShader->eTargetLanguage) ? "uint(0)" : "0") : "0.0"); // Old ES3.0 Adrenos treat 0u as const int.
}
else
{
@ -1099,6 +1089,12 @@ void ToGLSL::TranslateTextureSample(Instruction* psInst,
{
offset = "Offset";
}
if (psContext->IsSwitch() && psInst->eOpcode == OPCODE_GATHER4_PO)
{
// it seems that other GLSLCore compilers accept textureGather(sampler2D sampler, vec2 texCoord, ivec2 texelOffset, int component) with the "texelOffset" parameter,
// however this is not in the GLSL spec, and Switch's GLSLc compiler requires to use the textureGatherOffset version of the function
offset = "Offset";
}
switch (eResDim)
{
@ -1193,8 +1189,7 @@ void ToGLSL::TranslateTextureSample(Instruction* psInst,
//
// Here we create a temp texcoord var with the reference value embedded
if ((ui32Flags & TEXSMP_FLAG_DEPTHCOMPARE) &&
eResDim != RESOURCE_DIMENSION_TEXTURECUBEARRAY &&
!(ui32Flags & TEXSMP_FLAG_GATHER))
(eResDim != RESOURCE_DIMENSION_TEXTURECUBEARRAY && !(ui32Flags & TEXSMP_FLAG_GATHER)))
{
uniqueNameCounter = psContext->psShader->asPhases[psContext->currentPhase].m_NextTexCoordTemp++;
psContext->AddIndentation();
@ -1261,8 +1256,7 @@ void ToGLSL::TranslateTextureSample(Instruction* psInst,
// Texture coordinates, either from previously constructed temp
// or straight from the psDestAddr operand
if ((ui32Flags & TEXSMP_FLAG_DEPTHCOMPARE) &&
eResDim != RESOURCE_DIMENSION_TEXTURECUBEARRAY &&
!(ui32Flags & TEXSMP_FLAG_GATHER))
(eResDim != RESOURCE_DIMENSION_TEXTURECUBEARRAY && !(ui32Flags & TEXSMP_FLAG_GATHER)))
bformata(glsl, "txVec%d", uniqueNameCounter);
else
TranslateTexCoord(eResDim, psDestAddr);
@ -1270,8 +1264,7 @@ void ToGLSL::TranslateTextureSample(Instruction* psInst,
// If depth compare reference was not embedded to texcoord
// then insert it here as a separate param
if ((ui32Flags & TEXSMP_FLAG_DEPTHCOMPARE) &&
eResDim == RESOURCE_DIMENSION_TEXTURECUBEARRAY &&
(ui32Flags & TEXSMP_FLAG_GATHER))
(eResDim == RESOURCE_DIMENSION_TEXTURECUBEARRAY || (ui32Flags & TEXSMP_FLAG_GATHER)))
{
bcatcstr(glsl, ", ");
TranslateOperand(psSrcRef, TO_AUTO_BITCAST_TO_FLOAT);
@ -1365,7 +1358,7 @@ void ToGLSL::TranslateTextureSample(Instruction* psInst,
}
else
{
// Comp selection not supported with dephth compare gather
// Component selection not supported with depth compare gather
}
}
}
@ -1490,12 +1483,15 @@ void ToGLSL::TranslateShaderStorageStore(Instruction* psInst)
bcatcstr(glsl, "]");
//Dest type is currently always a uint array.
uint32_t srcFlag = TO_FLAG_UNSIGNED_INTEGER;
if (DeclareRWStructuredBufferTemplateTypeAsInteger(psContext, psDest))
srcFlag = TO_FLAG_INTEGER;
bcatcstr(glsl, " = ");
if (psSrc->GetNumSwizzleElements() > 1)
TranslateOperand(psSrc, TO_FLAG_UNSIGNED_INTEGER, 1 << (srcComponent++));
TranslateOperand(psSrc, srcFlag, 1 << (srcComponent++));
else
TranslateOperand(psSrc, TO_FLAG_UNSIGNED_INTEGER, OPERAND_4_COMPONENT_MASK_X);
TranslateOperand(psSrc, srcFlag, OPERAND_4_COMPONENT_MASK_X);
bcatcstr(glsl, ";\n");
}
@ -1606,6 +1602,7 @@ void ToGLSL::TranslateAtomicMemOp(Instruction* psInst)
{
bstring glsl = *psContext->currentGLSLString;
int numParenthesis = 0;
uint32_t ui32DstDataTypeFlag = TO_FLAG_DESTINATION | TO_FLAG_NAME_ONLY;
uint32_t ui32DataTypeFlag = TO_FLAG_INTEGER;
const char* func = "";
Operand* dest = 0;
@ -1900,6 +1897,14 @@ void ToGLSL::TranslateAtomicMemOp(Instruction* psInst)
break;
}
}
else if (psBinding->eType == RTYPE_UAV_RWSTRUCTURED)
{
if (DeclareRWStructuredBufferTemplateTypeAsInteger(psContext, dest))
{
isUint = false;
ui32DstDataTypeFlag |= TO_FLAG_INTEGER;
}
}
}
if (isUint && HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
@ -1918,7 +1923,8 @@ void ToGLSL::TranslateAtomicMemOp(Instruction* psInst)
bcatcstr(glsl, func);
bcatcstr(glsl, "(");
TranslateOperand(dest, TO_FLAG_DESTINATION | TO_FLAG_NAME_ONLY);
TranslateOperand(dest, ui32DstDataTypeFlag);
if (texDim > 0)
{
bcatcstr(glsl, ", ");
@ -2037,11 +2043,7 @@ void ToGLSL::TranslateConditional(
else
bcatcstr(glsl, " != ");
if (isInt)
bcatcstr(glsl, "0)");
else
bcatcstr(glsl, "uint(0u))");
bcatcstr(glsl, isInt ? "0)" : "uint(0))"); // Old ES3.0 Adrenos treat 0u as const int.
if (psInst->eOpcode != OPCODE_IF)
{
@ -2435,9 +2437,9 @@ void ToGLSL::TranslateInstruction(Instruction* psInst, bool isEmbedded /* = fals
bcatcstr(glsl, "(");
TranslateOperand(&psInst->asOperands[boolOp], TO_FLAG_BOOL, destMask);
if (HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
bcatcstr(glsl, ") * 0xffffffffu) & ");
bcatcstr(glsl, ") * 0xFFFFFFFFu) & ");
else
bcatcstr(glsl, ") * 0xffffffff) & ");
bcatcstr(glsl, ") * -1) & "); // GLSL ES 2 spec: high precision ints are guaranteed to have a range of at least (-2^16, 2^16)
TranslateOperand(&psInst->asOperands[otherOp], TO_FLAG_UNSIGNED_INTEGER, destMask);
}
@ -3095,6 +3097,8 @@ void ToGLSL::TranslateInstruction(Instruction* psInst, bool isEmbedded /* = fals
psContext->AddIndentation();
bcatcstr(glsl, "//BFI\n");
#endif
if (psContext->psShader->eTargetLanguage == LANG_ES_300)
UseExtraFunctionDependency("int_bitfieldInsert");
psContext->AddIndentation();
AddAssignToDest(&psInst->asOperands[0], SVT_INT, numoverall_elements, &numParenthesis);
@ -3111,7 +3115,10 @@ void ToGLSL::TranslateInstruction(Instruction* psInst, bool isEmbedded /* = fals
continue;
k++;
bcatcstr(glsl, "bitfieldInsert(");
if (psContext->psShader->eTargetLanguage == LANG_ES_300)
bcatcstr(glsl, "int_bitfieldInsert(");
else
bcatcstr(glsl, "bitfieldInsert(");
for (j = 4; j >= 1; --j)
{
@ -3428,7 +3435,8 @@ void ToGLSL::TranslateInstruction(Instruction* psInst, bool isEmbedded /* = fals
name = bformat(HLSLCC_TEMP_PREFIX "i_while_true_%d", m_NumDeclaredWhileTrueLoops++);
// Workaround limitation with WebGL 1.0 GLSL, as we're expecting something to break the loop in any case
int hardcoded_iteration_limit = 0x7FFFFFFF;
// Fragment shaders on some devices don't like too large integer constants (Adreno 3xx, for example)
int hardcoded_iteration_limit = (psContext->psShader->eShaderType == PIXEL_SHADER) ? 0x7FFF : 0x7FFFFFFF;
bformata(glsl, "for(int %s = 0 ; %s < 0x%X ; %s++){\n", name->data, name->data, hardcoded_iteration_limit, name->data);
}
@ -3856,6 +3864,10 @@ void ToGLSL::TranslateInstruction(Instruction* psInst, bool isEmbedded /* = fals
case RETURN_TYPE_UINT:
srcDataType = SVT_UINT;
break;
case RETURN_TYPE_SNORM:
case RETURN_TYPE_UNORM:
srcDataType = SVT_FLOAT;
break;
default:
ASSERT(0);
// Suppress uninitialised variable warning
@ -4219,7 +4231,8 @@ void ToGLSL::TranslateInstruction(Instruction* psInst, bool isEmbedded /* = fals
psContext->AddIndentation();
bcatcstr(glsl, "//NOT\n");
#endif
if (!HaveNativeBitwiseOps(psContext->psShader->eTargetLanguage))
// Adreno 3xx fails on ~a with "Internal compiler error: unexpected operator", use op_not instead
if (!HaveNativeBitwiseOps(psContext->psShader->eTargetLanguage) || psContext->psShader->eTargetLanguage == LANG_ES_300)
{
UseExtraFunctionDependency("op_not");

65
src/toGLSLOperand.cpp
View File

@ -374,12 +374,7 @@ static void printImmediate32(HLSLCrossCompilerContext *psContext, uint32_t value
case SVT_INT:
case SVT_INT16:
case SVT_INT12:
// Adreno bug (happens only on android 4.* GLES3) casting unsigned representation of negative values to signed int
// results in undefined value/fails to link shader, need to print as signed decimal
if (value > 0x7fffffff && psContext->psShader->eTargetLanguage == LANG_ES_300)
bformata(glsl, "%i", (int32_t)value);
// Need special handling for anything >= uint 0x3fffffff
else if (value > 0x3ffffffe)
if (value > 0x3ffffffe)
{
if (HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
bformata(glsl, "int(0x%Xu)", value);
@ -439,10 +434,17 @@ void ToGLSL::DeclareDynamicIndexWrapper(const char* psName, SHADER_VARIABLE_CLAS
if (m_FunctionDefinitions.find(suffix) == m_FunctionDefinitions.end())
{
m_FunctionDefinitions.insert(std::make_pair(suffix, "#define UNITY_DYNAMIC_INDEX_ES2 0\n"));
m_FunctionDefinitionsOrder.push_back(suffix);
}
bcatcstr(glsl, "\n");
char name[256];
if ((eClass == SVC_MATRIX_COLUMNS || eClass == SVC_MATRIX_ROWS) && psContext->flags & HLSLCC_FLAG_TRANSLATE_MATRICES)
sprintf(name, HLSLCC_TRANSLATE_MATRIX_FORMAT_STRING "%s", ui32Rows, ui32Columns, psName);
else
memcpy(name, psName, strlen(psName) + 1);
if (eClass == SVC_STRUCT)
{
bformata(glsl, "%s_Type %s%s", psName, psName, suffix);
@ -474,9 +476,9 @@ void ToGLSL::DeclareDynamicIndexWrapper(const char* psName, SHADER_VARIABLE_CLAS
}
bformata(glsl, "(int i){\n");
bcatcstr(glsl, "#if UNITY_DYNAMIC_INDEX_ES2\n");
bformata(glsl, " return %s[i];\n", psName);
bformata(glsl, " return %s[i];\n", name);
bcatcstr(glsl, "#else\n");
bformata(glsl, "#define d_ar %s\n", psName);
bformata(glsl, "#define d_ar %s\n", name);
bformata(glsl, " if (i <= 0) return d_ar[0];");
// Let's draw a line somewhere with this workaround
@ -489,6 +491,7 @@ void ToGLSL::DeclareDynamicIndexWrapper(const char* psName, SHADER_VARIABLE_CLAS
bcatcstr(glsl, "#endif\n");
bformata(glsl, "}\n\n");
m_FunctionDefinitions.insert(std::make_pair(psName, ""));
m_FunctionDefinitionsOrder.push_back(psName);
}
void ToGLSL::TranslateVariableNameWithMask(bstring glsl, const Operand* psOperand, uint32_t ui32TOFlag, uint32_t* pui32IgnoreSwizzle, uint32_t ui32CompMask, int *piRebase)
@ -578,11 +581,17 @@ void ToGLSL::TranslateVariableNameWithMask(bstring glsl, const Operand* psOperan
{
if (CanDoDirectCast(psContext, eType, requestedType) || !HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
{
bformata(glsl, "%s(", GetConstructorForType(psContext, requestedType, requestedComponents, false));
numParenthesis++;
hasCtor = 1;
if (eType == SVT_BOOL)
{
needsBoolUpscale = 1;
// make sure to wrap the whole thing in parens so the upscale
// multiply only applies to the bool
bcatcstr(glsl, "(");
numParenthesis++;
}
bformata(glsl, "%s(", GetConstructorForType(psContext, requestedType, requestedComponents, false));
numParenthesis++;
}
else
{
@ -756,9 +765,7 @@ void ToGLSL::TranslateVariableNameWithMask(bstring glsl, const Operand* psOperan
case OPERAND_TYPE_OUTPUT_DEPTH:
if (psContext->psShader->eTargetLanguage == LANG_ES_100 && !psContext->EnableExtension("GL_EXT_frag_depth"))
{
bcatcstr(psContext->extensions, "#ifdef GL_EXT_frag_depth\n");
bcatcstr(psContext->extensions, "#define gl_FragDepth gl_FragDepthEXT\n");
bcatcstr(psContext->extensions, "#endif\n");
}
// fall through
case OPERAND_TYPE_OUTPUT_DEPTH_GREATER_EQUAL:
@ -770,6 +777,13 @@ void ToGLSL::TranslateVariableNameWithMask(bstring glsl, const Operand* psOperan
case OPERAND_TYPE_TEMP:
{
SHADER_VARIABLE_TYPE eTempType = psOperand->GetDataType(psContext);
if (psOperand->eSpecialName == NAME_UNDEFINED && psOperand->specialName.length())
{
bcatcstr(glsl, psOperand->specialName.c_str());
break;
}
bcatcstr(glsl, HLSLCC_TEMP_PREFIX);
ASSERT(psOperand->ui32RegisterNumber < 0x10000); // Sanity check after temp splitting.
switch (eTempType)
@ -1160,7 +1174,7 @@ void ToGLSL::TranslateVariableNameWithMask(bstring glsl, const Operand* psOperan
if (((psVarType->Class == SVC_MATRIX_COLUMNS) || (psVarType->Class == SVC_MATRIX_ROWS)) && (psVarType->Elements > 1) && ((psContext->flags & HLSLCC_FLAG_TRANSLATE_MATRICES) == 0))
{
// Special handling for old matrix arrays
bformata(glsl, "%%s / 4%s", squareBrackets[squareBracketType][0], fullIndexOss.str().c_str(), squareBrackets[squareBracketType][1]);
bformata(glsl, "%s%s / 4%s", squareBrackets[squareBracketType][0], fullIndexOss.str().c_str(), squareBrackets[squareBracketType][1]);
bformata(glsl, "%s%s %% 4%s", squareBrackets[squareBracketType][0], fullIndexOss.str().c_str(), squareBrackets[squareBracketType][1]);
}
else // This path is atm the default
@ -1258,7 +1272,7 @@ void ToGLSL::TranslateVariableNameWithMask(bstring glsl, const Operand* psOperan
}
case OPERAND_TYPE_IMMEDIATE_CONSTANT_BUFFER:
{
if (psContext->IsVulkan())
if (psContext->IsVulkan() || psContext->IsSwitch())
{
bformata(glsl, "ImmCB_%d", psContext->currentPhase);
TranslateOperandIndex(psOperand, 0);
@ -1352,7 +1366,20 @@ void ToGLSL::TranslateVariableNameWithMask(bstring glsl, const Operand* psOperan
}
case OPERAND_TYPE_INPUT_THREAD_ID_IN_GROUP_FLATTENED://SV_GroupIndex
{
bcatcstr(glsl, "gl_LocalInvocationIndex");
if (requestedComponents > 1 && !hasCtor)
{
bcatcstr(glsl, GetConstructorForType(psContext, eType, requestedComponents, false));
bcatcstr(glsl, "(");
numParenthesis++;
hasCtor = 1;
}
for (uint32_t i = 0; i < requestedComponents; i++)
{
bcatcstr(glsl, "gl_LocalInvocationIndex");
if (i < requestedComponents - 1)
bcatcstr(glsl, ", ");
}
*pui32IgnoreSwizzle = 1; // No swizzle meaningful for scalar.
break;
}
@ -1458,9 +1485,9 @@ void ToGLSL::TranslateVariableNameWithMask(bstring glsl, const Operand* psOperan
break;
case NAME_IS_FRONT_FACE:
if (HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
bcatcstr(glsl, "(gl_FrontFacing ? 0xffffffffu : uint(0))");
bcatcstr(glsl, "(gl_FrontFacing ? 0xffffffffu : uint(0))"); // Old ES3.0 Adrenos treat 0u as const int
else
bcatcstr(glsl, "(gl_FrontFacing ? int(1) : int(0))");
bcatcstr(glsl, "(gl_FrontFacing ? 1 : 0)");
*pui32IgnoreSwizzle = 1;
break;
case NAME_PRIMITIVE_ID:
@ -1546,10 +1573,12 @@ void ToGLSL::TranslateVariableNameWithMask(bstring glsl, const Operand* psOperan
if (HaveUnsignedTypes(psContext->psShader->eTargetLanguage))
bcatcstr(glsl, ") * int(0xffffffffu)");
else
bcatcstr(glsl, ") * int(0xffffffff)");
bcatcstr(glsl, ") * int(0xffff)"); // GLSL ES 2 spec: high precision ints are guaranteed to have a range of (-2^16, 2^16)
}
numParenthesis--;
bcatcstr(glsl, ")");
numParenthesis--;
}
while (numParenthesis != 0)

85
src/toMetal.cpp
View File

@ -254,6 +254,8 @@ bool ToMetal::Translate()
if (m_StructDefinitions[GetInputStructName()].m_Members.size() > 0)
{
m_StructDefinitions[GetInputStructName()].m_Dependencies.push_back(GetInputStructName());
if (psContext->psDependencies)
psContext->psDependencies->m_SharedDependencies.push_back(GetInputStructName());
// Hack, we're reusing Mtl_VertexOut as an hull shader input array, so no need to declare original contents
m_StructDefinitions[GetInputStructName()].m_Members.clear();
@ -266,10 +268,14 @@ bool ToMetal::Translate()
if (psContext->psDependencies)
{
for (auto itr = psContext->psDependencies->m_SharedFunctionMembers.begin(); itr != psContext->psDependencies->m_SharedFunctionMembers.end(); itr++)
for (auto i = psContext->psDependencies->m_SharedFunctionMembers.begin(), in = psContext->psDependencies->m_SharedFunctionMembers.end(); i != in;)
{
tessVertexFunctionArguments += itr->first.c_str();
tessVertexFunctionArguments += ", ";
tessVertexFunctionArguments += i->first.c_str();
++i;
// we want to avoid trailing comma
if (i != in)
tessVertexFunctionArguments += ", ";
}
}
}
@ -301,7 +307,22 @@ bool ToMetal::Translate()
{
if (psShader->eShaderType == HULL_SHADER)
{
m_StructDefinitions[""].m_Members.push_back(std::make_pair("vertexInput", "Mtl_VertexIn vertexInput [[ stage_in ]]"));
if (psContext->psDependencies)
{
// if we go for fully procedural geometry we might end up without Mtl_VertexIn
for (std::vector<std::string>::const_iterator itr = psContext->psDependencies->m_SharedDependencies.begin(); itr != psContext->psDependencies->m_SharedDependencies.end(); itr++)
{
if (*itr == "Mtl_VertexIn")
{
m_StructDefinitions[""].m_Members.push_back(std::make_pair("vertexInput", "Mtl_VertexIn vertexInput [[ stage_in ]]"));
if (tessVertexFunctionArguments.length())
tessVertexFunctionArguments += ", ";
tessVertexFunctionArguments += "vertexInput";
break;
}
}
}
m_StructDefinitions[""].m_Members.push_back(std::make_pair("tID", "uint2 tID [[ thread_position_in_grid ]]"));
m_StructDefinitions[""].m_Members.push_back(std::make_pair("groupID", "ushort2 groupID [[ threadgroup_position_in_grid ]]"));
@ -345,22 +366,33 @@ bool ToMetal::Translate()
m_StructDefinitions[""].m_Members.push_back(std::make_pair("input", GetInputStructName() + " input [[ stage_in ]]"));
}
if ((psShader->eShaderType == VERTEX_SHADER || psShader->eShaderType == HULL_SHADER) && (psContext->flags & HLSLCC_FLAG_METAL_TESSELLATION) != 0)
{
// m_StructDefinitions is inherited between tessellation shader stages but some builtins need exceptions
std::for_each(m_StructDefinitions[""].m_Members.begin(), m_StructDefinitions[""].m_Members.end(), [&psShader](MemberDefinitions::value_type &mem)
{
if (mem.first == "mtl_InstanceID")
{
if (psShader->eShaderType == VERTEX_SHADER)
mem.second.assign("uint mtl_InstanceID");
else if (psShader->eShaderType == HULL_SHADER)
mem.second.assign("// mtl_InstanceID passed through groupID");
}
});
}
m_StructDefinitions[""].m_Dependencies.push_back(GetInputStructName());
if (psContext->psDependencies)
psContext->psDependencies->m_SharedDependencies.push_back(GetInputStructName());
}
if ((psShader->eShaderType == VERTEX_SHADER || psShader->eShaderType == HULL_SHADER || psShader->eShaderType == DOMAIN_SHADER) && (psContext->flags & HLSLCC_FLAG_METAL_TESSELLATION) != 0)
{
// m_StructDefinitions is inherited between tessellation shader stages but some builtins need exceptions
std::for_each(m_StructDefinitions[""].m_Members.begin(), m_StructDefinitions[""].m_Members.end(), [&psShader](MemberDefinitions::value_type &mem)
{
if (mem.first == "mtl_InstanceID")
{
if (psShader->eShaderType == VERTEX_SHADER)
mem.second.assign("uint mtl_InstanceID");
else if (psShader->eShaderType == HULL_SHADER)
mem.second.assign("// mtl_InstanceID passed through groupID");
}
else if (mem.first == "mtl_VertexID")
{
if (psShader->eShaderType == VERTEX_SHADER)
mem.second.assign("uint mtl_VertexID");
else if (psShader->eShaderType == HULL_SHADER)
mem.second.assign("// mtl_VertexID generated in compute kernel");
else if (psShader->eShaderType == DOMAIN_SHADER)
mem.second.assign("// mtl_VertexID unused");
}
});
}
if (psShader->eShaderType != COMPUTE_SHADER)
@ -368,6 +400,8 @@ bool ToMetal::Translate()
if (m_StructDefinitions[GetOutputStructName()].m_Members.size() > 0)
{
m_StructDefinitions[""].m_Dependencies.push_back(GetOutputStructName());
if (psContext->psDependencies)
psContext->psDependencies->m_SharedDependencies.push_back(GetOutputStructName());
}
}
@ -395,7 +429,10 @@ bool ToMetal::Translate()
case PIXEL_SHADER:
if (psShader->sInfo.bEarlyFragmentTests)
bcatcstr(bodyglsl, "[[early_fragment_tests]]\n");
bcatcstr(bodyglsl, "fragment Mtl_FragmentOut xlatMtlMain(\n");
if (m_StructDefinitions[GetOutputStructName()].m_Members.size() > 0)
bcatcstr(bodyglsl, "fragment Mtl_FragmentOut xlatMtlMain(\n");
else
bcatcstr(bodyglsl, "fragment void xlatMtlMain(\n");
break;
case COMPUTE_SHADER:
bcatcstr(bodyglsl, "kernel void computeMain(\n");
@ -470,7 +507,7 @@ bool ToMetal::Translate()
psContext->AddIndentation();
bcatcstr(bodyglsl, "const uint controlPointID = (tID.x % patchInfo.numControlPointsPerPatch);\n");
psContext->AddIndentation();
bcatcstr(bodyglsl, "const uint internalControlPointID = (mtl_InstanceID * (patchInfo.numControlPointsPerPatch * patchInfo.numPatches)) + tID.x;\n");
bcatcstr(bodyglsl, "const uint mtl_VertexID = (mtl_InstanceID * (patchInfo.numControlPointsPerPatch * patchInfo.numPatches)) + tID.x;\n");
psContext->AddIndentation();
bformata(bodyglsl, "threadgroup %s inputGroup[numPatchesInThreadGroup];\n", GetInputStructName().c_str());
@ -491,7 +528,7 @@ bool ToMetal::Translate()
// Passthrough control point phase, run the rest only once per patch
psContext->AddIndentation();
bformata(bodyglsl, "input.cp[controlPointID] = vertexFunction(%svertexInput);\n", tessVertexFunctionArguments.c_str());
bformata(bodyglsl, "input.cp[controlPointID] = vertexFunction(%s);\n", tessVertexFunctionArguments.c_str());
DoHullShaderPassthrough(psContext);
@ -554,7 +591,7 @@ bool ToMetal::Translate()
psContext->indent++;
psContext->AddIndentation();
bformata(bodyglsl, "input.cp[controlPointID] = vertexFunction(%svertexInput);\n", tessVertexFunctionArguments.c_str());
bformata(bodyglsl, "input.cp[controlPointID] = vertexFunction(%s);\n", tessVertexFunctionArguments.c_str());
}
else
{
@ -613,7 +650,7 @@ bool ToMetal::Translate()
if (hasControlPoint)
{
psContext->AddIndentation();
bcatcstr(bodyglsl, "controlPoints[internalControlPointID] = output;\n");
bcatcstr(bodyglsl, "controlPoints[mtl_VertexID] = output;\n");
}
psContext->AddIndentation();

157
src/toMetalDeclaration.cpp
View File

@ -6,6 +6,8 @@
#include <sstream>
#include <cmath>
using namespace HLSLcc;
#ifndef fpcheck
#ifdef _MSC_VER
#define fpcheck(x) (_isnan(x) || !_finite(x))
@ -64,7 +66,7 @@ bool ToMetal::TranslateSystemValue(const Operand *psOperand, const ShaderInfo::I
{
case NAME_POSITION:
if (psContext->psShader->eShaderType == PIXEL_SHADER)
result = "mtl_FragCoord";
result = "hlslcc_FragCoord";
else
result = "mtl_Position";
if (outSkipPrefix != NULL) *outSkipPrefix = true;
@ -211,11 +213,21 @@ void ToMetal::DeclareBuiltinInput(const Declaration *psDecl)
{
const SPECIAL_NAME eSpecialName = psDecl->asOperands[0].eSpecialName;
Shader* psShader = psContext->psShader;
const Operand* psOperand = &psDecl->asOperands[0];
const int regSpace = psOperand->GetRegisterSpace(psContext);
ASSERT(regSpace == 0);
// we need to at least mark if they are scalars or not (as we might need to use vector ctor)
if (psOperand->GetNumInputElements(psContext) == 1)
psShader->abScalarInput[regSpace][psOperand->ui32RegisterNumber] |= (int)psOperand->ui32CompMask;
switch (eSpecialName)
{
case NAME_POSITION:
ASSERT(psContext->psShader->eShaderType == PIXEL_SHADER);
m_StructDefinitions[""].m_Members.push_back(std::make_pair("mtl_FragCoord", "float4 mtl_FragCoord [[ position ]]"));
bcatcstr(GetEarlyMain(psContext), "float4 hlslcc_FragCoord = float4(mtl_FragCoord.xyz, 1.0/mtl_FragCoord.w);\n");
break;
case NAME_RENDER_TARGET_ARRAY_INDEX:
// Only supported on a Mac
@ -277,7 +289,6 @@ void ToMetal::DeclareClipPlanes(const Declaration* decl, unsigned declCount)
else if (psFirstClipSignature->ui32Mask & (1 << 1)) compCount = 2;
}
ShaderPhase* phase = &shader->asPhases[psContext->currentPhase];
for (unsigned i = 0, n = declCount; i < n; ++i)
{
const Operand* operand = &decl[i].asOperands[0];
@ -287,22 +298,19 @@ void ToMetal::DeclareClipPlanes(const Declaration* decl, unsigned declCount)
shader->sInfo.GetOutputSignatureFromRegister(operand->ui32RegisterNumber, operand->ui32CompMask, 0, &signature);
const int semanticIndex = signature->ui32SemanticIndex;
bformata(phase->earlyMain, " float4 phase%d_ClipDistance%d;\n", psContext->currentPhase, signature->ui32SemanticIndex);
bformata(GetEarlyMain(psContext), "float4 phase%d_ClipDistance%d;\n", psContext->currentPhase, signature->ui32SemanticIndex);
const char* swizzleStr[] = { "x", "y", "z", "w" };
phase->hasPostShaderCode = 1;
if (planeCount > 1)
{
for (int i = 0; i < compCount; ++i)
{
bformata(phase->postShaderCode, " %s.mtl_ClipDistance[%d] = phase%d_ClipDistance%d.%s;\n",
"output", semanticIndex * compCount + i, psContext->currentPhase, semanticIndex, swizzleStr[i]
);
bformata(GetPostShaderCode(psContext), "%s.mtl_ClipDistance[%d] = phase%d_ClipDistance%d.%s;\n", "output", semanticIndex * compCount + i, psContext->currentPhase, semanticIndex, swizzleStr[i]);
}
}
else
{
bformata(phase->postShaderCode, " %s.mtl_ClipDistance = phase%d_ClipDistance%d.x;\n", "output", psContext->currentPhase, semanticIndex);
bformata(GetPostShaderCode(psContext), "%s.mtl_ClipDistance = phase%d_ClipDistance%d.x;\n", "output", psContext->currentPhase, semanticIndex);
}
}
}
@ -311,10 +319,11 @@ void ToMetal::GenerateTexturesReflection(HLSLccReflection* refl)
{
for (unsigned i = 0, n = m_Textures.size(); i < n; ++i)
{
const std::string samplerName1 = m_Textures[i].name, samplerName2 = "sampler" + m_Textures[i].name;
// Match CheckSamplerAndTextureNameMatch behavior
const std::string samplerName1 = m_Textures[i].name, samplerName2 = "sampler" + m_Textures[i].name, samplerName3 = "sampler_" + m_Textures[i].name;
for (unsigned j = 0, m = m_Samplers.size(); j < m; ++j)
{
if (m_Samplers[j].name == samplerName1 || m_Samplers[j].name == samplerName2)
if (m_Samplers[j].name == samplerName1 || m_Samplers[j].name == samplerName2 || m_Samplers[j].name == samplerName3)
{
m_Textures[i].samplerBind = m_Samplers[j].slot;
break;
@ -513,11 +522,7 @@ void ToMetal::HandleOutputRedirect(const Declaration *psDecl, const std::string
ASSERT(psContext->psShader->aIndexedOutput[regSpace][psOperand->ui32RegisterNumber] == 0);
psContext->AddIndentation();
bformata(psPhase->earlyMain, "%s phase%d_Output%d_%d;\n", typeName.c_str(), psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
psPhase->hasPostShaderCode = 1;
psContext->currentGLSLString = &psPhase->postShaderCode;
bformata(GetEarlyMain(psContext), "%s phase%d_Output%d_%d;\n", typeName.c_str(), psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
while (comp < 4)
{
@ -541,39 +546,36 @@ void ToMetal::HandleOutputRedirect(const Declaration *psDecl, const std::string
mask = psSig->ui32Mask;
((Operand *)psOperand)->ui32CompMask = 1 << comp;
psContext->AddIndentation();
bcatcstr(psPhase->postShaderCode, TranslateOperand(psOperand, TO_FLAG_NAME_ONLY).c_str());
bcatcstr(psPhase->postShaderCode, " = ");
bstring str = GetPostShaderCode(psContext);
bcatcstr(str, TranslateOperand(psOperand, TO_FLAG_NAME_ONLY).c_str());
bcatcstr(str, " = ");
if (psSig->eComponentType == INOUT_COMPONENT_SINT32)
{
bformata(psPhase->postShaderCode, "as_type<int>(");
bformata(str, "as_type<int>(");
hasCast = 1;
}
else if (psSig->eComponentType == INOUT_COMPONENT_UINT32)
{
bformata(psPhase->postShaderCode, "as_type<uint>(");
bformata(str, "as_type<uint>(");
hasCast = 1;
}
bformata(psPhase->postShaderCode, "phase%d_Output%d_%d.", psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
bformata(str, "phase%d_Output%d_%d.", psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
// Print out mask
for (i = 0; i < 4; i++)
{
if ((mask & (1 << i)) == 0)
continue;
bformata(psPhase->postShaderCode, "%c", "xyzw"[i]);
bformata(str, "%c", "xyzw"[i]);
}
if (hasCast)
bcatcstr(psPhase->postShaderCode, ")");
bcatcstr(str, ")");
comp += numComps;
bcatcstr(psPhase->postShaderCode, ";\n");
bcatcstr(str, ";\n");
}
psContext->currentGLSLString = &psContext->glsl;
((Operand *)psOperand)->ui32CompMask = origMask;
if (regSpace == 0)
psShader->asPhases[psContext->currentPhase].acOutputNeedsRedirect[psOperand->ui32RegisterNumber] = 0xfe;
@ -612,33 +614,31 @@ void ToMetal::HandleInputRedirect(const Declaration *psDecl, const std::string &
ASSERT(psContext->psShader->aIndexedInput[regSpace][psOperand->ui32RegisterNumber] == 0);
psContext->currentGLSLString = &psPhase->earlyMain;
psContext->AddIndentation();
++psContext->indent;
bcatcstr(psPhase->earlyMain, " ");
// Does the input have multiple array components (such as geometry shader input, or domain shader control point input)
if ((psShader->eShaderType == DOMAIN_SHADER && regSpace == 0) || (psShader->eShaderType == GEOMETRY_SHADER))
{
// The count is actually stored in psOperand->aui32ArraySizes[0]
origArraySize = psOperand->aui32ArraySizes[0];
// bformata(glsl, "%s vec4 phase%d_Input%d_%d[%d];\n", Precision, psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber, origArraySize);
bformata(psPhase->earlyMain, "%s phase%d_Input%d_%d[%d];\n", typeName.c_str(), psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber, origArraySize);
bformata(GetEarlyMain(psContext), "%s phase%d_Input%d_%d[%d];\n", typeName.c_str(), psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber, origArraySize);
needsLooping = 1;
i = origArraySize - 1;
}
else
// bformata(glsl, "%s vec4 phase%d_Input%d_%d;\n", Precision, psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
bformata(psPhase->earlyMain, "%s phase%d_Input%d_%d;\n", typeName.c_str(), psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
bformata(GetEarlyMain(psContext), "%s phase%d_Input%d_%d;\n", typeName.c_str(), psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber);
// Do a conditional loop. In normal cases needsLooping == 0 so this is only run once.
do
{
int comp = 0;
bcatcstr(psPhase->earlyMain, " ");
bstring str = GetEarlyMain(psContext);
if (needsLooping)
bformata(psPhase->earlyMain, "phase%d_Input%d_%d[%d] = %s(", psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber, i, typeName.c_str());
bformata(str, "phase%d_Input%d_%d[%d] = %s(", psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber, i, typeName.c_str());
else
bformata(psPhase->earlyMain, "phase%d_Input%d_%d = %s(", psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber, typeName.c_str());
bformata(str, "phase%d_Input%d_%d = %s(", psContext->currentPhase, regSpace, psOperand->ui32RegisterNumber, typeName.c_str());
while (comp < 4)
{
@ -656,9 +656,9 @@ void ToMetal::HandleInputRedirect(const Declaration *psDecl, const std::string &
if (psSig->eComponentType != INOUT_COMPONENT_FLOAT32)
{
if (numComps > 1)
bformata(psPhase->earlyMain, "as_type<float%d>(", numComps);
bformata(str, "as_type<float%d>(", numComps);
else
bformata(psPhase->earlyMain, "as_type<float>(");
bformata(str, "as_type<float>(");
hasCast = 1;
}
@ -669,7 +669,7 @@ void ToMetal::HandleInputRedirect(const Declaration *psDecl, const std::string &
// And the component mask
psOperand->ui32CompMask = 1 << comp;
bformata(psPhase->earlyMain, TranslateOperand(psOperand, TO_FLAG_NAME_ONLY).c_str());
bformata(str, TranslateOperand(psOperand, TO_FLAG_NAME_ONLY).c_str());
// Restore the original array size value and mask
psOperand->ui32CompMask = origMask;
@ -677,23 +677,23 @@ void ToMetal::HandleInputRedirect(const Declaration *psDecl, const std::string &
psOperand->aui32ArraySizes[0] = origArraySize;
if (hasCast)
bcatcstr(psPhase->earlyMain, ")");
bcatcstr(str, ")");
comp += numComps;
}
else // no signature found -> fill with zero
{
bcatcstr(psPhase->earlyMain, "0");
bcatcstr(str, "0");
comp++;
}
if (comp < 4)
bcatcstr(psPhase->earlyMain, ", ");
bcatcstr(str, ", ");
}
bcatcstr(psPhase->earlyMain, ");\n");
bcatcstr(str, ");\n");
}
while ((--i) >= 0);
psContext->currentGLSLString = &psContext->glsl;
--psContext->indent;
if (regSpace == 0)
psShader->asPhases[psContext->currentPhase].acInputNeedsRedirect[psOperand->ui32RegisterNumber] = 0xfe;
@ -894,17 +894,20 @@ static std::string GetInterpolationString(INTERPOLATION_MODE eMode)
}
}
void ToMetal::DeclareStructVariable(const std::string &parentName, const ShaderVar &var, bool withinCB, uint32_t cumulativeOffset)
void ToMetal::DeclareStructVariable(const std::string &parentName, const ShaderVar &var, bool withinCB, uint32_t cumulativeOffset, bool isUsed)
{
DeclareStructVariable(parentName, var.sType, withinCB, cumulativeOffset + var.ui32StartOffset);
DeclareStructVariable(parentName, var.sType, withinCB, cumulativeOffset + var.ui32StartOffset, isUsed);
}
void ToMetal::DeclareStructVariable(const std::string &parentName, const ShaderVarType &var, bool withinCB, uint32_t cumulativeOffset)
void ToMetal::DeclareStructVariable(const std::string &parentName, const ShaderVarType &var, bool withinCB, uint32_t cumulativeOffset, bool isUsed)
{
// CB arrays need to be defined as 4 component vectors to match DX11 data layout
bool arrayWithinCB = (withinCB && (var.Elements > 1) && (psContext->psShader->eShaderType == COMPUTE_SHADER));
bool doDeclare = true;
if (isUsed == false && ((psContext->flags & HLSLCC_FLAG_REMOVE_UNUSED_GLOBALS)) == 0)
isUsed = true;
if (var.Class == SVC_STRUCT)
{
if (m_StructDefinitions.find(var.name + "_Type") == m_StructDefinitions.end())
@ -914,7 +917,7 @@ void ToMetal::DeclareStructVariable(const std::string &parentName, const ShaderV
if (var.Parent == NULL && var.Elements > 1 && withinCB)
{
// var.Type being SVT_VOID indicates it is a struct in this case.
psContext->m_Reflection.OnConstant(var.fullName, var.Offset + cumulativeOffset, var.Type, var.Rows, var.Columns, false, var.Elements);
psContext->m_Reflection.OnConstant(var.fullName, var.Offset + cumulativeOffset, var.Type, var.Rows, var.Columns, false, var.Elements, true);
}
std::ostringstream oss;
@ -950,9 +953,9 @@ void ToMetal::DeclareStructVariable(const std::string &parentName, const ShaderV
// On non-compute we can fake that we still have a matrix, as CB upload code will fill the data correctly on 4x4 matrices.
// That way we avoid the issues with mismatching types for builtins etc.
if (psContext->psShader->eShaderType == COMPUTE_SHADER)
doDeclare = psContext->m_Reflection.OnConstant(var.fullName, var.Offset + cumulativeOffset, var.Type, 1, 4, false, elemCount);
doDeclare = psContext->m_Reflection.OnConstant(var.fullName, var.Offset + cumulativeOffset, var.Type, 1, 4, false, elemCount, isUsed);
else
doDeclare = psContext->m_Reflection.OnConstant(var.fullName, var.Offset + cumulativeOffset, var.Type, var.Rows, var.Columns, true, var.Elements);
doDeclare = psContext->m_Reflection.OnConstant(var.fullName, var.Offset + cumulativeOffset, var.Type, var.Rows, var.Columns, true, var.Elements, isUsed);
}
}
else
@ -966,7 +969,7 @@ void ToMetal::DeclareStructVariable(const std::string &parentName, const ShaderV
// TODO Verify whether the offset is from the beginning of the CB or from the beginning of the struct
if (withinCB)
doDeclare = psContext->m_Reflection.OnConstant(var.fullName, var.Offset + cumulativeOffset, var.Type, var.Rows, var.Columns, true, var.Elements);
doDeclare = psContext->m_Reflection.OnConstant(var.fullName, var.Offset + cumulativeOffset, var.Type, var.Rows, var.Columns, true, var.Elements, isUsed);
}
if (doDeclare)
@ -983,7 +986,7 @@ void ToMetal::DeclareStructVariable(const std::string &parentName, const ShaderV
}
if (withinCB)
doDeclare = psContext->m_Reflection.OnConstant(var.fullName, var.Offset + cumulativeOffset, var.Type, 1, var.Columns, false, var.Elements);
doDeclare = psContext->m_Reflection.OnConstant(var.fullName, var.Offset + cumulativeOffset, var.Type, 1, var.Columns, false, var.Elements, isUsed);
if (doDeclare)
m_StructDefinitions[parentName].m_Members.push_back(std::make_pair(var.name, oss.str()));
@ -1008,7 +1011,7 @@ void ToMetal::DeclareStructVariable(const std::string &parentName, const ShaderV
}
if (withinCB)
doDeclare = psContext->m_Reflection.OnConstant(var.fullName, var.Offset + cumulativeOffset, var.Type, 1, 1, false, var.Elements);
doDeclare = psContext->m_Reflection.OnConstant(var.fullName, var.Offset + cumulativeOffset, var.Type, 1, 1, false, var.Elements, isUsed);
if (doDeclare)
m_StructDefinitions[parentName].m_Members.push_back(std::make_pair(var.name, oss.str()));
@ -1026,7 +1029,7 @@ void ToMetal::DeclareStructType(const std::string &name, const std::vector<Shade
if (stripUnused && !itr->sType.m_IsUsed)
continue;
DeclareStructVariable(name, *itr, withinCB, cumulativeOffset);
DeclareStructVariable(name, *itr, withinCB, cumulativeOffset, itr->sType.m_IsUsed);
}
}
@ -1118,23 +1121,18 @@ void ToMetal::DeclareBufferVariable(const Declaration *psDecl, bool isRaw, bool
// In addition to the actual declaration, we need pointer modification and possible counter declaration
// in early main:
std::ostringstream earlymainoss;
// Possible counter is always in the beginning of the buffer
if (isUAV && psDecl->sUAV.bCounter)
{
earlymainoss << " device atomic_uint *" << BufName << "_counter = reinterpret_cast<device atomic_uint *> (" << BufName << ");\n";
bformata(GetEarlyMain(psContext), "device atomic_uint *%s_counter = reinterpret_cast<device atomic_uint *> (%s);\n", BufName.c_str(), BufName.c_str());
}
// Some GPUs don't allow memory access below buffer binding offset in the shader so always bind compute buffer
// at offset 0 instead of GetDataOffset().
// We can't tell at shader compile time if the buffer actually has counter or not. Therefore we'll always reserve
// space for the counter and bump the data pointer to beginning of the actual data here.
earlymainoss << " " << BufName << " = reinterpret_cast<" << BufConst
<< "device " << (isRaw ? "uint" : BufType) << " *> (reinterpret_cast<device "
<< BufConst << "atomic_uint *> (" << BufName << ") + 1);\n";
bformata(psContext->psShader->asPhases[psContext->currentPhase].earlyMain, earlymainoss.str().c_str());
bformata(GetEarlyMain(psContext), "%s = reinterpret_cast<%sdevice %s *> (reinterpret_cast<device %satomic_uint *> (%s) + 1);\n", BufName.c_str(), BufConst.c_str(), (isRaw ? "uint" : BufType.c_str()), BufConst.c_str(), BufName.c_str());
}
static int ParseInlineSamplerWrapMode(const std::string& samplerName, const std::string& wrapName)
@ -1155,7 +1153,7 @@ static int ParseInlineSamplerWrapMode(const std::string& samplerName, const std:
return res;
}
static bool EmitInlineSampler(HLSLCrossCompilerContext* ctx, const std::string& name)
static bool EmitInlineSampler(HLSLCrossCompilerContext* psContext, const std::string& name)
{
// See if it's a sampler that goes with the texture, or an "inline" sampler
// where sampler states are hardcoded in the shader directly.
@ -1189,8 +1187,8 @@ static bool EmitInlineSampler(HLSLCrossCompilerContext* ctx, const std::string&
return false;
}
bstring str = ctx->psShader->asPhases[ctx->currentPhase].earlyMain;
bformata(str, "\tconstexpr sampler %s(", name.c_str());
bstring str = GetEarlyMain(psContext);
bformata(str, "constexpr sampler %s(", name.c_str());
if (hasCompare)
bformata(str, "compare_func::greater_equal,");
@ -1365,6 +1363,7 @@ void ToMetal::TranslateDeclaration(const Declaration* psDecl)
if (psDecl->eOpcode == OPCODE_DCL_INPUT_PS_SIV && psOperand->eSpecialName == NAME_POSITION)
{
m_StructDefinitions[""].m_Members.push_back(std::make_pair("mtl_FragCoord", "float4 mtl_FragCoord [[ position ]]"));
bcatcstr(GetEarlyMain(psContext), "float4 hlslcc_FragCoord = float4(mtl_FragCoord.xyz, 1.0/mtl_FragCoord.w);\n");
break;
}
@ -1477,29 +1476,28 @@ void ToMetal::TranslateDeclaration(const Declaration* psDecl)
{
uint32_t i = 0;
const uint32_t ui32NumTemps = psDecl->value.ui32NumTemps;
glsl = psContext->psShader->asPhases[psContext->currentPhase].earlyMain;
for (i = 0; i < ui32NumTemps; i++)
{
if (psShader->psFloatTempSizes[i] != 0)
bformata(glsl, " %s " HLSLCC_TEMP_PREFIX "%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_FLOAT, psShader->psFloatTempSizes[i]), i);
bformata(GetEarlyMain(psContext), "%s " HLSLCC_TEMP_PREFIX "%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_FLOAT, psShader->psFloatTempSizes[i]), i);
if (psShader->psFloat16TempSizes[i] != 0)
bformata(glsl, " %s " HLSLCC_TEMP_PREFIX "16_%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_FLOAT16, psShader->psFloat16TempSizes[i]), i);
bformata(GetEarlyMain(psContext), "%s " HLSLCC_TEMP_PREFIX "16_%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_FLOAT16, psShader->psFloat16TempSizes[i]), i);
if (psShader->psFloat10TempSizes[i] != 0)
bformata(glsl, " %s " HLSLCC_TEMP_PREFIX "10_%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_FLOAT10, psShader->psFloat10TempSizes[i]), i);
bformata(GetEarlyMain(psContext), "%s " HLSLCC_TEMP_PREFIX "10_%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_FLOAT10, psShader->psFloat10TempSizes[i]), i);
if (psShader->psIntTempSizes[i] != 0)
bformata(glsl, " %s " HLSLCC_TEMP_PREFIX "i%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_INT, psShader->psIntTempSizes[i]), i);
bformata(GetEarlyMain(psContext), "%s " HLSLCC_TEMP_PREFIX "i%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_INT, psShader->psIntTempSizes[i]), i);
if (psShader->psInt16TempSizes[i] != 0)
bformata(glsl, " %s " HLSLCC_TEMP_PREFIX "i16_%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_INT16, psShader->psInt16TempSizes[i]), i);
bformata(GetEarlyMain(psContext), "%s " HLSLCC_TEMP_PREFIX "i16_%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_INT16, psShader->psInt16TempSizes[i]), i);
if (psShader->psInt12TempSizes[i] != 0)
bformata(glsl, " %s " HLSLCC_TEMP_PREFIX "i12_%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_INT12, psShader->psInt12TempSizes[i]), i);
bformata(GetEarlyMain(psContext), "%s " HLSLCC_TEMP_PREFIX "i12_%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_INT12, psShader->psInt12TempSizes[i]), i);
if (psShader->psUIntTempSizes[i] != 0)
bformata(glsl, " %s " HLSLCC_TEMP_PREFIX "u%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_UINT, psShader->psUIntTempSizes[i]), i);
bformata(GetEarlyMain(psContext), "%s " HLSLCC_TEMP_PREFIX "u%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_UINT, psShader->psUIntTempSizes[i]), i);
if (psShader->psUInt16TempSizes[i] != 0)
bformata(glsl, " %s " HLSLCC_TEMP_PREFIX "u16_%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_UINT16, psShader->psUInt16TempSizes[i]), i);
bformata(GetEarlyMain(psContext), "%s " HLSLCC_TEMP_PREFIX "u16_%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_UINT16, psShader->psUInt16TempSizes[i]), i);
if (psShader->fp64 && (psShader->psDoubleTempSizes[i] != 0))
bformata(glsl, " %s " HLSLCC_TEMP_PREFIX "d%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_DOUBLE, psShader->psDoubleTempSizes[i]), i);
bformata(GetEarlyMain(psContext), "%s " HLSLCC_TEMP_PREFIX "d%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_DOUBLE, psShader->psDoubleTempSizes[i]), i);
if (psShader->psBoolTempSizes[i] != 0)
bformata(glsl, " %s " HLSLCC_TEMP_PREFIX "b%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_BOOL, psShader->psBoolTempSizes[i]), i);
bformata(GetEarlyMain(psContext), "%s " HLSLCC_TEMP_PREFIX "b%d;\n", HLSLcc::GetConstructorForType(psContext, SVT_BOOL, psShader->psBoolTempSizes[i]), i);
}
break;
}
@ -1701,7 +1699,7 @@ void ToMetal::TranslateDeclaration(const Declaration* psDecl)
const uint32_t ui32RegIndex = psDecl->sIdxTemp.ui32RegIndex;
const uint32_t ui32RegCount = psDecl->sIdxTemp.ui32RegCount;
const uint32_t ui32RegComponentSize = psDecl->sIdxTemp.ui32RegComponentSize;
bformata(psContext->psShader->asPhases[psContext->currentPhase].earlyMain, "float%d TempArray%d[%d];\n", ui32RegComponentSize, ui32RegIndex, ui32RegCount);
bformata(GetEarlyMain(psContext), "float%d TempArray%d[%d];\n", ui32RegComponentSize, ui32RegIndex, ui32RegCount);
break;
}
case OPCODE_DCL_INDEX_RANGE:
@ -1782,9 +1780,7 @@ void ToMetal::TranslateDeclaration(const Declaration* psDecl)
oldString = psContext->currentGLSLString;
psContext->currentGLSLString = &psContext->psShader->asPhases[psContext->currentPhase].earlyMain;
psContext->AddIndentation();
psContext->currentGLSLString = oldString;
bformata(psContext->psShader->asPhases[psContext->currentPhase].earlyMain, "%s4 phase%d_%sput%d_%d[%d];\n", type, psContext->currentPhase, isInput ? "In" : "Out", regSpace, startReg, psDecl->value.ui32IndexRange);
oldString = psContext->currentGLSLString;
glsl = isInput ? psContext->psShader->asPhases[psContext->currentPhase].earlyMain : psContext->psShader->asPhases[psContext->currentPhase].postShaderCode;
psContext->currentGLSLString = &glsl;
if (isInput == 0)
@ -1821,7 +1817,6 @@ void ToMetal::TranslateDeclaration(const Declaration* psDecl)
realName = psContext->GetDeclaredInputName(&psDecl->asOperands[0], &dummy, 1, NULL);
psContext->AddIndentation();
bformata(glsl, "phase%d_Input%d_%d[%d]", psContext->currentPhase, regSpace, startReg, i);
if (destMask != OPERAND_4_COMPONENT_MASK_ALL)
@ -2079,8 +2074,7 @@ void ToMetal::TranslateDeclaration(const Declaration* psDecl)
<< "_Type " << TranslateOperand(&psDecl->asOperands[0], TO_FLAG_NAME_ONLY)
<< "[" << psDecl->sTGSM.ui32Count << "]";
bformata(psContext->psShader->asPhases[psContext->currentPhase].earlyMain, "\t%s;\n", oss.str().c_str());
bformata(GetEarlyMain(psContext), "%s;\n", oss.str().c_str());
psVarType->name = "$Element";
psVarType->Columns = psDecl->sTGSM.ui32Stride / 4;
@ -2095,8 +2089,7 @@ void ToMetal::TranslateDeclaration(const Declaration* psDecl)
oss << "threadgroup uint " << TranslateOperand(&psDecl->asOperands[0], TO_FLAG_NAME_ONLY)
<< "[" << (psDecl->sTGSM.ui32Count / psDecl->sTGSM.ui32Stride) << "]";
bformata(psContext->psShader->asPhases[psContext->currentPhase].earlyMain, "\t%s;\n", oss.str().c_str());
bformata(GetEarlyMain(psContext), "%s;\n", oss.str().c_str());
psVarType->name = "$Element";
psVarType->Columns = 1;

70
src/toMetalInstruction.cpp
View File

@ -244,7 +244,7 @@ void ToMetal::AddComparison(Instruction* psInst, ComparisonType eType,
glsl << TranslateOperand(&psInst->asOperands[2], typeFlag, destMask);
if (!isBoolDest)
{
bcatcstr(glsl, ") ? 0xFFFFFFFFu : 0u");
bcatcstr(glsl, ") ? 0xFFFFFFFFu : uint(0)");
}
AddAssignPrologue(needsParenthesis);
}
@ -334,7 +334,7 @@ void ToMetal::AddMOVCBinaryOp(const Operand *pDest, const Operand *src0, Operand
else
{
if (s0Type == SVT_UINT || s0Type == SVT_UINT16)
bcatcstr(glsl, " != 0u) ? ");
bcatcstr(glsl, " != uint(0)) ? ");
else if (s0Type == SVT_BOOL)
bcatcstr(glsl, ") ? ");
else
@ -358,13 +358,18 @@ void ToMetal::AddMOVCBinaryOp(const Operand *pDest, const Operand *src0, Operand
{
// TODO: We can actually do this in one op using mix().
int srcElem = -1;
SHADER_VARIABLE_TYPE dstType = pDest->GetDataType(psContext);
SHADER_VARIABLE_TYPE s0Type = src0->GetDataType(psContext);
// Use an extra temp if dest is also one of the sources. Without this some swizzle combinations
// might alter the source before all components are handled.
const char* tempName = "hlslcc_movcTemp";
bool dstIsSrc1 = (pDest->eType == src1->eType) && (pDest->ui32RegisterNumber == src1->ui32RegisterNumber);
bool dstIsSrc2 = (pDest->eType == src2->eType) && (pDest->ui32RegisterNumber == src2->ui32RegisterNumber);
const std::string tempName = "hlslcc_movcTemp";
bool dstIsSrc1 = (pDest->eType == src1->eType)
&& (dstType == src1->GetDataType(psContext))
&& (pDest->ui32RegisterNumber == src1->ui32RegisterNumber);
bool dstIsSrc2 = (pDest->eType == src2->eType)
&& (dstType == src2->GetDataType(psContext))
&& (pDest->ui32RegisterNumber == src2->ui32RegisterNumber);
if (dstIsSrc1 || dstIsSrc2)
{
@ -375,7 +380,10 @@ void ToMetal::AddMOVCBinaryOp(const Operand *pDest, const Operand *src0, Operand
int numComponents = (pDest->eType == OPERAND_TYPE_TEMP) ?
psContext->psShader->GetTempComponentCount(eDestType, pDest->ui32RegisterNumber) :
pDest->iNumComponents;
bformata(glsl, "%s %s = %s;\n", HLSLcc::GetConstructorForType(psContext, eDestType, numComponents), tempName, TranslateOperand(pDest, TO_FLAG_NAME_ONLY).c_str());
bformata(glsl, "%s %s = %s;\n", HLSLcc::GetConstructorForType(psContext, eDestType, numComponents), tempName.c_str(), TranslateOperand(pDest, TO_FLAG_NAME_ONLY).c_str());
// Override OPERAND_TYPE_TEMP name temporarily
const_cast<Operand *>(pDest)->specialName.assign(tempName);
}
for (destElem = 0; destElem < 4; ++destElem)
@ -408,23 +416,20 @@ void ToMetal::AddMOVCBinaryOp(const Operand *pDest, const Operand *src0, Operand
}
}
if (!dstIsSrc1)
glsl << TranslateOperand(src1, SVTTypeToFlag(eDestType), 1 << srcElem);
else
bformata(glsl, "%s%s", tempName, TranslateOperandSwizzle(src1, 1 << srcElem, 0).c_str());
glsl << TranslateOperand(src1, SVTTypeToFlag(eDestType), 1 << srcElem);
bcatcstr(glsl, " : ");
if (!dstIsSrc2)
glsl << TranslateOperand(src2, SVTTypeToFlag(eDestType), 1 << srcElem);
else
bformata(glsl, "%s%s", tempName, TranslateOperandSwizzle(src2, 1 << srcElem, 0).c_str());
glsl << TranslateOperand(src2, SVTTypeToFlag(eDestType), 1 << srcElem);
AddAssignPrologue(numParenthesis);
}
if (dstIsSrc1 || dstIsSrc2)
{
const_cast<Operand *>(pDest)->specialName.clear();
psContext->AddIndentation();
glsl << TranslateOperand(pDest, TO_FLAG_NAME_ONLY);
bformata(glsl, " = %s;\n", tempName.c_str());
--psContext->indent;
psContext->AddIndentation();
bcatcstr(glsl, "}\n");
@ -967,7 +972,7 @@ void ToMetal::GetResInfoData(Instruction* psInst, int index, int destElem)
int dim = GetNumTextureDimensions(psInst->eResDim);
if (dim < (index + 1) && index != 3)
{
bcatcstr(glsl, eResInfoReturnType == RESINFO_INSTRUCTION_RETURN_UINT ? "0u" : "0.0");
bcatcstr(glsl, eResInfoReturnType == RESINFO_INSTRUCTION_RETURN_UINT ? "uint(0)" : "0.0");
}
else
{
@ -1892,7 +1897,7 @@ void ToMetal::TranslateConditional(
}
else if (psInst->eOpcode == OPCODE_RETC) // FIXME! Need to spew out shader epilogue
{
if (psContext->psShader->eShaderType == COMPUTE_SHADER)
if (psContext->psShader->eShaderType == COMPUTE_SHADER || (psContext->psShader->eShaderType == PIXEL_SHADER && m_StructDefinitions[GetOutputStructName()].m_Members.size() == 0))
statement = "return";
else
statement = "return output";
@ -1925,11 +1930,11 @@ void ToMetal::TranslateConditional(
if (psInst->eOpcode != OPCODE_IF)
{
bformata(glsl, ")==uint(0u)){%s;}\n", statement);
bformata(glsl, ")==uint(0)){%s;}\n", statement);
}
else
{
bcatcstr(glsl, ")==uint(0u)){\n");
bcatcstr(glsl, ")==uint(0)){\n");
}
}
else
@ -1940,11 +1945,11 @@ void ToMetal::TranslateConditional(
if (psInst->eOpcode != OPCODE_IF)
{
bformata(glsl, ")!=uint(0u)){%s;}\n", statement);
bformata(glsl, ")!=uint(0)){%s;}\n", statement);
}
else
{
bcatcstr(glsl, ")!=uint(0u)){\n");
bcatcstr(glsl, ")!=uint(0)){\n");
}
}
}
@ -2774,7 +2779,7 @@ void ToMetal::TranslateInstruction(Instruction* psInst)
#endif
}
psContext->AddIndentation();
if (psContext->psShader->eShaderType == COMPUTE_SHADER)
if (psContext->psShader->eShaderType == COMPUTE_SHADER || (psContext->psShader->eShaderType == PIXEL_SHADER && m_StructDefinitions[GetOutputStructName()].m_Members.size() == 0))
bcatcstr(glsl, "return;\n");
else
bcatcstr(glsl, "return output;\n");
@ -3589,15 +3594,26 @@ template <int N> vec<int, N> bitFieldExtractI(const vec<uint, N> width, const ve
psContext->AddIndentation();
uint32_t destMask = psInst->asOperands[0].GetAccessMask();
uint32_t src2SwizCount = psInst->asOperands[3].GetNumSwizzleElements(destMask);
uint32_t src1SwizCount = psInst->asOperands[2].GetNumSwizzleElements(destMask);
uint32_t src0SwizCount = psInst->asOperands[1].GetNumSwizzleElements(destMask);
uint32_t ui32Flags = 0;
if (src1SwizCount != src0SwizCount || src2SwizCount != src0SwizCount)
{
uint32_t maxElems = std::max(src2SwizCount, std::max(src1SwizCount, src0SwizCount));
ui32Flags |= (TO_AUTO_EXPAND_TO_VEC2 << (maxElems - 2));
}
AddAssignToDest(&psInst->asOperands[0], isUBFE ? SVT_UINT : SVT_INT, psInst->asOperands[0].GetNumSwizzleElements(), &numParenthesis);
bcatcstr(glsl, "bitFieldExtract");
bcatcstr(glsl, isUBFE ? "U" : "I");
bcatcstr(glsl, "(");
glsl << TranslateOperand(&psInst->asOperands[1], TO_FLAG_UNSIGNED_INTEGER, destMask);
glsl << TranslateOperand(&psInst->asOperands[1], ui32Flags | TO_FLAG_UNSIGNED_INTEGER, destMask);
bcatcstr(glsl, ", ");
glsl << TranslateOperand(&psInst->asOperands[2], TO_FLAG_UNSIGNED_INTEGER, destMask);
glsl << TranslateOperand(&psInst->asOperands[2], ui32Flags | TO_FLAG_UNSIGNED_INTEGER, destMask);
bcatcstr(glsl, ", ");
glsl << TranslateOperand(&psInst->asOperands[3], isUBFE ? TO_FLAG_UNSIGNED_INTEGER : TO_FLAG_INTEGER, destMask);
glsl << TranslateOperand(&psInst->asOperands[3], ui32Flags | (isUBFE ? TO_FLAG_UNSIGNED_INTEGER : TO_FLAG_INTEGER), destMask);
bcatcstr(glsl, ")");
AddAssignPrologue(numParenthesis);
break;

52
src/toMetalOperand.cpp
View File

@ -409,11 +409,17 @@ std::string ToMetal::TranslateVariableName(const Operand* psOperand, uint32_t ui
{
if (CanDoDirectCast(psContext, eType, requestedType))
{
oss << GetConstructorForType(psContext, requestedType, requestedComponents, false) << "(";
numParenthesis++;
hasCtor = 1;
if (eType == SVT_BOOL)
{
needsBoolUpscale = 1;
// make sure to wrap the whole thing in parens so the upscale
// multiply only applies to the bool
oss << "(";
numParenthesis++;
}
oss << GetConstructorForType(psContext, requestedType, requestedComponents, false) << "(";
numParenthesis++;
}
else
{
@ -554,6 +560,13 @@ std::string ToMetal::TranslateVariableName(const Operand* psOperand, uint32_t ui
case OPERAND_TYPE_TEMP:
{
SHADER_VARIABLE_TYPE eTempType = psOperand->GetDataType(psContext);
if (psOperand->eSpecialName == NAME_UNDEFINED && psOperand->specialName.length())
{
oss << psOperand->specialName;
break;
}
oss << HLSLCC_TEMP_PREFIX;
ASSERT(psOperand->ui32RegisterNumber < 0x10000); // Sanity check after temp splitting.
switch (eTempType)
@ -993,7 +1006,18 @@ std::string ToMetal::TranslateVariableName(const Operand* psOperand, uint32_t ui
}
case OPERAND_TYPE_INPUT_THREAD_ID_IN_GROUP_FLATTENED://SV_GroupIndex
{
oss << "mtl_ThreadIndexInThreadGroup";
if (requestedComponents > 1 && !hasCtor)
{
oss << GetConstructorForType(psContext, eType, requestedComponents, false) << "(";
numParenthesis++;
hasCtor = 1;
}
for (uint32_t i = 0; i < requestedComponents; i++)
{
oss << "mtl_ThreadIndexInThreadGroup";
if (i < requestedComponents - 1)
oss << ", ";
}
*pui32IgnoreSwizzle = 1; // No swizzle meaningful for scalar.
break;
}
@ -1087,42 +1111,45 @@ std::string ToMetal::TranslateVariableName(const Operand* psOperand, uint32_t ui
// Not on Metal
ASSERT(0);
break;
// as far as i understand tesselation factors are always coming from tessFactor variable (it is always declared in ToMetal::Translate)
case NAME_FINAL_QUAD_U_EQ_0_EDGE_TESSFACTOR:
case NAME_FINAL_TRI_U_EQ_0_EDGE_TESSFACTOR:
case NAME_FINAL_LINE_DENSITY_TESSFACTOR:
if (psContext->psShader->aIndexedOutput[1][psOperand->ui32RegisterNumber])
oss << "edgeTessellationFactor";
oss << "tessFactor.edgeTessellationFactor";
else
oss << "edgeTessellationFactor[0]";
oss << "tessFactor.edgeTessellationFactor[0]";
*pui32IgnoreSwizzle = 1;
break;
case NAME_FINAL_QUAD_V_EQ_0_EDGE_TESSFACTOR:
case NAME_FINAL_TRI_V_EQ_0_EDGE_TESSFACTOR:
case NAME_FINAL_LINE_DETAIL_TESSFACTOR:
oss << "edgeTessellationFactor[1]";
oss << "tessFactor.edgeTessellationFactor[1]";
*pui32IgnoreSwizzle = 1;
break;
case NAME_FINAL_QUAD_U_EQ_1_EDGE_TESSFACTOR:
case NAME_FINAL_TRI_W_EQ_0_EDGE_TESSFACTOR:
oss << "edgeTessellationFactor[2]";
oss << "tessFactor.edgeTessellationFactor[2]";
*pui32IgnoreSwizzle = 1;
break;
case NAME_FINAL_QUAD_V_EQ_1_EDGE_TESSFACTOR:
oss << "edgeTessellationFactor[3]";
oss << "tessFactor.edgeTessellationFactor[3]";
*pui32IgnoreSwizzle = 1;
break;
case NAME_FINAL_TRI_INSIDE_TESSFACTOR:
case NAME_FINAL_QUAD_U_INSIDE_TESSFACTOR:
if (psContext->psShader->aIndexedOutput[1][psOperand->ui32RegisterNumber])
oss << "insideTessellationFactor";
oss << "tessFactor.insideTessellationFactor";
else
oss << "insideTessellationFactor[0]";
oss << "tessFactor.insideTessellationFactor[0]";
*pui32IgnoreSwizzle = 1;
break;
case NAME_FINAL_QUAD_V_INSIDE_TESSFACTOR:
oss << "insideTessellationFactor[1]";
oss << "tessFactor.insideTessellationFactor[1]";
*pui32IgnoreSwizzle = 1;
break;
default:
const std::string patchPrefix = "patch.";
@ -1166,6 +1193,9 @@ std::string ToMetal::TranslateVariableName(const Operand* psOperand, uint32_t ui
else
oss << ") * int(0xffffffffu)";
numParenthesis--;
oss << ")";
numParenthesis--;
}
while (numParenthesis != 0)