Fixes to compute workgroup sizes and barriers.
Improved extraction of entry point name and workgroup size from SPIR-V. Consolidate to a single ThirdPartyConfig.md document. MSL enhancements to nested function use of globals. Support customizing MSL based on iOS or macOS platform. MSL threadgroup barrier memory scope only on iOS MSL 2.0. Update to latest SPIRV libraries.
This commit is contained in:
Bill Hollings
parent
097b9f29d1
commit
9b185a9c3a
@ -1,5 +1,5 @@
|
||||
<a class="site-logo" href="https://github.com/KhronosGroup/MoltenVK" title="MoltenVK">
|
||||
<img src="../Docs/images/MoltenVK-Logo-Banner.png" alt="MoltenV" style="width:256px;height:auto">
|
||||
<img src="images/MoltenVK-Logo-Banner.png" alt="MoltenVK" style="width:256px;height:auto">
|
||||
</a>
|
||||
|
||||
Copyright (c) 2014-2017 [The Brenwill Workshop Ltd.](http://www.brenwill.com)
|
||||
@ -11,6 +11,14 @@ For best results, use a Markdown reader.*
|
||||
Table of Contents
|
||||
-----------------
|
||||
|
||||
- *Vulkan-Hpp*
|
||||
- [Using the *Vulkan-Hpp* Spec Repository with **MoltenVK**](#install_vulkan_spec)
|
||||
- [Updating the *Vulkan-Hpp* library version](#update_vulkan_spec)
|
||||
|
||||
- *Vulkan-LoaderAndValidationLayers*
|
||||
- [Using the *Vulkan-LoaderAndValidationLayers* Repository with **MoltenVK**](#install_vulkan_lvl)
|
||||
- [Updating the *Vulkan-LoaderAndValidationLayers* library version](#update_vulkan_lvl)
|
||||
|
||||
- *SPIRV-Cross*
|
||||
- [Using the *SPIRV-Cross* library with **MoltenVKShaderConverter**](#install_spirv-cross)
|
||||
- [Updating the *SPIRV-Cross* library version](#update_spirv-cross)
|
||||
@ -27,6 +35,81 @@ Table of Contents
|
||||
- [Adding the *glslang* library to a new *Xcode* project](#add_glslang)
|
||||
|
||||
|
||||
<a name="install_vulkan_spec"></a>
|
||||
Using the *Vulkan-Hpp* Spec Repository with *MoltenVK*
|
||||
------------------------------------------------------
|
||||
|
||||
**MoltenVK** uses the official *Khronos Vulkan* specification repository to provide the standard
|
||||
*Vulkan* API header files and *Vulkan Specification* documentation.
|
||||
|
||||
To add the *Khronos Vulkan* specification repository to **MoltenVK**, open a *Terminal*
|
||||
session and perform the following command-line steps:
|
||||
|
||||
1. Ensure you have `python3` and `asciidoctor` installed:
|
||||
|
||||
brew install python3
|
||||
sudo gem install asciidoctor
|
||||
|
||||
2. If you used the `--recursive` option when cloning the `MoltenVK` repository, you should already
|
||||
have the `Vulkan-Hpp` submodule, and you can skip to *Step 3* below. If you did **_not_**
|
||||
use the `--recursive` option when cloning the `MoltenVK` repository, retrieve the `Vulkan-Hpp`
|
||||
submodule into the `External` directory as follows, from within the `MoltenVK` repository directory:
|
||||
|
||||
git submodule update --init --recursive External/Vulkan-Hpp
|
||||
|
||||
3. In the `Externals` folder within the `MoltenVK` repository, build the spec and header files
|
||||
as follows from the main directory of this `MoltenVK` repository:
|
||||
|
||||
cd External
|
||||
./makeVulkanSpec
|
||||
|
||||
|
||||
|
||||
<a name="update_vulkan_spec"></a>
|
||||
Updating the *Vulkan-Hpp* library version
|
||||
-----------------------------------------
|
||||
|
||||
If you are developing enhancements to **MoltenVK**, you can update the version of `Vulkan-Hpp`
|
||||
used by **MoltenVK** to the latest version available by re-cloning and re-building the
|
||||
`Vulkan-Hpp` submodule using the `getLatestVulkanSpec` script:
|
||||
|
||||
cd External
|
||||
./getLatestVulkanSpec
|
||||
|
||||
The updated version will then be "locked in" the next time the `MoltenVK` repository is committed to `git`.
|
||||
|
||||
|
||||
|
||||
<a name="install_vulkan_lvl"></a>
|
||||
Using the *Vulkan-LoaderAndValidationLayers* Spec Repository with *MoltenVK*
|
||||
----------------------------------------------------------------------------
|
||||
|
||||
**MoltenVK** uses the *Khronos Vulkan Loader and Validation Layers* repository to allow **MoltenVK**
|
||||
to act as an *Installable Client Driver* to support the *Vulkan Loader API*.
|
||||
|
||||
If you used the `--recursive` option when cloning the `MoltenVK` repository, you should already
|
||||
have the `Vulkan-LoaderAndValidationLayers` submodule. If you did **_not_** use the `--recursive`
|
||||
option when cloning the `MoltenVK` repository, retrieve the `Vulkan-LoaderAndValidationLayers`
|
||||
submodule into the `External` directory as follows, from within the `MoltenVK` repository directory:
|
||||
|
||||
git submodule update --init External/Vulkan-LoaderAndValidationLayers
|
||||
|
||||
|
||||
|
||||
<a name="update_vulkan_lvl"></a>
|
||||
Updating the *Vulkan-LoaderAndValidationLayers* library version
|
||||
---------------------------------------------------------------
|
||||
|
||||
If you are developing enhancements to **MoltenVK**, you can update the version of `Vulkan-LoaderAndValidationLayers`
|
||||
used by **MoltenVK** to the latest version available by re-cloning and re-building the `Vulkan-LoaderAndValidationLayers`
|
||||
submodule using the `getLatestVulkanLVL` script:
|
||||
|
||||
cd External
|
||||
./getLatestVulkanLVL
|
||||
|
||||
The updated version will then be "locked in" the next time the `MoltenVK` repository is committed to `git`.
|
||||
|
||||
|
||||
|
||||
<a name="install_spirv-cross"></a>
|
||||
Using the *SPIRV-Cross* library with *MoltenVKShaderConverter*
|
||||
@ -304,3 +387,4 @@ However, to add the `glslang` library to a new *Xcode* project:
|
||||
[above](#install_glslang)) with simply `glslang` (the name of the symlink). Be sure you only
|
||||
replace the part of the path that matches the `path-to-glslang-repo-folder`. Do not replace
|
||||
any part of the path that indicates a subfolder within that repository folder.
|
||||
|
||||
2
External/SPIRV-Cross
vendored
2
External/SPIRV-Cross
vendored
@ -1 +1 @@
|
||||
Subproject commit bcc96d8c7011d5be504174ddbbadcacad227ea89
|
||||
Subproject commit 95910ddd5aa03cbd7188fc7c107f9cc893136f10
|
||||
2
External/SPIRV-Tools
vendored
2
External/SPIRV-Tools
vendored
@ -1 +1 @@
|
||||
Subproject commit 726573a0e546e0d7ce5431db00d309542488406d
|
||||
Subproject commit 1acce99255ee28a582776372593e91b585d3c011
|
||||
2
External/glslang
vendored
2
External/glslang
vendored
@ -1 +1 @@
|
||||
Subproject commit 6a14f78061a3c4e6d9d881cadb63c252cb6e5c0a
|
||||
Subproject commit 046bae0babd17ecc19fc7cbe40c35aa13ac2ee65
|
||||
@ -295,7 +295,6 @@
|
||||
A95870F71C90D29F009EB096 /* MVKCommandResourceFactory.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKCommandResourceFactory.mm; sourceTree = "<group>"; };
|
||||
A95B7D671D3EE486003183D3 /* MVKCommandEncoderState.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKCommandEncoderState.h; sourceTree = "<group>"; };
|
||||
A95B7D681D3EE486003183D3 /* MVKCommandEncoderState.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKCommandEncoderState.mm; sourceTree = "<group>"; };
|
||||
A97773311F420A21006A3601 /* ThirdPartyConfig.md */ = {isa = PBXFileReference; lastKnownFileType = net.daringfireball.markdown; path = ThirdPartyConfig.md; sourceTree = "<group>"; };
|
||||
A98149411FB6A3F7005F00B4 /* MVKBaseObject.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = MVKBaseObject.cpp; sourceTree = "<group>"; };
|
||||
A98149421FB6A3F7005F00B4 /* MVKBaseObject.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKBaseObject.h; sourceTree = "<group>"; };
|
||||
A98149431FB6A3F7005F00B4 /* MVKEnvironment.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKEnvironment.h; sourceTree = "<group>"; };
|
||||
@ -513,7 +512,6 @@
|
||||
A9F0429B1FB4CF82009FCCB8 /* Common */,
|
||||
A9AC84381D061E7000E2CA97 /* include */,
|
||||
A9C86CB61C55B8350096CAF2 /* MoltenVKShaderConverter.xcodeproj */,
|
||||
A97773311F420A21006A3601 /* ThirdPartyConfig.md */,
|
||||
A9F55D2E198BE6A7004EC31B /* Products */,
|
||||
);
|
||||
sourceTree = "<group>";
|
||||
|
||||
@ -60,11 +60,9 @@ public:
|
||||
protected:
|
||||
void handleCompilationError(NSError* err, const char* opDesc);
|
||||
MTLFunctionConstant* getFunctionConstant(NSArray<MTLFunctionConstant*>* mtlFCs, NSUInteger mtlFCID);
|
||||
const std::string cleanMSLFunctionName(const std::string& name);
|
||||
|
||||
id<MTLLibrary> _mtlLibrary;
|
||||
std::unordered_map<std::string, std::string> _mtlFunctionNameMap;
|
||||
SPIRVLocalSizesByEntryPointName _localSizes;
|
||||
SPIRVEntryPointsByName _entryPoints;
|
||||
};
|
||||
|
||||
|
||||
|
||||
@ -35,8 +35,8 @@ MVKMTLFunction MVKShaderLibrary::getMTLFunction(const VkPipelineShaderStageCreat
|
||||
|
||||
// Ensure the function name is compatible with Metal (Metal does not allow main()
|
||||
// as a function name), and retrieve the unspecialized Metal function with that name.
|
||||
string funcName = cleanMSLFunctionName(pShaderStage->pName);
|
||||
NSString* mtlFuncName = @(funcName.c_str());
|
||||
SPIRVEntryPoint& ep = _entryPoints[pShaderStage->pName];
|
||||
NSString* mtlFuncName = @(ep.mtlFunctionName.c_str());
|
||||
|
||||
NSTimeInterval startTime = _device->getPerformanceTimestamp();
|
||||
id<MTLFunction> mtlFunc = [[_mtlLibrary newFunctionWithName: mtlFuncName] autorelease];
|
||||
@ -78,11 +78,10 @@ MVKMTLFunction MVKShaderLibrary::getMTLFunction(const VkPipelineShaderStageCreat
|
||||
}
|
||||
}
|
||||
} else {
|
||||
mvkNotifyErrorWithText(VK_ERROR_INITIALIZATION_FAILED, "Shader module does not contain an entry point named '%s'.", funcName.c_str());
|
||||
mvkNotifyErrorWithText(VK_ERROR_INITIALIZATION_FAILED, "Shader module does not contain an entry point named '%s'.", mtlFuncName.UTF8String);
|
||||
}
|
||||
|
||||
SPIRVLocalSize wgSize = _localSizes[funcName];
|
||||
return { mtlFunc, MTLSizeMake(wgSize.width, wgSize.height, wgSize.depth) };
|
||||
return { mtlFunc, MTLSizeMake(ep.workgroupSize.width, ep.workgroupSize.height, ep.workgroupSize.depth) };
|
||||
}
|
||||
|
||||
// Returns the MTLFunctionConstant with the specified ID from the specified array of function constants.
|
||||
@ -92,12 +91,6 @@ MTLFunctionConstant* MVKShaderLibrary::getFunctionConstant(NSArray<MTLFunctionCo
|
||||
return nil;
|
||||
}
|
||||
|
||||
// Cleans the specified shader function name so it can be used as as an MSL function name.
|
||||
const std::string MVKShaderLibrary::cleanMSLFunctionName(const std::string& funcName) {
|
||||
string cleanName = _mtlFunctionNameMap[funcName];
|
||||
return cleanName.empty() ? funcName : cleanName;
|
||||
}
|
||||
|
||||
MVKShaderLibrary::MVKShaderLibrary(MVKDevice* device, SPIRVToMSLConverter& mslConverter) : MVKBaseDeviceObject(device) {
|
||||
NSTimeInterval startTime = _device->getPerformanceTimestamp();
|
||||
@autoreleasepool {
|
||||
@ -110,8 +103,7 @@ MVKShaderLibrary::MVKShaderLibrary(MVKDevice* device, SPIRVToMSLConverter& mslCo
|
||||
}
|
||||
_device->addShaderCompilationEventPerformance(_device->_shaderCompilationPerformance.mslCompile, startTime);
|
||||
|
||||
_mtlFunctionNameMap = mslConverter.getEntryPointNameMap();
|
||||
_localSizes = mslConverter.getLocalSizes();
|
||||
_entryPoints = mslConverter.getEntryPoints();
|
||||
}
|
||||
|
||||
MVKShaderLibrary::MVKShaderLibrary(MVKDevice* device,
|
||||
@ -225,9 +217,8 @@ MVKShaderModule::MVKShaderModule(MVKDevice* device,
|
||||
}
|
||||
case kMVKMagicNumberMSLSourceCode: { // MSL source code
|
||||
uintptr_t pMSLCode = uintptr_t(pCreateInfo->pCode) + sizeof(MVKMSLSPIRVHeader);
|
||||
unordered_map<string, string> entryPointNameMap;
|
||||
SPIRVLocalSizesByEntryPointName localSizes;
|
||||
_converter.setMSL((char*)pMSLCode, entryPointNameMap, localSizes);
|
||||
SPIRVEntryPointsByName entryPoints;
|
||||
_converter.setMSL((char*)pMSLCode, entryPoints);
|
||||
_defaultLibrary = new MVKShaderLibrary(_device, _converter);
|
||||
break;
|
||||
}
|
||||
|
||||
@ -1,96 +0,0 @@
|
||||
<a class="site-logo" href="https://github.com/KhronosGroup/MoltenVK" title="MoltenVK">
|
||||
<img src="../Docs/images/MoltenVK-Logo-Banner.png" alt="MoltenVK" style="width:256px;height:auto">
|
||||
</a>
|
||||
|
||||
Copyright (c) 2014-2017 [The Brenwill Workshop Ltd.](http://www.brenwill.com)
|
||||
|
||||
*This document is written in [Markdown](http://en.wikipedia.org/wiki/Markdown) format.
|
||||
For best results, use a Markdown reader.*
|
||||
|
||||
|
||||
Table of Contents
|
||||
-----------------
|
||||
|
||||
- *Vulkan-Hpp*
|
||||
- [Using the *Vulkan-Hpp* Spec Repository with **MoltenVK**](#install_vulkan_spec)
|
||||
- [Updating the *Vulkan-Hpp* library version](#update_vulkan_spec)
|
||||
|
||||
- *Vulkan-LoaderAndValidationLayers*
|
||||
- [Using the *Vulkan-LoaderAndValidationLayers* Repository with **MoltenVK**](#install_vulkan_lvl)
|
||||
- [Updating the *Vulkan-LoaderAndValidationLayers* library version](#update_vulkan_lvl)
|
||||
|
||||
|
||||
<a name="install_vulkan_spec"></a>
|
||||
Using the *Vulkan-Hpp* Spec Repository with *MoltenVK*
|
||||
------------------------------------------------------
|
||||
|
||||
**MoltenVK** uses the official *Khronos Vulkan* specification repository to provide the standard
|
||||
*Vulkan* API header files and *Vulkan Specification* documentation.
|
||||
|
||||
To add the *Khronos Vulkan* specification repository to **MoltenVK**, open a *Terminal*
|
||||
session and perform the following command-line steps:
|
||||
|
||||
1. Ensure you have `python3` and `asciidoctor` installed:
|
||||
|
||||
brew install python3
|
||||
sudo gem install asciidoctor
|
||||
|
||||
2. If you used the `--recursive` option when cloning the `MoltenVK` repository, you should already
|
||||
have the `Vulkan-Hpp` submodule, and you can skip to *Step 3* below. If you did **_not_**
|
||||
use the `--recursive` option when cloning the `MoltenVK` repository, retrieve the `Vulkan-Hpp`
|
||||
submodule into the `External` directory as follows, from within the `MoltenVK` repository directory:
|
||||
|
||||
git submodule update --init --recursive External/Vulkan-Hpp
|
||||
|
||||
3. In the `Externals` folder within the `MoltenVK` repository, build the spec and header files
|
||||
as follows from the main directory of this `MoltenVK` repository:
|
||||
|
||||
cd External
|
||||
./makeVulkanSpec
|
||||
|
||||
|
||||
|
||||
<a name="update_vulkan_spec"></a>
|
||||
Updating the *Vulkan-Hpp* library version
|
||||
-----------------------------------------
|
||||
|
||||
If you are developing enhancements to **MoltenVK**, you can update the version of `Vulkan-Hpp`
|
||||
used by **MoltenVK** to the latest version available by re-cloning and re-building the
|
||||
`Vulkan-Hpp` submodule using the `getLatestVulkanSpec` script:
|
||||
|
||||
cd External
|
||||
./getLatestVulkanSpec
|
||||
|
||||
The updated version will then be "locked in" the next time the `MoltenVK` repository is committed to `git`.
|
||||
|
||||
|
||||
|
||||
<a name="install_vulkan_lvl"></a>
|
||||
Using the *Vulkan-LoaderAndValidationLayers* Spec Repository with *MoltenVK*
|
||||
----------------------------------------------------------------------------
|
||||
|
||||
**MoltenVK** uses the *Khronos Vulkan Loader and Validation Layers* repository to allow **MoltenVK**
|
||||
to act as an *Installable Client Driver* to support the *Vulkan Loader API*.
|
||||
|
||||
If you used the `--recursive` option when cloning the `MoltenVK` repository, you should already
|
||||
have the `Vulkan-LoaderAndValidationLayers` submodule. If you did **_not_** use the `--recursive`
|
||||
option when cloning the `MoltenVK` repository, retrieve the `Vulkan-LoaderAndValidationLayers`
|
||||
submodule into the `External` directory as follows, from within the `MoltenVK` repository directory:
|
||||
|
||||
git submodule update --init External/Vulkan-LoaderAndValidationLayers
|
||||
|
||||
|
||||
|
||||
<a name="update_vulkan_lvl"></a>
|
||||
Updating the *Vulkan-LoaderAndValidationLayers* library version
|
||||
---------------------------------------------------------------
|
||||
|
||||
If you are developing enhancements to **MoltenVK**, you can update the version of `Vulkan-LoaderAndValidationLayers`
|
||||
used by **MoltenVK** to the latest version available by re-cloning and re-building the `Vulkan-LoaderAndValidationLayers`
|
||||
submodule using the `getLatestVulkanLVL` script:
|
||||
|
||||
cd External
|
||||
./getLatestVulkanLVL
|
||||
|
||||
The updated version will then be "locked in" the next time the `MoltenVK` repository is committed to `git`.
|
||||
|
||||
@ -132,6 +132,7 @@
|
||||
/* End PBXContainerItemProxy section */
|
||||
|
||||
/* Begin PBXFileReference section */
|
||||
A9077A4F1FF2D17700BE1757 /* ThirdPartyConfig.md */ = {isa = PBXFileReference; lastKnownFileType = net.daringfireball.markdown; name = ThirdPartyConfig.md; path = Docs/ThirdPartyConfig.md; sourceTree = "<group>"; };
|
||||
A92DB3E41CE0F37D00FBC835 /* README.md */ = {isa = PBXFileReference; lastKnownFileType = net.daringfireball.markdown; path = README.md; sourceTree = "<group>"; };
|
||||
A92DB3E51CE0F37D00FBC835 /* LICENSE */ = {isa = PBXFileReference; lastKnownFileType = text; path = LICENSE; sourceTree = "<group>"; };
|
||||
A92DB3E61CE0F37D00FBC835 /* Whats_New.md */ = {isa = PBXFileReference; lastKnownFileType = net.daringfireball.markdown; name = Whats_New.md; path = Docs/Whats_New.md; sourceTree = "<group>"; };
|
||||
@ -167,6 +168,7 @@
|
||||
A92DB3E41CE0F37D00FBC835 /* README.md */,
|
||||
A92DB3E51CE0F37D00FBC835 /* LICENSE */,
|
||||
A98149E51FB78829005F00B4 /* MoltenVK_Runtime_UserGuide.md */,
|
||||
A9077A4F1FF2D17700BE1757 /* ThirdPartyConfig.md */,
|
||||
A92DB3E61CE0F37D00FBC835 /* Whats_New.md */,
|
||||
);
|
||||
name = Docs;
|
||||
|
||||
@ -117,9 +117,7 @@ MVK_PUBLIC_SYMBOL void SPIRVToMSLConverterContext::alignUsageWith(SPIRVToMSLConv
|
||||
#pragma mark SPIRVToMSLConverter
|
||||
|
||||
/** Populates content extracted from the SPRI-V compiler. */
|
||||
void populateFromCompiler(spirv_cross::Compiler& compiler,
|
||||
unordered_map<string, string>& entryPointNameMap,
|
||||
SPIRVLocalSizesByEntryPointName& localSizes);
|
||||
void populateFromCompiler(spirv_cross::Compiler& compiler, SPIRVEntryPointsByName& entryPoints);
|
||||
|
||||
MVK_PUBLIC_SYMBOL void SPIRVToMSLConverter::setSPIRV(const vector<uint32_t>& spirv) { _spirv = spirv; }
|
||||
|
||||
@ -175,6 +173,14 @@ MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverter::convert(SPIRVToMSLConverterContext&
|
||||
// Establish the MSL options for the compiler
|
||||
// This needs to be done in two steps...for CompilerMSL and its superclass.
|
||||
auto mslOpts = mslCompiler.get_options();
|
||||
|
||||
#if MVK_MACOS
|
||||
mslOpts.platform = spirv_cross::CompilerMSL::Options::macOS;
|
||||
#endif
|
||||
#if MVK_IOS
|
||||
mslOpts.platform = spirv_cross::CompilerMSL::Options::iOS;
|
||||
#endif
|
||||
|
||||
mslOpts.msl_version = context.options.mslVersion;
|
||||
mslOpts.enable_point_size_builtin = context.options.isRenderingPoints;
|
||||
mslOpts.resolve_specialized_array_lengths = true;
|
||||
@ -198,7 +204,7 @@ MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverter::convert(SPIRVToMSLConverterContext&
|
||||
}
|
||||
|
||||
// Populate content extracted from the SPRI-V compiler.
|
||||
populateFromCompiler(mslCompiler, _entryPointNameMap, _localSizes);
|
||||
populateFromCompiler(mslCompiler, _entryPoints);
|
||||
|
||||
// To check GLSL conversion
|
||||
if (shouldLogGLSL) {
|
||||
@ -280,24 +286,20 @@ void SPIRVToMSLConverter::logSource(string& src, const char* srcLang, const char
|
||||
|
||||
#pragma mark Support functions
|
||||
|
||||
void populateFromCompiler(spirv_cross::Compiler& compiler,
|
||||
unordered_map<string, string>& entryPointNameMap,
|
||||
SPIRVLocalSizesByEntryPointName& localSizes) {
|
||||
void populateFromCompiler(spirv_cross::Compiler& compiler, SPIRVEntryPointsByName& entryPoints) {
|
||||
|
||||
uint32_t minDim = 1;
|
||||
entryPointNameMap.clear();
|
||||
localSizes.clear();
|
||||
entryPoints.clear();
|
||||
for (string& epOrigName : compiler.get_entry_points()) {
|
||||
auto& ep = compiler.get_entry_point(epOrigName);
|
||||
auto& spvEP = compiler.get_entry_point(epOrigName);
|
||||
auto& wgSize = spvEP.workgroup_size;
|
||||
|
||||
entryPointNameMap[epOrigName] = ep.name;
|
||||
|
||||
auto& wgSize = ep.workgroup_size;
|
||||
SPIRVLocalSize spvLS;
|
||||
spvLS.width = max(wgSize.x, minDim);
|
||||
spvLS.height = max(wgSize.y, minDim);
|
||||
spvLS.depth = max(wgSize.z, minDim);
|
||||
localSizes[epOrigName] = spvLS;
|
||||
SPIRVEntryPoint mvkEP;
|
||||
mvkEP.mtlFunctionName = spvEP.name;
|
||||
mvkEP.workgroupSize.width = max(wgSize.x, minDim);
|
||||
mvkEP.workgroupSize.height = max(wgSize.y, minDim);
|
||||
mvkEP.workgroupSize.depth = max(wgSize.z, minDim);
|
||||
entryPoints[epOrigName] = mvkEP;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@ -129,15 +129,22 @@ namespace mvk {
|
||||
|
||||
} SPIRVToMSLConverterContext;
|
||||
|
||||
/** Specifies the SPIRV LocalSize, which is the number of threads in a compute shader workgroup. */
|
||||
/**
|
||||
* Describes a SPIRV entry point, including the Metal function name (which may be
|
||||
* different than the Vulkan entry point name if the original name was illegal in Metal),
|
||||
* and the number of threads in each workgroup, if the shader is a compute shader.
|
||||
*/
|
||||
typedef struct {
|
||||
uint32_t width = 1;
|
||||
uint32_t height = 1;
|
||||
uint32_t depth = 1;
|
||||
} SPIRVLocalSize;
|
||||
std::string mtlFunctionName;
|
||||
struct {
|
||||
uint32_t width = 1;
|
||||
uint32_t height = 1;
|
||||
uint32_t depth = 1;
|
||||
} workgroupSize;
|
||||
} SPIRVEntryPoint;
|
||||
|
||||
/** Holds a map of the LocalSize value for each compute function, indexed by SPIRV entry point name. */
|
||||
typedef std::unordered_map<std::string, SPIRVLocalSize> SPIRVLocalSizesByEntryPointName;
|
||||
/** Holds a map of entry point info, indexed by the SPIRV entry point name. */
|
||||
typedef std::unordered_map<std::string, SPIRVEntryPoint> SPIRVEntryPointsByName;
|
||||
|
||||
/** Special constant used in a MSLResourceBinding descriptorSet element to indicate the bindings for the push constants. */
|
||||
static const uint32_t kPushConstDescSet = std::numeric_limits<uint32_t>::max();
|
||||
@ -184,19 +191,8 @@ namespace mvk {
|
||||
*/
|
||||
const std::string& getMSL() { return _msl; }
|
||||
|
||||
/**
|
||||
* Returns a mapping between the original entry point name in the SPIR-V and a
|
||||
* possibly modified name as required to bypass restrictions in naming entry
|
||||
* points within MSL. Specifically, the entry point name "main" is illegal in MSL,
|
||||
* and is replaced by "main0" in this mapping.
|
||||
*/
|
||||
const std::unordered_map<std::string, std::string>& getEntryPointNameMap() { return _entryPointNameMap; }
|
||||
|
||||
/**
|
||||
* Returns a mapping of the local size of each entry point.
|
||||
* This is only meaningful for compute shaders.
|
||||
*/
|
||||
const SPIRVLocalSizesByEntryPointName& getLocalSizes() { return _localSizes; }
|
||||
/** Returns a mapping of entry point info, indexed by SPIR-V entry point name. */
|
||||
const SPIRVEntryPointsByName& getEntryPoints() { return _entryPoints; }
|
||||
|
||||
/**
|
||||
* Returns whether the most recent conversion was successful.
|
||||
@ -212,12 +208,9 @@ namespace mvk {
|
||||
const std::string& getResultLog() { return _resultLog; }
|
||||
|
||||
/** Sets MSL source code. This can be used when MSL is supplied directly. */
|
||||
void setMSL(const std::string& msl,
|
||||
const std::unordered_map<std::string, std::string>& entryPointNameMap,
|
||||
const SPIRVLocalSizesByEntryPointName& localSizes) {
|
||||
void setMSL(const std::string& msl, const SPIRVEntryPointsByName& entryPoints) {
|
||||
_msl = msl;
|
||||
_entryPointNameMap = entryPointNameMap;
|
||||
_localSizes = localSizes;
|
||||
_entryPoints = entryPoints;
|
||||
}
|
||||
|
||||
protected:
|
||||
@ -230,8 +223,7 @@ namespace mvk {
|
||||
std::vector<uint32_t> _spirv;
|
||||
std::string _msl;
|
||||
std::string _resultLog;
|
||||
std::unordered_map<std::string, std::string> _entryPointNameMap;
|
||||
SPIRVLocalSizesByEntryPointName _localSizes;
|
||||
SPIRVEntryPointsByName _entryPoints;
|
||||
bool _wasConverted = false;
|
||||
};
|
||||
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@ -77,7 +77,7 @@
|
||||
isEnabled = "NO">
|
||||
</CommandLineArgument>
|
||||
<CommandLineArgument
|
||||
argument = "/Users/bill/Documents/Dev/iOSProjects/Molten/MoltenVK/External/SPIRV-Cross/shaders-msl/vert/forum_test.vert"
|
||||
argument = "/Users/bill/Documents/Dev/iOSProjects/Molten/MoltenVK-bh/External/SPIRV-Cross/shaders-msl/vert/dynamic.flatten.vert"
|
||||
isEnabled = "NO">
|
||||
</CommandLineArgument>
|
||||
<CommandLineArgument
|
||||
@ -85,7 +85,7 @@
|
||||
isEnabled = "YES">
|
||||
</CommandLineArgument>
|
||||
<CommandLineArgument
|
||||
argument = "/Users/bill/Documents/Dev/iOSProjects/Molten/functions_nested-opt.spv"
|
||||
argument = "/Users/bill/Documents/Dev/iOSProjects/Molten/dynamic.flatten.vert.spv"
|
||||
isEnabled = "YES">
|
||||
</CommandLineArgument>
|
||||
<CommandLineArgument
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user