Merge pull request #2342 from erinmaus/relax-shader-storage-requirements

Relax shader storage requirements

src/modules/graphics/Graphics.cpp

@@ -566,7 +566,7 @@ bool Graphics::validateShader(bool gles, const std::vector<std::string> &stagess
 		}
 	}
 
-	return Shader::validate(stages, err);
+	return Shader::validate(stages, err, options);
 }
 
 Texture *Graphics::getDefaultTexture(TextureType type, DataBaseType dataType, bool depthSample)
@@ -2988,6 +2988,8 @@ STRINGMAP_CLASS_BEGIN(Graphics, Graphics::Feature, Graphics::FEATURE_MAX_ENUM, f
 	{ "texelbuffer",              Graphics::FEATURE_TEXEL_BUFFER         },
 	{ "copytexturetobuffer",      Graphics::FEATURE_COPY_TEXTURE_TO_BUFFER },
 	{ "indirectdraw",             Graphics::FEATURE_INDIRECT_DRAW        },
+	{ "vertexwrite",              Graphics::FEATURE_VERTEX_WRITE         },
+	{ "pixelwrite",               Graphics::FEATURE_PIXEL_WRITE          },
 }
 STRINGMAP_CLASS_END(Graphics, Graphics::Feature, Graphics::FEATURE_MAX_ENUM, feature)
 

src/modules/graphics/Graphics.h

@@ -163,6 +163,8 @@ public:
 		FEATURE_TEXEL_BUFFER,
 		FEATURE_COPY_TEXTURE_TO_BUFFER,
 		FEATURE_INDIRECT_DRAW,
+		FEATURE_VERTEX_WRITE,
+		FEATURE_PIXEL_WRITE,
 		FEATURE_MAX_ENUM
 	};
 

src/modules/graphics/Shader.cpp

@@ -650,7 +650,7 @@ Shader::Shader(StrongRef<ShaderStage> _stages[], const CompileOptions &options)
 	, debugName(options.debugName)
 {
 	std::string err;
-	if (!validateInternal(_stages, err, reflection))
+	if (!validateInternal(_stages, err, reflection, options))
 		throw love::Exception("%s", err.c_str());
 
 	std::vector<std::string> unsetVertexInputLocations;
@@ -1029,10 +1029,10 @@ bool Shader::isUsingDeprecatedTextureUniform() const
 	return it != reflection.allUniforms.end() && it->second->stageMask != 0;
 }
 
-bool Shader::validate(StrongRef<ShaderStage> stages[], std::string& err)
+bool Shader::validate(StrongRef<ShaderStage> stages[], std::string& err, const CompileOptions &options)
 {
 	Reflection reflection;
-	return validateInternal(stages, err, reflection);
+	return validateInternal(stages, err, reflection, options);
 }
 
 static DataBaseType getBaseType(glslang::TBasicType basictype)
@@ -1246,7 +1246,7 @@ static bool AddFieldsToFormat(std::vector<Buffer::DataDeclaration> &format, int
 	return true;
 }
 
-bool Shader::validateInternal(StrongRef<ShaderStage> stages[], std::string &err, Reflection &reflection)
+bool Shader::validateInternal(StrongRef<ShaderStage> stages[], std::string &err, Reflection &reflection, const CompileOptions &options)
 {
 	glslang::TProgram program;
 
@@ -1307,6 +1307,7 @@ bool Shader::validateInternal(StrongRef<ShaderStage> stages[], std::string &err,
 	reflection.textureCount = 0;
 	reflection.bufferCount = 0;
 
+	auto &capabilities = Module::getInstance<Graphics>(Module::M_GRAPHICS)->getCapabilities();
 	for (int i = 0; i < program.getNumUniformVariables(); i++)
 	{
 		const glslang::TObjectReflection &info = program.getUniform(i);
@@ -1349,9 +1350,21 @@ bool Shader::validateInternal(StrongRef<ShaderStage> stages[], std::string &err,
 		}
 		else if (type->isImage())
 		{
-			if ((info.stages & (~EShLangComputeMask)) != 0)
+			if ((info.stages & (~EShLangComputeMask)) != 0 && !options.features[FEATURE_WRITE])
 			{
-				err = "Shader validation error:\nStorage Texture uniform variables (image2D, etc) are only allowed in compute shaders.";
+				err = "Shader validation error:\nStorage Texture uniform variables (image2D, etc) are only allowed in compute shaders unless explicitly enabled.";
+				return false;
+			}
+
+			if ((!qualifiers.isReadOnly() || qualifiers.isWriteOnly()) && ((info.stages & (~EShLangFragmentMask)) != 0) && !capabilities.features[Graphics::FEATURE_PIXEL_WRITE])
+			{
+				err = "Shader validation error:\nPlatform does not have writable Storage Texture uniform variables (image2D, etc) capabilities in pixel shaders.";
+				return false;
+			}
+
+			if ((!qualifiers.isReadOnly() || qualifiers.isWriteOnly()) && ((info.stages & (~EShLangVertexMask)) != 0) && !capabilities.features[Graphics::FEATURE_VERTEX_WRITE])
+			{
+				err = "Shader validation error:\nPlatform does not have writable Storage Texture uniform variables (image2D, etc) capabilities in vertex shaders.";
 				return false;
 			}
 
@@ -1467,9 +1480,21 @@ bool Shader::validateInternal(StrongRef<ShaderStage> stages[], std::string &err,
 		{
 			const glslang::TQualifier &qualifiers = type->getQualifier();
 
-			if ((!qualifiers.isReadOnly() || qualifiers.isWriteOnly()) && ((info.stages & (~EShLangComputeMask)) != 0))
+			if ((!qualifiers.isReadOnly() || qualifiers.isWriteOnly()) && ((info.stages & (~EShLangComputeMask)) != 0) && !options.features[FEATURE_WRITE])
 			{
-				err = "Shader validation error:\nStorage Buffer block '" + info.name + "' must be marked as readonly in vertex and pixel shaders.";
+				err = "Shader validation error:\nStorage Buffer block '" + info.name + "' must be marked as readonly in vertex and pixel shaders unless explicitly enabled.";
+				return false;
+			}
+			
+			if ((!qualifiers.isReadOnly() || qualifiers.isWriteOnly()) && ((info.stages & (~EShLangFragmentMask)) != 0) && !capabilities.features[Graphics::FEATURE_PIXEL_WRITE])
+			{
+				err = "Shader validation error:\nPlatform does not have writable Storage Buffer blocks capabilities in pixel shaders.";
+				return false;
+			}
+
+			if ((!qualifiers.isReadOnly() || qualifiers.isWriteOnly()) && ((info.stages & (~EShLangVertexMask)) != 0) && !capabilities.features[Graphics::FEATURE_VERTEX_WRITE])
+			{
+				err = "Shader validation error:\nPlatform does not have writable Storage Buffer blocks capabilities in vertex shaders.";
 				return false;
 			}
 
@@ -1872,5 +1897,11 @@ bool Shader::getConstant(BuiltinUniform in, const char *&out)
 	return builtinNames.find(in, out);
 }
 
+STRINGMAP_CLASS_BEGIN(Shader, Shader::Feature, Shader::FEATURE_MAX_ENUM, feature)
+{
+	{ "write", Shader::FEATURE_WRITE },
+}
+STRINGMAP_CLASS_END(Shader, Shader::Feature, Shader::FEATURE_MAX_ENUM, feature)
+
 } // graphics
 } // love

src/modules/graphics/Shader.h

@@ -56,6 +56,12 @@ public:
 		LANGUAGE_MAX_ENUM
 	};
 
+	enum Feature
+	{
+		FEATURE_WRITE,
+		FEATURE_MAX_ENUM
+	};
+
 	// Built-in uniform variables.
 	enum BuiltinUniform
 	{
@@ -119,6 +125,7 @@ public:
 	{
 		std::map<std::string, std::string> defines;
 		std::string debugName;
+		bool features[FEATURE_MAX_ENUM] = {};
 	};
 
 	struct SourceInfo
@@ -275,7 +282,7 @@ public:
 	static SourceInfo getSourceInfo(const std::string &src);
 	static std::string createShaderStageCode(Graphics *gfx, ShaderStageType stage, const std::string &code, const CompileOptions &options, const SourceInfo &info, bool gles, bool checksystemfeatures);
 
-	static bool validate(StrongRef<ShaderStage> stages[], std::string &err);
+	static bool validate(StrongRef<ShaderStage> stages[], std::string &err, const CompileOptions &options);
 
 	static bool initialize();
 	static void deinitialize();
@@ -288,6 +295,8 @@ public:
 	static bool getConstant(const char *in, BuiltinUniform &out);
 	static bool getConstant(BuiltinUniform in, const char *&out);
 
+	STRINGMAP_CLASS_DECLARE(Feature);
+
 protected:
 
 	struct Reflection
@@ -330,7 +339,7 @@ protected:
 
 	static std::string canonicaliizeUniformName(const std::string &name);
 	static size_t getUniformDataSizePacked(const UniformInfo &u);
-	static bool validateInternal(StrongRef<ShaderStage> stages[], std::string& err, Reflection &reflection);
+	static bool validateInternal(StrongRef<ShaderStage> stages[], std::string& err, Reflection &reflection, const CompileOptions &options);
 	static DataBaseType getDataBaseType(PixelFormat format);
 	static bool isResourceBaseTypeCompatible(DataBaseType a, DataBaseType b);
 

src/modules/graphics/metal/Graphics.h

@@ -210,7 +210,7 @@ private:
 	id<MTLDepthStencilState> getCachedDepthStencilState(const DepthState &depth, const StencilState &stencil);
 	void applyRenderState(id<MTLRenderCommandEncoder> renderEncoder, VertexAttributesID attributesID);
 	bool applyShaderUniforms(id<MTLComputeCommandEncoder> encoder, love::graphics::Shader *shader);
-	void applyShaderUniforms(id<MTLRenderCommandEncoder> renderEncoder, love::graphics::Shader *shader, Texture *maintex);
+	bool applyShaderUniforms(id<MTLRenderCommandEncoder> renderEncoder, love::graphics::Shader *shader, Texture *maintex);
 
 	id<MTLCommandQueue> commandQueue;
 

src/modules/graphics/metal/Graphics.mm

@@ -1105,7 +1105,7 @@ bool Graphics::applyShaderUniforms(id<MTLComputeCommandEncoder> encoder, love::g
 	return allWritableVariablesSet;
 }
 
-void Graphics::applyShaderUniforms(id<MTLRenderCommandEncoder> renderEncoder, love::graphics::Shader *shader, love::graphics::Texture *maintex)
+bool Graphics::applyShaderUniforms(id<MTLRenderCommandEncoder> renderEncoder, love::graphics::Shader *shader, love::graphics::Texture *maintex)
 {
 	Shader *s = (Shader *)shader;
 
@@ -1168,6 +1168,8 @@ void Graphics::applyShaderUniforms(id<MTLRenderCommandEncoder> renderEncoder, lo
 
 	uniformBufferOffset += alignUp(size, alignment);
 
+	bool allWritableVariablesSet = true;
+
 	for (const Shader::TextureBinding &b : s->getTextureBindings())
 	{
 		id<MTLTexture> texture = b.texture;
@@ -1183,7 +1185,13 @@ void Graphics::applyShaderUniforms(id<MTLRenderCommandEncoder> renderEncoder, lo
 		uint8 sampindex = b.samplerStages[SHADERSTAGE_VERTEX];
 
 		if (texindex != LOVE_UINT8_MAX)
+		{
 			setTexture(renderEncoder, bindings, SHADERSTAGE_VERTEX, texindex, texture);
+
+			if ((b.access & Shader::ACCESS_WRITE) != 0 && texture == nil)
+				allWritableVariablesSet = false;
+		}
+
 		if (sampindex != LOVE_UINT8_MAX)
 			setSampler(renderEncoder, bindings, SHADERSTAGE_VERTEX, sampindex, samplertex);
 
@@ -1191,7 +1199,13 @@ void Graphics::applyShaderUniforms(id<MTLRenderCommandEncoder> renderEncoder, lo
 		sampindex = b.samplerStages[SHADERSTAGE_PIXEL];
 
 		if (texindex != LOVE_UINT8_MAX)
+		{
 			setTexture(renderEncoder, bindings, SHADERSTAGE_PIXEL, texindex, texture);
+
+			if ((b.access & Shader::ACCESS_WRITE) != 0 && texture == nil)
+				allWritableVariablesSet = false;
+		}
+
 		if (sampindex != LOVE_UINT8_MAX)
 			setSampler(renderEncoder, bindings, SHADERSTAGE_PIXEL, sampindex, samplertex);
 	}
@@ -1200,11 +1214,24 @@ void Graphics::applyShaderUniforms(id<MTLRenderCommandEncoder> renderEncoder, lo
 	{
 		uint8 index = b.stages[SHADERSTAGE_VERTEX];
 		if (index != LOVE_UINT8_MAX)
+		{
 			setBuffer(renderEncoder, bindings, SHADERSTAGE_VERTEX, index, b.buffer, 0);
+
+			if ((b.access & Shader::ACCESS_WRITE) != 0 && b.buffer == nil)
+				allWritableVariablesSet = false;
+		}
+
 		index = b.stages[SHADERSTAGE_PIXEL];
 		if (index != LOVE_UINT8_MAX)
+		{
 			setBuffer(renderEncoder, bindings, SHADERSTAGE_PIXEL, index, b.buffer, 0);
+
+			if ((b.access & Shader::ACCESS_WRITE) != 0 && b.buffer == nil)
+				allWritableVariablesSet = false;
+		}
 	}
+
+	return allWritableVariablesSet;
 }
 
 static void setVertexBuffers(id<MTLRenderCommandEncoder> encoder, love::graphics::Shader *shader, const BufferBindings *buffers, Graphics::RenderEncoderBindings &bindings)
@@ -1240,7 +1267,8 @@ void Graphics::draw(const DrawCommand &cmd)
 	}
 
 	applyRenderState(encoder, cmd.attributesID);
-	applyShaderUniforms(encoder, Shader::current, cmd.texture);
+	if (!applyShaderUniforms(encoder, Shader::current, cmd.texture))
+		return;
 
 	setVertexBuffers(encoder, Shader::current, cmd.buffers, renderBindings);
 
@@ -1272,7 +1300,8 @@ void Graphics::draw(const DrawIndexedCommand &cmd)
 	}
 
 	applyRenderState(encoder, cmd.attributesID);
-	applyShaderUniforms(encoder, Shader::current, cmd.texture);
+	if (!applyShaderUniforms(encoder, Shader::current, cmd.texture))
+		return;
 
 	setVertexBuffers(encoder, Shader::current, cmd.buffers, renderBindings);
 
@@ -1337,7 +1366,8 @@ void Graphics::drawQuads(int start, int count, VertexAttributesID attributesID,
 	}
 
 	applyRenderState(encoder, attributesID);
-	applyShaderUniforms(encoder, Shader::current, texture);
+	if (!applyShaderUniforms(encoder, Shader::current, texture))
+		return;
 
 	id<MTLBuffer> ib = getMTLBuffer(quadIndexBuffer);
 
@@ -2255,8 +2285,21 @@ void Graphics::initCapabilities()
 		capabilities.features[FEATURE_INDIRECT_DRAW] = true;
 	else
 		capabilities.features[FEATURE_INDIRECT_DRAW] = false;
+
+	// Apple 3 devices support read/write to buffers in functions, while Apple 4 supports read/write to images.
+	// So let's err on the safe side and check support for Apple 4.
+	if (families.apple[4])
+	{
+		capabilities.features[FEATURE_VERTEX_WRITE] = true;
+		capabilities.features[FEATURE_PIXEL_WRITE] = true;
+	}
+	else
+	{
+		capabilities.features[FEATURE_VERTEX_WRITE] = false;
+		capabilities.features[FEATURE_PIXEL_WRITE] = false;
+	}
 	
-	static_assert(FEATURE_MAX_ENUM == 13, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
+	static_assert(FEATURE_MAX_ENUM == 15, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
 
 	// https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf
 	capabilities.limits[LIMIT_POINT_SIZE] = 511;

src/modules/graphics/opengl/Graphics.cpp

@@ -444,7 +444,7 @@ void Graphics::setActive(bool enable)
 	active = enable;
 }
 
-static bool computeDispatchBarriers(Shader *shader, GLbitfield &preDispatchBarriers, GLbitfield &postDispatchBarriers)
+static bool shaderBarriers(Shader *shader, GLbitfield &preDispatchBarriers, GLbitfield &postDispatchBarriers)
 {
 	for (auto buffer : shader->getActiveWritableStorageBuffers())
 	{
@@ -505,7 +505,7 @@ bool Graphics::dispatch(love::graphics::Shader *s, int x, int y, int z)
 	GLbitfield preDispatchBarriers = 0;
 	GLbitfield postDispatchBarriers = 0;
 
-	if (!computeDispatchBarriers(shader, preDispatchBarriers, postDispatchBarriers))
+	if (!shaderBarriers(shader, preDispatchBarriers, postDispatchBarriers))
 		return false;
 
 	// glMemoryBarrier before dispatch to make sure non-compute-read ->
@@ -534,7 +534,7 @@ bool Graphics::dispatch(love::graphics::Shader *s, love::graphics::Buffer *indir
 	GLbitfield preDispatchBarriers = 0;
 	GLbitfield postDispatchBarriers = 0;
 
-	if (!computeDispatchBarriers(shader, preDispatchBarriers, postDispatchBarriers))
+	if (!shaderBarriers(shader, preDispatchBarriers, postDispatchBarriers))
 		return false;
 
 	if (preDispatchBarriers != 0)
@@ -559,6 +559,15 @@ void Graphics::draw(const DrawCommand &cmd)
 	VertexAttributes attributes;
 	findVertexAttributes(cmd.attributesID, attributes);
 
+	GLbitfield preDrawBarriers = 0;
+	GLbitfield postDrawBarriers = 0;
+
+	if (!shaderBarriers((Shader *)Shader::current, preDrawBarriers, postDrawBarriers))
+		return;
+
+	if (preDrawBarriers != 0)
+		glMemoryBarrier(preDrawBarriers);
+
 	gl.prepareDraw(this);
 	gl.setVertexAttributes(attributes, *cmd.buffers);
 	gl.bindTextureToUnit(cmd.texture, 0, false);
@@ -576,6 +585,9 @@ void Graphics::draw(const DrawCommand &cmd)
 	else
 		glDrawArrays(glprimitivetype, cmd.vertexStart, cmd.vertexCount);
 
+	if (postDrawBarriers != 0)
+		glMemoryBarrier(postDrawBarriers);
+
 	++drawCalls;
 }
 
@@ -584,6 +596,15 @@ void Graphics::draw(const DrawIndexedCommand &cmd)
 	VertexAttributes attributes;
 	findVertexAttributes(cmd.attributesID, attributes);
 
+	GLbitfield preDrawBarriers = 0;
+	GLbitfield postDrawBarriers = 0;
+
+	if (!shaderBarriers((Shader *)Shader::current, preDrawBarriers, postDrawBarriers))
+		return;
+
+	if (preDrawBarriers != 0)
+		glMemoryBarrier(preDrawBarriers);
+
 	gl.prepareDraw(this);
 	gl.setVertexAttributes(attributes, *cmd.buffers);
 	gl.bindTextureToUnit(cmd.texture, 0, false);
@@ -606,6 +627,9 @@ void Graphics::draw(const DrawIndexedCommand &cmd)
 		glDrawElementsInstanced(glprimitivetype, cmd.indexCount, gldatatype, gloffset, cmd.instanceCount);
 	else
 		glDrawElements(glprimitivetype, cmd.indexCount, gldatatype, gloffset);
+	
+	if (postDrawBarriers != 0)
+		glMemoryBarrier(postDrawBarriers);
 
 	++drawCalls;
 }
@@ -638,6 +662,12 @@ void Graphics::drawQuads(int start, int count, VertexAttributesID attributesID,
 	const int MAX_VERTICES_PER_DRAW = LOVE_UINT16_MAX;
 	const int MAX_QUADS_PER_DRAW    = MAX_VERTICES_PER_DRAW / 4;
 
+	GLbitfield preDrawBarriers = 0;
+	GLbitfield postDrawBarriers = 0;
+
+	if (!shaderBarriers((Shader *)Shader::current, preDrawBarriers, postDrawBarriers))
+		return;
+
 	VertexAttributes attributes;
 	findVertexAttributes(attributesID, attributes);
 
@@ -655,9 +685,16 @@ void Graphics::drawQuads(int start, int count, VertexAttributesID attributesID,
 
 		for (int quadindex = 0; quadindex < count; quadindex += MAX_QUADS_PER_DRAW)
 		{
+			if (preDrawBarriers != 0)
+				glMemoryBarrier(preDrawBarriers);
+
 			int quadcount = std::min(MAX_QUADS_PER_DRAW, count - quadindex);
 
 			glDrawElementsBaseVertex(GL_TRIANGLES, quadcount * 6, GL_UNSIGNED_SHORT, BUFFER_OFFSET(0), basevertex);
+
+			if (postDrawBarriers != 0)
+				glMemoryBarrier(postDrawBarriers);
+
 			++drawCalls;
 
 			basevertex += quadcount * 4;
@@ -671,11 +708,18 @@ void Graphics::drawQuads(int start, int count, VertexAttributesID attributesID,
 
 		for (int quadindex = 0; quadindex < count; quadindex += MAX_QUADS_PER_DRAW)
 		{
+			if (preDrawBarriers != 0)
+				glMemoryBarrier(preDrawBarriers);
+
 			gl.setVertexAttributes(attributes, bufferscopy);
 
 			int quadcount = std::min(MAX_QUADS_PER_DRAW, count - quadindex);
 
 			glDrawElements(GL_TRIANGLES, quadcount * 6, GL_UNSIGNED_SHORT, BUFFER_OFFSET(0));
+
+			if (postDrawBarriers != 0)
+				glMemoryBarrier(postDrawBarriers);
+
 			++drawCalls;
 
 			if (count > MAX_QUADS_PER_DRAW)
@@ -1593,7 +1637,9 @@ void Graphics::initCapabilities()
 	capabilities.features[FEATURE_TEXEL_BUFFER] = gl.isBufferUsageSupported(BUFFERUSAGE_TEXEL);
 	capabilities.features[FEATURE_COPY_TEXTURE_TO_BUFFER] = gl.isCopyTextureToBufferSupported();
 	capabilities.features[FEATURE_INDIRECT_DRAW] = capabilities.features[FEATURE_GLSL4];
-	static_assert(FEATURE_MAX_ENUM == 13, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
+	capabilities.features[FEATURE_VERTEX_WRITE] = capabilities.features[FEATURE_GLSL4];
+	capabilities.features[FEATURE_PIXEL_WRITE] = capabilities.features[FEATURE_GLSL4];
+	static_assert(FEATURE_MAX_ENUM == 15, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
 
 	capabilities.limits[LIMIT_POINT_SIZE] = gl.getMaxPointSize();
 	capabilities.limits[LIMIT_TEXTURE_SIZE] = gl.getMax2DTextureSize();

src/modules/graphics/vulkan/Graphics.cpp

@@ -821,6 +821,8 @@ bool Graphics::setMode(void *context, const BackbufferSettings &settings)
 
 void Graphics::initCapabilities()
 {
+	VkPhysicalDeviceFeatures features;
+	vkGetPhysicalDeviceFeatures(physicalDevice, &features);
 	capabilities.features[FEATURE_MULTI_RENDER_TARGET_FORMATS] = true;
 	capabilities.features[FEATURE_CLAMP_ZERO] = true;
 	capabilities.features[FEATURE_CLAMP_ONE] = true;
@@ -834,7 +836,9 @@ void Graphics::initCapabilities()
 	capabilities.features[FEATURE_TEXEL_BUFFER] = true;
 	capabilities.features[FEATURE_COPY_TEXTURE_TO_BUFFER] = true;
 	capabilities.features[FEATURE_INDIRECT_DRAW] = true;
-	static_assert(FEATURE_MAX_ENUM == 13, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
+	capabilities.features[FEATURE_VERTEX_WRITE] = features.vertexPipelineStoresAndAtomics;
+	capabilities.features[FEATURE_PIXEL_WRITE] = features.fragmentStoresAndAtomics;
+	static_assert(FEATURE_MAX_ENUM == 15, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
 
 	VkPhysicalDeviceProperties properties;
 	vkGetPhysicalDeviceProperties(physicalDevice, &properties);
@@ -953,6 +957,11 @@ void Graphics::draw(const DrawCommand &cmd)
 {
 	prepareDraw(cmd.attributesID, *cmd.buffers, cmd.texture, cmd.primitiveType, cmd.cullMode);
 
+	VkAccessFlags dstAccessMask;
+	VkPipelineStageFlags dstStageMask;
+	if (!prepareBarrier(dstAccessMask, dstStageMask))
+		return;
+
 	if (cmd.indirectBuffer != nullptr)
 	{
 		vkCmdDrawIndirect(
@@ -972,6 +981,7 @@ void Graphics::draw(const DrawCommand &cmd)
 			0);
 	}
 
+	tryBarrier(dstAccessMask, dstStageMask);
 	drawCalls++;
 }
 
@@ -979,6 +989,11 @@ void Graphics::draw(const DrawIndexedCommand &cmd)
 {
 	prepareDraw(cmd.attributesID, *cmd.buffers, cmd.texture, cmd.primitiveType, cmd.cullMode);
 
+	VkAccessFlags dstAccessMask;
+	VkPipelineStageFlags dstStageMask;
+	if (!prepareBarrier(dstAccessMask, dstStageMask))
+		return;
+
 	vkCmdBindIndexBuffer(
 		commandBuffers.at(currentFrame),
 		(VkBuffer) cmd.indexBuffer->getHandle(),
@@ -1005,6 +1020,7 @@ void Graphics::draw(const DrawIndexedCommand &cmd)
 			0);
 	}
 
+	tryBarrier(dstAccessMask, dstStageMask);
 	drawCalls++;
 }
 
@@ -1015,6 +1031,12 @@ void Graphics::drawQuads(int start, int count, VertexAttributesID attributesID,
 
 	prepareDraw(attributesID, buffers, texture, PRIMITIVE_TRIANGLES, CULL_NONE);
 
+	VkAccessFlags dstAccessMask;
+	VkPipelineStageFlags dstStageMask;
+	if (!prepareBarrier(dstAccessMask, dstStageMask))
+		return;
+
+
 	vkCmdBindIndexBuffer(
 		commandBuffers.at(currentFrame),
 		(VkBuffer)quadIndexBuffer->getHandle(),
@@ -1036,6 +1058,7 @@ void Graphics::drawQuads(int start, int count, VertexAttributesID attributesID,
 			0);
 		baseVertex += quadcount * 4;
 
+		tryBarrier(dstAccessMask, dstStageMask);
 		drawCalls++;
 	}
 }
@@ -1291,7 +1314,7 @@ graphics::StreamBuffer *Graphics::newStreamBuffer(BufferUsage type, size_t size)
 	return new StreamBuffer(this, type, size);
 }
 
-static bool computeDispatchBarrierFlags(Shader *shader, VkAccessFlags &dstAccessFlags, VkPipelineStageFlags &dstStageFlags)
+static bool shaderBarrierFlags(Shader *shader, VkAccessFlags &dstAccessFlags, VkPipelineStageFlags &dstStageFlags)
 {
 	for (const auto &info : shader->getActiveTextureInfo())
 	{
@@ -1333,7 +1356,7 @@ bool Graphics::dispatch(love::graphics::Shader *shader, int x, int y, int z)
 	barrier.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER;
 	barrier.srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT;
 	VkPipelineStageFlags dstStageMask = 0;
-	if (!computeDispatchBarrierFlags(computeShader, barrier.dstAccessMask, dstStageMask))
+	if (!shaderBarrierFlags(computeShader, barrier.dstAccessMask, dstStageMask))
 		return false;
 
 	usedShadersInFrame.insert(computeShader);
@@ -1362,7 +1385,7 @@ bool Graphics::dispatch(love::graphics::Shader *shader, love::graphics::Buffer *
 	barrier.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER;
 	barrier.srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT;
 	VkPipelineStageFlags dstStageMask = 0;
-	if (!computeDispatchBarrierFlags(computeShader, barrier.dstAccessMask, dstStageMask))
+	if (!shaderBarrierFlags(computeShader, barrier.dstAccessMask, dstStageMask))
 		return false;
 
 	usedShadersInFrame.insert(computeShader);
@@ -2788,6 +2811,32 @@ void Graphics::prepareDraw(VertexAttributesID attributesID, const BufferBindings
 		vkCmdBindVertexBuffers(commandBuffers.at(currentFrame), VERTEX_BUFFER_BINDING_START, buffercount, vkbuffers, vkoffsets);
 }
 
+bool Graphics::prepareBarrier(VkAccessFlags &dstAccessMask, VkPipelineStageFlags &dstStageMask)
+{
+	auto shader = dynamic_cast<Shader *>(Shader::current);
+	if (!shader)
+		return false;
+
+	if (!shaderBarrierFlags(shader, dstAccessMask, dstStageMask))
+		return false;
+
+	return true;
+}
+
+void Graphics::tryBarrier(VkAccessFlags dstAccessMask, VkPipelineStageFlags dstStageMask)
+{
+	if (dstAccessMask == 0 && dstStageMask == 0)
+		return;
+
+	VkMemoryBarrier barrier{};
+	barrier.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER;
+	barrier.srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT;
+	barrier.dstAccessMask = dstAccessMask;
+
+	if (barrier.dstAccessMask != 0 || dstStageMask != 0)
+		vkCmdPipelineBarrier(commandBuffers.at(currentFrame), VK_PIPELINE_STAGE_VERTEX_SHADER_BIT | VK_PIPELINE_STAGE_FRAGMENT_SHADER_BIT, dstStageMask, 0, 1, &barrier, 0, nullptr, 0, nullptr);
+}
+
 void Graphics::setDefaultRenderPass()
 {
 	uint32_t numClearValues = 2;

src/modules/graphics/vulkan/Graphics.h

@@ -372,6 +372,8 @@ private:
 		VertexAttributesID attributesID,
 		const BufferBindings &buffers, graphics::Texture *texture,
 		PrimitiveType, CullMode);
+	bool prepareBarrier(VkAccessFlags &dstAccessMask, VkPipelineStageFlags &dstStageMask);
+	void tryBarrier(VkAccessFlags dstAccessMask, VkPipelineStageFlags dstStageMask);
 	void setRenderPass(const RenderTargets &rts, int pixelw, int pixelh);
 	void setDefaultRenderPass();
 	void startRenderPass();

src/modules/graphics/wrap_Graphics.cpp

@@ -1586,6 +1586,20 @@ static int w_getShaderSource(lua_State *L, int startidx, std::vector<std::string
 		if (!lua_isnoneornil(L, -1))
 			options.debugName = luax_checkstring(L, -1);
 		lua_pop(L, 1);
+
+		for (int feature = 0; feature < Shader::FEATURE_MAX_ENUM; ++feature)
+		{
+			const char *featureKey;
+			if (Shader::getConstant((Shader::Feature)feature, featureKey))
+			{
+				lua_getfield(L, optionsidx, featureKey);
+				if (!lua_isnoneornil(L, -1))
+				{
+					options.features[feature] = lua_toboolean(L, -1);
+				}
+				lua_pop(L, 1);
+			}
+		}
 	}
 
 	return 0;

testing/tests/graphics.lua

@@ -1003,6 +1003,90 @@ love.test.graphics.Shader = function(test)
   else
     test:assertTrue(true, "skip shader IO test")
   end
+
+  if love.graphics.getSupported().glsl4 and love.graphics.getSupported().vertexwrite and love.graphics.getSupported().pixelwrite then
+    local name = love.graphics.getRendererInfo()
+    local isGLES = name:match("OpenGL ES") ~= nil
+
+    local success, message = love.graphics.validateShader(isGLES, [[
+      #pragma language glsl4
+
+      #ifdef PIXEL
+      buffer ColorBuffer
+      {
+        vec4 colors[];
+      };
+
+      void effect()
+      {
+        colors[0] = VaryingColor;
+        love_PixelColor = colors[0];
+      }
+      #endif
+
+      #ifdef VERTEX
+      vec4 position(mat4 m, vec4 p)
+      {
+        return m * p;
+      }
+      #endif
+    ]])
+
+    test:assertFalse(success, "shader should not validate (SSBO)")
+    test:assertEquals("Shader validation error:\nStorage Buffer block 'ColorBuffer' must be marked as readonly in vertex and pixel shaders unless explicitly enabled.", message)
+
+    success, message = love.graphics.validateShader(isGLES, [[
+      #pragma language glsl4
+
+      buffer ColorBuffer
+      {
+        vec4 colors[];
+      };
+
+      void effect()
+      {
+        colors[0] = VaryingColor;
+        love_PixelColor = colors[0];
+      }
+    ]], { write = true })
+
+    test:assertTrue(success, "shader should validate (SSBO)")
+    test:assertEquals(nil, message)
+
+    success, message = love.graphics.validateShader(isGLES, [[
+        #pragma language glsl4
+
+        layout(rgba32f) uniform readonly highp image2D RGBAImage;
+
+        void effect()
+        {
+          love_PixelColor = VaryingColor * imageLoad(RGBAImage, ivec2(gl_FragCoord.xy));
+        }
+    ]])
+      
+    test:assertFalse(success, "shader should not validate (image)")
+    test:assertEquals("Shader validation error:\nStorage Texture uniform variables (image2D, etc) are only allowed in compute shaders unless explicitly enabled.", message)
+    
+    success, message = love.graphics.validateShader(isGLES, [[
+      #pragma language glsl4
+
+      #ifdef GL_ES
+        precision highp float;
+      #endif
+
+      layout(rgba32f) uniform readonly highp image2D RGBAImage;
+
+      void effect()
+      {
+        love_PixelColor = VaryingColor * imageLoad(RGBAImage, ivec2(gl_FragCoord.xy));
+      }
+    ]], { write = true })
+
+    test:assertTrue(success, "shader should validate (image)")
+    test:assertEquals(nil, message)
+  else
+    test:assertTrue(true, "skip feature test")
+  end
 end