Add functions for indirect draws and compute dispatches.

Resolves #1879.

- Add an 'indirectdraw' boolean field to the graphics feature support table returned by love.graphics.getSupported. This is almost always supported when compute shaders are supported, except on a few older phones.

- Add an 'indirectarguments' boolean field to the settings table in love.graphics.newBuffer.

- Add love.graphics.dispatchIndirect(shader, argumentsbuffer [, argumentsindex = 1]).
    The compute dispatch's threadgroup width, height, and depth values are fetched from the buffer (as 3 uints) instead of coming from function parameters.

- Add love.graphics.drawIndirect(mesh, argumentsbuffer, argumentsindex, x, y, ....).
    Vertex or index count, instance count, and other related parameters for drawing the mesh are fetched from the buffer (as 4 or 5 uints depending on whether the mesh has an index buffer) instead of coming from function parameters. It's usually a good idea to keep parameters other than the instance count in sync with what the mesh should be using.

- Add love.graphics.drawFromShaderIndirect(drawmode, argumentsbuffer [, argumentsindex = 1] [, maintexture = nil]) and drawFromShaderIndirect(indexbuffer, argumentsbuffer [, argumentsindex] [, maintexture = nil]).
    Vertex or index count, instance count, and other related parameters are fetched from the buffer as 4 or 5 uints, as above.

For the dispatch indirect arguments buffer, it has to have 3 uint32 elements: { uint threadgroupsX, uint threadgroupsY, uint threadgroupsZ }.

For non-indexed draws, the arguments buffer has to have 4 uint32 elements: { uint vertexCount, uint instanceCount, uint baseVertex, uint baseInstance }. Note that baseInstance should always be set to 0 as many drivers don't support non-zero values.

For draws which use an index buffer, the arguments buffer has to have 5 uint32 elements: { uint indexCount, uint instanceCount, uint firstIndex, uint baseVertex, uint baseInstance }. As above, the baseInstance value should always be 0.

A buffer can be created to have an array of those structures, which can be used with the argumentsindex parameter of the Indirect dispatch/draw functions.

src/modules/graphics/Buffer.cpp

@@ -52,6 +52,7 @@ Buffer::Buffer(Graphics *gfx, const Settings &settings, const std::vector<DataDe
 	bool vertexbuffer = usageFlags & BUFFERUSAGEFLAG_VERTEX;
 	bool texelbuffer = usageFlags & BUFFERUSAGEFLAG_TEXEL;
 	bool storagebuffer = usageFlags & BUFFERUSAGEFLAG_SHADER_STORAGE;
+	bool indirectbuffer = usageFlags & BUFFERUSAGEFLAG_INDIRECT_ARGUMENTS;
 
 	if (texelbuffer && !caps.features[Graphics::FEATURE_TEXEL_BUFFER])
 		throw love::Exception("Texel buffers are not supported on this system.");
@@ -62,8 +63,11 @@ Buffer::Buffer(Graphics *gfx, const Settings &settings, const std::vector<DataDe
 	if (storagebuffer && dataUsage == BUFFERDATAUSAGE_STREAM)
 		throw love::Exception("Buffers created with 'stream' data usage cannot be used as a shader storage buffer.");
 
-	if (dataUsage == BUFFERDATAUSAGE_READBACK && (indexbuffer || vertexbuffer || texelbuffer || storagebuffer))
-		throw love::Exception("Buffers created with 'readback' data usage cannot be index, vertex, texel, or shaderstorage buffer types.");
+	if (indirectbuffer && !caps.features[Graphics::FEATURE_INDIRECT_DRAW])
+		throw love::Exception("Indirect argument buffers are not supported on this system.");
+
+	if (dataUsage == BUFFERDATAUSAGE_READBACK && (indexbuffer || vertexbuffer || texelbuffer || storagebuffer || indirectbuffer))
+		throw love::Exception("Buffers created with 'readback' data usage cannot be index, vertex, texel, shaderstorage, or indirectarguments buffer types.");
 
 	size_t offset = 0;
 	size_t stride = 0;
@@ -184,6 +188,18 @@ Buffer::Buffer(Graphics *gfx, const Settings &settings, const std::vector<DataDe
 					member.decl.name.c_str(), memberoffset, offset);
 		}
 
+		if (indirectbuffer)
+		{
+			if (info.isMatrix || info.components != 1
+				|| (info.baseType != DATA_BASETYPE_UINT && info.baseType != DATA_BASETYPE_INT))
+			{
+				throw love::Exception("Indirect argument buffers must use single-component int or uint types.");
+			}
+
+			if (bufferformat.size() > 5)
+				throw love::Exception("Indirect argument buffers only support up to 5 values per array element.");
+		}
+
 		member.offset = memberoffset;
 		member.size = membersize;
 

src/modules/graphics/Graphics.cpp

@@ -1542,7 +1542,49 @@ void Graphics::copyBufferToTexture(Buffer *source, Texture *dest, size_t sourceo
 	dest->copyFromBuffer(source, sourceoffset, sourcewidth, size, slice, mipmap, rect);
 }
 
-void Graphics::dispatchThreadgroups(Shader* shader, int x, int y, int z)
+static const char *getIndirectArgsTypeName(Graphics::IndirectArgsType argstype)
+{
+	switch (argstype)
+	{
+		case Graphics::INDIRECT_ARGS_DISPATCH: return "Compute shader threadgroup argument data";
+		case Graphics::INDIRECT_ARGS_DRAW_VERTICES: return "Draw vertices argument data";
+		case Graphics::INDIRECT_ARGS_DRAW_INDICES: return "Draw indices argument data";
+	}
+
+	return "(Unknown argument data)";
+}
+
+void Graphics::validateIndirectArgsBuffer(IndirectArgsType argstype, Buffer *indirectargs, int argsindex)
+{
+	if (!capabilities.features[FEATURE_INDIRECT_DRAW])
+		throw love::Exception("Indirect draws and compute dispatches are not supported on this system.");
+
+	if ((indirectargs->getUsageFlags() & BUFFERUSAGEFLAG_INDIRECT_ARGUMENTS) == 0)
+		throw love::Exception("The given Buffer must be created with the indirectarguments usage flag set, to be used for indirect arguments.");
+
+	if (argsindex < 0)
+		throw love::Exception("The given indirect argument index cannot be negative.");
+
+	size_t argelements = 0;
+	if (argstype == INDIRECT_ARGS_DISPATCH)
+		argelements = 3;
+	else if (argstype == INDIRECT_ARGS_DRAW_VERTICES)
+		argelements = 4;
+	else if (argstype == INDIRECT_ARGS_DRAW_INDICES)
+		argelements = 5;
+
+	size_t totalmembers = indirectargs->getArrayLength() * indirectargs->getDataMembers().size();
+
+	if (totalmembers % argelements != 0)
+		throw love::Exception("%s requires the given indirect argument Buffer to have a multiple of %ld int or uint values.", getIndirectArgsTypeName(argstype), argelements);
+
+	size_t argsoffset = argsindex * indirectargs->getArrayStride();
+
+	if (indirectargs->getSize() < argsoffset + sizeof(uint32) * argelements)
+		throw love::Exception("The given index into the indirect argument Buffer does not fit within the Buffer's size.");
+}
+
+void Graphics::dispatchThreadgroups(Shader *shader, int x, int y, int z)
 {
 	if (!shader->hasStage(SHADERSTAGE_COMPUTE))
 		throw love::Exception("Only compute shaders can have threads dispatched.");
@@ -1562,7 +1604,28 @@ void Graphics::dispatchThreadgroups(Shader* shader, int x, int y, int z)
 	auto prevshader = Shader::current;
 	shader->attach();
 
-	bool success = dispatch(x, y, z);
+	bool success = dispatch(shader, x, y, z);
+
+	if (prevshader != nullptr)
+		prevshader->attach();
+
+	if (!success)
+		throw love::Exception("Compute shader must have resources bound to all writable texture and buffer variables.");
+}
+
+void Graphics::dispatchIndirect(Shader *shader, Buffer *indirectargs, int argsindex)
+{
+	if (!shader->hasStage(SHADERSTAGE_COMPUTE))
+		throw love::Exception("Only compute shaders can have threads dispatched.");
+
+	validateIndirectArgsBuffer(INDIRECT_ARGS_DISPATCH, indirectargs, argsindex);
+
+	flushBatchedDraws();
+
+	auto prevshader = Shader::current;
+	shader->attach();
+
+	bool success = dispatch(shader, indirectargs, argsindex * indirectargs->getArrayStride());
 
 	if (prevshader != nullptr)
 		prevshader->attach();
@@ -1819,6 +1882,11 @@ void Graphics::drawInstanced(Mesh *mesh, const Matrix4 &m, int instancecount)
 	mesh->drawInstanced(this, m, instancecount);
 }
 
+void Graphics::drawIndirect(Mesh *mesh, const Matrix4 &m, Buffer *indirectargs, int argsindex)
+{
+	mesh->drawIndirect(this, m, indirectargs, argsindex);
+}
+
 void Graphics::drawFromShader(PrimitiveType primtype, int vertexcount, int instancecount, Texture *maintexture)
 {
 	if (primtype == PRIMITIVE_TRIANGLE_FAN && vertexcount > LOVE_UINT16_MAX)
@@ -1899,6 +1967,61 @@ void Graphics::drawFromShader(Buffer *indexbuffer, int indexcount, int instancec
 	draw(cmd);
 }
 
+void Graphics::drawFromShaderIndirect(PrimitiveType primtype, Buffer *indirectargs, int argsindex, Texture *maintexture)
+{
+	flushBatchedDraws();
+
+	if (primtype == PRIMITIVE_TRIANGLE_FAN)
+		throw love::Exception("The fan draw mode is not supported in indirect draws.");
+
+	if (Shader::isDefaultActive() || !Shader::current)
+		throw love::Exception("drawFromShaderIndirect can only be used with a custom shader.");
+
+	validateIndirectArgsBuffer(INDIRECT_ARGS_DRAW_VERTICES, indirectargs, argsindex);
+
+	Shader::current->validateDrawState(primtype, maintexture);
+
+	VertexAttributes attributes;
+	BufferBindings buffers;
+
+	DrawCommand cmd(&attributes, &buffers);
+
+	cmd.primitiveType = primtype;
+	cmd.indirectBuffer = indirectargs;
+	cmd.indirectBufferOffset = argsindex * indirectargs->getArrayStride();
+	cmd.texture = maintexture;
+
+	draw(cmd);
+}
+
+void Graphics::drawFromShaderIndirect(Buffer *indexbuffer, Buffer *indirectargs, int argsindex, Texture *maintexture)
+{
+	flushBatchedDraws();
+
+	if (!(indexbuffer->getUsageFlags() & BUFFERUSAGEFLAG_INDEX))
+		throw love::Exception("The buffer passed to the indexed variant of drawFromShaderIndirect must be an index buffer.");
+
+	if (Shader::isDefaultActive() || !Shader::current)
+		throw love::Exception("drawFromShaderIndirect can only be used with a custom shader.");
+
+	validateIndirectArgsBuffer(INDIRECT_ARGS_DRAW_INDICES, indirectargs, argsindex);
+
+	Shader::current->validateDrawState(PRIMITIVE_TRIANGLES, maintexture);
+
+	VertexAttributes attributes;
+	BufferBindings buffers;
+
+	DrawIndexedCommand cmd(&attributes, &buffers, indexbuffer);
+
+	cmd.primitiveType = PRIMITIVE_TRIANGLES;
+	cmd.indexType = getIndexDataType(indexbuffer->getDataMember(0).decl.format);
+	cmd.indirectBuffer = indirectargs;
+	cmd.indexBufferOffset = argsindex * indirectargs->getArrayStride();
+	cmd.texture = maintexture;
+
+	draw(cmd);
+}
+
 void Graphics::print(const std::vector<love::font::ColoredString> &str, const Matrix4 &m)
 {
 	checkSetDefaultFont();
@@ -2533,6 +2656,7 @@ STRINGMAP_CLASS_BEGIN(Graphics, Graphics::Feature, Graphics::FEATURE_MAX_ENUM, f
 	{ "copytexturetobuffer",      Graphics::FEATURE_COPY_TEXTURE_TO_BUFFER },
 	{ "copyrendertargettobuffer", Graphics::FEATURE_COPY_RENDER_TARGET_TO_BUFFER },
 	{ "mipmaprange",              Graphics::FEATURE_MIPMAP_RANGE         },
+	{ "indirectdraw",             Graphics::FEATURE_INDIRECT_DRAW        },
 }
 STRINGMAP_CLASS_END(Graphics, Graphics::Feature, Graphics::FEATURE_MAX_ENUM, feature)
 

src/modules/graphics/Graphics.h

@@ -165,6 +165,7 @@ public:
 		FEATURE_COPY_TEXTURE_TO_BUFFER,
 		FEATURE_COPY_RENDER_TARGET_TO_BUFFER,
 		FEATURE_MIPMAP_RANGE,
+		FEATURE_INDIRECT_DRAW,
 		FEATURE_MAX_ENUM
 	};
 
@@ -199,6 +200,13 @@ public:
 		TEMPORARY_RT_STENCIL = (1 << 1),
 	};
 
+	enum IndirectArgsType
+	{
+		INDIRECT_ARGS_DISPATCH,
+		INDIRECT_ARGS_DRAW_VERTICES,
+		INDIRECT_ARGS_DRAW_INDICES,
+	};
+
 	struct Capabilities
 	{
 		double limits[LIMIT_MAX_ENUM];
@@ -236,6 +244,9 @@ public:
 		int vertexCount = 0;
 		int instanceCount = 1;
 
+		Buffer *indirectBuffer = nullptr;
+		size_t indirectBufferOffset = 0;
+
 		Texture *texture = nullptr;
 
 		// TODO: This should be moved out to a state transition API?
@@ -261,6 +272,9 @@ public:
 		Resource *indexBuffer;
 		size_t indexBufferOffset = 0;
 
+		Buffer *indirectBuffer = nullptr;
+		size_t indirectBufferOffset = 0;
+
 		Texture *texture = nullptr;
 
 		// TODO: This should be moved out to a state transition API?
@@ -689,16 +703,20 @@ public:
 	void copyTextureToBuffer(Texture *source, Buffer *dest, int slice, int mipmap, const Rect &rect, size_t destoffset, int destwidth);
 	void copyBufferToTexture(Buffer *source, Texture *dest, size_t sourceoffset, int sourcewidth, int slice, int mipmap, const Rect &rect);
 
-	void dispatchThreadgroups(Shader* shader, int x, int y, int z);
+	void dispatchThreadgroups(Shader *shader, int x, int y, int z);
+	void dispatchIndirect(Shader *shader, Buffer *indirectargs, int argsindex);
 
 	void draw(Drawable *drawable, const Matrix4 &m);
 	void draw(Texture *texture, Quad *quad, const Matrix4 &m);
 	void drawLayer(Texture *texture, int layer, const Matrix4 &m);
 	void drawLayer(Texture *texture, int layer, Quad *quad, const Matrix4 &m);
 	void drawInstanced(Mesh *mesh, const Matrix4 &m, int instancecount);
+	void drawIndirect(Mesh *mesh, const Matrix4 &m, Buffer *indirectargs, int argsindex);
 
 	void drawFromShader(PrimitiveType primtype, int vertexcount, int instancecount, Texture *maintexture);
 	void drawFromShader(Buffer *indexbuffer, int indexcount, int instancecount, int startindex, Texture *maintexture);
+	void drawFromShaderIndirect(PrimitiveType primtype, Buffer *indirectargs, int argsindex, Texture *maintexture);
+	void drawFromShaderIndirect(Buffer *indexbuffer, Buffer *indirectargs, int argsindex, Texture *maintexture);
 
 	/**
 	 * Draws text at the specified coordinates
@@ -873,6 +891,8 @@ public:
 
 	void cleanupCachedShaderStage(ShaderStageType type, const std::string &cachekey);
 
+	void validateIndirectArgsBuffer(IndirectArgsType argstype, Buffer *indirectargs, int argsindex);
+
 	template <typename T>
 	T *getScratchBuffer(size_t count)
 	{
@@ -1001,7 +1021,8 @@ protected:
 	virtual GraphicsReadback *newReadbackInternal(ReadbackMethod method, Buffer *buffer, size_t offset, size_t size, data::ByteData *dest, size_t destoffset) = 0;
 	virtual GraphicsReadback *newReadbackInternal(ReadbackMethod method, Texture *texture, int slice, int mipmap, const Rect &rect, image::ImageData *dest, int destx, int desty) = 0;
 
-	virtual bool dispatch(int x, int y, int z) = 0;
+	virtual bool dispatch(Shader *shader, int x, int y, int z) = 0;
+	virtual bool dispatch(Shader *shader, Buffer *indirectargs, size_t argsoffset) = 0;
 
 	virtual void setRenderTargetsInternal(const RenderTargets &rts, int pixelw, int pixelh, bool hasSRGBtexture) = 0;
 

src/modules/graphics/Mesh.cpp

@@ -525,17 +525,38 @@ bool Mesh::getDrawRange(int &start, int &count) const
 
 void Mesh::draw(Graphics *gfx, const love::Matrix4 &m)
 {
-	drawInstanced(gfx, m, 1);
+	drawInternal(gfx, m, 1, nullptr, 0);
 }
 
 void Mesh::drawInstanced(Graphics *gfx, const Matrix4 &m, int instancecount)
 {
-	if (vertexCount <= 0 || instancecount <= 0)
+	drawInternal(gfx, m, instancecount, nullptr, 0);
+}
+
+void Mesh::drawIndirect(Graphics *gfx, const Matrix4 &m, Buffer *indirectargs, int argsindex)
+{
+	drawInternal(gfx, m, 0, indirectargs, argsindex);
+}
+
+void Mesh::drawInternal(Graphics *gfx, const Matrix4 &m, int instancecount, Buffer *indirectargs, int argsindex)
+{
+	if (vertexCount <= 0 || (instancecount <= 0 && indirectargs == nullptr))
 		return;
 
 	if (instancecount > 1 && !gfx->getCapabilities().features[Graphics::FEATURE_INSTANCING])
 		throw love::Exception("Instancing is not supported on this system.");
 
+	if (indirectargs != nullptr)
+	{
+		if (primitiveType == PRIMITIVE_TRIANGLE_FAN)
+			throw love::Exception("The fan draw mode is not supported in indirect draws.");
+
+		if (useIndexBuffer && indexBuffer != nullptr)
+			gfx->validateIndirectArgsBuffer(Graphics::INDIRECT_ARGS_DRAW_INDICES, indirectargs, argsindex);
+		else
+			gfx->validateIndirectArgsBuffer(Graphics::INDIRECT_ARGS_DRAW_VERTICES, indirectargs, argsindex);
+	}
+
 	// Some graphics backends don't natively support triangle fans. So we'd
 	// have to emulate them with triangles plus an index buffer... which doesn't
 	// work so well when there's already a custom index buffer.
@@ -616,7 +637,7 @@ void Mesh::drawInstanced(Graphics *gfx, const Matrix4 &m, int instancecount)
 		}
 	}
 
-	if (indexbuffer != nullptr && indexcount > 0)
+	if (indexbuffer != nullptr && (indexcount > 0 || indirectargs != nullptr))
 	{
 		Range r(0, indexcount);
 		if (range.isValid())
@@ -633,10 +654,13 @@ void Mesh::drawInstanced(Graphics *gfx, const Matrix4 &m, int instancecount)
 		cmd.indexBufferOffset = r.getOffset() * indexbuffer->getArrayStride();
 		cmd.indexCount = (int) r.getSize();
 
+		cmd.indirectBuffer = indirectargs;
+		cmd.indirectBufferOffset = argsindex * (indirectargs != nullptr ? indirectargs->getArrayStride() : 0);
+
 		if (cmd.indexCount > 0)
 			gfx->draw(cmd);
 	}
-	else if (vertexCount > 0)
+	else if (vertexCount > 0 || indirectargs != nullptr)
 	{
 		Range r(0, vertexCount);
 		if (range.isValid())
@@ -651,6 +675,9 @@ void Mesh::drawInstanced(Graphics *gfx, const Matrix4 &m, int instancecount)
 		cmd.texture = texture;
 		cmd.cullMode = gfx->getMeshCullMode();
 
+		cmd.indirectBuffer = indirectargs;
+		cmd.indirectBufferOffset = argsindex * (indirectargs != nullptr ? indirectargs->getArrayStride() : 0);
+
 		if (cmd.vertexCount > 0)
 			gfx->draw(cmd);
 	}

src/modules/graphics/Mesh.h

@@ -176,6 +176,7 @@ public:
 	void draw(Graphics *gfx, const Matrix4 &m) override;
 
 	void drawInstanced(Graphics *gfx, const Matrix4 &m, int instancecount);
+	void drawIndirect(Graphics *gfx, const Matrix4 &m, Buffer *indirectargs, int argsindex);
 
 	static std::vector<Buffer::DataDeclaration> getDefaultVertexFormat();
 
@@ -186,6 +187,8 @@ private:
 	void setupAttachedAttributes();
 	int getAttachedAttributeIndex(const std::string &name) const;
 
+	void drawInternal(Graphics *gfx, const Matrix4 &m, int instancecount, Buffer *indirectargs, int argsindex);
+
 	std::vector<Buffer::DataMember> vertexFormat;
 
 	std::vector<BufferAttribute> attachedAttributes;

src/modules/graphics/metal/Graphics.h

@@ -74,7 +74,8 @@ public:
 
 	void setActive(bool active) override;
 
-	bool dispatch(int x, int y, int z) override;
+	bool dispatch(love::graphics::Shader *shader, int x, int y, int z) override;
+	bool dispatch(love::graphics::Shader *shader, love::graphics::Buffer *indirectargs, size_t argsoffset) override;
 
 	void draw(const DrawCommand &cmd) override;
 	void draw(const DrawIndexedCommand &cmd) override;

src/modules/graphics/metal/Graphics.mm

@@ -1204,10 +1204,19 @@ void Graphics::draw(const DrawCommand &cmd)
 
 	setVertexBuffers(encoder, Shader::current, cmd.buffers, renderBindings);
 
-	[encoder drawPrimitives:getMTLPrimitiveType(cmd.primitiveType)
-				vertexStart:cmd.vertexStart
-				vertexCount:cmd.vertexCount
-			  instanceCount:cmd.instanceCount];
+	if (cmd.indirectBuffer != nullptr)
+	{
+		[encoder drawPrimitives:getMTLPrimitiveType(cmd.primitiveType)
+				 indirectBuffer:getMTLBuffer(cmd.indirectBuffer)
+		   indirectBufferOffset:cmd.indirectBufferOffset];
+	}
+	else
+	{
+		[encoder drawPrimitives:getMTLPrimitiveType(cmd.primitiveType)
+					vertexStart:cmd.vertexStart
+					vertexCount:cmd.vertexCount
+				  instanceCount:cmd.instanceCount];
+	}
 
 	++drawCalls;
 }}
@@ -1229,12 +1238,24 @@ void Graphics::draw(const DrawIndexedCommand &cmd)
 
 	auto indexType = cmd.indexType == INDEX_UINT32 ? MTLIndexTypeUInt32 : MTLIndexTypeUInt16;
 
-	[encoder drawIndexedPrimitives:getMTLPrimitiveType(cmd.primitiveType)
-						indexCount:cmd.indexCount
-						 indexType:indexType
-					   indexBuffer:getMTLBuffer(cmd.indexBuffer)
-				 indexBufferOffset:cmd.indexBufferOffset
-					 instanceCount:cmd.instanceCount];
+	if (cmd.indirectBuffer != nullptr)
+	{
+		[encoder drawIndexedPrimitives:getMTLPrimitiveType(cmd.primitiveType)
+							 indexType:indexType
+						   indexBuffer:getMTLBuffer(cmd.indexBuffer)
+					 indexBufferOffset:cmd.indexBufferOffset
+						indirectBuffer:getMTLBuffer(cmd.indirectBuffer)
+				  indirectBufferOffset:cmd.indexBufferOffset];
+	}
+	else
+	{
+		[encoder drawIndexedPrimitives:getMTLPrimitiveType(cmd.primitiveType)
+							indexCount:cmd.indexCount
+							 indexType:indexType
+						   indexBuffer:getMTLBuffer(cmd.indexBuffer)
+					 indexBufferOffset:cmd.indexBufferOffset
+						 instanceCount:cmd.instanceCount];
+	}
 
 	++drawCalls;
 }}
@@ -1330,10 +1351,9 @@ void Graphics::drawQuads(int start, int count, const VertexAttributes &attribute
 	}
 }}
 
-bool Graphics::dispatch(int x, int y, int z)
+bool Graphics::dispatch(love::graphics::Shader *s, int x, int y, int z)
 { @autoreleasepool {
-	// Set by higher level code before calling dispatch(x, y, z).
-	auto shader = (Shader *) Shader::current;
+	auto shader = (Shader *) s;
 
 	int tX, tY, tZ;
 	shader->getLocalThreadgroupSize(&tX, &tY, &tZ);
@@ -1356,6 +1376,32 @@ bool Graphics::dispatch(int x, int y, int z)
 	return true;
 }}
 
+bool Graphics::dispatch(love::graphics::Shader *s, love::graphics::Buffer *indirectargs, size_t argsoffset)
+{
+	auto shader = (Shader *) s;
+
+	int tX, tY, tZ;
+	shader->getLocalThreadgroupSize(&tX, &tY, &tZ);
+
+	id<MTLComputePipelineState> pipeline = shader->getComputePipeline();
+	if (pipeline == nil)
+		return false;
+
+	id<MTLComputeCommandEncoder> computeEncoder = useComputeEncoder();
+
+	if (!applyShaderUniforms(computeEncoder, shader))
+		return false;
+
+	// TODO: track this state?
+	[computeEncoder setComputePipelineState:pipeline];
+
+	[computeEncoder dispatchThreadgroupsWithIndirectBuffer:getMTLBuffer(indirectargs)
+									  indirectBufferOffset:argsoffset
+									 threadsPerThreadgroup:MTLSizeMake(tX, tY, tZ)];
+
+	return true;
+}
+
 void Graphics::setRenderTargetsInternal(const RenderTargets &rts, int /*pixelw*/, int /*pixelh*/, bool /*hasSRGBtexture*/)
 { @autoreleasepool {
 	endPass(false);
@@ -2173,7 +2219,13 @@ void Graphics::initCapabilities()
 	capabilities.features[FEATURE_COPY_TEXTURE_TO_BUFFER] = true;
 	capabilities.features[FEATURE_COPY_RENDER_TARGET_TO_BUFFER] = true;
 	capabilities.features[FEATURE_MIPMAP_RANGE] = true;
-	static_assert(FEATURE_MAX_ENUM == 18, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
+
+	if (families.mac[1] || families.macCatalyst[1] || families.apple[3])
+		capabilities.features[FEATURE_INDIRECT_DRAW] = true;
+	else
+		capabilities.features[FEATURE_INDIRECT_DRAW] = false;
+	
+	static_assert(FEATURE_MAX_ENUM == 19, "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/Buffer.cpp

@@ -78,8 +78,10 @@ Buffer::Buffer(love::graphics::Graphics *gfx, const Settings &settings, const st
 		mapUsage = BUFFERUSAGE_VERTEX;
 	else if (usageFlags & BUFFERUSAGEFLAG_INDEX)
 		mapUsage = BUFFERUSAGE_INDEX;
-	else  if (usageFlags & BUFFERUSAGEFLAG_SHADER_STORAGE)
+	else if (usageFlags & BUFFERUSAGEFLAG_SHADER_STORAGE)
 		mapUsage = BUFFERUSAGE_SHADER_STORAGE;
+	else if (usageFlags & BUFFERUSAGEFLAG_INDIRECT_ARGUMENTS)
+		mapUsage = BUFFERUSAGE_INDIRECT_ARGUMENTS;
 
 	target = OpenGL::getGLBufferType(mapUsage);
 

src/modules/graphics/opengl/Graphics.cpp

@@ -505,7 +505,6 @@ void Graphics::setActive(bool enable)
 
 static bool computeDispatchBarriers(Shader *shader, GLbitfield &preDispatchBarriers, GLbitfield &postDispatchBarriers)
 {
-	// TODO: handle indirect argument buffer types, when those are added.
 	for (auto buffer : shader->getActiveWritableStorageBuffers())
 	{
 		if (buffer == nullptr)
@@ -521,6 +520,10 @@ static bool computeDispatchBarriers(Shader *shader, GLbitfield &preDispatchBarri
 			postDispatchBarriers |= GL_SHADER_STORAGE_BARRIER_BIT;
 		}
 
+		// TODO: does this need a pre dispatch barrier too?
+		if (usage & BUFFERUSAGEFLAG_INDIRECT_ARGUMENTS)
+			postDispatchBarriers |= GL_COMMAND_BARRIER_BIT;
+
 		if (usage & BUFFERUSAGEFLAG_TEXEL)
 			postDispatchBarriers |= GL_TEXTURE_FETCH_BARRIER_BIT;
 
@@ -554,10 +557,9 @@ static bool computeDispatchBarriers(Shader *shader, GLbitfield &preDispatchBarri
 	return true;
 }
 
-bool Graphics::dispatch(int x, int y, int z)
+bool Graphics::dispatch(love::graphics::Shader *s, int x, int y, int z)
 {
-	// Set by higher level code before calling dispatch(x, y, z).
-	auto shader = (Shader *) Shader::current;
+	auto shader = (Shader *) s;
 
 	GLbitfield preDispatchBarriers = 0;
 	GLbitfield postDispatchBarriers = 0;
@@ -584,6 +586,33 @@ bool Graphics::dispatch(int x, int y, int z)
 	return true;
 }
 
+bool Graphics::dispatch(love::graphics::Shader *s, love::graphics::Buffer *indirectargs, size_t argsoffset)
+{
+	auto shader = (Shader *) s;
+
+	GLbitfield preDispatchBarriers = 0;
+	GLbitfield postDispatchBarriers = 0;
+
+	if (!computeDispatchBarriers(shader, preDispatchBarriers, postDispatchBarriers))
+		return false;
+
+	if (preDispatchBarriers != 0)
+		glMemoryBarrier(preDispatchBarriers);
+
+	// Note: OpenGL has separate bind points for draw versus dispatch indirect
+	// buffers. Our gl.bindBuffer wrapper uses the draw bind point, so we can't
+	// use it here.
+	glBindBuffer(GL_DISPATCH_INDIRECT_BUFFER, (GLuint)indirectargs->getHandle());
+	glDispatchComputeIndirect(argsoffset);
+
+	// Not as (theoretically) efficient as issuing the barrier right before
+	// they're used later, but much less complicated.
+	if (postDispatchBarriers != 0)
+		glMemoryBarrier(postDispatchBarriers);
+
+	return true;
+}
+
 void Graphics::draw(const DrawCommand &cmd)
 {
 	gl.prepareDraw(this);
@@ -593,7 +622,12 @@ void Graphics::draw(const DrawCommand &cmd)
 
 	GLenum glprimitivetype = OpenGL::getGLPrimitiveType(cmd.primitiveType);
 
-	if (cmd.instanceCount > 1)
+	if (cmd.indirectBuffer != nullptr)
+	{
+		gl.bindBuffer(BUFFERUSAGE_INDIRECT_ARGUMENTS, (GLuint) cmd.indirectBuffer->getHandle());
+		glDrawArraysIndirect(glprimitivetype, BUFFER_OFFSET(cmd.indirectBufferOffset));
+	}
+	else if (cmd.instanceCount > 1)
 		glDrawArraysInstanced(glprimitivetype, cmd.vertexStart, cmd.vertexCount, cmd.instanceCount);
 	else
 		glDrawArrays(glprimitivetype, cmd.vertexStart, cmd.vertexCount);
@@ -614,7 +648,14 @@ void Graphics::draw(const DrawIndexedCommand &cmd)
 
 	gl.bindBuffer(BUFFERUSAGE_INDEX, cmd.indexBuffer->getHandle());
 
-	if (cmd.instanceCount > 1)
+	if (cmd.indirectBuffer != nullptr)
+	{
+		// Note: OpenGL doesn't support indirect indexed draws with a non-zero
+		// index buffer offset.
+		gl.bindBuffer(BUFFERUSAGE_INDIRECT_ARGUMENTS, (GLuint) cmd.indirectBuffer->getHandle());
+		glDrawElementsIndirect(glprimitivetype, gldatatype, BUFFER_OFFSET(cmd.indirectBufferOffset));
+	}
+	else if (cmd.instanceCount > 1)
 		glDrawElementsInstanced(glprimitivetype, cmd.indexCount, gldatatype, gloffset, cmd.instanceCount);
 	else
 		glDrawElements(glprimitivetype, cmd.indexCount, gldatatype, gloffset);
@@ -1662,7 +1703,8 @@ void Graphics::initCapabilities()
 	capabilities.features[FEATURE_COPY_TEXTURE_TO_BUFFER] = gl.isCopyTextureToBufferSupported();
 	capabilities.features[FEATURE_COPY_RENDER_TARGET_TO_BUFFER] = gl.isCopyRenderTargetToBufferSupported();
 	capabilities.features[FEATURE_MIPMAP_RANGE] = GLAD_VERSION_1_2 || GLAD_ES_VERSION_3_0;
-	static_assert(FEATURE_MAX_ENUM == 18, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
+	capabilities.features[FEATURE_INDIRECT_DRAW] = capabilities.features[FEATURE_GLSL4];
+	static_assert(FEATURE_MAX_ENUM == 19, "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/opengl/Graphics.h

@@ -70,7 +70,8 @@ public:
 
 	void setActive(bool active) override;
 
-	bool dispatch(int x, int y, int z) override;
+	bool dispatch(love::graphics::Shader *shader, int x, int y, int z) override;
+	bool dispatch(love::graphics::Shader *shader, love::graphics::Buffer *indirectargs, size_t argsoffset) override;
 
 	void draw(const DrawCommand &cmd) override;
 	void draw(const DrawIndexedCommand &cmd) override;

src/modules/graphics/opengl/OpenGL.cpp

@@ -667,6 +667,7 @@ GLenum OpenGL::getGLBufferType(BufferUsage usage)
 		case BUFFERUSAGE_INDEX: return GL_ELEMENT_ARRAY_BUFFER;
 		case BUFFERUSAGE_TEXEL: return GL_TEXTURE_BUFFER;
 		case BUFFERUSAGE_SHADER_STORAGE: return GL_SHADER_STORAGE_BUFFER;
+		case BUFFERUSAGE_INDIRECT_ARGUMENTS: return GL_DRAW_INDIRECT_BUFFER;
 		case BUFFERUSAGE_MAX_ENUM: return GL_ZERO;
 	}
 
@@ -1491,6 +1492,8 @@ bool OpenGL::isBufferUsageSupported(BufferUsage usage) const
 		return GLAD_VERSION_3_1 || GLAD_ES_VERSION_3_2;
 	case BUFFERUSAGE_SHADER_STORAGE:
 		return (GLAD_VERSION_4_3 && isCoreProfile()) || GLAD_ES_VERSION_3_1;
+	case BUFFERUSAGE_INDIRECT_ARGUMENTS:
+		return (GLAD_VERSION_4_0 && isCoreProfile()) || GLAD_ES_VERSION_3_1;
 	case BUFFERUSAGE_MAX_ENUM:
 		return false;
 	}

src/modules/graphics/vertex.cpp

@@ -367,10 +367,11 @@ const char *getConstant(BuiltinVertexAttribute attrib)
 
 STRINGMAP_BEGIN(BufferUsage, BUFFERUSAGE_MAX_ENUM, bufferUsageName)
 {
-	{ "vertex",        BUFFERUSAGE_VERTEX         },
-	{ "index",         BUFFERUSAGE_INDEX          },
-	{ "texel",         BUFFERUSAGE_TEXEL          },
-	{ "shaderstorage", BUFFERUSAGE_SHADER_STORAGE },
+	{ "vertex",            BUFFERUSAGE_VERTEX             },
+	{ "index",             BUFFERUSAGE_INDEX              },
+	{ "texel",             BUFFERUSAGE_TEXEL              },
+	{ "shaderstorage",     BUFFERUSAGE_SHADER_STORAGE     },
+	{ "indirectarguments", BUFFERUSAGE_INDIRECT_ARGUMENTS },
 }
 STRINGMAP_END(BufferUsage, BUFFERUSAGE_MAX_ENUM, bufferUsageName)
 

src/modules/graphics/vertex.h

@@ -61,6 +61,7 @@ enum BufferUsage
 	BUFFERUSAGE_TEXEL,
 	BUFFERUSAGE_UNIFORM,
 	BUFFERUSAGE_SHADER_STORAGE,
+	BUFFERUSAGE_INDIRECT_ARGUMENTS,
 	BUFFERUSAGE_MAX_ENUM
 };
 
@@ -71,6 +72,7 @@ enum BufferUsageFlags
 	BUFFERUSAGEFLAG_INDEX = 1 << BUFFERUSAGE_INDEX,
 	BUFFERUSAGEFLAG_TEXEL = 1 << BUFFERUSAGE_TEXEL,
 	BUFFERUSAGEFLAG_SHADER_STORAGE = 1 << BUFFERUSAGE_SHADER_STORAGE,
+	BUFFERUSAGEFLAG_INDIRECT_ARGUMENTS = 1 << BUFFERUSAGE_INDIRECT_ARGUMENTS,
 };
 
 enum IndexDataType

src/modules/graphics/vulkan/Buffer.cpp

@@ -37,6 +37,7 @@ static VkBufferUsageFlags getUsageBit(BufferUsage mode)
 	case BUFFERUSAGE_UNIFORM: return VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT;
 	case BUFFERUSAGE_TEXEL: return VK_BUFFER_USAGE_STORAGE_TEXEL_BUFFER_BIT;
 	case BUFFERUSAGE_SHADER_STORAGE: return VK_BUFFER_USAGE_STORAGE_BUFFER_BIT;
+	case BUFFERUSAGE_INDIRECT_ARGUMENTS: return VK_BUFFER_USAGE_INDIRECT_BUFFER_BIT;
 	default:
 		throw love::Exception("unsupported BufferUsage mode");
 	}
@@ -56,10 +57,10 @@ static VkBufferUsageFlags getVulkanUsageFlags(BufferUsageFlags flags)
 
 Buffer::Buffer(love::graphics::Graphics *gfx, const Settings &settings, const std::vector<DataDeclaration> &format, const void *data, size_t size, size_t arraylength)
 	: love::graphics::Buffer(gfx, settings, format, size, arraylength)
-	, usageFlags(settings.usageFlags)
-	, vgfx(dynamic_cast<Graphics*>(gfx))
 	, zeroInitialize(settings.zeroInitialize)
 	, initialData(data)
+	, vgfx(dynamic_cast<Graphics*>(gfx))
+	, usageFlags(settings.usageFlags)
 {
 	loadVolatile();
 }

src/modules/graphics/vulkan/Graphics.cpp

@@ -567,7 +567,8 @@ void Graphics::initCapabilities()
 	capabilities.features[FEATURE_COPY_TEXTURE_TO_BUFFER] = true;
 	capabilities.features[FEATURE_COPY_RENDER_TARGET_TO_BUFFER] = true;
 	capabilities.features[FEATURE_MIPMAP_RANGE] = true;
-	static_assert(FEATURE_MAX_ENUM == 18, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
+	capabilities.features[FEATURE_INDIRECT_DRAW] = true;
+	static_assert(FEATURE_MAX_ENUM == 19, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
 
 	VkPhysicalDeviceProperties properties;
 	vkGetPhysicalDeviceProperties(physicalDevice, &properties);
@@ -689,12 +690,25 @@ void Graphics::draw(const DrawCommand &cmd)
 {
 	prepareDraw(*cmd.attributes, *cmd.buffers, cmd.texture, cmd.primitiveType, cmd.cullMode);
 
-	vkCmdDraw(
-		commandBuffers.at(currentFrame),
-		static_cast<uint32_t>(cmd.vertexCount),
-		static_cast<uint32_t>(cmd.instanceCount),
-		static_cast<uint32_t>(cmd.vertexStart),
-		0);
+	if (cmd.indirectBuffer != nullptr)
+	{
+		vkCmdDrawIndirect(
+			commandBuffers.at(currentFrame),
+			(VkBuffer) cmd.indirectBuffer->getHandle(),
+			cmd.indirectBufferOffset,
+			1,
+			0);
+	}
+	else
+	{
+		vkCmdDraw(
+			commandBuffers.at(currentFrame),
+			(uint32) cmd.vertexCount,
+			(uint32) cmd.instanceCount,
+			(uint32) cmd.vertexStart,
+			0);
+	}
+
 	drawCalls++;
 }
 
@@ -704,16 +718,30 @@ void Graphics::draw(const DrawIndexedCommand &cmd)
 
 	vkCmdBindIndexBuffer(
 		commandBuffers.at(currentFrame),
-		(VkBuffer)cmd.indexBuffer->getHandle(),
-		static_cast<VkDeviceSize>(cmd.indexBufferOffset),
+		(VkBuffer) cmd.indexBuffer->getHandle(),
+		(VkDeviceSize) cmd.indexBufferOffset,
 		Vulkan::getVulkanIndexBufferType(cmd.indexType));
-	vkCmdDrawIndexed(
-		commandBuffers.at(currentFrame),
-		static_cast<uint32_t>(cmd.indexCount),
-		static_cast<uint32_t>(cmd.instanceCount),
-		0,
-		0,
-		0);
+
+	if (cmd.indirectBuffer != nullptr)
+	{
+		vkCmdDrawIndexedIndirect(
+			commandBuffers.at(currentFrame),
+			(VkBuffer) cmd.indirectBuffer->getHandle(),
+			cmd.indirectBufferOffset,
+			1,
+			0);
+	}
+	else
+	{
+		vkCmdDrawIndexed(
+			commandBuffers.at(currentFrame),
+			(uint32) cmd.indexCount,
+			(uint32) cmd.instanceCount,
+			0,
+			0,
+			0);
+	}
+
 	drawCalls++;
 }
 
@@ -1005,7 +1033,22 @@ graphics::StreamBuffer *Graphics::newStreamBuffer(BufferUsage type, size_t size)
 	return new StreamBuffer(this, type, size);
 }
 
-bool Graphics::dispatch(int x, int y, int z)
+bool Graphics::dispatch(love::graphics::Shader *shader, int x, int y, int z)
+{
+	if (renderPassState.active)
+		endRenderPass();
+
+	vkCmdBindPipeline(commandBuffers.at(currentFrame), VK_PIPELINE_BIND_POINT_COMPUTE, computeShader->getComputePipeline());
+
+	computeShader->cmdPushDescriptorSets(commandBuffers.at(currentFrame), VK_PIPELINE_BIND_POINT_COMPUTE);
+
+	// TODO: does this need any layout transitions?
+	vkCmdDispatch(commandBuffers.at(currentFrame), (uint32) x, (uint32) y, (uint32) z);
+
+	return true;
+}
+
+bool Graphics::dispatch(love::graphics::Shader *shader, love::graphics::Buffer *indirectargs, size_t argsoffset)
 {
 	if (renderPassState.active)
 		endRenderPass();
@@ -1014,7 +1057,8 @@ bool Graphics::dispatch(int x, int y, int z)
 
 	computeShader->cmdPushDescriptorSets(commandBuffers.at(currentFrame), VK_PIPELINE_BIND_POINT_COMPUTE);
 
-	vkCmdDispatch(commandBuffers.at(currentFrame), static_cast<uint32_t>(x), static_cast<uint32_t>(y), static_cast<uint32_t>(z));
+	// TODO: does this need any layout transitions?
+	vkCmdDispatchIndirect(commandBuffers.at(currentFrame), (VkBuffer) indirectargs->getHandle(), argsoffset);
 
 	return true;
 }

src/modules/graphics/vulkan/Graphics.h

@@ -304,7 +304,8 @@ protected:
 	graphics::ShaderStage *newShaderStageInternal(ShaderStageType stage, const std::string &cachekey, const std::string &source, bool gles) override;
 	graphics::Shader *newShaderInternal(StrongRef<love::graphics::ShaderStage> stages[SHADERSTAGE_MAX_ENUM]) override;
 	graphics::StreamBuffer *newStreamBuffer(BufferUsage type, size_t size) override;
-	bool dispatch(int x, int y, int z) override;
+	bool dispatch(love::graphics::Shader *shader, int x, int y, int z) override;
+	bool dispatch(love::graphics::Shader *shader, love::graphics::Buffer *indirectargs, size_t argsoffset) override;
 	void initCapabilities() override;
 	void getAPIStats(int &shaderswitches) const override;
 	void setRenderTargetsInternal(const RenderTargets &rts, int pixelw, int pixelh, bool hasSRGBtexture) override;

src/modules/graphics/wrap_Graphics.cpp

@@ -3099,6 +3099,20 @@ int w_drawInstanced(lua_State *L)
 	return 0;
 }
 
+int w_drawIndirect(lua_State *L)
+{
+	Mesh *t = luax_checkmesh(L, 1);
+	Buffer *argsbuffer = luax_checkbuffer(L, 2);
+	int argsindex = (int) luaL_checkinteger(L, 3) - 1;
+
+	luax_checkstandardtransform(L, 4, [&](const Matrix4 &m)
+	{
+		luax_catchexcept(L, [&]() { instance()->drawIndirect(t, m, argsbuffer, argsindex); });
+	});
+
+	return 0;
+}
+
 int w_drawFromShader(lua_State *L)
 {
 	if (luax_istype(L, 1, Buffer::type))
@@ -3135,6 +3149,40 @@ int w_drawFromShader(lua_State *L)
 	return 0;
 }
 
+int w_drawFromShaderIndirect(lua_State *L)
+{
+	if (luax_istype(L, 1, Buffer::type))
+	{
+		// Indexed drawing.
+		Buffer *t = luax_checkbuffer(L, 1);
+		Buffer *argsbuffer = luax_checkbuffer(L, 2);
+		int argsindex = (int) luaL_optinteger(L, 3, 1) - 1;
+
+		Texture *tex = nullptr;
+		if (!lua_isnoneornil(L, 4))
+			tex = luax_checktexture(L, 4);
+
+		luax_catchexcept(L, [&]() { instance()->drawFromShaderIndirect(t, argsbuffer, argsindex, tex); });
+	}
+	else
+	{
+		const char *primstr = luaL_checkstring(L, 1);
+		PrimitiveType primtype = PRIMITIVE_TRIANGLES;
+		if (!getConstant(primstr, primtype))
+			return luax_enumerror(L, "primitive type", getConstants(primtype), primstr);
+
+		Buffer *argsbuffer = luax_checkbuffer(L, 2);
+		int argsindex = (int) luaL_optinteger(L, 3, 1) - 1;
+
+		Texture *tex = nullptr;
+		if (!lua_isnoneornil(L, 4))
+			tex = luax_checktexture(L, 4);
+
+		luax_catchexcept(L, [&]() { instance()->drawFromShaderIndirect(primtype, argsbuffer, argsindex, tex); });
+	}
+	return 0;
+}
+
 int w_print(lua_State *L)
 {
 	std::vector<love::font::ColoredString> str;
@@ -3529,6 +3577,15 @@ int w_dispatchThreadgroups(lua_State* L)
 	return 0;
 }
 
+int w_dispatchIndirect(lua_State *L)
+{
+	Shader *shader = luax_checkshader(L, 1);
+	Buffer *argsbuffer = luax_checkbuffer(L, 2);
+	int argsindex = (int) luaL_optinteger(L, 3, 1) - 1;
+	luax_catchexcept(L, [&]() { instance()->dispatchIndirect(shader, argsbuffer, argsindex); });
+	return 0;
+}
+
 int w_copyBuffer(lua_State *L)
 {
 	Buffer *source = luax_checkbuffer(L, 1);
@@ -3844,12 +3901,15 @@ static const luaL_Reg functions[] =
 	{ "draw", w_draw },
 	{ "drawLayer", w_drawLayer },
 	{ "drawInstanced", w_drawInstanced },
+	{ "drawIndirect", w_drawIndirect },
 	{ "drawFromShader", w_drawFromShader },
+	{ "drawFromShaderIndirect", w_drawFromShaderIndirect },
 
 	{ "print", w_print },
 	{ "printf", w_printf },
 
 	{ "dispatchThreadgroups", w_dispatchThreadgroups },
+	{ "dispatchIndirect", w_dispatchIndirect },
 
 	{ "copyBuffer", w_copyBuffer },
 	{ "copyBufferToTexture", w_copyBufferToTexture },