From 8706f0ef6a414fddaf6a2494383f7da7820b7bec Mon Sep 17 00:00:00 2001
From: Sasha Szpakowski <slime73@gmail.com>
Date: Sun, 29 Jan 2023 15:58:31 -0400
Subject: [PATCH 1/6] Rename Texture:isRenderTarget to Texture:isCanvas.

---
 src/modules/graphics/wrap_Texture.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/src/modules/graphics/wrap_Texture.cpp b/src/modules/graphics/wrap_Texture.cpp
index 3e09da10e..00f6058af 100644
--- a/src/modules/graphics/wrap_Texture.cpp
+++ b/src/modules/graphics/wrap_Texture.cpp
@@ -279,7 +279,7 @@ int w_Texture_getFormat(lua_State *L)
 	return 1;
 }
 
-int w_Texture_isRenderTarget(lua_State *L)
+int w_Texture_isCanvas(lua_State *L)
 {
 	Texture *t = luax_checktexture(L, 1);
 	luax_pushboolean(L, t->isRenderTarget());
@@ -497,7 +497,7 @@ const luaL_Reg w_Texture_functions[] =
 	{ "setWrap", w_Texture_setWrap },
 	{ "getWrap", w_Texture_getWrap },
 	{ "getFormat", w_Texture_getFormat },
-	{ "isRenderTarget", w_Texture_isRenderTarget },
+	{ "isCanvas", w_Texture_isCanvas },
 	{ "isComputeWritable", w_Texture_isComputeWritable },
 	{ "isReadable", w_Texture_isReadable },
 	{ "getMipmapMode", w_Texture_getMipmapMode },

From defe811f6266ac20221a0ce9f6f39932fb625ade Mon Sep 17 00:00:00 2001
From: Sasha Szpakowski <slime73@gmail.com>
Date: Sun, 29 Jan 2023 16:02:39 -0400
Subject: [PATCH 2/6] Allow creating a mipmapped texture with less than the
 full mip range.

When creating a texture with user-specified mipmaps, it no longer errors if all mipmaps down to 1x1 aren't present.

Added a 'mipmapcount' field to the settings table in newImage and friends (only used when mipmaps are enabled for the texture.) If it's set, only mipmap levels up to the specified count will be created instead of always creating the full range.

Old OpenGL ES 2 devices don't support loading a mipmapped texture without the full mipmap range, so there's a new 'mipmaprange' boolean field in the table returned by love.graphics.getSupported.
---
 src/modules/graphics/Graphics.cpp        |  1 +
 src/modules/graphics/Graphics.h          |  1 +
 src/modules/graphics/Texture.cpp         | 31 ++++++++++++++++++------
 src/modules/graphics/Texture.h           |  2 ++
 src/modules/graphics/metal/Graphics.mm   |  3 ++-
 src/modules/graphics/opengl/Graphics.cpp |  3 ++-
 src/modules/graphics/opengl/OpenGL.cpp   |  6 +++++
 src/modules/graphics/opengl/Texture.cpp  |  4 +++
 src/modules/graphics/vulkan/Graphics.cpp |  3 ++-
 src/modules/graphics/wrap_Graphics.cpp   |  5 ++++
 10 files changed, 48 insertions(+), 11 deletions(-)

diff --git a/src/modules/graphics/Graphics.cpp b/src/modules/graphics/Graphics.cpp
index 85733e0be..3d688001f 100644
--- a/src/modules/graphics/Graphics.cpp
+++ b/src/modules/graphics/Graphics.cpp
@@ -2532,6 +2532,7 @@ STRINGMAP_CLASS_BEGIN(Graphics, Graphics::Feature, Graphics::FEATURE_MAX_ENUM, f
 	{ "copybuffertotexture",      Graphics::FEATURE_COPY_BUFFER_TO_TEXTURE },
 	{ "copytexturetobuffer",      Graphics::FEATURE_COPY_TEXTURE_TO_BUFFER },
 	{ "copyrendertargettobuffer", Graphics::FEATURE_COPY_RENDER_TARGET_TO_BUFFER },
+	{ "mipmaprange",              Graphics::FEATURE_MIPMAP_RANGE         },
 }
 STRINGMAP_CLASS_END(Graphics, Graphics::Feature, Graphics::FEATURE_MAX_ENUM, feature)
 
diff --git a/src/modules/graphics/Graphics.h b/src/modules/graphics/Graphics.h
index 328665ac9..df9917051 100644
--- a/src/modules/graphics/Graphics.h
+++ b/src/modules/graphics/Graphics.h
@@ -164,6 +164,7 @@ public:
 		FEATURE_COPY_BUFFER_TO_TEXTURE,
 		FEATURE_COPY_TEXTURE_TO_BUFFER,
 		FEATURE_COPY_RENDER_TARGET_TO_BUFFER,
+		FEATURE_MIPMAP_RANGE,
 		FEATURE_MAX_ENUM
 	};
 
diff --git a/src/modules/graphics/Texture.cpp b/src/modules/graphics/Texture.cpp
index 3cdd379ff..e9b7ab1f7 100644
--- a/src/modules/graphics/Texture.cpp
+++ b/src/modules/graphics/Texture.cpp
@@ -180,6 +180,9 @@ Texture::Texture(Graphics *gfx, const Settings &settings, const Slices *slices)
 	, samplerState()
 	, graphicsMemorySize(0)
 {
+	const auto &caps = gfx->getCapabilities();
+	int requestedMipmapCount = settings.mipmapCount;
+
 	if (slices != nullptr && slices->getMipmapCount() > 0 && slices->getSliceCount() > 0)
 	{
 		texType = slices->getTextureType();
@@ -189,8 +192,15 @@ Texture::Texture(Graphics *gfx, const Settings &settings, const Slices *slices)
 
 		int dataMipmaps = 1;
 		if (slices->validate() && slices->getMipmapCount() > 1)
+		{
 			dataMipmaps = slices->getMipmapCount();
 
+			if (requestedMipmapCount > 0)
+				requestedMipmapCount = std::min(requestedMipmapCount, dataMipmaps);
+			else
+				requestedMipmapCount = dataMipmaps;
+		}
+
 		love::image::ImageDataBase *slice = slices->get(0, 0);
 
 		format = slice->getFormat();
@@ -231,7 +241,17 @@ Texture::Texture(Graphics *gfx, const Settings &settings, const Slices *slices)
 		mipmapsMode = MIPMAPS_MANUAL;
 
 	if (mipmapsMode != MIPMAPS_NONE)
-		mipmapCount = getTotalMipmapCount(pixelWidth, pixelHeight, depth);
+	{
+		int totalMipmapCount = getTotalMipmapCount(pixelWidth, pixelHeight, depth);
+
+		if (requestedMipmapCount > 0)
+			mipmapCount = std::min(totalMipmapCount, requestedMipmapCount);
+		else
+			mipmapCount = totalMipmapCount;
+
+		if (mipmapCount != totalMipmapCount && !caps.features[Graphics::FEATURE_MIPMAP_RANGE])
+			throw love::Exception("Custom mipmap ranges for a texture are not supported on this system (%d mipmap levels are required but only %d levels were provided.)", totalMipmapCount, mipmapCount);
+	}
 
 	const char *miperr = nullptr;
 	if (mipmapsMode == MIPMAPS_AUTO && !supportsGenerateMipmaps(miperr))
@@ -288,7 +308,7 @@ Texture::Texture(Graphics *gfx, const Settings &settings, const Slices *slices)
 		throw love::Exception("The %s%s pixel format is not supported%s on this system.", fstr, readablestr, rtstr);
 	}
 
-	if (!gfx->getCapabilities().textureTypes[texType])
+	if (!caps.textureTypes[texType])
 	{
 		const char *textypestr = "unknown";
 		Texture::getConstant(texType, textypestr);
@@ -864,14 +884,8 @@ bool Texture::Slices::validate() const
 
 	int w = firstdata->getWidth();
 	int h = firstdata->getHeight();
-	int depth = textureType == TEXTURE_VOLUME ? slicecount : 1;
 	PixelFormat format = firstdata->getFormat();
 
-	int expectedmips = Texture::getTotalMipmapCount(w, h, depth);
-
-	if (mipcount != expectedmips && mipcount != 1)
-		throw love::Exception("Texture does not have all required mipmap levels (expected %d, got %d)", expectedmips, mipcount);
-
 	if (textureType == TEXTURE_CUBE && w != h)
 		throw love::Exception("Cube textures must have equal widths and heights for each cube face.");
 
@@ -947,6 +961,7 @@ static StringMap<Texture::SettingType, Texture::SETTING_MAX_ENUM>::Entry setting
 	{ "height",       Texture::SETTING_HEIGHT        },
 	{ "layers",       Texture::SETTING_LAYERS        },
 	{ "mipmaps",      Texture::SETTING_MIPMAPS       },
+	{ "mipmapcount",  Texture::SETTING_MIPMAP_COUNT  },
 	{ "format",       Texture::SETTING_FORMAT        },
 	{ "linear",       Texture::SETTING_LINEAR        },
 	{ "type",         Texture::SETTING_TYPE          },
diff --git a/src/modules/graphics/Texture.h b/src/modules/graphics/Texture.h
index b9da6a44d..2132bf741 100644
--- a/src/modules/graphics/Texture.h
+++ b/src/modules/graphics/Texture.h
@@ -166,6 +166,7 @@ public:
 		SETTING_HEIGHT,
 		SETTING_LAYERS,
 		SETTING_MIPMAPS,
+		SETTING_MIPMAP_COUNT,
 		SETTING_FORMAT,
 		SETTING_LINEAR,
 		SETTING_TYPE,
@@ -185,6 +186,7 @@ public:
 		int layers = 1; // depth for 3D textures
 		TextureType type = TEXTURE_2D;
 		MipmapsMode mipmaps = MIPMAPS_NONE;
+		int mipmapCount = 0; // only used when > 0.
 		PixelFormat format = PIXELFORMAT_NORMAL;
 		bool linear = false;
 		float dpiScale = 1.0f;
diff --git a/src/modules/graphics/metal/Graphics.mm b/src/modules/graphics/metal/Graphics.mm
index acbbe3d3b..f6440affa 100644
--- a/src/modules/graphics/metal/Graphics.mm
+++ b/src/modules/graphics/metal/Graphics.mm
@@ -2172,7 +2172,8 @@ void Graphics::initCapabilities()
 	capabilities.features[FEATURE_COPY_BUFFER_TO_TEXTURE] = true;
 	capabilities.features[FEATURE_COPY_TEXTURE_TO_BUFFER] = true;
 	capabilities.features[FEATURE_COPY_RENDER_TARGET_TO_BUFFER] = true;
-	static_assert(FEATURE_MAX_ENUM == 17, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
+	capabilities.features[FEATURE_MIPMAP_RANGE] = true;
+	static_assert(FEATURE_MAX_ENUM == 18, "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;
diff --git a/src/modules/graphics/opengl/Graphics.cpp b/src/modules/graphics/opengl/Graphics.cpp
index b37bcdca2..33b25a887 100644
--- a/src/modules/graphics/opengl/Graphics.cpp
+++ b/src/modules/graphics/opengl/Graphics.cpp
@@ -1661,7 +1661,8 @@ void Graphics::initCapabilities()
 	capabilities.features[FEATURE_COPY_BUFFER_TO_TEXTURE] = gl.isCopyBufferToTextureSupported();
 	capabilities.features[FEATURE_COPY_TEXTURE_TO_BUFFER] = gl.isCopyTextureToBufferSupported();
 	capabilities.features[FEATURE_COPY_RENDER_TARGET_TO_BUFFER] = gl.isCopyRenderTargetToBufferSupported();
-	static_assert(FEATURE_MAX_ENUM == 17, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
+	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.limits[LIMIT_POINT_SIZE] = gl.getMaxPointSize();
 	capabilities.limits[LIMIT_TEXTURE_SIZE] = gl.getMax2DTextureSize();
diff --git a/src/modules/graphics/opengl/OpenGL.cpp b/src/modules/graphics/opengl/OpenGL.cpp
index 3bd3ed0d8..e96a15765 100644
--- a/src/modules/graphics/opengl/OpenGL.cpp
+++ b/src/modules/graphics/opengl/OpenGL.cpp
@@ -1388,6 +1388,12 @@ bool OpenGL::rawTexStorage(TextureType target, int levels, PixelFormat pixelform
 	GLenum gltarget = getGLTextureType(target);
 	TextureFormat fmt = convertPixelFormat(pixelformat, false, isSRGB);
 
+	// This shouldn't be needed for glTexStorage, but some drivers don't follow
+	// the spec apparently.
+	// https://stackoverflow.com/questions/13859061/does-an-immutable-texture-need-a-gl-texture-max-level
+	if (GLAD_VERSION_1_2 || GLAD_ES_VERSION_3_0)
+		glTexParameteri(gltarget, GL_TEXTURE_MAX_LEVEL, levels - 1);
+
 	if (fmt.swizzled)
 	{
 		glTexParameteri(gltarget, GL_TEXTURE_SWIZZLE_R, fmt.swizzle[0]);
diff --git a/src/modules/graphics/opengl/Texture.cpp b/src/modules/graphics/opengl/Texture.cpp
index c8eef069b..70fdbd344 100644
--- a/src/modules/graphics/opengl/Texture.cpp
+++ b/src/modules/graphics/opengl/Texture.cpp
@@ -264,6 +264,10 @@ void Texture::createTexture()
 	if (!isCompressed())
 		gl.rawTexStorage(texType, mipcount, format, sRGB, pixelWidth, pixelHeight, texType == TEXTURE_VOLUME ? depth : layers);
 
+	// rawTexStorage handles this for uncompressed textures.
+	if (isCompressed() && (GLAD_VERSION_1_1 || GLAD_ES_VERSION_3_0))
+		glTexParameteri(gltype, GL_TEXTURE_MAX_LEVEL, mipcount - 1);
+
 	int w = pixelWidth;
 	int h = pixelHeight;
 	int d = depth;
diff --git a/src/modules/graphics/vulkan/Graphics.cpp b/src/modules/graphics/vulkan/Graphics.cpp
index 768d4056e..05b1d2e23 100644
--- a/src/modules/graphics/vulkan/Graphics.cpp
+++ b/src/modules/graphics/vulkan/Graphics.cpp
@@ -566,7 +566,8 @@ void Graphics::initCapabilities()
 	capabilities.features[FEATURE_COPY_BUFFER_TO_TEXTURE] = true;
 	capabilities.features[FEATURE_COPY_TEXTURE_TO_BUFFER] = true;
 	capabilities.features[FEATURE_COPY_RENDER_TARGET_TO_BUFFER] = true;
-	static_assert(FEATURE_MAX_ENUM == 17, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
+	capabilities.features[FEATURE_MIPMAP_RANGE] = true;
+	static_assert(FEATURE_MAX_ENUM == 18, "Graphics::initCapabilities must be updated when adding a new graphics feature!");
 
 	VkPhysicalDeviceProperties properties;
 	vkGetPhysicalDeviceProperties(physicalDevice, &properties);
diff --git a/src/modules/graphics/wrap_Graphics.cpp b/src/modules/graphics/wrap_Graphics.cpp
index d50c19d24..bec8225b7 100644
--- a/src/modules/graphics/wrap_Graphics.cpp
+++ b/src/modules/graphics/wrap_Graphics.cpp
@@ -781,6 +781,11 @@ static void luax_checktexturesettings(lua_State *L, int idx, bool opt, bool chec
 	}
 	lua_pop(L, 1);
 
+	lua_getfield(L, idx, Texture::getConstant(Texture::SETTING_MIPMAP_COUNT));
+	if (!lua_isnoneornil(L, -1))
+		s.mipmapCount = (int) luaL_checkinteger(L, -1);
+	lua_pop(L, 1);
+
 	s.linear = luax_boolflag(L, idx, Texture::getConstant(Texture::SETTING_LINEAR), s.linear);
 	s.msaa = luax_intflag(L, idx, Texture::getConstant(Texture::SETTING_MSAA), s.msaa);
 

From e13b0078a20f4ea67fe4fca722e659b1566da0c6 Mon Sep 17 00:00:00 2001
From: Sasha Szpakowski <slime73@gmail.com>
Date: Sat, 4 Feb 2023 10:25:13 -0400
Subject: [PATCH 3/6] shaders: ignore code in comments when parsing entry point
 functions.

Fixes #1581.
---
 src/modules/graphics/Shader.cpp | 141 ++++++++++++++++++++++++++++----
 1 file changed, 123 insertions(+), 18 deletions(-)

diff --git a/src/modules/graphics/Shader.cpp b/src/modules/graphics/Shader.cpp
index b91a9c694..eea1e4a95 100644
--- a/src/modules/graphics/Shader.cpp
+++ b/src/modules/graphics/Shader.cpp
@@ -22,6 +22,7 @@
 #include "Shader.h"
 #include "Graphics.h"
 #include "math/MathModule.h"
+#include "common/Range.h"
 
 // glslang
 #include "libraries/glslang/glslang/Public/ShaderLang.h"
@@ -425,6 +426,111 @@ static const Version versions[] =
 	{ "#version 430 core", "#version 320 es" },
 };
 
+enum CommentType
+{
+	COMMENT_NONE,
+	COMMENT_LINE,
+	COMMENT_BLOCK,
+};
+
+static void parseComments(const std::string &src, std::vector<Range> &comments)
+{
+	CommentType commenttype = COMMENT_NONE;
+	Range comment;
+
+	const char *srcbytes = src.data();
+	size_t len = src.length();
+
+	for (size_t i = 0; i < len; i++)
+	{
+		char curchar = srcbytes[i];
+
+		if (commenttype == COMMENT_NONE)
+		{
+			if (curchar == '/' && i + 1 < len)
+			{
+				char nextchar = srcbytes[i + 1];
+				if (nextchar == '/')
+				{
+					commenttype = COMMENT_LINE;
+					comment = Range(i, 1);
+				}
+				else if (nextchar == '*')
+				{
+					commenttype = COMMENT_BLOCK;
+					comment = Range(i, 1);
+				}
+			}
+		}
+		else if (commenttype == COMMENT_LINE)
+		{
+			if (curchar == '\n')
+			{
+				commenttype = COMMENT_NONE;
+				comment.last = i;
+				comments.push_back(comment);
+			}
+		}
+		else if (commenttype == COMMENT_BLOCK)
+		{
+			if (curchar == '/' && i > 0 && srcbytes[i - 1] == '*')
+			{
+				commenttype = COMMENT_NONE;
+				comment.last = i;
+				comments.push_back(comment);
+			}
+		}
+	}
+
+	if (commenttype == COMMENT_LINE)
+	{
+		comment.last = len - 1;
+		comments.push_back(comment);
+	}
+}
+
+static bool inComment(size_t i, const std::vector<Range> &comments)
+{
+	Range r(i, 1);
+
+	for (const Range &comment : comments)
+	{
+		if (comment.contains(r))
+			return true;
+	}
+
+	return false;
+}
+
+static bool textSearch(const std::string &src, const std::string &str, const std::vector<Range> &comments)
+{
+	size_t start = 0;
+	size_t found = std::string::npos;
+
+	while ((found = src.find(str, start)) != std::string::npos)
+	{
+		if (!inComment(found, comments))
+			return true;
+		start = found + str.size();
+	}
+
+	return false;
+}
+
+static bool regexSearch(const std::string &src, const std::string &rstr, const std::vector<Range> &comments)
+{
+	std::regex r(rstr);
+
+	for (auto it = std::sregex_iterator(src.begin(), src.end(), r); it != std::sregex_iterator(); it++)
+	{
+		const std::smatch &m = *it;
+		if (!inComment(m.position(), comments))
+			return true;
+	}
+
+	return false;
+}
+
 static Shader::Language getTargetLanguage(const std::string &src)
 {
 	std::regex r("^\\s*#pragma language (\\w+)");
@@ -435,33 +541,30 @@ static Shader::Language getTargetLanguage(const std::string &src)
 	return lang;
 }
 
-static Shader::EntryPoint getVertexEntryPoint(const std::string &src)
+static Shader::EntryPoint getVertexEntryPoint(const std::string &src, const std::vector<Range> &comments)
 {
-	std::smatch m;
-
-	if (std::regex_search(src, m, std::regex("void\\s+vertexmain\\s*\\(")))
+	if (regexSearch(src, "void\\s+vertexmain\\s*\\(", comments))
 		return Shader::ENTRYPOINT_RAW;
 
-	if (std::regex_search(src, m, std::regex("vec4\\s+position\\s*\\(")))
+	if (regexSearch(src, "vec4\\s+position\\s*\\(", comments))
 		return Shader::ENTRYPOINT_HIGHLEVEL;
 
 	return Shader::ENTRYPOINT_NONE;
 }
 
-static Shader::EntryPoint getPixelEntryPoint(const std::string &src, bool &mrt)
+static Shader::EntryPoint getPixelEntryPoint(const std::string &src, const std::vector<Range> &comments, bool &mrt)
 {
 	mrt = false;
-	std::smatch m;
 
-	if (std::regex_search(src, m, std::regex("void\\s+pixelmain\\s*\\(")))
+	if (regexSearch(src, "void\\s+pixelmain\\s*\\(", comments))
 		return Shader::ENTRYPOINT_RAW;
 
-	if (std::regex_search(src, m, std::regex("vec4\\s+effect\\s*\\(")))
+	if (regexSearch(src, "vec4\\s+effect\\s*\\(", comments))
 		return Shader::ENTRYPOINT_HIGHLEVEL;
 
-	if (std::regex_search(src, m, std::regex("void\\s+effect\\s*\\(")))
+	if (regexSearch(src, "void\\s+effect\\s*\\(", comments))
 	{
-		if (src.find("love_RenderTargets") != std::string::npos || src.find("love_Canvases") != std::string::npos)
+		if (textSearch(src, "love_RenderTargets", comments) || textSearch(src, "love_Canvases", comments))
 			mrt = true;
 		return Shader::ENTRYPOINT_CUSTOM;
 	}
@@ -469,10 +572,9 @@ static Shader::EntryPoint getPixelEntryPoint(const std::string &src, bool &mrt)
 	return Shader::ENTRYPOINT_NONE;
 }
 
-static Shader::EntryPoint getComputeEntryPoint(const std::string &src) {
-	std::smatch m;
-
-	if (std::regex_search(src, m, std::regex("void\\s+computemain\\s*\\(")))
+static Shader::EntryPoint getComputeEntryPoint(const std::string &src, const std::vector<Range> &comments)
+{
+	if (regexSearch(src, "void\\s+computemain\\s*\\(", comments))
 		return Shader::ENTRYPOINT_RAW;
 
 	return Shader::ENTRYPOINT_NONE;
@@ -489,11 +591,14 @@ Shader *Shader::standardShaders[Shader::STANDARD_MAX_ENUM] = {nullptr};
 
 Shader::SourceInfo Shader::getSourceInfo(const std::string &src)
 {
+	std::vector<Range> comments;
+	glsl::parseComments(src, comments);
+
 	SourceInfo info = {};
 	info.language = glsl::getTargetLanguage(src);
-	info.stages[SHADERSTAGE_VERTEX] = glsl::getVertexEntryPoint(src);
-	info.stages[SHADERSTAGE_PIXEL] = glsl::getPixelEntryPoint(src, info.usesMRT);
-	info.stages[SHADERSTAGE_COMPUTE] = glsl::getComputeEntryPoint(src);
+	info.stages[SHADERSTAGE_VERTEX] = glsl::getVertexEntryPoint(src, comments);
+	info.stages[SHADERSTAGE_PIXEL] = glsl::getPixelEntryPoint(src, comments, info.usesMRT);
+	info.stages[SHADERSTAGE_COMPUTE] = glsl::getComputeEntryPoint(src, comments);
 	if (info.stages[SHADERSTAGE_COMPUTE])
 		info.language = LANGUAGE_GLSL4;
 	return info;

From d3747cebb1e876c2d9dbcaafcf2fa41da5934103 Mon Sep 17 00:00:00 2001
From: Sasha Szpakowski <slime73@gmail.com>
Date: Sun, 5 Feb 2023 15:16:10 -0400
Subject: [PATCH 4/6] Rename love.graphics.drawShaderVertices to
 drawFromShader.

---
 changes.txt                            |  2 +-
 src/modules/graphics/Graphics.cpp      | 26 +++++++++++++-------------
 src/modules/graphics/Graphics.h        |  4 ++--
 src/modules/graphics/wrap_Graphics.cpp |  8 ++++----
 4 files changed, 20 insertions(+), 20 deletions(-)

diff --git a/changes.txt b/changes.txt
index 1ca00130b..00a5e6060 100644
--- a/changes.txt
+++ b/changes.txt
@@ -55,7 +55,7 @@ Released: N/A
 * Added Compute Shader support via new 'computemain' shader entry point.
 * Added love.graphics.dispatchThreadgroups for running compute shaders.
 * Added Shader:hasStage.
-* Added love.graphics.drawShaderVertices.
+* Added love.graphics.drawFromShader.
 * Added love.graphics.getQuadIndexBuffer.
 * Added variants of love.graphics.applyTransform and replaceTransform which accept x,y,angle,sx,sy,ox,oy parameters.
 * Added APIs to override the default orthographic projection: love.graphics.setOrthoProjection, setPerspectiveProjection, and resetProjection.
diff --git a/src/modules/graphics/Graphics.cpp b/src/modules/graphics/Graphics.cpp
index 3d688001f..318333106 100644
--- a/src/modules/graphics/Graphics.cpp
+++ b/src/modules/graphics/Graphics.cpp
@@ -1819,29 +1819,29 @@ void Graphics::drawInstanced(Mesh *mesh, const Matrix4 &m, int instancecount)
 	mesh->drawInstanced(this, m, instancecount);
 }
 
-void Graphics::drawShaderVertices(PrimitiveType primtype, int vertexcount, int instancecount, Texture *maintexture)
+void Graphics::drawFromShader(PrimitiveType primtype, int vertexcount, int instancecount, Texture *maintexture)
 {
 	if (primtype == PRIMITIVE_TRIANGLE_FAN && vertexcount > LOVE_UINT16_MAX)
-		throw love::Exception("drawShaderVertices cannot draw more than %d vertices when the 'fan' draw mode is used.", LOVE_UINT16_MAX);
+		throw love::Exception("drawFromShader cannot draw more than %d vertices when the 'fan' draw mode is used.", LOVE_UINT16_MAX);
 
 	// Emulated triangle fan via an index buffer.
 	if (primtype == PRIMITIVE_TRIANGLE_FAN && getFanIndexBuffer())
 	{
 		int indexcount = getIndexCount(TRIANGLEINDEX_FAN, vertexcount);
-		drawShaderVertices(getFanIndexBuffer(), indexcount, instancecount, 0, maintexture);
+		drawFromShader(getFanIndexBuffer(), indexcount, instancecount, 0, maintexture);
 		return;
 	}
 
 	flushBatchedDraws();
 
 	if (!capabilities.features[FEATURE_GLSL3])
-		throw love::Exception("drawShaderVertices is not supported on this system (GLSL3 support is required.)");
+		throw love::Exception("drawFromShader is not supported on this system (GLSL3 support is required.)");
 
 	if (Shader::isDefaultActive() || !Shader::current)
-		throw love::Exception("drawShaderVertices can only be used with a custom shader.");
+		throw love::Exception("drawFromShader can only be used with a custom shader.");
 
 	if (vertexcount < 0 || instancecount < 0)
-		throw love::Exception("drawShaderVertices vertex and instance count parameters must not be negative.");
+		throw love::Exception("drawFromShader vertex and instance count parameters must not be negative.");
 
 	Shader::current->validateDrawState(primtype, maintexture);
 
@@ -1858,27 +1858,27 @@ void Graphics::drawShaderVertices(PrimitiveType primtype, int vertexcount, int i
 	draw(cmd);
 }
 
-void Graphics::drawShaderVertices(Buffer *indexbuffer, int indexcount, int instancecount, int startindex, Texture *maintexture)
+void Graphics::drawFromShader(Buffer *indexbuffer, int indexcount, int instancecount, int startindex, Texture *maintexture)
 {
 	flushBatchedDraws();
 
 	if (!capabilities.features[FEATURE_GLSL3])
-		throw love::Exception("drawShaderVertices is not supported on this system (GLSL3 support is required.)");
+		throw love::Exception("drawFromShader is not supported on this system (GLSL3 support is required.)");
 
 	if (!(indexbuffer->getUsageFlags() & BUFFERUSAGEFLAG_INDEX))
-		throw love::Exception("The buffer passed to drawShaderVertices must be an index buffer.");
+		throw love::Exception("The buffer passed to drawFromShader must be an index buffer.");
 
 	if (startindex < 0)
-		throw love::Exception("drawShaderVertices startindex parameter must not be negative.");
+		throw love::Exception("drawFromShader startindex parameter must not be negative.");
 
 	if (indexcount < 0 || instancecount < 0)
-		throw love::Exception("drawShaderVertices index and instance count parameters must not be negative.");
+		throw love::Exception("drawFromShader index and instance count parameters must not be negative.");
 
 	if ((size_t)(startindex + indexcount) > indexbuffer->getArrayLength() * indexbuffer->getDataMembers().size())
-		throw love::Exception("drawShaderVertices startindex and index count parameters do not fit in the given index buffer.");
+		throw love::Exception("drawFromShader startindex and index count parameters do not fit in the given index buffer.");
 
 	if (Shader::isDefaultActive() || !Shader::current)
-		throw love::Exception("drawShaderVertices can only be used with a custom shader.");
+		throw love::Exception("drawFromShader can only be used with a custom shader.");
 
 	Shader::current->validateDrawState(PRIMITIVE_TRIANGLES, maintexture);
 
diff --git a/src/modules/graphics/Graphics.h b/src/modules/graphics/Graphics.h
index df9917051..622996c2e 100644
--- a/src/modules/graphics/Graphics.h
+++ b/src/modules/graphics/Graphics.h
@@ -697,8 +697,8 @@ public:
 	void drawLayer(Texture *texture, int layer, Quad *quad, const Matrix4 &m);
 	void drawInstanced(Mesh *mesh, const Matrix4 &m, int instancecount);
 
-	void drawShaderVertices(PrimitiveType primtype, int vertexcount, int instancecount, Texture *maintexture);
-	void drawShaderVertices(Buffer *indexbuffer, int indexcount, int instancecount, int startindex, Texture *maintexture);
+	void drawFromShader(PrimitiveType primtype, int vertexcount, int instancecount, Texture *maintexture);
+	void drawFromShader(Buffer *indexbuffer, int indexcount, int instancecount, int startindex, Texture *maintexture);
 
 	/**
 	 * Draws text at the specified coordinates
diff --git a/src/modules/graphics/wrap_Graphics.cpp b/src/modules/graphics/wrap_Graphics.cpp
index bec8225b7..496ea2205 100644
--- a/src/modules/graphics/wrap_Graphics.cpp
+++ b/src/modules/graphics/wrap_Graphics.cpp
@@ -3099,7 +3099,7 @@ int w_drawInstanced(lua_State *L)
 	return 0;
 }
 
-int w_drawShaderVertices(lua_State *L)
+int w_drawFromShader(lua_State *L)
 {
 	if (luax_istype(L, 1, Buffer::type))
 	{
@@ -3114,7 +3114,7 @@ int w_drawShaderVertices(lua_State *L)
 		if (!lua_isnoneornil(L, 5))
 			tex = luax_checktexture(L, 5);
 
-		luax_catchexcept(L, [&]() { instance()->drawShaderVertices(t, indexcount, instancecount, indexstart, tex); });
+		luax_catchexcept(L, [&]() { instance()->drawFromShader(t, indexcount, instancecount, indexstart, tex); });
 	}
 	else
 	{
@@ -3130,7 +3130,7 @@ int w_drawShaderVertices(lua_State *L)
 		if (!lua_isnoneornil(L, 4))
 			tex = luax_checktexture(L, 4);
 
-		luax_catchexcept(L, [&]() { instance()->drawShaderVertices(primtype, vertexcount, instancecount, tex); });
+		luax_catchexcept(L, [&]() { instance()->drawFromShader(primtype, vertexcount, instancecount, tex); });
 	}
 	return 0;
 }
@@ -3844,7 +3844,7 @@ static const luaL_Reg functions[] =
 	{ "draw", w_draw },
 	{ "drawLayer", w_drawLayer },
 	{ "drawInstanced", w_drawInstanced },
-	{ "drawShaderVertices", w_drawShaderVertices },
+	{ "drawFromShader", w_drawFromShader },
 
 	{ "print", w_print },
 	{ "printf", w_printf },

From 2178b634c80f8184c6ce6888eb0cd878b185a392 Mon Sep 17 00:00:00 2001
From: Sasha Szpakowski <slime73@gmail.com>
Date: Sun, 5 Feb 2023 15:55:19 -0400
Subject: [PATCH 5/6] 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          |  20 +++-
 src/modules/graphics/Graphics.cpp        | 128 ++++++++++++++++++++++-
 src/modules/graphics/Graphics.h          |  25 ++++-
 src/modules/graphics/Mesh.cpp            |  35 ++++++-
 src/modules/graphics/Mesh.h              |   3 +
 src/modules/graphics/metal/Graphics.h    |   3 +-
 src/modules/graphics/metal/Graphics.mm   |  80 +++++++++++---
 src/modules/graphics/opengl/Buffer.cpp   |   4 +-
 src/modules/graphics/opengl/Graphics.cpp |  56 ++++++++--
 src/modules/graphics/opengl/Graphics.h   |   3 +-
 src/modules/graphics/opengl/OpenGL.cpp   |   3 +
 src/modules/graphics/vertex.cpp          |   9 +-
 src/modules/graphics/vertex.h            |   2 +
 src/modules/graphics/vulkan/Buffer.cpp   |   5 +-
 src/modules/graphics/vulkan/Graphics.cpp |  80 ++++++++++----
 src/modules/graphics/vulkan/Graphics.h   |   3 +-
 src/modules/graphics/wrap_Graphics.cpp   |  60 +++++++++++
 17 files changed, 460 insertions(+), 59 deletions(-)

diff --git a/src/modules/graphics/Buffer.cpp b/src/modules/graphics/Buffer.cpp
index 84d3256bd..3ebbc61d1 100644
--- a/src/modules/graphics/Buffer.cpp
+++ b/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;
 
diff --git a/src/modules/graphics/Graphics.cpp b/src/modules/graphics/Graphics.cpp
index 318333106..cfd915a49 100644
--- a/src/modules/graphics/Graphics.cpp
+++ b/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)
 
diff --git a/src/modules/graphics/Graphics.h b/src/modules/graphics/Graphics.h
index 622996c2e..456e5d28f 100644
--- a/src/modules/graphics/Graphics.h
+++ b/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;
 
diff --git a/src/modules/graphics/Mesh.cpp b/src/modules/graphics/Mesh.cpp
index 6acfa51ef..78abd2d96 100644
--- a/src/modules/graphics/Mesh.cpp
+++ b/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);
 	}
diff --git a/src/modules/graphics/Mesh.h b/src/modules/graphics/Mesh.h
index 0ca9ecdaa..b164b7cfb 100644
--- a/src/modules/graphics/Mesh.h
+++ b/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;
diff --git a/src/modules/graphics/metal/Graphics.h b/src/modules/graphics/metal/Graphics.h
index 9b7f30a04..a0568f2f0 100644
--- a/src/modules/graphics/metal/Graphics.h
+++ b/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;
diff --git a/src/modules/graphics/metal/Graphics.mm b/src/modules/graphics/metal/Graphics.mm
index f6440affa..4c9ba328f 100644
--- a/src/modules/graphics/metal/Graphics.mm
+++ b/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;
diff --git a/src/modules/graphics/opengl/Buffer.cpp b/src/modules/graphics/opengl/Buffer.cpp
index e4258324c..913eaa816 100644
--- a/src/modules/graphics/opengl/Buffer.cpp
+++ b/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);
 
diff --git a/src/modules/graphics/opengl/Graphics.cpp b/src/modules/graphics/opengl/Graphics.cpp
index 33b25a887..664d26fa7 100644
--- a/src/modules/graphics/opengl/Graphics.cpp
+++ b/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();
diff --git a/src/modules/graphics/opengl/Graphics.h b/src/modules/graphics/opengl/Graphics.h
index f22137dad..997239c15 100644
--- a/src/modules/graphics/opengl/Graphics.h
+++ b/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;
diff --git a/src/modules/graphics/opengl/OpenGL.cpp b/src/modules/graphics/opengl/OpenGL.cpp
index e96a15765..b169a06a9 100644
--- a/src/modules/graphics/opengl/OpenGL.cpp
+++ b/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;
 	}
diff --git a/src/modules/graphics/vertex.cpp b/src/modules/graphics/vertex.cpp
index 3082bf41c..d916ebe12 100644
--- a/src/modules/graphics/vertex.cpp
+++ b/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)
 
diff --git a/src/modules/graphics/vertex.h b/src/modules/graphics/vertex.h
index f9bf41774..6578548a8 100644
--- a/src/modules/graphics/vertex.h
+++ b/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
diff --git a/src/modules/graphics/vulkan/Buffer.cpp b/src/modules/graphics/vulkan/Buffer.cpp
index db109b233..5eaf03f6d 100644
--- a/src/modules/graphics/vulkan/Buffer.cpp
+++ b/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();
 }
diff --git a/src/modules/graphics/vulkan/Graphics.cpp b/src/modules/graphics/vulkan/Graphics.cpp
index 05b1d2e23..c9fe696d1 100644
--- a/src/modules/graphics/vulkan/Graphics.cpp
+++ b/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,7 @@ 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();
@@ -1014,7 +1042,23 @@ 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?
+	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();
+
+	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?
+	vkCmdDispatchIndirect(commandBuffers.at(currentFrame), (VkBuffer) indirectargs->getHandle(), argsoffset);
 
 	return true;
 }
diff --git a/src/modules/graphics/vulkan/Graphics.h b/src/modules/graphics/vulkan/Graphics.h
index 7d88e58f6..d699fbbc8 100644
--- a/src/modules/graphics/vulkan/Graphics.h
+++ b/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;
diff --git a/src/modules/graphics/wrap_Graphics.cpp b/src/modules/graphics/wrap_Graphics.cpp
index 496ea2205..313c4f0e8 100644
--- a/src/modules/graphics/wrap_Graphics.cpp
+++ b/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 },

From 6ac7c09d9083b7fe1261a3c566dcc4a7af061fc1 Mon Sep 17 00:00:00 2001
From: Sasha Szpakowski <slime73@gmail.com>
Date: Sun, 5 Feb 2023 22:16:28 -0400
Subject: [PATCH 6/6] shaders: default to std430 layout for storage buffers.

---
 src/modules/graphics/Shader.cpp | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/src/modules/graphics/Shader.cpp b/src/modules/graphics/Shader.cpp
index eea1e4a95..f6c314420 100644
--- a/src/modules/graphics/Shader.cpp
+++ b/src/modules/graphics/Shader.cpp
@@ -136,6 +136,10 @@ static const char global_functions[] = R"(
 	#endif
 #endif
 
+#if __VERSION__ >= 430 || (defined(GL_ES) && __VERSION__ >= 310)
+	layout (std430) buffer;
+#endif
+
 #if __VERSION__ >= 130 && !defined(LOVE_GLSL1_ON_GLSL3)
 	#define Texel texture
 #else