Merge pull request #6 from billhollings/master
Fixes to compute workgroup sizes and barriers.
This commit is contained in:
Bill Hollings
GitHub
commit
51193c7805
@ -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 {
|
||||
std::string mtlFunctionName;
|
||||
struct {
|
||||
uint32_t width = 1;
|
||||
uint32_t height = 1;
|
||||
uint32_t depth = 1;
|
||||
} SPIRVLocalSize;
|
||||
} 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