AvaloniaUI
/
angle
зеркало из https://github.com/AvaloniaUI/angle.git
1
0
Форкнуть 0
Код Задачи Пакеты Проекты Релизы Вики Активность

Upstream Apple's direct-to-Metal backend: compile translator. 

This change is meant to merge the translator changes from Apple's
direct-to-Metal backend. Taken from Kyle Piddington's CL:
https://chromium-review.googlesource.com/c/angle/angle/+/2857366/
The goal of this CL is to merge the translator code in a state that
compiles, but not to switch the Metal backend over to use this
translator backend yet.

Bug: angleproject:5505
Change-Id: I68a6354604498cd5fd1eb96c13fc56f3b38f2bd0
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/2897536
Reviewed-by: Jonah Ryan-Davis <jonahr@google.com>
Commit-Queue: Jonah Ryan-Davis <jonahr@google.com>
This commit is contained in:
Kyle Piddington 2021-04-26 16:56:15 -07:00 коммит произвёл Angle LUCI CQ
Родитель 9459456b09
Коммит d7aa013091
129 изменённых файлов: 21146 добавлений и 202 удалений
Показать статистику Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы

5
BUILD.gn
Просмотреть файл

@ -560,7 +560,10 @@ template("translator_lib") {
if (angle_enable_metal) {
sources += angle_translator_lib_metal_sources
defines += [ "ANGLE_ENABLE_METAL" ]
defines += [
"ANGLE_ENABLE_METAL",
"ANGLE_ENABLE_METAL_SPIRV",
]
}
}

14
include/GLSLANG/ShaderLang.h
Просмотреть файл

@ -26,7 +26,7 @@
// Version number for shader translation API.
// It is incremented every time the API changes.
#define ANGLE_SH_VERSION 258
#define ANGLE_SH_VERSION 259
enum ShShaderSpec
{
@ -74,6 +74,9 @@ enum ShShaderOutput
// Output SPIR-V to be cross compiled to Metal.
SH_SPIRV_METAL_OUTPUT = 0x8B4C,
// Output for MSL
SH_MSL_METAL_OUTPUT = 0x8B4D,
};
// Compile options.
@ -587,6 +590,15 @@ struct ShBuiltInResources
int MaxClipDistances;
int MaxCullDistances;
int MaxCombinedClipAndCullDistances;
// Direct-to-metal backend constants:
// Binding index for driver uniforms:
int DriverUniformsBindingIndex;
// Binding index for default uniforms:
int DefaultUniformsBindingIndex;
// Binding index for UBO's argument buffer
int UBOArgumentBufferBindingIndex;
};
//

1
scripts/export_targets.py
Просмотреть файл

@ -185,6 +185,7 @@ IGNORED_INCLUDES = {
b'compiler/translator/TranslatorGLSL.h',
b'compiler/translator/TranslatorHLSL.h',
b'compiler/translator/TranslatorMetal.h',
b'compiler/translator/TranslatorMetalDirect.h',
b'compiler/translator/TranslatorVulkan.h',
b'contrib/optimizations/slide_hash_neon.h',
b'dirent_on_windows.h',

2
src/common/PoolAlloc.cpp
Просмотреть файл

@ -335,7 +335,7 @@ void PoolAllocator::unlock()
void Allocation::checkAllocList() const
{
for (const Allocation *alloc = this; alloc != 0; alloc = alloc->mPrevAlloc)
alloc->check();
alloc->checkAlloc();
}
} // namespace angle

2
src/common/PoolAlloc.h
Просмотреть файл

@ -65,7 +65,7 @@ class Allocation
#endif
}
void check() const
void checkAlloc() const
{
checkGuardBlock(preGuard(), kGuardBlockBeginVal, "before");
checkGuardBlock(postGuard(), kGuardBlockEndVal, "after");

16
src/common/apple_platform_utils.h
Просмотреть файл

@ -15,24 +15,32 @@
// TARGET_OS_MACCATALYST only available in MacSDK 10.15
#if TARGET_OS_MACCATALYST
// ANGLE_APPLE_AVAILABLE_XCI: check if either of the 3 platforms (OSX/Catalyst/iOS) min verions is
// available:
#if TARGET_OS_MACCATALYST
# define ANGLE_APPLE_AVAILABLE_XCI(macVer, macCatalystVer, iOSVer) \
@available(macOS macVer, macCatalyst macCatalystVer, iOS iOSVer, *)
// ANGLE_APPLE_AVAILABLE_XC: check if either of the 2 platforms (OSX/Catalyst) min verions is
// available:
# define ANGLE_APPLE_AVAILABLE_XC(macVer, macCatalystVer) \
@available(macOS macVer, macCatalyst macCatalystVer, *)
// ANGLE_APPLE_AVAILABLE_CI: check if either of the 2 platforms (Catalyst/iOS) min verions is
// available:
# define ANGLE_APPLE_AVAILABLE_CI(macCatalystVer, iOSVer) \
@available(macCatalyst macCatalystVer, iOS iOSVer, *)
#else
# define ANGLE_APPLE_AVAILABLE_XCI(macVer, macCatalystVer, iOSVer) \
ANGLE_APPLE_AVAILABLE_XI(macVer, iOSVer)
// ANGLE_APPLE_AVAILABLE_XC: check if either of the 2 platforms (OSX/Catalyst) min verions is
// available:
# define ANGLE_APPLE_AVAILABLE_XC(macVer, macCatalystVer) @available(macOS macVer, *)
# define ANGLE_APPLE_AVAILABLE_CI(macCatalystVer, iOSVer) @available(iOS iOSVer, tvOS iOSVer, *)
#endif
// ANGLE_APPLE_AVAILABLE_XI: check if either of the 2 platforms (OSX/iOS) min verions is available:
#define ANGLE_APPLE_AVAILABLE_XI(macVer, iOSVer) @available(macOS macVer, iOS iOSVer, *)
#define ANGLE_APPLE_AVAILABLE_XI(macVer, iOSVer) \
@available(macOS macVer, iOS iOSVer, tvOS iOSVer, *)
// ANGLE_APPLE_AVAILABLE_I: check if a particular iOS version is available
#define ANGLE_APPLE_AVAILABLE_I(iOSVer) @available(iOS iOSVer, tvOS iOSVer, *)
#endif

7
src/common/platform.h
Просмотреть файл

@ -140,11 +140,16 @@
# if defined(__arm64__) || defined(__aarch64__)
# define ANGLE_CPU_ARM64 1
# endif
// EAGL should be enabled on iOS, but not Mac Catalyst unless it is running on Apple Silicon.
# // EAGL should be enabled on iOS, but not Mac Catalyst unless it is running on Apple Silicon.
# if (defined(ANGLE_PLATFORM_IOS) && !defined(ANGLE_PLATFORM_MACCATALYST)) || \
(defined(ANGLE_PLATFORM_MACCATALYST) && defined(ANGLE_CPU_ARM64))
# define ANGLE_ENABLE_EAGL
# endif
# // Identify Metal API >= what shipped on macOS Catalina.
# if (defined(ANGLE_PLATFORM_MACOS) && __MAC_OS_X_VERSION_MAX_ALLOWED >= 101500) || \
(defined(ANGLE_PLATFORM_IOS) && __IPHONE_OS_VERSION_MAX_ALLOWED >= 130000)
# define ANGLE_WITH_MODERN_METAL_API 1
# endif
#endif
// Define ANGLE_WITH_ASAN macro.

53
src/common/system_utils_ios.mm Normal file
Просмотреть файл

@ -0,0 +1,53 @@
//
// Copyright 2015 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
// system_utils_ios.mm: Implementation of iOS-specific functions for OSX
#include "system_utils.h"
#include <unistd.h>
#include <CoreServices/CoreServices.h>
#include <mach-o/dyld.h>
#include <mach/mach.h>
#include <mach/mach_time.h>
#include <array>
#include <cstdlib>
#include <vector>
#import <Foundation/Foundation.h>
namespace angle
{
std::string GetExecutablePath()
{
NSString *executableString = [[NSBundle mainBundle] executablePath];
std::string result([executableString UTF8String]);
return result;
}
std::string GetExecutableDirectory()
{
std::string executablePath = GetExecutablePath();
size_t lastPathSepLoc = executablePath.find_last_of("/");
return (lastPathSepLoc != std::string::npos) ? executablePath.substr(0, lastPathSepLoc) : "";
}
const char *GetSharedLibraryExtension()
{
return "dylib";
}
double GetCurrentTime()
{
mach_timebase_info_data_t timebaseInfo;
mach_timebase_info(&timebaseInfo);
double secondCoeff = timebaseInfo.numer * 1e-9 / timebaseInfo.denom;
return secondCoeff * mach_absolute_time();
}
} // namespace angle

1
src/common/utilities.h
Просмотреть файл

@ -37,6 +37,7 @@ size_t VariableInternalSize(GLenum type);
size_t VariableExternalSize(GLenum type);
int VariableRowCount(GLenum type);
int VariableColumnCount(GLenum type);
int VariableAttributeCount(GLenum type);
bool IsSamplerType(GLenum type);
bool IsSamplerCubeType(GLenum type);
bool IsSamplerYUVType(GLenum type);

102
src/compiler.gni
Просмотреть файл

@ -92,6 +92,7 @@ angle_translator_sources = [
"src/compiler/translator/TranslatorGLSL.h",
"src/compiler/translator/TranslatorHLSL.h",
"src/compiler/translator/TranslatorMetal.h",
"src/compiler/translator/TranslatorMetalDirect.h",
"src/compiler/translator/TranslatorVulkan.h",
"src/compiler/translator/Types.cpp",
"src/compiler/translator/Types.h",
@ -140,18 +141,34 @@ angle_translator_sources = [
"src/compiler/translator/tree_ops/ForcePrecisionQualifier.h",
"src/compiler/translator/tree_ops/InitializeVariables.cpp",
"src/compiler/translator/tree_ops/InitializeVariables.h",
"src/compiler/translator/tree_ops/NameNamelessUniformBuffers.cpp",
"src/compiler/translator/tree_ops/NameNamelessUniformBuffers.h",
"src/compiler/translator/tree_ops/PruneEmptyCases.cpp",
"src/compiler/translator/tree_ops/PruneEmptyCases.h",
"src/compiler/translator/tree_ops/PruneNoOps.cpp",
"src/compiler/translator/tree_ops/PruneNoOps.h",
"src/compiler/translator/tree_ops/RemoveArrayLengthMethod.cpp",
"src/compiler/translator/tree_ops/RemoveArrayLengthMethod.h",
"src/compiler/translator/tree_ops/RemoveAtomicCounterBuiltins.cpp",
"src/compiler/translator/tree_ops/RemoveAtomicCounterBuiltins.h",
"src/compiler/translator/tree_ops/RemoveDynamicIndexing.cpp",
"src/compiler/translator/tree_ops/RemoveDynamicIndexing.h",
"src/compiler/translator/tree_ops/RemoveInactiveInterfaceVariables.cpp",
"src/compiler/translator/tree_ops/RemoveInactiveInterfaceVariables.h",
"src/compiler/translator/tree_ops/RemoveInvariantDeclaration.cpp",
"src/compiler/translator/tree_ops/RemoveInvariantDeclaration.h",
"src/compiler/translator/tree_ops/RemoveUnreferencedVariables.cpp",
"src/compiler/translator/tree_ops/RemoveUnreferencedVariables.h",
"src/compiler/translator/tree_ops/RewriteAtomicCounters.cpp",
"src/compiler/translator/tree_ops/RewriteAtomicCounters.h",
"src/compiler/translator/tree_ops/RewriteCubeMapSamplersAs2DArray.cpp",
"src/compiler/translator/tree_ops/RewriteCubeMapSamplersAs2DArray.h",
"src/compiler/translator/tree_ops/RewriteDfdy.cpp",
"src/compiler/translator/tree_ops/RewriteDfdy.h",
"src/compiler/translator/tree_ops/RewriteRowMajorMatrices.cpp",
"src/compiler/translator/tree_ops/RewriteRowMajorMatrices.h",
"src/compiler/translator/tree_ops/RewriteStructSamplers.cpp",
"src/compiler/translator/tree_ops/RewriteStructSamplers.h",
"src/compiler/translator/tree_ops/RewriteTexelFetchOffset.cpp",
"src/compiler/translator/tree_ops/RewriteTexelFetchOffset.h",
"src/compiler/translator/tree_ops/ScalarizeVecAndMatConstructorArgs.cpp",
@ -172,6 +189,7 @@ angle_translator_sources = [
"src/compiler/translator/tree_ops/gl/mac/UnfoldShortCircuitAST.h",
"src/compiler/translator/tree_ops/vulkan/DeclarePerVertexBlocks.h",
"src/compiler/translator/tree_ops/vulkan/EarlyFragmentTestsOptimization.h",
"src/compiler/translator/tree_util/AsNode.h",
"src/compiler/translator/tree_util/BuiltIn.h",
"src/compiler/translator/tree_util/BuiltIn_ESSL_autogen.h",
"src/compiler/translator/tree_util/BuiltIn_complete_autogen.h",
@ -242,13 +260,11 @@ angle_translator_glsl_sources = [
"src/compiler/translator/tree_ops/gl/RewriteRepeatedAssignToSwizzled.cpp",
"src/compiler/translator/tree_ops/gl/UseInterfaceBlockFields.cpp",
"src/compiler/translator/tree_ops/gl/VectorizeVectorScalarArithmetic.cpp",
"src/compiler/translator/tree_ops/gl/mac/RewriteRowMajorMatrices.h",
"src/compiler/translator/tree_ops/gl/mac/RewriteUnaryMinusOperatorFloat.h",
]
angle_translator_glsl_mac_sources = [
"src/compiler/translator/tree_ops/gl/mac/AddAndTrueToLoopCondition.cpp",
"src/compiler/translator/tree_ops/gl/mac/RewriteDoWhile.cpp",
"src/compiler/translator/tree_ops/gl/mac/RewriteRowMajorMatrices.cpp",
"src/compiler/translator/tree_ops/gl/mac/RewriteUnaryMinusOperatorFloat.cpp",
"src/compiler/translator/tree_ops/gl/mac/UnfoldShortCircuitAST.cpp",
]
@ -324,26 +340,14 @@ angle_translator_lib_vulkan_sources = [
"src/compiler/translator/tree_ops/vulkan/FlagSamplersWithTexelFetch.h",
"src/compiler/translator/tree_ops/vulkan/MonomorphizeUnsupportedFunctionsInVulkanGLSL.cpp",
"src/compiler/translator/tree_ops/vulkan/MonomorphizeUnsupportedFunctionsInVulkanGLSL.h",
"src/compiler/translator/tree_ops/vulkan/RemoveAtomicCounterBuiltins.cpp",
"src/compiler/translator/tree_ops/vulkan/RemoveAtomicCounterBuiltins.h",
"src/compiler/translator/tree_ops/vulkan/RemoveInactiveInterfaceVariables.cpp",
"src/compiler/translator/tree_ops/vulkan/RemoveInactiveInterfaceVariables.h",
"src/compiler/translator/tree_ops/vulkan/ReplaceForShaderFramebufferFetch.cpp",
"src/compiler/translator/tree_ops/vulkan/ReplaceForShaderFramebufferFetch.h",
"src/compiler/translator/tree_ops/vulkan/RewriteArrayOfArrayOfOpaqueUniforms.cpp",
"src/compiler/translator/tree_ops/vulkan/RewriteArrayOfArrayOfOpaqueUniforms.h",
"src/compiler/translator/tree_ops/vulkan/RewriteAtomicCounters.cpp",
"src/compiler/translator/tree_ops/vulkan/RewriteAtomicCounters.h",
"src/compiler/translator/tree_ops/vulkan/RewriteCubeMapSamplersAs2DArray.cpp",
"src/compiler/translator/tree_ops/vulkan/RewriteCubeMapSamplersAs2DArray.h",
"src/compiler/translator/tree_ops/vulkan/RewriteDfdy.cpp",
"src/compiler/translator/tree_ops/vulkan/RewriteDfdy.h",
"src/compiler/translator/tree_ops/vulkan/RewriteInterpolateAtOffset.cpp",
"src/compiler/translator/tree_ops/vulkan/RewriteInterpolateAtOffset.h",
"src/compiler/translator/tree_ops/vulkan/RewriteR32fImages.cpp",
"src/compiler/translator/tree_ops/vulkan/RewriteR32fImages.h",
"src/compiler/translator/tree_ops/vulkan/RewriteStructSamplers.cpp",
"src/compiler/translator/tree_ops/vulkan/RewriteStructSamplers.h",
"src/compiler/translator/tree_ops/vulkan/SeparateStructFromUniformDeclarations.cpp",
"src/compiler/translator/tree_ops/vulkan/SeparateStructFromUniformDeclarations.h",
]
@ -363,6 +367,76 @@ angle_translator_lib_metal_sources = [
"src/compiler/translator/OutputVulkanGLSLForMetal.h",
"src/compiler/translator/OutputVulkanGLSLForMetal.mm",
"src/compiler/translator/TranslatorMetal.cpp",
"src/compiler/translator/TranslatorMetalDirect.cpp",
"src/compiler/translator/TranslatorMetalDirect/AddExplicitTypeCasts.cpp",
"src/compiler/translator/TranslatorMetalDirect/AddExplicitTypeCasts.h",
"src/compiler/translator/TranslatorMetalDirect/AstHelpers.cpp",
"src/compiler/translator/TranslatorMetalDirect/AstHelpers.h",
"src/compiler/translator/TranslatorMetalDirect/ConstantNames.h",
"src/compiler/translator/TranslatorMetalDirect/DebugSink.h",
"src/compiler/translator/TranslatorMetalDirect/DiscoverDependentFunctions.cpp",
"src/compiler/translator/TranslatorMetalDirect/DiscoverDependentFunctions.h",
"src/compiler/translator/TranslatorMetalDirect/DiscoverEnclosingFunctionTraverser.cpp",
"src/compiler/translator/TranslatorMetalDirect/DiscoverEnclosingFunctionTraverser.h",
"src/compiler/translator/TranslatorMetalDirect/EmitMetal.cpp",
"src/compiler/translator/TranslatorMetalDirect/EmitMetal.h",
"src/compiler/translator/TranslatorMetalDirect/FixTypeConstructors.cpp",
"src/compiler/translator/TranslatorMetalDirect/FixTypeConstructors.h",
"src/compiler/translator/TranslatorMetalDirect/HoistConstants.cpp",
"src/compiler/translator/TranslatorMetalDirect/HoistConstants.h",
"src/compiler/translator/TranslatorMetalDirect/IdGen.cpp",
"src/compiler/translator/TranslatorMetalDirect/IdGen.h",
"src/compiler/translator/TranslatorMetalDirect/IntermRebuild.cpp",
"src/compiler/translator/TranslatorMetalDirect/IntermRebuild.h",
"src/compiler/translator/TranslatorMetalDirect/IntroduceVertexIndexID.cpp",
"src/compiler/translator/TranslatorMetalDirect/IntroduceVertexIndexID.h",
"src/compiler/translator/TranslatorMetalDirect/Layout.cpp",
"src/compiler/translator/TranslatorMetalDirect/Layout.h",
"src/compiler/translator/TranslatorMetalDirect/MapFunctionsToDefinitions.cpp",
"src/compiler/translator/TranslatorMetalDirect/MapFunctionsToDefinitions.h",
"src/compiler/translator/TranslatorMetalDirect/MapSymbols.cpp",
"src/compiler/translator/TranslatorMetalDirect/MapSymbols.h",
"src/compiler/translator/TranslatorMetalDirect/ModifyStruct.cpp",
"src/compiler/translator/TranslatorMetalDirect/ModifyStruct.h",
"src/compiler/translator/TranslatorMetalDirect/Name.cpp",
"src/compiler/translator/TranslatorMetalDirect/Name.h",
"src/compiler/translator/TranslatorMetalDirect/NameEmbeddedUniformStructsMetal.cpp",
"src/compiler/translator/TranslatorMetalDirect/NameEmbeddedUniformStructsMetal.h",
"src/compiler/translator/TranslatorMetalDirect/NodeType.h",
"src/compiler/translator/TranslatorMetalDirect/Pipeline.cpp",
"src/compiler/translator/TranslatorMetalDirect/Pipeline.h",
"src/compiler/translator/TranslatorMetalDirect/ProgramPrelude.cpp",
"src/compiler/translator/TranslatorMetalDirect/ProgramPrelude.h",
"src/compiler/translator/TranslatorMetalDirect/ReduceInterfaceBlocks.cpp",
"src/compiler/translator/TranslatorMetalDirect/ReduceInterfaceBlocks.h",
"src/compiler/translator/TranslatorMetalDirect/Reference.h",
"src/compiler/translator/TranslatorMetalDirect/RewriteCaseDeclarations.cpp",
"src/compiler/translator/TranslatorMetalDirect/RewriteCaseDeclarations.h",
"src/compiler/translator/TranslatorMetalDirect/RewriteGlobalQualifierDecls.cpp",
"src/compiler/translator/TranslatorMetalDirect/RewriteGlobalQualifierDecls.h",
"src/compiler/translator/TranslatorMetalDirect/RewriteKeywords.cpp",
"src/compiler/translator/TranslatorMetalDirect/RewriteKeywords.h",
"src/compiler/translator/TranslatorMetalDirect/RewriteOutArgs.cpp",
"src/compiler/translator/TranslatorMetalDirect/RewriteOutArgs.h",
"src/compiler/translator/TranslatorMetalDirect/RewritePipelines.cpp",
"src/compiler/translator/TranslatorMetalDirect/RewritePipelines.h",
"src/compiler/translator/TranslatorMetalDirect/RewriteUnaddressableReferences.cpp",
"src/compiler/translator/TranslatorMetalDirect/RewriteUnaddressableReferences.h",
"src/compiler/translator/TranslatorMetalDirect/SeparateCompoundExpressions.cpp",
"src/compiler/translator/TranslatorMetalDirect/SeparateCompoundExpressions.h",
"src/compiler/translator/TranslatorMetalDirect/SeparateCompoundStructDeclarations.cpp",
"src/compiler/translator/TranslatorMetalDirect/SeparateCompoundStructDeclarations.h",
"src/compiler/translator/TranslatorMetalDirect/SkippingTraverser.h",
"src/compiler/translator/TranslatorMetalDirect/SymbolEnv.cpp",
"src/compiler/translator/TranslatorMetalDirect/SymbolEnv.h",
"src/compiler/translator/TranslatorMetalDirect/ToposortStructs.cpp",
"src/compiler/translator/TranslatorMetalDirect/ToposortStructs.h",
"src/compiler/translator/TranslatorMetalDirect/TranslatorMetalUtils.cpp",
"src/compiler/translator/TranslatorMetalDirect/TranslatorMetalUtils.h",
"src/compiler/translator/TranslatorMetalDirect/TransposeRowMajorMatrices.cpp",
"src/compiler/translator/TranslatorMetalDirect/TransposeRowMajorMatrices.h",
"src/compiler/translator/TranslatorMetalDirect/WrapMain.cpp",
"src/compiler/translator/TranslatorMetalDirect/WrapMain.h",
]
angle_preprocessor_sources = [

14
src/compiler/translator/CodeGen.cpp
Просмотреть файл

@ -21,9 +21,13 @@
#endif // ANGLE_ENABLE_VULKAN
#ifdef ANGLE_ENABLE_METAL
# include "compiler/translator/TranslatorMetal.h"
# include "compiler/translator/TranslatorMetalDirect.h"
#endif // ANGLE_ENABLE_METAL
#ifdef ANGLE_ENABLE_METAL_SPIRV
# include "compiler/translator/TranslatorMetal.h"
#endif // ANGLE_ENABLE_METAL_SPIRV
#include "compiler/translator/util.h"
namespace sh
@ -64,11 +68,17 @@ TCompiler *ConstructCompiler(sh::GLenum type, ShShaderSpec spec, ShShaderOutput
}
#endif // ANGLE_ENABLE_VULKAN
#ifdef ANGLE_ENABLE_METAL
#ifdef ANGLE_ENABLE_METAL_SPIRV
if (IsOutputMetal(output))
{
return new TranslatorMetal(type, spec);
}
#endif
#ifdef ANGLE_ENABLE_METAL
if (IsOutputMetalDirect(output))
{
return new TranslatorMetalDirect(type, spec, output);
}
#endif // ANGLE_ENABLE_METAL
// Unsupported compiler or unknown format. Return nullptr per the sh::ConstructCompiler API.

2
src/compiler/translator/Common.h
Просмотреть файл

@ -31,6 +31,8 @@ struct TSourceLoc
int last_line;
};
constexpr TSourceLoc kNoSourceLoc{-1, -1, -1, -1};
//
// Put POOL_ALLOCATOR_NEW_DELETE in base classes to make them use this scheme.
//

6
src/compiler/translator/Compiler.h
Просмотреть файл

@ -36,6 +36,9 @@ class TParseContext;
#ifdef ANGLE_ENABLE_HLSL
class TranslatorHLSL;
#endif // ANGLE_ENABLE_HLSL
#ifdef ANGLE_ENABLE_METAL
class TranslatorMetalDirect;
#endif // ANGLE_ENABLE_METAL
using SpecConstUsageBits = angle::PackedEnumBitSet<vk::SpecConstUsage, uint32_t>;
@ -66,6 +69,9 @@ class TShHandleBase
#ifdef ANGLE_ENABLE_HLSL
virtual TranslatorHLSL *getAsTranslatorHLSL() { return 0; }
#endif // ANGLE_ENABLE_HLSL
#ifdef ANGLE_ENABLE_METAL
virtual TranslatorMetalDirect *getAsTranslatorMetalDirect() { return nullptr; }
#endif // ANGLE_ENABLE_METAL
protected:
// Memory allocator. Allocates and tracks memory required by the compiler.

14
src/compiler/translator/ConstantUnion.cpp
Просмотреть файл

@ -67,11 +67,15 @@ bool IsValidShiftOffset(const TConstantUnion &rhs)
} // anonymous namespace
TConstantUnion::TConstantUnion()
{
iConst = 0;
type = EbtVoid;
}
TConstantUnion::TConstantUnion() : iConst(0), type(EbtVoid) {}
TConstantUnion::TConstantUnion(int i) : iConst(i), type(EbtInt) {}
TConstantUnion::TConstantUnion(unsigned int u) : uConst(u), type(EbtUInt) {}
TConstantUnion::TConstantUnion(float f) : fConst(f), type(EbtFloat) {}
TConstantUnion::TConstantUnion(bool b) : bConst(b), type(EbtBool) {}
int TConstantUnion::getIConst() const
{

4
src/compiler/translator/ConstantUnion.h
Просмотреть файл

@ -20,6 +20,10 @@ class TConstantUnion
public:
POOL_ALLOCATOR_NEW_DELETE
TConstantUnion();
TConstantUnion(int i);
TConstantUnion(unsigned int u);
TConstantUnion(float f);
TConstantUnion(bool b);
bool cast(TBasicType newType, const TConstantUnion &constant);

41
src/compiler/translator/IntermNode.cpp
Просмотреть файл

@ -41,7 +41,7 @@ TPrecision GetHigherPrecision(TPrecision left, TPrecision right)
TConstantUnion *Vectorize(const TConstantUnion &constant, size_t size)
{
TConstantUnion *constUnion = new TConstantUnion[size];
for (unsigned int i = 0; i < size; ++i)
for (size_t i = 0; i < size; ++i)
constUnion[i] = constant;
return constUnion;
@ -418,6 +418,14 @@ TIntermBlock::TIntermBlock(const TIntermBlock &node)
mIsTreeRoot = false;
}
TIntermBlock::TIntermBlock(std::initializer_list<TIntermNode *> stmts)
{
for (TIntermNode *stmt : stmts)
{
appendStatement(stmt);
}
}
size_t TIntermBlock::getChildCount() const
{
return mStatements.size();
@ -455,6 +463,37 @@ bool TIntermFunctionPrototype::replaceChildNode(TIntermNode *original, TIntermNo
return false;
}
TIntermDeclaration::TIntermDeclaration(const TVariable *var, TIntermTyped *initExpr)
{
if (initExpr)
{
appendDeclarator(
new TIntermBinary(TOperator::EOpInitialize, new TIntermSymbol(var), initExpr));
}
else
{
appendDeclarator(new TIntermSymbol(var));
}
}
TIntermDeclaration::TIntermDeclaration(std::initializer_list<const TVariable *> declarators)
: TIntermDeclaration()
{
for (const TVariable *d : declarators)
{
appendDeclarator(new TIntermSymbol(d));
}
}
TIntermDeclaration::TIntermDeclaration(std::initializer_list<TIntermTyped *> declarators)
: TIntermDeclaration()
{
for (TIntermTyped *d : declarators)
{
appendDeclarator(d);
}
}
size_t TIntermDeclaration::getChildCount() const
{
return mDeclarators.size();

4
src/compiler/translator/IntermNode.h
Просмотреть файл

@ -674,6 +674,7 @@ class TIntermBlock : public TIntermNode, public TIntermAggregateBase
{
public:
TIntermBlock() : TIntermNode(), mIsTreeRoot(false) {}
TIntermBlock(std::initializer_list<TIntermNode *> stmts);
~TIntermBlock() override {}
TIntermBlock *getAsBlock() override { return this; }
@ -784,6 +785,9 @@ class TIntermDeclaration : public TIntermNode, public TIntermAggregateBase
{
public:
TIntermDeclaration() : TIntermNode() {}
TIntermDeclaration(const TVariable *var, TIntermTyped *initExpr);
TIntermDeclaration(std::initializer_list<const TVariable *> declarators);
TIntermDeclaration(std::initializer_list<TIntermTyped *> declarators);
~TIntermDeclaration() override {}
TIntermDeclaration *getAsDeclarationNode() override { return this; }

3
src/compiler/translator/Symbol.cpp
Просмотреть файл

@ -197,7 +197,8 @@ void TFunction::shareParameters(const TFunction &parametersSource)
ImmutableString TFunction::buildMangledName() const
{
std::string newName(name().data(), name().length());
ImmutableString name = this->name();
std::string newName(name.data(), name.length());
newName += kFunctionMangledNameSeparator;
for (size_t i = 0u; i < mParamCount; ++i)

2
src/compiler/translator/Symbol.h
Просмотреть файл

@ -222,7 +222,7 @@ class TFunction : public TSymbol
TOperator getBuiltInOp() const { return mOp; }
void setDefined() { defined = true; }
bool isDefined() { return defined; }
bool isDefined() const { return defined; }
void setHasPrototypeDeclaration() { mHasPrototypeDeclaration = true; }
bool hasPrototypeDeclaration() const { return mHasPrototypeDeclaration; }

7
src/compiler/translator/SymbolTable.h
Просмотреть файл

@ -297,6 +297,13 @@ class TSymbolTable : angle::NonCopyable, TSymbolTableBase
const ShBuiltInResources &resources);
void clearCompilationResults();
int getDefaultUniformsBindingIndex() const { return mResources.DriverUniformsBindingIndex; }
int getDriverUniformsBindingIndex() const { return mResources.DefaultUniformsBindingIndex; }
int getUBOArgumentBufferBindingIndex() const
{
return mResources.UBOArgumentBufferBindingIndex;
}
private:
friend class TSymbolUniqueId;

2
src/compiler/translator/TranslatorGLSL.cpp
Просмотреть файл

@ -11,8 +11,8 @@
#include "compiler/translator/ExtensionGLSL.h"
#include "compiler/translator/OutputGLSL.h"
#include "compiler/translator/VersionGLSL.h"
#include "compiler/translator/tree_ops/RewriteRowMajorMatrices.h"
#include "compiler/translator/tree_ops/RewriteTexelFetchOffset.h"
#include "compiler/translator/tree_ops/gl/mac/RewriteRowMajorMatrices.h"
#include "compiler/translator/tree_ops/gl/mac/RewriteUnaryMinusOperatorFloat.h"
namespace sh

1227
src/compiler/translator/TranslatorMetalDirect.cpp Normal file
Просмотреть файл

Разница между файлами не показана из-за своего большого размера Загрузить разницу

197
src/compiler/translator/TranslatorMetalDirect.h Normal file
Просмотреть файл

@ -0,0 +1,197 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_H_
#include "compiler/translator/Compiler.h"
namespace sh
{
constexpr const char kUniformsVar[] = "angleUniforms";
constexpr const char kViewport[] = "viewport";
constexpr const char kHalfRenderArea[] = "halfRenderArea";
constexpr const char kFlipXY[] = "flipXY";
constexpr const char kNegFlipXY[] = "negFlipXY";
constexpr const char kClipDistancesEnabled[] = "clipDistancesEnabled";
constexpr const char kXfbActiveUnpaused[] = "xfbActiveUnpaused";
constexpr const char kXfbVerticesPerDraw[] = "xfbVerticesPerDraw";
constexpr const char kXfbBufferOffsets[] = "xfbBufferOffsets";
constexpr const char kAcbBufferOffsets[] = "acbBufferOffsets";
constexpr const char kDepthRange[] = "depthRange";
constexpr const char kPreRotation[] = "preRotation";
constexpr const char kFragRotation[] = "fragRotation";
constexpr const char kUnassignedAttributeString[] = " __unassigned_attribute__";
class DriverUniform;
class SpecConst;
class TOutputMSL;
typedef std::unordered_map<size_t, std::string> originalNamesMap;
typedef std::unordered_map<std::string, size_t> samplerBindingMap;
typedef std::unordered_map<std::string, size_t> textureBindingMap;
typedef std::unordered_map<std::string, size_t> userUniformBufferBindingMap;
typedef std::pair<size_t, size_t> uboBindingInfo;
struct UBOBindingInfo
{
size_t bindIndex = 0;
size_t arraySize = 0;
};
typedef std::unordered_map<std::string, UBOBindingInfo> uniformBufferBindingMap;
class TranslatorMetalReflection
{
public:
TranslatorMetalReflection() : hasUBOs(false), hasFlatInput(false) {}
~TranslatorMetalReflection() {}
void addOriginalName(const size_t id, const std::string &name)
{
originalNames.insert({id, name});
}
void addSamplerBinding(const std::string &name, size_t samplerBinding)
{
samplerBindings.insert({name, samplerBinding});
}
void addTextureBinding(const std::string &name, size_t textureBinding)
{
textureBindings.insert({name, textureBinding});
}
void addUserUniformBufferBinding(const std::string &name, size_t userUniformBufferBinding)
{
userUniformBufferBindings.insert({name, userUniformBufferBinding});
}
void addUniformBufferBinding(const std::string &name, UBOBindingInfo bindingInfo)
{
uniformBufferBindings.insert({name, bindingInfo});
}
std::string getOriginalName(const size_t id) { return originalNames.at(id); }
samplerBindingMap getSamplerBindings() const { return samplerBindings; }
textureBindingMap getTextureBindings() const { return textureBindings; }
userUniformBufferBindingMap getUserUniformBufferBindings() const
{
return userUniformBufferBindings;
}
uniformBufferBindingMap getUniformBufferBindings() const { return uniformBufferBindings; }
size_t getSamplerBinding(const std::string &name) const
{
auto it = samplerBindings.find(name);
if (it != samplerBindings.end())
{
return it->second;
}
// If we can't find a matching sampler, assert out on Debug, and return an invalid value on
// release.
ASSERT(0);
return std::numeric_limits<size_t>::max();
}
size_t getTextureBinding(const std::string &name) const
{
auto it = textureBindings.find(name);
if (it != textureBindings.end())
{
return it->second;
}
// If we can't find a matching texture, assert out on Debug, and return an invalid value on
// release.
ASSERT(0);
return std::numeric_limits<size_t>::max();
}
size_t getUserUniformBufferBinding(const std::string &name) const
{
auto it = userUniformBufferBindings.find(name);
if (it != userUniformBufferBindings.end())
{
return it->second;
}
// If we can't find a matching Uniform binding, assert out on Debug, and return an invalid
// value.
ASSERT(0);
return std::numeric_limits<size_t>::max();
}
UBOBindingInfo getUniformBufferBinding(const std::string &name) const
{
auto it = uniformBufferBindings.find(name);
if (it != uniformBufferBindings.end())
{
return it->second;
}
// If we can't find a matching UBO binding by name, assert out on Debug, and return an
// invalid value.
ASSERT(0);
return {.bindIndex = std::numeric_limits<size_t>::max(),
.arraySize = std::numeric_limits<size_t>::max()};
}
void reset()
{
hasUBOs = false;
hasFlatInput = false;
originalNames.clear();
samplerBindings.clear();
textureBindings.clear();
userUniformBufferBindings.clear();
uniformBufferBindings.clear();
}
bool hasUBOs = false;
bool hasFlatInput = false;
private:
originalNamesMap originalNames;
samplerBindingMap samplerBindings;
textureBindingMap textureBindings;
userUniformBufferBindingMap userUniformBufferBindings;
uniformBufferBindingMap uniformBufferBindings;
};
class TranslatorMetalDirect : public TCompiler
{
public:
TranslatorMetalDirect(sh::GLenum type, ShShaderSpec spec, ShShaderOutput output);
#ifdef ANGLE_ENABLE_METAL
TranslatorMetalDirect *getAsTranslatorMetalDirect() override { return this; }
#endif
static const char *GetCoverageMaskEnabledConstName();
static const char *GetRasterizationDiscardEnabledConstName();
void enableEmulatedInstanceID(bool e) { mEmulatedInstanceID = e; }
TranslatorMetalReflection *getTranslatorMetalReflection() { return &translatorMetalReflection; }
protected:
bool translate(TIntermBlock *root,
ShCompileOptions compileOptions,
PerformanceDiagnostics *perfDiagnostics) override;
ANGLE_NO_DISCARD bool translateImpl(TInfoSinkBase &sink,
TIntermBlock *root,
ShCompileOptions compileOptions,
PerformanceDiagnostics *perfDiagnostics,
SpecConst *specConst,
DriverUniform *driverUniforms);
ANGLE_NO_DISCARD bool shouldFlattenPragmaStdglInvariantAll() override;
ANGLE_NO_DISCARD bool transformDepthBeforeCorrection(TIntermBlock *root,
const DriverUniform *driverUniforms);
ANGLE_NO_DISCARD bool insertSampleMaskWritingLogic(TIntermBlock &root,
DriverUniform &driverUniforms);
ANGLE_NO_DISCARD bool insertRasterizationDiscardLogic(TIntermBlock &root);
ANGLE_NO_DISCARD TIntermSwizzle *getDriverUniformNegFlipYRef(
DriverUniform &driverUniforms) const;
bool mEmulatedInstanceID = false;
TranslatorMetalReflection translatorMetalReflection = {};
};
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_H_

95
src/compiler/translator/TranslatorMetalDirect/AddExplicitTypeCasts.cpp Normal file
Просмотреть файл

@ -0,0 +1,95 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/AddExplicitTypeCasts.h"
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
using namespace sh;
namespace
{
class Rewriter : public TIntermRebuild
{
SymbolEnv &mSymbolEnv;
public:
Rewriter(TCompiler &compiler, SymbolEnv &symbolEnv)
: TIntermRebuild(compiler, false, true), mSymbolEnv(symbolEnv)
{}
PostResult visitAggregatePost(TIntermAggregate &callNode) override
{
const size_t argCount = callNode.getChildCount();
const TType &retType = callNode.getType();
if (callNode.isConstructor())
{
if (IsScalarBasicType(retType))
{
if (argCount == 1)
{
TIntermTyped &arg = GetArg(callNode, 0);
const TType argType = arg.getType();
if (argType.isVector())
{
return CoerceSimple(retType, SubVector(arg, 0, 1));
}
}
}
else if (retType.isVector())
{
if (argCount == 1)
{
TIntermTyped &arg = GetArg(callNode, 0);
const TType argType = arg.getType();
if (argType.isVector())
{
return CoerceSimple(retType, SubVector(arg, 0, retType.getNominalSize()));
}
}
for (size_t i = 0; i < argCount; ++i)
{
TIntermTyped &arg = GetArg(callNode, i);
SetArg(callNode, i, CoerceSimple(retType.getBasicType(), arg));
}
}
else if (retType.isMatrix())
{
if (argCount == 1)
{
TIntermTyped &arg = GetArg(callNode, 0);
const TType argType = arg.getType();
if (argType.isMatrix())
{
if (retType.getCols() != argType.getCols() ||
retType.getRows() != argType.getRows())
{
TemplateArg templateArgs[] = {retType.getCols(), retType.getRows()};
return mSymbolEnv.callFunctionOverload(
Name("cast"), retType, *new TIntermSequence{&arg}, 2, templateArgs);
}
}
}
}
}
return callNode;
}
};
} // anonymous namespace
bool sh::AddExplicitTypeCasts(TCompiler &compiler, TIntermBlock &root, SymbolEnv &symbolEnv)
{
Rewriter rewriter(compiler, symbolEnv);
if (!rewriter.rebuildRoot(root))
{
return false;
}
return true;
}

24
src/compiler/translator/TranslatorMetalDirect/AddExplicitTypeCasts.h Normal file
Просмотреть файл

@ -0,0 +1,24 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_ADDEXPLICITTYPECASTS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_ADDEXPLICITTYPECASTS_H_
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/ProgramPrelude.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
namespace sh
{
// Adds explicit type casts into the AST where casting is done implicitly.
ANGLE_NO_DISCARD bool AddExplicitTypeCasts(TCompiler &compiler,
TIntermBlock &root,
SymbolEnv &symbolEnv);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_ADDEXPLICITTYPECASTS_H_

474
src/compiler/translator/TranslatorMetalDirect/AstHelpers.cpp Normal file
Просмотреть файл

@ -0,0 +1,474 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <cstring>
#include <numeric>
#include <unordered_map>
#include <unordered_set>
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
Declaration sh::ViewDeclaration(TIntermDeclaration &declNode)
{
ASSERT(declNode.getChildCount() == 1);
TIntermNode *childNode = declNode.getChildNode(0);
ASSERT(childNode);
TIntermSymbol *symbolNode;
if ((symbolNode = childNode->getAsSymbolNode()))
{
return {*symbolNode, nullptr};
}
else
{
TIntermBinary *initNode = childNode->getAsBinaryNode();
ASSERT(initNode);
ASSERT(initNode->getOp() == TOperator::EOpInitialize);
symbolNode = initNode->getLeft()->getAsSymbolNode();
ASSERT(symbolNode);
return {*symbolNode, initNode->getRight()};
}
}
const TVariable &sh::CreateStructTypeVariable(TSymbolTable &symbolTable,
const TStructure &structure)
{
auto *type = new TType(&structure, true);
auto *var = new TVariable(&symbolTable, ImmutableString(""), type, SymbolType::Empty);
return *var;
}
const TVariable &sh::CreateInstanceVariable(TSymbolTable &symbolTable,
const TStructure &structure,
const Name &name,
TQualifier qualifier,
const TSpan<const unsigned int> *arraySizes)
{
auto *type = new TType(&structure, false);
type->setQualifier(qualifier);
if (arraySizes)
{
type->makeArrays(*arraySizes);
}
auto *var = new TVariable(&symbolTable, name.rawName(), type, name.symbolType());
return *var;
}
static void AcquireFunctionExtras(TFunction &dest, const TFunction &src)
{
if (src.isDefined())
{
dest.setDefined();
}
if (src.hasPrototypeDeclaration())
{
dest.setHasPrototypeDeclaration();
}
}
TIntermSequence &sh::CloneSequenceAndPrepend(const TIntermSequence &seq, TIntermNode &node)
{
auto *newSeq = new TIntermSequence();
newSeq->push_back(&node);
for (TIntermNode *oldNode : seq)
{
newSeq->push_back(oldNode);
}
return *newSeq;
}
void sh::AddParametersFrom(TFunction &dest, const TFunction &src)
{
const size_t paramCount = src.getParamCount();
for (size_t i = 0; i < paramCount; ++i)
{
const TVariable *var = src.getParam(i);
dest.addParameter(var);
}
}
const TFunction &sh::CloneFunction(TSymbolTable &symbolTable,
IdGen &idGen,
const TFunction &oldFunc)
{
ASSERT(oldFunc.symbolType() == SymbolType::UserDefined);
Name newName = idGen.createNewName(Name(oldFunc));
auto &newFunc = *new TFunction(&symbolTable, newName.rawName(), newName.symbolType(),
&oldFunc.getReturnType(), oldFunc.isKnownToNotHaveSideEffects());
AcquireFunctionExtras(newFunc, oldFunc);
AddParametersFrom(newFunc, oldFunc);
return newFunc;
}
const TFunction &sh::CloneFunctionAndPrependParam(TSymbolTable &symbolTable,
IdGen *idGen,
const TFunction &oldFunc,
const TVariable &newParam)
{
ASSERT(oldFunc.symbolType() == SymbolType::UserDefined ||
oldFunc.symbolType() == SymbolType::AngleInternal);
Name newName = idGen ? idGen->createNewName(Name(oldFunc)) : Name(oldFunc);
auto &newFunc = *new TFunction(&symbolTable, newName.rawName(), newName.symbolType(),
&oldFunc.getReturnType(), oldFunc.isKnownToNotHaveSideEffects());
AcquireFunctionExtras(newFunc, oldFunc);
newFunc.addParameter(&newParam);
AddParametersFrom(newFunc, oldFunc);
return newFunc;
}
const TFunction &sh::CloneFunctionAndAppendParams(TSymbolTable &symbolTable,
IdGen *idGen,
const TFunction &oldFunc,
const std::vector<const TVariable *> &newParams)
{
ASSERT(oldFunc.symbolType() == SymbolType::UserDefined ||
oldFunc.symbolType() == SymbolType::AngleInternal);
Name newName = idGen ? idGen->createNewName(Name(oldFunc)) : Name(oldFunc);
auto &newFunc = *new TFunction(&symbolTable, newName.rawName(), newName.symbolType(),
&oldFunc.getReturnType(), oldFunc.isKnownToNotHaveSideEffects());
AcquireFunctionExtras(newFunc, oldFunc);
AddParametersFrom(newFunc, oldFunc);
for (auto *param : newParams)
{
newFunc.addParameter(param);
}
return newFunc;
}
const TFunction &sh::CloneFunctionAndChangeReturnType(TSymbolTable &symbolTable,
IdGen *idGen,
const TFunction &oldFunc,
const TStructure &newReturn)
{
ASSERT(oldFunc.symbolType() == SymbolType::UserDefined);
Name newName = idGen ? idGen->createNewName(Name(oldFunc)) : Name(oldFunc);
auto *newReturnType = new TType(&newReturn, true);
auto &newFunc = *new TFunction(&symbolTable, newName.rawName(), newName.symbolType(),
newReturnType, oldFunc.isKnownToNotHaveSideEffects());
AcquireFunctionExtras(newFunc, oldFunc);
AddParametersFrom(newFunc, oldFunc);
return newFunc;
}
TIntermTyped &sh::GetArg(const TIntermAggregate &call, size_t index)
{
ASSERT(index < call.getChildCount());
auto *arg = call.getChildNode(index);
ASSERT(arg);
auto *targ = arg->getAsTyped();
ASSERT(targ);
return *targ;
}
void sh::SetArg(TIntermAggregate &call, size_t index, TIntermTyped &arg)
{
ASSERT(index < call.getChildCount());
(*call.getSequence())[index] = &arg;
}
int sh::GetFieldIndex(const TStructure &structure, const ImmutableString &fieldName)
{
const TFieldList &fieldList = structure.fields();
int i = 0;
for (TField *field : fieldList)
{
if (field->name() == fieldName)
{
return i;
}
++i;
}
return -1;
}
TIntermBinary &sh::AccessField(const TVariable &structInstanceVar, const ImmutableString &fieldName)
{
return AccessField(*new TIntermSymbol(&structInstanceVar), fieldName);
}
TIntermBinary &sh::AccessField(TIntermTyped &object, const ImmutableString &fieldName)
{
const TStructure *structure = object.getType().getStruct();
ASSERT(structure);
const int index = GetFieldIndex(*structure, fieldName);
ASSERT(index >= 0);
return AccessFieldByIndex(object, index);
}
TIntermBinary &sh::AccessFieldByIndex(TIntermTyped &object, int index)
{
#if defined(ANGLE_ENABLE_ASSERTS)
const TType &type = object.getType();
ASSERT(!type.isArray());
const TStructure *structure = type.getStruct();
ASSERT(structure);
ASSERT(0 <= index);
ASSERT(static_cast<size_t>(index) < structure->fields().size());
#endif
return *new TIntermBinary(
TOperator::EOpIndexDirectStruct, &object,
new TIntermConstantUnion(new TConstantUnion(index), *new TType(TBasicType::EbtInt)));
}
TIntermBinary &sh::AccessIndex(TIntermTyped &indexableNode, int index)
{
#if defined(ANGLE_ENABLE_ASSERTS)
const TType &type = indexableNode.getType();
ASSERT(type.isArray() || type.isVector() || type.isMatrix());
#endif
auto *accessNode = new TIntermBinary(
TOperator::EOpIndexDirect, &indexableNode,
new TIntermConstantUnion(new TConstantUnion(index), *new TType(TBasicType::EbtInt)));
return *accessNode;
}
TIntermTyped &sh::AccessIndex(TIntermTyped &node, const int *index)
{
if (index)
{
return AccessIndex(node, *index);
}
return node;
}
TIntermTyped &sh::SubVector(TIntermTyped &vectorNode, int begin, int end)
{
ASSERT(vectorNode.getType().isVector());
ASSERT(0 <= begin);
ASSERT(end <= 4);
ASSERT(begin <= end);
if (begin == 0 && end == vectorNode.getType().getNominalSize())
{
return vectorNode;
}
TVector<int> offsets(static_cast<size_t>(end - begin));
std::iota(offsets.begin(), offsets.end(), begin);
auto *swizzle = new TIntermSwizzle(&vectorNode, offsets);
return *swizzle;
}
bool sh::IsScalarBasicType(const TType &type)
{
if (!type.isScalar())
{
return false;
}
return HasScalarBasicType(type);
}
bool sh::IsVectorBasicType(const TType &type)
{
if (!type.isVector())
{
return false;
}
return HasScalarBasicType(type);
}
bool sh::HasScalarBasicType(TBasicType type)
{
switch (type)
{
case TBasicType::EbtFloat:
case TBasicType::EbtDouble:
case TBasicType::EbtInt:
case TBasicType::EbtUInt:
case TBasicType::EbtBool:
return true;
default:
return false;
}
}
bool sh::HasScalarBasicType(const TType &type)
{
return HasScalarBasicType(type.getBasicType());
}
static void InitType(TType &type)
{
if (type.isArray())
{
auto sizes = type.getArraySizes();
type.toArrayBaseType();
type.makeArrays(sizes);
}
}
TType &sh::CloneType(const TType &type)
{
auto &clone = *new TType(type);
InitType(clone);
return clone;
}
TType &sh::InnermostType(const TType &type)
{
auto &inner = *new TType(type);
inner.toArrayBaseType();
InitType(inner);
return inner;
}
TType &sh::DropColumns(const TType &matrixType)
{
ASSERT(matrixType.isMatrix());
ASSERT(HasScalarBasicType(matrixType));
const char *mangledName = nullptr;
auto &vectorType =
*new TType(matrixType.getBasicType(), matrixType.getPrecision(), matrixType.getQualifier(),
matrixType.getRows(), 1, matrixType.getArraySizes(), mangledName);
InitType(vectorType);
return vectorType;
}
TType &sh::DropOuterDimension(const TType &arrayType)
{
ASSERT(arrayType.isArray());
const char *mangledName = nullptr;
const auto &arraySizes = arrayType.getArraySizes();
auto &innerType =
*new TType(arrayType.getBasicType(), arrayType.getPrecision(), arrayType.getQualifier(),
arrayType.getNominalSize(), arrayType.getSecondarySize(),
arraySizes.subspan(0, arraySizes.size() - 1), mangledName);
InitType(innerType);
return innerType;
}
static TType &SetTypeDimsImpl(const TType &type, int primary, int secondary)
{
ASSERT(1 < primary && primary <= 4);
ASSERT(1 <= secondary && secondary <= 4);
ASSERT(HasScalarBasicType(type));
const char *mangledName = nullptr;
auto &newType = *new TType(type.getBasicType(), type.getPrecision(), type.getQualifier(),
primary, secondary, type.getArraySizes(), mangledName);
InitType(newType);
return newType;
}
TType &sh::SetVectorDim(const TType &type, int newDim)
{
ASSERT(type.isRank0() || type.isVector());
return SetTypeDimsImpl(type, newDim, 1);
}
TType &sh::SetMatrixRowDim(const TType &matrixType, int newDim)
{
ASSERT(matrixType.isMatrix());
ASSERT(1 < newDim && newDim <= 4);
return SetTypeDimsImpl(matrixType, matrixType.getCols(), newDim);
}
bool sh::HasMatrixField(const TStructure &structure)
{
for (const TField *field : structure.fields())
{
const TType &type = *field->type();
if (type.isMatrix())
{
return true;
}
}
return false;
}
bool sh::HasArrayField(const TStructure &structure)
{
for (const TField *field : structure.fields())
{
const TType &type = *field->type();
if (type.isArray())
{
return true;
}
}
return false;
}
TIntermTyped &sh::CoerceSimple(TBasicType toType, TIntermTyped &fromNode)
{
const TType &fromType = fromNode.getType();
ASSERT(HasScalarBasicType(toType));
ASSERT(HasScalarBasicType(fromType));
ASSERT(!fromType.isArray());
if (toType != fromType.getBasicType())
{
return *TIntermAggregate::CreateConstructor(
*new TType(toType, fromType.getNominalSize(), fromType.getSecondarySize()),
new TIntermSequence{&fromNode});
}
return fromNode;
}
TIntermTyped &sh::CoerceSimple(const TType &toType, TIntermTyped &fromNode)
{
const TType &fromType = fromNode.getType();
ASSERT(HasScalarBasicType(toType));
ASSERT(HasScalarBasicType(fromType));
ASSERT(toType.getNominalSize() == fromType.getNominalSize());
ASSERT(toType.getSecondarySize() == fromType.getSecondarySize());
ASSERT(!toType.isArray());
ASSERT(!fromType.isArray());
if (toType.getBasicType() != fromType.getBasicType())
{
return *TIntermAggregate::CreateConstructor(toType, new TIntermSequence{&fromNode});
}
return fromNode;
}
TIntermTyped &sh::AsType(SymbolEnv &symbolEnv, const TType &toType, TIntermTyped &fromNode)
{
const TType &fromType = fromNode.getType();
ASSERT(HasScalarBasicType(toType));
ASSERT(HasScalarBasicType(fromType));
ASSERT(!toType.isArray());
ASSERT(!fromType.isArray());
if (toType == fromType)
{
return fromNode;
}
TemplateArg targ(toType);
return symbolEnv.callFunctionOverload(Name("as_type", SymbolType::BuiltIn), toType,
*new TIntermSequence{&fromNode}, 1, &targ);
}

157
src/compiler/translator/TranslatorMetalDirect/AstHelpers.h Normal file
Просмотреть файл

@ -0,0 +1,157 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_ASTHELPERS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_ASTHELPERS_H_
#include <cstring>
#include <unordered_map>
#include <unordered_set>
#include "compiler/translator/TranslatorMetalDirect/IdGen.h"
#include "compiler/translator/TranslatorMetalDirect/Name.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
namespace sh
{
// A convenience view of a TIntermDeclaration node's children.
struct Declaration
{
TIntermSymbol &symbol;
TIntermTyped *initExpr; // Non-null iff declaration is initialized.
};
// Returns a `Declaration` view of the given node.
Declaration ViewDeclaration(TIntermDeclaration &declNode);
// Creates a variable for a struct type.
const TVariable &CreateStructTypeVariable(TSymbolTable &symbolTable, const TStructure &structure);
// Creates a variable for a struct instance.
const TVariable &CreateInstanceVariable(TSymbolTable &symbolTable,
const TStructure &structure,
const Name &name,
TQualifier qualifier = TQualifier::EvqTemporary,
const TSpan<const unsigned int> *arraySizes = nullptr);
// The input sequence should be discarded from AST after this is called.
TIntermSequence &CloneSequenceAndPrepend(const TIntermSequence &seq, TIntermNode &node);
// Appends parameters from `src` function to `dest` function.
void AddParametersFrom(TFunction &dest, const TFunction &src);
// Clones a function.
const TFunction &CloneFunction(TSymbolTable &symbolTable, IdGen &idGen, const TFunction &oldFunc);
// Clones a function and prepends the provided extr parameter.
// If `idGen` is null, the original function must be discarded from the AST.
const TFunction &CloneFunctionAndPrependParam(TSymbolTable &symbolTable,
IdGen *idGen,
const TFunction &oldFunc,
const TVariable &newParam);
// Clones a function and appends the provided extra parameters.
// If `idGen` is null, the original function must be discarded from the AST.
const TFunction &CloneFunctionAndAppendParams(TSymbolTable &symbolTable,
IdGen *idGen,
const TFunction &oldFunc,
const std::vector<const TVariable *> &newParam);
// Clones a function and changes its return type.
// If `idGen` is null, the original function must be discarded from the AST.
const TFunction &CloneFunctionAndChangeReturnType(TSymbolTable &symbolTable,
IdGen *idGen,
const TFunction &oldFunc,
const TStructure &newReturn);
// Gets the argument of a function call at the given index.
TIntermTyped &GetArg(const TIntermAggregate &call, size_t index);
// Sets the argument of a function call at the given index.
void SetArg(TIntermAggregate &call, size_t index, TIntermTyped &arg);
// Returns the field index within the given struct for the given field name.
// Returns -1 if the struct has no field with the given name.
int GetFieldIndex(const TStructure &structure, const ImmutableString &fieldName);
// Accesses a field for the given variable with the given field name.
// The variable must be a struct instance.
TIntermBinary &AccessField(const TVariable &structInstanceVar, const ImmutableString &fieldName);
// Accesses a field for the given node with the given field name.
// The node must be a struct instance.
TIntermBinary &AccessField(TIntermTyped &object, const ImmutableString &fieldName);
// Accesses a field for the given node by its field index.
// The node must be a struct instance.
TIntermBinary &AccessFieldByIndex(TIntermTyped &object, int index);
// Accesses an element by index for the given node.
// The node must be an array, vector, or matrix.
TIntermBinary &AccessIndex(TIntermTyped &indexableNode, int index);
// Accesses an element by index for the given node if `index` is non-null.
// Returns the original node if `index` is null.
// The node must be an array, vector, or matrix if `index` is non-null.
TIntermTyped &AccessIndex(TIntermTyped &node, const int *index);
// Returns a subvector based on the input slice range.
// This returns the original node if the slice is an identity for the node.
TIntermTyped &SubVector(TIntermTyped &vectorNode, int begin, int end);
// Matches scalar bool, int, uint, float, double.
bool IsScalarBasicType(const TType &type);
// Matches vector bool, int, uint, float, double.
bool IsVectorBasicType(const TType &type);
// Matches bool, int, uint, float, double.
// Type does not need to be a scalar.
bool HasScalarBasicType(const TType &type);
// Matches bool, int, uint, float, double.
bool HasScalarBasicType(TBasicType type);
// Clones a type.
TType &CloneType(const TType &type);
// Clones a type and drops all array dimensions.
TType &InnermostType(const TType &type);
// Creates a vector type by dropping the columns off of a matrix type.
TType &DropColumns(const TType &matrixType);
// Creates a type by dropping the outer dimension off of an array type.
TType &DropOuterDimension(const TType &arrayType);
// Creates a scalar or vector type by changing the dimensions of a vector type.
TType &SetVectorDim(const TType &type, int newDim);
// Creates a matrix type by changing the row dimensions of a matrix type.
TType &SetMatrixRowDim(const TType &matrixType, int newDim);
// Returns true iff the structure directly contains a field with matrix type.
bool HasMatrixField(const TStructure &structure);
// Returns true iff the structure directly contains a field with array type.
bool HasArrayField(const TStructure &structure);
// Coerces `fromNode` to `toType` by a constructor call of `toType` if their types differ.
// Vector and matrix dimensions are retained.
// Array types are not allowed.
TIntermTyped &CoerceSimple(TBasicType toType, TIntermTyped &fromNode);
// Coerces `fromNode` to `toType` by a constructor call of `toType` if their types differ.
// Vector and matrix dimensions must coincide between to and from.
// Array types are not allowed.
TIntermTyped &CoerceSimple(const TType &toType, TIntermTyped &fromNode);
TIntermTyped &AsType(SymbolEnv &symbolEnv, const TType &toType, TIntermTyped &fromNode);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_ASTHELPERS_H_

26
src/compiler/translator/TranslatorMetalDirect/ConstantNames.h Normal file
Просмотреть файл

@ -0,0 +1,26 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_CONSTANT_NAMES_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_CONSTANT_NAMES_H_
#include "compiler/translator/TranslatorMetalDirect/Name.h"
namespace sh
{
namespace constant_names
{
constexpr Name kCoverageMaskEnabled("ANGLE_CoverageMaskEnabled", SymbolType::AngleInternal);
constexpr Name kRasterizationDiscardEnabled("ANGLE_RasterizationDiscard",
SymbolType::AngleInternal);
} // namespace constant_names
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_CONSTANT_NAMES_H_

159
src/compiler/translator/TranslatorMetalDirect/DebugSink.h Normal file
Просмотреть файл

@ -0,0 +1,159 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <iostream>
#include "compiler/translator/InfoSink.h"
namespace sh
{
class StringObserver
{
public:
StringObserver(const std::string &needle) : needle(needle) { ASSERT(!needle.empty()); }
bool observe(char c)
{
if (needle[currPos] == c)
{
++currPos;
if (currPos == needle.size())
{
reset();
return true;
}
}
else
{
reset();
}
return false;
}
const std::string &getNeedle() const { return needle; }
void reset() { currPos = 0; }
private:
std::string needle;
size_t currPos = 0;
};
class DebugSink : angle::NonCopyable
{
public:
friend class EscapedSink;
class EscapedSink : angle::NonCopyable
{
friend class DebugSink;
private:
EscapedSink(DebugSink &owner) : mOwner(owner), mBegin(owner.size()) {}
public:
EscapedSink(EscapedSink &&other) : mOwner(other.mOwner), mBegin(other.mBegin) {}
~EscapedSink()
{
const char *p = mOwner.c_str();
const int end = mOwner.size();
mOwner.onWrite(p + mBegin, p + end);
}
TInfoSinkBase &get() { return mOwner.mParent; }
operator TInfoSinkBase &() { return get(); }
private:
DebugSink &mOwner;
const int mBegin;
};
public:
DebugSink(TInfoSinkBase &parent, bool alsoLogToStdout)
: mParent(parent), mAlsoLogToStdout(alsoLogToStdout)
{}
void watch(std::string const &needle)
{
if (!needle.empty())
{
mObservers.emplace_back(needle);
}
}
void erase()
{
mParent.erase();
for (StringObserver &observer : mObservers)
{
observer.reset();
}
}
int size() { return mParent.size(); }
const TPersistString &str() const { return mParent.str(); }
const char *c_str() const { return mParent.c_str(); }
EscapedSink escape() { return EscapedSink(*this); }
template <typename T>
DebugSink &operator<<(const T &value)
{
const size_t begin = mParent.size();
mParent << value;
const size_t end = mParent.size();
const char *p = mParent.c_str();
onWrite(p + begin, p + end);
return *this;
}
private:
void onWrite(const char *begin, char const *end)
{
const char *p = begin;
while (p != end)
{
if (mAlsoLogToStdout)
{
std::cout << *p;
}
for (StringObserver &observer : mObservers)
{
if (observer.observe(*p))
{
if (mAlsoLogToStdout)
{
std::cout.flush();
}
const std::string &needle = observer.getNeedle();
(void)needle;
ASSERT(true); // place your breakpoint here
}
}
++p;
}
if (mAlsoLogToStdout)
{
std::cout.flush();
}
}
private:
TInfoSinkBase &mParent;
std::vector<StringObserver> mObservers;
bool mAlsoLogToStdout;
};
} // namespace sh

131
src/compiler/translator/TranslatorMetalDirect/DiscoverDependentFunctions.cpp Normal file
Просмотреть файл

@ -0,0 +1,131 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <cstring>
#include <unordered_map>
#include <unordered_set>
#include "compiler/translator/TranslatorMetalDirect/DiscoverDependentFunctions.h"
#include "compiler/translator/TranslatorMetalDirect/DiscoverEnclosingFunctionTraverser.h"
#include "compiler/translator/TranslatorMetalDirect/MapFunctionsToDefinitions.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
class Discoverer : public DiscoverEnclosingFunctionTraverser
{
private:
const std::function<bool(const TVariable &)> &mVars;
const FunctionToDefinition &mFuncToDef;
std::unordered_set<const TFunction *> mNonDepFunctions;
public:
std::unordered_set<const TFunction *> mDepFunctions;
public:
Discoverer(const std::function<bool(const TVariable &)> &vars,
const FunctionToDefinition &funcToDef)
: DiscoverEnclosingFunctionTraverser(true, false, true), mVars(vars), mFuncToDef(funcToDef)
{}
void visitSymbol(TIntermSymbol *symbolNode) override
{
const TVariable &var = symbolNode->variable();
if (!mVars(var))
{
return;
}
const TFunction *owner = discoverEnclosingFunction(symbolNode);
ASSERT(owner);
mDepFunctions.insert(owner);
}
bool visitAggregate(Visit visit, TIntermAggregate *aggregateNode) override
{
if (visit != Visit::PreVisit)
{
return true;
}
if (!aggregateNode->isConstructor())
{
const TFunction *func = aggregateNode->getFunction();
if (mNonDepFunctions.find(func) != mNonDepFunctions.end())
{
return true;
}
if (mDepFunctions.find(func) == mDepFunctions.end())
{
auto it = mFuncToDef.find(func);
if (it == mFuncToDef.end())
{
return true;
}
// Recursion is banned in GLSL, so I believe AngleIR has this property too.
// This implementation assumes (direct and mutual) recursion is prohibited.
TIntermFunctionDefinition &funcDefNode = *it->second;
funcDefNode.traverse(this);
if (mNonDepFunctions.find(func) != mNonDepFunctions.end())
{
return true;
}
ASSERT(mDepFunctions.find(func) != mDepFunctions.end());
}
const TFunction *owner = discoverEnclosingFunction(aggregateNode);
ASSERT(owner);
mDepFunctions.insert(owner);
}
return true;
}
bool visitFunctionDefinition(Visit visit, TIntermFunctionDefinition *funcDefNode) override
{
const TFunction *func = funcDefNode->getFunction();
if (visit != Visit::PostVisit)
{
if (mDepFunctions.find(func) != mDepFunctions.end())
{
return false;
}
if (mNonDepFunctions.find(func) != mNonDepFunctions.end())
{
return false;
}
return true;
}
if (mDepFunctions.find(func) == mDepFunctions.end())
{
mNonDepFunctions.insert(func);
}
return true;
}
};
} // namespace
std::unordered_set<const TFunction *> sh::DiscoverDependentFunctions(
TIntermBlock &root,
const std::function<bool(const TVariable &)> &vars)
{
const FunctionToDefinition funcToDef = MapFunctionsToDefinitions(root);
Discoverer discoverer(vars, funcToDef);
root.traverse(&discoverer);
return std::move(discoverer.mDepFunctions);
}

25
src/compiler/translator/TranslatorMetalDirect/DiscoverDependentFunctions.h Normal file
Просмотреть файл

@ -0,0 +1,25 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_DISCOVERDEPENDENTFUNCTIONS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_DISCOVERDEPENDENTFUNCTIONS_H_
#include <unordered_set>
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
namespace sh
{
// Finds and returns all functions that contain the provided variables.
ANGLE_NO_DISCARD std::unordered_set<const TFunction *> DiscoverDependentFunctions(
TIntermBlock &root,
const std::function<bool(const TVariable &)> &vars);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_DISCOVERDEPENDENTFUNCTIONS_H_

33
src/compiler/translator/TranslatorMetalDirect/DiscoverEnclosingFunctionTraverser.cpp Normal file
Просмотреть файл

@ -0,0 +1,33 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/DiscoverEnclosingFunctionTraverser.h"
using namespace sh;
DiscoverEnclosingFunctionTraverser::DiscoverEnclosingFunctionTraverser(bool preVisit_,
bool inVisit_,
bool postVisit_,
TSymbolTable *symbolTable)
: TIntermTraverser(preVisit_, inVisit_, postVisit_, symbolTable)
{}
const TFunction *DiscoverEnclosingFunctionTraverser::discoverEnclosingFunction(TIntermNode *node)
{
ASSERT(!node->getAsFunctionDefinition());
unsigned height = 0;
while (TIntermNode *ancestor = getAncestorNode(height))
{
if (TIntermFunctionDefinition *funcDefNode = ancestor->getAsFunctionDefinition())
{
return funcDefNode->getFunction();
}
++height;
}
return nullptr;
}

31
src/compiler/translator/TranslatorMetalDirect/DiscoverEnclosingFunctionTraverser.h Normal file
Просмотреть файл

@ -0,0 +1,31 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_DISCOVERENCLOSINGFUNCTIONTRAVERSER_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_DISCOVERENCLOSINGFUNCTIONTRAVERSER_H_
#include "compiler/translator/tree_util/IntermTraverse.h"
namespace sh
{
// A TIntermTraverser that supports discovery of the function a node belongs to.
class DiscoverEnclosingFunctionTraverser : public TIntermTraverser
{
public:
DiscoverEnclosingFunctionTraverser(bool preVisit,
bool inVisit,
bool postVisit,
TSymbolTable *symbolTable = nullptr);
// Returns the function a node belongs inside.
// Returns null if the node does not belong inside a function.
const TFunction *discoverEnclosingFunction(TIntermNode *node);
};
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_DISCOVERENCLOSINGFUNCTIONTRAVERSER_H_

2492
src/compiler/translator/TranslatorMetalDirect/EmitMetal.cpp Normal file
Просмотреть файл

Разница между файлами не показана из-за своего большого размера Загрузить разницу

33
src/compiler/translator/TranslatorMetalDirect/EmitMetal.h Normal file
Просмотреть файл

@ -0,0 +1,33 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_EMITMETAL_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_EMITMETAL_H_
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/IdGen.h"
#include "compiler/translator/TranslatorMetalDirect/ProgramPrelude.h"
#include "compiler/translator/TranslatorMetalDirect/RewriteGlobalQualifierDecls.h"
#include "compiler/translator/TranslatorMetalDirect/RewritePipelines.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
namespace sh
{
// Walks the AST and emits Metal code.
ANGLE_NO_DISCARD bool EmitMetal(TCompiler &compiler,
TIntermBlock &root,
IdGen &idGen,
const PipelineStructs &pipelineStructs,
const Invariants &invariants,
SymbolEnv &symbolEnv,
const ProgramPreludeConfig &ppc,
TSymbolTable *symbolTable);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_EMITMETAL_H_

102
src/compiler/translator/TranslatorMetalDirect/FixTypeConstructors.cpp Normal file
Просмотреть файл

@ -0,0 +1,102 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/FixTypeConstructors.h"
#include <unordered_map>
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
#include "compiler/translator/tree_ops/SimplifyLoopConditions.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
class FixTypeTraverser : public TIntermTraverser
{
public:
FixTypeTraverser() : TIntermTraverser(false, false, true) {}
bool visitAggregate(Visit visit, TIntermAggregate *aggregateNode) override
{
if (visit != Visit::PostVisit)
{
return true;
}
if (aggregateNode->isConstructor())
{
const TType &retType = aggregateNode->getType();
if (retType.isScalar())
{
// No-op.
}
else if (retType.isVector())
{
size_t primarySize = retType.getNominalSize() * retType.getArraySizeProduct();
TIntermSequence *args = aggregateNode->getSequence();
size_t argsSize = 0;
size_t beforeSize = 0;
TIntermNode *lastArg = nullptr;
for (TIntermNode *&arg : *args)
{
TIntermTyped *targ = arg->getAsTyped();
lastArg = arg;
if (targ)
{
argsSize += targ->getNominalSize();
}
if (argsSize <= primarySize)
{
beforeSize += targ->getNominalSize();
}
}
if (argsSize > primarySize)
{
size_t swizzleSize = primarySize - beforeSize;
TIntermTyped *targ = lastArg->getAsTyped();
TIntermSwizzle *newSwizzle = nullptr;
switch (swizzleSize)
{
case 1:
newSwizzle = new TIntermSwizzle(targ->deepCopy(), {0});
break;
case 2:
newSwizzle = new TIntermSwizzle(targ->deepCopy(), {0, 1});
break;
case 3:
newSwizzle = new TIntermSwizzle(targ->deepCopy(), {0, 1, 2});
break;
default:
UNREACHABLE(); // Should not be reached in case of 0, or 4
}
if (newSwizzle)
{
this->queueReplacementWithParent(aggregateNode, lastArg, newSwizzle,
OriginalNode::IS_DROPPED);
}
}
}
else if (retType.isMatrix())
{
// TBD if issues
}
}
return true;
}
};
} // anonymous namespace
////////////////////////////////////////////////////////////////////////////////
bool sh::FixTypeConstructors(TCompiler &compiler, SymbolEnv &symbolEnv, TIntermBlock &root)
{
FixTypeTraverser traverser;
root.traverse(&traverser);
if (!traverser.updateTree(&compiler, &root))
{
return false;
}
return true;
}

21
src/compiler/translator/TranslatorMetalDirect/FixTypeConstructors.h Normal file
Просмотреть файл

@ -0,0 +1,21 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_FIXTYPECONSTRUCTORS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_FIXTYPECONSTRUCTORS_H_
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
namespace sh
{
ANGLE_NO_DISCARD bool FixTypeConstructors(TCompiler &compiler,
SymbolEnv &SymbolEnv,
TIntermBlock &root);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_FIXTYPECONSTRUCTORS_H_

96
src/compiler/translator/TranslatorMetalDirect/HoistConstants.cpp Normal file
Просмотреть файл

@ -0,0 +1,96 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/HoistConstants.h"
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
#include "compiler/translator/TranslatorMetalDirect/Layout.h"
#include "compiler/translator/tree_util/FindFunction.h"
#include "compiler/translator/tree_util/ReplaceVariable.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
class Rewriter : private TIntermRebuild
{
private:
const size_t mMinRequiredSize;
TIntermSequence mHoistedDeclNodes;
public:
Rewriter(TCompiler &compiler, size_t minRequiredSize)
: TIntermRebuild(compiler, true, false), mMinRequiredSize(minRequiredSize)
{}
PreResult visitDeclarationPre(TIntermDeclaration &declNode) override
{
if (getParentFunction())
{
Declaration decl = ViewDeclaration(declNode);
const TType &type = decl.symbol.getType();
if (type.getQualifier() == TQualifier::EvqConst)
{
if (decl.initExpr && decl.initExpr->hasConstantValue())
{
const size_t size = MetalLayoutOf(type).sizeOf;
if (size >= mMinRequiredSize)
{
mHoistedDeclNodes.push_back(&declNode);
return nullptr;
}
}
}
}
return {declNode, VisitBits::Neither};
}
bool rewrite(TIntermBlock &root, IdGen &idGen)
{
if (!rebuildRoot(root))
{
return false;
}
if (mHoistedDeclNodes.empty())
{
return true;
}
root.insertChildNodes(FindFirstFunctionDefinitionIndex(&root), mHoistedDeclNodes);
for (TIntermNode *opaqueDeclNode : mHoistedDeclNodes)
{
TIntermDeclaration *declNode = opaqueDeclNode->getAsDeclarationNode();
ASSERT(declNode);
const TVariable &oldVar = ViewDeclaration(*declNode).symbol.variable();
const Name newName = idGen.createNewName(oldVar.name());
auto *newVar = new TVariable(&mSymbolTable, newName.rawName(), &oldVar.getType(),
newName.symbolType());
if (!ReplaceVariable(&mCompiler, &root, &oldVar, newVar))
{
return false;
}
}
return true;
}
};
} // anonymous namespace
////////////////////////////////////////////////////////////////////////////////
bool sh::HoistConstants(TCompiler &compiler,
TIntermBlock &root,
IdGen &idGen,
size_t minRequiredSize)
{
return Rewriter(compiler, minRequiredSize).rewrite(root, idGen);
}

26
src/compiler/translator/TranslatorMetalDirect/HoistConstants.h Normal file
Просмотреть файл

@ -0,0 +1,26 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_HOISTCONSTANTS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_HOISTCONSTANTS_H_
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/IdGen.h"
namespace sh
{
// Hoists function-local constants to the global scope if their Metal sizeof meets
// `minRequiredSize`.
ANGLE_NO_DISCARD bool HoistConstants(TCompiler &compiler,
TIntermBlock &root,
IdGen &idGen,
size_t minRequiredSize);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_HOISTCONSTANTS_H_

97
src/compiler/translator/TranslatorMetalDirect/IdGen.cpp Normal file
Просмотреть файл

@ -0,0 +1,97 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <cctype>
#include <cstring>
#include <limits>
#include <unordered_map>
#include <unordered_set>
#include "compiler/translator/TranslatorMetalDirect/RewriteKeywords.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
IdGen::IdGen() {}
template <typename String, typename StringToImmutable>
Name IdGen::createNewName(size_t count,
const String *baseNames,
const StringToImmutable &toImmutable)
{
const unsigned id = mNext++;
char idBuffer[std::numeric_limits<unsigned>::digits10 + 1];
sprintf(idBuffer, "%u", id);
mNewNameBuffer.clear();
mNewNameBuffer += '_';
mNewNameBuffer += idBuffer;
// Note:
// Double underscores are only allowed in C++ (and thus Metal) vendor identifiers, so here we
// take care not to introduce any.
for (size_t i = 0; i < count; ++i)
{
const ImmutableString baseName = toImmutable(baseNames[i]);
if (!baseName.empty())
{
const char *base = baseName.data();
if (baseName.beginsWith(kAngleInternalPrefix))
{
base += sizeof(kAngleInternalPrefix) - 1;
}
if (*base == '_')
{
++base;
}
ASSERT(*base != '_');
if (mNewNameBuffer.back() != '_')
{
mNewNameBuffer += '_';
}
mNewNameBuffer += base;
}
}
return Name(ImmutableString(mNewNameBuffer), SymbolType::AngleInternal);
}
Name IdGen::createNewName(const ImmutableString &baseName)
{
return createNewName({baseName});
}
Name IdGen::createNewName(const Name &baseName)
{
return createNewName(baseName.rawName());
}
Name IdGen::createNewName(const char *baseName)
{
return createNewName(ImmutableString(baseName));
}
Name IdGen::createNewName(std::initializer_list<ImmutableString> baseNames)
{
return createNewName(baseNames.size(), baseNames.begin(),
[](const ImmutableString &s) { return s; });
}
Name IdGen::createNewName(std::initializer_list<Name> baseNames)
{
return createNewName(baseNames.size(), baseNames.begin(),
[](const Name &s) { return s.rawName(); });
}
Name IdGen::createNewName(std::initializer_list<const char *> baseNames)
{
return createNewName(baseNames.size(), baseNames.begin(),
[](const char *s) { return ImmutableString(s); });
}

41
src/compiler/translator/TranslatorMetalDirect/IdGen.h Normal file
Просмотреть файл

@ -0,0 +1,41 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_IDGEN_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_IDGEN_H_
#include "common/angleutils.h"
#include "compiler/translator/TranslatorMetalDirect/Name.h"
namespace sh
{
// For creating new fresh names.
// All names created are marked as SymbolType::AngleInternal.
class IdGen : angle::NonCopyable
{
public:
IdGen();
Name createNewName(const ImmutableString &baseName);
Name createNewName(const Name &baseName);
Name createNewName(const char *baseName);
Name createNewName(std::initializer_list<ImmutableString> baseNames);
Name createNewName(std::initializer_list<Name> baseNames);
Name createNewName(std::initializer_list<const char *> baseNames);
private:
template <typename String, typename StringToImmutable>
Name createNewName(size_t count, const String *baseNames, const StringToImmutable &toImmutable);
private:
unsigned mNext = 0; // `unsigned` because of "%u" use in sprintf
std::string mNewNameBuffer; // reusable buffer to avoid tons of reallocations
};
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_IDGEN_H_

1028
src/compiler/translator/TranslatorMetalDirect/IntermRebuild.cpp Normal file
Просмотреть файл

Разница между файлами не показана из-за своего большого размера Загрузить разницу

328
src/compiler/translator/TranslatorMetalDirect/IntermRebuild.h Normal file
Просмотреть файл

@ -0,0 +1,328 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_INTERMREBUILD_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_INTERMREBUILD_H_
#include "compiler/translator/TranslatorMetalDirect/NodeType.h"
#include "compiler/translator/tree_util/IntermTraverse.h"
namespace sh
{
// Walks the tree to rebuild nodes.
// This class is intended to be derived with overridden visitXXX functions.
//
// Each visitXXX function that does not have a Visit parameter simply has the visitor called
// exactly once, regardless of (preVisit) or (postVisit) values.
// Each visitXXX function that has a Visit parameter behaves as follows:
// * If (preVisit):
// - The node is visited before children are traversed.
// - The returned value is used to replace the visited node. The returned value may be the same
// as the original node.
// - If multiple nodes are returned, children and post visits of the returned nodes are not
// preformed, even if it is a singleton collection.
// * If (childVisit)
// - If any new children are returned, the node is automatically rebuilt with the new children
// before post visit.
// - Depending on the type of the node, null children may be discarded.
// - Ill-typed children cause rebuild errors. Ill-typed means the node to automatically rebuild
// cannot accept a child of a certain type as input to its constructor.
// - Only instances of TIntermAggregateBase can accept Multi results for any of its children.
// If supplied, the nodes are spliced children at the spot of the original child.
// * If (postVisit)
// - The node is visited after any children are traversed.
// - Only after such a rebuild (or lack thereof), the post-visit is performed.
//
// Nodes in visit functions are allowed to be modified in place, including TIntermAggregateBase
// child sequences.
//
// The default implementations of all the visitXXX functions support full pre and post traversal
// without modifying the visited nodes.
//
class TIntermRebuild : angle::NonCopyable
{
enum class Action
{
ReplaceSingle,
ReplaceMulti,
Drop,
Fail,
};
public:
struct Fail
{};
enum VisitBits : size_t
{
// No bits are set.
Empty = 0u,
// Allow visit of returned node's children.
Children = 1u << 0u,
// Allow post visit of returned node.
Post = 1u << 1u,
// If (Children) bit, only visit if the returned node is the same as the original node.
ChildrenRequiresSame = 1u << 2u,
// If (Post) bit, only visit if the returned node is the same as the original node.
PostRequiresSame = 1u << 3u,
RequireSame = ChildrenRequiresSame | PostRequiresSame,
Neither = Empty,
Both = Children | Post,
BothWhenSame = Both | RequireSame,
};
private:
struct NodeStackGuard;
template <typename T>
struct ConsList
{
T value;
ConsList<T> *tail;
};
class BaseResult
{
BaseResult(const BaseResult &) = delete;
BaseResult &operator=(const BaseResult &) = delete;
public:
BaseResult(BaseResult &&other) = default;
BaseResult(BaseResult &other); // For subclass move constructor impls
BaseResult(TIntermNode &node, VisitBits visit);
BaseResult(TIntermNode *node, VisitBits visit);
BaseResult(nullptr_t);
BaseResult(Fail);
BaseResult(std::vector<TIntermNode *> &&nodes);
void moveAssignImpl(BaseResult &other); // For subclass move assign impls
static BaseResult Multi(std::vector<TIntermNode *> &&nodes);
template <typename Iter>
static BaseResult Multi(Iter nodesBegin, Iter nodesEnd)
{
std::vector<TIntermNode *> nodes;
for (Iter nodesCurr = nodesBegin; nodesCurr != nodesEnd; ++nodesCurr)
{
nodes.push_back(*nodesCurr);
}
return std::move(nodes);
}
bool isFail() const;
bool isDrop() const;
TIntermNode *single() const;
const std::vector<TIntermNode *> *multi() const;
public:
Action mAction;
VisitBits mVisit;
TIntermNode *mSingle;
std::vector<TIntermNode *> mMulti;
};
public:
class PreResult : private BaseResult
{
friend class TIntermRebuild;
public:
PreResult(PreResult &&other);
PreResult(TIntermNode &node, VisitBits visit = VisitBits::BothWhenSame);
PreResult(TIntermNode *node, VisitBits visit = VisitBits::BothWhenSame);
PreResult(nullptr_t); // Used to drop a node.
PreResult(Fail); // Used to signal failure.
void operator=(PreResult &&other);
static PreResult Multi(std::vector<TIntermNode *> &&nodes)
{
return BaseResult::Multi(std::move(nodes));
}
template <typename Iter>
static PreResult Multi(Iter nodesBegin, Iter nodesEnd)
{
return BaseResult::Multi(nodesBegin, nodesEnd);
}
using BaseResult::isDrop;
using BaseResult::isFail;
using BaseResult::multi;
using BaseResult::single;
private:
PreResult(BaseResult &&other);
};
class PostResult : private BaseResult
{
friend class TIntermRebuild;
public:
PostResult(PostResult &&other);
PostResult(TIntermNode &node);
PostResult(TIntermNode *node);
PostResult(nullptr_t); // Used to drop a node
PostResult(Fail); // Used to signal failure.
void operator=(PostResult &&other);
static PostResult Multi(std::vector<TIntermNode *> &&nodes)
{
return BaseResult::Multi(std::move(nodes));
}
template <typename Iter>
static PostResult Multi(Iter nodesBegin, Iter nodesEnd)
{
return BaseResult::Multi(nodesBegin, nodesEnd);
}
using BaseResult::isDrop;
using BaseResult::isFail;
using BaseResult::multi;
using BaseResult::single;
private:
PostResult(BaseResult &&other);
};
public:
TIntermRebuild(TCompiler &compiler, bool preVisit, bool postVisit);
virtual ~TIntermRebuild();
// Rebuilds the tree starting at the provided root. If a new node would be returned for the
// root, the root node's children become that of the new node instead. Returns false if failure
// occurred.
ANGLE_NO_DISCARD bool rebuildRoot(TIntermBlock &root);
protected:
virtual PreResult visitSymbolPre(TIntermSymbol &node);
virtual PreResult visitConstantUnionPre(TIntermConstantUnion &node);
virtual PreResult visitFunctionPrototypePre(TIntermFunctionPrototype &node);
virtual PreResult visitPreprocessorDirectivePre(TIntermPreprocessorDirective &node);
virtual PreResult visitUnaryPre(TIntermUnary &node);
virtual PreResult visitBinaryPre(TIntermBinary &node);
virtual PreResult visitTernaryPre(TIntermTernary &node);
virtual PreResult visitSwizzlePre(TIntermSwizzle &node);
virtual PreResult visitIfElsePre(TIntermIfElse &node);
virtual PreResult visitSwitchPre(TIntermSwitch &node);
virtual PreResult visitCasePre(TIntermCase &node);
virtual PreResult visitLoopPre(TIntermLoop &node);
virtual PreResult visitBranchPre(TIntermBranch &node);
virtual PreResult visitDeclarationPre(TIntermDeclaration &node);
virtual PreResult visitBlockPre(TIntermBlock &node);
virtual PreResult visitAggregatePre(TIntermAggregate &node);
virtual PreResult visitFunctionDefinitionPre(TIntermFunctionDefinition &node);
virtual PreResult visitGlobalQualifierDeclarationPre(TIntermGlobalQualifierDeclaration &node);
virtual PostResult visitSymbolPost(TIntermSymbol &node);
virtual PostResult visitConstantUnionPost(TIntermConstantUnion &node);
virtual PostResult visitFunctionPrototypePost(TIntermFunctionPrototype &node);
virtual PostResult visitPreprocessorDirectivePost(TIntermPreprocessorDirective &node);
virtual PostResult visitUnaryPost(TIntermUnary &node);
virtual PostResult visitBinaryPost(TIntermBinary &node);
virtual PostResult visitTernaryPost(TIntermTernary &node);
virtual PostResult visitSwizzlePost(TIntermSwizzle &node);
virtual PostResult visitIfElsePost(TIntermIfElse &node);
virtual PostResult visitSwitchPost(TIntermSwitch &node);
virtual PostResult visitCasePost(TIntermCase &node);
virtual PostResult visitLoopPost(TIntermLoop &node);
virtual PostResult visitBranchPost(TIntermBranch &node);
virtual PostResult visitDeclarationPost(TIntermDeclaration &node);
virtual PostResult visitBlockPost(TIntermBlock &node);
virtual PostResult visitAggregatePost(TIntermAggregate &node);
virtual PostResult visitFunctionDefinitionPost(TIntermFunctionDefinition &node);
virtual PostResult visitGlobalQualifierDeclarationPost(TIntermGlobalQualifierDeclaration &node);
// Can be used to rebuild a specific node during a traversal. Useful for fine control of
// rebuilding a node's children.
ANGLE_NO_DISCARD PostResult rebuild(TIntermNode &node);
// Rebuilds the provided node in place. If a new node would be returned, the old node's children
// become that of the new node instead. Returns false if failure occurred.
ANGLE_NO_DISCARD bool rebuildInPlace(TIntermAggregate &node);
// Rebuilds the provided node in place. If a new node would be returned, the old node's children
// become that of the new node instead. Returns false if failure occurred.
ANGLE_NO_DISCARD bool rebuildInPlace(TIntermBlock &node);
// Rebuilds the provided node in place. If a new node would be returned, the old node's children
// become that of the new node instead. Returns false if failure occurred.
ANGLE_NO_DISCARD bool rebuildInPlace(TIntermDeclaration &node);
// If currently at or below a function declaration body, this returns the function that encloses
// the currently visited node. (This returns null if at a function declaration node.)
const TFunction *getParentFunction() const;
TIntermNode *getParentNode(size_t offset = 0) const;
private:
template <typename Node>
ANGLE_NO_DISCARD bool rebuildInPlaceImpl(Node &node);
PostResult traverseAny(TIntermNode &node);
template <typename Node>
Node *traverseAnyAs(TIntermNode &node);
template <typename Node>
bool traverseAnyAs(TIntermNode &node, Node *&out);
PreResult traversePre(TIntermNode &originalNode);
TIntermNode *traverseChildren(NodeType currNodeType,
const TIntermNode &originalNode,
TIntermNode &currNode,
VisitBits visit);
PostResult traversePost(NodeType nodeType,
const TIntermNode &originalNode,
TIntermNode &currNode,
VisitBits visit);
bool traverseAggregateBaseChildren(TIntermAggregateBase &node);
TIntermNode *traverseUnaryChildren(TIntermUnary &node);
TIntermNode *traverseBinaryChildren(TIntermBinary &node);
TIntermNode *traverseTernaryChildren(TIntermTernary &node);
TIntermNode *traverseSwizzleChildren(TIntermSwizzle &node);
TIntermNode *traverseIfElseChildren(TIntermIfElse &node);
TIntermNode *traverseSwitchChildren(TIntermSwitch &node);
TIntermNode *traverseCaseChildren(TIntermCase &node);
TIntermNode *traverseLoopChildren(TIntermLoop &node);
TIntermNode *traverseBranchChildren(TIntermBranch &node);
TIntermNode *traverseDeclarationChildren(TIntermDeclaration &node);
TIntermNode *traverseBlockChildren(TIntermBlock &node);
TIntermNode *traverseAggregateChildren(TIntermAggregate &node);
TIntermNode *traverseFunctionDefinitionChildren(TIntermFunctionDefinition &node);
TIntermNode *traverseGlobalQualifierDeclarationChildren(
TIntermGlobalQualifierDeclaration &node);
protected:
TCompiler &mCompiler;
TSymbolTable &mSymbolTable;
const TFunction *mParentFunc = nullptr;
GetNodeType getNodeType;
private:
ConsList<TIntermNode *> mNodeStack{nullptr, nullptr};
bool mPreVisit;
bool mPostVisit;
};
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_INTERMREBUILD_H_

89
src/compiler/translator/TranslatorMetalDirect/IntroduceVertexIndexID.cpp Normal file
Просмотреть файл

@ -0,0 +1,89 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/IntroduceVertexIndexID.h"
#include "compiler/translator/StaticType.h"
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
#include "compiler/translator/tree_util/BuiltIn.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
constexpr const TVariable kgl_VertexIDMetal(BuiltInId::gl_VertexID,
ImmutableString("gl_VertexID"),
SymbolType::BuiltIn,
TExtension::UNDEFINED,
StaticType::Get<EbtUInt, EbpHigh, EvqVertexID, 1, 1>());
constexpr const TVariable kgl_instanceIdMetal(
BuiltInId::gl_VertexID,
ImmutableString("instanceIdMod"),
SymbolType::AngleInternal,
TExtension::UNDEFINED,
StaticType::Get<EbtUInt, EbpHigh, EvqInstanceID, 1, 1>());
class Rewriter : public TIntermRebuild
{
public:
Rewriter(TCompiler &compiler) : TIntermRebuild(compiler, true, true) {}
private:
PreResult visitFunctionDefinitionPre(TIntermFunctionDefinition &node) override
{
if (node.getFunction()->isMain())
{
const TFunction *mainFunction = node.getFunction();
bool needsVertexId = true;
bool needsInstanceId = true;
std::vector<const TVariable *> mVariablesToIntroduce;
for (size_t i = 0; i < mainFunction->getParamCount(); ++i)
{
const TVariable *param = mainFunction->getParam(i);
Name instanceIDName =
Pipeline{Pipeline::Type::InstanceId, nullptr}.getStructInstanceName(
Pipeline::Variant::Modified);
if (Name(*param) == instanceIDName)
{
needsInstanceId = false;
}
else if (param->getType().getQualifier() == TQualifier::EvqVertexID)
{
needsVertexId = false;
}
}
if (needsInstanceId)
{
mVariablesToIntroduce.push_back(&kgl_instanceIdMetal);
}
if (needsVertexId)
{
mVariablesToIntroduce.push_back(&kgl_VertexIDMetal);
}
const TFunction &newFunction = CloneFunctionAndAppendParams(
mSymbolTable, nullptr, *node.getFunction(), mVariablesToIntroduce);
TIntermFunctionPrototype *newProto = new TIntermFunctionPrototype(&newFunction);
return new TIntermFunctionDefinition(newProto, node.getBody());
}
return node;
}
};
} // anonymous namespace
////////////////////////////////////////////////////////////////////////////////
bool sh::IntroduceVertexAndInstanceIndex(TCompiler &compiler, TIntermBlock &root)
{
if (!Rewriter(compiler).rebuildRoot(root))
{
return false;
}
return true;
}

20
src/compiler/translator/TranslatorMetalDirect/IntroduceVertexIndexID.h Normal file
Просмотреть файл

@ -0,0 +1,20 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_INTRODUCEVERTEXINDEX_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_INTRODUCEVERTEXINDEX_H_
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/Pipeline.h"
namespace sh
{
ANGLE_NO_DISCARD bool IntroduceVertexAndInstanceIndex(TCompiler &compiler, TIntermBlock &root);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_INTRODUCEVERTEXINDEX_H_

401
src/compiler/translator/TranslatorMetalDirect/Layout.cpp Normal file
Просмотреть файл

@ -0,0 +1,401 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <algorithm>
#include <cctype>
#include <cstring>
#include <limits>
#include <unordered_map>
#include <unordered_set>
#include "compiler/translator/Symbol.h"
#include "compiler/translator/TranslatorMetalDirect/Layout.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
enum class Language
{
Metal,
GLSL,
};
static size_t RoundUpToMultipleOf(size_t x, size_t multiple)
{
const size_t rem = x % multiple;
return rem == 0 ? x : x + (multiple - rem);
}
void Layout::requireAlignment(size_t align, bool pad)
{
alignOf = std::max(alignOf, align);
if (pad)
{
sizeOf = RoundUpToMultipleOf(sizeOf, align);
}
}
bool Layout::operator==(const Layout &other) const
{
return sizeOf == other.sizeOf && alignOf == other.alignOf;
}
void Layout::operator+=(const Layout &other)
{
requireAlignment(other.alignOf, true);
sizeOf += other.sizeOf;
}
void Layout::operator*=(size_t scale)
{
sizeOf *= scale;
}
Layout Layout::operator+(const Layout &other) const
{
auto self = *this;
self += other;
return self;
}
Layout Layout::operator*(size_t scale) const
{
auto self = *this;
self *= scale;
return self;
}
static Layout ScalarLayoutOf(const TType &type, Language language)
{
const TBasicType basicType = type.getBasicType();
switch (basicType)
{
case TBasicType::EbtBool:
return {1, 1};
case TBasicType::EbtInt:
case TBasicType::EbtUInt:
case TBasicType::EbtFloat:
return {4, 4};
default:
if (IsSampler(basicType))
{
switch (language)
{
case Language::Metal:
return {8, 8};
case Language::GLSL:
return {4, 4};
}
}
UNIMPLEMENTED();
return Layout::Invalid();
}
}
static const size_t innerScalesPacked[] = {0, 1, 2, 3, 4};
static const size_t innerScalesUnpacked[] = {0, 1, 2, 4, 4};
Layout sh::MetalLayoutOf(const TType &type, MetalLayoutOfConfig config)
{
ASSERT(type.getNominalSize() <= 4);
ASSERT(type.getSecondarySize() <= 4);
const TLayoutBlockStorage storage = type.getLayoutQualifier().blockStorage;
const bool isPacked = !config.disablePacking && (storage == TLayoutBlockStorage::EbsPacked ||
storage == TLayoutBlockStorage::EbsShared);
if (type.isArray() && !config.maskArray)
{
config.maskArray = true;
const Layout layout = MetalLayoutOf(type, config);
config.maskArray = false;
const size_t vol = type.getArraySizeProduct();
return layout * vol;
}
if (const TStructure *structure = type.getStruct())
{
ASSERT(type.getNominalSize() == 1);
ASSERT(type.getSecondarySize() == 1);
auto config2 = config;
config2.maskArray = false;
auto layout = Layout::Identity();
const TFieldList &fields = structure->fields();
for (const TField *field : fields)
{
layout += MetalLayoutOf(*field->type(), config2);
}
if (config.assumeStructsAreTailPadded)
{
size_t pad =
(kDefaultStructAlignmentSize - layout.sizeOf) % kDefaultStructAlignmentSize;
layout.sizeOf += pad;
}
layout.sizeOf = RoundUpToMultipleOf(layout.sizeOf, layout.alignOf);
return layout;
}
if (config.treatSamplersAsTextureEnv && IsSampler(type.getBasicType()))
{
return {16, 8}; // pointer{8,8} and metal::sampler{8,8}
}
const Layout scalarLayout = ScalarLayoutOf(type, Language::Metal);
if (type.isRank0())
{
return scalarLayout;
}
else if (type.isVector())
{
if (isPacked)
{
const size_t innerScale = innerScalesPacked[type.getNominalSize()];
auto layout = Layout{scalarLayout.sizeOf * innerScale, scalarLayout.alignOf};
return layout;
}
else
{
const size_t innerScale = innerScalesUnpacked[type.getNominalSize()];
auto layout = Layout::Both(scalarLayout.sizeOf * innerScale);
return layout;
}
}
else
{
ASSERT(type.isMatrix());
ASSERT(type.getBasicType() == TBasicType::EbtFloat);
// typeCxR <=> typeR[C]
const size_t innerScale = innerScalesUnpacked[type.getRows()];
const size_t outerScale = static_cast<size_t>(type.getCols());
const size_t n = scalarLayout.sizeOf * innerScale;
return {n * outerScale, n};
}
}
TLayoutBlockStorage sh::Overlay(TLayoutBlockStorage oldStorage, const TType &type)
{
const TLayoutBlockStorage newStorage = type.getLayoutQualifier().blockStorage;
switch (newStorage)
{
case TLayoutBlockStorage::EbsUnspecified:
return oldStorage == TLayoutBlockStorage::EbsUnspecified ? kDefaultLayoutBlockStorage
: oldStorage;
default:
return newStorage;
}
}
TLayoutMatrixPacking sh::Overlay(TLayoutMatrixPacking oldPacking, const TType &type)
{
const TLayoutMatrixPacking newPacking = type.getLayoutQualifier().matrixPacking;
switch (newPacking)
{
case TLayoutMatrixPacking::EmpUnspecified:
return oldPacking == TLayoutMatrixPacking::EmpUnspecified ? kDefaultLayoutMatrixPacking
: oldPacking;
default:
return newPacking;
}
}
bool sh::CanBePacked(TLayoutBlockStorage storage)
{
switch (storage)
{
case TLayoutBlockStorage::EbsPacked:
case TLayoutBlockStorage::EbsShared:
return true;
case TLayoutBlockStorage::EbsStd140:
case TLayoutBlockStorage::EbsStd430:
return false;
case TLayoutBlockStorage::EbsUnspecified:
UNREACHABLE();
return false;
}
}
bool sh::CanBePacked(TLayoutQualifier layoutQualifier)
{
return CanBePacked(layoutQualifier.blockStorage);
}
bool sh::CanBePacked(const TType &type)
{
if (!type.isVector())
{
return false;
}
return CanBePacked(type.getLayoutQualifier());
}
void sh::SetBlockStorage(TType &type, TLayoutBlockStorage storage)
{
auto qual = type.getLayoutQualifier();
qual.blockStorage = storage;
type.setLayoutQualifier(qual);
}
static Layout CommonGlslStructLayoutOf(TField const *const *begin,
TField const *const *end,
const TLayoutBlockStorage storage,
const TLayoutMatrixPacking matrixPacking,
const bool maskArray,
const size_t baseAlignment)
{
const bool isPacked =
storage == TLayoutBlockStorage::EbsPacked || storage == TLayoutBlockStorage::EbsShared;
auto layout = Layout::Identity();
for (auto iter = begin; iter != end; ++iter)
{
layout += GlslLayoutOf(*(*iter)->type(), storage, matrixPacking, false);
}
if (!isPacked) // XXX: Correct?
{
layout.sizeOf = RoundUpToMultipleOf(layout.sizeOf, layout.alignOf);
}
layout.requireAlignment(baseAlignment, true);
return layout;
}
static Layout CommonGlslLayoutOf(const TType &type,
const TLayoutBlockStorage storage,
const TLayoutMatrixPacking matrixPacking,
const bool maskArray,
const size_t baseAlignment)
{
ASSERT(storage != TLayoutBlockStorage::EbsUnspecified);
const bool isPacked =
storage == TLayoutBlockStorage::EbsPacked || storage == TLayoutBlockStorage::EbsShared;
if (type.isArray() && !type.isMatrix() && !maskArray)
{
auto layout = GlslLayoutOf(type, storage, matrixPacking, true);
layout *= type.getArraySizeProduct();
layout.requireAlignment(baseAlignment, true);
return layout;
}
if (const TStructure *structure = type.getStruct())
{
ASSERT(type.getNominalSize() == 1);
ASSERT(type.getSecondarySize() == 1);
const TFieldList &fields = structure->fields();
return CommonGlslStructLayoutOf(fields.data(), fields.data() + fields.size(), storage,
matrixPacking, maskArray, baseAlignment);
}
const auto scalarLayout = ScalarLayoutOf(type, Language::GLSL);
if (type.isRank0())
{
return scalarLayout;
}
else if (type.isVector())
{
if (isPacked)
{
const size_t sizeScale = innerScalesPacked[type.getNominalSize()];
const size_t alignScale = innerScalesUnpacked[type.getNominalSize()];
auto layout =
Layout{scalarLayout.sizeOf * sizeScale, scalarLayout.alignOf * alignScale};
return layout;
}
else
{
const size_t innerScale = innerScalesUnpacked[type.getNominalSize()];
auto layout = Layout::Both(scalarLayout.sizeOf * innerScale);
return layout;
}
}
else
{
ASSERT(type.isMatrix());
size_t innerDim;
size_t outerDim;
switch (matrixPacking)
{
case TLayoutMatrixPacking::EmpColumnMajor:
innerDim = static_cast<size_t>(type.getRows());
outerDim = static_cast<size_t>(type.getCols());
break;
case TLayoutMatrixPacking::EmpRowMajor:
innerDim = static_cast<size_t>(type.getCols());
outerDim = static_cast<size_t>(type.getRows());
break;
case TLayoutMatrixPacking::EmpUnspecified:
UNREACHABLE();
innerDim = 0;
outerDim = 0;
}
outerDim *= type.getArraySizeProduct();
const size_t innerScale = innerScalesUnpacked[innerDim];
const size_t n = innerScale * scalarLayout.sizeOf;
Layout layout = {outerDim * n, n};
layout.requireAlignment(baseAlignment, true);
return layout;
}
}
Layout sh::GlslLayoutOf(const TType &type,
TLayoutBlockStorage storage,
TLayoutMatrixPacking matrixPacking,
bool maskArray)
{
ASSERT(type.getNominalSize() <= 4);
ASSERT(type.getSecondarySize() <= 4);
storage = Overlay(storage, type);
matrixPacking = Overlay(matrixPacking, type);
switch (storage)
{
case TLayoutBlockStorage::EbsPacked:
return CommonGlslLayoutOf(type, storage, matrixPacking, maskArray, 1);
case TLayoutBlockStorage::EbsShared:
return CommonGlslLayoutOf(type, storage, matrixPacking, maskArray, 16);
case TLayoutBlockStorage::EbsStd140:
return CommonGlslLayoutOf(type, storage, matrixPacking, maskArray, 16);
case TLayoutBlockStorage::EbsStd430:
return CommonGlslLayoutOf(type, storage, matrixPacking, maskArray, 1);
case TLayoutBlockStorage::EbsUnspecified:
UNREACHABLE();
return Layout::Invalid();
}
}
ANGLE_NO_DISCARD Layout sh::GlslStructLayoutOf(TField const *const *begin,
TField const *const *end,
TLayoutBlockStorage storage,
TLayoutMatrixPacking matrixPacking,
bool maskArray)
{
ASSERT(storage != TLayoutBlockStorage::EbsUnspecified);
ASSERT(matrixPacking != TLayoutMatrixPacking::EmpUnspecified);
switch (storage)
{
case TLayoutBlockStorage::EbsPacked:
return CommonGlslStructLayoutOf(begin, end, storage, matrixPacking, maskArray, 1);
case TLayoutBlockStorage::EbsShared:
return CommonGlslStructLayoutOf(begin, end, storage, matrixPacking, maskArray, 16);
case TLayoutBlockStorage::EbsStd140:
return CommonGlslStructLayoutOf(begin, end, storage, matrixPacking, maskArray, 16);
case TLayoutBlockStorage::EbsStd430:
return CommonGlslStructLayoutOf(begin, end, storage, matrixPacking, maskArray, 1);
case TLayoutBlockStorage::EbsUnspecified:
UNREACHABLE();
return Layout::Invalid();
}
}

89
src/compiler/translator/TranslatorMetalDirect/Layout.h Normal file
Просмотреть файл

@ -0,0 +1,89 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_LAYOUT_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_LAYOUT_H_
#include "common/angleutils.h"
#include "compiler/translator/Types.h"
namespace sh
{
constexpr const auto kDefaultLayoutBlockStorage = TLayoutBlockStorage::EbsShared;
constexpr const auto kDefaultLayoutMatrixPacking = TLayoutMatrixPacking::EmpColumnMajor;
constexpr const auto kDefaultStructAlignmentSize = 16;
// Returns `oldStorage` if `type` has unspecified block storage.
// Otherwise returns block storage of `type`.
TLayoutBlockStorage Overlay(TLayoutBlockStorage oldStorage, const TType &type);
// Returns `oldPacking` if `type` has unspecified matrix packing.
// Otherwise returns matrix packing of `type`.
TLayoutMatrixPacking Overlay(TLayoutMatrixPacking oldPacking, const TType &type);
// Returns whether or not it is permissable for the block storage to use a packed representation.
bool CanBePacked(TLayoutBlockStorage storage);
// Returns whether or not it is permissable for the layout qualifier to use a packed representation.
bool CanBePacked(TLayoutQualifier layoutQualifier);
// Returns whether or not it is permissable for the type to use a packed representation.
bool CanBePacked(const TType &type);
// Sets the block storage for a type.
void SetBlockStorage(TType &type, TLayoutBlockStorage storage);
// Contains `sizeof` and `alignof` information.
struct Layout
{
size_t sizeOf;
size_t alignOf;
static Layout Identity() { return {0, 1}; }
static Layout Invalid() { return {0, 0}; }
static Layout Both(size_t n) { return {n, n}; }
void requireAlignment(size_t align, bool pad);
bool operator==(const Layout &other) const;
void operator+=(const Layout &other);
void operator*=(size_t scale);
Layout operator+(const Layout &other) const;
Layout operator*(size_t scale) const;
};
struct MetalLayoutOfConfig
{
bool disablePacking = false;
bool maskArray = false;
bool treatSamplersAsTextureEnv = false;
bool assumeStructsAreTailPadded = false; // Pad to multiple of 16
};
// Returns the layout of a type if it were to be represented in a Metal program.
// This deliberately ignores the TLayoutBlockStorage and TLayoutMatrixPacking of any type.
ANGLE_NO_DISCARD Layout MetalLayoutOf(const TType &type, MetalLayoutOfConfig config = {});
// Returns the layout of a type if it were to be represented in a GLSL program.
ANGLE_NO_DISCARD Layout
GlslLayoutOf(const TType &type,
TLayoutBlockStorage storage = TLayoutBlockStorage::EbsUnspecified,
TLayoutMatrixPacking matrixPacking = TLayoutMatrixPacking::EmpUnspecified,
bool maskArray = false);
// Returns the layout of a structure if it were to be represented in a GLSL program.
ANGLE_NO_DISCARD Layout GlslStructLayoutOf(TField const *const *begin,
TField const *const *end,
TLayoutBlockStorage storage,
TLayoutMatrixPacking matrixPacking,
bool maskArray = false);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_LAYOUT_H_

34
src/compiler/translator/TranslatorMetalDirect/MapFunctionsToDefinitions.cpp Normal file
Просмотреть файл

@ -0,0 +1,34 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/MapFunctionsToDefinitions.h"
#include "compiler/translator/Symbol.h"
using namespace sh;
class Mapper : public TIntermTraverser
{
public:
FunctionToDefinition mFuncToDef;
public:
Mapper() : TIntermTraverser(true, false, false) {}
bool visitFunctionDefinition(Visit, TIntermFunctionDefinition *funcDefNode) override
{
const TFunction *func = funcDefNode->getFunction();
ASSERT(func->getBuiltInOp() == TOperator::EOpNull);
mFuncToDef[func] = funcDefNode;
return false;
}
};
FunctionToDefinition sh::MapFunctionsToDefinitions(TIntermBlock &root)
{
Mapper mapper;
root.traverse(&mapper);
return std::move(mapper.mFuncToDef);
}

25
src/compiler/translator/TranslatorMetalDirect/MapFunctionsToDefinitions.h Normal file
Просмотреть файл

@ -0,0 +1,25 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_MAPFUNCTIONSTODEFINITIONS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_MAPFUNCTIONSTODEFINITIONS_H_
#include <unordered_map>
#include "compiler/translator/tree_util/IntermTraverse.h"
namespace sh
{
// A map from functions to their corresponding definitions.
using FunctionToDefinition = std::unordered_map<const TFunction *, TIntermFunctionDefinition *>;
// Maps functions to their corresponding definitions.
ANGLE_NO_DISCARD FunctionToDefinition MapFunctionsToDefinitions(TIntermBlock &root);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_MAPFUNCTIONSTODEFINITIONS_H_

46
src/compiler/translator/TranslatorMetalDirect/MapSymbols.cpp Normal file
Просмотреть файл

@ -0,0 +1,46 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/MapSymbols.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
class Rewriter : public TIntermRebuild
{
private:
std::function<TIntermNode &(const TFunction *, TIntermSymbol &)> mMap;
public:
Rewriter(TCompiler &compiler,
std::function<TIntermNode &(const TFunction *, TIntermSymbol &)> map)
: TIntermRebuild(compiler, false, true), mMap(map)
{}
PostResult visitSymbolPost(TIntermSymbol &symbolNode) override
{
return mMap(getParentFunction(), symbolNode);
}
};
} // namespace
bool sh::MapSymbols(TCompiler &compiler,
TIntermBlock &root,
std::function<TIntermNode &(const TFunction *, TIntermSymbol &)> map)
{
Rewriter rewriter(compiler, std::move(map));
if (!rewriter.rebuildRoot(root))
{
return false;
}
return true;
}

27
src/compiler/translator/TranslatorMetalDirect/MapSymbols.h Normal file
Просмотреть файл

@ -0,0 +1,27 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_MAPVARIABLESTOMEMBERACCESS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_MAPVARIABLESTOMEMBERACCESS_H_
#include <functional>
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
namespace sh
{
// Maps TIntermSymbol nodes to TIntermNode nodes.
// The parent function of a symbol is provided to the mapping when applicable.
ANGLE_NO_DISCARD bool MapSymbols(
TCompiler &compiler,
TIntermBlock &root,
std::function<TIntermNode &(const TFunction *, TIntermSymbol &)> map);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_MAPVARIABLESTOMEMBERACCESS_H_

1059
src/compiler/translator/TranslatorMetalDirect/ModifyStruct.cpp Normal file
Просмотреть файл

Разница между файлами не показана из-за своего большого размера Загрузить разницу

139
src/compiler/translator/TranslatorMetalDirect/ModifyStruct.h Normal file
Просмотреть файл

@ -0,0 +1,139 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_MODIFYSTRUCT_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_MODIFYSTRUCT_H_
#include <cstring>
#include <unordered_map>
#include <unordered_set>
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/IdGen.h"
#include "compiler/translator/TranslatorMetalDirect/Layout.h"
#include "compiler/translator/TranslatorMetalDirect/Name.h"
#include "compiler/translator/TranslatorMetalDirect/ProgramPrelude.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
namespace sh
{
enum class ConvertType
{
OriginalToModified,
ModifiedToOriginal,
};
// Configures how struct modification is performed.
class ModifyStructConfig
{
public:
struct Predicate
{
using Func = bool (*)(const TField &);
static bool False(const TField &) { return false; }
static bool True(const TField &) { return true; }
};
struct SaturateVector
{
// Valid return values are [0, 1, 2, 3, 4].
// If original dim >= return value, the field remains untouched.
using Func = int (*)(const TField &);
static int DontSaturate(const TField &) { return 0; }
static int FullySaturate(const TField &) { return 4; }
};
public:
ModifyStructConfig(ConvertType convertType, bool allowPacking, bool allowPadding)
: convertType(convertType), allowPacking(allowPacking), allowPadding(allowPadding)
{}
// Matrix field is split into multiple fields of row vectors.
Predicate::Func splitMatrixColumns = Predicate::False;
// Array fields are split into multiple fields of element type.
Predicate::Func inlineArray = Predicate::False;
// Struct fields have their subfields inlined directly.
Predicate::Func inlineStruct = Predicate::False;
// Struct fields are modified.
Predicate::Func recurseStruct = Predicate::False;
// Vector and scalar bool fields are promoted to uint fields.
Predicate::Func promoteBoolToUint = Predicate::False;
// Creates a new structure where scalar or vector fields are saturated vectors.
// e.g. `float -> float4`.
// e.g. `float2 -> float4`.
SaturateVector::Func saturateScalarOrVector = SaturateVector::DontSaturate;
// Creates a new structure where scalar or vector array fields are saturated.
// e.g. `float[10] -> float4[10]`
// e.g. `float2[10] -> float4[10]`
SaturateVector::Func saturateScalarOrVectorArrays = SaturateVector::DontSaturate;
// Creates a new structure where matrix fields are row-saturated.
// e.g. `float2x2 -> float2x4`.
SaturateVector::Func saturateMatrixRows = SaturateVector::DontSaturate;
TLayoutBlockStorage initialBlockStorage = kDefaultLayoutBlockStorage;
TLayoutMatrixPacking initialMatrixPacking = kDefaultLayoutMatrixPacking;
ConvertType convertType;
bool allowPacking;
bool allowPadding;
AddressSpace externalAddressSpace;
};
struct ModifiedStructMachinery
{
const TStructure *modifiedStruct = nullptr;
TIntermFunctionDefinition *funcOriginalToModified = nullptr;
TIntermFunctionDefinition *funcModifiedToOriginal = nullptr;
TIntermFunctionDefinition *&getConverter(ConvertType convertType)
{
if (convertType == ConvertType::OriginalToModified)
{
return funcOriginalToModified;
}
else
{
return funcModifiedToOriginal;
}
}
};
// Indexed by topological order.
class ModifiedStructMachineries
{
public:
size_t size() const;
const ModifiedStructMachinery &at(size_t index) const;
const ModifiedStructMachinery *find(const TStructure &s) const;
void insert(const TStructure &s, const ModifiedStructMachinery &machinery);
private:
std::unordered_map<const TStructure *, ModifiedStructMachinery> originalToMachinery;
std::vector<const TStructure *> ordering;
};
// Returns true and `outMachinery` populated if modifications were performed.
// Returns false otherwise.
bool TryCreateModifiedStruct(TCompiler &compiler,
SymbolEnv &symbolEnv,
IdGen &idGen,
const ModifyStructConfig &config,
const TStructure &originalStruct,
const Name &modifiedStructName,
ModifiedStructMachineries &outMachineries,
const bool isUBO,
const bool allowPadding);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_MODIFYSTRUCT_H_

219
src/compiler/translator/TranslatorMetalDirect/Name.cpp Normal file
Просмотреть файл

@ -0,0 +1,219 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/Name.h"
#include "common/debug.h"
#include "compiler/translator/tree_util/IntermTraverse.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
template <typename T>
static ImmutableString GetName(T const &object)
{
if (object.symbolType() == SymbolType::Empty)
{
return kEmptyImmutableString;
}
return object.name();
}
Name::Name(const TField &field) : Name(GetName(field), field.symbolType()) {}
Name::Name(const TSymbol &symbol) : Name(GetName(symbol), symbol.symbolType()) {}
bool Name::operator==(const Name &other) const
{
return mRawName == other.mRawName && mSymbolType == other.mSymbolType;
}
bool Name::operator!=(const Name &other) const
{
return !(*this == other);
}
bool Name::operator<(const Name &other) const
{
if (mRawName < other.mRawName)
{
return true;
}
if (other.mRawName < mRawName)
{
return false;
}
return mSymbolType < other.mSymbolType;
}
bool Name::empty() const
{
return mSymbolType == SymbolType::Empty;
}
bool Name::beginsWith(const Name &prefix) const
{
if (mSymbolType != prefix.mSymbolType)
{
return false;
}
return mRawName.beginsWith(prefix.mRawName);
}
void Name::emit(TInfoSinkBase &out) const
{
switch (mSymbolType)
{
case SymbolType::BuiltIn:
case SymbolType::UserDefined:
ASSERT(!mRawName.empty());
out << mRawName;
break;
case SymbolType::AngleInternal:
ASSERT(!mRawName.empty());
if (mRawName.beginsWith(kAngleInternalPrefix))
{
out << mRawName;
}
else if (mRawName[0] != '_')
{
out << kAngleInternalPrefix << '_' << mRawName;
}
else
{
out << kAngleInternalPrefix << mRawName;
}
break;
case SymbolType::Empty:
UNREACHABLE();
break;
}
}
////////////////////////////////////////////////////////////////////////////////
namespace
{
// NOTE: This matches more things than FindSymbolNode.
class ExpressionContainsNameVisitor : public TIntermTraverser
{
Name mName;
bool mFoundName = false;
public:
ExpressionContainsNameVisitor(const Name &name)
: TIntermTraverser(true, false, false), mName(name)
{}
bool foundName() const { return mFoundName; }
void visitSymbol(TIntermSymbol *node) override
{
if (Name(node->variable()) == mName)
{
mFoundName = true;
}
}
bool visitSwizzle(Visit, TIntermSwizzle *) override { return !mFoundName; }
bool visitBinary(Visit visit, TIntermBinary *node) override { return !mFoundName; }
bool visitUnary(Visit visit, TIntermUnary *node) override { return !mFoundName; }
bool visitTernary(Visit visit, TIntermTernary *node) override { return !mFoundName; }
bool visitAggregate(Visit visit, TIntermAggregate *node) override
{
if (node->isConstructor())
{
const TType &type = node->getType();
const TStructure *structure = type.getStruct();
if (structure && Name(*structure) == mName)
{
mFoundName = true;
}
}
else
{
const TFunction *func = node->getFunction();
if (func && Name(*func) == mName)
{
mFoundName = true;
}
}
return !mFoundName;
}
bool visitIfElse(Visit visit, TIntermIfElse *node) override
{
UNREACHABLE();
return false;
}
bool visitSwitch(Visit, TIntermSwitch *) override
{
UNREACHABLE();
return false;
}
bool visitCase(Visit, TIntermCase *) override
{
UNREACHABLE();
return false;
}
void visitFunctionPrototype(TIntermFunctionPrototype *) override { UNREACHABLE(); }
bool visitFunctionDefinition(Visit, TIntermFunctionDefinition *) override
{
UNREACHABLE();
return false;
}
bool visitBlock(Visit, TIntermBlock *) override
{
UNREACHABLE();
return false;
}
bool visitGlobalQualifierDeclaration(Visit, TIntermGlobalQualifierDeclaration *) override
{
UNREACHABLE();
return false;
}
bool visitDeclaration(Visit, TIntermDeclaration *) override
{
UNREACHABLE();
return false;
}
bool visitLoop(Visit, TIntermLoop *) override
{
UNREACHABLE();
return false;
}
bool visitBranch(Visit, TIntermBranch *) override
{
UNREACHABLE();
return false;
}
void visitPreprocessorDirective(TIntermPreprocessorDirective *) override { UNREACHABLE(); }
};
} // anonymous namespace
bool sh::ExpressionContainsName(const Name &name, TIntermTyped &node)
{
ExpressionContainsNameVisitor visitor(name);
node.traverse(&visitor);
return visitor.foundName();
}

68
src/compiler/translator/TranslatorMetalDirect/Name.h Normal file
Просмотреть файл

@ -0,0 +1,68 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_NAME_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_NAME_H_
#include "compiler/translator/ImmutableString.h"
#include "compiler/translator/InfoSink.h"
#include "compiler/translator/IntermNode.h"
#include "compiler/translator/Symbol.h"
#include "compiler/translator/Types.h"
namespace sh
{
constexpr char kAngleInternalPrefix[] = "ANGLE";
// Represents the name of a symbol.
class Name
{
public:
constexpr Name(const Name &) = default;
constexpr Name() : Name(kEmptyImmutableString, SymbolType::Empty) {}
explicit constexpr Name(ImmutableString rawName, SymbolType symbolType)
: mRawName(rawName), mSymbolType(symbolType)
{
ASSERT(rawName.empty() == (symbolType == SymbolType::Empty));
}
explicit constexpr Name(const char *rawName, SymbolType symbolType = SymbolType::AngleInternal)
: Name(ImmutableString(rawName), symbolType)
{}
explicit Name(const std::string &rawName, SymbolType symbolType)
: Name(ImmutableString(rawName), symbolType)
{}
explicit Name(const TField &field);
explicit Name(const TSymbol &symbol);
Name &operator=(const Name &) = default;
bool operator==(const Name &other) const;
bool operator!=(const Name &other) const;
bool operator<(const Name &other) const;
constexpr const ImmutableString &rawName() const { return mRawName; }
constexpr SymbolType symbolType() const { return mSymbolType; }
bool empty() const;
bool beginsWith(const Name &prefix) const;
void emit(TInfoSinkBase &out) const;
private:
ImmutableString mRawName;
SymbolType mSymbolType;
};
ANGLE_NO_DISCARD bool ExpressionContainsName(const Name &name, TIntermTyped &node);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_NAME_H_

115
src/compiler/translator/TranslatorMetalDirect/NameEmbeddedUniformStructsMetal.cpp Normal file
Просмотреть файл

@ -0,0 +1,115 @@
//
// Copyright 2018 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
// NameEmbeddedUniformStructsMetal: Gives nameless uniform struct internal names.
//
#include "compiler/translator/TranslatorMetalDirect/NameEmbeddedUniformStructsMetal.h"
#include "compiler/translator/SymbolTable.h"
#include "compiler/translator/tree_util/IntermTraverse.h"
namespace sh
{
namespace
{
// This traverser translates embedded uniform structs into a specifier and declaration.
// This makes the declarations easier to move into uniform blocks.
class Traverser : public TIntermTraverser
{
public:
std::unordered_map<int, TIntermSymbol *> replacements;
explicit Traverser(TSymbolTable *symbolTable)
: TIntermTraverser(true, false, false, symbolTable)
{}
bool visitDeclaration(Visit visit, TIntermDeclaration *decl) override
{
ASSERT(visit == PreVisit);
if (!mInGlobalScope)
{
return false;
}
const TIntermSequence &sequence = *(decl->getSequence());
ASSERT(sequence.size() == 1);
TIntermTyped *declarator = sequence.front()->getAsTyped();
const TType &type = declarator->getType();
if (type.isStructSpecifier() && type.getQualifier() == EvqUniform)
{
const TStructure *structure = type.getStruct();
if (structure->symbolType() == SymbolType::Empty)
{
doReplacement(decl, declarator, structure);
}
}
return false;
}
void visitSymbol(TIntermSymbol *decl) override
{
auto symbol = replacements.find(decl->uniqueId().get());
if (symbol != replacements.end())
{
queueReplacement(symbol->second->deepCopy(), OriginalNode::IS_DROPPED);
}
}
private:
void doReplacement(TIntermDeclaration *decl,
TIntermTyped *declarator,
const TStructure *oldStructure)
{
// struct <structName> { ... };
TStructure *structure = new TStructure(mSymbolTable, kEmptyImmutableString,
&oldStructure->fields(), SymbolType::AngleInternal);
TType *namedType = new TType(structure, true);
namedType->setQualifier(EvqGlobal);
TVariable *structVariable =
new TVariable(mSymbolTable, kEmptyImmutableString, namedType, SymbolType::Empty);
TIntermSymbol *structDeclarator = new TIntermSymbol(structVariable);
TIntermDeclaration *structDeclaration = new TIntermDeclaration;
structDeclaration->appendDeclarator(structDeclarator);
TIntermSequence *newSequence = new TIntermSequence;
newSequence->push_back(structDeclaration);
// uniform <structName> <structUniformName>;
TIntermSymbol *asSymbol = declarator->getAsSymbolNode();
if (asSymbol && asSymbol->variable().symbolType() != SymbolType::Empty)
{
TIntermDeclaration *namedDecl = new TIntermDeclaration;
TType *uniformType = new TType(structure, false);
uniformType->setQualifier(EvqUniform);
TVariable *newVar = new TVariable(mSymbolTable, asSymbol->getName(), uniformType,
asSymbol->variable().symbolType());
TIntermSymbol *newSymbol = new TIntermSymbol(newVar);
replacements[asSymbol->uniqueId().get()] = newSymbol;
namedDecl->appendDeclarator(newSymbol);
newSequence->push_back(namedDecl);
}
mMultiReplacements.emplace_back(getParentNode()->getAsBlock(), decl,
std::move(*newSequence));
}
};
} // anonymous namespace
bool NameEmbeddedStructUniformsMetal(TCompiler *compiler,
TIntermBlock *root,
TSymbolTable *symbolTable)
{
Traverser nameStructs(symbolTable);
root->traverse(&nameStructs);
return nameStructs.updateTree(compiler, root);
}
} // namespace sh

31
src/compiler/translator/TranslatorMetalDirect/NameEmbeddedUniformStructsMetal.h Normal file
Просмотреть файл

@ -0,0 +1,31 @@
//
// Copyright 2018 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
// NameEmbeddedUniformStructs: Gives nameless uniform struct internal names.
//
// For example:
// uniform struct { int a; } uni;
// becomes:
// struct s1 { int a; };
// uniform s1 uni;
//
#ifndef COMPILER_TRANSLATOR_TREEOPS_NAMEEMBEDDEDUNIFORMSTRUCTSMETAL_H_
#define COMPILER_TRANSLATOR_TREEOPS_NAMEEMBEDDEDUNIFORMSTRUCTSMETAL_H_
#include "common/angleutils.h"
namespace sh
{
class TCompiler;
class TIntermBlock;
class TSymbolTable;
ANGLE_NO_DISCARD bool NameEmbeddedStructUniformsMetal(TCompiler *compiler,
TIntermBlock *root,
TSymbolTable *symbolTable);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TREEOPS_NAMEEMBEDDEDUNIFORMSTRUCTSMETAL_H_

155
src/compiler/translator/TranslatorMetalDirect/NodeType.h Normal file
Просмотреть файл

@ -0,0 +1,155 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_NODETYPE_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_NODETYPE_H_
#include "compiler/translator/tree_util/IntermTraverse.h"
namespace sh
{
enum class NodeType
{
Unknown,
Symbol,
ConstantUnion,
FunctionPrototype,
PreprocessorDirective,
Unary,
Binary,
Ternary,
Swizzle,
IfElse,
Switch,
Case,
FunctionDefinition,
Aggregate,
Block,
GlobalQualifierDeclaration,
Declaration,
Loop,
Branch,
};
// This is a function like object instead of a function that stack allocates this because
// TIntermTraverser is a heavy object to construct.
class GetNodeType : private TIntermTraverser
{
NodeType nodeType;
public:
GetNodeType() : TIntermTraverser(true, false, false) {}
NodeType operator()(TIntermNode &node)
{
node.visit(Visit::PreVisit, this);
return nodeType;
}
private:
void visitSymbol(TIntermSymbol *) override { nodeType = NodeType::Symbol; }
void visitConstantUnion(TIntermConstantUnion *) override { nodeType = NodeType::ConstantUnion; }
void visitFunctionPrototype(TIntermFunctionPrototype *) override
{
nodeType = NodeType::FunctionPrototype;
}
void visitPreprocessorDirective(TIntermPreprocessorDirective *) override
{
nodeType = NodeType::PreprocessorDirective;
}
bool visitSwizzle(Visit, TIntermSwizzle *) override
{
nodeType = NodeType::Swizzle;
return false;
}
bool visitBinary(Visit, TIntermBinary *) override
{
nodeType = NodeType::Binary;
return false;
}
bool visitUnary(Visit, TIntermUnary *) override
{
nodeType = NodeType::Unary;
return false;
}
bool visitTernary(Visit, TIntermTernary *) override
{
nodeType = NodeType::Ternary;
return false;
}
bool visitIfElse(Visit, TIntermIfElse *) override
{
nodeType = NodeType::IfElse;
return false;
}
bool visitSwitch(Visit, TIntermSwitch *) override
{
nodeType = NodeType::Switch;
return false;
}
bool visitCase(Visit, TIntermCase *) override
{
nodeType = NodeType::Case;
return false;
}
bool visitFunctionDefinition(Visit, TIntermFunctionDefinition *) override
{
nodeType = NodeType::FunctionDefinition;
return false;
}
bool visitAggregate(Visit, TIntermAggregate *) override
{
nodeType = NodeType::Aggregate;
return false;
}
bool visitBlock(Visit, TIntermBlock *) override
{
nodeType = NodeType::Block;
return false;
}
bool visitGlobalQualifierDeclaration(Visit, TIntermGlobalQualifierDeclaration *) override
{
nodeType = NodeType::GlobalQualifierDeclaration;
return false;
}
bool visitDeclaration(Visit, TIntermDeclaration *) override
{
nodeType = NodeType::Declaration;
return false;
}
bool visitLoop(Visit, TIntermLoop *) override
{
nodeType = NodeType::Loop;
return false;
}
bool visitBranch(Visit, TIntermBranch *) override
{
nodeType = NodeType::Branch;
return false;
}
};
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_NODETYPE_H_

503
src/compiler/translator/TranslatorMetalDirect/Pipeline.cpp Normal file
Просмотреть файл

@ -0,0 +1,503 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/Pipeline.h"
#include "compiler/translator/tree_util/BuiltIn.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
#define VARIANT_NAME(variant, base) (variant == Variant::Modified ? base "Mod" : base)
bool Pipeline::uses(const TVariable &var) const
{
if (var.symbolType() == SymbolType::Empty)
{
return false;
}
if (globalInstanceVar)
{
return &var == globalInstanceVar;
}
const TType &nodeType = var.getType();
const TQualifier qualifier = nodeType.getQualifier();
switch (type)
{
case Type::VertexIn:
switch (qualifier)
{
case TQualifier::EvqAttribute:
case TQualifier::EvqVertexIn:
return true;
default:
return false;
}
case Type::VertexOut:
switch (qualifier)
{
case TQualifier::EvqVertexOut:
case TQualifier::EvqPosition:
case TQualifier::EvqFlatOut:
case TQualifier::EvqPointSize:
case TQualifier::EvqSmoothOut:
case TQualifier::EvqCentroidOut:
case TQualifier::EvqNoPerspectiveOut:
case TQualifier::EvqVaryingOut:
return true;
default:
return false;
}
case Type::FragmentIn:
switch (qualifier)
{
case TQualifier::EvqFragmentIn:
case TQualifier::EvqFlatIn:
case TQualifier::EvqSmoothIn:
case TQualifier::EvqCentroidIn:
case TQualifier::EvqNoPerspectiveIn:
case TQualifier::EvqVaryingIn:
return true;
default:
return false;
}
case Type::FragmentOut:
switch (qualifier)
{
case TQualifier::EvqFragmentOut:
case TQualifier::EvqFragColor:
case TQualifier::EvqFragData:
case TQualifier::EvqFragDepth:
case TQualifier::EvqSampleMask:
return true;
default:
return false;
}
case Type::UserUniforms:
switch (qualifier)
{
case TQualifier::EvqUniform:
return true;
default:
return false;
}
case Type::NonConstantGlobals:
switch (qualifier)
{
case TQualifier::EvqGlobal:
return true;
default:
return false;
}
case Type::InvocationVertexGlobals:
switch (qualifier)
{
case TQualifier::EvqVertexID:
return true;
default:
return false;
}
case Type::InvocationFragmentGlobals:
switch (qualifier)
{
case TQualifier::EvqFragCoord:
case TQualifier::EvqPointCoord:
case TQualifier::EvqFrontFacing:
return true;
default:
return false;
}
case Type::UniformBuffer:
switch (qualifier)
{
case TQualifier::EvqBuffer:
return true;
default:
return false;
}
case Type::AngleUniforms:
UNREACHABLE(); // globalInstanceVar should be non-null and thus never reach here.
return false;
case Type::Texture:
return IsSampler(nodeType.getBasicType());
case Type::InstanceId:
return Name(var) == Name(*BuiltInVariable::gl_InstanceID());
}
}
Name Pipeline::getStructTypeName(Variant variant) const
{
const char *name;
switch (type)
{
case Type::VertexIn:
name = VARIANT_NAME(variant, "VertexIn");
break;
case Type::VertexOut:
name = VARIANT_NAME(variant, "VertexOut");
break;
case Type::FragmentIn:
name = VARIANT_NAME(variant, "FragmentIn");
break;
case Type::FragmentOut:
name = VARIANT_NAME(variant, "FragmentOut");
break;
case Type::UserUniforms:
name = VARIANT_NAME(variant, "UserUniforms");
break;
case Type::AngleUniforms:
name = VARIANT_NAME(variant, "AngleUniforms");
break;
case Type::NonConstantGlobals:
name = VARIANT_NAME(variant, "NonConstGlobals");
break;
case Type::InvocationVertexGlobals:
name = VARIANT_NAME(variant, "InvocationVertexGlobals");
break;
case Type::InvocationFragmentGlobals:
name = VARIANT_NAME(variant, "InvocationFragmentGlobals");
break;
case Type::Texture:
name = VARIANT_NAME(variant, "TextureEnvs");
break;
case Type::InstanceId:
name = VARIANT_NAME(variant, "InstanceId");
break;
case Type::UniformBuffer:
name = VARIANT_NAME(variant, "UniformBuffer");
}
return Name(name);
}
Name Pipeline::getStructInstanceName(Variant variant) const
{
const char *name;
switch (type)
{
case Type::VertexIn:
name = VARIANT_NAME(variant, "vertexIn");
break;
case Type::VertexOut:
name = VARIANT_NAME(variant, "vertexOut");
break;
case Type::FragmentIn:
name = VARIANT_NAME(variant, "fragmentIn");
break;
case Type::FragmentOut:
name = VARIANT_NAME(variant, "fragmentOut");
break;
case Type::UserUniforms:
name = VARIANT_NAME(variant, "userUniforms");
break;
case Type::AngleUniforms:
name = VARIANT_NAME(variant, "angleUniforms");
break;
case Type::NonConstantGlobals:
name = VARIANT_NAME(variant, "nonConstGlobals");
break;
case Type::InvocationVertexGlobals:
name = VARIANT_NAME(variant, "invocationVertexGlobals");
break;
case Type::InvocationFragmentGlobals:
name = VARIANT_NAME(variant, "invocationFragmentGlobals");
break;
case Type::Texture:
name = VARIANT_NAME(variant, "textureEnvs");
break;
case Type::InstanceId:
name = VARIANT_NAME(variant, "instanceId");
break;
case Type::UniformBuffer:
name = VARIANT_NAME(variant, "uniformBuffer");
}
return Name(name);
}
static bool AllowPacking(Pipeline::Type type)
{
using Type = Pipeline::Type;
switch (type)
{
case Type::UniformBuffer:
case Type::UserUniforms:
return true;
case Type::VertexIn:
case Type::VertexOut:
case Type::FragmentIn:
case Type::FragmentOut:
case Type::AngleUniforms:
case Type::NonConstantGlobals:
case Type::InvocationVertexGlobals:
case Type::InvocationFragmentGlobals:
case Type::Texture:
case Type::InstanceId:
return false;
}
}
static bool AllowPadding(Pipeline::Type type)
{
using Type = Pipeline::Type;
switch (type)
{
case Type::UserUniforms:
case Type::VertexIn:
case Type::VertexOut:
case Type::FragmentIn:
case Type::FragmentOut:
case Type::AngleUniforms:
case Type::NonConstantGlobals:
case Type::InvocationVertexGlobals:
case Type::InvocationFragmentGlobals:
case Type::UniformBuffer:
return true;
case Type::Texture:
case Type::InstanceId:
return false;
}
}
enum Compare
{
LT,
LTE,
EQ,
GTE,
GT,
};
template <typename T>
static bool CompareBy(Compare op, const T &x, const T &y)
{
switch (op)
{
case LT:
return x < y;
case LTE:
return x <= y;
case EQ:
return x == y;
case GTE:
return x >= y;
case GT:
return x > y;
}
}
template <TBasicType BT, Compare Cmp, int MatchDim, int NewDim>
static int SaturateVectorOf(const TField &field)
{
static_assert(NewDim >= MatchDim, "");
const TType &type = *field.type();
ASSERT(type.isScalar() || type.isVector());
const bool cond = type.getBasicType() == BT && !type.isArray() &&
CompareBy(Cmp, type.getNominalSize(), MatchDim) &&
type.getQualifier() != TQualifier::EvqFragDepth;
if (cond)
{
return NewDim;
}
return 0;
}
ModifyStructConfig Pipeline::externalStructModifyConfig() const
{
using Pred = ModifyStructConfig::Predicate;
using SatVec = ModifyStructConfig::SaturateVector;
ModifyStructConfig config(
isPipelineOut() ? ConvertType::OriginalToModified : ConvertType::ModifiedToOriginal,
AllowPacking(type), AllowPadding(type));
config.externalAddressSpace = externalAddressSpace();
switch (type)
{
case Type::VertexIn:
config.inlineArray = Pred::True;
config.splitMatrixColumns = Pred::True;
config.inlineStruct = Pred::True;
break;
case Type::VertexOut:
config.inlineArray = Pred::True;
config.splitMatrixColumns = Pred::True;
config.inlineStruct = Pred::True;
break;
case Type::FragmentIn:
config.inlineArray = Pred::True;
config.splitMatrixColumns = Pred::True;
config.inlineStruct = Pred::True;
break;
case Type::FragmentOut:
config.inlineArray = Pred::True;
config.splitMatrixColumns = Pred::True;
config.inlineStruct = Pred::True;
config.saturateScalarOrVector = [](const TField &field) {
if (field.type()->getQualifier() == TQualifier::EvqSampleMask)
{
return 1;
}
if (int s = SaturateVectorOf<TBasicType::EbtInt, LT, 4, 4>(field))
{
return s;
}
if (int s = SaturateVectorOf<TBasicType::EbtUInt, LT, 4, 4>(field))
{
return s;
}
if (int s = SaturateVectorOf<TBasicType::EbtFloat, LT, 4, 4>(field))
{
return s;
}
return 0;
};
break;
case Type::UserUniforms:
config.promoteBoolToUint = Pred::True;
config.saturateMatrixRows = SatVec::FullySaturate;
config.saturateScalarOrVectorArrays = SatVec::FullySaturate;
config.recurseStruct = Pred::True;
break;
case Type::AngleUniforms:
config.initialBlockStorage = TLayoutBlockStorage::EbsStd430; // XXX: Correct?
break;
case Type::NonConstantGlobals:
break;
case Type::UniformBuffer:
config.promoteBoolToUint = Pred::True;
config.saturateMatrixRows = SatVec::FullySaturate;
config.saturateScalarOrVectorArrays = SatVec::FullySaturate;
config.recurseStruct = Pred::True;
break;
case Type::InvocationVertexGlobals:
case Type::InvocationFragmentGlobals:
case Type::Texture:
case Type::InstanceId:
break;
}
return config;
}
bool Pipeline::alwaysRequiresLocalVariableDeclarationInMain() const
{
switch (type)
{
case Type::VertexIn:
case Type::FragmentIn:
case Type::UserUniforms:
case Type::AngleUniforms:
case Type::UniformBuffer:
return false;
case Type::VertexOut:
case Type::FragmentOut:
case Type::NonConstantGlobals:
case Type::InvocationVertexGlobals:
case Type::InvocationFragmentGlobals:
case Type::Texture:
case Type::InstanceId:
return true;
}
}
bool Pipeline::isPipelineOut() const
{
switch (type)
{
case Type::VertexIn:
case Type::FragmentIn:
case Type::UserUniforms:
case Type::AngleUniforms:
case Type::NonConstantGlobals:
case Type::InvocationVertexGlobals:
case Type::InvocationFragmentGlobals:
case Type::Texture:
case Type::InstanceId:
case Type::UniformBuffer:
return false;
case Type::VertexOut:
case Type::FragmentOut:
return true;
}
}
AddressSpace Pipeline::externalAddressSpace() const
{
switch (type)
{
case Type::VertexIn:
case Type::FragmentIn:
case Type::NonConstantGlobals:
case Type::InvocationVertexGlobals:
case Type::InvocationFragmentGlobals:
case Type::Texture:
case Type::InstanceId:
case Type::FragmentOut:
case Type::VertexOut:
return AddressSpace::Thread;
case Type::UserUniforms:
case Type::AngleUniforms:
case Type::UniformBuffer:
return AddressSpace::Constant;
}
}
bool PipelineStructs::matches(const TStructure &s, bool internal, bool external) const
{
PipelineScoped<TStructure> ps[] = {
fragmentIn,
fragmentOut,
vertexIn,
vertexOut,
userUniforms,
/* angleUniforms, */
nonConstantGlobals,
invocationVertexGlobals,
invocationFragmentGlobals,
uniformBuffers,
texture,
instanceId,
};
for (const auto &p : ps)
{
if (internal && p.internal == &s)
{
return true;
}
if (external && p.external == &s)
{
return true;
}
}
return false;
}

125
src/compiler/translator/TranslatorMetalDirect/Pipeline.h Normal file
Просмотреть файл

@ -0,0 +1,125 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_PIPELINE_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_PIPELINE_H_
#include "compiler/translator/Symbol.h"
#include "compiler/translator/TranslatorMetalDirect/ModifyStruct.h"
#include "compiler/translator/TranslatorMetalDirect/Name.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
namespace sh
{
// Data that is scoped as `external` and `internal` for a given pipeline.
template <typename T>
struct PipelineScoped
{
// Data that is configured to talk externally to the program.
// May coincide with `internal`, but may also diverge from `internal`.
const T *external = nullptr;
// Data that is configured to talk internally within the program.
// May coincide with `external`, but may also diverge from `external`.
const T *internal = nullptr;
// Returns true iff the input coincides with either `external` or `internal` data.
bool matches(const T &object) const { return external == &object || internal == &object; }
// Both `external` and `internal` representations are non-null.
bool isTotallyFull() const { return external && internal; }
// Both `external` and `internal` representations are null.
bool isTotallyEmpty() const { return !external && !internal; }
// Both `external` and `internal` representations are the same.
bool isUniform() const { return external == internal; }
};
// Represents a high-level program pipeline.
class Pipeline
{
public:
enum class Type
{
VertexIn,
VertexOut,
FragmentIn,
FragmentOut,
UserUniforms,
AngleUniforms,
NonConstantGlobals,
InvocationVertexGlobals,
InvocationFragmentGlobals,
UniformBuffer,
Texture,
InstanceId,
};
enum class Variant
{
// For all internal pipeline uses.
// For external pipeline uses if pipeline does not require splitting or saturation.
Original,
// Only for external pipeline uses if the pipeline was split or saturated.
Modified,
};
public:
// The type of the pipeline.
Type type;
// Non-null if a global instance of the pipeline struct already exists.
// If non-null struct splitting should not be needed.
const TVariable *globalInstanceVar;
public:
// Returns true iff the variable belongs to the pipeline.
bool uses(const TVariable &var) const;
// Returns the name for the struct type that stores variables of this pipeline.
Name getStructTypeName(Variant variant) const;
// Returns the name for the struct instance that stores variables of this pipeline.
Name getStructInstanceName(Variant variant) const;
ModifyStructConfig externalStructModifyConfig() const;
// Returns true if the pipeline always requires a non-parameter local instance declaration of
// the pipeline structures.
bool alwaysRequiresLocalVariableDeclarationInMain() const;
// Returns true iff the pipeline is an output pipeline. The external pipeline structure should
// be returned from `main`.
bool isPipelineOut() const;
AddressSpace externalAddressSpace() const;
};
// A collection of various pipeline structures.
struct PipelineStructs : angle::NonCopyable
{
PipelineScoped<TStructure> fragmentIn;
PipelineScoped<TStructure> fragmentOut;
PipelineScoped<TStructure> vertexIn;
PipelineScoped<TStructure> vertexOut;
PipelineScoped<TStructure> userUniforms;
PipelineScoped<TStructure> angleUniforms;
PipelineScoped<TStructure> nonConstantGlobals;
PipelineScoped<TStructure> invocationVertexGlobals;
PipelineScoped<TStructure> invocationFragmentGlobals;
PipelineScoped<TStructure> uniformBuffers;
PipelineScoped<TStructure> texture;
PipelineScoped<TStructure> instanceId;
bool matches(const TStructure &s, bool internal, bool external) const;
};
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_PIPELINE_H_

3922
src/compiler/translator/TranslatorMetalDirect/ProgramPrelude.cpp Normal file
Просмотреть файл

Разница между файлами не показана из-за своего большого размера Загрузить разницу

48
src/compiler/translator/TranslatorMetalDirect/ProgramPrelude.h Normal file
Просмотреть файл

@ -0,0 +1,48 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_PROGRAMPRELUDE_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_PROGRAMPRELUDE_H_
#include <unordered_set>
#include "common/angleutils.h"
namespace sh
{
class TInfoSinkBase;
class TIntermBlock;
enum class MetalShaderType
{
None,
Vertex,
Fragment,
Compute, // Unused currently
Count,
};
struct ProgramPreludeConfig
{
public:
ProgramPreludeConfig() {}
explicit ProgramPreludeConfig(MetalShaderType shaderType) : shaderType(shaderType) {}
bool hasStructEq = false;
MetalShaderType shaderType = MetalShaderType::None;
};
// This emits fixed helper Metal code directly without adding code to the AST. This walks the AST to
// figure out the required what prelude MSL code is needed for various things in the AST. You can
// think of this as effectively inlining various portions of a helper library into the emitted
// Metal program.
ANGLE_NO_DISCARD bool EmitProgramPrelude(TIntermBlock &root,
TInfoSinkBase &out,
const ProgramPreludeConfig &ppc);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_PROGRAMPRELUDE_H_

129
src/compiler/translator/TranslatorMetalDirect/ReduceInterfaceBlocks.cpp Normal file
Просмотреть файл

@ -0,0 +1,129 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <algorithm>
#include <unordered_map>
#include "compiler/translator/TranslatorMetalDirect.h"
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
#include "compiler/translator/TranslatorMetalDirect/ReduceInterfaceBlocks.h"
#include "compiler/translator/tree_ops/SeparateDeclarations.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
class Reducer : public TIntermRebuild
{
std::unordered_map<const TInterfaceBlock *, const TVariable *> mLiftedMap;
std::unordered_map<const TVariable *, const TVariable *> mInstanceMap;
IdGen &mIdGen;
public:
Reducer(TCompiler &compiler, IdGen &idGen)
: TIntermRebuild(compiler, true, false), mIdGen(idGen)
{}
PreResult visitDeclarationPre(TIntermDeclaration &declNode) override
{
ASSERT(declNode.getChildCount() == 1);
TIntermNode &node = *declNode.getChildNode(0);
if (TIntermSymbol *symbolNode = node.getAsSymbolNode())
{
const TVariable &var = symbolNode->variable();
const TType &type = var.getType();
const SymbolType symbolType = var.symbolType();
if (const TInterfaceBlock *interfaceBlock = type.getInterfaceBlock())
{
if (symbolType == SymbolType::Empty)
{
// Create instance variable
auto &structure =
*new TStructure(&mSymbolTable, interfaceBlock->name(),
&interfaceBlock->fields(), interfaceBlock->symbolType());
auto &structVar = CreateStructTypeVariable(mSymbolTable, structure);
auto &instanceVar = CreateInstanceVariable(
mSymbolTable, structure, mIdGen.createNewName(interfaceBlock->name()),
TQualifier::EvqBuffer, &type.getArraySizes());
mLiftedMap[interfaceBlock] = &instanceVar;
TIntermNode *replacements[] = {
new TIntermDeclaration{new TIntermSymbol(&structVar)},
new TIntermDeclaration{new TIntermSymbol(&instanceVar)}};
return PreResult::Multi(std::begin(replacements), std::end(replacements));
}
else
{
ASSERT(type.getQualifier() == TQualifier::EvqUniform);
auto &structure =
*new TStructure(&mSymbolTable, interfaceBlock->name(),
&interfaceBlock->fields(), interfaceBlock->symbolType());
auto &structVar = CreateStructTypeVariable(mSymbolTable, structure);
auto &instanceVar =
CreateInstanceVariable(mSymbolTable, structure, Name(var),
TQualifier::EvqBuffer, &type.getArraySizes());
mInstanceMap[&var] = &instanceVar;
TIntermNode *replacements[] = {
new TIntermDeclaration{new TIntermSymbol(&structVar)},
new TIntermDeclaration{new TIntermSymbol(&instanceVar)}};
return PreResult::Multi(std::begin(replacements), std::end(replacements));
}
}
}
return {declNode, VisitBits::Both};
}
PreResult visitSymbolPre(TIntermSymbol &symbolNode) override
{
const TVariable &var = symbolNode.variable();
{
auto it = mInstanceMap.find(&var);
if (it != mInstanceMap.end())
{
return *new TIntermSymbol(it->second);
}
}
if (const TInterfaceBlock *ib = var.getType().getInterfaceBlock())
{
auto it = mLiftedMap.find(ib);
if (it != mLiftedMap.end())
{
return AccessField(*(it->second), var.name());
}
}
return symbolNode;
}
};
} // anonymous namespace
////////////////////////////////////////////////////////////////////////////////
bool sh::ReduceInterfaceBlocks(TCompiler &compiler, TIntermBlock &root, IdGen &idGen)
{
Reducer reducer(compiler, idGen);
if (!reducer.rebuildRoot(root))
{
return false;
}
if (!SeparateDeclarations(&compiler, &root))
{
return false;
}
return true;
}

35
src/compiler/translator/TranslatorMetalDirect/ReduceInterfaceBlocks.h Normal file
Просмотреть файл

@ -0,0 +1,35 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REDUCEINTERFACEBLOCKS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REDUCEINTERFACEBLOCKS_H_
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
namespace sh
{
// This rewrites interface block declarations only.
//
// Access of interface blocks is not rewritten (e.g. TOperator::EOpIndexDirectInterfaceBlock). //
// XXX: ^ Still true?
//
// Example:
// uniform Foo { int x; };
// Becomes:
// uniform int x;
//
// Example:
// uniform Foo { int x; } foo;
// Becomes:
// struct Foo { int x; }; uniform Foo x;
//
ANGLE_NO_DISCARD bool ReduceInterfaceBlocks(TCompiler &compiler, TIntermBlock &root, IdGen &idGen);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REDUCEINTERFACEBLOCKS_H_

50
src/compiler/translator/TranslatorMetalDirect/Reference.h Normal file
Просмотреть файл

@ -0,0 +1,50 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REFERENCE_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REFERENCE_H_
namespace sh
{
// Similar to std::reference_wrapper, but also lifts comparison operators.
template <typename T>
class Ref
{
public:
Ref(const Ref &) = default;
Ref(Ref &&) = default;
Ref(T &ref) : mPtr(&ref) {}
Ref &operator=(const Ref &) = default;
Ref &operator=(Ref &&) = default;
bool operator==(const Ref &other) const { return *mPtr == *other.mPtr; }
bool operator!=(const Ref &other) const { return *mPtr != *other.mPtr; }
bool operator<=(const Ref &other) const { return *mPtr <= *other.mPtr; }
bool operator>=(const Ref &other) const { return *mPtr >= *other.mPtr; }
bool operator<(const Ref &other) const { return *mPtr < *other.mPtr; }
bool operator>(const Ref &other) const { return *mPtr > *other.mPtr; }
T &get() { return *mPtr; }
T const &get() const { return *mPtr; }
operator T &() { return *mPtr; }
operator T const &() const { return *mPtr; }
operator T *() { return *mPtr; }
operator T const *() const { return *mPtr; }
private:
T *mPtr;
};
template <typename T>
using CRef = Ref<T const>;
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REFERENCE_H_

91
src/compiler/translator/TranslatorMetalDirect/RewriteCaseDeclarations.cpp Normal file
Просмотреть файл

@ -0,0 +1,91 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/RewriteCaseDeclarations.h"
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
class Rewriter : public TIntermRebuild
{
std::vector<std::vector<const TVariable *>> mDeclaredVarStack;
public:
Rewriter(TCompiler &compiler) : TIntermRebuild(compiler, true, true) {}
~Rewriter() override { ASSERT(mDeclaredVarStack.empty()); }
private:
PreResult visitSwitchPre(TIntermSwitch &node) override
{
mDeclaredVarStack.emplace_back();
return node;
}
PostResult visitSwitchPost(TIntermSwitch &node) override
{
ASSERT(!mDeclaredVarStack.empty());
const auto vars = std::move(mDeclaredVarStack.back());
mDeclaredVarStack.pop_back();
if (!vars.empty())
{
auto &block = *new TIntermBlock();
for (const TVariable *var : vars)
{
block.appendStatement(new TIntermDeclaration{var});
}
block.appendStatement(&node);
return block;
}
return node;
}
PreResult visitDeclarationPre(TIntermDeclaration &node) override
{
if (!mDeclaredVarStack.empty())
{
TIntermNode *parent = getParentNode();
if (parent->getAsBlock())
{
TIntermNode *grandparent = getParentNode(1);
if (grandparent && grandparent->getAsSwitchNode())
{
Declaration decl = ViewDeclaration(node);
mDeclaredVarStack.back().push_back(&decl.symbol.variable());
if (decl.initExpr)
{
return *new TIntermBinary(TOperator::EOpAssign, &decl.symbol,
decl.initExpr);
}
else
{
return nullptr;
}
}
}
}
return node;
}
};
} // anonymous namespace
////////////////////////////////////////////////////////////////////////////////
bool sh::RewriteCaseDeclarations(TCompiler &compiler, TIntermBlock &root)
{
if (!Rewriter(compiler).rebuildRoot(root))
{
return false;
}
return true;
}

50
src/compiler/translator/TranslatorMetalDirect/RewriteCaseDeclarations.h Normal file
Просмотреть файл

@ -0,0 +1,50 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITECASEDECLARATIONS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITECASEDECLARATIONS_H_
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
namespace sh
{
// EXAMPLE
// switch (expr)
// {
// case 0:
// int x = 0;
// break;
// case 1:
// int y = 0;
// {
// int z = 0;
// }
// break;
// }
// Becomes
// {
// int x;
// int y;
// switch (expr)
// {
// case 0:
// x = 0;
// break;
// case 1:
// y = 0;
// {
// int z = 0;
// }
// break;
// }
// }
ANGLE_NO_DISCARD bool RewriteCaseDeclarations(TCompiler &compiler, TIntermBlock &root);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITECASEDECLARATIONS_H_

113
src/compiler/translator/TranslatorMetalDirect/RewriteGlobalQualifierDecls.cpp Normal file
Просмотреть файл

@ -0,0 +1,113 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/RewriteGlobalQualifierDecls.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
using namespace sh;
namespace
{
class FindDeclaredGlobals : public TIntermRebuild
{
public:
std::unordered_set<const TVariable *> mDeclaredGlobals;
FindDeclaredGlobals(TCompiler &compiler) : TIntermRebuild(compiler, true, false) {}
PreResult visitDeclarationPre(TIntermDeclaration &declNode) override
{
TIntermNode *declaratorNode = declNode.getChildNode(0);
TIntermSymbol *symbolNode = nullptr;
if (TIntermBinary *initNode = declaratorNode->getAsBinaryNode())
{
symbolNode = initNode->getLeft()->getAsSymbolNode();
}
else
{
symbolNode = declaratorNode->getAsSymbolNode();
}
ASSERT(symbolNode);
const TVariable &var = symbolNode->variable();
mDeclaredGlobals.insert(&var);
return {declNode, VisitBits::Neither};
}
PreResult visitFunctionDefinitionPre(TIntermFunctionDefinition &node) override
{
return {node, VisitBits::Neither};
}
};
class Rewriter : public TIntermRebuild
{
const std::unordered_set<const TVariable *> &mDeclaredGlobals;
Invariants &mOutInvariants;
public:
Rewriter(TCompiler &compiler,
const std::unordered_set<const TVariable *> &declaredGlobals,
Invariants &outInvariants)
: TIntermRebuild(compiler, true, false),
mDeclaredGlobals(declaredGlobals),
mOutInvariants(outInvariants)
{}
PreResult visitGlobalQualifierDeclarationPre(
TIntermGlobalQualifierDeclaration &gqDeclNode) override
{
TIntermSymbol &symbolNode = *gqDeclNode.getSymbol();
const TVariable &var = symbolNode.variable();
if (gqDeclNode.isInvariant())
{
mOutInvariants.insert(var);
}
if (mDeclaredGlobals.find(&var) == mDeclaredGlobals.end())
{
return *new TIntermDeclaration{&symbolNode};
}
return nullptr;
}
PreResult visitDeclarationPre(TIntermDeclaration &node) override
{
return {node, VisitBits::Neither};
}
PreResult visitFunctionDefinitionPre(TIntermFunctionDefinition &node) override
{
return {node, VisitBits::Neither};
}
};
} // anonymous namespace
bool sh::RewriteGlobalQualifierDecls(TCompiler &compiler,
TIntermBlock &root,
Invariants &outInvariants)
{
FindDeclaredGlobals fdg(compiler);
if (!fdg.rebuildRoot(root))
{
UNREACHABLE();
return false;
}
Rewriter rewriter(compiler, fdg.mDeclaredGlobals, outInvariants);
if (!rewriter.rebuildRoot(root))
{
return false;
}
return true;
}

48
src/compiler/translator/TranslatorMetalDirect/RewriteGlobalQualifierDecls.h Normal file
Просмотреть файл

@ -0,0 +1,48 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEGLOBALQUALIFIERDECLS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEGLOBALQUALIFIERDECLS_H_
#include <unordered_set>
#include "compiler/translator/Compiler.h"
namespace sh
{
// Tracks TVariables and TFields that are marked as "invariant".
class Invariants
{
public:
void insert(const TVariable &var) { mInvariants.insert(&var); }
void insert(const TField &field) { mInvariants.insert(&field); }
bool contains(const TVariable &var) const
{
return mInvariants.find(&var) != mInvariants.end();
}
bool contains(const TField &field) const
{
return mInvariants.find(&field) != mInvariants.end();
}
private:
std::unordered_set<const void *> mInvariants;
};
// This rewrites TIntermGlobalQualifierDeclarations as TIntermDeclarations.
// `outInvariants` is populated with the information that would otherwise be lost by such a
// transform.
ANGLE_NO_DISCARD bool RewriteGlobalQualifierDecls(TCompiler &compiler,
TIntermBlock &root,
Invariants &outInvariants);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEGLOBALQUALIFIERDECLS_H_

462
src/compiler/translator/TranslatorMetalDirect/RewriteKeywords.cpp Normal file
Просмотреть файл

@ -0,0 +1,462 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <cctype>
#include <cstring>
#include <limits>
#include <map>
#include <unordered_map>
#include <unordered_set>
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
#include "compiler/translator/TranslatorMetalDirect/RewriteKeywords.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
template <typename T>
using Remapping = std::unordered_map<const T *, const T *>;
class Rewriter : public TIntermRebuild
{
private:
const std::set<ImmutableString> &mKeywords;
IdGen &mIdGen;
Remapping<TField> modifiedFields;
Remapping<TFieldList> mFieldLists;
Remapping<TFunction> mFunctions;
Remapping<TInterfaceBlock> mInterfaceBlocks;
Remapping<TStructure> mStructures;
Remapping<TVariable> mVariables;
std::map<ImmutableString, std::string> mPredefinedNames;
std::string mNewNameBuffer;
private:
template <typename T>
ImmutableString maybeCreateNewName(T const &object)
{
if (needsRenaming(object, false))
{
auto it = mPredefinedNames.find(Name(object).rawName());
if (it != mPredefinedNames.end())
{
return ImmutableString(it->second);
}
return mIdGen.createNewName(Name(object)).rawName();
}
return Name(object).rawName();
}
const TField *createRenamed(const TField &field)
{
auto *renamed =
new TField(const_cast<TType *>(&getRenamedOrOriginal(*field.type())),
maybeCreateNewName(field), field.line(), SymbolType::AngleInternal);
return renamed;
}
const TFieldList *createRenamed(const TFieldList &fieldList)
{
auto *renamed = new TFieldList();
for (const TField *field : fieldList)
{
renamed->push_back(const_cast<TField *>(&getRenamedOrOriginal(*field)));
}
return renamed;
}
const TFunction *createRenamed(const TFunction &function)
{
auto *renamed =
new TFunction(&mSymbolTable, maybeCreateNewName(function), SymbolType::AngleInternal,
&getRenamedOrOriginal(function.getReturnType()),
function.isKnownToNotHaveSideEffects());
const size_t paramCount = function.getParamCount();
for (size_t i = 0; i < paramCount; ++i)
{
const TVariable &param = *function.getParam(i);
renamed->addParameter(&getRenamedOrOriginal(param));
}
if (function.isDefined())
{
renamed->setDefined();
}
if (function.hasPrototypeDeclaration())
{
renamed->setHasPrototypeDeclaration();
}
return renamed;
}
const TInterfaceBlock *createRenamed(const TInterfaceBlock &interfaceBlock)
{
TLayoutQualifier layoutQualifier = TLayoutQualifier::Create();
layoutQualifier.blockStorage = interfaceBlock.blockStorage();
layoutQualifier.binding = interfaceBlock.blockBinding();
auto *renamed =
new TInterfaceBlock(&mSymbolTable, maybeCreateNewName(interfaceBlock),
&getRenamedOrOriginal(interfaceBlock.fields()), layoutQualifier,
SymbolType::AngleInternal, interfaceBlock.extension());
return renamed;
}
const TStructure *createRenamed(const TStructure &structure)
{
auto *renamed =
new TStructure(&mSymbolTable, maybeCreateNewName(structure),
&getRenamedOrOriginal(structure.fields()), SymbolType::AngleInternal);
renamed->setAtGlobalScope(structure.atGlobalScope());
return renamed;
}
const TType *createRenamed(const TType &type)
{
TType *renamed;
if (const TStructure *structure = type.getStruct())
{
renamed = new TType(&getRenamedOrOriginal(*structure), type.isStructSpecifier());
}
else if (const TInterfaceBlock *interfaceBlock = type.getInterfaceBlock())
{
renamed = new TType(&getRenamedOrOriginal(*interfaceBlock), type.getQualifier(),
type.getLayoutQualifier());
}
else
{
UNREACHABLE(); // Can't rename built-in types.
renamed = nullptr;
}
if (type.isArray())
{
renamed->makeArrays(type.getArraySizes());
}
renamed->setPrecise(type.isPrecise());
renamed->setInvariant(type.isInvariant());
renamed->setMemoryQualifier(type.getMemoryQualifier());
renamed->setLayoutQualifier(type.getLayoutQualifier());
return renamed;
}
const TVariable *createRenamed(const TVariable &variable)
{
auto *renamed = new TVariable(&mSymbolTable, maybeCreateNewName(variable),
&getRenamedOrOriginal(variable.getType()),
SymbolType::AngleInternal, variable.extension());
return renamed;
}
template <typename T>
const T *tryGetRenamedImpl(const T &object, Remapping<T> *remapping)
{
if (!needsRenaming(object, true))
{
return nullptr;
}
if (remapping)
{
auto it = remapping->find(&object);
if (it != remapping->end())
{
return it->second;
}
}
const T *renamedObject = createRenamed(object);
if (remapping)
{
(*remapping)[&object] = renamedObject;
}
return renamedObject;
}
const TField *tryGetRenamed(const TField &field)
{
return tryGetRenamedImpl(field, &modifiedFields);
}
const TFieldList *tryGetRenamed(const TFieldList &fieldList)
{
return tryGetRenamedImpl(fieldList, &mFieldLists);
}
const TFunction *tryGetRenamed(const TFunction &func)
{
return tryGetRenamedImpl(func, &mFunctions);
}
const TInterfaceBlock *tryGetRenamed(const TInterfaceBlock &interfaceBlock)
{
return tryGetRenamedImpl(interfaceBlock, &mInterfaceBlocks);
}
const TStructure *tryGetRenamed(const TStructure &structure)
{
return tryGetRenamedImpl(structure, &mStructures);
}
const TType *tryGetRenamed(const TType &type)
{
return tryGetRenamedImpl(type, static_cast<Remapping<TType> *>(nullptr));
}
const TVariable *tryGetRenamed(const TVariable &variable)
{
return tryGetRenamedImpl(variable, &mVariables);
}
template <typename T>
const T &getRenamedOrOriginal(const T &object)
{
const T *renamed = tryGetRenamed(object);
if (renamed)
{
return *renamed;
}
return object;
}
template <typename T>
bool needsRenamingImpl(const T &object) const
{
const SymbolType symbolType = object.symbolType();
switch (symbolType)
{
case SymbolType::BuiltIn:
case SymbolType::AngleInternal:
case SymbolType::Empty:
return false;
case SymbolType::UserDefined:
break;
}
const ImmutableString name = Name(object).rawName();
if (mKeywords.find(name) != mKeywords.end())
{
return true;
}
if (name.beginsWith(kAngleInternalPrefix))
{
return true;
}
return false;
}
bool needsRenaming(const TField &field, bool recursive) const
{
return needsRenamingImpl(field) || (recursive && needsRenaming(*field.type(), true));
}
bool needsRenaming(const TFieldList &fieldList, bool recursive) const
{
ASSERT(recursive);
for (const TField *field : fieldList)
{
if (needsRenaming(*field, true))
{
return true;
}
}
return false;
}
bool needsRenaming(const TFunction &function, bool recursive) const
{
if (needsRenamingImpl(function))
{
return true;
}
if (!recursive)
{
return false;
}
const size_t paramCount = function.getParamCount();
for (size_t i = 0; i < paramCount; ++i)
{
const TVariable &param = *function.getParam(i);
if (needsRenaming(param, true))
{
return true;
}
}
return false;
}
bool needsRenaming(const TInterfaceBlock &interfaceBlock, bool recursive) const
{
return needsRenamingImpl(interfaceBlock) ||
(recursive && needsRenaming(interfaceBlock.fields(), true));
}
bool needsRenaming(const TStructure &structure, bool recursive) const
{
return needsRenamingImpl(structure) ||
(recursive && needsRenaming(structure.fields(), true));
}
bool needsRenaming(const TType &type, bool recursive) const
{
if (const TStructure *structure = type.getStruct())
{
return needsRenaming(*structure, recursive);
}
else if (const TInterfaceBlock *interfaceBlock = type.getInterfaceBlock())
{
return needsRenaming(*interfaceBlock, recursive);
}
else
{
return false;
}
}
bool needsRenaming(const TVariable &variable, bool recursive) const
{
return needsRenamingImpl(variable) ||
(recursive && needsRenaming(variable.getType(), true));
}
public:
Rewriter(TCompiler &compiler, IdGen &idGen, const std::set<ImmutableString> &keywords)
: TIntermRebuild(compiler, false, true), mKeywords(keywords), mIdGen(idGen)
{}
PostResult visitSymbolPost(TIntermSymbol &symbolNode) override
{
const TVariable &var = symbolNode.variable();
if (needsRenaming(var, true))
{
const TVariable &rVar = getRenamedOrOriginal(var);
return *new TIntermSymbol(&rVar);
}
return symbolNode;
}
PostResult visitFunctionPrototype(TIntermFunctionPrototype &funcProtoNode)
{
const TFunction &func = *funcProtoNode.getFunction();
if (needsRenaming(func, true))
{
const TFunction &rFunc = getRenamedOrOriginal(func);
return *new TIntermFunctionPrototype(&rFunc);
}
return funcProtoNode;
}
PostResult visitDeclarationPost(TIntermDeclaration &declNode) override
{
Declaration decl = ViewDeclaration(declNode);
const TVariable &var = decl.symbol.variable();
if (needsRenaming(var, true))
{
const TVariable &rVar = getRenamedOrOriginal(var);
return *new TIntermDeclaration(&rVar, decl.initExpr);
}
return declNode;
}
PostResult visitFunctionDefinitionPost(TIntermFunctionDefinition &funcDefNode) override
{
TIntermFunctionPrototype &funcProtoNode = *funcDefNode.getFunctionPrototype();
const TFunction &func = *funcProtoNode.getFunction();
if (needsRenaming(func, true))
{
const TFunction &rFunc = getRenamedOrOriginal(func);
auto *rFuncProtoNode = new TIntermFunctionPrototype(&rFunc);
return *new TIntermFunctionDefinition(rFuncProtoNode, funcDefNode.getBody());
}
return funcDefNode;
}
PostResult visitAggregatePost(TIntermAggregate &aggregateNode) override
{
if (aggregateNode.isConstructor())
{
const TType &type = aggregateNode.getType();
if (needsRenaming(type, true))
{
const TType &rType = getRenamedOrOriginal(type);
return TIntermAggregate::CreateConstructor(rType, aggregateNode.getSequence());
}
}
else
{
const TFunction &func = *aggregateNode.getFunction();
if (needsRenaming(func, true))
{
const TFunction &rFunc = getRenamedOrOriginal(func);
switch (aggregateNode.getOp())
{
case TOperator::EOpCallFunctionInAST:
return TIntermAggregate::CreateFunctionCall(rFunc,
aggregateNode.getSequence());
case TOperator::EOpCallInternalRawFunction:
return TIntermAggregate::CreateRawFunctionCall(rFunc,
aggregateNode.getSequence());
default:
return TIntermAggregate::CreateBuiltInFunctionCall(
rFunc, aggregateNode.getSequence());
}
}
}
return aggregateNode;
}
void predefineName(const ImmutableString name, std::string prePopulatedName)
{
mPredefinedNames[name] = prePopulatedName;
}
};
} // anonymous namespace
////////////////////////////////////////////////////////////////////////////////
bool sh::RewriteKeywords(TCompiler &compiler,
TIntermBlock &root,
IdGen &idGen,
const std::set<ImmutableString> &keywords)
{
Rewriter rewriter(compiler, idGen, keywords);
const auto &inputAttrs = compiler.getAttributes();
for (const auto &var : inputAttrs)
{
rewriter.predefineName(ImmutableString(var.name), var.mappedName);
}
if (!rewriter.rebuildRoot(root))
{
return false;
}
return true;
}

27
src/compiler/translator/TranslatorMetalDirect/RewriteKeywords.h Normal file
Просмотреть файл

@ -0,0 +1,27 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEKEYWORDS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEKEYWORDS_H_
#include <set>
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/IdGen.h"
namespace sh
{
// This walks the tree and renames all names that conflict with the input `keywords`.
ANGLE_NO_DISCARD bool RewriteKeywords(TCompiler &compiler,
TIntermBlock &root,
IdGen &idGen,
const std::set<ImmutableString> &keywords);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEKEYWORDS_H_

221
src/compiler/translator/TranslatorMetalDirect/RewriteOutArgs.cpp Normal file
Просмотреть файл

@ -0,0 +1,221 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/RewriteOutArgs.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
using namespace sh;
namespace
{
template <typename T>
class SmallMultiSet
{
public:
struct Entry
{
T elem;
size_t count;
};
const Entry *find(const T &x) const
{
for (auto &entry : mEntries)
{
if (x == entry.elem)
{
return &entry;
}
}
return nullptr;
}
size_t multiplicity(const T &x) const
{
const Entry *entry = find(x);
return entry ? entry->count : 0;
}
const Entry &insert(const T &x)
{
Entry *entry = findMutable(x);
if (entry)
{
++entry->count;
return *entry;
}
else
{
mEntries.push_back({x, 1});
return mEntries.back();
}
}
void clear() { mEntries.clear(); }
bool empty() const { return mEntries.empty(); }
size_t uniqueSize() const { return mEntries.size(); }
private:
ANGLE_INLINE Entry *findMutable(const T &x) { return const_cast<Entry *>(find(x)); }
private:
std::vector<Entry> mEntries;
};
const TVariable *GetVariable(TIntermNode &node)
{
TIntermTyped *tyNode = node.getAsTyped();
ASSERT(tyNode);
if (TIntermSymbol *symbol = tyNode->getAsSymbolNode())
{
return &symbol->variable();
}
return nullptr;
}
class Rewriter : public TIntermRebuild
{
SmallMultiSet<const TVariable *> mVarBuffer; // reusable buffer
SymbolEnv &mSymbolEnv;
public:
~Rewriter() override { ASSERT(mVarBuffer.empty()); }
Rewriter(TCompiler &compiler, SymbolEnv &symbolEnv)
: TIntermRebuild(compiler, false, true), mSymbolEnv(symbolEnv)
{}
static bool argAlreadyProcessed(TIntermTyped *arg)
{
if (arg->getAsAggregate())
{
const TFunction *func = arg->getAsAggregate()->getFunction();
if (func && func->symbolType() == SymbolType::AngleInternal &&
func->name() == "swizzle_ref")
{
return true;
}
}
return false;
}
PostResult visitAggregatePost(TIntermAggregate &aggregateNode) override
{
ASSERT(mVarBuffer.empty());
const TFunction *func = aggregateNode.getFunction();
if (!func)
{
return aggregateNode;
}
TIntermSequence &args = *aggregateNode.getSequence();
size_t argCount = args.size();
auto getParamQualifier = [&](size_t i) {
const TVariable &param = *func->getParam(i);
const TType &paramType = param.getType();
const TQualifier paramQual = paramType.getQualifier();
switch (paramQual)
{
case TQualifier::EvqOut:
case TQualifier::EvqInOut:
if (!mSymbolEnv.isReference(param))
{
mSymbolEnv.markAsReference(param, AddressSpace::Thread);
}
break;
default:
break;
}
return paramQual;
};
bool mightAlias = false;
for (size_t i = 0; i < argCount; ++i)
{
const TQualifier paramQual = getParamQualifier(i);
switch (paramQual)
{
case TQualifier::EvqOut:
case TQualifier::EvqInOut:
{
const TVariable *var = GetVariable(*args[i]);
if (mVarBuffer.insert(var).count > 1)
{
mightAlias = true;
i = argCount;
}
}
break;
default:
break;
}
}
const bool hasIndeterminateVar = mVarBuffer.find(nullptr);
if (!mightAlias)
{
mightAlias = hasIndeterminateVar && mVarBuffer.uniqueSize() > 1;
}
if (mightAlias)
{
for (size_t i = 0; i < argCount; ++i)
{
TIntermTyped *arg = args[i]->getAsTyped();
ASSERT(arg);
if (!argAlreadyProcessed(arg))
{
const TVariable *var = GetVariable(*arg);
const TQualifier paramQual = getParamQualifier(i);
if (hasIndeterminateVar || mVarBuffer.multiplicity(var) > 1)
{
switch (paramQual)
{
case TQualifier::EvqOut:
args[i] = &mSymbolEnv.callFunctionOverload(
Name("out"), arg->getType(), *new TIntermSequence{arg});
break;
case TQualifier::EvqInOut:
args[i] = &mSymbolEnv.callFunctionOverload(
Name("inout"), arg->getType(), *new TIntermSequence{arg});
break;
default:
break;
}
}
}
}
}
mVarBuffer.clear();
return aggregateNode;
}
};
} // anonymous namespace
bool sh::RewriteOutArgs(TCompiler &compiler, TIntermBlock &root, SymbolEnv &symbolEnv)
{
Rewriter rewriter(compiler, symbolEnv);
if (!rewriter.rebuildRoot(root))
{
return false;
}
return true;
}

33
src/compiler/translator/TranslatorMetalDirect/RewriteOutArgs.h Normal file
Просмотреть файл

@ -0,0 +1,33 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEOUTARGS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEOUTARGS_H_
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/ProgramPrelude.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
namespace sh
{
// e.g.:
// /*void foo(out int x, inout int y)*/
// foo(z, w);
// becomes
// foo(Out(z), InOut(w));
// unless `z` and `w` are detected to never alias.
// The translated example effectively behaves the same as:
// int _1;
// int _2 = w;
// foo(_1, _2);
// z = _1;
// w = _2;
ANGLE_NO_DISCARD bool RewriteOutArgs(TCompiler &compiler, TIntermBlock &root, SymbolEnv &symbolEnv);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEOUTARGS_H_

980
src/compiler/translator/TranslatorMetalDirect/RewritePipelines.cpp Normal file
Просмотреть файл

@ -0,0 +1,980 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <cstring>
#include <unordered_map>
#include <unordered_set>
#include "compiler/translator/TranslatorMetalDirect.h"
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/DiscoverDependentFunctions.h"
#include "compiler/translator/TranslatorMetalDirect/IdGen.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
#include "compiler/translator/TranslatorMetalDirect/MapSymbols.h"
#include "compiler/translator/TranslatorMetalDirect/Pipeline.h"
#include "compiler/translator/TranslatorMetalDirect/RewritePipelines.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
#include "compiler/translator/tree_ops/PruneNoOps.h"
#include "compiler/translator/tree_util/DriverUniform.h"
#include "compiler/translator/tree_util/FindMain.h"
#include "compiler/translator/tree_util/IntermTraverse.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
using VariableSet = std::unordered_set<const TVariable *>;
using VariableList = std::vector<const TVariable *>;
////////////////////////////////////////////////////////////////////////////////
struct PipelineStructInfo
{
VariableSet pipelineVariables;
PipelineScoped<TStructure> pipelineStruct;
const TFunction *funcOriginalToModified = nullptr;
const TFunction *funcModifiedToOriginal = nullptr;
bool isEmpty() const
{
if (pipelineStruct.isTotallyEmpty())
{
ASSERT(pipelineVariables.empty());
return true;
}
else
{
ASSERT(pipelineStruct.isTotallyFull());
ASSERT(!pipelineVariables.empty());
return false;
}
}
};
class GeneratePipelineStruct : private TIntermRebuild
{
private:
const Pipeline &mPipeline;
SymbolEnv &mSymbolEnv;
Invariants &mInvariants;
VariableList mPipelineVariableList;
IdGen &mIdGen;
PipelineStructInfo mInfo;
public:
static bool Exec(PipelineStructInfo &out,
TCompiler &compiler,
TIntermBlock &root,
IdGen &idGen,
const Pipeline &pipeline,
SymbolEnv &symbolEnv,
Invariants &invariants)
{
GeneratePipelineStruct self(compiler, idGen, pipeline, symbolEnv, invariants);
if (!self.exec(root))
{
return false;
}
out = self.mInfo;
return true;
}
private:
GeneratePipelineStruct(TCompiler &compiler,
IdGen &idGen,
const Pipeline &pipeline,
SymbolEnv &symbolEnv,
Invariants &invariants)
: TIntermRebuild(compiler, true, true),
mPipeline(pipeline),
mSymbolEnv(symbolEnv),
mInvariants(invariants),
mIdGen(idGen)
{}
bool exec(TIntermBlock &root)
{
if (!rebuildRoot(root))
{
return false;
}
if (mInfo.pipelineVariables.empty())
{
return true;
}
TIntermSequence seq;
const TStructure &pipelineStruct = [&]() -> const TStructure & {
if (mPipeline.globalInstanceVar)
{
return *mPipeline.globalInstanceVar->getType().getStruct();
}
else
{
return createInternalPipelineStruct(root, seq);
}
}();
ModifiedStructMachineries modifiedMachineries;
const bool isUBO = mPipeline.type == Pipeline::Type::UniformBuffer;
const bool modified = TryCreateModifiedStruct(
mCompiler, mSymbolEnv, mIdGen, mPipeline.externalStructModifyConfig(), pipelineStruct,
mPipeline.getStructTypeName(Pipeline::Variant::Modified), modifiedMachineries, isUBO,
!isUBO);
if (modified)
{
ASSERT(mPipeline.type != Pipeline::Type::Texture);
ASSERT(mPipeline.type == Pipeline::Type::AngleUniforms ||
!mPipeline.globalInstanceVar); // This shouldn't happen by construction.
auto getFunction = [](sh::TIntermFunctionDefinition *funcDecl) {
return funcDecl ? funcDecl->getFunction() : nullptr;
};
const size_t size = modifiedMachineries.size();
ASSERT(size > 0);
for (size_t i = 0; i < size; ++i)
{
const ModifiedStructMachinery &machinery = modifiedMachineries.at(i);
ASSERT(machinery.modifiedStruct);
seq.push_back(new TIntermDeclaration{
&CreateStructTypeVariable(mSymbolTable, *machinery.modifiedStruct)});
if (mPipeline.isPipelineOut())
{
ASSERT(machinery.funcOriginalToModified);
ASSERT(!machinery.funcModifiedToOriginal);
seq.push_back(machinery.funcOriginalToModified);
}
else
{
ASSERT(machinery.funcModifiedToOriginal);
ASSERT(!machinery.funcOriginalToModified);
seq.push_back(machinery.funcModifiedToOriginal);
}
if (i == size - 1)
{
mInfo.funcOriginalToModified = getFunction(machinery.funcOriginalToModified);
mInfo.funcModifiedToOriginal = getFunction(machinery.funcModifiedToOriginal);
mInfo.pipelineStruct.internal = &pipelineStruct;
mInfo.pipelineStruct.external =
modified ? machinery.modifiedStruct : &pipelineStruct;
}
}
}
else
{
mInfo.pipelineStruct.internal = &pipelineStruct;
mInfo.pipelineStruct.external = &pipelineStruct;
}
root.insertChildNodes(FindMainIndex(&root), seq);
return true;
}
private:
PreResult visitFunctionDefinitionPre(TIntermFunctionDefinition &node) override
{
return {node, VisitBits::Neither};
}
PostResult visitDeclarationPost(TIntermDeclaration &declNode) override
{
Declaration decl = ViewDeclaration(declNode);
const TVariable &var = decl.symbol.variable();
if (mPipeline.uses(var))
{
ASSERT(mInfo.pipelineVariables.find(&var) == mInfo.pipelineVariables.end());
mInfo.pipelineVariables.insert(&var);
mPipelineVariableList.push_back(&var);
return nullptr;
}
return declNode;
}
const TStructure &createInternalPipelineStruct(TIntermBlock &root, TIntermSequence &outDeclSeq)
{
auto &fields = *new TFieldList();
switch (mPipeline.type)
{
case Pipeline::Type::Texture:
{
for (const TVariable *var : mPipelineVariableList)
{
ASSERT(!mInvariants.contains(*var));
const TType &varType = var->getType();
const TBasicType samplerType = varType.getBasicType();
const TStructure &textureEnv = mSymbolEnv.getTextureEnv(samplerType);
auto *textureEnvType = new TType(&textureEnv, false);
if (varType.isArray())
{
textureEnvType->makeArrays(varType.getArraySizes());
}
fields.push_back(
new TField(textureEnvType, var->name(), kNoSourceLoc, var->symbolType()));
}
}
break;
case Pipeline::Type::UniformBuffer:
{
for (const TVariable *var : mPipelineVariableList)
{
auto &type = CloneType(var->getType());
auto *field = new TField(&type, var->name(), kNoSourceLoc, var->symbolType());
mSymbolEnv.markAsPointer(*field, AddressSpace::Constant);
mSymbolEnv.markAsUBO(*field);
mSymbolEnv.markAsPointer(*var, AddressSpace::Constant);
fields.push_back(field);
}
}
break;
default:
{
for (const TVariable *var : mPipelineVariableList)
{
auto &type = CloneType(var->getType());
auto *field = new TField(&type, var->name(), kNoSourceLoc, var->symbolType());
fields.push_back(field);
if (mInvariants.contains(*var))
{
mInvariants.insert(*field);
}
}
}
break;
}
Name pipelineStructName = mPipeline.getStructTypeName(Pipeline::Variant::Original);
auto &s = *new TStructure(&mSymbolTable, pipelineStructName.rawName(), &fields,
pipelineStructName.symbolType());
outDeclSeq.push_back(new TIntermDeclaration{&CreateStructTypeVariable(mSymbolTable, s)});
return s;
}
};
////////////////////////////////////////////////////////////////////////////////
PipelineScoped<TVariable> CreatePipelineMainLocalVar(TSymbolTable &symbolTable,
const Pipeline &pipeline,
PipelineScoped<TStructure> pipelineStruct)
{
ASSERT(pipelineStruct.isTotallyFull());
PipelineScoped<TVariable> pipelineMainLocalVar;
auto populateExternalMainLocalVar = [&]() {
ASSERT(!pipelineMainLocalVar.external);
pipelineMainLocalVar.external = &CreateInstanceVariable(
symbolTable, *pipelineStruct.external,
pipeline.getStructInstanceName(pipelineStruct.isUniform()
? Pipeline::Variant::Original
: Pipeline::Variant::Modified));
};
auto populateDistinctInternalMainLocalVar = [&]() {
ASSERT(!pipelineMainLocalVar.internal);
pipelineMainLocalVar.internal =
&CreateInstanceVariable(symbolTable, *pipelineStruct.internal,
pipeline.getStructInstanceName(Pipeline::Variant::Original));
};
if (pipeline.type == Pipeline::Type::InstanceId)
{
populateDistinctInternalMainLocalVar();
}
else if (pipeline.alwaysRequiresLocalVariableDeclarationInMain())
{
populateExternalMainLocalVar();
if (pipelineStruct.isUniform())
{
pipelineMainLocalVar.internal = pipelineMainLocalVar.external;
}
else
{
populateDistinctInternalMainLocalVar();
}
}
else if (!pipelineStruct.isUniform())
{
populateDistinctInternalMainLocalVar();
}
return pipelineMainLocalVar;
}
class PipelineFunctionEnv
{
private:
TCompiler &mCompiler;
SymbolEnv &mSymbolEnv;
TSymbolTable &mSymbolTable;
IdGen &mIdGen;
const Pipeline &mPipeline;
const std::unordered_set<const TFunction *> &mPipelineFunctions;
const PipelineScoped<TStructure> mPipelineStruct;
PipelineScoped<TVariable> &mPipelineMainLocalVar;
std::unordered_map<const TFunction *, const TFunction *> mFuncMap;
public:
PipelineFunctionEnv(TCompiler &compiler,
SymbolEnv &symbolEnv,
IdGen &idGen,
const Pipeline &pipeline,
const std::unordered_set<const TFunction *> &pipelineFunctions,
PipelineScoped<TStructure> pipelineStruct,
PipelineScoped<TVariable> &pipelineMainLocalVar)
: mCompiler(compiler),
mSymbolEnv(symbolEnv),
mSymbolTable(symbolEnv.symbolTable()),
mIdGen(idGen),
mPipeline(pipeline),
mPipelineFunctions(pipelineFunctions),
mPipelineStruct(pipelineStruct),
mPipelineMainLocalVar(pipelineMainLocalVar)
{}
bool isOriginalPipelineFunction(const TFunction &func) const
{
return mPipelineFunctions.find(&func) != mPipelineFunctions.end();
}
bool isUpdatedPipelineFunction(const TFunction &func) const
{
auto it = mFuncMap.find(&func);
if (it == mFuncMap.end())
{
return false;
}
return &func == it->second;
}
const TFunction &getUpdatedFunction(const TFunction &func)
{
ASSERT(isOriginalPipelineFunction(func) || isUpdatedPipelineFunction(func));
const TFunction *newFunc;
auto it = mFuncMap.find(&func);
if (it == mFuncMap.end())
{
const bool isMain = func.isMain();
if (isMain && mPipeline.isPipelineOut())
{
ASSERT(func.getReturnType().getBasicType() == TBasicType::EbtVoid);
newFunc = &CloneFunctionAndChangeReturnType(mSymbolTable, nullptr, func,
*mPipelineStruct.external);
}
else if (isMain && (mPipeline.type == Pipeline::Type::InvocationVertexGlobals ||
mPipeline.type == Pipeline::Type::InvocationFragmentGlobals))
{
std::vector<const TVariable *> variables;
for (const TField *field : mPipelineStruct.external->fields())
{
variables.push_back(new TVariable(&mSymbolTable, field->name(), field->type(),
field->symbolType()));
}
newFunc = &CloneFunctionAndAppendParams(mSymbolTable, nullptr, func, variables);
}
else if (isMain && mPipeline.type == Pipeline::Type::Texture)
{
std::vector<const TVariable *> variables;
TranslatorMetalReflection *reflection =
((sh::TranslatorMetalDirect *)&mCompiler)->getTranslatorMetalReflection();
for (const TField *field : mPipelineStruct.external->fields())
{
const TStructure *textureEnv = field->type()->getStruct();
ASSERT(textureEnv && textureEnv->fields().size() == 2);
for (const TField *subfield : textureEnv->fields())
{
const Name name = mIdGen.createNewName({field->name(), subfield->name()});
TType &type = *new TType(*subfield->type());
ASSERT(!type.isArray());
type.makeArrays(field->type()->getArraySizes());
auto *var =
new TVariable(&mSymbolTable, name.rawName(), &type, name.symbolType());
variables.push_back(var);
reflection->addOriginalName(var->uniqueId().get(), field->name().data());
}
}
newFunc = &CloneFunctionAndAppendParams(mSymbolTable, nullptr, func, variables);
}
else if (isMain && mPipeline.type == Pipeline::Type::InstanceId)
{
Name name = mPipeline.getStructInstanceName(Pipeline::Variant::Modified);
auto *var = new TVariable(&mSymbolTable, name.rawName(),
new TType(TBasicType::EbtUInt), name.symbolType());
newFunc = &CloneFunctionAndPrependParam(mSymbolTable, nullptr, func, *var);
mPipelineMainLocalVar.external = var;
}
else if (isMain && mPipeline.alwaysRequiresLocalVariableDeclarationInMain())
{
ASSERT(mPipelineMainLocalVar.isTotallyFull());
newFunc = &func;
}
else
{
const TVariable *var;
AddressSpace addressSpace;
if (isMain && !mPipelineMainLocalVar.isUniform())
{
var = &CreateInstanceVariable(
mSymbolTable, *mPipelineStruct.external,
mPipeline.getStructInstanceName(Pipeline::Variant::Modified));
addressSpace = mPipeline.externalAddressSpace();
}
else
{
if (mPipeline.type == Pipeline::Type::UniformBuffer)
{
TranslatorMetalReflection *reflection =
((sh::TranslatorMetalDirect *)&mCompiler)
->getTranslatorMetalReflection();
// TODO: need more checks to make sure they line up? Could be reordered?
ASSERT(mPipelineStruct.external->fields().size() ==
mPipelineStruct.internal->fields().size());
for (size_t i = 0; i < mPipelineStruct.external->fields().size(); i++)
{
const TField *externalField = mPipelineStruct.external->fields()[i];
const TField *internalField = mPipelineStruct.internal->fields()[i];
const TType &externalType = *externalField->type();
const TType &internalType = *internalField->type();
ASSERT(externalType.getBasicType() == internalType.getBasicType());
if (externalType.getBasicType() == TBasicType::EbtStruct)
{
const TStructure *externalEnv = externalType.getStruct();
const TStructure *internalEnv = internalType.getStruct();
const std::string internalName =
reflection->getOriginalName(internalEnv->uniqueId().get());
reflection->addOriginalName(externalEnv->uniqueId().get(),
internalName);
}
}
}
var = &CreateInstanceVariable(
mSymbolTable, *mPipelineStruct.internal,
mPipeline.getStructInstanceName(Pipeline::Variant::Original));
addressSpace = mPipelineMainLocalVar.isUniform()
? mPipeline.externalAddressSpace()
: AddressSpace::Thread;
}
bool markAsReference = true;
if (isMain)
{
switch (mPipeline.type)
{
case Pipeline::Type::VertexIn:
case Pipeline::Type::FragmentIn:
markAsReference = false;
break;
default:
break;
}
}
if (markAsReference)
{
mSymbolEnv.markAsReference(*var, addressSpace);
}
newFunc = &CloneFunctionAndPrependParam(mSymbolTable, nullptr, func, *var);
}
mFuncMap[&func] = newFunc;
mFuncMap[newFunc] = newFunc;
}
else
{
newFunc = it->second;
}
return *newFunc;
}
TIntermFunctionPrototype *createUpdatedFunctionPrototype(
TIntermFunctionPrototype &funcProtoNode)
{
const TFunction &func = *funcProtoNode.getFunction();
if (!isOriginalPipelineFunction(func) && !isUpdatedPipelineFunction(func))
{
return nullptr;
}
const TFunction &newFunc = getUpdatedFunction(func);
return new TIntermFunctionPrototype(&newFunc);
}
};
class UpdatePipelineFunctions : private TIntermRebuild
{
private:
const Pipeline &mPipeline;
const PipelineScoped<TStructure> mPipelineStruct;
PipelineScoped<TVariable> &mPipelineMainLocalVar;
SymbolEnv &mSymbolEnv;
PipelineFunctionEnv mEnv;
const TFunction *mFuncOriginalToModified;
const TFunction *mFuncModifiedToOriginal;
public:
static bool ThreadPipeline(TCompiler &compiler,
TIntermBlock &root,
const Pipeline &pipeline,
const std::unordered_set<const TFunction *> &pipelineFunctions,
PipelineScoped<TStructure> pipelineStruct,
PipelineScoped<TVariable> &pipelineMainLocalVar,
IdGen &idGen,
SymbolEnv &symbolEnv,
const TFunction *funcOriginalToModified,
const TFunction *funcModifiedToOriginal)
{
UpdatePipelineFunctions self(compiler, pipeline, pipelineFunctions, pipelineStruct,
pipelineMainLocalVar, idGen, symbolEnv, funcOriginalToModified,
funcModifiedToOriginal);
if (!self.rebuildRoot(root))
{
return false;
}
return true;
}
private:
UpdatePipelineFunctions(TCompiler &compiler,
const Pipeline &pipeline,
const std::unordered_set<const TFunction *> &pipelineFunctions,
PipelineScoped<TStructure> pipelineStruct,
PipelineScoped<TVariable> &pipelineMainLocalVar,
IdGen &idGen,
SymbolEnv &symbolEnv,
const TFunction *funcOriginalToModified,
const TFunction *funcModifiedToOriginal)
: TIntermRebuild(compiler, false, true),
mPipeline(pipeline),
mPipelineStruct(pipelineStruct),
mPipelineMainLocalVar(pipelineMainLocalVar),
mSymbolEnv(symbolEnv),
mEnv(compiler,
symbolEnv,
idGen,
pipeline,
pipelineFunctions,
pipelineStruct,
mPipelineMainLocalVar),
mFuncOriginalToModified(funcOriginalToModified),
mFuncModifiedToOriginal(funcModifiedToOriginal)
{
ASSERT(mPipelineStruct.isTotallyFull());
}
const TVariable &getInternalPipelineVariable(const TFunction &pipelineFunc)
{
if (pipelineFunc.isMain() && (mPipeline.alwaysRequiresLocalVariableDeclarationInMain() ||
!mPipelineMainLocalVar.isUniform()))
{
ASSERT(mPipelineMainLocalVar.internal);
return *mPipelineMainLocalVar.internal;
}
else
{
ASSERT(pipelineFunc.getParamCount() > 0);
return *pipelineFunc.getParam(0);
}
}
const TVariable &getExternalPipelineVariable(const TFunction &mainFunc)
{
ASSERT(mainFunc.isMain());
if (mPipelineMainLocalVar.external)
{
return *mPipelineMainLocalVar.external;
}
else
{
ASSERT(mainFunc.getParamCount() > 0);
return *mainFunc.getParam(0);
}
}
PostResult visitAggregatePost(TIntermAggregate &callNode) override
{
if (callNode.isConstructor())
{
return callNode;
}
else
{
const TFunction &oldCalledFunc = *callNode.getFunction();
if (!mEnv.isOriginalPipelineFunction(oldCalledFunc))
{
return callNode;
}
const TFunction &newCalledFunc = mEnv.getUpdatedFunction(oldCalledFunc);
const TFunction *oldOwnerFunc = getParentFunction();
ASSERT(oldOwnerFunc);
const TFunction &newOwnerFunc = mEnv.getUpdatedFunction(*oldOwnerFunc);
return *TIntermAggregate::CreateFunctionCall(
newCalledFunc, &CloneSequenceAndPrepend(
*callNode.getSequence(),
*new TIntermSymbol(&getInternalPipelineVariable(newOwnerFunc))));
}
}
PostResult visitFunctionPrototypePost(TIntermFunctionPrototype &funcProtoNode) override
{
TIntermFunctionPrototype *newFuncProtoNode =
mEnv.createUpdatedFunctionPrototype(funcProtoNode);
if (newFuncProtoNode == nullptr)
{
return funcProtoNode;
}
return *newFuncProtoNode;
}
PostResult visitFunctionDefinitionPost(TIntermFunctionDefinition &funcDefNode) override
{
if (funcDefNode.getFunction()->isMain())
{
return visitMain(funcDefNode);
}
else
{
return visitNonMain(funcDefNode);
}
}
TIntermNode &visitNonMain(TIntermFunctionDefinition &funcDefNode)
{
TIntermFunctionPrototype &funcProtoNode = *funcDefNode.getFunctionPrototype();
ASSERT(!funcProtoNode.getFunction()->isMain());
TIntermFunctionPrototype *newFuncProtoNode =
mEnv.createUpdatedFunctionPrototype(funcProtoNode);
if (newFuncProtoNode == nullptr)
{
return funcDefNode;
}
const TFunction &func = *newFuncProtoNode->getFunction();
ASSERT(!func.isMain());
TIntermBlock *body = funcDefNode.getBody();
return *new TIntermFunctionDefinition(newFuncProtoNode, body);
}
TIntermNode &visitMain(TIntermFunctionDefinition &funcDefNode)
{
TIntermFunctionPrototype &funcProtoNode = *funcDefNode.getFunctionPrototype();
ASSERT(funcProtoNode.getFunction()->isMain());
TIntermFunctionPrototype *newFuncProtoNode =
mEnv.createUpdatedFunctionPrototype(funcProtoNode);
if (newFuncProtoNode == nullptr)
{
return funcDefNode;
}
const TFunction &func = *newFuncProtoNode->getFunction();
ASSERT(func.isMain());
auto callModifiedToOriginal = [&](TIntermBlock &body) {
ASSERT(mPipelineMainLocalVar.internal);
if (!mPipeline.isPipelineOut())
{
ASSERT(mFuncModifiedToOriginal);
auto *m = new TIntermSymbol(&getExternalPipelineVariable(func));
auto *o = new TIntermSymbol(mPipelineMainLocalVar.internal);
body.appendStatement(TIntermAggregate::CreateFunctionCall(
*mFuncModifiedToOriginal, new TIntermSequence{m, o}));
}
};
auto callOriginalToModified = [&](TIntermBlock &body) {
ASSERT(mPipelineMainLocalVar.internal);
if (mPipeline.isPipelineOut())
{
ASSERT(mFuncOriginalToModified);
auto *o = new TIntermSymbol(mPipelineMainLocalVar.internal);
auto *m = new TIntermSymbol(&getExternalPipelineVariable(func));
body.appendStatement(TIntermAggregate::CreateFunctionCall(
*mFuncOriginalToModified, new TIntermSequence{o, m}));
}
};
TIntermBlock *body = funcDefNode.getBody();
if (mPipeline.alwaysRequiresLocalVariableDeclarationInMain())
{
ASSERT(mPipelineMainLocalVar.isTotallyFull());
auto *newBody = new TIntermBlock();
newBody->appendStatement(new TIntermDeclaration{mPipelineMainLocalVar.internal});
if (mPipeline.type == Pipeline::Type::InvocationVertexGlobals ||
mPipeline.type == Pipeline::Type::InvocationFragmentGlobals)
{
// Populate struct instance with references to global pipeline variables.
for (const TField *field : mPipelineStruct.external->fields())
{
auto *var = new TVariable(&mSymbolTable, field->name(), field->type(),
field->symbolType());
auto *symbol = new TIntermSymbol(var);
auto &accessNode = AccessField(*mPipelineMainLocalVar.internal, var->name());
auto *assignNode = new TIntermBinary(TOperator::EOpAssign, &accessNode, symbol);
newBody->appendStatement(assignNode);
}
}
else if (mPipeline.type == Pipeline::Type::Texture)
{
const TFieldList &fields = mPipelineStruct.external->fields();
ASSERT(func.getParamCount() >= 2 * fields.size());
size_t paramIndex = func.getParamCount() - 2 * fields.size();
for (const TField *field : fields)
{
const TVariable &textureParam = *func.getParam(paramIndex++);
const TVariable &samplerParam = *func.getParam(paramIndex++);
auto go = [&](TIntermTyped &env, const int *index) {
TIntermTyped &textureField = AccessField(
AccessIndex(*env.deepCopy(), index), ImmutableString("texture"));
TIntermTyped &samplerField = AccessField(
AccessIndex(*env.deepCopy(), index), ImmutableString("sampler"));
auto mkAssign = [&](TIntermTyped &field, const TVariable &param) {
return new TIntermBinary(TOperator::EOpAssign, &field,
&mSymbolEnv.callFunctionOverload(
Name("addressof"), field.getType(),
*new TIntermSequence{&AccessIndex(
*new TIntermSymbol(&param), index)}));
};
newBody->appendStatement(mkAssign(textureField, textureParam));
newBody->appendStatement(mkAssign(samplerField, samplerParam));
};
TIntermTyped &env = AccessField(*mPipelineMainLocalVar.internal, field->name());
const TType &envType = env.getType();
if (envType.isArray())
{
ASSERT(!envType.isArrayOfArrays());
const auto n = static_cast<int>(envType.getArraySizeProduct());
for (int i = 0; i < n; ++i)
{
go(env, &i);
}
}
else
{
go(env, nullptr);
}
}
}
else if (mPipeline.type == Pipeline::Type::InstanceId)
{
newBody->appendStatement(new TIntermBinary(
TOperator::EOpAssign,
&AccessFieldByIndex(*new TIntermSymbol(&getInternalPipelineVariable(func)), 0),
&AsType(mSymbolEnv, *new TType(TBasicType::EbtInt),
*new TIntermSymbol(&getExternalPipelineVariable(func)))));
}
else if (!mPipelineMainLocalVar.isUniform())
{
newBody->appendStatement(new TIntermDeclaration{mPipelineMainLocalVar.external});
callModifiedToOriginal(*newBody);
}
newBody->appendStatement(body);
if (!mPipelineMainLocalVar.isUniform())
{
callOriginalToModified(*newBody);
}
if (mPipeline.isPipelineOut())
{
newBody->appendStatement(new TIntermBranch(
TOperator::EOpReturn, new TIntermSymbol(mPipelineMainLocalVar.external)));
}
body = newBody;
}
else if (!mPipelineMainLocalVar.isUniform())
{
ASSERT(!mPipelineMainLocalVar.external);
ASSERT(mPipelineMainLocalVar.internal);
auto *newBody = new TIntermBlock();
newBody->appendStatement(new TIntermDeclaration{mPipelineMainLocalVar.internal});
callModifiedToOriginal(*newBody);
newBody->appendStatement(body);
callOriginalToModified(*newBody);
body = newBody;
}
return *new TIntermFunctionDefinition(newFuncProtoNode, body);
}
};
////////////////////////////////////////////////////////////////////////////////
bool UpdatePipelineSymbols(Pipeline::Type pipelineType,
TCompiler &compiler,
TIntermBlock &root,
SymbolEnv &symbolEnv,
const VariableSet &pipelineVariables,
PipelineScoped<TVariable> pipelineMainLocalVar)
{
auto map = [&](const TFunction *owner, TIntermSymbol &symbol) -> TIntermNode & {
const TVariable &var = symbol.variable();
if (pipelineVariables.find(&var) == pipelineVariables.end())
{
return symbol;
}
ASSERT(owner);
const TVariable *structInstanceVar;
if (owner->isMain())
{
ASSERT(pipelineMainLocalVar.internal);
structInstanceVar = pipelineMainLocalVar.internal;
}
else
{
ASSERT(owner->getParamCount() > 0);
structInstanceVar = owner->getParam(0);
}
ASSERT(structInstanceVar);
return AccessField(*structInstanceVar, var.name());
};
return MapSymbols(compiler, root, map);
}
////////////////////////////////////////////////////////////////////////////////
bool RewritePipeline(TCompiler &compiler,
TIntermBlock &root,
IdGen &idGen,
const Pipeline &pipeline,
SymbolEnv &symbolEnv,
Invariants &invariants,
PipelineScoped<TStructure> &outStruct)
{
ASSERT(outStruct.isTotallyEmpty());
TSymbolTable &symbolTable = compiler.getSymbolTable();
PipelineStructInfo psi;
if (!GeneratePipelineStruct::Exec(psi, compiler, root, idGen, pipeline, symbolEnv, invariants))
{
return false;
}
if (psi.isEmpty())
{
return true;
}
const auto pipelineFunctions = DiscoverDependentFunctions(root, [&](const TVariable &var) {
return psi.pipelineVariables.find(&var) != psi.pipelineVariables.end();
});
auto pipelineMainLocalVar =
CreatePipelineMainLocalVar(symbolTable, pipeline, psi.pipelineStruct);
if (!UpdatePipelineFunctions::ThreadPipeline(
compiler, root, pipeline, pipelineFunctions, psi.pipelineStruct, pipelineMainLocalVar,
idGen, symbolEnv, psi.funcOriginalToModified, psi.funcModifiedToOriginal))
{
return false;
}
if (!pipeline.globalInstanceVar)
{
if (!UpdatePipelineSymbols(pipeline.type, compiler, root, symbolEnv, psi.pipelineVariables,
pipelineMainLocalVar))
{
return false;
}
}
if (!PruneNoOps(&compiler, &root, &compiler.getSymbolTable()))
{
return false;
}
outStruct = psi.pipelineStruct;
return true;
}
} // anonymous namespace
bool sh::RewritePipelines(TCompiler &compiler,
TIntermBlock &root,
IdGen &idGen,
DriverUniform &angleUniformsGlobalInstanceVar,
SymbolEnv &symbolEnv,
Invariants &invariants,
PipelineStructs &outStructs)
{
struct Info
{
Pipeline::Type pipelineType;
PipelineScoped<TStructure> &outStruct;
const TVariable *globalInstanceVar;
};
Info infos[] = {
{Pipeline::Type::InstanceId, outStructs.instanceId, nullptr},
{Pipeline::Type::Texture, outStructs.texture, nullptr},
{Pipeline::Type::NonConstantGlobals, outStructs.nonConstantGlobals, nullptr},
{Pipeline::Type::AngleUniforms, outStructs.angleUniforms,
angleUniformsGlobalInstanceVar.getDriverUniformsVariable()},
{Pipeline::Type::UserUniforms, outStructs.userUniforms, nullptr},
{Pipeline::Type::VertexIn, outStructs.vertexIn, nullptr},
{Pipeline::Type::VertexOut, outStructs.vertexOut, nullptr},
{Pipeline::Type::FragmentIn, outStructs.fragmentIn, nullptr},
{Pipeline::Type::FragmentOut, outStructs.fragmentOut, nullptr},
{Pipeline::Type::InvocationVertexGlobals, outStructs.invocationVertexGlobals, nullptr},
{Pipeline::Type::InvocationFragmentGlobals, outStructs.invocationFragmentGlobals, nullptr},
{Pipeline::Type::UniformBuffer, outStructs.uniformBuffers, nullptr},
};
for (Info &info : infos)
{
Pipeline pipeline{info.pipelineType, info.globalInstanceVar};
if (!RewritePipeline(compiler, root, idGen, pipeline, symbolEnv, invariants,
info.outStruct))
{
return false;
}
}
return true;
}

45
src/compiler/translator/TranslatorMetalDirect/RewritePipelines.h Normal file
Просмотреть файл

@ -0,0 +1,45 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEPIPELINES_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEPIPELINES_H_
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/IdGen.h"
#include "compiler/translator/TranslatorMetalDirect/Pipeline.h"
#include "compiler/translator/TranslatorMetalDirect/RewriteGlobalQualifierDecls.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
namespace sh
{
// This rewrites all pipelines.
//
// For each pipeline:
// - Discover all variables that are used by the pipeline
// - Move the variables into an internal pipeline struct instance and update old variables to be
// member access instead.
// - Dependency inject the internal pipeline struct to all functions that access variables from
// the struct.
// - A new external pipeline struct is created if needed for impedance reasons. Otherwise the
// external and internal pipeline structs are the same.
// - Add `main` parameter or return value for the external pipeline struct as needed.
// - Inside `main`, map the external struct to the internal struct if they differ and is provided
// as a parameter to `main`.
// - Inside `main`, map the internal struct to the external struct if they differ and is returned
// from `main`.
ANGLE_NO_DISCARD bool RewritePipelines(TCompiler &compiler,
TIntermBlock &root,
IdGen &idGen,
DriverUniform &angleUniformsGlobalInstanceVar,
SymbolEnv &symbolEnv,
Invariants &invariants,
PipelineStructs &outStructs);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEPIPELINES_H_

375
src/compiler/translator/TranslatorMetalDirect/RewriteUnaddressableReferences.cpp Normal file
Просмотреть файл

@ -0,0 +1,375 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/RewriteUnaddressableReferences.h"
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
#include "compiler/translator/tree_util/AsNode.h"
using namespace sh;
namespace
{
bool IsOutParam(const TType &paramType)
{
const TQualifier qual = paramType.getQualifier();
switch (qual)
{
case TQualifier::EvqInOut:
case TQualifier::EvqOut:
return true;
default:
return false;
}
}
bool IsVectorAccess(TIntermBinary &binary)
{
TOperator op = binary.getOp();
switch (op)
{
case TOperator::EOpIndexDirect:
case TOperator::EOpIndexIndirect:
break;
default:
return false;
}
const TType &leftType = binary.getLeft()->getType();
if (!leftType.isVector() || leftType.isArray())
{
return false;
}
ASSERT(IsScalarBasicType(binary.getType()));
return true;
}
bool IsVectorAccess(TIntermNode &node)
{
if (auto *bin = node.getAsBinaryNode())
{
return IsVectorAccess(*bin);
}
return false;
}
// Differs from IsAssignment in that it does not include (++) or (--).
bool IsAssignEqualsSign(TOperator op)
{
switch (op)
{
case TOperator::EOpAssign:
case TOperator::EOpInitialize:
case TOperator::EOpAddAssign:
case TOperator::EOpSubAssign:
case TOperator::EOpMulAssign:
case TOperator::EOpVectorTimesMatrixAssign:
case TOperator::EOpVectorTimesScalarAssign:
case TOperator::EOpMatrixTimesScalarAssign:
case TOperator::EOpMatrixTimesMatrixAssign:
case TOperator::EOpDivAssign:
case TOperator::EOpIModAssign:
case TOperator::EOpBitShiftLeftAssign:
case TOperator::EOpBitShiftRightAssign:
case TOperator::EOpBitwiseAndAssign:
case TOperator::EOpBitwiseXorAssign:
case TOperator::EOpBitwiseOrAssign:
return true;
default:
return false;
}
}
// Only includes ($=) style assigns, where ($) is a binary op.
bool IsCompoundAssign(TOperator op)
{
switch (op)
{
case TOperator::EOpAddAssign:
case TOperator::EOpSubAssign:
case TOperator::EOpMulAssign:
case TOperator::EOpVectorTimesMatrixAssign:
case TOperator::EOpVectorTimesScalarAssign:
case TOperator::EOpMatrixTimesScalarAssign:
case TOperator::EOpMatrixTimesMatrixAssign:
case TOperator::EOpDivAssign:
case TOperator::EOpIModAssign:
case TOperator::EOpBitShiftLeftAssign:
case TOperator::EOpBitShiftRightAssign:
case TOperator::EOpBitwiseAndAssign:
case TOperator::EOpBitwiseXorAssign:
case TOperator::EOpBitwiseOrAssign:
return true;
default:
return false;
}
}
bool ReturnsReference(TOperator op)
{
switch (op)
{
case TOperator::EOpAssign:
case TOperator::EOpInitialize:
case TOperator::EOpAddAssign:
case TOperator::EOpSubAssign:
case TOperator::EOpMulAssign:
case TOperator::EOpVectorTimesMatrixAssign:
case TOperator::EOpVectorTimesScalarAssign:
case TOperator::EOpMatrixTimesScalarAssign:
case TOperator::EOpMatrixTimesMatrixAssign:
case TOperator::EOpDivAssign:
case TOperator::EOpIModAssign:
case TOperator::EOpBitShiftLeftAssign:
case TOperator::EOpBitShiftRightAssign:
case TOperator::EOpBitwiseAndAssign:
case TOperator::EOpBitwiseXorAssign:
case TOperator::EOpBitwiseOrAssign:
case TOperator::EOpPostIncrement:
case TOperator::EOpPostDecrement:
case TOperator::EOpPreIncrement:
case TOperator::EOpPreDecrement:
case TOperator::EOpIndexDirect:
case TOperator::EOpIndexIndirect:
case TOperator::EOpIndexDirectStruct:
case TOperator::EOpIndexDirectInterfaceBlock:
return true;
default:
return false;
}
}
TIntermTyped &DecomposeCompoundAssignment(TIntermBinary &node)
{
TOperator op = node.getOp();
switch (op)
{
case TOperator::EOpAddAssign:
op = TOperator::EOpAdd;
break;
case TOperator::EOpSubAssign:
op = TOperator::EOpSub;
break;
case TOperator::EOpMulAssign:
op = TOperator::EOpMul;
break;
case TOperator::EOpVectorTimesMatrixAssign:
op = TOperator::EOpVectorTimesMatrix;
break;
case TOperator::EOpVectorTimesScalarAssign:
op = TOperator::EOpVectorTimesScalar;
break;
case TOperator::EOpMatrixTimesScalarAssign:
op = TOperator::EOpMatrixTimesScalar;
break;
case TOperator::EOpMatrixTimesMatrixAssign:
op = TOperator::EOpMatrixTimesMatrix;
break;
case TOperator::EOpDivAssign:
op = TOperator::EOpDiv;
break;
case TOperator::EOpIModAssign:
op = TOperator::EOpIMod;
break;
case TOperator::EOpBitShiftLeftAssign:
op = TOperator::EOpBitShiftLeft;
break;
case TOperator::EOpBitShiftRightAssign:
op = TOperator::EOpBitShiftRight;
break;
case TOperator::EOpBitwiseAndAssign:
op = TOperator::EOpBitwiseAnd;
break;
case TOperator::EOpBitwiseXorAssign:
op = TOperator::EOpBitwiseXor;
break;
case TOperator::EOpBitwiseOrAssign:
op = TOperator::EOpBitwiseOr;
break;
default:
UNREACHABLE();
}
// This assumes SeparateCompoundExpressions has already been called.
// This assumption allows this code to not need to introduce temporaries.
//
// e.g. dont have to worry about:
// vec[hasSideEffect()] *= 4
// becoming
// vec[hasSideEffect()] = vec[hasSideEffect()] * 4
TIntermTyped *left = node.getLeft();
TIntermTyped *right = node.getRight();
return *new TIntermBinary(TOperator::EOpAssign, left->deepCopy(),
new TIntermBinary(op, left, right));
}
class Rewriter1 : public TIntermRebuild
{
public:
Rewriter1(TCompiler &compiler) : TIntermRebuild(compiler, false, true) {}
PostResult visitBinaryPost(TIntermBinary &binaryNode) override
{
const TOperator op = binaryNode.getOp();
if (IsCompoundAssign(op))
{
TIntermTyped &left = *binaryNode.getLeft();
if (left.getAsSwizzleNode() || IsVectorAccess(left))
{
return DecomposeCompoundAssignment(binaryNode);
}
}
return binaryNode;
}
};
class Rewriter2 : public TIntermRebuild
{
std::vector<bool> mRequiresAddressingStack;
SymbolEnv &mSymbolEnv;
private:
bool requiresAddressing() const
{
if (mRequiresAddressingStack.empty())
{
return false;
}
return mRequiresAddressingStack.back();
}
public:
~Rewriter2() override { ASSERT(mRequiresAddressingStack.empty()); }
Rewriter2(TCompiler &compiler, SymbolEnv &symbolEnv)
: TIntermRebuild(compiler, true, true), mSymbolEnv(symbolEnv)
{}
PreResult visitAggregatePre(TIntermAggregate &aggregateNode) override
{
const TFunction *func = aggregateNode.getFunction();
if (!func)
{
return aggregateNode;
}
TIntermSequence &args = *aggregateNode.getSequence();
size_t argCount = args.size();
for (size_t i = 0; i < argCount; ++i)
{
const TVariable &param = *func->getParam(i);
const TType &paramType = param.getType();
TIntermNode *arg = args[i];
ASSERT(arg);
mRequiresAddressingStack.push_back(IsOutParam(paramType));
args[i] = rebuild(*arg).single();
ASSERT(args[i]);
ASSERT(!mRequiresAddressingStack.empty());
mRequiresAddressingStack.pop_back();
}
return {aggregateNode, VisitBits::Neither};
}
PostResult visitSwizzlePost(TIntermSwizzle &swizzleNode) override
{
if (!requiresAddressing())
{
return swizzleNode;
}
TIntermTyped &vecNode = *swizzleNode.getOperand();
const TQualifierList &offsets = swizzleNode.getSwizzleOffsets();
ASSERT(!offsets.empty());
ASSERT(offsets.size() <= 4);
auto &args = *new TIntermSequence();
args.reserve(offsets.size() + 1);
args.push_back(&vecNode);
for (int offset : offsets)
{
args.push_back(new TIntermConstantUnion(new TConstantUnion(offset),
*new TType(TBasicType::EbtInt)));
}
return mSymbolEnv.callFunctionOverload(Name("swizzle_ref"), swizzleNode.getType(), args);
}
PreResult visitBinaryPre(TIntermBinary &binaryNode) override
{
const TOperator op = binaryNode.getOp();
const bool isAccess = IsVectorAccess(binaryNode);
const bool disableTop = !ReturnsReference(op) || !requiresAddressing();
const bool disableLeft = disableTop;
const bool disableRight = disableTop || isAccess || IsAssignEqualsSign(op);
auto traverse = [&](TIntermTyped &node, const bool disable) -> TIntermTyped & {
if (disable)
{
mRequiresAddressingStack.push_back(false);
}
auto *newNode = asNode<TIntermTyped>(rebuild(node).single());
ASSERT(newNode);
if (disable)
{
mRequiresAddressingStack.pop_back();
}
return *newNode;
};
TIntermTyped &leftNode = *binaryNode.getLeft();
TIntermTyped &rightNode = *binaryNode.getRight();
TIntermTyped &newLeft = traverse(leftNode, disableLeft);
TIntermTyped &newRight = traverse(rightNode, disableRight);
if (!isAccess || disableTop)
{
if (&leftNode == &newLeft && &rightNode == &newRight)
{
return {&binaryNode, VisitBits::Neither};
}
return {*new TIntermBinary(op, &newLeft, &newRight), VisitBits::Neither};
}
return {mSymbolEnv.callFunctionOverload(Name("elem_ref"), binaryNode.getType(),
*new TIntermSequence{&newLeft, &newRight}),
VisitBits::Neither};
}
};
} // anonymous namespace
bool sh::RewriteUnaddressableReferences(TCompiler &compiler,
TIntermBlock &root,
SymbolEnv &symbolEnv)
{
if (!Rewriter1(compiler).rebuildRoot(root))
{
return false;
}
if (!Rewriter2(compiler, symbolEnv).rebuildRoot(root))
{
return false;
}
return true;
}

31
src/compiler/translator/TranslatorMetalDirect/RewriteUnaddressableReferences.h Normal file
Просмотреть файл

@ -0,0 +1,31 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEUNADDRESSABLEREFERENCES_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEUNADDRESSABLEREFERENCES_H_
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/ProgramPrelude.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
namespace sh
{
// Given:
// void foo(out x);
// It is possible for the following to be legal in GLSL but not in Metal:
// foo(expression);
// This can happen in cases where `expression` is a vector swizzle or vector element access.
// This rewrite functionality introduces temporaries that serve as proxies to be passed to the
// out/inout parameters as needed. The corresponding expressions get populated with their
// proxies after the function call.
ANGLE_NO_DISCARD bool RewriteUnaddressableReferences(TCompiler &compiler,
TIntermBlock &root,
SymbolEnv &symbolEnv);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_REWRITEUNADDRESSABLEREFERENCES_H_

657
src/compiler/translator/TranslatorMetalDirect/SeparateCompoundExpressions.cpp Normal file
Просмотреть файл

@ -0,0 +1,657 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <unordered_map>
#include "common/system_utils.h"
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
#include "compiler/translator/TranslatorMetalDirect/SeparateCompoundExpressions.h"
#include "compiler/translator/tree_ops/SimplifyLoopConditions.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
bool IsIndex(TOperator op)
{
switch (op)
{
case TOperator::EOpIndexDirect:
case TOperator::EOpIndexDirectInterfaceBlock:
case TOperator::EOpIndexDirectStruct:
case TOperator::EOpIndexIndirect:
return true;
default:
return false;
}
}
bool IsIndex(TIntermTyped &expr)
{
if (auto *binary = expr.getAsBinaryNode())
{
return IsIndex(binary->getOp());
}
return expr.getAsSwizzleNode();
}
bool ViewBinaryChain(TOperator op, TIntermTyped &node, std::vector<TIntermTyped *> &out)
{
TIntermBinary *binary = node.getAsBinaryNode();
if (!binary || binary->getOp() != op)
{
return false;
}
TIntermTyped *left = binary->getLeft();
TIntermTyped *right = binary->getRight();
if (!ViewBinaryChain(op, *left, out))
{
out.push_back(left);
}
if (!ViewBinaryChain(op, *right, out))
{
out.push_back(right);
}
return true;
}
std::vector<TIntermTyped *> ViewBinaryChain(TIntermBinary &node)
{
std::vector<TIntermTyped *> chain;
ViewBinaryChain(node.getOp(), node, chain);
ASSERT(chain.size() >= 2);
return chain;
}
class PrePass : public TIntermRebuild
{
public:
PrePass(TCompiler &compiler) : TIntermRebuild(compiler, true, true) {}
private:
// Change chains of
// x OP y OP z
// to
// x OP (y OP z)
// regardless of original parenthesization.
TIntermTyped &reassociateRight(TIntermBinary &node)
{
const TOperator op = node.getOp();
std::vector<TIntermTyped *> chain = ViewBinaryChain(node);
TIntermTyped *result = chain.back();
chain.pop_back();
ASSERT(result);
const auto begin = chain.rbegin();
const auto end = chain.rend();
for (auto iter = begin; iter != end; ++iter)
{
TIntermTyped *part = *iter;
ASSERT(part);
TIntermNode *temp = rebuild(*part).single();
ASSERT(temp);
part = temp->getAsTyped();
ASSERT(part);
result = new TIntermBinary(op, part, result);
}
return *result;
}
private:
PreResult visitBinaryPre(TIntermBinary &node) override
{
const TOperator op = node.getOp();
if (op == TOperator::EOpLogicalAnd || op == TOperator::EOpLogicalOr)
{
return {reassociateRight(node), VisitBits::Neither};
}
return node;
}
};
class Separator : public TIntermRebuild
{
IdGen &mIdGen;
std::vector<std::vector<TIntermNode *>> mStmtsStack;
std::vector<std::unordered_map<const TVariable *, TIntermDeclaration *>> mBindingMapStack;
std::unordered_map<TIntermTyped *, TIntermTyped *> mExprMap;
std::unordered_set<TIntermDeclaration *> mMaskedDecls;
public:
Separator(TCompiler &compiler, SymbolEnv &symbolEnv, IdGen &idGen)
: TIntermRebuild(compiler, true, true), mIdGen(idGen)
{}
~Separator() override
{
ASSERT(mStmtsStack.empty());
ASSERT(mExprMap.empty());
ASSERT(mBindingMapStack.empty());
}
private:
std::vector<TIntermNode *> &getCurrStmts()
{
ASSERT(!mStmtsStack.empty());
return mStmtsStack.back();
}
std::unordered_map<const TVariable *, TIntermDeclaration *> &getCurrBindingMap()
{
ASSERT(!mBindingMapStack.empty());
return mBindingMapStack.back();
}
void pushStmt(TIntermNode &node) { getCurrStmts().push_back(&node); }
bool isTerminalExpr(TIntermNode &node)
{
NodeType nodeType = getNodeType(node);
switch (nodeType)
{
case NodeType::Symbol:
case NodeType::ConstantUnion:
return true;
default:
return false;
}
}
TIntermTyped *pullMappedExpr(TIntermTyped *node, bool allowBacktrack)
{
TIntermTyped *expr;
{
auto iter = mExprMap.find(node);
if (iter == mExprMap.end())
{
return node;
}
ASSERT(node);
expr = iter->second;
ASSERT(expr);
mExprMap.erase(iter);
}
if (allowBacktrack)
{
auto &bindingMap = getCurrBindingMap();
while (TIntermSymbol *symbol = expr->getAsSymbolNode())
{
const TVariable &var = symbol->variable();
auto iter = bindingMap.find(&var);
if (iter == bindingMap.end())
{
return expr;
}
ASSERT(var.symbolType() == SymbolType::AngleInternal);
TIntermDeclaration *decl = iter->second;
ASSERT(decl);
expr = ViewDeclaration(*decl).initExpr;
ASSERT(expr);
bindingMap.erase(iter);
mMaskedDecls.insert(decl);
}
}
return expr;
}
bool isStandaloneExpr(TIntermTyped &expr)
{
if (getParentNode()->getAsBlock())
{
return true;
}
ASSERT(expr.getType().getBasicType() != TBasicType::EbtVoid);
return false;
}
void pushBinding(TIntermTyped &oldExpr, TIntermTyped &newExpr)
{
if (isStandaloneExpr(newExpr))
{
pushStmt(newExpr);
return;
}
if (IsIndex(newExpr))
{
mExprMap[&oldExpr] = &newExpr;
return;
}
auto &bindingMap = getCurrBindingMap();
const Name name = mIdGen.createNewName("");
auto *var =
new TVariable(&mSymbolTable, name.rawName(), &newExpr.getType(), name.symbolType());
auto *decl = new TIntermDeclaration(var, &newExpr);
pushStmt(*decl);
mExprMap[&oldExpr] = new TIntermSymbol(var);
bindingMap[var] = decl;
}
void pushStacks()
{
mStmtsStack.emplace_back();
mBindingMapStack.emplace_back();
}
void popStacks()
{
ASSERT(!mBindingMapStack.empty());
ASSERT(!mStmtsStack.empty());
ASSERT(mStmtsStack.back().empty());
mBindingMapStack.pop_back();
mStmtsStack.pop_back();
}
void pushStmtsIntoBlock(TIntermBlock &block, std::vector<TIntermNode *> &stmts)
{
TIntermSequence &seq = *block.getSequence();
for (TIntermNode *stmt : stmts)
{
if (TIntermDeclaration *decl = stmt->getAsDeclarationNode())
{
auto iter = mMaskedDecls.find(decl);
if (iter != mMaskedDecls.end())
{
mMaskedDecls.erase(iter);
continue;
}
}
seq.push_back(stmt);
}
}
TIntermBlock &buildBlockWithTailAssign(const TVariable &var, TIntermTyped &newExpr)
{
std::vector<TIntermNode *> stmts = std::move(getCurrStmts());
popStacks();
auto &block = *new TIntermBlock();
auto &seq = *block.getSequence();
seq.reserve(1 + stmts.size());
pushStmtsIntoBlock(block, stmts);
seq.push_back(new TIntermBinary(TOperator::EOpAssign, new TIntermSymbol(&var), &newExpr));
return block;
}
private:
PreResult visitBlockPre(TIntermBlock &node) override
{
pushStacks();
return node;
}
PostResult visitBlockPost(TIntermBlock &node) override
{
std::vector<TIntermNode *> stmts = std::move(getCurrStmts());
popStacks();
TIntermSequence &seq = *node.getSequence();
seq.clear();
seq.reserve(stmts.size());
pushStmtsIntoBlock(node, stmts);
TIntermNode *parent = getParentNode();
if (parent && parent->getAsBlock())
{
pushStmt(node);
}
return node;
}
PreResult visitDeclarationPre(TIntermDeclaration &node) override
{
Declaration decl = ViewDeclaration(node);
if (!decl.initExpr || isTerminalExpr(*decl.initExpr))
{
pushStmt(node);
return {node, VisitBits::Neither};
}
return node;
}
PostResult visitDeclarationPost(TIntermDeclaration &node) override
{
Declaration decl = ViewDeclaration(node);
ASSERT(decl.symbol.variable().symbolType() != SymbolType::Empty);
ASSERT(!decl.symbol.variable().getType().isStructSpecifier());
TIntermTyped *newInitExpr = pullMappedExpr(decl.initExpr, true);
if (decl.initExpr == newInitExpr)
{
pushStmt(node);
}
else
{
auto &newNode = *new TIntermDeclaration();
newNode.appendDeclarator(
new TIntermBinary(TOperator::EOpInitialize, &decl.symbol, newInitExpr));
pushStmt(newNode);
}
return node;
}
PostResult visitUnaryPost(TIntermUnary &node) override
{
TIntermTyped *expr = node.getOperand();
TIntermTyped *newExpr = pullMappedExpr(expr, false);
if (expr == newExpr)
{
pushBinding(node, node);
}
else
{
pushBinding(node, *new TIntermUnary(node.getOp(), newExpr, node.getFunction()));
}
return node;
}
PreResult visitBinaryPre(TIntermBinary &node) override
{
const TOperator op = node.getOp();
if (op == TOperator::EOpLogicalAnd || op == TOperator::EOpLogicalOr)
{
TIntermTyped *left = node.getLeft();
TIntermTyped *right = node.getRight();
PostResult leftResult = rebuild(*left);
ASSERT(leftResult.single());
pushStacks();
PostResult rightResult = rebuild(*right);
ASSERT(rightResult.single());
return {node, VisitBits::Post};
}
return node;
}
PostResult visitBinaryPost(TIntermBinary &node) override
{
const TOperator op = node.getOp();
if (op == TOperator::EOpInitialize && getParentNode()->getAsDeclarationNode())
{
// Special case is handled by visitDeclarationPost
return node;
}
TIntermTyped *left = node.getLeft();
TIntermTyped *right = node.getRight();
if (op == TOperator::EOpLogicalAnd || op == TOperator::EOpLogicalOr)
{
const Name name = mIdGen.createNewName("");
auto *var = new TVariable(&mSymbolTable, name.rawName(), new TType(TBasicType::EbtBool),
name.symbolType());
TIntermTyped *newRight = pullMappedExpr(right, true);
TIntermBlock *rightBlock = &buildBlockWithTailAssign(*var, *newRight);
TIntermTyped *newLeft = pullMappedExpr(left, true);
TIntermTyped *cond = new TIntermSymbol(var);
if (op == TOperator::EOpLogicalOr)
{
cond = new TIntermUnary(TOperator::EOpLogicalNot, cond, nullptr);
}
pushStmt(*new TIntermDeclaration(var, newLeft));
pushStmt(*new TIntermIfElse(cond, rightBlock, nullptr));
if (!isStandaloneExpr(node))
{
mExprMap[&node] = new TIntermSymbol(var);
}
return node;
}
const bool isAssign = IsAssignment(op);
TIntermTyped *newLeft = pullMappedExpr(left, false);
TIntermTyped *newRight = pullMappedExpr(right, isAssign);
if (op == TOperator::EOpComma)
{
pushBinding(node, *newRight);
return node;
}
else
{
TIntermBinary *newNode;
if (left == newLeft && right == newRight)
{
newNode = &node;
}
else
{
newNode = new TIntermBinary(op, newLeft, newRight);
}
pushBinding(node, *newNode);
return node;
}
}
PreResult visitTernaryPre(TIntermTernary &node) override
{
PostResult condResult = rebuild(*node.getCondition());
ASSERT(condResult.single());
pushStacks();
PostResult thenResult = rebuild(*node.getTrueExpression());
ASSERT(thenResult.single());
pushStacks();
PostResult elseResult = rebuild(*node.getFalseExpression());
ASSERT(elseResult.single());
return {node, VisitBits::Post};
}
PostResult visitTernaryPost(TIntermTernary &node) override
{
TIntermTyped *cond = node.getCondition();
TIntermTyped *then = node.getTrueExpression();
TIntermTyped *else_ = node.getFalseExpression();
const Name name = mIdGen.createNewName("");
auto *var =
new TVariable(&mSymbolTable, name.rawName(), &node.getType(), name.symbolType());
TIntermTyped *newElse = pullMappedExpr(else_, false);
TIntermBlock *elseBlock = &buildBlockWithTailAssign(*var, *newElse);
TIntermTyped *newThen = pullMappedExpr(then, true);
TIntermBlock *thenBlock = &buildBlockWithTailAssign(*var, *newThen);
TIntermTyped *newCond = pullMappedExpr(cond, true);
pushStmt(*new TIntermDeclaration{var});
pushStmt(*new TIntermIfElse(newCond, thenBlock, elseBlock));
if (!isStandaloneExpr(node))
{
mExprMap[&node] = new TIntermSymbol(var);
}
return node;
}
PostResult visitSwizzlePost(TIntermSwizzle &node) override
{
TIntermTyped *expr = node.getOperand();
TIntermTyped *newExpr = pullMappedExpr(expr, false);
if (expr == newExpr)
{
pushBinding(node, node);
}
else
{
pushBinding(node, *new TIntermSwizzle(newExpr, node.getSwizzleOffsets()));
}
return node;
}
PostResult visitAggregatePost(TIntermAggregate &node) override
{
TIntermSequence &args = *node.getSequence();
for (TIntermNode *&arg : args)
{
TIntermTyped *targ = arg->getAsTyped();
ASSERT(targ);
arg = pullMappedExpr(targ, false);
}
pushBinding(node, node);
return node;
}
PostResult visitPreprocessorDirectivePost(TIntermPreprocessorDirective &node) override
{
pushStmt(node);
return node;
}
PostResult visitFunctionPrototypePost(TIntermFunctionPrototype &node) override
{
if (!getParentFunction())
{
pushStmt(node);
}
return node;
}
PreResult visitCasePre(TIntermCase &node) override
{
if (TIntermTyped *cond = node.getCondition())
{
ASSERT(isTerminalExpr(*cond));
}
pushStmt(node);
return {node, VisitBits::Neither};
}
PostResult visitSwitchPost(TIntermSwitch &node) override
{
TIntermTyped *init = node.getInit();
TIntermTyped *newInit = pullMappedExpr(init, false);
if (init == newInit)
{
pushStmt(node);
}
else
{
pushStmt(*new TIntermSwitch(newInit, node.getStatementList()));
}
return node;
}
PostResult visitFunctionDefinitionPost(TIntermFunctionDefinition &node) override
{
pushStmt(node);
return node;
}
PostResult visitIfElsePost(TIntermIfElse &node) override
{
TIntermTyped *cond = node.getCondition();
TIntermTyped *newCond = pullMappedExpr(cond, false);
if (cond == newCond)
{
pushStmt(node);
}
else
{
pushStmt(*new TIntermIfElse(newCond, node.getTrueBlock(), node.getFalseBlock()));
}
return node;
}
PostResult visitBranchPost(TIntermBranch &node) override
{
TIntermTyped *expr = node.getExpression();
TIntermTyped *newExpr = pullMappedExpr(expr, false);
if (expr == newExpr)
{
pushStmt(node);
}
else
{
pushStmt(*new TIntermBranch(node.getFlowOp(), newExpr));
}
return node;
}
PreResult visitLoopPre(TIntermLoop &node) override
{
if (!rebuildInPlace(*node.getBody()))
{
UNREACHABLE();
}
pushStmt(node);
return {node, VisitBits::Neither};
}
PostResult visitConstantUnionPost(TIntermConstantUnion &node) override
{
const TType &type = node.getType();
if (!type.isScalar())
{
pushBinding(node, node);
}
return node;
}
PostResult visitGlobalQualifierDeclarationPost(TIntermGlobalQualifierDeclaration &node) override
{
ASSERT(false); // These should be scrubbed from AST before rewriter is called.
pushStmt(node);
return node;
}
};
} // anonymous namespace
////////////////////////////////////////////////////////////////////////////////
bool sh::SeparateCompoundExpressions(TCompiler &compiler,
SymbolEnv &symbolEnv,
IdGen &idGen,
TIntermBlock &root)
{
std::string disableSeparateString =
angle::GetEnvironmentVar("GMT_DISABLE_SEPARATE_COMPOUND_EXPRESSIONS");
bool disableSeparateBool = !disableSeparateString.empty() && (disableSeparateString == "1");
if (disableSeparateBool)
{
return true;
}
if (!SimplifyLoopConditions(&compiler, &root, &compiler.getSymbolTable()))
{
return false;
}
if (!PrePass(compiler).rebuildRoot(root))
{
return false;
}
if (!Separator(compiler, symbolEnv, idGen).rebuildRoot(root))
{
return false;
}
return true;
}

47
src/compiler/translator/TranslatorMetalDirect/SeparateCompoundExpressions.h Normal file
Просмотреть файл

@ -0,0 +1,47 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_SEPARATECOMPOUNDEXPRESSIONS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_SEPARATECOMPOUNDEXPRESSIONS_H_
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/IdGen.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
namespace sh
{
// Transforms code to (usually) have most one non-terminal expression per statement.
// This also rewrites (&&), (||), and (?:) into raw if/if-not/if-else statements, respectively.
//
// e.g.
// int x = 6 + foo(y, bar());
// becomes
// auto _1 = bar();
// auto _2 = foo(y, _1);
// auto _3 = 6 + _2;
// int x = _3;
//
// WARNING:
// - This does not rewrite object indexing operators as a whole (e.g. foo.x, foo[x]), but will
// rewrite the arguments to the operator (when applicable).
// e.g.
// foo(getVec()[i + 2] + 1);
// becomes
// auto _1 = getVec();
// auto _2 = i + 2;
// auto _3 = _1[_2] + 1; // Index operator remains in (+) expr here.
// foo(_3);
//
ANGLE_NO_DISCARD bool SeparateCompoundExpressions(TCompiler &compiler,
SymbolEnv &symbolEnv,
IdGen &idGen,
TIntermBlock &root);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_SEPARATECOMPOUNDEXPRESSIONS_H_

111
src/compiler/translator/TranslatorMetalDirect/SeparateCompoundStructDeclarations.cpp Normal file
Просмотреть файл

@ -0,0 +1,111 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <algorithm>
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/SeparateCompoundStructDeclarations.h"
#include "compiler/translator/tree_ops/SeparateDeclarations.h"
#include "compiler/translator/tree_util/IntermTraverse.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
class Separator : public TIntermTraverser
{
public:
std::unordered_map<int, TIntermSymbol *> replacementMap;
Separator(TSymbolTable &symbolTable, IdGen &idGen)
: TIntermTraverser(false, false, true, &symbolTable), mIdGen(idGen)
{}
IdGen &mIdGen;
bool visitDeclaration(Visit, TIntermDeclaration *declNode) override
{
ASSERT(declNode->getChildCount() == 1);
Declaration declaration = ViewDeclaration(*declNode);
const TVariable &var = declaration.symbol.variable();
const TType &type = var.getType();
const SymbolType symbolType = var.symbolType();
if (type.isStructSpecifier() && symbolType != SymbolType::Empty)
{
const TStructure *structure = type.getStruct();
TVariable *structVar = nullptr;
TType *instanceType = nullptr;
// Name unnamed inline structs
if (structure->symbolType() == SymbolType::Empty)
{
const TStructure *structDefn =
new TStructure(mSymbolTable, mIdGen.createNewName("__unnamed").rawName(),
&(structure->fields()), SymbolType::AngleInternal);
structVar = new TVariable(mSymbolTable, ImmutableString(""),
new TType(structDefn, true), SymbolType::Empty);
instanceType = new TType(structDefn, false);
}
else
{
structVar = new TVariable(mSymbolTable, ImmutableString(""),
new TType(structure, true), SymbolType::Empty);
instanceType = new TType(structure, false);
}
instanceType->setQualifier(type.getQualifier());
auto *instanceVar =
new TVariable(mSymbolTable, var.name(), instanceType, symbolType, var.extension());
TIntermSequence replacements;
replacements.push_back(new TIntermSymbol(structVar));
TIntermSymbol *instanceSymbol = new TIntermSymbol(instanceVar);
TIntermNode *instanceReplacement = instanceSymbol;
if (declaration.initExpr)
{
instanceReplacement =
new TIntermBinary(EOpInitialize, instanceSymbol, declaration.initExpr);
}
replacements.push_back(instanceReplacement);
replacementMap[declaration.symbol.uniqueId().get()] = instanceSymbol;
mMultiReplacements.push_back(NodeReplaceWithMultipleEntry(
declNode, declNode->getChildNode(0), std::move(replacements)));
}
return false;
}
void visitSymbol(TIntermSymbol *decl) override
{
auto symbol = replacementMap.find(decl->uniqueId().get());
if (symbol != replacementMap.end())
{
queueReplacement(symbol->second->deepCopy(), OriginalNode::IS_DROPPED);
}
}
};
} // anonymous namespace
////////////////////////////////////////////////////////////////////////////////
bool sh::SeparateCompoundStructDeclarations(TCompiler &compiler, IdGen &idGen, TIntermBlock &root)
{
Separator separator(compiler.getSymbolTable(), idGen);
root.traverse(&separator);
if (!separator.updateTree(&compiler, &root))
{
return false;
}
if (!SeparateDeclarations(&compiler, &root))
{
return false;
}
return true;
}

26
src/compiler/translator/TranslatorMetalDirect/SeparateCompoundStructDeclarations.h Normal file
Просмотреть файл

@ -0,0 +1,26 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_SEPARATECOMPOUNDSTRUCTDECLARATIONS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_SEPARATECOMPOUNDSTRUCTDECLARATIONS_H_
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/IdGen.h"
namespace sh
{
// Example:
// struct Foo { int x; } foo;
// Becomes:
// struct Foo {int x; }; Foo foo;
ANGLE_NO_DISCARD bool SeparateCompoundStructDeclarations(TCompiler &compiler,
IdGen &idGen,
TIntermBlock &root);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_SEPARATECOMPOUNDSTRUCTDECLARATIONS_H_

47
src/compiler/translator/TranslatorMetalDirect/SkippingTraverser.h Normal file
Просмотреть файл

@ -0,0 +1,47 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_SKIPPINGTRAVERSER_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_SKIPPINGTRAVERSER_H_
#include "compiler/translator/tree_util/IntermTraverse.h"
namespace sh
{
// A TIntermTraverser that skips traversing childen by default.
class SkippingTraverser : public TIntermTraverser
{
public:
SkippingTraverser(bool preVisit_,
bool inVisit_,
bool postVisit_,
TSymbolTable *symbolTable = nullptr)
: TIntermTraverser(preVisit_, inVisit_, postVisit_, symbolTable)
{}
bool visitSwizzle(Visit, TIntermSwizzle *) { return false; }
bool visitUnary(Visit, TIntermUnary *) { return false; }
bool visitBinary(Visit, TIntermBinary *) { return false; }
bool visitTernary(Visit, TIntermTernary *) { return false; }
bool visitIfElse(Visit, TIntermIfElse *) { return false; }
bool visitSwitch(Visit, TIntermSwitch *) { return false; }
bool visitCase(Visit, TIntermCase *) { return false; }
bool visitFunctionDefinition(Visit, TIntermFunctionDefinition *) { return false; }
bool visitAggregate(Visit, TIntermAggregate *) { return false; }
bool visitBlock(Visit, TIntermBlock *) { return false; }
bool visitDeclaration(Visit, TIntermDeclaration *) { return false; }
bool visitLoop(Visit, TIntermLoop *) { return false; }
bool visitBranch(Visit, TIntermBranch *) { return false; }
bool visitGlobalQualifierDeclaration(Visit, TIntermGlobalQualifierDeclaration *)
{
return false;
}
};
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_SKIPPINGTRAVERSER_H_

702
src/compiler/translator/TranslatorMetalDirect/SymbolEnv.cpp Normal file
Просмотреть файл

@ -0,0 +1,702 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <algorithm>
#include <limits>
#include "compiler/translator/ImmutableStringBuilder.h"
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
constexpr AddressSpace kAddressSpaces[] = {
AddressSpace::Constant,
AddressSpace::Device,
AddressSpace::Thread,
};
char const *sh::toString(AddressSpace space)
{
switch (space)
{
case AddressSpace::Constant:
return "constant";
case AddressSpace::Device:
return "device";
case AddressSpace::Thread:
return "thread";
}
}
////////////////////////////////////////////////////////////////////////////////
using NameToStruct = std::map<Name, const TStructure *>;
class StructFinder : TIntermRebuild
{
NameToStruct nameToStruct;
StructFinder(TCompiler &compiler) : TIntermRebuild(compiler, true, false) {}
PreResult visitDeclarationPre(TIntermDeclaration &node) override
{
Declaration decl = ViewDeclaration(node);
const TVariable &var = decl.symbol.variable();
const TType &type = var.getType();
if (var.symbolType() == SymbolType::Empty && type.isStructSpecifier())
{
const TStructure *s = type.getStruct();
ASSERT(s);
const Name name(*s);
const TStructure *&z = nameToStruct[name];
ASSERT(!z);
z = s;
}
return node;
}
PreResult visitFunctionDefinitionPre(TIntermFunctionDefinition &node) override
{
return {node, VisitBits::Neither};
}
public:
static NameToStruct FindStructs(TCompiler &compiler, TIntermBlock &root)
{
StructFinder finder(compiler);
if (!finder.rebuildRoot(root))
{
UNREACHABLE();
}
return std::move(finder.nameToStruct);
}
};
////////////////////////////////////////////////////////////////////////////////
TemplateArg::TemplateArg(bool value) : mKind(Kind::Bool), mValue(value) {}
TemplateArg::TemplateArg(int value) : mKind(Kind::Int), mValue(value) {}
TemplateArg::TemplateArg(unsigned value) : mKind(Kind::UInt), mValue(value) {}
TemplateArg::TemplateArg(const TType &value) : mKind(Kind::Type), mValue(value) {}
bool TemplateArg::operator==(const TemplateArg &other) const
{
if (mKind != other.mKind)
{
return false;
}
switch (mKind)
{
case Kind::Bool:
return mValue.b == other.mValue.b;
case Kind::Int:
return mValue.i == other.mValue.i;
case Kind::UInt:
return mValue.u == other.mValue.u;
case Kind::Type:
return *mValue.t == *other.mValue.t;
}
}
bool TemplateArg::operator<(const TemplateArg &other) const
{
if (mKind < other.mKind)
{
return true;
}
if (mKind > other.mKind)
{
return false;
}
switch (mKind)
{
case Kind::Bool:
return mValue.b < other.mValue.b;
case Kind::Int:
return mValue.i < other.mValue.i;
case Kind::UInt:
return mValue.u < other.mValue.u;
case Kind::Type:
return *mValue.t < *other.mValue.t;
}
}
////////////////////////////////////////////////////////////////////////////////
bool SymbolEnv::TemplateName::operator==(const TemplateName &other) const
{
return baseName == other.baseName && templateArgs == other.templateArgs;
}
bool SymbolEnv::TemplateName::operator<(const TemplateName &other) const
{
if (baseName < other.baseName)
{
return true;
}
if (other.baseName < baseName)
{
return false;
}
return templateArgs < other.templateArgs;
}
bool SymbolEnv::TemplateName::empty() const
{
return baseName.empty() && templateArgs.empty();
}
void SymbolEnv::TemplateName::clear()
{
baseName = Name();
templateArgs.clear();
}
Name SymbolEnv::TemplateName::fullName(std::string &buffer) const
{
ASSERT(buffer.empty());
if (templateArgs.empty())
{
return baseName;
}
static constexpr size_t n = std::max({
std::numeric_limits<unsigned>::digits10, //
std::numeric_limits<int>::digits10, //
5, // max_length("true", "false")
});
buffer.reserve(baseName.rawName().length() + (n + 2) * templateArgs.size() + 1);
buffer += baseName.rawName().data();
if (!templateArgs.empty())
{
buffer += "<";
bool first = true;
char argBuffer[n + 1];
for (const TemplateArg &arg : templateArgs)
{
if (first)
{
first = false;
}
else
{
buffer += ", ";
}
const TemplateArg::Value value = arg.value();
const TemplateArg::Kind kind = arg.kind();
switch (kind)
{
case TemplateArg::Kind::Bool:
if (value.b)
{
buffer += "true";
}
else
{
buffer += "false";
}
break;
case TemplateArg::Kind::Int:
sprintf(argBuffer, "%i", value.i);
buffer += argBuffer;
break;
case TemplateArg::Kind::UInt:
sprintf(argBuffer, "%u", value.u);
buffer += argBuffer;
break;
case TemplateArg::Kind::Type:
{
const TType &type = *value.t;
if (const TStructure *s = type.getStruct())
{
buffer += s->name().data();
}
else if (HasScalarBasicType(type))
{
ASSERT(!type.isArray()); // TODO
buffer += type.getBasicString();
if (type.isVector())
{
sprintf(argBuffer, "%i", type.getNominalSize());
buffer += argBuffer;
}
else if (type.isMatrix())
{
sprintf(argBuffer, "%i", type.getCols());
buffer += argBuffer;
buffer += "x";
sprintf(argBuffer, "%i", type.getRows());
buffer += argBuffer;
}
}
}
break;
}
}
buffer += ">";
}
const ImmutableString name(buffer);
buffer.clear();
return Name(name, baseName.symbolType());
}
void SymbolEnv::TemplateName::assign(const Name &name, size_t argCount, const TemplateArg *args)
{
baseName = name;
templateArgs.clear();
for (size_t i = 0; i < argCount; ++i)
{
templateArgs.push_back(args[i]);
}
}
////////////////////////////////////////////////////////////////////////////////
SymbolEnv::SymbolEnv(TCompiler &compiler, TIntermBlock &root)
: mSymbolTable(compiler.getSymbolTable()),
mNameToStruct(StructFinder::FindStructs(compiler, root))
{}
const TStructure &SymbolEnv::remap(const TStructure &s) const
{
const Name name(s);
auto iter = mNameToStruct.find(name);
if (iter == mNameToStruct.end())
{
return s;
}
const TStructure &z = *iter->second;
return z;
}
const TStructure *SymbolEnv::remap(const TStructure *s) const
{
if (s)
{
return &remap(*s);
}
return nullptr;
}
const TFunction &SymbolEnv::getFunctionOverloadImpl()
{
ASSERT(!mReusableSigBuffer.empty());
SigToFunc &sigToFunc = mOverloads[mReusableTemplateNameBuffer];
TFunction *&func = sigToFunc[mReusableSigBuffer];
if (!func)
{
const TType &returnType = mReusableSigBuffer.back();
mReusableSigBuffer.pop_back();
const Name name = mReusableTemplateNameBuffer.fullName(mReusableStringBuffer);
func = new TFunction(&mSymbolTable, name.rawName(), name.symbolType(), &returnType, false);
for (const TType &paramType : mReusableSigBuffer)
{
func->addParameter(
new TVariable(&mSymbolTable, kEmptyImmutableString, &paramType, SymbolType::Empty));
}
}
mReusableSigBuffer.clear();
mReusableTemplateNameBuffer.clear();
return *func;
}
const TFunction &SymbolEnv::getFunctionOverload(const Name &name,
const TType &returnType,
size_t paramCount,
const TType **paramTypes,
size_t templateArgCount,
const TemplateArg *templateArgs)
{
ASSERT(mReusableSigBuffer.empty());
ASSERT(mReusableTemplateNameBuffer.empty());
for (size_t i = 0; i < paramCount; ++i)
{
mReusableSigBuffer.push_back(*paramTypes[i]);
}
mReusableSigBuffer.push_back(returnType);
mReusableTemplateNameBuffer.assign(name, templateArgCount, templateArgs);
return getFunctionOverloadImpl();
}
TIntermAggregate &SymbolEnv::callFunctionOverload(const Name &name,
const TType &returnType,
TIntermSequence &args,
size_t templateArgCount,
const TemplateArg *templateArgs)
{
ASSERT(mReusableSigBuffer.empty());
ASSERT(mReusableTemplateNameBuffer.empty());
for (TIntermNode *arg : args)
{
TIntermTyped *targ = arg->getAsTyped();
ASSERT(targ);
mReusableSigBuffer.push_back(targ->getType());
}
mReusableSigBuffer.push_back(returnType);
mReusableTemplateNameBuffer.assign(name, templateArgCount, templateArgs);
const TFunction &func = getFunctionOverloadImpl();
return *TIntermAggregate::CreateRawFunctionCall(func, &args);
}
const TStructure &SymbolEnv::newStructure(const Name &name, TFieldList &fields)
{
ASSERT(name.symbolType() == SymbolType::AngleInternal);
TStructure *&s = mAngleStructs[name.rawName()];
ASSERT(!s);
s = new TStructure(&mSymbolTable, name.rawName(), &fields, name.symbolType());
return *s;
}
const TStructure &SymbolEnv::getTextureEnv(TBasicType samplerType)
{
ASSERT(IsSampler(samplerType));
const TStructure *&env = mTextureEnvs[samplerType];
if (env == nullptr)
{
auto *textureType = new TType(samplerType);
auto *texture = new TField(textureType, ImmutableString("texture"), kNoSourceLoc,
SymbolType::UserDefined);
markAsPointer(*texture, AddressSpace::Thread);
auto *sampler =
new TField(new TType(&getSamplerStruct(), false), ImmutableString("sampler"),
kNoSourceLoc, SymbolType::UserDefined);
markAsPointer(*sampler, AddressSpace::Thread);
std::string envName;
envName += "TextureEnv<";
envName += GetTextureTypeName(samplerType).rawName().data();
envName += ">";
env = &newStructure(Name(envName, SymbolType::AngleInternal),
*new TFieldList{texture, sampler});
}
return *env;
}
const TStructure &SymbolEnv::getSamplerStruct()
{
if (!mSampler)
{
mSampler = new TStructure(&mSymbolTable, ImmutableString("metal::sampler"),
new TFieldList(), SymbolType::UserDefined);
}
return *mSampler;
}
void SymbolEnv::markSpace(VarField x,
AddressSpace space,
std::unordered_map<VarField, AddressSpace> &map)
{
// It is in principle permissible to have references to pointers or multiple pointers, but this
// is not required for now and would require code changes to get right.
ASSERT(!isPointer(x));
ASSERT(!isReference(x));
map[x] = space;
}
void SymbolEnv::removeSpace(VarField x, std::unordered_map<VarField, AddressSpace> &map)
{
// It is in principle permissible to have references to pointers or multiple pointers, but this
// is not required for now and would require code changes to get right.
map.erase(x);
}
const AddressSpace *SymbolEnv::isSpace(VarField x,
const std::unordered_map<VarField, AddressSpace> &map) const
{
const auto iter = map.find(x);
if (iter == map.end())
{
return nullptr;
}
const AddressSpace space = iter->second;
const auto index = static_cast<std::underlying_type_t<AddressSpace>>(space);
return &kAddressSpaces[index];
}
void SymbolEnv::markAsPointer(VarField x, AddressSpace space)
{
return markSpace(x, space, mPointers);
}
void SymbolEnv::removePointer(VarField x)
{
return removeSpace(x, mPointers);
}
void SymbolEnv::markAsReference(VarField x, AddressSpace space)
{
return markSpace(x, space, mReferences);
}
const AddressSpace *SymbolEnv::isPointer(VarField x) const
{
return isSpace(x, mPointers);
}
const AddressSpace *SymbolEnv::isReference(VarField x) const
{
return isSpace(x, mReferences);
}
void SymbolEnv::markAsPacked(const TField &field)
{
mPackedFields.insert(&field);
}
bool SymbolEnv::isPacked(const TField &field) const
{
return mPackedFields.find(&field) != mPackedFields.end();
}
void SymbolEnv::markAsUBO(VarField x)
{
mUboFields.insert(x);
}
bool SymbolEnv::isUBO(VarField x) const
{
return mUboFields.find(x) != mUboFields.end();
}
static TBasicType GetTextureBasicType(TBasicType basicType)
{
ASSERT(IsSampler(basicType));
switch (basicType)
{
case EbtSampler2D:
case EbtSampler3D:
case EbtSamplerCube:
case EbtSampler2DArray:
case EbtSamplerExternalOES:
case EbtSamplerExternal2DY2YEXT:
case EbtSampler2DRect:
case EbtSampler2DMS:
case EbtSampler2DMSArray:
case EbtSamplerVideoWEBGL:
case EbtSampler2DShadow:
case EbtSamplerCubeShadow:
case EbtSampler2DArrayShadow:
case EbtSampler1D:
case EbtSampler1DArray:
case EbtSampler1DArrayShadow:
case EbtSamplerBuffer:
case EbtSamplerCubeArray:
case EbtSamplerCubeArrayShadow:
case EbtSampler1DShadow:
case EbtSampler2DRectShadow:
return TBasicType::EbtFloat;
case EbtISampler2D:
case EbtISampler3D:
case EbtISamplerCube:
case EbtISampler2DArray:
case EbtISampler2DMS:
case EbtISampler2DMSArray:
case EbtISampler1D:
case EbtISampler1DArray:
case EbtISampler2DRect:
case EbtISamplerBuffer:
case EbtISamplerCubeArray:
return TBasicType::EbtInt;
case EbtUSampler2D:
case EbtUSampler3D:
case EbtUSamplerCube:
case EbtUSampler2DArray:
case EbtUSampler2DMS:
case EbtUSampler2DMSArray:
case EbtUSampler1D:
case EbtUSampler1DArray:
case EbtUSampler2DRect:
case EbtUSamplerBuffer:
case EbtUSamplerCubeArray:
return TBasicType::EbtUInt;
default:
UNREACHABLE();
return TBasicType::EbtVoid;
}
}
Name sh::GetTextureTypeName(TBasicType samplerType)
{
ASSERT(IsSampler(samplerType));
const TBasicType textureType = GetTextureBasicType(samplerType);
const char *name;
#define HANDLE_TEXTURE_NAME(baseName) \
do \
{ \
switch (textureType) \
{ \
case TBasicType::EbtFloat: \
name = "metal::" baseName "<float>"; \
break; \
case TBasicType::EbtInt: \
name = "metal::" baseName "<int>"; \
break; \
case TBasicType::EbtUInt: \
name = "metal::" baseName "<uint>"; \
break; \
default: \
UNREACHABLE(); \
name = nullptr; \
break; \
} \
} while (false)
switch (samplerType)
{
// 1d
case EbtSampler1D: // Desktop GLSL sampler type:
case EbtISampler1D:
case EbtUSampler1D:
HANDLE_TEXTURE_NAME("texture1d");
break;
// 1d array
case EbtSampler1DArray:
case EbtISampler1DArray:
case EbtUSampler1DArray:
HANDLE_TEXTURE_NAME("texture1d_array");
break;
// Buffer textures
case EbtSamplerBuffer:
case EbtISamplerBuffer:
case EbtUSamplerBuffer:
HANDLE_TEXTURE_NAME("texture_buffer");
break;
// 2d textures
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
case EbtSampler2DRect:
case EbtUSampler2DRect:
case EbtISampler2DRect:
HANDLE_TEXTURE_NAME("texture2d");
break;
// 3d textures
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
HANDLE_TEXTURE_NAME("texture3d");
break;
// Cube textures
case EbtSamplerCube:
case EbtISamplerCube:
case EbtUSamplerCube:
HANDLE_TEXTURE_NAME("texturecube");
break;
// 2d array textures
case EbtSampler2DArray:
case EbtUSampler2DArray:
case EbtISampler2DArray:
HANDLE_TEXTURE_NAME("texture2d_array");
break;
case EbtSampler2DMS:
case EbtISampler2DMS:
case EbtUSampler2DMS:
HANDLE_TEXTURE_NAME("texture2d_ms");
break;
case EbtSampler2DMSArray:
case EbtISampler2DMSArray:
case EbtUSampler2DMSArray:
HANDLE_TEXTURE_NAME("texture2d_ms_array");
break;
// cube array
case EbtSamplerCubeArray:
case EbtISamplerCubeArray:
case EbtUSamplerCubeArray:
HANDLE_TEXTURE_NAME("texturecube_array");
break;
// Shadow
case EbtSampler1DShadow:
case EbtSampler1DArrayShadow:
UNIMPLEMENTED();
HANDLE_TEXTURE_NAME("TODO");
break;
case EbtSampler2DRectShadow:
case EbtSampler2DShadow:
HANDLE_TEXTURE_NAME("depth2d");
break;
case EbtSamplerCubeShadow:
HANDLE_TEXTURE_NAME("depthcube");
break;
case EbtSampler2DArrayShadow:
HANDLE_TEXTURE_NAME("depth2d_array");
break;
case EbtSamplerCubeArrayShadow:
HANDLE_TEXTURE_NAME("depthcube_array");
break;
// Extentions
case EbtSamplerExternalOES: // Only valid if OES_EGL_image_external exists:
case EbtSamplerExternal2DY2YEXT: // Only valid if GL_EXT_YUV_target exists:
case EbtSamplerVideoWEBGL:
UNIMPLEMENTED();
HANDLE_TEXTURE_NAME("TODO");
break;
default:
UNREACHABLE();
name = nullptr;
break;
}
#undef HANDLE_TEXTURE_NAME
return Name(name, SymbolType::UserDefined);
}

225
src/compiler/translator/TranslatorMetalDirect/SymbolEnv.h Normal file
Просмотреть файл

@ -0,0 +1,225 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_SYMBOLENV_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_SYMBOLENV_H_
#include <unordered_set>
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/Name.h"
#include "compiler/translator/TranslatorMetalDirect/Reference.h"
#include "compiler/translator/Types.h"
namespace sh
{
enum class AddressSpace
{
Constant,
Device,
Thread,
};
char const *toString(AddressSpace space);
class VarField
{
public:
VarField(const TVariable &var) : mVariable(&var) {}
VarField(const TField &field) : mField(&field) {}
ANGLE_INLINE const TVariable *variable() const { return mVariable; }
ANGLE_INLINE const TField *field() const { return mField; }
ANGLE_INLINE bool operator==(const VarField &other) const
{
return mVariable == other.mVariable && mField == other.mField;
}
private:
const TVariable *mVariable = nullptr;
const TField *mField = nullptr;
};
} // namespace sh
namespace std
{
template <>
struct hash<sh::VarField>
{
size_t operator()(const sh::VarField &x) const
{
const sh::TVariable *var = x.variable();
if (var)
{
return std::hash<const sh::TVariable *>()(var);
}
const sh::TField *field = x.field();
return std::hash<const sh::TField *>()(field);
}
};
} // namespace std
namespace sh
{
class TemplateArg
{
public:
enum class Kind
{
Bool,
Int,
UInt,
Type,
};
union Value
{
Value(bool value) : b(value) {}
Value(int value) : i(value) {}
Value(unsigned value) : u(value) {}
Value(const TType &value) : t(&value) {}
bool b;
int i;
unsigned u;
const TType *t;
};
TemplateArg(bool value);
TemplateArg(int value);
TemplateArg(unsigned value);
TemplateArg(const TType &value);
bool operator==(const TemplateArg &other) const;
bool operator<(const TemplateArg &other) const;
Kind kind() const { return mKind; }
Value value() const { return mValue; }
public:
Kind mKind;
Value mValue;
};
// An environment for creating and uniquely sharing TSymbol objects.
class SymbolEnv
{
class TemplateName
{
Name baseName;
std::vector<TemplateArg> templateArgs;
public:
bool operator==(const TemplateName &other) const;
bool operator<(const TemplateName &other) const;
bool empty() const;
void clear();
Name fullName(std::string &buffer) const;
void assign(const Name &name, size_t argCount, const TemplateArg *args);
};
using TypeRef = CRef<TType>;
using Sig = std::vector<TypeRef>; // Param0, Param1, ... ParamN, Return
using SigToFunc = std::map<Sig, TFunction *>;
using Overloads = std::map<TemplateName, SigToFunc>;
using AngleStructs = std::map<ImmutableString, TStructure *>;
using TextureStructs = std::map<TBasicType, TStructure *>;
public:
SymbolEnv(TCompiler &compiler, TIntermBlock &root);
TSymbolTable &symbolTable() { return mSymbolTable; }
// There exist Angle rewrites that can lead to incoherent structs
//
// Example
// struct A { ... }
// struct B { A' a; } // A' has same name as A but is not identical.
// becomes
// struct A { ... }
// struct B { A a; }
//
// This notably happens when A contains a sampler in the original GLSL code but is rewritten to
// not have a sampler, yet the A' struct field still contains the sampler field.
const TStructure &remap(const TStructure &s) const;
// Like `TStructure &remap(const TStructure &s)` but additionally maps null to null.
const TStructure *remap(const TStructure *s) const;
const TFunction &getFunctionOverload(const Name &name,
const TType &returnType,
size_t paramCount,
const TType **paramTypes,
size_t templateArgCount = 0,
const TemplateArg *templateArgs = nullptr);
TIntermAggregate &callFunctionOverload(const Name &name,
const TType &returnType,
TIntermSequence &args,
size_t templateArgCount = 0,
const TemplateArg *templateArgs = nullptr);
const TStructure &newStructure(const Name &name, TFieldList &fields);
const TStructure &getTextureEnv(TBasicType samplerType);
const TStructure &getSamplerStruct();
void markAsPointer(VarField x, AddressSpace space);
void removePointer(VarField x);
const AddressSpace *isPointer(VarField x) const;
void markAsReference(VarField x, AddressSpace space);
void removeAsReference(VarField x);
const AddressSpace *isReference(VarField x) const;
void markAsPacked(const TField &field);
bool isPacked(const TField &field) const;
void markAsUBO(VarField x);
bool isUBO(VarField x) const;
private:
const TFunction &getFunctionOverloadImpl();
void markSpace(VarField x, AddressSpace space, std::unordered_map<VarField, AddressSpace> &map);
void removeSpace(VarField x, std::unordered_map<VarField, AddressSpace> &map);
const AddressSpace *isSpace(VarField x,
const std::unordered_map<VarField, AddressSpace> &map) const;
private:
TSymbolTable &mSymbolTable;
std::map<Name, const TStructure *> mNameToStruct;
Overloads mOverloads;
Sig mReusableSigBuffer;
TemplateName mReusableTemplateNameBuffer;
std::string mReusableStringBuffer;
AngleStructs mAngleStructs;
std::unordered_map<TBasicType, const TStructure *> mTextureEnvs;
const TStructure *mSampler = nullptr;
std::unordered_map<VarField, AddressSpace> mPointers;
std::unordered_map<VarField, AddressSpace> mReferences;
std::unordered_set<const TField *> mPackedFields;
std::unordered_set<VarField> mUboFields;
};
Name GetTextureTypeName(TBasicType samplerType);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_SYMBOLENV_H_

382
src/compiler/translator/TranslatorMetalDirect/ToposortStructs.cpp Normal file
Просмотреть файл

@ -0,0 +1,382 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <functional>
#include <unordered_map>
#include <unordered_set>
#include <vector>
#include "compiler/translator/ImmutableStringBuilder.h"
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/ToposortStructs.h"
#include "compiler/translator/tree_util/IntermNode_util.h"
#include "compiler/translator/tree_util/IntermTraverse.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
template <typename T>
using Edges = std::unordered_set<T>;
template <typename T>
using Graph = std::unordered_map<T, Edges<T>>;
void BuildGraphImpl(SymbolEnv &symbolEnv, Graph<const TStructure *> &g, const TStructure *s)
{
if (g.find(s) != g.end())
{
return;
}
Edges<const TStructure *> &es = g[s];
const TFieldList &fs = s->fields();
for (const TField *f : fs)
{
if (const TStructure *z = symbolEnv.remap(f->type()->getStruct()))
{
es.insert(z);
BuildGraphImpl(symbolEnv, g, z);
Edges<const TStructure *> &ez = g[z];
es.insert(ez.begin(), ez.end());
}
}
}
Graph<const TStructure *> BuildGraph(SymbolEnv &symbolEnv,
const std::vector<const TStructure *> &structs)
{
Graph<const TStructure *> g;
for (const TStructure *s : structs)
{
BuildGraphImpl(symbolEnv, g, s);
}
return g;
}
// Algorthm: https://en.wikipedia.org/wiki/Topological_sorting#Depth-first_search
template <typename T>
std::vector<T> Toposort(const Graph<T> &g)
{
// nodes with temporary mark
std::unordered_set<T> temps;
// nodes without permanent mark
std::unordered_set<T> invPerms;
for (const auto &entry : g)
{
invPerms.insert(entry.first);
}
// L <- Empty list that will contain the sorted elements
std::vector<T> L;
// function visit(node n)
std::function<void(T)> visit = [&](T n) -> void {
// if n has a permanent mark then
if (invPerms.find(n) == invPerms.end())
{
// return
return;
}
// if n has a temporary mark then
if (temps.find(n) != temps.end())
{
// stop (not a DAG)
UNREACHABLE();
}
// mark n with a temporary mark
temps.insert(n);
// for each node m with an edge from n to m do
auto enIter = g.find(n);
ASSERT(enIter != g.end());
const Edges<T> &en = enIter->second;
for (T m : en)
{
// visit(m)
visit(m);
}
// remove temporary mark from n
temps.erase(n);
// mark n with a permanent mark
invPerms.erase(n);
// add n to head of L
L.push_back(n);
};
// while exists nodes without a permanent mark do
while (!invPerms.empty())
{
// select an unmarked node n
T n = *invPerms.begin();
// visit(n)
visit(n);
}
return L;
}
TIntermFunctionDefinition *CreateStructEqualityFunction(TSymbolTable &symbolTable,
const TStructure &aStructType)
{
////////////////////
auto &funcEquality =
*new TFunction(&symbolTable, ImmutableString("equal"), SymbolType::AngleInternal,
new TType(TBasicType::EbtBool), true);
auto &aStruct = CreateInstanceVariable(symbolTable, aStructType, Name("a"));
auto &bStruct = CreateInstanceVariable(symbolTable, aStructType, Name("b"));
funcEquality.addParameter(&aStruct);
funcEquality.addParameter(&bStruct);
auto &bodyEquality = *new TIntermBlock();
std::vector<TIntermTyped *> andNodes;
////////////////////
const TFieldList &aFields = aStructType.fields();
const size_t size = aFields.size();
auto testEquality = [&](TIntermTyped &a, TIntermTyped &b) -> TIntermTyped * {
ASSERT(a.getType() == b.getType());
const TType &type = a.getType();
if (type.isVector() || type.isMatrix() || type.getStruct())
{
auto *func =
new TFunction(&symbolTable, ImmutableString("equal"), SymbolType::AngleInternal,
new TType(TBasicType::EbtBool), true);
return TIntermAggregate::CreateFunctionCall(*func, new TIntermSequence{&a, &b});
}
else
{
return new TIntermBinary(TOperator::EOpEqual, &a, &b);
}
};
for (size_t idx = 0; idx < size; ++idx)
{
const TField &aField = *aFields[idx];
const TType &aFieldType = *aField.type();
auto &aFieldName = aField.name();
if (aFieldType.isArray())
{
ASSERT(!aFieldType.isArrayOfArrays()); // TODO
int dim = aFieldType.getOutermostArraySize();
for (int d = 0; d < dim; ++d)
{
auto &aAccess = AccessIndex(AccessField(aStruct, aFieldName), d);
auto &bAccess = AccessIndex(AccessField(bStruct, aFieldName), d);
auto *eqNode = testEquality(bAccess, aAccess);
andNodes.push_back(eqNode);
}
}
else
{
auto &aAccess = AccessField(aStruct, aFieldName);
auto &bAccess = AccessField(bStruct, aFieldName);
auto *eqNode = testEquality(bAccess, aAccess);
andNodes.push_back(eqNode);
}
}
ASSERT(andNodes.size() > 0); // Empty structs are not allowed in GLSL
TIntermTyped *outNode = andNodes.back();
andNodes.pop_back();
for (TIntermTyped *andNode : andNodes)
{
outNode = new TIntermBinary(TOperator::EOpLogicalAnd, andNode, outNode);
}
bodyEquality.appendStatement(new TIntermBranch(TOperator::EOpReturn, outNode));
auto *funcProtoEquality = new TIntermFunctionPrototype(&funcEquality);
return new TIntermFunctionDefinition(funcProtoEquality, &bodyEquality);
}
struct DeclaredStructure
{
TIntermDeclaration *declNode;
TIntermFunctionDefinition *equalityFunctionDefinition;
const TStructure *structure;
};
bool GetAsDeclaredStructure(SymbolEnv &symbolEnv,
TIntermNode &node,
DeclaredStructure &out,
TSymbolTable &symbolTable,
const std::unordered_set<const TStructure *> &usedStructs)
{
if (TIntermDeclaration *declNode = node.getAsDeclarationNode())
{
ASSERT(declNode->getChildCount() == 1);
TIntermNode &node = *declNode->getChildNode(0);
if (TIntermSymbol *symbolNode = node.getAsSymbolNode())
{
const TVariable &var = symbolNode->variable();
const TType &type = var.getType();
if (const TStructure *structure = symbolEnv.remap(type.getStruct()))
{
if (type.isStructSpecifier())
{
out.declNode = declNode;
out.structure = structure;
out.equalityFunctionDefinition =
usedStructs.find(structure) == usedStructs.end()
? nullptr
: CreateStructEqualityFunction(symbolTable, *structure);
return true;
}
}
}
}
return false;
}
class FindStructEqualityUse : public TIntermTraverser
{
public:
SymbolEnv &mSymbolEnv;
std::unordered_set<const TStructure *> mUsedStructs;
FindStructEqualityUse(SymbolEnv &symbolEnv)
: TIntermTraverser(false, false, true), mSymbolEnv(symbolEnv)
{}
bool visitBinary(Visit, TIntermBinary *binary) override
{
const TOperator op = binary->getOp();
switch (op)
{
case TOperator::EOpEqual:
case TOperator::EOpNotEqual:
{
const TType &leftType = binary->getLeft()->getType();
const TType &rightType = binary->getRight()->getType();
ASSERT(leftType.getStruct() == rightType.getStruct());
if (const TStructure *structure = mSymbolEnv.remap(leftType.getStruct()))
{
useStruct(*structure);
}
}
break;
default:
break;
}
return true;
}
private:
void useStruct(const TStructure &structure)
{
if (mUsedStructs.insert(&structure).second)
{
for (const TField *field : structure.fields())
{
if (const TStructure *subStruct = mSymbolEnv.remap(field->type()->getStruct()))
{
useStruct(*subStruct);
}
}
}
}
};
} // anonymous namespace
////////////////////////////////////////////////////////////////////////////////
bool sh::ToposortStructs(TCompiler &compiler,
SymbolEnv &symbolEnv,
TIntermBlock &root,
ProgramPreludeConfig &ppc)
{
FindStructEqualityUse finder(symbolEnv);
root.traverse(&finder);
ppc.hasStructEq = !finder.mUsedStructs.empty();
DeclaredStructure declaredStruct;
std::vector<DeclaredStructure> declaredStructs;
std::vector<TIntermNode *> nonStructStmtNodes;
{
const size_t stmtCount = root.getChildCount();
for (size_t i = 0; i < stmtCount; ++i)
{
TIntermNode &stmtNode = *root.getChildNode(i);
if (GetAsDeclaredStructure(symbolEnv, stmtNode, declaredStruct,
compiler.getSymbolTable(), finder.mUsedStructs))
{
declaredStructs.push_back(declaredStruct);
}
else
{
nonStructStmtNodes.push_back(&stmtNode);
}
}
}
{
std::vector<const TStructure *> structs;
std::unordered_map<const TStructure *, DeclaredStructure> rawToDeclared;
for (const DeclaredStructure &d : declaredStructs)
{
structs.push_back(d.structure);
ASSERT(rawToDeclared.find(d.structure) == rawToDeclared.end());
rawToDeclared[d.structure] = d;
}
// Note: Graph may contain more than only explicitly declared structures.
Graph<const TStructure *> g = BuildGraph(symbolEnv, structs);
std::vector<const TStructure *> sortedStructs = Toposort(g);
ASSERT(declaredStructs.size() <= sortedStructs.size());
declaredStructs.clear();
for (const TStructure *s : sortedStructs)
{
auto it = rawToDeclared.find(s);
if (it != rawToDeclared.end())
{
auto &d = it->second;
ASSERT(d.declNode);
declaredStructs.push_back(d);
}
}
}
{
TIntermSequence newStmtNodes;
for (DeclaredStructure &declaredStruct : declaredStructs)
{
ASSERT(declaredStruct.declNode);
newStmtNodes.push_back(declaredStruct.declNode);
if (declaredStruct.equalityFunctionDefinition)
{
newStmtNodes.push_back(declaredStruct.equalityFunctionDefinition);
}
}
for (TIntermNode *stmtNode : nonStructStmtNodes)
{
ASSERT(stmtNode);
newStmtNodes.push_back(stmtNode);
}
*root.getSequence() = newStmtNodes;
}
return compiler.validateAST(&root);
}

27
src/compiler/translator/TranslatorMetalDirect/ToposortStructs.h Normal file
Просмотреть файл

@ -0,0 +1,27 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_TOPOSORTSTRUCTS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_TOPOSORTSTRUCTS_H_
#include "common/angleutils.h"
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/ProgramPrelude.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
namespace sh
{
// Does a toposort on structs based on type dependencies.
// Struct type declarations are moved to the top of the root block.
ANGLE_NO_DISCARD bool ToposortStructs(TCompiler &compiler,
SymbolEnv &symbolEnv,
TIntermBlock &root,
ProgramPreludeConfig &ppc);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_TOPOSORTSTRUCTS_H_

305
src/compiler/translator/TranslatorMetalDirect/TranslatorMetalUtils.cpp Normal file
Просмотреть файл

@ -0,0 +1,305 @@
//
// Copyright 2002 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#if defined(_MSC_VER)
# pragma warning(disable : 4718)
#endif
#include "compiler/translator/TranslatorMetalDirect/TranslatorMetalUtils.h"
#include <algorithm>
#include <climits>
#include "compiler/translator/HashNames.h"
#include "compiler/translator/ImmutableString.h"
#include "compiler/translator/InfoSink.h"
#include "compiler/translator/IntermNode.h"
#include "compiler/translator/SymbolTable.h"
#include "compiler/translator/Types.h"
namespace sh
{
const char *getBasicMetalString(const TType *t)
{
switch (t->getBasicType())
{
case EbtVoid:
return "void";
case EbtFloat:
return "float";
case EbtInt:
return "int";
case EbtUInt:
return "uint";
case EbtBool:
return "bool";
case EbtYuvCscStandardEXT:
return "yuvCscStandardEXT";
case EbtSampler2D:
return "sampler2D";
case EbtSampler3D:
return "sampler3D";
case EbtSamplerCube:
return "samplerCube";
case EbtSamplerExternalOES:
return "samplerExternalOES";
case EbtSamplerExternal2DY2YEXT:
return "__samplerExternal2DY2YEXT";
case EbtSampler2DRect:
return "sampler2DRect";
case EbtSampler2DArray:
return "sampler2DArray";
case EbtSampler2DMS:
return "sampler2DMS";
case EbtSampler2DMSArray:
return "sampler2DMSArray";
case EbtSamplerCubeArray:
return "samplerCubeArray";
case EbtISampler2D:
return "isampler2D";
case EbtISampler3D:
return "isampler3D";
case EbtISamplerCube:
return "isamplerCube";
case EbtISampler2DArray:
return "isampler2DArray";
case EbtISampler2DMS:
return "isampler2DMS";
case EbtISampler2DMSArray:
return "isampler2DMSArray";
case EbtISamplerCubeArray:
return "isamplerCubeArray";
case EbtUSampler2D:
return "usampler2D";
case EbtUSampler3D:
return "usampler3D";
case EbtUSamplerCube:
return "usamplerCube";
case EbtUSampler2DArray:
return "usampler2DArray";
case EbtUSampler2DMS:
return "usampler2DMS";
case EbtUSampler2DMSArray:
return "usampler2DMSArray";
case EbtUSamplerCubeArray:
return "usamplerCubeArray";
case EbtSampler2DShadow:
return "sampler2DShadow";
case EbtSamplerCubeShadow:
return "samplerCubeShadow";
case EbtSampler2DArrayShadow:
return "sampler2DArrayShadow";
case EbtSamplerCubeArrayShadow:
return "samplerCubeArrayShadow";
case EbtStruct:
return "structure";
case EbtInterfaceBlock:
return "interface block";
case EbtImage2D:
return "texture2d";
case EbtIImage2D:
return "texture2d<int>";
case EbtUImage2D:
return "texture2d<uint>";
case EbtImage3D:
return "texture3d";
case EbtIImage3D:
return "texture3d<int>";
case EbtUImage3D:
return "texture3d<uint>";
case EbtImage2DArray:
if (t->isUnsizedArray())
{
return "array_ref<texture2d>";
}
else
{
return "NOT_IMPLEMENTED";
}
case EbtIImage2DArray:
if (t->isUnsizedArray())
{
return "array_ref<texture2d<int>>";
}
else
{
return "NOT_IMPLEMENTED";
}
case EbtUImage2DArray:
if (t->isUnsizedArray())
{
return "array_ref<texture2d<uint>>";
}
else
{
return "NOT_IMPLEMENTED";
}
case EbtImageCube:
return "texturecube";
case EbtIImageCube:
return "texturecube<int>";
case EbtUImageCube:
return "texturecube<uint>";
case EbtImageCubeArray:
if (t->isUnsizedArray())
{
return "array_ref<texturecube>";
}
else
{
return "NOT_IMPLEMENTED";
}
case EbtIImageCubeArray:
if (t->isUnsizedArray())
{
return "array_ref<texturecube<int>>";
}
else
{
return "NOT_IMPLEMENTED";
}
case EbtUImageCubeArray:
if (t->isUnsizedArray())
{
return "array_ref<texturecube<uint>>";
}
else
{
return "NOT_IMPLEMENTED";
}
case EbtAtomicCounter:
return "atomic_uint";
case EbtSamplerVideoWEBGL:
return "$samplerVideoWEBGL";
default:
UNREACHABLE();
return "unknown type";
}
}
const char *getBuiltInMetalTypeNameString(const TType *t)
{
if (t->isMatrix())
{
switch (t->getCols())
{
case 2:
switch (t->getRows())
{
case 2:
return "float2";
case 3:
return "float2x3";
case 4:
return "float2x4";
default:
UNREACHABLE();
return nullptr;
}
case 3:
switch (t->getRows())
{
case 2:
return "float3x2";
case 3:
return "float3";
case 4:
return "float3x4";
default:
UNREACHABLE();
return nullptr;
}
case 4:
switch (t->getRows())
{
case 2:
return "float4x2";
case 3:
return "float4x3";
case 4:
return "float4";
default:
UNREACHABLE();
return nullptr;
}
default:
UNREACHABLE();
return nullptr;
}
}
if (t->isVector())
{
switch (t->getBasicType())
{
case EbtFloat:
switch (t->getNominalSize())
{
case 2:
return "float2";
case 3:
return "float3";
case 4:
return "float4";
default:
UNREACHABLE();
return nullptr;
}
case EbtInt:
switch (t->getNominalSize())
{
case 2:
return "int2";
case 3:
return "int3";
case 4:
return "int4";
default:
UNREACHABLE();
return nullptr;
}
case EbtBool:
switch (t->getNominalSize())
{
case 2:
return "bool2";
case 3:
return "bool3";
case 4:
return "bool4";
default:
UNREACHABLE();
return nullptr;
}
case EbtUInt:
switch (t->getNominalSize())
{
case 2:
return "uint2";
case 3:
return "uint3";
case 4:
return "uint4";
default:
UNREACHABLE();
return nullptr;
}
default:
UNREACHABLE();
return nullptr;
}
}
ASSERT(t->getBasicType() != EbtStruct);
ASSERT(t->getBasicType() != EbtInterfaceBlock);
return getBasicMetalString(t);
}
ImmutableString GetMetalTypeName(const TType &type, ShHashFunction64 hashFunction, NameMap *nameMap)
{
if (type.getBasicType() == EbtStruct)
return HashName(type.getStruct(), hashFunction, nameMap);
else
return ImmutableString(getBuiltInMetalTypeNameString(&type));
}
} // namespace sh

31
src/compiler/translator/TranslatorMetalDirect/TranslatorMetalUtils.h Normal file
Просмотреть файл

@ -0,0 +1,31 @@
//
// Copyright 2002 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_TRANSLATOR_METAL_UTILS_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_TRANSLATOR_METAL_UTILS_H_
#include "common/angleutils.h"
#include "common/debug.h"
#include "compiler/translator/BaseTypes.h"
#include "compiler/translator/Common.h"
#include "compiler/translator/HashNames.h"
#include "compiler/translator/ImmutableString.h"
#include "compiler/translator/SymbolUniqueId.h"
#include "compiler/translator/Types.h"
namespace sh
{
const char *getBasicMetalString(const TType *t);
const char *getBuiltInMetalTypeNameString(const TType *t);
ImmutableString GetMetalTypeName(const TType &type,
ShHashFunction64 hashFunction,
NameMap *nameMap);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_TRANSLATOR_METAL_UTILS_H_

51
src/compiler/translator/TranslatorMetalDirect/TransposeRowMajorMatrices.cpp Normal file
Просмотреть файл

@ -0,0 +1,51 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <algorithm>
#include <unordered_map>
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
#include "compiler/translator/TranslatorMetalDirect/IntermRebuild.h"
#include "compiler/translator/TranslatorMetalDirect/TransposeRowMajorMatrices.h"
#include "compiler/translator/tree_ops/SeparateDeclarations.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
class Rewriter : public TIntermRebuild
{
public:
Rewriter(TCompiler &compiler) : TIntermRebuild(compiler, true, false) {}
// PreResult visitConstantUnionPre(TIntermSymbol &symbolNode) override
// {
// const TVariable &var = symbolNode.variable();
// const TType &type = var.getType();
// //const TLayoutQualifier &layoutQualifier = type.getLayoutQualifier();
// if(type.isMatrix())
// {
// return *(new TIntermUnary(EOpTranspose, symbolNode.deepCopy(), nullptr));
// }
// return symbolNode;
// }
};
} // anonymous namespace
////////////////////////////////////////////////////////////////////////////////
bool sh::TransposeRowMajorMatricies(TCompiler &compiler, TIntermBlock &root)
{
Rewriter rewriter(compiler);
if (!rewriter.rebuildRoot(root))
{
return false;
}
return true;
}

21
src/compiler/translator/TranslatorMetalDirect/TransposeRowMajorMatrices.h Normal file
Просмотреть файл

@ -0,0 +1,21 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_TRANSPOSEROWMAJORMATRICIES_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_TRANSPOSEROWMAJORMATRICIES_H_
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/ProgramPrelude.h"
#include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
namespace sh
{
ANGLE_NO_DISCARD bool TransposeRowMajorMatricies(TCompiler &compiler, TIntermBlock &root);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_TRANSPOSEROWMAJORMATRICIES_H_

94
src/compiler/translator/TranslatorMetalDirect/WrapMain.cpp Normal file
Просмотреть файл

@ -0,0 +1,94 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include "compiler/translator/TranslatorMetalDirect/WrapMain.h"
#include "compiler/translator/Compiler.h"
#include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
class Wrapper : public TIntermTraverser
{
private:
IdGen &mIdGen;
public:
Wrapper(TSymbolTable &symbolTable, IdGen &idGen)
: TIntermTraverser(false, false, true, &symbolTable), mIdGen(idGen)
{}
bool visitBlock(Visit, TIntermBlock *blockNode) override
{
if (blockNode != getRootNode())
{
return true;
}
for (TIntermNode *node : *blockNode->getSequence())
{
if (TIntermFunctionDefinition *funcDefNode = node->getAsFunctionDefinition())
{
const TFunction &func = *funcDefNode->getFunction();
if (func.isMain())
{
visitMain(*blockNode, funcDefNode);
break;
}
}
}
return true;
}
private:
void visitMain(TIntermBlock &root, TIntermFunctionDefinition *funcDefNode)
{
const TFunction &func = *funcDefNode->getFunction();
ASSERT(func.isMain());
ASSERT(func.getReturnType().getBasicType() == TBasicType::EbtVoid);
ASSERT(func.getParamCount() == 0);
const TFunction &externalMainFunc = *funcDefNode->getFunction();
const TFunction &internalMainFunc = CloneFunction(*mSymbolTable, mIdGen, externalMainFunc);
TIntermFunctionPrototype *externalMainProto = funcDefNode->getFunctionPrototype();
TIntermFunctionPrototype *internalMainProto =
new TIntermFunctionPrototype(&internalMainFunc);
TIntermBlock *externalMainBody = new TIntermBlock();
externalMainBody->appendStatement(
TIntermAggregate::CreateFunctionCall(internalMainFunc, new TIntermSequence()));
TIntermBlock *internalMainBody = funcDefNode->getBody();
TIntermFunctionDefinition *externalMainDef =
new TIntermFunctionDefinition(externalMainProto, externalMainBody);
TIntermFunctionDefinition *internalMainDef =
new TIntermFunctionDefinition(internalMainProto, internalMainBody);
mMultiReplacements.push_back(NodeReplaceWithMultipleEntry(
&root, funcDefNode, TIntermSequence{internalMainDef, externalMainDef}));
}
};
} // namespace
bool sh::WrapMain(TCompiler &compiler, IdGen &idGen, TIntermBlock &root)
{
TSymbolTable &symbolTable = compiler.getSymbolTable();
Wrapper wrapper(symbolTable, idGen);
root.traverse(&wrapper);
if (!wrapper.updateTree(&compiler, &root))
{
return false;
}
return true;
}

28
src/compiler/translator/TranslatorMetalDirect/WrapMain.h Normal file
Просмотреть файл

@ -0,0 +1,28 @@
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#ifndef COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_WRAPMAIN_H_
#define COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_WRAPMAIN_H_
#include "compiler/translator/TranslatorMetalDirect/IdGen.h"
#include "compiler/translator/tree_util/IntermTraverse.h"
namespace sh
{
// Changes
// void main(args) { main-body }
// To
// void FRESH_NAME(args) { main-body }
// void main(args) { FRESH_NAME(args); }
//
// This transformation is useful if the original `main` has multiple return paths because this
// reduces down to a single path in the new `main`. Nice for inserting cleanup code in `main`.
ANGLE_NO_DISCARD bool WrapMain(TCompiler &compiler, IdGen &idGen, TIntermBlock &root);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TRANSLATORMETALDIRECT_WRAPMAIN_H_

12
src/compiler/translator/TranslatorVulkan.cpp
Просмотреть файл

@ -21,19 +21,19 @@
#include "compiler/translator/OutputVulkanGLSL.h"
#include "compiler/translator/StaticType.h"
#include "compiler/translator/glslang_wrapper.h"
#include "compiler/translator/tree_ops/RemoveAtomicCounterBuiltins.h"
#include "compiler/translator/tree_ops/RemoveInactiveInterfaceVariables.h"
#include "compiler/translator/tree_ops/RewriteAtomicCounters.h"
#include "compiler/translator/tree_ops/RewriteCubeMapSamplersAs2DArray.h"
#include "compiler/translator/tree_ops/RewriteDfdy.h"
#include "compiler/translator/tree_ops/RewriteStructSamplers.h"
#include "compiler/translator/tree_ops/vulkan/DeclarePerVertexBlocks.h"
#include "compiler/translator/tree_ops/vulkan/FlagSamplersWithTexelFetch.h"
#include "compiler/translator/tree_ops/vulkan/MonomorphizeUnsupportedFunctionsInVulkanGLSL.h"
#include "compiler/translator/tree_ops/vulkan/RemoveAtomicCounterBuiltins.h"
#include "compiler/translator/tree_ops/vulkan/RemoveInactiveInterfaceVariables.h"
#include "compiler/translator/tree_ops/vulkan/ReplaceForShaderFramebufferFetch.h"
#include "compiler/translator/tree_ops/vulkan/RewriteArrayOfArrayOfOpaqueUniforms.h"
#include "compiler/translator/tree_ops/vulkan/RewriteAtomicCounters.h"
#include "compiler/translator/tree_ops/vulkan/RewriteCubeMapSamplersAs2DArray.h"
#include "compiler/translator/tree_ops/vulkan/RewriteDfdy.h"
#include "compiler/translator/tree_ops/vulkan/RewriteInterpolateAtOffset.h"
#include "compiler/translator/tree_ops/vulkan/RewriteR32fImages.h"
#include "compiler/translator/tree_ops/vulkan/RewriteStructSamplers.h"
#include "compiler/translator/tree_ops/vulkan/SeparateStructFromUniformDeclarations.h"
#include "compiler/translator/tree_util/BuiltIn.h"
#include "compiler/translator/tree_util/DriverUniform.h"

1
src/compiler/translator/Types.h
Просмотреть файл

@ -237,6 +237,7 @@ class TType
bool isVector() const { return primarySize > 1 && secondarySize == 1; }
bool isVectorArray() const { return primarySize > 1 && secondarySize == 1 && isArray(); }
bool isRank0() const { return primarySize == 1 && secondarySize == 1; }
bool isScalar() const
{
return primarySize == 1 && secondarySize == 1 && !mStructure && !isArray();

127
src/compiler/translator/tree_ops/NameNamelessUniformBuffers.cpp Normal file
Просмотреть файл

@ -0,0 +1,127 @@
//
// Copyright 2019 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
// NameNamelessUniformBuffers: Gives nameless uniform buffer variables internal names.
//
#include "compiler/translator/tree_ops/NameNamelessUniformBuffers.h"
#include "compiler/translator/SymbolTable.h"
#include "compiler/translator/tree_util/IntermNode_util.h"
#include "compiler/translator/tree_util/IntermTraverse.h"
namespace sh
{
namespace
{
// Traverse uniform buffer declarations and give name to nameless declarations. Keeps track of
// the interface fields which will be used in the source without the interface block variable name
// and replaces them with name.field.
class NameUniformBufferVariablesTraverser : public TIntermTraverser
{
public:
explicit NameUniformBufferVariablesTraverser(TSymbolTable *symbolTable)
: TIntermTraverser(true, false, false, symbolTable)
{}
bool visitDeclaration(Visit visit, TIntermDeclaration *decl) override
{
ASSERT(visit == PreVisit);
const TIntermSequence &sequence = *(decl->getSequence());
TIntermTyped *variableNode = sequence.front()->getAsTyped();
const TType &type = variableNode->getType();
// If it's an interface block, it may have to be converted if it contains any row-major
// fields.
if (!type.isInterfaceBlock())
{
return true;
}
// Multi declaration statements are already separated, so there can only be one variable
// here.
ASSERT(sequence.size() == 1);
const TVariable *variable = &variableNode->getAsSymbolNode()->variable();
if (variable->symbolType() != SymbolType::Empty)
{
return false;
}
TIntermDeclaration *newDeclaration = new TIntermDeclaration;
TVariable *newVariable = new TVariable(mSymbolTable, kEmptyImmutableString, &type,
SymbolType::AngleInternal, variable->extension());
newDeclaration->appendDeclarator(new TIntermSymbol(newVariable));
queueReplacement(newDeclaration, OriginalNode::IS_DROPPED);
// It's safe to key the map with the interface block, as there couldn't have been multiple
// declarations with this interface block (as the variable is nameless), so for nameless
// uniform buffers, the interface block is unique.
mNamelessUniformBuffersMap[type.getInterfaceBlock()] = newVariable;
return false;
}
void visitSymbol(TIntermSymbol *symbol) override
{
const TType &type = symbol->getType();
// The symbols we are looking for have the interface block pointer set, but are not
// interface blocks. These are references to fields of nameless uniform buffers.
if (type.isInterfaceBlock() || type.getInterfaceBlock() == nullptr)
{
return;
}
const TInterfaceBlock *block = type.getInterfaceBlock();
// If block variable is not nameless, there's nothing to do.
if (mNamelessUniformBuffersMap.count(block) == 0)
{
return;
}
const ImmutableString symbolName = symbol->getName();
// Find which field it is
const TVector<TField *> fields = block->fields();
for (size_t fieldIndex = 0; fieldIndex < fields.size(); ++fieldIndex)
{
const TField *field = fields[fieldIndex];
if (field->name() != symbolName)
{
continue;
}
// Replace this node with a binary node that indexes the named uniform buffer.
TIntermSymbol *namedUniformBuffer =
new TIntermSymbol(mNamelessUniformBuffersMap[block]);
TIntermBinary *replacement =
new TIntermBinary(EOpIndexDirectInterfaceBlock, namedUniformBuffer,
CreateIndexNode(static_cast<uint32_t>(fieldIndex)));
queueReplacement(replacement, OriginalNode::IS_DROPPED);
return;
}
UNREACHABLE();
}
private:
// A map from nameless uniform buffers to their named replacements.
std::unordered_map<const TInterfaceBlock *, const TVariable *> mNamelessUniformBuffersMap;
};
} // anonymous namespace
bool NameNamelessUniformBuffers(TCompiler *compiler, TIntermBlock *root, TSymbolTable *symbolTable)
{
NameUniformBufferVariablesTraverser nameUniformBufferVariables(symbolTable);
root->traverse(&nameUniformBufferVariables);
return nameUniformBufferVariables.updateTree(compiler, root);
}
} // namespace sh

32
src/compiler/translator/tree_ops/NameNamelessUniformBuffers.h Normal file
Просмотреть файл

@ -0,0 +1,32 @@
//
// Copyright 2019 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
// NameNamelessUniformBuffers: Gives nameless uniform buffer variables internal names.
//
// For example:
// uniform UniformBuffer { int a; };
// x = a;
// becomes:
// uniform UniformBuffer { int a; } s123;
// x = s123.a;
//
#ifndef COMPILER_TRANSLATOR_TREEOPS_NAMENAMELESSUNIFORMBUFFERS_H_
#define COMPILER_TRANSLATOR_TREEOPS_NAMENAMELESSUNIFORMBUFFERS_H_
#include "common/angleutils.h"
namespace sh
{
class TCompiler;
class TIntermBlock;
class TSymbolTable;
ANGLE_NO_DISCARD bool NameNamelessUniformBuffers(TCompiler *compiler,
TIntermBlock *root,
TSymbolTable *symbolTable);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TREEOPS_NAMENAMELESSUNIFORMBUFFERS_H_

2
src/compiler/translator/tree_ops/vulkan/RemoveAtomicCounterBuiltins.cpp → src/compiler/translator/tree_ops/RemoveAtomicCounterBuiltins.cpp
Просмотреть файл

@ -6,7 +6,7 @@
// RemoveAtomicCounterBuiltins: Remove atomic counter builtins.
//
#include "compiler/translator/tree_ops/vulkan/RemoveAtomicCounterBuiltins.h"
#include "compiler/translator/tree_ops/RemoveAtomicCounterBuiltins.h"
#include "compiler/translator/Compiler.h"
#include "compiler/translator/tree_util/IntermTraverse.h"

6
src/compiler/translator/tree_ops/vulkan/RemoveAtomicCounterBuiltins.h → src/compiler/translator/tree_ops/RemoveAtomicCounterBuiltins.h
Просмотреть файл

@ -8,8 +8,8 @@
// atomic counters are actually in use. This pass removes the builtins and
// asserts no atomic counters are declared.
#ifndef COMPILER_TRANSLATOR_TREEOPS_VULKAN_REMOVEATOMICCOUNTERBUILTINS_H_
#define COMPILER_TRANSLATOR_TREEOPS_VULKAN_REMOVEATOMICCOUNTERBUILTINS_H_
#ifndef COMPILER_TRANSLATOR_TREEOPS_REMOVEATOMICCOUNTERBUILTINS_H_
#define COMPILER_TRANSLATOR_TREEOPS_REMOVEATOMICCOUNTERBUILTINS_H_
#include "common/angleutils.h"
@ -21,4 +21,4 @@ class TIntermBlock;
ANGLE_NO_DISCARD bool RemoveAtomicCounterBuiltins(TCompiler *compiler, TIntermBlock *root);
} // namespace sh
#endif // COMPILER_TRANSLATOR_TREEOPS_VULKAN_REMOVEATOMICCOUNTERBUILTINS_H_
#endif // COMPILER_TRANSLATOR_TREEOPS_REMOVEATOMICCOUNTERBUILTINS_H_

2
src/compiler/translator/tree_ops/vulkan/RemoveInactiveInterfaceVariables.cpp → src/compiler/translator/tree_ops/RemoveInactiveInterfaceVariables.cpp
Просмотреть файл

@ -7,7 +7,7 @@
// Drop shader interface variable declarations for those that are inactive.
//
#include "compiler/translator/tree_ops/vulkan/RemoveInactiveInterfaceVariables.h"
#include "compiler/translator/tree_ops/RemoveInactiveInterfaceVariables.h"
#include "compiler/translator/SymbolTable.h"
#include "compiler/translator/tree_util/IntermTraverse.h"

6
src/compiler/translator/tree_ops/vulkan/RemoveInactiveInterfaceVariables.h → src/compiler/translator/tree_ops/RemoveInactiveInterfaceVariables.h
Просмотреть файл

@ -12,8 +12,8 @@
// instead of relying on what was gathered during CollectVariables.
//
#ifndef COMPILER_TRANSLATOR_TREEOPS_VULKAN_REMOVEINACTIVEVARIABLES_H_
#define COMPILER_TRANSLATOR_TREEOPS_VULKAN_REMOVEINACTIVEVARIABLES_H_
#ifndef COMPILER_TRANSLATOR_TREEOPS_REMOVEINACTIVEVARIABLES_H_
#define COMPILER_TRANSLATOR_TREEOPS_REMOVEINACTIVEVARIABLES_H_
#include "common/angleutils.h"
@ -38,4 +38,4 @@ ANGLE_NO_DISCARD bool RemoveInactiveInterfaceVariables(
} // namespace sh
#endif // COMPILER_TRANSLATOR_TREEOPS_VULKAN_REMOVEINACTIVEVARIABLES_H_
#endif // COMPILER_TRANSLATOR_TREEOPS_REMOVEINACTIVEVARIABLES_H_

Некоторые файлы не были показаны из-за слишком большого количества измененных файлов Показать больше