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/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 85733e0be..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,29 +1882,34 @@ 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::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)
-		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 +1926,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);
 
@@ -1899,6 +1967,61 @@ void Graphics::drawShaderVertices(Buffer *indexbuffer, int indexcount, int insta
 	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();
@@ -2532,6 +2655,8 @@ 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         },
+	{ "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 328665ac9..456e5d28f 100644
--- a/src/modules/graphics/Graphics.h
+++ b/src/modules/graphics/Graphics.h
@@ -164,6 +164,8 @@ public:
 		FEATURE_COPY_BUFFER_TO_TEXTURE,
 		FEATURE_COPY_TEXTURE_TO_BUFFER,
 		FEATURE_COPY_RENDER_TARGET_TO_BUFFER,
+		FEATURE_MIPMAP_RANGE,
+		FEATURE_INDIRECT_DRAW,
 		FEATURE_MAX_ENUM
 	};
 
@@ -198,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];
@@ -235,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?
@@ -260,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?
@@ -688,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 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);
+	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
@@ -872,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)
 	{
@@ -1000,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/Shader.cpp b/src/modules/graphics/Shader.cpp
index b91a9c694..f6c314420 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"
@@ -135,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
@@ -425,6 +430,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 +545,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 +576,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 +595,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;
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.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 acbbe3d3b..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);
@@ -2172,7 +2218,14 @@ 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;
+
+	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 b37bcdca2..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);
@@ -1661,7 +1702,9 @@ 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;
+	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 3bd3ed0d8..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;
 	}
 
@@ -1388,6 +1389,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]);
@@ -1485,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/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/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 7166b22f9..ac865262a 100644
--- a/src/modules/graphics/vulkan/Graphics.cpp
+++ b/src/modules/graphics/vulkan/Graphics.cpp
@@ -567,7 +567,9 @@ 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;
+	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 +691,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 +719,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 +1034,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)
 {
 	usedShadersInFrame.insert(computeShader);
 
@@ -1016,7 +1045,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 3606cf842..458f22c4d 100644
--- a/src/modules/graphics/vulkan/Graphics.h
+++ b/src/modules/graphics/vulkan/Graphics.h
@@ -303,7 +303,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 d50c19d24..313c4f0e8 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);
 
@@ -3094,7 +3099,21 @@ int w_drawInstanced(lua_State *L)
 	return 0;
 }
 
-int w_drawShaderVertices(lua_State *L)
+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))
 	{
@@ -3109,7 +3128,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
 	{
@@ -3125,7 +3144,41 @@ 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;
+}
+
+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;
 }
@@ -3524,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);
@@ -3839,12 +3901,15 @@ static const luaL_Reg functions[] =
 	{ "draw", w_draw },
 	{ "drawLayer", w_drawLayer },
 	{ "drawInstanced", w_drawInstanced },
-	{ "drawShaderVertices", w_drawShaderVertices },
+	{ "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 },
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 },