MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.cpp - external/github.com/KhronosGroup/MoltenVK - Git at Google

 /*
  * SPIRVToMSLConverter.cpp
  *
  * Copyright (c) 2014-2018 The Brenwill Workshop Ltd. (http://www.brenwill.com)
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
  * You may obtain a copy of the License at
  *
  *     http://www.apache.org/licenses/LICENSE-2.0
  *
  * Unless required by applicable law or agreed to in writing, software
  * distributed under the License is distributed on an "AS IS" BASIS,
  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  * See the License for the specific language governing permissions and
  * limitations under the License.
  */

 #include "SPIRVToMSLConverter.h"
 #include "MVKCommonEnvironment.h"
 #include "MVKStrings.h"
 #include "FileSupport.h"
 #include "spirv_msl.hpp"
 #include <spirv-tools/libspirv.h>
 #import <CoreFoundation/CFByteOrder.h>

 using namespace mvk;
 using namespace std;


 #pragma mark -
 #pragma mark SPIRVToMSLConverterContext

 // Returns whether the vector contains the value (using a matches(T&) comparison member function). */
 template<class T>
 bool contains(const vector<T>& vec, const T& val) {
     for (const T& vecVal : vec) { if (vecVal.matches(val)) { return true; } }
     return false;
 }

 MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverterOptions::matches(const SPIRVToMSLConverterOptions& other) const {
 	if (entryPointStage != other.entryPointStage) { return false; }
     if (mslVersion != other.mslVersion) { return false; }
 	if (texelBufferTextureWidth != other.texelBufferTextureWidth) { return false; }
 	if (auxBufferIndex != other.auxBufferIndex) { return false; }
     if (!!shouldFlipVertexY != !!other.shouldFlipVertexY) { return false; }
     if (!!isRenderingPoints != !!other.isRenderingPoints) { return false; }
 	if (entryPointName != other.entryPointName) { return false; }
     return true;
 }

 MVK_PUBLIC_SYMBOL bool MSLVertexAttribute::matches(const MSLVertexAttribute& other) const {
     if (location != other.location) { return false; }
     if (mslBuffer != other.mslBuffer) { return false; }
     if (mslOffset != other.mslOffset) { return false; }
     if (mslStride != other.mslStride) { return false; }
     if (!!isPerInstance != !!other.isPerInstance) { return false; }
     return true;
 }

 MVK_PUBLIC_SYMBOL bool MSLResourceBinding::matches(const MSLResourceBinding& other) const {
     if (stage != other.stage) { return false; }
     if (descriptorSet != other.descriptorSet) { return false; }
     if (binding != other.binding) { return false; }
     if (mslBuffer != other.mslBuffer) { return false; }
     if (mslTexture != other.mslTexture) { return false; }
     if (mslSampler != other.mslSampler) { return false; }
     return true;
 }

 // Check them all in case inactive VA's duplicate locations used by active VA's.
 MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverterContext::isVertexAttributeLocationUsed(uint32_t location) const {
     for (auto& va : vertexAttributes) {
         if ((va.location == location) && va.isUsedByShader) { return true; }
     }
     return false;
 }

 // Check them all in case inactive VA's duplicate buffers used by active VA's.
 MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverterContext::isVertexBufferUsed(uint32_t mslBuffer) const {
     for (auto& va : vertexAttributes) {
         if ((va.mslBuffer == mslBuffer) && va.isUsedByShader) { return true; }
     }
     return false;
 }

 MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverterContext::matches(const SPIRVToMSLConverterContext& other) const {

     if ( !options.matches(other.options) ) { return false; }

 	if (options.entryPointStage == spv::ExecutionModelVertex) {
 		for (const auto& va : vertexAttributes) {
 			if (va.isUsedByShader && !contains(other.vertexAttributes, va)) { return false; }
 		}
 	}

     for (const auto& rb : resourceBindings) {
         if (rb.isUsedByShader && !contains(other.resourceBindings, rb)) { return false; }
     }

     return true;
 }

 MVK_PUBLIC_SYMBOL void SPIRVToMSLConverterContext::alignWith(const SPIRVToMSLConverterContext& srcContext) {

 	options.isRasterizationDisabled = srcContext.options.isRasterizationDisabled;
 	options.needsAuxBuffer = srcContext.options.needsAuxBuffer;

 	if (options.entryPointStage == spv::ExecutionModelVertex) {
 		for (auto& va : vertexAttributes) {
 			va.isUsedByShader = false;
 			for (auto& srcVA : srcContext.vertexAttributes) {
 				if (va.matches(srcVA)) { va.isUsedByShader = srcVA.isUsedByShader; }
 			}
 		}
 	}

     for (auto& rb : resourceBindings) {
         rb.isUsedByShader = false;
         for (auto& srcRB : srcContext.resourceBindings) {
             if (rb.matches(srcRB)) { rb.isUsedByShader = srcRB.isUsedByShader; }
         }
     }
 }


 #pragma mark -
 #pragma mark SPIRVToMSLConverter

 // Populates the entry point with info extracted from the SPRI-V compiler.
 void populateEntryPoint(SPIRVEntryPoint& entryPoint, spirv_cross::Compiler* pCompiler, SPIRVToMSLConverterOptions& options);

 MVK_PUBLIC_SYMBOL void SPIRVToMSLConverter::setSPIRV(const vector<uint32_t>& spirv) { _spirv = spirv; }

 MVK_PUBLIC_SYMBOL void SPIRVToMSLConverter::setSPIRV(const uint32_t* spirvCode, size_t length) {
 	_spirv.clear();			// Clear for reuse
 	_spirv.reserve(length);
 	for (size_t i = 0; i < length; i++) {
 		_spirv.push_back(spirvCode[i]);
 	}
 }

 MVK_PUBLIC_SYMBOL const vector<uint32_t>& SPIRVToMSLConverter::getSPIRV() { return _spirv; }

 MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverter::convert(SPIRVToMSLConverterContext& context,
 													bool shouldLogSPIRV,
 													bool shouldLogMSL,
                                                     bool shouldLogGLSL) {
 	_wasConverted = true;
 	_resultLog.clear();
 	_msl.clear();

 	if (shouldLogSPIRV) { logSPIRV("Converting"); }

 	// Add vertex attributes
 	vector<spirv_cross::MSLVertexAttr> vtxAttrs;
 	spirv_cross::MSLVertexAttr va;
 	for (auto& ctxVA : context.vertexAttributes) {
 		va.location = ctxVA.location;
         va.msl_buffer = ctxVA.mslBuffer;
         va.msl_offset = ctxVA.mslOffset;
         va.msl_stride = ctxVA.mslStride;
         va.per_instance = ctxVA.isPerInstance;
         va.used_by_shader = ctxVA.isUsedByShader;
 		vtxAttrs.push_back(va);
 	}

 	// Add resource bindings
 	vector<spirv_cross::MSLResourceBinding> resBindings;
 	spirv_cross::MSLResourceBinding rb;
 	for (auto& ctxRB : context.resourceBindings) {
 		rb.desc_set = ctxRB.descriptorSet;
 		rb.binding = ctxRB.binding;
 		rb.stage = ctxRB.stage;
 		rb.msl_buffer = ctxRB.mslBuffer;
 		rb.msl_texture = ctxRB.mslTexture;
 		rb.msl_sampler = ctxRB.mslSampler;
         rb.used_by_shader = ctxRB.isUsedByShader;
 		resBindings.push_back(rb);
 	}

 	spirv_cross::CompilerMSL* pMSLCompiler = nullptr;

 #ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
 	try {
 #endif
 		pMSLCompiler = new spirv_cross::CompilerMSL(_spirv);

 		if (context.options.hasEntryPoint()) {
 			pMSLCompiler->set_entry_point(context.options.entryPointName, context.options.entryPointStage);
 		}

 		// Establish the MSL options for the compiler
 		// This needs to be done in two steps...for CompilerMSL and its superclass.
 		auto mslOpts = pMSLCompiler->get_msl_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.texel_buffer_texture_width = context.options.texelBufferTextureWidth;
 		mslOpts.aux_buffer_index = context.options.auxBufferIndex;
 		mslOpts.enable_point_size_builtin = context.options.isRenderingPoints;
 		mslOpts.disable_rasterization = context.options.isRasterizationDisabled;
 		mslOpts.swizzle_texture_samples = true;
 		pMSLCompiler->set_msl_options(mslOpts);

 		auto scOpts = pMSLCompiler->get_common_options();
 		scOpts.vertex.flip_vert_y = context.options.shouldFlipVertexY;
 		pMSLCompiler->set_common_options(scOpts);

 		_msl = pMSLCompiler->compile(&vtxAttrs, &resBindings);

         if (shouldLogMSL) { logSource(_msl, "MSL", "Converted"); }

 #ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
 	} catch (spirv_cross::CompilerError& ex) {
 		string errMsg("MSL conversion error: ");
 		errMsg += ex.what();
 		logError(errMsg.data());
         if (shouldLogMSL && pMSLCompiler) {
             _msl = pMSLCompiler->get_partial_source();
             logSource(_msl, "MSL", "Partially converted");
         }
 	}
 #endif

     // Populate content extracted from the SPRI-V compiler.
 	populateEntryPoint(_entryPoint, pMSLCompiler, context.options);
 	context.options.isRasterizationDisabled = pMSLCompiler && pMSLCompiler->get_is_rasterization_disabled();
 	context.options.needsAuxBuffer = pMSLCompiler && pMSLCompiler->needs_aux_buffer();
 	delete pMSLCompiler;

 	// Copy whether the vertex attributes and resource bindings are used by the shader
 	uint32_t vaCnt = (uint32_t)vtxAttrs.size();
 	for (uint32_t vaIdx = 0; vaIdx < vaCnt; vaIdx++) {
 		context.vertexAttributes[vaIdx].isUsedByShader = vtxAttrs[vaIdx].used_by_shader;
 	}
 	uint32_t rbCnt = (uint32_t)resBindings.size();
 	for (uint32_t rbIdx = 0; rbIdx < rbCnt; rbIdx++) {
 		context.resourceBindings[rbIdx].isUsedByShader = resBindings[rbIdx].used_by_shader;
 	}

     // To check GLSL conversion
     if (shouldLogGLSL) {
 		spirv_cross::CompilerGLSL* pGLSLCompiler = nullptr;

 #ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
 		try {
 #endif
 			pGLSLCompiler = new spirv_cross::CompilerGLSL(_spirv);
 			string glsl = pGLSLCompiler->compile();
             logSource(glsl, "GLSL", "Estimated original");
 #ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
         } catch (spirv_cross::CompilerError& ex) {
             string errMsg("Original GLSL extraction error: ");
             errMsg += ex.what();
             logMsg(errMsg.data());
 			if (pGLSLCompiler) {
 				string glsl = pGLSLCompiler->get_partial_source();
 				logSource(glsl, "GLSL", "Partially converted");
 			}
         }
 #endif
 		delete pGLSLCompiler;
 	}

 	return _wasConverted;
 }

 /** Appends the message text to the result log. */
 void SPIRVToMSLConverter::logMsg(const char* logMsg) {
 	string trimMsg = trim(logMsg);
 	if ( !trimMsg.empty() ) {
 		_resultLog += trimMsg;
 		_resultLog += "\n\n";
 	}
 }

 /** Appends the error text to the result log, sets the wasConverted property to false, and returns it. */
 bool SPIRVToMSLConverter::logError(const char* errMsg) {
 	logMsg(errMsg);
 	_wasConverted = false;
 	return _wasConverted;
 }

 /** Appends the SPIR-V to the result log, indicating whether it is being converted or was converted. */
 void SPIRVToMSLConverter::logSPIRV(const char* opDesc) {

 	string spvLog;
 	mvk::logSPIRV(_spirv, spvLog);

 	_resultLog += opDesc;
 	_resultLog += " SPIR-V:\n";
 	_resultLog += spvLog;
 	_resultLog += "\nEnd SPIR-V\n\n";

 	// Uncomment one or both of the following lines to get additional debugging and tracability capabilities.
 	// The SPIR-V can be written in binary form to a file, and/or logged in human readable form to the console.
 	// These can be helpful if errors occur during conversion of SPIR-V to MSL.
 //	writeSPIRVToFile("spvout.spv");
 //	printf("\n%s\n", getResultLog().c_str());
 }

 /**
  * Writes the SPIR-V code to a file. This can be useful for debugging
  * when the SPRIR-V did not originally come from a known file
  */
 void SPIRVToMSLConverter::writeSPIRVToFile(string spvFilepath) {
 	vector<char> fileContents;
 	spirvToBytes(_spirv, fileContents);
 	string errMsg;
 	if (writeFile(spvFilepath, fileContents, errMsg)) {
 		_resultLog += "Saved SPIR-V to file: " + absolutePath(spvFilepath) + "\n\n";
 	} else {
 		_resultLog += "Could not write SPIR-V file. " + errMsg + "\n\n";
 	}
 }

 /** Validates that the SPIR-V code will disassemble during logging. */
 bool SPIRVToMSLConverter::validateSPIRV() {
 	if (_spirv.size() < 5) { return false; }
 	if (_spirv[0] != spv::MagicNumber) { return false; }
 	if (_spirv[4] != 0) { return false; }
 	return true;
 }

 /** Appends the source to the result log, prepending with the operation. */
 void SPIRVToMSLConverter::logSource(string& src, const char* srcLang, const char* opDesc) {
     _resultLog += opDesc;
     _resultLog += " ";
     _resultLog += srcLang;
     _resultLog += ":\n";
     _resultLog += src;
     _resultLog += "\nEnd ";
     _resultLog += srcLang;
     _resultLog += "\n\n";
 }


 #pragma mark Support functions

 // Populate a workgroup size dimension.
 void populateWorkgroupDimension(SPIRVWorkgroupSizeDimension& wgDim, uint32_t size, spirv_cross::SpecializationConstant& spvSpecConst) {
 	wgDim.size = max(size, 1u);
 	wgDim.isSpecialized = (spvSpecConst.id != 0);
 	wgDim.specializationID = spvSpecConst.constant_id;
 }

 void populateEntryPoint(SPIRVEntryPoint& entryPoint, spirv_cross::Compiler* pCompiler, SPIRVToMSLConverterOptions& options) {

 	if ( !pCompiler ) { return; }

 	spirv_cross::SPIREntryPoint spvEP;
 	if (options.hasEntryPoint()) {
 		spvEP = pCompiler->get_entry_point(options.entryPointName, options.entryPointStage);
 	} else {
 		const auto& entryPoints = pCompiler->get_entry_points_and_stages();
 		if ( !entryPoints.empty() ) {
 			auto& ep = entryPoints[0];
 			spvEP = pCompiler->get_entry_point(ep.name, ep.execution_model);
 		}
 	}

 	spirv_cross::SpecializationConstant widthSC, heightSC, depthSC;
 	pCompiler->get_work_group_size_specialization_constants(widthSC, heightSC, depthSC);

 	entryPoint.mtlFunctionName = spvEP.name;
 	populateWorkgroupDimension(entryPoint.workgroupSize.width, spvEP.workgroup_size.x, widthSC);
 	populateWorkgroupDimension(entryPoint.workgroupSize.height, spvEP.workgroup_size.y, heightSC);
 	populateWorkgroupDimension(entryPoint.workgroupSize.depth, spvEP.workgroup_size.z, depthSC);
 }

 MVK_PUBLIC_SYMBOL void mvk::logSPIRV(vector<uint32_t>& spirv, string& spvLog) {
 	if ( !((spirv.size() > 4) &&
 		   (spirv[0] == spv::MagicNumber) &&
 		   (spirv[4] == 0)) ) { return; }

 	uint32_t options = (SPV_BINARY_TO_TEXT_OPTION_INDENT);
 	spv_text text;
 	spv_diagnostic diagnostic = nullptr;
 	spv_context context = spvContextCreate(SPV_ENV_VULKAN_1_0);
 	spv_result_t error = spvBinaryToText(context, spirv.data(), spirv.size(), options, &text, &diagnostic);
 	spvContextDestroy(context);
 	if (error) {
 		spvDiagnosticPrint(diagnostic);
 		spvDiagnosticDestroy(diagnostic);
 		return;
 	}
 	spvLog.append(text->str, text->length);
 	spvTextDestroy(text);
 }

 MVK_PUBLIC_SYMBOL void mvk::spirvToBytes(const vector<uint32_t>& spv, vector<char>& bytes) {
 	// Assumes desired endianness.
 	size_t byteCnt = spv.size() * sizeof(uint32_t);
 	char* cBytes = (char*)spv.data();
 	bytes.clear();
 	bytes.insert(bytes.end(), cBytes, cBytes + byteCnt);
 }

 MVK_PUBLIC_SYMBOL void mvk::bytesToSPIRV(const vector<char>& bytes, vector<uint32_t>& spv) {
 	size_t spvCnt = bytes.size() / sizeof(uint32_t);
 	uint32_t* cSPV = (uint32_t*)bytes.data();
 	spv.clear();
 	spv.insert(spv.end(), cSPV, cSPV + spvCnt);
 	ensureSPIRVEndianness(spv);
 }

 MVK_PUBLIC_SYMBOL bool mvk::ensureSPIRVEndianness(vector<uint32_t>& spv) {
 	if (spv.empty()) { return false; }					// Nothing to convert

 	uint32_t magNum = spv.front();
 	if (magNum == spv::MagicNumber) { return false; }	// No need to convert

 	if (CFSwapInt32(magNum) == spv::MagicNumber) {		// Yep, it's SPIR-V, but wrong endianness
 		for (auto& elem : spv) { elem = CFSwapInt32(elem); }
 		return true;
 	}
 	return false;		// Not SPIR-V, so don't convert
 }
	/*
	* SPIRVToMSLConverter.cpp
	*
	* Copyright (c) 2014-2018 The Brenwill Workshop Ltd. (http://www.brenwill.com)
	*
	* Licensed under the Apache License, Version 2.0 (the "License");
	* you may not use this file except in compliance with the License.
	* You may obtain a copy of the License at
	*
	* http://www.apache.org/licenses/LICENSE-2.0
	*
	* Unless required by applicable law or agreed to in writing, software
	* distributed under the License is distributed on an "AS IS" BASIS,
	* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
	* See the License for the specific language governing permissions and
	* limitations under the License.
	*/

	#include "SPIRVToMSLConverter.h"
	#include "MVKCommonEnvironment.h"
	#include "MVKStrings.h"
	#include "FileSupport.h"
	#include "spirv_msl.hpp"
	#include <spirv-tools/libspirv.h>
	#import <CoreFoundation/CFByteOrder.h>

	using namespace mvk;
	using namespace std;


	#pragma mark -
	#pragma mark SPIRVToMSLConverterContext

	// Returns whether the vector contains the value (using a matches(T&) comparison member function). */
	template<class T>
	bool contains(const vector<T>& vec, const T& val) {
	for (const T& vecVal : vec) { if (vecVal.matches(val)) { return true; } }
	return false;
	}

	MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverterOptions::matches(const SPIRVToMSLConverterOptions& other) const {
	if (entryPointStage != other.entryPointStage) { return false; }
	if (mslVersion != other.mslVersion) { return false; }
	if (texelBufferTextureWidth != other.texelBufferTextureWidth) { return false; }
	if (auxBufferIndex != other.auxBufferIndex) { return false; }
	if (!!shouldFlipVertexY != !!other.shouldFlipVertexY) { return false; }
	if (!!isRenderingPoints != !!other.isRenderingPoints) { return false; }
	if (entryPointName != other.entryPointName) { return false; }
	return true;
	}

	MVK_PUBLIC_SYMBOL bool MSLVertexAttribute::matches(const MSLVertexAttribute& other) const {
	if (location != other.location) { return false; }
	if (mslBuffer != other.mslBuffer) { return false; }
	if (mslOffset != other.mslOffset) { return false; }
	if (mslStride != other.mslStride) { return false; }
	if (!!isPerInstance != !!other.isPerInstance) { return false; }
	return true;
	}

	MVK_PUBLIC_SYMBOL bool MSLResourceBinding::matches(const MSLResourceBinding& other) const {
	if (stage != other.stage) { return false; }
	if (descriptorSet != other.descriptorSet) { return false; }
	if (binding != other.binding) { return false; }
	if (mslBuffer != other.mslBuffer) { return false; }
	if (mslTexture != other.mslTexture) { return false; }
	if (mslSampler != other.mslSampler) { return false; }
	return true;
	}

	// Check them all in case inactive VA's duplicate locations used by active VA's.
	MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverterContext::isVertexAttributeLocationUsed(uint32_t location) const {
	for (auto& va : vertexAttributes) {
	if ((va.location == location) && va.isUsedByShader) { return true; }
	}
	return false;
	}

	// Check them all in case inactive VA's duplicate buffers used by active VA's.
	MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverterContext::isVertexBufferUsed(uint32_t mslBuffer) const {
	for (auto& va : vertexAttributes) {
	if ((va.mslBuffer == mslBuffer) && va.isUsedByShader) { return true; }
	}
	return false;
	}

	MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverterContext::matches(const SPIRVToMSLConverterContext& other) const {

	if ( !options.matches(other.options) ) { return false; }

	if (options.entryPointStage == spv::ExecutionModelVertex) {
	for (const auto& va : vertexAttributes) {
	if (va.isUsedByShader && !contains(other.vertexAttributes, va)) { return false; }
	}
	}

	for (const auto& rb : resourceBindings) {
	if (rb.isUsedByShader && !contains(other.resourceBindings, rb)) { return false; }
	}

	return true;
	}

	MVK_PUBLIC_SYMBOL void SPIRVToMSLConverterContext::alignWith(const SPIRVToMSLConverterContext& srcContext) {

	options.isRasterizationDisabled = srcContext.options.isRasterizationDisabled;
	options.needsAuxBuffer = srcContext.options.needsAuxBuffer;

	if (options.entryPointStage == spv::ExecutionModelVertex) {
	for (auto& va : vertexAttributes) {
	va.isUsedByShader = false;
	for (auto& srcVA : srcContext.vertexAttributes) {
	if (va.matches(srcVA)) { va.isUsedByShader = srcVA.isUsedByShader; }
	}
	}
	}

	for (auto& rb : resourceBindings) {
	rb.isUsedByShader = false;
	for (auto& srcRB : srcContext.resourceBindings) {
	if (rb.matches(srcRB)) { rb.isUsedByShader = srcRB.isUsedByShader; }
	}
	}
	}


	#pragma mark -
	#pragma mark SPIRVToMSLConverter

	// Populates the entry point with info extracted from the SPRI-V compiler.
	void populateEntryPoint(SPIRVEntryPoint& entryPoint, spirv_cross::Compiler* pCompiler, SPIRVToMSLConverterOptions& options);

	MVK_PUBLIC_SYMBOL void SPIRVToMSLConverter::setSPIRV(const vector<uint32_t>& spirv) { _spirv = spirv; }

	MVK_PUBLIC_SYMBOL void SPIRVToMSLConverter::setSPIRV(const uint32_t* spirvCode, size_t length) {
	_spirv.clear(); // Clear for reuse
	_spirv.reserve(length);
	for (size_t i = 0; i < length; i++) {
	_spirv.push_back(spirvCode[i]);
	}
	}

	MVK_PUBLIC_SYMBOL const vector<uint32_t>& SPIRVToMSLConverter::getSPIRV() { return _spirv; }

	MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverter::convert(SPIRVToMSLConverterContext& context,
	bool shouldLogSPIRV,
	bool shouldLogMSL,
	bool shouldLogGLSL) {
	_wasConverted = true;
	_resultLog.clear();
	_msl.clear();

	if (shouldLogSPIRV) { logSPIRV("Converting"); }

	// Add vertex attributes
	vector<spirv_cross::MSLVertexAttr> vtxAttrs;
	spirv_cross::MSLVertexAttr va;
	for (auto& ctxVA : context.vertexAttributes) {
	va.location = ctxVA.location;
	va.msl_buffer = ctxVA.mslBuffer;
	va.msl_offset = ctxVA.mslOffset;
	va.msl_stride = ctxVA.mslStride;
	va.per_instance = ctxVA.isPerInstance;
	va.used_by_shader = ctxVA.isUsedByShader;
	vtxAttrs.push_back(va);
	}

	// Add resource bindings
	vector<spirv_cross::MSLResourceBinding> resBindings;
	spirv_cross::MSLResourceBinding rb;
	for (auto& ctxRB : context.resourceBindings) {
	rb.desc_set = ctxRB.descriptorSet;
	rb.binding = ctxRB.binding;
	rb.stage = ctxRB.stage;
	rb.msl_buffer = ctxRB.mslBuffer;
	rb.msl_texture = ctxRB.mslTexture;
	rb.msl_sampler = ctxRB.mslSampler;
	rb.used_by_shader = ctxRB.isUsedByShader;
	resBindings.push_back(rb);
	}

	spirv_cross::CompilerMSL* pMSLCompiler = nullptr;

	#ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
	try {
	#endif
	pMSLCompiler = new spirv_cross::CompilerMSL(_spirv);

	if (context.options.hasEntryPoint()) {
	pMSLCompiler->set_entry_point(context.options.entryPointName, context.options.entryPointStage);
	}

	// Establish the MSL options for the compiler
	// This needs to be done in two steps...for CompilerMSL and its superclass.
	auto mslOpts = pMSLCompiler->get_msl_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.texel_buffer_texture_width = context.options.texelBufferTextureWidth;
	mslOpts.aux_buffer_index = context.options.auxBufferIndex;
	mslOpts.enable_point_size_builtin = context.options.isRenderingPoints;
	mslOpts.disable_rasterization = context.options.isRasterizationDisabled;
	mslOpts.swizzle_texture_samples = true;
	pMSLCompiler->set_msl_options(mslOpts);

	auto scOpts = pMSLCompiler->get_common_options();
	scOpts.vertex.flip_vert_y = context.options.shouldFlipVertexY;
	pMSLCompiler->set_common_options(scOpts);

	_msl = pMSLCompiler->compile(&vtxAttrs, &resBindings);

	if (shouldLogMSL) { logSource(_msl, "MSL", "Converted"); }

	#ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
	} catch (spirv_cross::CompilerError& ex) {
	string errMsg("MSL conversion error: ");
	errMsg += ex.what();
	logError(errMsg.data());
	if (shouldLogMSL && pMSLCompiler) {
	_msl = pMSLCompiler->get_partial_source();
	logSource(_msl, "MSL", "Partially converted");
	}
	}
	#endif

	// Populate content extracted from the SPRI-V compiler.
	populateEntryPoint(_entryPoint, pMSLCompiler, context.options);
	context.options.isRasterizationDisabled = pMSLCompiler && pMSLCompiler->get_is_rasterization_disabled();
	context.options.needsAuxBuffer = pMSLCompiler && pMSLCompiler->needs_aux_buffer();
	delete pMSLCompiler;

	// Copy whether the vertex attributes and resource bindings are used by the shader
	uint32_t vaCnt = (uint32_t)vtxAttrs.size();
	for (uint32_t vaIdx = 0; vaIdx < vaCnt; vaIdx++) {
	context.vertexAttributes[vaIdx].isUsedByShader = vtxAttrs[vaIdx].used_by_shader;
	}
	uint32_t rbCnt = (uint32_t)resBindings.size();
	for (uint32_t rbIdx = 0; rbIdx < rbCnt; rbIdx++) {
	context.resourceBindings[rbIdx].isUsedByShader = resBindings[rbIdx].used_by_shader;
	}

	// To check GLSL conversion
	if (shouldLogGLSL) {
	spirv_cross::CompilerGLSL* pGLSLCompiler = nullptr;

	#ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
	try {
	#endif
	pGLSLCompiler = new spirv_cross::CompilerGLSL(_spirv);
	string glsl = pGLSLCompiler->compile();
	logSource(glsl, "GLSL", "Estimated original");
	#ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
	} catch (spirv_cross::CompilerError& ex) {
	string errMsg("Original GLSL extraction error: ");
	errMsg += ex.what();
	logMsg(errMsg.data());
	if (pGLSLCompiler) {
	string glsl = pGLSLCompiler->get_partial_source();
	logSource(glsl, "GLSL", "Partially converted");
	}
	}
	#endif
	delete pGLSLCompiler;
	}

	return _wasConverted;
	}

	/** Appends the message text to the result log. */
	void SPIRVToMSLConverter::logMsg(const char* logMsg) {
	string trimMsg = trim(logMsg);
	if ( !trimMsg.empty() ) {
	_resultLog += trimMsg;
	_resultLog += "\n\n";
	}
	}

	/** Appends the error text to the result log, sets the wasConverted property to false, and returns it. */
	bool SPIRVToMSLConverter::logError(const char* errMsg) {
	logMsg(errMsg);
	_wasConverted = false;
	return _wasConverted;
	}

	/** Appends the SPIR-V to the result log, indicating whether it is being converted or was converted. */
	void SPIRVToMSLConverter::logSPIRV(const char* opDesc) {

	string spvLog;
	mvk::logSPIRV(_spirv, spvLog);

	_resultLog += opDesc;
	_resultLog += " SPIR-V:\n";
	_resultLog += spvLog;
	_resultLog += "\nEnd SPIR-V\n\n";

	// Uncomment one or both of the following lines to get additional debugging and tracability capabilities.
	// The SPIR-V can be written in binary form to a file, and/or logged in human readable form to the console.
	// These can be helpful if errors occur during conversion of SPIR-V to MSL.
	// writeSPIRVToFile("spvout.spv");
	// printf("\n%s\n", getResultLog().c_str());
	}

	/**
	* Writes the SPIR-V code to a file. This can be useful for debugging
	* when the SPRIR-V did not originally come from a known file
	*/
	void SPIRVToMSLConverter::writeSPIRVToFile(string spvFilepath) {
	vector<char> fileContents;
	spirvToBytes(_spirv, fileContents);
	string errMsg;
	if (writeFile(spvFilepath, fileContents, errMsg)) {
	_resultLog += "Saved SPIR-V to file: " + absolutePath(spvFilepath) + "\n\n";
	} else {
	_resultLog += "Could not write SPIR-V file. " + errMsg + "\n\n";
	}
	}

	/** Validates that the SPIR-V code will disassemble during logging. */
	bool SPIRVToMSLConverter::validateSPIRV() {
	if (_spirv.size() < 5) { return false; }
	if (_spirv[0] != spv::MagicNumber) { return false; }
	if (_spirv[4] != 0) { return false; }
	return true;
	}

	/** Appends the source to the result log, prepending with the operation. */
	void SPIRVToMSLConverter::logSource(string& src, const char* srcLang, const char* opDesc) {
	_resultLog += opDesc;
	_resultLog += " ";
	_resultLog += srcLang;
	_resultLog += ":\n";
	_resultLog += src;
	_resultLog += "\nEnd ";
	_resultLog += srcLang;
	_resultLog += "\n\n";
	}


	#pragma mark Support functions

	// Populate a workgroup size dimension.
	void populateWorkgroupDimension(SPIRVWorkgroupSizeDimension& wgDim, uint32_t size, spirv_cross::SpecializationConstant& spvSpecConst) {
	wgDim.size = max(size, 1u);
	wgDim.isSpecialized = (spvSpecConst.id != 0);
	wgDim.specializationID = spvSpecConst.constant_id;
	}

	void populateEntryPoint(SPIRVEntryPoint& entryPoint, spirv_cross::Compiler* pCompiler, SPIRVToMSLConverterOptions& options) {

	if ( !pCompiler ) { return; }

	spirv_cross::SPIREntryPoint spvEP;
	if (options.hasEntryPoint()) {
	spvEP = pCompiler->get_entry_point(options.entryPointName, options.entryPointStage);
	} else {
	const auto& entryPoints = pCompiler->get_entry_points_and_stages();
	if ( !entryPoints.empty() ) {
	auto& ep = entryPoints[0];
	spvEP = pCompiler->get_entry_point(ep.name, ep.execution_model);
	}
	}

	spirv_cross::SpecializationConstant widthSC, heightSC, depthSC;
	pCompiler->get_work_group_size_specialization_constants(widthSC, heightSC, depthSC);

	entryPoint.mtlFunctionName = spvEP.name;
	populateWorkgroupDimension(entryPoint.workgroupSize.width, spvEP.workgroup_size.x, widthSC);
	populateWorkgroupDimension(entryPoint.workgroupSize.height, spvEP.workgroup_size.y, heightSC);
	populateWorkgroupDimension(entryPoint.workgroupSize.depth, spvEP.workgroup_size.z, depthSC);
	}

	MVK_PUBLIC_SYMBOL void mvk::logSPIRV(vector<uint32_t>& spirv, string& spvLog) {
	if ( !((spirv.size() > 4) &&
	(spirv[0] == spv::MagicNumber) &&
	(spirv[4] == 0)) ) { return; }

	uint32_t options = (SPV_BINARY_TO_TEXT_OPTION_INDENT);
	spv_text text;
	spv_diagnostic diagnostic = nullptr;
	spv_context context = spvContextCreate(SPV_ENV_VULKAN_1_0);
	spv_result_t error = spvBinaryToText(context, spirv.data(), spirv.size(), options, &text, &diagnostic);
	spvContextDestroy(context);
	if (error) {
	spvDiagnosticPrint(diagnostic);
	spvDiagnosticDestroy(diagnostic);
	return;
	}
	spvLog.append(text->str, text->length);
	spvTextDestroy(text);
	}

	MVK_PUBLIC_SYMBOL void mvk::spirvToBytes(const vector<uint32_t>& spv, vector<char>& bytes) {
	// Assumes desired endianness.
	size_t byteCnt = spv.size() * sizeof(uint32_t);
	char* cBytes = (char*)spv.data();
	bytes.clear();
	bytes.insert(bytes.end(), cBytes, cBytes + byteCnt);
	}

	MVK_PUBLIC_SYMBOL void mvk::bytesToSPIRV(const vector<char>& bytes, vector<uint32_t>& spv) {
	size_t spvCnt = bytes.size() / sizeof(uint32_t);
	uint32_t* cSPV = (uint32_t*)bytes.data();
	spv.clear();
	spv.insert(spv.end(), cSPV, cSPV + spvCnt);
	ensureSPIRVEndianness(spv);
	}

	MVK_PUBLIC_SYMBOL bool mvk::ensureSPIRVEndianness(vector<uint32_t>& spv) {
	if (spv.empty()) { return false; } // Nothing to convert

	uint32_t magNum = spv.front();
	if (magNum == spv::MagicNumber) { return false; } // No need to convert

	if (CFSwapInt32(magNum) == spv::MagicNumber) { // Yep, it's SPIR-V, but wrong endianness
	for (auto& elem : spv) { elem = CFSwapInt32(elem); }
	return true;
	}
	return false; // Not SPIR-V, so don't convert
	}