diff --git a/Source/engine/render/blit_impl.hpp b/Source/engine/render/blit_impl.hpp
index f352b5a19fd..8bd76897c11 100644
--- a/Source/engine/render/blit_impl.hpp
+++ b/Source/engine/render/blit_impl.hpp
@@ -8,13 +8,9 @@
 #include "engine/palette.h"
 #include "utils/attributes.h"
 
-namespace devilution {
+#define DEVILUTIONX_BLIT_EXECUTION_POLICY DVL_EXECUTION_UNSEQ
 
-#if __cpp_lib_execution >= 201902L
-#define DEVILUTIONX_BLIT_EXECUTION_POLICY std::execution::unseq,
-#else
-#define DEVILUTIONX_BLIT_EXECUTION_POLICY
-#endif
+namespace devilution {
 
 DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void BlitFillDirect(uint8_t *dst, unsigned length, uint8_t color)
 {
diff --git a/Source/engine/render/dun_render.cpp b/Source/engine/render/dun_render.cpp
index 2fda02f407d..b2f9599f467 100644
--- a/Source/engine/render/dun_render.cpp
+++ b/Source/engine/render/dun_render.cpp
@@ -14,14 +14,12 @@
 
 #include <SDL_endian.h>
 
-#include <algorithm>
 #include <climits>
 #include <cstdint>
+#include <cstring>
 
 #include "engine/render/blit_impl.hpp"
 #include "levels/dun_tile.hpp"
-#include "lighting.h"
-#include "options.h"
 #include "utils/attributes.h"
 #ifdef DEBUG_STR
 #include "engine/render/text_render.hpp"
@@ -54,13 +52,6 @@ constexpr int_fast16_t TriangleHeight = LowerHeight + TriangleUpperHeight;
 /** For triangles, for each pixel drawn vertically, this many pixels are drawn horizontally. */
 constexpr int_fast16_t XStep = 2;
 
-int_fast16_t GetTileHeight(TileType tile)
-{
-	if (tile == TileType::LeftTriangle || tile == TileType::RightTriangle)
-		return TriangleHeight;
-	return Height;
-}
-
 #ifdef DEBUG_STR
 std::pair<std::string_view, UiFlags> GetTileDebugStr(TileType tile)
 {
@@ -261,37 +252,6 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLine(uint8_t *DVL_RESTRICT dst, c
 	}
 }
 
-struct Clip {
-	int_fast16_t top;
-	int_fast16_t bottom;
-	int_fast16_t left;
-	int_fast16_t right;
-	int_fast16_t width;
-	int_fast16_t height;
-};
-
-DVL_ALWAYS_INLINE Clip CalculateClip(int_fast16_t x, int_fast16_t y, int_fast16_t w, int_fast16_t h, const Surface &out)
-{
-	Clip clip;
-	clip.top = y + 1 < h ? h - (y + 1) : 0;
-	clip.bottom = y + 1 > out.h() ? (y + 1) - out.h() : 0;
-	clip.left = x < 0 ? -x : 0;
-	clip.right = x + w > out.w() ? x + w - out.w() : 0;
-	clip.width = w - clip.left - clip.right;
-	clip.height = h - clip.top - clip.bottom;
-	return clip;
-}
-
-DVL_ALWAYS_INLINE bool IsFullyDark(const uint8_t *DVL_RESTRICT tbl)
-{
-	return tbl == FullyDarkLightTable;
-}
-
-DVL_ALWAYS_INLINE bool IsFullyLit(const uint8_t *DVL_RESTRICT tbl)
-{
-	return tbl == FullyLitLightTable;
-}
-
 template <LightType Light, bool Transparent>
 DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderSquareFull(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl)
 {
@@ -302,7 +262,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderSquareFull(uint8_t *DVL_RESTRICT
 }
 
 template <LightType Light, bool Transparent>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderSquareClipped(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderSquareClipped(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	src += clip.bottom * Height + clip.left;
 	for (auto i = 0; i < clip.height; ++i, dst -= dstPitch) {
@@ -344,7 +304,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTransparentSquareFull(uint8_t *DV
 
 template <LightType Light, bool OpaquePrefix, int8_t PrefixIncrement>
 // NOLINTNEXTLINE(readability-function-cognitive-complexity): Actually complex and has to be fast.
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTransparentSquareClipped(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTransparentSquareClipped(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const auto skipRestOfTheLine = [&src](int_fast16_t remainingWidth) {
 		while (remainingWidth > 0) {
@@ -443,7 +403,7 @@ struct DiamondClipY {
 };
 
 template <int_fast16_t UpperHeight = TriangleUpperHeight>
-DVL_ALWAYS_INLINE DiamondClipY CalculateDiamondClipY(const Clip &clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT DiamondClipY CalculateDiamondClipY(const DunTileClip &clip)
 {
 	DiamondClipY result;
 	if (clip.bottom > LowerHeight) {
@@ -462,12 +422,12 @@ DVL_ALWAYS_INLINE DiamondClipY CalculateDiamondClipY(const Clip &clip)
 	return result;
 }
 
-DVL_ALWAYS_INLINE std::size_t CalculateTriangleSourceSkipLowerBottom(int_fast16_t numLines)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT std::size_t CalculateTriangleSourceSkipLowerBottom(int_fast16_t numLines)
 {
 	return XStep * numLines * (numLines + 1) / 2 + 2 * ((numLines + 1) / 2);
 }
 
-DVL_ALWAYS_INLINE std::size_t CalculateTriangleSourceSkipUpperBottom(int_fast16_t numLines)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT std::size_t CalculateTriangleSourceSkipUpperBottom(int_fast16_t numLines)
 {
 	return 2 * TriangleUpperHeight * numLines - numLines * (numLines - 1) + 2 * ((numLines + 1) / 2);
 }
@@ -544,7 +504,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTriangleFull(uint8_t *DVL_RES
 }
 
 template <LightType Light, bool Transparent>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTriangleClipVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTriangleClipVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const DiamondClipY clipY = CalculateDiamondClipY(clip);
 	RenderLeftTriangleLowerClipVertical<Light, Transparent>(clipY, dst, dstPitch, src, tbl);
@@ -560,7 +520,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTriangleClipVertical(uint8_t
 }
 
 template <LightType Light, bool Transparent>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTriangleClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTriangleClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const DiamondClipY clipY = CalculateDiamondClipY(clip);
 	const int_fast16_t clipLeft = clip.left;
@@ -579,7 +539,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTriangleClipLeftAndVertical(u
 }
 
 template <LightType Light, bool Transparent>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTriangleClipRightAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTriangleClipRightAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const DiamondClipY clipY = CalculateDiamondClipY(clip);
 	const int_fast16_t clipRight = clip.right;
@@ -674,7 +634,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTriangleFull(uint8_t *DVL_RE
 }
 
 template <LightType Light, bool Transparent>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTriangleClipVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTriangleClipVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const DiamondClipY clipY = CalculateDiamondClipY(clip);
 	RenderRightTriangleLowerClipVertical<Light, Transparent>(clipY, dst, dstPitch, src, tbl);
@@ -688,7 +648,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTriangleClipVertical(uint8_t
 }
 
 template <LightType Light, bool Transparent>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTriangleClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTriangleClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const DiamondClipY clipY = CalculateDiamondClipY(clip);
 	const int_fast16_t clipLeft = clip.left;
@@ -705,7 +665,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTriangleClipLeftAndVertical(
 }
 
 template <LightType Light, bool Transparent>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTriangleClipRightAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTriangleClipRightAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const DiamondClipY clipY = CalculateDiamondClipY(clip);
 	const int_fast16_t clipRight = clip.right;
@@ -769,7 +729,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTrapezoidUpperHalf(uint8_t *DVL_R
 }
 
 template <LightType Light, bool OpaquePrefix, int8_t PrefixIncrement>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTrapezoidUpperHalfClipVertical(const Clip &clip, const DiamondClipY &clipY, uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl) {
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTrapezoidUpperHalfClipVertical(const DunTileClip &clip, const DiamondClipY &clipY, uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl) {
 	const auto upperMax = TrapezoidUpperHeight - clipY.upperTop;
 	int8_t prefix = InitPrefix<PrefixIncrement>(clip.bottom);
 	for (auto i = 1 + clipY.upperBottom; i <= upperMax; ++i, dst -= dstPitch) {
@@ -780,7 +740,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTrapezoidUpperHalfClipVertical(co
 }
 
 template <LightType Light, bool OpaquePrefix, int8_t PrefixIncrement>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTrapezoidUpperHalfClipLeftAndVertical(const Clip &clip, const DiamondClipY &clipY, uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl) {
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTrapezoidUpperHalfClipLeftAndVertical(const DunTileClip &clip, const DiamondClipY &clipY, uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl) {
 	const auto upperMax = TrapezoidUpperHeight - clipY.upperTop;
 	int8_t prefix = InitPrefix<PrefixIncrement>(clip.bottom);
 	for (auto i = 1 + clipY.upperBottom; i <= upperMax; ++i, dst -= dstPitch) {
@@ -791,7 +751,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTrapezoidUpperHalfClipLeftAndVert
 }
 
 template <LightType Light, bool OpaquePrefix, int8_t PrefixIncrement>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTrapezoidUpperHalfClipRightAndVertical(const Clip &clip, const DiamondClipY &clipY, uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl) {
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTrapezoidUpperHalfClipRightAndVertical(const DunTileClip &clip, const DiamondClipY &clipY, uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl) {
 	const auto upperMax = TrapezoidUpperHeight - clipY.upperTop;
 	int8_t prefix = InitPrefix<PrefixIncrement>(clip.bottom);
 	for (auto i = 1 + clipY.upperBottom; i <= upperMax; ++i, dst -= dstPitch) {
@@ -810,7 +770,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTrapezoidFull(uint8_t *DVL_RE
 }
 
 template <LightType Light, bool OpaquePrefix, int8_t PrefixIncrement>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTrapezoidClipVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTrapezoidClipVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const DiamondClipY clipY = CalculateDiamondClipY<TrapezoidUpperHeight>(clip);
 	RenderLeftTriangleLowerClipVertical<Light, LowerHalfTransparent<OpaquePrefix, PrefixIncrement>>(clipY, dst, dstPitch, src, tbl);
@@ -820,7 +780,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTrapezoidClipVertical(uint8_t
 }
 
 template <LightType Light, bool OpaquePrefix, int8_t PrefixIncrement>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTrapezoidClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTrapezoidClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const DiamondClipY clipY = CalculateDiamondClipY<TrapezoidUpperHeight>(clip);
 	RenderLeftTriangleLowerClipLeftAndVertical<Light, LowerHalfTransparent<OpaquePrefix, PrefixIncrement>>(clip.left, clipY, dst, dstPitch, src, tbl);
@@ -830,7 +790,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTrapezoidClipLeftAndVertical(
 }
 
 template <LightType Light, bool OpaquePrefix, int8_t PrefixIncrement>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTrapezoidClipRightAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderLeftTrapezoidClipRightAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const DiamondClipY clipY = CalculateDiamondClipY<TrapezoidUpperHeight>(clip);
 	RenderLeftTriangleLowerClipRightAndVertical<Light, LowerHalfTransparent<OpaquePrefix, PrefixIncrement>>(clip.right, clipY, dst, dstPitch, src, tbl);
@@ -863,7 +823,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTrapezoidFull(uint8_t *DVL_R
 }
 
 template <LightType Light, bool OpaquePrefix, int8_t PrefixIncrement>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTrapezoidClipVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTrapezoidClipVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const DiamondClipY clipY = CalculateDiamondClipY<TrapezoidUpperHeight>(clip);
 	RenderRightTriangleLowerClipVertical<Light, LowerHalfTransparent<OpaquePrefix, PrefixIncrement>>(clipY, dst, dstPitch, src, tbl);
@@ -872,7 +832,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTrapezoidClipVertical(uint8_
 }
 
 template <LightType Light, bool OpaquePrefix, int8_t PrefixIncrement>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTrapezoidClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTrapezoidClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const DiamondClipY clipY = CalculateDiamondClipY<TrapezoidUpperHeight>(clip);
 	RenderRightTriangleLowerClipLeftAndVertical<Light, LowerHalfTransparent<OpaquePrefix, PrefixIncrement>>(clip.left, clipY, dst, dstPitch, src, tbl);
@@ -881,7 +841,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTrapezoidClipLeftAndVertical
 }
 
 template <LightType Light, bool OpaquePrefix, int8_t PrefixIncrement>
-DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTrapezoidClipRightAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, Clip clip)
+DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderRightTrapezoidClipRightAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, const uint8_t *DVL_RESTRICT src, const uint8_t *DVL_RESTRICT tbl, DunTileClip clip)
 {
 	const DiamondClipY clipY = CalculateDiamondClipY<TrapezoidUpperHeight>(clip);
 	RenderRightTriangleLowerClipRightAndVertical<Light, LowerHalfTransparent<OpaquePrefix, PrefixIncrement>>(clip.right, clipY, dst, dstPitch, src, tbl);
@@ -1009,7 +969,7 @@ DVL_ALWAYS_INLINE DVL_ATTRIBUTE_HOT void RenderTileDispatch(TileType tile, uint8
 }
 
 // Blit with left and vertical clipping.
-void RenderBlackTileClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, int sx, DiamondClipY clipY)
+void RenderSingleColorTileClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, int sx, DiamondClipY clipY, uint8_t color)
 {
 	dst += XStep * (LowerHeight - clipY.lowerBottom - 1);
 	// Lower triangle (drawn bottom to top):
@@ -1018,9 +978,9 @@ void RenderBlackTileClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstP
 		const auto w = 2 * XStep * i;
 		const auto curX = sx + TILE_WIDTH / 2 - XStep * i;
 		if (curX >= 0) {
-			memset(dst, 0, w);
+			memset(dst, color, w);
 		} else if (-curX <= w) {
-			memset(dst - curX, 0, w + curX);
+			memset(dst - curX, color, w + curX);
 		}
 	}
 	dst += 2 * XStep + XStep * clipY.upperBottom;
@@ -1030,9 +990,9 @@ void RenderBlackTileClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstP
 		const auto w = 2 * XStep * (TriangleUpperHeight - i);
 		const auto curX = sx + TILE_WIDTH / 2 - XStep * (TriangleUpperHeight - i);
 		if (curX >= 0) {
-			memset(dst, 0, w);
+			memset(dst, color, w);
 		} else if (-curX <= w) {
-			memset(dst - curX, 0, w + curX);
+			memset(dst - curX, color, w + curX);
 		} else {
 			break;
 		}
@@ -1040,7 +1000,7 @@ void RenderBlackTileClipLeftAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstP
 }
 
 // Blit with right and vertical clipping.
-void RenderBlackTileClipRightAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, int_fast16_t maxWidth, DiamondClipY clipY)
+void RenderSingleColorTileClipRightAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, int_fast16_t maxWidth, DiamondClipY clipY, uint8_t color)
 {
 	dst += XStep * (LowerHeight - clipY.lowerBottom - 1);
 	// Lower triangle (drawn bottom to top):
@@ -1050,7 +1010,7 @@ void RenderBlackTileClipRightAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dst
 		const auto endX = TILE_WIDTH / 2 + XStep * i;
 		const auto skip = endX > maxWidth ? endX - maxWidth : 0;
 		if (width > skip)
-			memset(dst, 0, width - skip);
+			memset(dst, color, width - skip);
 	}
 	dst += 2 * XStep + XStep * clipY.upperBottom;
 	// Upper triangle (drawn bottom to top):
@@ -1066,35 +1026,35 @@ void RenderBlackTileClipRightAndVertical(uint8_t *DVL_RESTRICT dst, uint16_t dst
 }
 
 // Blit with vertical clipping only.
-void RenderBlackTileClipY(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, DiamondClipY clipY)
+void RenderSingleColorTileClipY(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, DiamondClipY clipY, uint8_t color)
 {
 	dst += XStep * (LowerHeight - clipY.lowerBottom - 1);
 	// Lower triangle (drawn bottom to top):
 	const auto lowerMax = LowerHeight - clipY.lowerTop;
 	for (auto i = 1 + clipY.lowerBottom; i <= lowerMax; ++i, dst -= dstPitch + XStep) {
-		memset(dst, 0, 2 * XStep * i);
+		memset(dst, color, 2 * XStep * i);
 	}
 	dst += 2 * XStep + XStep * clipY.upperBottom;
 	// Upper triangle (drawn bottom to top):
 	const auto upperMax = TriangleUpperHeight - clipY.upperTop;
 	for (auto i = 1 + clipY.upperBottom; i <= upperMax; ++i, dst -= dstPitch - XStep) {
-		memset(dst, 0, TILE_WIDTH - 2 * XStep * i);
+		memset(dst, color, TILE_WIDTH - 2 * XStep * i);
 	}
 }
 
-// Blit a black tile without clipping (must be fully in bounds).
-void RenderBlackTileFull(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch)
+// Blit a single color tile without clipping (must be fully in bounds).
+void RenderSingleColorTileFull(uint8_t *DVL_RESTRICT dst, uint16_t dstPitch, uint8_t color)
 {
 	dst += XStep * (LowerHeight - 1);
 	// Tile is fully in bounds, can use constant loop boundaries.
 	// Lower triangle (drawn bottom to top):
 	for (unsigned i = 1; i <= LowerHeight; ++i, dst -= dstPitch + XStep) {
-		memset(dst, 0, 2 * XStep * i);
+		memset(dst, color, 2 * XStep * i);
 	}
 	dst += 2 * XStep;
 	// Upper triangle (drawn bottom to to top):
 	for (unsigned i = 1; i <= TriangleUpperHeight; ++i, dst -= dstPitch - XStep) {
-		memset(dst, 0, TILE_WIDTH - 2 * XStep * i);
+		memset(dst, color, TILE_WIDTH - 2 * XStep * i);
 	}
 }
 
@@ -1134,11 +1094,9 @@ std::string_view MaskTypeToString(MaskType maskType)
 }
 #endif
 
-void RenderTile(const Surface &out, Point position,
-    LevelCelBlock levelCelBlock, MaskType maskType, const uint8_t *tbl)
+void RenderTile(uint8_t *dst, uint16_t dstPitch,
+    TileType tile, const uint8_t *src, MaskType maskType, const uint8_t *tbl, const DunTileClip &clip)
 {
-	const TileType tile = levelCelBlock.type();
-
 #ifdef DEBUG_RENDER_OFFSET_X
 	position.x += DEBUG_RENDER_OFFSET_X;
 #endif
@@ -1149,15 +1107,6 @@ void RenderTile(const Surface &out, Point position,
 	DBGCOLOR = GetTileDebugColor(tile);
 #endif
 
-	const Clip clip = CalculateClip(position.x, position.y, Width, GetTileHeight(tile), out);
-	if (clip.width <= 0 || clip.height <= 0)
-		return;
-
-	const auto *pFrameTable = reinterpret_cast<const uint32_t *>(pDungeonCels.get());
-	const auto *src = reinterpret_cast<const uint8_t *>(&pDungeonCels[SDL_SwapLE32(pFrameTable[levelCelBlock.frame()])]);
-	uint8_t *dst = out.at(static_cast<int>(position.x + clip.left), static_cast<int>(position.y - clip.bottom));
-	const uint16_t dstPitch = out.pitch();
-
 #ifdef DUN_RENDER_STATS
 	++DunRenderStats[DunRenderType { tile, maskType }];
 #endif
@@ -1189,7 +1138,7 @@ void RenderTile(const Surface &out, Point position,
 #endif
 }
 
-void world_draw_black_tile(const Surface &out, int sx, int sy)
+void RenderSingleColorTile(const Surface &out, int sx, int sy, uint8_t color)
 {
 #ifdef DEBUG_RENDER_OFFSET_X
 	sx += DEBUG_RENDER_OFFSET_X;
@@ -1197,25 +1146,63 @@ void world_draw_black_tile(const Surface &out, int sx, int sy)
 #ifdef DEBUG_RENDER_OFFSET_Y
 	sy += DEBUG_RENDER_OFFSET_Y;
 #endif
-	auto clip = CalculateClip(sx, sy, TILE_WIDTH, TriangleHeight, out);
-	if (clip.width <= 0 || clip.height <= 0)
-		return;
+	auto clip = CalculateDunTileClip(sx, sy, TILE_WIDTH, TriangleHeight, out);
+	if (clip.width <= 0 || clip.height <= 0) return;
 
 	auto clipY = CalculateDiamondClipY(clip);
 	uint8_t *dst = out.at(sx, static_cast<int>(sy - clip.bottom));
 	if (clip.width == TILE_WIDTH) {
 		if (clip.height == TriangleHeight) {
-			RenderBlackTileFull(dst, out.pitch());
+			RenderSingleColorTileFull(dst, out.pitch(), color);
 		} else {
-			RenderBlackTileClipY(dst, out.pitch(), clipY);
+			RenderSingleColorTileClipY(dst, out.pitch(), clipY, color);
 		}
 	} else {
 		if (clip.right == 0) {
-			RenderBlackTileClipLeftAndVertical(dst, out.pitch(), sx, clipY);
+			RenderSingleColorTileClipLeftAndVertical(dst, out.pitch(), sx, clipY, color);
 		} else {
-			RenderBlackTileClipRightAndVertical(dst, out.pitch(), clip.width, clipY);
+			RenderSingleColorTileClipRightAndVertical(dst, out.pitch(), clip.width, clipY, color);
 		}
 	}
 }
 
+void DunTileApplyTrans(LevelCelBlock levelCelBlock, uint8_t *dst, const uint8_t *tbl)
+{
+	const TileType tile = levelCelBlock.type();
+	const uint8_t *src = GetDungeonTileSrc(levelCelBlock.frame());
+
+	size_t size;
+	switch (tile) {
+	case TileType::TransparentSquare:
+		for (size_t i = 0; i < Height; ++i) {
+			uint_fast8_t drawWidth = Width;
+			while (drawWidth > 0) {
+				auto v = static_cast<int8_t>(*src++);
+				*dst++ = v;
+				if (v > 0) {
+					BlitPixelsWithMap(dst, src, v, tbl);
+					src += v;
+					dst += v;
+				} else {
+					v = static_cast<int8_t>(-v);
+				}
+				drawWidth -= v;
+			}
+		}
+		return;
+	case TileType::LeftTriangle:
+	case TileType::RightTriangle:
+		size = 544;
+		break;
+	case TileType::LeftTrapezoid:
+	case TileType::RightTrapezoid:
+		size = 800;
+		break;
+	case TileType::Square:
+		size = Width * Height;
+		break;
+	}
+	BlitPixelsWithMap(dst, src, size, tbl);
+}
+
 } // namespace devilution
diff --git a/Source/engine/render/dun_render.hpp b/Source/engine/render/dun_render.hpp
index 684c8322094..825450aa2e8 100644
--- a/Source/engine/render/dun_render.hpp
+++ b/Source/engine/render/dun_render.hpp
@@ -12,6 +12,7 @@
 #include "engine/point.hpp"
 #include "engine/surface.hpp"
 #include "levels/dun_tile.hpp"
+#include "lighting.h"
 
 // #define DUN_RENDER_STATS
 #ifdef DUN_RENDER_STATS
@@ -163,6 +164,41 @@ std::string_view TileTypeToString(TileType tileType);
 std::string_view MaskTypeToString(MaskType maskType);
 #endif
 
+struct DunTileClip {
+	int_fast16_t top;
+	int_fast16_t bottom;
+	int_fast16_t left;
+	int_fast16_t right;
+	int_fast16_t width;
+	int_fast16_t height;
+};
+
+void RenderTile(uint8_t *dst, uint16_t dstPitch,
+    TileType tile, const uint8_t *src, MaskType maskType, const uint8_t *tbl, const DunTileClip &clip);
+
+DVL_ALWAYS_INLINE const uint8_t *GetDungeonTileSrc(uint16_t frame)
+{
+	const auto *pFrameTable = reinterpret_cast<const uint32_t *>(pDungeonCels.get());
+	return reinterpret_cast<const uint8_t *>(&pDungeonCels[SDL_SwapLE32(pFrameTable[frame])]);
+}
+
+DVL_ALWAYS_INLINE DunTileClip CalculateDunTileClip(int_fast16_t x, int_fast16_t y, int_fast16_t w, int_fast16_t h, const Surface &out)
+{
+	DunTileClip clip;
+	clip.top = y + 1 < h ? h - (y + 1) : 0;
+	clip.bottom = y + 1 > out.h() ? (y + 1) - out.h() : 0;
+	clip.left = x < 0 ? -x : 0;
+	clip.right = x + w > out.w() ? x + w - out.w() : 0;
+	clip.width = w - clip.left - clip.right;
+	clip.height = h - clip.top - clip.bottom;
+	return clip;
+}
+
+DVL_ALWAYS_INLINE int_fast16_t GetTileHeight(TileType tile)
+{
+	return tile == TileType::LeftTriangle || tile == TileType::RightTriangle ? 31 : 32;
+}
+
 /**
  * @brief Blit current world CEL to the given buffer
  * @param out Target buffer
@@ -171,15 +207,55 @@ std::string_view MaskTypeToString(MaskType maskType);
  * @param maskType The mask to use,
  * @param tbl LightTable or TRN for a tile.
  */
-void RenderTile(const Surface &out, Point position,
-    LevelCelBlock levelCelBlock, MaskType maskType, const uint8_t *tbl);
+DVL_ALWAYS_INLINE void RenderTile(const Surface &out, Point position,
+    LevelCelBlock levelCelBlock, MaskType maskType, const uint8_t *tbl)
+{
+	const TileType tile = levelCelBlock.type();
+	const DunTileClip clip = CalculateDunTileClip(position.x, position.y, DunFrameWidth, GetTileHeight(tile), out);
+	if (clip.width <= 0 || clip.height <= 0) return;
+	uint8_t *dst = out.at(static_cast<int>(position.x + clip.left), static_cast<int>(position.y - clip.bottom));
+	RenderTile(dst, out.pitch(), tile, GetDungeonTileSrc(levelCelBlock.frame()), maskType, tbl, clip);
+}
 
 /**
- * @brief Render a black 64x31 tile ◆
+ * @brief Render a single color 64x31 tile ◆
  * @param out Target buffer
  * @param sx Target buffer coordinate (left corner of the tile)
  * @param sy Target buffer coordinate (bottom corner of the tile)
+ * @param color Color index
+ */
+void RenderSingleColorTile(const Surface &out, int sx, int sy, uint8_t color = 0);
+
+DVL_ALWAYS_INLINE bool IsFullyDark(const uint8_t *DVL_RESTRICT tbl)
+{
+	return tbl == FullyDarkLightTable;
+}
+
+DVL_ALWAYS_INLINE bool IsFullyLit(const uint8_t *DVL_RESTRICT tbl)
+{
+	return tbl == FullyLitLightTable;
+}
+
+/**
+ * @brief Renders a tile without masking and without lighting.
+ */
+DVL_ALWAYS_INLINE void RenderFullyLitOpaqueTile(TileType tile, const Surface &out, Point position, const uint8_t *src)
+{
+	const DunTileClip clip = CalculateDunTileClip(position.x, position.y, DunFrameWidth, GetTileHeight(tile), out);
+	if (clip.width <= 0 || clip.height <= 0) return;
+	uint8_t *dst = out.at(static_cast<int>(position.x + clip.left), static_cast<int>(position.y - clip.bottom));
+	const uint16_t dstPitch = out.pitch();
+
+	// Doesn't matter what `FullyLitLightTable` light table points to, as long as it's not `nullptr`.
+	uint8_t *fullyLitBefore = FullyLitLightTable;
+	FullyLitLightTable = LightTables[0].data();
+	RenderTile(dst, dstPitch, tile, src, MaskType::Solid, FullyLitLightTable, clip);
+	FullyLitLightTable = fullyLitBefore;
+}
+
+/**
+ * @brief Writes a tile with the color swaps from `tbl` to `dst`.
  */
-void world_draw_black_tile(const Surface &out, int sx, int sy);
+void DunTileApplyTrans(LevelCelBlock levelCelBlock, uint8_t *DVL_RESTRICT dst, const uint8_t *tbl);
 
 } // namespace devilution
diff --git a/Source/engine/render/scrollrt.cpp b/Source/engine/render/scrollrt.cpp
index 32680101db9..913639f5b2e 100644
--- a/Source/engine/render/scrollrt.cpp
+++ b/Source/engine/render/scrollrt.cpp
@@ -5,6 +5,8 @@
  */
 #include "engine/render/scrollrt.h"
 
+#include <algorithm>
+#include <array>
 #include <cmath>
 #include <cstdint>
 
@@ -19,6 +21,7 @@
 #include "doom.h"
 #include "engine/backbuffer_state.hpp"
 #include "engine/dx.h"
+#include "engine/point.hpp"
 #include "engine/render/clx_render.hpp"
 #include "engine/render/dun_render.hpp"
 #include "engine/render/text_render.hpp"
@@ -29,6 +32,7 @@
 #include "hwcursor.hpp"
 #include "init.h"
 #include "inv.h"
+#include "levels/dun_tile.hpp"
 #include "lighting.h"
 #include "lua/lua.hpp"
 #include "minitext.h"
@@ -48,9 +52,7 @@
 #include "stores.h"
 #include "towners.h"
 #include "utils/attributes.h"
-#include "utils/bitset2d.hpp"
 #include "utils/display.h"
-#include "utils/endian.hpp"
 #include "utils/log.hpp"
 #include "utils/str_cat.hpp"
 
@@ -521,8 +523,7 @@ void DrawCell(const Surface &out, Point tilePosition, Point targetBufferPosition
 			const MaskType maskType = getFirstTileMaskLeft(tileType);
 			if (levelCelBlock.hasValue()) {
 				if (maskType != MaskType::LeftFoliage || tileType == TileType::TransparentSquare) {
-					RenderTile(out, targetBufferPosition,
-					    levelCelBlock, maskType, tbl);
+						RenderTile(out, targetBufferPosition, levelCelBlock, maskType, tbl);
 				}
 			}
 		}
@@ -533,8 +534,7 @@ void DrawCell(const Surface &out, Point tilePosition, Point targetBufferPosition
 			if (levelCelBlock.hasValue()) {
 				if (transparency || !foliage || levelCelBlock.type() == TileType::TransparentSquare) {
 					if (maskType != MaskType::RightFoliage || tileType == TileType::TransparentSquare) {
-						RenderTile(out, targetBufferPosition + Displacement { TILE_WIDTH / 2, 0 },
-						    levelCelBlock, maskType, tbl);
+						RenderTile(out, targetBufferPosition + Displacement { TILE_WIDTH / 2, 0 }, levelCelBlock, maskType, tbl);
 					}
 				}
 			}
@@ -543,56 +543,14 @@ void DrawCell(const Surface &out, Point tilePosition, Point targetBufferPosition
 	}
 
 	for (uint_fast8_t i = 2, n = MicroTileLen; i < n; i += 2) {
-		{
-			const LevelCelBlock levelCelBlock { pMap->mt[i] };
-			if (levelCelBlock.hasValue()) {
-				RenderTile(out, targetBufferPosition,
-				    levelCelBlock,
-				    transparency ? MaskType::Transparent : MaskType::Solid, tbl);
-			}
-		}
-		{
-			const LevelCelBlock levelCelBlock { pMap->mt[i + 1] };
-			if (levelCelBlock.hasValue()) {
-				RenderTile(out, targetBufferPosition + Displacement { TILE_WIDTH / 2, 0 },
-				    levelCelBlock,
-				    transparency ? MaskType::Transparent : MaskType::Solid, tbl);
-			}
-		}
-		targetBufferPosition.y -= TILE_HEIGHT;
-	}
-}
-
-/**
- * @brief Render a floor tile.
- * @param out Target buffer
- * @param tilePosition dPiece coordinates
- * @param targetBufferPosition Target buffer coordinate
- */
-void DrawFloorTile(const Surface &out, Point tilePosition, Point targetBufferPosition)
-{
-	const int lightTableIndex = dLight[tilePosition.x][tilePosition.y];
-
-	const uint8_t *tbl = LightTables[lightTableIndex].data();
-#ifdef _DEBUG
-	if (DebugPath && MyPlayer->IsPositionInPath(tilePosition))
-		tbl = GetPauseTRN();
-#endif
-
-	const uint16_t levelPieceId = dPiece[tilePosition.x][tilePosition.y];
-	{
-		const LevelCelBlock levelCelBlock { DPieceMicros[levelPieceId].mt[0] };
-		if (levelCelBlock.hasValue()) {
-			RenderTile(out, targetBufferPosition,
-			    levelCelBlock, MaskType::Solid, tbl);
+		if (const LevelCelBlock levelCelBlock { pMap->mt[i] }; levelCelBlock.hasValue()) {
+			RenderTile(out, targetBufferPosition, levelCelBlock, transparency ? MaskType::Transparent : MaskType::Solid, tbl);
 		}
-	}
-	{
-		const LevelCelBlock levelCelBlock { DPieceMicros[levelPieceId].mt[1] };
-		if (levelCelBlock.hasValue()) {
+		if (const LevelCelBlock levelCelBlock { pMap->mt[i + 1] }; levelCelBlock.hasValue()) {
 			RenderTile(out, targetBufferPosition + Displacement { TILE_WIDTH / 2, 0 },
-			    levelCelBlock, MaskType::Solid, tbl);
+					levelCelBlock, transparency ? MaskType::Transparent : MaskType::Solid, tbl);
 		}
+		targetBufferPosition.y -= TILE_HEIGHT;
 	}
 }
 
@@ -823,8 +781,90 @@ void DrawDungeon(const Surface &out, Point tilePosition, Point targetBufferPosit
 	}
 }
 
+constexpr int MinFloorTileToBakeLight = 2;
+constexpr size_t FloorTilesPerLightBufferSize = 96;
+constexpr size_t FloorTilesCachePerLightLevel = 32;
+constexpr size_t FloorBufferNumLightingLevels = NumLightingLevels
+#ifdef DEBUG_
+    + 1
+#endif
+    ;
+
+struct FloorTilesBufferEntry {
+	PointOf<int16_t> targetBufferPosition;
+	uint16_t levelCelBlock;
+
+	bool operator<(const FloorTilesBufferEntry &other) const
+	{
+		return levelCelBlock < other.levelCelBlock;
+	}
+};
+
+struct FloorTilesBuffer {
+	std::array<StaticVector<FloorTilesBufferEntry, FloorTilesPerLightBufferSize>, FloorBufferNumLightingLevels> perLightLevel;
+
+	void add(PointOf<int16_t> position,
+	    uint16_t levelCelBlock,
+	    uint8_t lightTableIndex)
+	{
+		perLightLevel[lightTableIndex].emplace_back(FloorTilesBufferEntry { position, levelCelBlock });
+	}
+};
+
+const uint8_t *GetLightTableFromIndex(uint8_t lightTableIndex)
+{
+	return
+#ifdef _DEBUG
+	    lightTableIndex == NumLightingLevels
+	    ? GetPauseTRN()
+	    :
+#endif
+	    LightTables[lightTableIndex].data();
+}
+
+void DrawIdenticalFloorTiles(const Surface &out, uint8_t lightTableIndex, FloorTilesBufferEntry *begin, FloorTilesBufferEntry *end)
+{
+	if (begin == end) return;
+	const LevelCelBlock levelCelBlock { begin->levelCelBlock };
+	const uint8_t *lightTable = GetLightTableFromIndex(lightTableIndex);
+	if (begin + MinFloorTileToBakeLight >= end) {
+		for (FloorTilesBufferEntry *it = begin; it != end; ++it) {
+			RenderTile(out, it->targetBufferPosition, levelCelBlock, MaskType::Solid, lightTable);
+		}
+	} else {
+		uint8_t bakedLightTile[DunFrameHeight * DunFrameWidth];
+		DunTileApplyTrans(levelCelBlock, bakedLightTile, lightTable);
+		const TileType type = levelCelBlock.type();
+		for (FloorTilesBufferEntry *it = begin; it != end; ++it) {
+			RenderFullyLitOpaqueTile(type, out, it->targetBufferPosition, bakedLightTile);
+		}
+	}
+}
+
+void DrawFloorBuffer(const Surface &out, uint8_t lightTableIndex, FloorTilesBuffer &perLightBuffers)
+{
+	StaticVector<FloorTilesBufferEntry, FloorTilesPerLightBufferSize> &buffer = perLightBuffers.perLightLevel[lightTableIndex];
+	if (buffer.empty()) return;
+	if (buffer.size() > 2) std::sort(buffer.begin(), buffer.end());
+	uint16_t prevLevelCelBlock = buffer[0].levelCelBlock;
+	size_t prevBegin = 0;
+	FloorTilesBufferEntry *arr = buffer.begin();
+	for (size_t i = 1, n = buffer.size(); i < n; ++i) {
+		const uint16_t levelCelBlock = buffer[i].levelCelBlock;
+		if (prevLevelCelBlock != levelCelBlock) {
+			DrawIdenticalFloorTiles(out, lightTableIndex, &arr[prevBegin], &arr[i]);
+			prevLevelCelBlock = levelCelBlock;
+			prevBegin = i;
+		}
+	}
+	if (prevBegin != buffer.size()) {
+		DrawIdenticalFloorTiles(out, lightTableIndex, &arr[prevBegin], buffer.end());
+	}
+	buffer.clear();
+}
+
 /**
- * @brief Render a row of tiles
+ * @brief Renders the floor tiles
  * @param out Buffer to render to
  * @param tilePosition dPiece coordinates
  * @param targetBufferPosition Target buffer coordinates
@@ -833,33 +873,67 @@ void DrawDungeon(const Surface &out, Point tilePosition, Point targetBufferPosit
  */
 void DrawFloor(const Surface &out, Point tilePosition, Point targetBufferPosition, int rows, int columns)
 {
+	FloorTilesBuffer buffer;
+	PointOf<int16_t> position = targetBufferPosition;
 	for (int i = 0; i < rows; i++) {
-		for (int j = 0; j < columns; j++) {
-			if (InDungeonBounds(tilePosition)) {
-				if (!TileHasAny(tilePosition, TileProperties::Solid))
-					DrawFloorTile(out, tilePosition, targetBufferPosition);
-			} else {
-				world_draw_black_tile(out, targetBufferPosition.x, targetBufferPosition.y);
+		for (int j = 0; j < columns; ++j) {
+			if (!InDungeonBounds(tilePosition)) {
+				RenderSingleColorTile(out, position.x, position.y);
+			} else if (!TileHasAny(tilePosition, TileProperties::Solid)) {
+				const uint16_t levelPieceId = dPiece[tilePosition.x][tilePosition.y];
+				const uint8_t lightTableIndex =
+#ifdef _DEBUG
+				    DebugPath && MyPlayer->IsPositionInPath(tilePosition)
+				    ? NumLightingLevels
+				    :
+#endif
+				    dLight[tilePosition.x][tilePosition.y];
+				const auto *lightTable = GetLightTableFromIndex(lightTableIndex);
+				const LevelCelBlock levelCelBlock0 = DPieceMicros[levelPieceId].mt[0];
+				const LevelCelBlock levelCelBlock1 = DPieceMicros[levelPieceId].mt[1];
+				if (IsFullyLit(lightTable) || IsFullyDark(lightTable)) {
+					// Fully lit / dark light tables do not need baked light, render them directly:
+					if (levelCelBlock0.hasValue()) {
+						RenderTile(out, position, levelCelBlock0, MaskType::Solid, lightTable);
+					}
+					if (levelCelBlock1.hasValue()) {
+						RenderTile(out, position + Displacement { TILE_WIDTH / 2, 0 }, levelCelBlock1, MaskType::Solid, lightTable);
+					}
+				} else {
+					if (buffer.perLightLevel[lightTableIndex].size() + 2 > FloorTilesPerLightBufferSize) {
+						DrawFloorBuffer(out, lightTableIndex, buffer);
+					}
+					if (levelCelBlock0.hasValue()) {
+						buffer.add(position, levelCelBlock0.data, lightTableIndex);
+					}
+					if (levelCelBlock1.hasValue()) {
+						buffer.add(position + Displacement { TILE_WIDTH / 2, 0 }, levelCelBlock1.data, lightTableIndex);
+					}
+				}
 			}
 			tilePosition += Direction::East;
-			targetBufferPosition.x += TILE_WIDTH;
+			position.x += TILE_WIDTH;
 		}
+
 		// Return to start of row
 		tilePosition += Displacement(Direction::West) * columns;
-		targetBufferPosition.x -= columns * TILE_WIDTH;
+		position.x -= columns * TILE_WIDTH;
 
 		// Jump to next row
-		targetBufferPosition.y += TILE_HEIGHT / 2;
+		position.y += TILE_HEIGHT / 2;
 		if ((i & 1) != 0) {
 			tilePosition.x++;
 			columns--;
-			targetBufferPosition.x += TILE_WIDTH / 2;
+			position.x += TILE_WIDTH / 2;
 		} else {
 			tilePosition.y++;
 			columns++;
-			targetBufferPosition.x -= TILE_WIDTH / 2;
+			position.x -= TILE_WIDTH / 2;
 		}
 	}
+	for (uint8_t lightTableIndex = 0; lightTableIndex < FloorBufferNumLightingLevels; ++lightTableIndex) {
+		DrawFloorBuffer(out, lightTableIndex, buffer);
+	}
 }
 
 [[nodiscard]] DVL_ALWAYS_INLINE bool IsWall(Point position)
@@ -987,7 +1061,7 @@ void Zoom(const Surface &out)
 
 Displacement tileOffset;
 Displacement tileShift;
-int tileColums;
+int tileColumns;
 int tileRows;
 
 void CalcFirstTilePosition(Point &position, Displacement &offset)
@@ -1048,7 +1122,7 @@ void DrawGame(const Surface &fullOut, Point position, Displacement offset)
 	    ? fullOut.subregionY(0, gnViewportHeight)
 	    : fullOut.subregionY(0, (gnViewportHeight + 1) / 2);
 
-	int columns = tileColums;
+	int columns = tileColumns;
 	int rows = tileRows;
 
 	// Skip rendering parts covered by the panels
@@ -1511,7 +1585,7 @@ void CalcViewportGeometry()
 	const int viewportHeight = GetViewportHeight() / zoomFactor;
 	const Point renderStart = startPosition - Displacement { TILE_WIDTH / 2, TILE_HEIGHT / 2 };
 	tileRows = (viewportHeight - renderStart.y + TILE_HEIGHT / 2 - 1) / (TILE_HEIGHT / 2);
-	tileColums = (screenWidth - renderStart.x + TILE_WIDTH - 1) / TILE_WIDTH;
+	tileColumns = (screenWidth - renderStart.x + TILE_WIDTH - 1) / TILE_WIDTH;
 }
 
 Point GetScreenPosition(Point tile)
diff --git a/Source/utils/attributes.h b/Source/utils/attributes.h
index 2f40a9bcb6e..cd8e95beb73 100644
--- a/Source/utils/attributes.h
+++ b/Source/utils/attributes.h
@@ -5,6 +5,10 @@
  */
 #pragma once
 
+#ifdef __cplusplus
+#include <execution>
+#endif
+
 #ifdef __has_attribute
 #define DVL_HAVE_ATTRIBUTE(x) __has_attribute(x)
 #else
@@ -32,6 +36,14 @@
 #define DVL_ALWAYS_INLINE inline
 #endif
 
+#if DVL_HAVE_ATTRIBUTE(noinline)
+#define DVL_NO_INLINE __attribute__((noinline))
+#elif defined(_MSC_VER)
+#define DVL_NO_INLINE __declspec(noinline)
+#else
+#define DVL_NO_INLINE
+#endif
+
 #if DVL_HAVE_ATTRIBUTE(hot)
 #define DVL_ATTRIBUTE_HOT __attribute__((hot))
 #else
@@ -102,3 +114,11 @@
 #else
 #define DVL_UNREACHABLE()
 #endif
+
+#ifdef __cplusplus
+#if __cpp_lib_execution >= 201902L
+#define DVL_EXECUTION_UNSEQ std::execution::unseq,
+#else
+#define DVL_EXECUTION_UNSEQ
+#endif
+#endif