Add new renderer and prepare for render overhaul

7 files changed, 1122 insertions, 214 deletions

diff --git a/src/main.c b/src/main.c
index e694057..f89e743 100644
--- a/src/main.c
+++ b/src/main.c
@@ -12,17 +12,33 @@ WITH REGARD TO THIS SOFTWARE.
 #include "gba/gba.h"
 
 #include "filesystem.c"
-#include "renderer.c"
+#include "renderer_m0.c"
 #include "sequencer.c"
 
-#define PROF_ENABLE 0
+#define PROF_ENABLE 1
 #include "profiling.c"
 
-//
-// Config parameters.
-//
 
-int main(void) {
+void
+render(void) {
+    PROF(screen_fill(0), clear_cycles);
+    PROF(draw_triggers(), draw_trigs_cycles);
+    PROF(draw_channels(), draw_btn_cycles);
+    PROF(draw_pattern_buttons(), draw_btn_cycles);
+    PROF(draw_bank_buttons(), draw_btn_cycles);
+    PROF(draw_bpm(), draw_btn_cycles);
+    PROF(draw_play(), draw_btn_cycles);
+    PROF(draw_stop(), draw_btn_cycles);
+    PROF(draw_piano(), draw_piano_cycles);
+    PROF(draw_parameters(), draw_param_cycles);
+    PROF(draw_trig_cursor(trig_selection_loc, COL_CURSOR), draw_cursor_cycles);
+    PROF(draw_channel_cursor(channel_selection_loc, COL_GREY), draw_cursor_cycles);
+    PROF(draw_pattern_cursor(pattern_selection_loc, COL_GREY), draw_cursor_cycles);
+    PROF(draw_current_step(COL_RED), draw_cursor_cycles);
+}
+
+int
+main(void) {
     // Adjust system wait times.
     SYSTEM_WAIT = SYSTEM_WAIT_CARTRIDGE;
 
@@ -38,15 +54,17 @@ int main(void) {
 
     // Initialize sequencer.
     sequencer_init();
+    txt_spacing(6);
 
     // Main loop.
     while (true) {
         poll_keys();
         bios_vblank_wait();
-        FRAME_START();
-        PROF(handle_sequencer_input(), input_cycles);
         PROF_SHOW();
+        FRAME_START();
         PROF(flip_buffer(), flip_cycles);
+        PROF(handle_sequencer_input(), input_cycles);
+        PROF(render(), render_cycles);
         FRAME_END();
     }
 
diff --git a/src/profiling.c b/src/profiling.c
index de969d2..7a4f4ad 100644
--- a/src/profiling.c
+++ b/src/profiling.c
@@ -52,22 +52,23 @@ static bool profile_bg_show = true;
             profile_bg_show ^= 1;\
         }\
         if (profile_show) {\
+            txt_color(1);\
             txt_position((PROF_SHOW_X), (PROF_SHOW_Y));\
-            draw_filled_rect((PROF_SHOW_X), (PROF_SHOW_X), 8 * 18, 8 * 16, 0);\
+            if (profile_bg_show) {\
+                draw_filled_rect((PROF_SHOW_X), (PROF_SHOW_X), 8 * 18, 8 * 14, 0);\
+            }\
             txt_printf("VIDEO\n");\
             txt_printf(">CLEAR  %.8lu\n", avg_clear_cycles);\
-            txt_printf(">LINES  %.8lu\n", avg_line_cycles);\
-            txt_printf(">RECT   %.8lu\n", avg_rect_cycles);\
-            txt_printf(">FRECT  %.8lu\n", avg_fill_rect_cycles);\
-            txt_printf(">1BPP   %.8lu\n", avg_icn_cycles);\
-            txt_printf(">2BPP   %.8lu\n", avg_chr_cycles);\
             txt_printf(">FLIP   %.8lu\n", avg_flip_cycles);\
-            txt_printf("TEXT\n");\
-            txt_printf(">DRAWF  %.8lu\n", avg_txt_drawf_cycles);\
-            txt_printf(">PRINTF %.8lu\n", avg_txt_printf_cycles);\
-            txt_printf(">RENDER %.8lu\n", avg_txt_render_cycles);\
-            txt_printf(">CLEAR  %.8lu\n", avg_txt_clear_cycles);\
+            txt_printf("SEQUENCER RENDER\n");\
+            txt_printf(">TRIGS  %.8lu\n", avg_draw_trigs_cycles);\
+            txt_printf(">BTNS   %.8lu\n", avg_draw_btns_cycles);\
+            txt_printf(">PARAM  %.8lu\n", avg_draw_param_cycles);\
+            txt_printf(">PIANO  %.8lu\n", avg_draw_piano_cycles);\
+            txt_printf(">CURSOR %.8lu\n", avg_draw_cursor_cycles);\
+            txt_printf(">RENDER %.8lu\n", avg_render_cycles);\
             txt_printf("TOTAL   %.8lu\n", avg_frame_cycles);\
+            txt_render();\
         }\
         if (profile_bg_show) {\
             u32 frame_time =\
@@ -80,40 +81,35 @@ static bool profile_bg_show = true;
                     FP_NUM(280896 * 60, 2),\
                     FP_NUM(avg_frame_cycles + 1, 2),\
                     2);\
-            txt_printf("TIME    %.8lu\n", frame_time >> 2);\
-            txt_printf("FPS     %.8lu\n", (fps >> 2) + 1);\
+            draw_filled_rect(8 * 18, 0, 239, 16, 0);\
+            txt_drawf("TIME: %.6lu", 8 * 18, 0, 1, frame_time >> 2);\
+            txt_drawf("MAX FPS:%.4lu", 8 * 18, 8, 1, (fps >> 2) + 1);\
         }\
     } while (0)
 
 static u32 prof_frame_counter = 0;
 
-static u32 frame_cycles      = 0;
-static u32 flip_cycles       = 0;
-static u32 clear_cycles      = 0;
-static u32 line_cycles       = 0;
-static u32 rect_cycles       = 0;
-static u32 fill_rect_cycles  = 0;
-static u32 chr_cycles        = 0;
-static u32 icn_cycles        = 0;
-static u32 txt_drawf_cycles  = 0;
-static u32 txt_printf_cycles = 0;
-static u32 txt_render_cycles = 0;
-static u32 txt_clear_cycles  = 0;
-static u32 input_cycles      = 0;
+static u32 frame_cycles        = 0;
+static u32 flip_cycles         = 0;
+static u32 clear_cycles        = 0;
+static u32 input_cycles        = 0;
+static u32 draw_trigs_cycles   = 0;
+static u32 draw_btn_cycles = 0;
+static u32 draw_piano_cycles   = 0;
+static u32 draw_param_cycles   = 0;
+static u32 draw_cursor_cycles   = 0;
+static u32 render_cycles       = 0;
 
-static u32 avg_frame_cycles      = 0;
-static u32 avg_flip_cycles       = 0;
-static u32 avg_clear_cycles      = 0;
-static u32 avg_line_cycles       = 0;
-static u32 avg_rect_cycles       = 0;
-static u32 avg_fill_rect_cycles  = 0;
-static u32 avg_chr_cycles        = 0;
-static u32 avg_icn_cycles        = 0;
-static u32 avg_txt_drawf_cycles  = 0;
-static u32 avg_txt_printf_cycles = 0;
-static u32 avg_txt_render_cycles = 0;
-static u32 avg_txt_clear_cycles  = 0;
-static u32 avg_input_cycles      = 0;
+static u32 avg_frame_cycles         = 0;
+static u32 avg_flip_cycles          = 0;
+static u32 avg_clear_cycles         = 0;
+static u32 avg_input_cycles         = 0;
+static u32 avg_draw_trigs_cycles    = 0;
+static u32 avg_draw_btns_cycles = 0;
+static u32 avg_draw_piano_cycles    = 0;
+static u32 avg_draw_param_cycles   = 0;
+static u32 avg_draw_cursor_cycles   = 0;
+static u32 avg_render_cycles        = 0;
 
 #if PROF_ENABLE == 1
 #define FRAME_START()\
@@ -122,29 +118,23 @@ static u32 avg_input_cycles      = 0;
             avg_frame_cycles      = frame_cycles      / prof_frame_counter;\
             avg_flip_cycles       = flip_cycles       / prof_frame_counter;\
             avg_clear_cycles      = clear_cycles      / prof_frame_counter;\
-            avg_line_cycles       = line_cycles       / prof_frame_counter;\
-            avg_rect_cycles       = rect_cycles       / prof_frame_counter;\
-            avg_fill_rect_cycles  = fill_rect_cycles  / prof_frame_counter;\
-            avg_chr_cycles        = chr_cycles        / prof_frame_counter;\
-            avg_icn_cycles        = icn_cycles        / prof_frame_counter;\
-            avg_txt_drawf_cycles  = txt_drawf_cycles  / prof_frame_counter;\
-            avg_txt_printf_cycles = txt_printf_cycles / prof_frame_counter;\
-            avg_txt_render_cycles = txt_render_cycles / prof_frame_counter;\
-            avg_txt_clear_cycles  = txt_clear_cycles  / prof_frame_counter;\
-            avg_input_cycles      = input_cycles  / prof_frame_counter;\
+            avg_draw_trigs_cycles        = draw_trigs_cycles        / prof_frame_counter;\
+            avg_draw_btns_cycles     = draw_btn_cycles     / prof_frame_counter;\
+            avg_draw_piano_cycles = draw_piano_cycles / prof_frame_counter;\
+            avg_draw_param_cycles      = draw_param_cycles     / prof_frame_counter;\
+            avg_draw_cursor_cycles      = draw_cursor_cycles     / prof_frame_counter;\
+            avg_input_cycles      = input_cycles     / prof_frame_counter;\
+            avg_render_cycles      = render_cycles     / prof_frame_counter;\
             frame_cycles          = 0;\
             flip_cycles           = 0;\
             clear_cycles          = 0;\
-            line_cycles           = 0;\
-            rect_cycles           = 0;\
-            fill_rect_cycles      = 0;\
-            chr_cycles            = 0;\
-            icn_cycles            = 0;\
-            txt_drawf_cycles      = 0;\
-            txt_printf_cycles     = 0;\
-            txt_render_cycles     = 0;\
-            txt_clear_cycles      = 0;\
             input_cycles          = 0;\
+            render_cycles          = 0;\
+            draw_trigs_cycles     = 0;\
+            draw_param_cycles     = 0;\
+            draw_cursor_cycles     = 0;\
+            draw_btn_cycles   = 0;\
+            draw_piano_cycles     = 0;\
             prof_frame_counter    = 0;\
         }\
         profile_start();\
diff --git a/src/renderer.h b/src/renderer.h
index 4620c27..e6637ef 100644
--- a/src/renderer.h
+++ b/src/renderer.h
@@ -3,31 +3,6 @@
 
 #include "gba/gba.h"
 
-// The frontbuffer is located at the beginning of the VRAM, and requires 20KB of
-// video memory for 32 * 20 tiles at 4bpp.
-#define FRONTBUF ((u32*)(MEM_VRAM))
-
-// Adjust both of these if the location of the map changes. Each screnblock
-// requires less than 2KB.
-#define FRONTBUF_TILEMAP ((u16*)(MEM_VRAM + KB(20)))
-#define FRONTBUF_SB      10
-
-// The backbuffer is located at the end of the VRAM. This can allow us to use
-// more backgrounds but eats into the available memory for sprites. This should
-// be fine for non sprite intensive applications. If more sprite memory is
-// needed, the backbuffer can be located at the end of the background memory
-// instead (64KB - 20KB).
-#define BACKBUF  ((u32*)(MEM_VRAM + KB(96) - KB(20)))
-
-// The font data is located at the end of the frontbuffer memory, after the tile
-// map and requires 8KB for 256 8x8 characters at 4bpp. This, along with the
-// tilemap information allow us to store the frontbuffer and font for a text
-// background in the first 2 charblocks (32KB).
-#define FONT_DATA    ((u32*)(MEM_VRAM + KB(22)))
-#define FONT_TILEMAP ((u16*)(MEM_VRAM + KB(30)))
-#define FONT_SB      15
-#define FONT_OFFSET  192
-
 // Draws a pixel to the given (x, y) position on the framebuffer. All drawing
 // functions use paletted colors (clr: 0-15).
 void draw_pixel(size_t x, size_t y, u8 clr);
@@ -41,17 +16,19 @@ void draw_rect(size_t x0, size_t y0, size_t x1, size_t y1, u8 clr);
 // Draw a filled rectangle between (x0, y0) and (x1, y1) (x0 <= x1 and y0 <= y1).
 void draw_filled_rect(size_t x0, size_t y0, size_t x1, size_t y1, u8 clr);
 
-// Draw a 8x8 tile starting at the (x, y) position. If the merge parameter is
-// set, colors will be added together instead of replaced. This could lead to
-// some merging issues if we are not careful with the chosen colors. The tile
-// color will be multiplied by the given clr parameter, which is useful to
-// change the color of flat tiles.
-void draw_tile(size_t x, size_t y, Tile *tile, u8 clr, bool merge);
+// Fills the framebuffer with the given color.
+void screen_fill(u8 clr);
+
+// Draws a chr sprite (16 * u8). The first 8 bytes correspond to ch0 and the
+// last 8 to ch1. If clr is 0 the regular 4bit color will be used, from clr 1-14
+// the given color will overwrite the existing one. Color 15 will "clear" the
+// sprite instead.
+void draw_chr(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y);
 
-// Fills the framebuffer with color 0.
-void clear_screen(void);
+// Draws a 1bpp icn sprite in the given color.
+void draw_icn(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y);
 
-// Copies the content of dirty tiles from the backbuffer into the frontbuffer.
+// Copies data and performs page flipping if needed.
 // To be called exactly once at the beginning of the VBlank.
 void flip_buffer(void);
 
diff --git a/src/renderer_m0.c b/src/renderer_m0.c
new file mode 100644
index 0000000..8bd4263
--- /dev/null
+++ b/src/renderer_m0.c
@@ -0,0 +1,811 @@
+//
+// This Mode 0 renderer provides a way of drawing directly to a framebuffer
+// (similar to Mode 3 and 4) while retaining the flexibility of using other
+// backgrounds if needed. It also performs double buffering to avoid tearing
+// artifacts and tries to only draw tiles that changed on each frame.
+//
+
+#include "renderer.h"
+#include "text.h"
+
+//
+// Parameters.
+//
+
+#define SUBPIXEL_LINES 1
+#define DEC_BIG_LUT 1
+#define FLIP_TYPE 3
+
+// Front/back buffers for double buffering.
+#define BUF_0 ((u32*)(MEM_VRAM))
+#define BUF_1 ((u32*)(MEM_VRAM + KB(20)))
+
+// Pointer to the backbuffer.
+static u32 *backbuf = BUF_1;
+
+// Tracking which tiles are "dirty" and need refreshing.
+static u32 dirty_tiles[21] = {0};
+
+// Position of the tilemap.
+#define TILE_MAP  ((u32*)(MEM_VRAM + KB(40)))
+
+// Charblock and screenblock for both render buffers.
+#define CB_0 0
+#define CB_1 1
+#define SB_0 20
+#define SB_1 22
+
+// Boundchecks can be disable at compile time but this will not always improve
+// the performance and can in fact make it worse. It is possible that this is
+// due to some aliasing optimizations but not sure at this moment.
+#ifdef DISABLE_BOUNDCHECK_SCREEN
+#define BOUNDCHECK_SCREEN(X,Y)
+#else
+#define BOUNDCHECK_SCREEN(X,Y) if ((X) >= SCREEN_WIDTH || (Y) >= SCREEN_HEIGHT) return;
+#endif
+
+// Swap A and B values without a tmp variable.
+#define SWAP(A, B) (((A) ^= (B)), ((B) ^= (A)), ((A) ^= (B)))
+
+// Swap A and B values to make sure A <= B.
+#define MAYBE_SWAP(A,B) if ((A) > (B)) { SWAP(A,B); }
+
+//
+// Basic primitives.
+//
+
+static inline
+void
+redraw(void) {
+    for (size_t i = 0; i < 21; i++) {
+        dirty_tiles[i] = 0xFFFFFFFF;
+    }
+}
+
+IWRAM_CODE
+void screen_fill(u8 clr)  {
+    // We have to make sure we leave the last tile blank to use as alpha channel
+    // when moving the BG during double buffering.
+    dma_fill(backbuf, 0x11111111 * clr, KB(20) - 32, 3);
+    redraw();
+}
+
+IWRAM_CODE
+void
+draw_pixel(size_t x, size_t y, u8 clr) {
+    BOUNDCHECK_SCREEN(x, y);
+
+    // Find row position for the given x/y coordinates.
+    size_t tile_x = x / 8;
+    size_t tile_y = y / 8;
+    size_t start_col = x % 8;
+    size_t start_row = y % 8;
+    u32 *dst = &backbuf[start_row + (tile_x + tile_y * 32) * 8];
+
+    // Update backbuffer.
+    size_t shift = start_col * sizeof(u32);
+    u32 mask = 0xF << shift;
+    u32 row = clr << shift;
+    *dst = (*dst & ~mask) | row;
+    dirty_tiles[tile_y] |= 1 << tile_x;
+}
+
+IWRAM_CODE
+static inline
+void
+draw_hline(size_t x0, size_t x1, size_t y0, u8 clr) {
+    BOUNDCHECK_SCREEN(x0, y0);
+    BOUNDCHECK_SCREEN(x1, y0);
+    // Find row positions for the given x/y coordinates.
+    size_t tile_x0 = x0 / 8;
+    size_t tile_x1 = x1 / 8;
+    size_t tile_y = y0 / 8;
+    size_t start_col = x0 % 8;
+    size_t end_col = x1 % 8;
+    size_t start_row = y0 % 8;
+    u32 dirty = (1 << tile_x0) | (1 << tile_x1);
+
+    // Horizontal line. There are 3 cases:
+    //     1. Lines fit on a single tile.
+    //     2. Lines go through 2 tiles, both require partial row updates.
+    //     3. Lines go through 3 or more tiles, first and last tiles use
+    //        partial row updates, rows in the middle can write the entire
+    //        row.
+    size_t dtx = tile_x1 - tile_x0;
+    u32 *dst = &backbuf[start_row + (tile_x0 + tile_y * 32) * 8];
+    if (dtx < 1) {
+        size_t shift_left = start_col * 4;
+        size_t shift_right = (7 - end_col) * 4;
+        u32 mask = (0xFFFFFFFF >> shift_right) & (0xFFFFFFFF << shift_left);
+        u32 row = (0x11111111 * clr) & mask;
+        *dst = (*dst & ~mask) | row;
+    } else {
+        size_t shift_left = start_col * 4;
+        size_t shift_right = (7 - end_col) * 4;
+        u32 mask = 0xFFFFFFFF;
+        u32 row = 0x11111111 * clr;
+        *dst = (*dst & ~(mask << shift_left)) | (row << shift_left);
+        dst += 8;
+        for (size_t i = 1; i < dtx; i++) {
+            dirty |= (1 << (tile_x0 + i));
+            *dst = row;
+            dst += 8;
+        }
+        *dst = (*dst & ~(mask >> shift_right)) | (row >> shift_right);
+    }
+    dirty_tiles[tile_y] |= dirty;
+}
+
+IWRAM_CODE
+UNROLL_LOOPS
+static inline
+void
+draw_vline(size_t x0, size_t y0, size_t y1, u8 clr) {
+    BOUNDCHECK_SCREEN(x0, y0);
+    BOUNDCHECK_SCREEN(x0, y1);
+    size_t tile_x = x0 / 8;
+    size_t tile_y = y0 / 8;
+    size_t tile_y0 = y0 / 8;
+    size_t tile_y1 = y1 / 8;
+    size_t start_col = x0 % 8;
+    size_t start_row0 = y0 % 8;
+    size_t start_row1 = y1 % 8;
+
+    size_t shift_left = start_col * 4;
+    u32 dirty = (1 << tile_x);
+
+    u32 *dst = &backbuf[start_row0 + (tile_x + tile_y * 32) * 8];
+    u32 mask = 0x0000000F << shift_left;
+    u32 row = (0x11111111 * clr) & mask;
+    u32 dty = tile_y1 - tile_y0;
+    if (dty < 1) {
+        for (size_t i = 0; i <= (y1 - y0); i++, dst++) {
+            dst[0] = (dst[0] & ~mask) | row;
+        }
+    } else {
+        for (size_t i = 0; i < (8 - start_row0); i++, dst++) {
+            dst[0] = (dst[0] & ~mask) | row;
+        }
+        dst += 8 * 31;
+        for (size_t j = 1; j < dty; j++) {
+            dirty_tiles[tile_y0 + j] |= dirty;
+            for (size_t i = 0; i < 8; i++, dst++) {
+                dst[0] = (dst[0] & ~mask) | row;
+            }
+            dst += 8 * 31;
+        }
+        for (size_t i = 0; i <= start_row1; i++, dst++) {
+            dst[0] = (dst[0] & ~mask) | row;
+        }
+    }
+    dirty_tiles[tile_y0] |= dirty;
+    dirty_tiles[tile_y1] |= dirty;
+}
+
+IWRAM_CODE
+void
+draw_line(size_t x0, size_t y0, size_t x1, size_t y1, u8 clr) {
+    BOUNDCHECK_SCREEN(x0, y0);
+    BOUNDCHECK_SCREEN(x1, y1);
+    if (y0 == y1) {
+        MAYBE_SWAP(x0, x1);
+        draw_hline(x0, x1, y0, clr);
+    } else if (x0 == x1) {
+        MAYBE_SWAP(y0, y1);
+        draw_vline(x0, y0, y1, clr);
+    } else {
+        // Fixed Precision constants.
+        const int fp_bit = 6;
+        const int fp_one = FP_NUM(1, fp_bit);
+        const int fp_half = fp_one >> 1;
+
+        int dx = x0 > x1 ? x0 - x1 : x1 - x0;
+        int dy = y0 > y1 ? y0 - y1 : y1 - y0;
+
+        if ((dx >= dy && x0 > x1) || (dx < dy && y0 > y1)) {
+            SWAP(x0, x1);
+            SWAP(y0, y1);
+        }
+
+#if SUBPIXEL_LINES == 1
+        int dxf = (dx << fp_bit);
+        int dyf = (dy << fp_bit);
+        int frac_x = x0 > x1 ? FP_NUM(x0 - x1, fp_bit) : FP_NUM(x1 - x0, fp_bit);
+        int frac_y = y0 > y1 ? FP_NUM(y0 - y1, fp_bit) : FP_NUM(y1 - y0, fp_bit);
+        int x_step = x0 > x1 ? -1 : 1;
+        int y_step = y0 > y1 ? -1 : 1;
+        int distance = (frac_y - fp_one) * dx - (frac_x - fp_half) * dy;
+        if (dx >= dy) {
+            int step = dxf / dyf;
+            int remaining = dx;
+            while (remaining > (step - 1)) {
+                distance += step * 2 * dyf;
+                if (distance >= 0) {
+                    draw_hline(x0, x0 + step - 1, y0, clr);
+                    x0 += x_step * step;
+                    remaining -= step;
+                } else {
+                    if (remaining < step) {
+                        break;
+                    }
+                    draw_hline(x0, x0 + step, y0, clr);
+                    distance += 2 * dyf;
+                    x0 += x_step * (step + 1);
+                    remaining -= step + 1;
+                }
+                distance -= 2 * dxf;
+                y0 += y_step;
+            }
+            if (remaining >= 0) {
+                draw_hline(x0, x0 + remaining, y0, clr);
+            }
+        } else {
+            int step = dyf / dxf;
+            int remaining = dy;
+            while (remaining > (step - 1)) {
+                    distance += step * 2 * dxf;
+                    if (distance >= 0) {
+                        draw_vline(x0, y0, y0 + step - 1, clr);
+                        y0 += y_step * step;
+                        remaining -= step;
+                    } else {
+                        draw_vline(x0, y0, y0 + step, clr);
+                        distance += 2 * dxf;
+                        y0 += y_step * (step + 1);
+                        remaining -= step + 1;
+                    }
+                    distance -= 2 * dyf;
+                    x0 += x_step;
+            }
+            if (remaining >= 0) {
+                draw_vline(x0, y0, y0 + remaining, clr);
+            }
+        }
+#else
+        int x_step = x0 > x1 ? -1 : 1;
+        int y_step = y0 > y1 ? -1 : 1;
+        if (dx >= dy) {
+            int diff = 2 * dy - dx;
+            for (int i = 0; i < dx + 1; i++) {
+                draw_pixel(x0, y0, clr);
+                if (diff >= 0) {
+                    diff -= 2 * dx;
+                    y0 += y_step;
+                }
+                diff += 2 * dy;
+                x0 += x_step;
+            }
+        } else {
+            int diff = 2 * dx - dy;
+            for (int i = 0; i < dy + 1; i++) {
+                draw_pixel(x0, y0, clr);
+                if (diff >= 0) {
+                    diff -= 2 * dy;
+                    x0 += x_step;
+                }
+                diff += 2 * dx;
+                y0 += y_step;
+            }
+        }
+#endif
+    }
+}
+
+IWRAM_CODE
+void
+draw_rect(size_t x0, size_t y0, size_t x1, size_t y1, u8 clr) {
+    BOUNDCHECK_SCREEN(x0, y0);
+    BOUNDCHECK_SCREEN(x1, y1);
+    MAYBE_SWAP(x0, x1);
+    MAYBE_SWAP(y0, y1);
+
+    draw_hline(x0, x1, y0, clr);
+    draw_hline(x0, x1, y1, clr);
+    draw_vline(x0, y0, y1, clr);
+    draw_vline(x1, y0, y1, clr);
+}
+
+IWRAM_CODE
+void
+draw_filled_rect(size_t x0, size_t y0, size_t x1, size_t y1, u8 clr) {
+    BOUNDCHECK_SCREEN(x0, y0);
+    BOUNDCHECK_SCREEN(x1, y1);
+    MAYBE_SWAP(x0, x1);
+    MAYBE_SWAP(y0, y1);
+
+    // Special condition. If the screen is to be completely filled, use the DMA
+    // instead.
+    if (x0 == 0 && x1 >= (SCREEN_WIDTH - 1) && y0 == 0 && y1 >= (SCREEN_HEIGHT - 1)) {
+        screen_fill(clr);
+        return;
+    }
+
+    for (size_t y = y0; y <= y1; y++) {
+        draw_hline(x0, x1, y, clr);
+    }
+}
+
+//
+// Sprites (chr/icn).
+//
+
+#if DEC_BIG_LUT == 1
+static u32 dec_byte_flip_x[256] = {
+    0x00000000, 0x00000001, 0x00000010, 0x00000011, 0x00000100,
+    0x00000101, 0x00000110, 0x00000111, 0x00001000, 0x00001001,
+    0x00001010, 0x00001011, 0x00001100, 0x00001101, 0x00001110,
+    0x00001111, 0x00010000, 0x00010001, 0x00010010, 0x00010011,
+    0x00010100, 0x00010101, 0x00010110, 0x00010111, 0x00011000,
+    0x00011001, 0x00011010, 0x00011011, 0x00011100, 0x00011101,
+    0x00011110, 0x00011111, 0x00100000, 0x00100001, 0x00100010,
+    0x00100011, 0x00100100, 0x00100101, 0x00100110, 0x00100111,
+    0x00101000, 0x00101001, 0x00101010, 0x00101011, 0x00101100,
+    0x00101101, 0x00101110, 0x00101111, 0x00110000, 0x00110001,
+    0x00110010, 0x00110011, 0x00110100, 0x00110101, 0x00110110,
+    0x00110111, 0x00111000, 0x00111001, 0x00111010, 0x00111011,
+    0x00111100, 0x00111101, 0x00111110, 0x00111111, 0x01000000,
+    0x01000001, 0x01000010, 0x01000011, 0x01000100, 0x01000101,
+    0x01000110, 0x01000111, 0x01001000, 0x01001001, 0x01001010,
+    0x01001011, 0x01001100, 0x01001101, 0x01001110, 0x01001111,
+    0x01010000, 0x01010001, 0x01010010, 0x01010011, 0x01010100,
+    0x01010101, 0x01010110, 0x01010111, 0x01011000, 0x01011001,
+    0x01011010, 0x01011011, 0x01011100, 0x01011101, 0x01011110,
+    0x01011111, 0x01100000, 0x01100001, 0x01100010, 0x01100011,
+    0x01100100, 0x01100101, 0x01100110, 0x01100111, 0x01101000,
+    0x01101001, 0x01101010, 0x01101011, 0x01101100, 0x01101101,
+    0x01101110, 0x01101111, 0x01110000, 0x01110001, 0x01110010,
+    0x01110011, 0x01110100, 0x01110101, 0x01110110, 0x01110111,
+    0x01111000, 0x01111001, 0x01111010, 0x01111011, 0x01111100,
+    0x01111101, 0x01111110, 0x01111111, 0x10000000, 0x10000001,
+    0x10000010, 0x10000011, 0x10000100, 0x10000101, 0x10000110,
+    0x10000111, 0x10001000, 0x10001001, 0x10001010, 0x10001011,
+    0x10001100, 0x10001101, 0x10001110, 0x10001111, 0x10010000,
+    0x10010001, 0x10010010, 0x10010011, 0x10010100, 0x10010101,
+    0x10010110, 0x10010111, 0x10011000, 0x10011001, 0x10011010,
+    0x10011011, 0x10011100, 0x10011101, 0x10011110, 0x10011111,
+    0x10100000, 0x10100001, 0x10100010, 0x10100011, 0x10100100,
+    0x10100101, 0x10100110, 0x10100111, 0x10101000, 0x10101001,
+    0x10101010, 0x10101011, 0x10101100, 0x10101101, 0x10101110,
+    0x10101111, 0x10110000, 0x10110001, 0x10110010, 0x10110011,
+    0x10110100, 0x10110101, 0x10110110, 0x10110111, 0x10111000,
+    0x10111001, 0x10111010, 0x10111011, 0x10111100, 0x10111101,
+    0x10111110, 0x10111111, 0x11000000, 0x11000001, 0x11000010,
+    0x11000011, 0x11000100, 0x11000101, 0x11000110, 0x11000111,
+    0x11001000, 0x11001001, 0x11001010, 0x11001011, 0x11001100,
+    0x11001101, 0x11001110, 0x11001111, 0x11010000, 0x11010001,
+    0x11010010, 0x11010011, 0x11010100, 0x11010101, 0x11010110,
+    0x11010111, 0x11011000, 0x11011001, 0x11011010, 0x11011011,
+    0x11011100, 0x11011101, 0x11011110, 0x11011111, 0x11100000,
+    0x11100001, 0x11100010, 0x11100011, 0x11100100, 0x11100101,
+    0x11100110, 0x11100111, 0x11101000, 0x11101001, 0x11101010,
+    0x11101011, 0x11101100, 0x11101101, 0x11101110, 0x11101111,
+    0x11110000, 0x11110001, 0x11110010, 0x11110011, 0x11110100,
+    0x11110101, 0x11110110, 0x11110111, 0x11111000, 0x11111001,
+    0x11111010, 0x11111011, 0x11111100, 0x11111101, 0x11111110,
+    0x11111111
+};
+
+static u32 dec_byte[256] = {
+    0x00000000, 0x10000000, 0x01000000, 0x11000000, 0x00100000,
+    0x10100000, 0x01100000, 0x11100000, 0x00010000, 0x10010000,
+    0x01010000, 0x11010000, 0x00110000, 0x10110000, 0x01110000,
+    0x11110000, 0x00001000, 0x10001000, 0x01001000, 0x11001000,
+    0x00101000, 0x10101000, 0x01101000, 0x11101000, 0x00011000,
+    0x10011000, 0x01011000, 0x11011000, 0x00111000, 0x10111000,
+    0x01111000, 0x11111000, 0x00000100, 0x10000100, 0x01000100,
+    0x11000100, 0x00100100, 0x10100100, 0x01100100, 0x11100100,
+    0x00010100, 0x10010100, 0x01010100, 0x11010100, 0x00110100,
+    0x10110100, 0x01110100, 0x11110100, 0x00001100, 0x10001100,
+    0x01001100, 0x11001100, 0x00101100, 0x10101100, 0x01101100,
+    0x11101100, 0x00011100, 0x10011100, 0x01011100, 0x11011100,
+    0x00111100, 0x10111100, 0x01111100, 0x11111100, 0x00000010,
+    0x10000010, 0x01000010, 0x11000010, 0x00100010, 0x10100010,
+    0x01100010, 0x11100010, 0x00010010, 0x10010010, 0x01010010,
+    0x11010010, 0x00110010, 0x10110010, 0x01110010, 0x11110010,
+    0x00001010, 0x10001010, 0x01001010, 0x11001010, 0x00101010,
+    0x10101010, 0x01101010, 0x11101010, 0x00011010, 0x10011010,
+    0x01011010, 0x11011010, 0x00111010, 0x10111010, 0x01111010,
+    0x11111010, 0x00000110, 0x10000110, 0x01000110, 0x11000110,
+    0x00100110, 0x10100110, 0x01100110, 0x11100110, 0x00010110,
+    0x10010110, 0x01010110, 0x11010110, 0x00110110, 0x10110110,
+    0x01110110, 0x11110110, 0x00001110, 0x10001110, 0x01001110,
+    0x11001110, 0x00101110, 0x10101110, 0x01101110, 0x11101110,
+    0x00011110, 0x10011110, 0x01011110, 0x11011110, 0x00111110,
+    0x10111110, 0x01111110, 0x11111110, 0x00000001, 0x10000001,
+    0x01000001, 0x11000001, 0x00100001, 0x10100001, 0x01100001,
+    0x11100001, 0x00010001, 0x10010001, 0x01010001, 0x11010001,
+    0x00110001, 0x10110001, 0x01110001, 0x11110001, 0x00001001,
+    0x10001001, 0x01001001, 0x11001001, 0x00101001, 0x10101001,
+    0x01101001, 0x11101001, 0x00011001, 0x10011001, 0x01011001,
+    0x11011001, 0x00111001, 0x10111001, 0x01111001, 0x11111001,
+    0x00000101, 0x10000101, 0x01000101, 0x11000101, 0x00100101,
+    0x10100101, 0x01100101, 0x11100101, 0x00010101, 0x10010101,
+    0x01010101, 0x11010101, 0x00110101, 0x10110101, 0x01110101,
+    0x11110101, 0x00001101, 0x10001101, 0x01001101, 0x11001101,
+    0x00101101, 0x10101101, 0x01101101, 0x11101101, 0x00011101,
+    0x10011101, 0x01011101, 0x11011101, 0x00111101, 0x10111101,
+    0x01111101, 0x11111101, 0x00000011, 0x10000011, 0x01000011,
+    0x11000011, 0x00100011, 0x10100011, 0x01100011, 0x11100011,
+    0x00010011, 0x10010011, 0x01010011, 0x11010011, 0x00110011,
+    0x10110011, 0x01110011, 0x11110011, 0x00001011, 0x10001011,
+    0x01001011, 0x11001011, 0x00101011, 0x10101011, 0x01101011,
+    0x11101011, 0x00011011, 0x10011011, 0x01011011, 0x11011011,
+    0x00111011, 0x10111011, 0x01111011, 0x11111011, 0x00000111,
+    0x10000111, 0x01000111, 0x11000111, 0x00100111, 0x10100111,
+    0x01100111, 0x11100111, 0x00010111, 0x10010111, 0x01010111,
+    0x11010111, 0x00110111, 0x10110111, 0x01110111, 0x11110111,
+    0x00001111, 0x10001111, 0x01001111, 0x11001111, 0x00101111,
+    0x10101111, 0x01101111, 0x11101111, 0x00011111, 0x10011111,
+    0x01011111, 0x11011111, 0x00111111, 0x10111111, 0x01111111,
+    0x11111111
+};
+
+IWRAM_CODE
+static inline
+u32
+decode_1bpp(u8 row, u8 flip_x) {
+    if (flip_x) {
+        return dec_byte_flip_x[row];
+    }
+    return dec_byte[row];
+}
+#else
+static u16 dec_nibble[] = {
+    0x0000, 0x1000, 0x0100, 0x1100,
+    0x0010, 0x1010, 0x0110, 0x1110,
+    0x0001, 0x1001, 0x0101, 0x1101,
+    0x0011, 0x1011, 0x0111, 0x1111,
+};
+
+static u16 dec_nibble_flip_x[] = {
+    0x0000, 0x0001, 0x0010, 0x0011,
+    0x0100, 0x0101, 0x0110, 0x0111,
+    0x1000, 0x1001, 0x1010, 0x1011,
+    0x1100, 0x1101, 0x1110, 0x1111,
+};
+
+IWRAM_CODE
+static inline
+u32
+decode_1bpp(u8 row, u8 flip_x) {
+    if (flip_x) {
+        u16 *lut = dec_nibble_flip_x;
+        return (u32)lut[(row >> 4) & 0xF] << 16 | (u32)lut[(row >> 0) & 0xF];
+    }
+    u16 *lut = dec_nibble;
+    return (u32)lut[(row >> 0) & 0xF] << 16 | (u32)lut[(row >> 4) & 0xF];
+}
+#endif
+
+IWRAM_CODE
+UNROLL_LOOPS
+void
+draw_chr(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y) {
+    BOUNDCHECK_SCREEN(x, y);
+    size_t tile_x0 = x / 8;
+    size_t tile_x1 = (x + 7) / 8;
+    size_t tile_y = y / 8;
+    size_t start_col = x % 8;
+    size_t start_row = y % 8;
+    size_t shift_left = start_col * 4;
+    size_t shift_right = (8 - start_col) * 4;
+    u32 dirty = (1 << tile_x0) | (1 << tile_x1);
+    u32 *dst = &backbuf[start_row + (tile_x0 + tile_y * 32) * 8];
+#if DEC_BIG_LUT
+    u32 *lut = flip_x ? dec_byte_flip_x : dec_byte;
+#endif
+    if (!flip_y) {
+        for(size_t v = 0; v < 8; v++, dst++) {
+            if ((y + v) >= SCREEN_HEIGHT) break;
+            u8 ch1 = sprite[v + 0];
+            u8 ch2 = sprite[v + 8];
+#if DEC_BIG_LUT
+            u32 clr_a = lut[ch1];
+            u32 clr_b = lut[ch2];
+#else
+            u32 clr_a = decode_1bpp(ch1, flip_x);
+            u32 clr_b = decode_1bpp(ch2, flip_x);
+#endif
+            u32 mask_a = (clr_a * 0xF);
+            u32 mask_b = (clr_b * 0xF);
+            u32 mask = (mask_a | mask_b);
+            u32 color;
+            if (clr == 0) {
+                color = clr_a + (clr_b << 1);
+            } else if (clr == 15) {
+                color = 0;
+            } else {
+                color = (clr_a | clr_b) * clr;
+            }
+            dst[0] = (dst[0] & ~(mask << shift_left))  | (color << shift_left);
+            dst[8] = (dst[8] & ~(mask >> shift_right)) | (color >> shift_right);
+            if ((start_row + v)  == 7) {
+                dirty_tiles[tile_y + 1] |= dirty;
+                dst += (32 - 1) * 8;
+            }
+        }
+    } else {
+        for(size_t v = 0; v < 8; v++, dst++) {
+            if ((y + v) >= SCREEN_HEIGHT) break;
+            u8 ch1 = sprite[(7 - v) + 0];
+            u8 ch2 = sprite[(7 - v) + 8];
+#if DEC_BIG_LUT
+            u32 clr_a = lut[ch1];
+            u32 clr_b = lut[ch2];
+#else
+            u32 clr_a = decode_1bpp(ch1, flip_x);
+            u32 clr_b = decode_1bpp(ch2, flip_x);
+#endif
+            u32 mask_a = (clr_a * 0xF);
+            u32 mask_b = (clr_b * 0xF);
+            u32 mask = (mask_a | mask_b);
+            u32 color;
+            if (clr == 0) {
+                color = clr_a + (clr_b << 1);
+            } else if (clr == 15) {
+                color = 0;
+            } else {
+                color = (clr_a | clr_b) * clr;
+            }
+            dst[0] = (dst[0] & ~(mask << shift_left))  | (color << shift_left);
+            dst[8] = (dst[8] & ~(mask >> shift_right)) | (color >> shift_right);
+            if ((start_row + v)  == 7) {
+                dirty_tiles[tile_y + 1] |= dirty;
+                dst += (32 - 1) * 8;
+            }
+        }
+    }
+    dirty_tiles[tile_y] |= dirty;
+}
+
+IWRAM_CODE
+UNROLL_LOOPS
+void
+draw_icn(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y) {
+    BOUNDCHECK_SCREEN(x, y);
+    size_t tile_x0 = x / 8;
+    size_t tile_x1 = (x + 7) / 8;
+    size_t tile_y = y / 8;
+    size_t start_col = x % 8;
+    size_t start_row = y % 8;
+    size_t shift_left = start_col * 4;
+    size_t shift_right = (8 - start_col) * 4;
+    u32 dirty = (1 << tile_x0) | (1 << tile_x1);
+    u32 *dst = &backbuf[start_row + (tile_x0 + tile_y * 32) * 8];
+#if DEC_BIG_LUT
+    u32 *lut = flip_x ? dec_byte_flip_x : dec_byte;
+#endif
+    if (!flip_y) {
+        for(size_t v = 0; v < 8; v++, dst++) {
+            if ((y + v) >= SCREEN_HEIGHT) break;
+            u8 ch1 = sprite[v + 0];
+#if DEC_BIG_LUT
+            u32 color = lut[ch1];
+#else
+            u32 color = decode_1bpp(ch1, flip_x);
+#endif
+            u32 mask = (color * 0xF);
+            color *= clr;
+            dst[0] = (dst[0] & ~(mask << shift_left))  | (color << shift_left);
+            dst[8] = (dst[8] & ~(mask >> shift_right)) | (color >> shift_right);
+            if ((start_row + v)  == 7) {
+                dirty_tiles[tile_y + 1] |= dirty;
+                dst += (32 - 1) * 8;
+            }
+        }
+    } else {
+        for(size_t v = 0; v < 8; v++, dst++) {
+            if ((y + v) >= SCREEN_HEIGHT) break;
+            u8 ch1 = sprite[(7 - v) + 0];
+#if DEC_BIG_LUT
+            u32 color = lut[ch1];
+#else
+            u32 color = decode_1bpp(ch1, flip_x);
+#endif
+            u32 mask = (color * 0xF);
+            color *= clr;
+            dst[0] = (dst[0] & ~(mask << shift_left))  | (color << shift_left);
+            dst[8] = (dst[8] & ~(mask >> shift_right)) | (color >> shift_right);
+            if ((start_row + v)  == 7) {
+                dirty_tiles[tile_y + 1] |= dirty;
+                dst += (32 - 1) * 8;
+            }
+        }
+    }
+    dirty_tiles[tile_y] |= dirty;
+}
+
+//
+// Flipping buffers/copying memory.
+//
+
+IWRAM_CODE
+void
+flip_buffer(void) {
+// Mode 0: double buffering without dirty tiles.
+#if FLIP_TYPE == 0
+    if (backbuf == BUF_0) {
+        backbuf = BUF_1;
+        BG_H_SCROLL_0 = 0;
+        BG_H_SCROLL_1 = -240;
+    } else {
+        backbuf = BUF_0;
+        BG_H_SCROLL_0 = -240;
+        BG_H_SCROLL_1 = 0;
+    }
+
+// Mode 1: single buffer, copy the dirty lines from backbuffer (BUF_1) to
+// frontbuffer (BUF_0) using the DMA.
+#elif FLIP_TYPE == 1
+    u32 *front = BUF_0;
+    u32 *back = BUF_1;
+    BG_H_SCROLL_0 = 0;
+    BG_H_SCROLL_1 = -240;
+    for (size_t j = 0; j < 20; ++j) {
+        if (dirty_tiles[j] == 0) {
+            continue;
+        }
+        u32 offset = j * 32 * 8;
+        dma_copy(front + offset, back + offset, (30 * 8 * 4), 3);
+        dirty_tiles[j] = 0;
+    }
+
+// Mode 2: single buffer, copy the dirty tiles from backbuffer (BUF_1) to
+// frontbuffer (BUF_0).
+#elif FLIP_TYPE == 2
+    u32 *front = BUF_0;
+    u32 *back = BUF_1;
+    BG_H_SCROLL_0 = 0;
+    BG_H_SCROLL_1 = -240;
+    for (size_t j = 0; j < 20; ++j) {
+        if (dirty_tiles[j] == 0) {
+            continue;
+        }
+        size_t k = 1;
+        for (size_t i = 0; i < 30; ++i, k <<= 1) {
+            if (dirty_tiles[j] & k) {
+                Tile *mem_front = front;
+                Tile *mem_back = back;
+                mem_front[i + j * 32] = mem_back[i + j * 32];
+            }
+        }
+        dirty_tiles[j] = 0;
+    }
+
+// Mode 3: Double buffering with dirty line, copying the dirty lines if needed
+// after flipping buffers with the DMA.
+#elif FLIP_TYPE == 3
+    bool should_flip = false;
+    for (size_t j = 0; j < 20; ++j) {
+        if (dirty_tiles[j] == 0) {
+            continue;
+        }
+        should_flip = true;
+        break;
+    }
+    if (!should_flip) {
+        return;
+    }
+    u32 *frontbuf = backbuf;
+    if (backbuf == BUF_0) {
+        backbuf = BUF_1;
+        BG_H_SCROLL_0 = 0;
+        BG_H_SCROLL_1 = -240;
+    } else {
+        backbuf = BUF_0;
+        BG_H_SCROLL_0 = -240;
+        BG_H_SCROLL_1 = 0;
+    }
+    for (size_t j = 0; j < 20; ++j) {
+        if (dirty_tiles[j] == 0) {
+            continue;
+        }
+        u32 offset = j * 32 * 8;
+        dma_copy(backbuf + offset, frontbuf + offset, (30 * 8 * 4), 3);
+        dirty_tiles[j] = 0;
+    }
+
+// Mode 4: Double buffering with dirty tiles, copying the dirty tiles if needed
+// after flipping buffers.
+#elif FLIP_TYPE == 4
+    bool should_flip = false;
+    for (size_t j = 0; j < 20; ++j) {
+        if (dirty_tiles[j] == 0) {
+            continue;
+        }
+        should_flip = true;
+        break;
+    }
+    if (!should_flip) {
+        return;
+    }
+    u32 *frontbuf = backbuf;
+    if (backbuf == BUF_0) {
+        backbuf = BUF_1;
+        BG_H_SCROLL_0 = 0;
+        BG_H_SCROLL_1 = -240;
+    } else {
+        backbuf = BUF_0;
+        BG_H_SCROLL_0 = -240;
+        BG_H_SCROLL_1 = 0;
+    }
+    for (size_t j = 0; j < 20; ++j) {
+        if (dirty_tiles[j] == 0) {
+            continue;
+        }
+        size_t k = 1;
+        for (size_t i = 0; i < 30; ++i, k <<= 1) {
+            if (dirty_tiles[j] & k) {
+                Tile *mem_front = frontbuf;
+                Tile *mem_back = backbuf;
+                mem_back[i + j * 32] = mem_front[i + j * 32];
+            }
+        }
+        dirty_tiles[j] = 0;
+    }
+#endif
+}
+
+//
+// Text rendering.
+//
+
+#include "font.h"
+
+// Font rendering function for the text engine.
+void
+txt_drawc(char c, size_t x, size_t y, u8 clr) {
+    u8 *tile = font_icn;
+    draw_icn(x, y, tile + 8 * c, clr, 1, 0);
+}
+
+//
+// Initialization.
+//
+
+void
+renderer_init(void) {
+    // Initialize display mode and bg palette.
+    DISP_CTRL = DISP_MODE_0 | DISP_BG_0 | DISP_BG_1;
+
+    // Clear VRAM.
+    dma_fill((u32*)MEM_VRAM, 0, KB(96), 3);
+
+    // Initialize backgrounds.
+    BG_CTRL(0) = BG_CHARBLOCK(CB_0) | BG_SCREENBLOCK(SB_0) | BG_PRIORITY(0) | BG_SIZE(1);
+    BG_CTRL(1) = BG_CHARBLOCK(CB_1) | BG_SCREENBLOCK(SB_1) | BG_PRIORITY(1) | BG_SIZE(1);
+
+    // Initialize background memory map for the render buffers. The backgrounds
+    // are 64x32 each, with the second screenblock pointing to a zeroed tile.
+    // This makes it so while scrolling the backgrounds to the second screen we
+    // effectively disabling them. Thanks to this we can perform double
+    // buffering with mode 0 rendering.
+    u16 *mem_map_fg = SCREENBLOCK_MEM[SB_0];
+    u16 *mem_map_fg_blank = SCREENBLOCK_MEM[SB_0 + 1];
+    u16 *mem_map_bg = SCREENBLOCK_MEM[SB_1];
+    u16 *mem_map_bg_blank = SCREENBLOCK_MEM[SB_1 + 1];
+    for (size_t i = 0; i < 32 * 20; ++i) {
+        mem_map_fg[i] = i;
+        mem_map_fg_blank[i] = 32 * 20 - 1;
+        mem_map_bg[i] = i + 32 * 4;
+        mem_map_bg_blank[i] = (32 * 20 - 1) +  32 * 4;
+    }
+
+    // Setup initial background state.
+    BG_H_SCROLL_0 = -240;
+    BG_H_SCROLL_1 = -240;
+
+    // Initialize default palette.
+    PAL_BUFFER_BG[0] = COLOR_BLACK;
+    PAL_BUFFER_BG[1] = COLOR_WHITE;
+    PAL_BUFFER_BG[2] = COLOR_RED;
+    PAL_BUFFER_BG[3] = COLOR_BLUE;
+    PAL_BUFFER_BG[4] = COLOR_CYAN;
+    PAL_BUFFER_BG[5] = COLOR_GREY;
+    PAL_BUFFER_BG[6] = COLOR_WHITE;
+    PAL_BUFFER_BG[7] = COLOR_GREEN;
+    PAL_BUFFER_BG[8] = COLOR_PURPLE;
+
+    // Initialize text engine.
+    txt_init(txt_drawc);
+}
diff --git a/src/sequencer.c b/src/sequencer.c
index b582cf4..ff3904d 100644
--- a/src/sequencer.c
+++ b/src/sequencer.c
@@ -527,9 +527,9 @@ draw_channels(void) {
         }
         u8 clr = active ? COL_FG : COL_GREY;
         size_t y = CHAN_START_Y + i * CHAN_OFFSET_Y;
-        draw_tile(CHAN_START_X, y, channel_tiles + k++, clr, false);
-        draw_tile(CHAN_START_X + 8, y, channel_tiles + k++, clr, false);
-        draw_tile(CHAN_START_X + 16, y, channel_tiles + k++, clr, false);
+        // draw_tile(CHAN_START_X, y, channel_tiles + k++, clr, false);
+        // draw_tile(CHAN_START_X + 8, y, channel_tiles + k++, clr, false);
+        // draw_tile(CHAN_START_X + 16, y, channel_tiles + k++, clr, false);
     }
 }
 
@@ -582,8 +582,8 @@ draw_trigger(size_t chan, size_t i) {
         size_t y = TRIG_START_Y + offset_y;
         Tile *tiles = ASSETS_NOTE_NAMES;
         tiles += 2 * trig.note;
-        draw_tile(x, y, tiles, COL_FG, true);
-        draw_tile(x + 8, y, tiles + 1, COL_FG, true);
+        // draw_tile(x, y, tiles, COL_FG, true);
+        // draw_tile(x + 8, y, tiles + 1, COL_FG, true);
     } else {
         clear_trigger(i);
     }
@@ -658,7 +658,7 @@ void
 draw_bank_buttons() {
     size_t x = BANK_START_X;
     size_t y = BANK_START_Y;
-    txt_drawf_small("BANK", x - 2, y - 10, 4, COL_FG);
+    // txt_drawf_small("BANK", x - 2, y - 10, 4, COL_FG);
     char bank_names[] = {
         'A', 'B', 'C', 'D',
     };
@@ -669,7 +669,7 @@ draw_bank_buttons() {
         }
         draw_filled_rect(x, y, x + PAT_W, y + PAT_H, COL_BG);
         draw_rect(x, y, x + PAT_W, y + PAT_H, color);
-        txt_drawc(bank_names[i], x + 4, y + 2, 6, color);
+        // txt_drawc(bank_names[i], x + 4, y + 2, 6, color);
         y += PAT_OFFSET_Y;
     }
 }
@@ -678,7 +678,7 @@ void
 draw_pattern_buttons() {
     size_t x = PAT_START_X;
     size_t y = PAT_START_Y;
-    txt_drawf_small("PAT", x, y - 10, 4, COL_FG);
+    // txt_drawf_small("PAT", x, y - 10, 4, COL_FG);
     char pat_names[] = {
         'A', 'B', 'C', 'D',
         'E', 'F', 'G', 'H',
@@ -693,7 +693,7 @@ draw_pattern_buttons() {
         }
         draw_filled_rect(x, y, x + PAT_W, y + PAT_H, COL_BG);
         draw_rect(x, y, x + PAT_W, y + PAT_H, color);
-        txt_drawc(pat_names[i], x + 4, y + 2, 6, color);
+        // txt_drawc(pat_names[i], x + 4, y + 2, 6, color);
         y += PAT_OFFSET_Y;
     }
 }
@@ -745,7 +745,7 @@ draw_bpm() {
     draw_filled_rect(x, y, x + R_COL_W, y + BPM_H, COL_BG);
     draw_rect(x, y, x + R_COL_W, y + BPM_H, COL_FG);
     draw_line(x + 5, y, x + 19, y, COL_BG);
-    txt_drawf_small("BPM", x + 5, y - 4, 4, COL_FG);
+    // txt_drawf_small("BPM", x + 5, y - 4, 4, COL_FG);
 
     // Make sure its horizontally centered if only 2 digits
     int bpm = patterns[pattern_selection_loc].bpm;
@@ -1069,24 +1069,24 @@ draw_parameters_wave(void) {
 
         // Wave text.
         x -= 2;
-        txt_drawf_small("%02x%02x%02x%02x", x, y + 20, 4, COL_FG,
-                wave_a[0], wave_a[1], wave_a[2], wave_a[3]);
-        txt_drawf_small("%02x%02x%02x%02x", x + 34, y + 20, 4, COL_FG,
-                wave_a[4], wave_a[5], wave_a[6], wave_a[7]);
-        txt_drawf_small("%02x%02x%02x%02x", x, y + 28, 4, COL_FG,
-                wave_a[8], wave_a[9], wave_a[10], wave_a[11]);
-        txt_drawf_small("%02x%02x%02x%02x", x + 34, y + 28, 4, COL_FG,
-                wave_a[12], wave_a[13], wave_a[14], wave_a[15]);
+        // txt_drawf_small("%02x%02x%02x%02x", x, y + 20, 4, COL_FG,
+        //         wave_a[0], wave_a[1], wave_a[2], wave_a[3]);
+        // txt_drawf_small("%02x%02x%02x%02x", x + 34, y + 20, 4, COL_FG,
+        //         wave_a[4], wave_a[5], wave_a[6], wave_a[7]);
+        // txt_drawf_small("%02x%02x%02x%02x", x, y + 28, 4, COL_FG,
+        //         wave_a[8], wave_a[9], wave_a[10], wave_a[11]);
+        // txt_drawf_small("%02x%02x%02x%02x", x + 34, y + 28, 4, COL_FG,
+        //         wave_a[12], wave_a[13], wave_a[14], wave_a[15]);
 
         x += 70;
-        txt_drawf_small("%02x%02x%02x%02x", x, y + 20, 4, COL_FG,
-                wave_b[0], wave_b[1], wave_b[2], wave_b[3]);
-        txt_drawf_small("%02x%02x%02x%02x", x + 34, y + 20, 4, COL_FG,
-                wave_b[4], wave_b[5], wave_b[6], wave_b[7]);
-        txt_drawf_small("%02x%02x%02x%02x", x, y + 28, 4, COL_FG,
-                wave_b[8], wave_b[9], wave_b[10], wave_b[11]);
-        txt_drawf_small("%02x%02x%02x%02x", x + 34, y + 28, 4, COL_FG,
-                wave_b[12], wave_b[13], wave_b[14], wave_b[15]);
+        // txt_drawf_small("%02x%02x%02x%02x", x, y + 20, 4, COL_FG,
+        //         wave_b[0], wave_b[1], wave_b[2], wave_b[3]);
+        // txt_drawf_small("%02x%02x%02x%02x", x + 34, y + 20, 4, COL_FG,
+        //         wave_b[4], wave_b[5], wave_b[6], wave_b[7]);
+        // txt_drawf_small("%02x%02x%02x%02x", x, y + 28, 4, COL_FG,
+        //         wave_b[8], wave_b[9], wave_b[10], wave_b[11]);
+        // txt_drawf_small("%02x%02x%02x%02x", x + 34, y + 28, 4, COL_FG,
+        //         wave_b[12], wave_b[13], wave_b[14], wave_b[15]);
     }
 
     // Draw default wave buttons.
@@ -1095,12 +1095,12 @@ draw_parameters_wave(void) {
         size_t x = PARAMS_START_X;
         size_t y = PARAMS_START_Y + PARAMS_H - 12;
         for (size_t i = 0, k = 0; i < 4 * 2; i += 2, k++) {
-            draw_tile(x + 17 * k,     y, wave_tiles + i, COL_FG, true);
-            draw_tile(x + 17 * k + 8, y, wave_tiles + i + 1, COL_FG, true);
+            // draw_tile(x + 17 * k,     y, wave_tiles + i, COL_FG, true);
+            // draw_tile(x + 17 * k + 8, y, wave_tiles + i + 1, COL_FG, true);
         }
         for (size_t i = 0, k = 0; i < 4 * 2; i += 2, k++) {
-            draw_tile(x + 17 * k + 70, y, wave_tiles + i, COL_FG, true);
-            draw_tile(x + 17 * k + 8 + 70, y, wave_tiles + i + 1, COL_FG, true);
+            // draw_tile(x + 17 * k + 70, y, wave_tiles + i, COL_FG, true);
+            // draw_tile(x + 17 * k + 8 + 70, y, wave_tiles + i + 1, COL_FG, true);
         }
     }
 
@@ -1113,7 +1113,7 @@ draw_parameters_wave(void) {
         draw_line(x, y + 5, x, y + 16, COL_FG);
         draw_line(x + 30, y + 5, x + 30, y + 17, COL_FG);
         draw_line(x, y + 17, x + 30, y + 17, COL_FG);
-        txt_drawf_small("mode", x + 6, y, 4, COL_FG);
+        // txt_drawf_small("mode", x + 6, y, 4, COL_FG);
 
         switch (pat->ch3.params[trig_selection_loc].wave_mode) {
             case 0: {
@@ -1137,7 +1137,7 @@ draw_parameters_wave(void) {
         draw_line(x, y + 8, x, y + 19, COL_FG);
         draw_line(x + 30, y + 8, x + 30, y + 19, COL_FG);
         draw_line(x, y + 20, x + 30, y + 20, COL_FG);
-        txt_drawf_small("vol", x + 8, y + 3, 4, COL_FG);
+        // txt_drawf_small("vol", x + 8, y + 3, 4, COL_FG);
 
         switch (pat->ch3.params[trig_selection_loc].wave_volume) {
             case 0: {
@@ -1232,7 +1232,7 @@ draw_parameters_square(ChannelSquareParams *params, bool sweep) {
             draw_line(x, y + 8, x, y + 19, COL_FG);
             draw_line(x + 24, y + 8, x + 24, y + 19, COL_FG);
             draw_line(x, y + 20, x + 24, y + 20, COL_FG);
-            txt_drawf_small("duty", x + 3, y + 3, 4, COL_FG);
+            // txt_drawf_small("duty", x + 3, y + 3, 4, COL_FG);
 
             switch (params->duty_cycle) {
                 case 0: {
@@ -1291,7 +1291,7 @@ draw_parameters_square(ChannelSquareParams *params, bool sweep) {
             draw_line(x, y + 8, x, y + 19, COL_FG);
             draw_line(x + 24, y + 8, x + 24, y + 19, COL_FG);
             draw_line(x, y + 20, x + 24, y + 20, COL_FG);
-            txt_drawf_small("vol", x + 5, y + 3, 4, COL_FG);
+            // txt_drawf_small("vol", x + 5, y + 3, 4, COL_FG);
 
             switch (params->env_volume) {
                 case 0: {
@@ -1354,7 +1354,7 @@ draw_parameters_square(ChannelSquareParams *params, bool sweep) {
             draw_line(x, y + 8, x, y + 19, COL_FG);
             draw_line(x + 24, y + 8, x + 24, y + 19, COL_FG);
             draw_line(x, y + 20, x + 24, y + 20, COL_FG);
-            txt_drawf_small("dir", x + 5, y + 3, 4, COL_FG);
+            // txt_drawf_small("dir", x + 5, y + 3, 4, COL_FG);
 
             char arr_up[2] = { 0x19, 0 };
             char arr_down[2] = { 0x18, 0 };
@@ -1377,7 +1377,7 @@ draw_parameters_square(ChannelSquareParams *params, bool sweep) {
             draw_line(x, y + 8, x, y + 19, COL_FG);
             draw_line(x + 24, y + 8, x + 24, y + 19, COL_FG);
             draw_line(x, y + 20, x + 24, y + 20, COL_FG);
-            txt_drawf_small("time", x + 3, y + 3, 4, COL_FG);
+            // txt_drawf_small("time", x + 3, y + 3, 4, COL_FG);
 
             switch (params->env_time) {
                 case 0: {
@@ -1417,7 +1417,7 @@ draw_parameters_square(ChannelSquareParams *params, bool sweep) {
         draw_line(x, y + 8, x, y + 19, COL_FG);
         draw_line(x + 24, y + 8, x + 24, y + 19, COL_FG);
         draw_line(x, y + 20, x + 24, y + 20, COL_FG);
-        txt_drawf_small("num", x + 5, y + 3, 4, COL_FG);
+        // txt_drawf_small("num", x + 5, y + 3, 4, COL_FG);
 
         switch (params->sweep_number) {
             case 0: {
@@ -1456,7 +1456,7 @@ draw_parameters_square(ChannelSquareParams *params, bool sweep) {
         draw_line(x, y + 8, x, y + 19, COL_FG);
         draw_line(x + 24, y + 8, x + 24, y + 19, COL_FG);
         draw_line(x, y + 20, x + 24, y + 20, COL_FG);
-        txt_drawf_small("time", x + 3, y + 3, 4, COL_FG);
+        // txt_drawf_small("time", x + 3, y + 3, 4, COL_FG);
 
         switch (params->sweep_time) {
             case 0: {
@@ -1495,7 +1495,7 @@ draw_parameters_square(ChannelSquareParams *params, bool sweep) {
         draw_line(x, y + 8, x, y + 19, COL_FG);
         draw_line(x + 24, y + 8, x + 24, y + 19, COL_FG);
         draw_line(x, y + 20, x + 24, y + 20, COL_FG);
-        txt_drawf_small("dir", x + 5, y + 3, 4, COL_FG);
+        // txt_drawf_small("dir", x + 5, y + 3, 4, COL_FG);
 
         char arr_up[2] = { 0x19, 0 };
         char arr_down[2] = { 0x18, 0 };
@@ -1513,10 +1513,10 @@ draw_parameters_square(ChannelSquareParams *params, bool sweep) {
     {
         size_t x = PARAMS_START_X + x_offset;
         size_t y = PARAMS_START_Y + PARAMS_H - 45;
-        txt_drawf_small("shape", x + 1, y - 12, 4, COL_FG);
-        txt_drawf_small("envelope", x + 54, y - 12, 4, COL_FG);
+        // txt_drawf_small("shape", x + 1, y - 12, 4, COL_FG);
+        // txt_drawf_small("envelope", x + 54, y - 12, 4, COL_FG);
         if (sweep) {
-            txt_drawf_small("sweep", x + 133, y - 12, 4, COL_FG);
+            // txt_drawf_small("sweep", x + 133, y - 12, 4, COL_FG);
         }
     }
 }
@@ -1538,7 +1538,7 @@ draw_parameters_noise(void) {
             draw_line(x, y + 8, x, y + 19, COL_FG);
             draw_line(x + 24, y + 8, x + 24, y + 19, COL_FG);
             draw_line(x, y + 20, x + 24, y + 20, COL_FG);
-            txt_drawf_small("mode", x + 3, y + 3, 4, COL_FG);
+            // txt_drawf_small("mode", x + 3, y + 3, 4, COL_FG);
 
             switch (params->bit_mode) {
                 case 0: {
@@ -1591,7 +1591,7 @@ draw_parameters_noise(void) {
             draw_line(x, y + 8, x, y + 19, COL_FG);
             draw_line(x + 24, y + 8, x + 24, y + 19, COL_FG);
             draw_line(x, y + 20, x + 24, y + 20, COL_FG);
-            txt_drawf_small("vol", x + 5, y + 3, 4, COL_FG);
+            // txt_drawf_small("vol", x + 5, y + 3, 4, COL_FG);
 
             switch (params->env_volume) {
                 case 0: {
@@ -1654,7 +1654,7 @@ draw_parameters_noise(void) {
             draw_line(x, y + 8, x, y + 19, COL_FG);
             draw_line(x + 24, y + 8, x + 24, y + 19, COL_FG);
             draw_line(x, y + 20, x + 24, y + 20, COL_FG);
-            txt_drawf_small("dir", x + 5, y + 3, 4, COL_FG);
+            // txt_drawf_small("dir", x + 5, y + 3, 4, COL_FG);
 
             char arr_up[2] = { 0x19, 0 };
             char arr_down[2] = { 0x18, 0 };
@@ -1677,7 +1677,7 @@ draw_parameters_noise(void) {
             draw_line(x, y + 8, x, y + 19, COL_FG);
             draw_line(x + 24, y + 8, x + 24, y + 19, COL_FG);
             draw_line(x, y + 20, x + 24, y + 20, COL_FG);
-            txt_drawf_small("time", x + 3, y + 3, 4, COL_FG);
+            // txt_drawf_small("time", x + 3, y + 3, 4, COL_FG);
 
             switch (params->env_time) {
                 case 0: {
@@ -1712,7 +1712,7 @@ draw_parameters_noise(void) {
     {
         size_t x = PARAMS_START_X + x_offset;
         size_t y = PARAMS_START_Y + PARAMS_H - 45;
-        txt_drawf_small("envelope", x + 54, y - 12, 4, COL_FG);
+        // txt_drawf_small("envelope", x + 54, y - 12, 4, COL_FG);
     }
 }
 
diff --git a/src/text/font.h b/src/text/font.h
index daf0c24..e4b089c 100644
--- a/src/text/font.h
+++ b/src/text/font.h
@@ -128,3 +128,134 @@ static const u32 font[] = {
     0x00000000, 0x00000000, 0x00000000, 0x00000000,
     0x00000000, 0x00000000, 0x00000000, 0x00000000,
 };
+
+u32 font_icn[] = {
+    0x00000000, 0x00000000, 0x00240000, 0x00003c42,
+    0x00240000, 0x0000423c, 0x3e360000, 0x00081c3e,
+    0x3e1c0800, 0x00081c3e, 0x3e1c1c00, 0x1c083e36,
+    0x3e1c0800, 0x1c083e3e, 0x18000000, 0x00000018,
+    0x667e7e7e, 0x7e7e7e66, 0x24180000, 0x00001824,
+    0x5a667e7e, 0x7e7e665a, 0x3e1c0800, 0x1c221c08,
+    0x1c221c00, 0x083e0808, 0x28381800, 0x0c0c0808,
+    0x24243c00, 0x06363424, 0x08220800, 0x0822081c,
+    0x3c1c0c04, 0x00040c1c, 0x3c383020, 0x00203038,
+    0x083e1c08, 0x081c3e08, 0x14141400, 0x00140014,
+    0x2a2a3c00, 0x2828282c, 0x3c043800, 0x1e203c42,
+    0x00000000, 0x0000007e, 0x083e1c08, 0x3e081c3e,
+    0x3e1c0800, 0x00080808, 0x08080800, 0x00081c3e,
+    0x30100000, 0x0010307e, 0x0c080000, 0x00080c7e,
+    0x00000000, 0x3e020200, 0x36140000, 0x0014367f,
+    0x1c080800, 0x003e3e1c, 0x1c3e3e00, 0x0008081c,
+    0x00000000, 0x00000000, 0x08080800, 0x00080008,
+    0x14141400, 0x00000000, 0x3e140000, 0x00143e14,
+    0x021c0800, 0x081e201c, 0x16260000, 0x00323408,
+    0x08140800, 0x002c1234, 0x08080800, 0x00000000,
+    0x04080000, 0x00080404, 0x10080000, 0x00081010,
+    0x08140000, 0x0014083e, 0x08080000, 0x0008083e,
+    0x00000000, 0x04080c00, 0x00000000, 0x0000003e,
+    0x00000000, 0x000c0c00, 0x18302000, 0x0002060c,
+    0x221c0000, 0x001c222a, 0x0a0c0000, 0x003e0808,
+    0x201c0000, 0x003e021c, 0x221c0000, 0x001c2218,
+    0x12120000, 0x00103e12, 0x023e0000, 0x001e201e,
+    0x021c0000, 0x001c221e, 0x203e0000, 0x00040810,
+    0x221c0000, 0x001c221c, 0x221c0000, 0x001c203c,
+    0x0c0c0000, 0x000c0c00, 0x0c0c0000, 0x04080c00,
+    0x08100000, 0x00100804, 0x3e000000, 0x00003e00,
+    0x08040000, 0x00040810, 0x10221c00, 0x00080008,
+    0x2a1c0000, 0x001c023a, 0x221c0000, 0x00223e22,
+    0x221e0000, 0x001e221e, 0x221c0000, 0x001c2202,
+    0x221e0000, 0x001e2222, 0x023e0000, 0x003e021e,
+    0x023e0000, 0x0002021e, 0x021c0000, 0x001c2232,
+    0x22220000, 0x0022223e, 0x083e0000, 0x003e0808,
+    0x20200000, 0x001c2220, 0x12220000, 0x0022120e,
+    0x02020000, 0x003e0202, 0x36220000, 0x0022222a,
+    0x2a260000, 0x00222232, 0x221c0000, 0x001c2222,
+    0x221e0000, 0x00021e22, 0x221c0000, 0x002c1222,
+    0x221e0000, 0x00221e22, 0x021c0000, 0x001e201c,
+    0x083e0000, 0x00080808, 0x22220000, 0x001c2222,
+    0x22220000, 0x00081422, 0x2a220000, 0x00142a2a,
+    0x14220000, 0x00221408, 0x22220000, 0x00080814,
+    0x103e0000, 0x003e0408, 0x041c0000, 0x001c0404,
+    0x0c060200, 0x00203018, 0x101c0000, 0x001c1010,
+    0x22140800, 0x00000000, 0x00000000, 0x003e0000,
+    0x10080400, 0x00000000, 0x201c0000, 0x003c223c,
+    0x1e020200, 0x001e2222, 0x021c0000, 0x001c0202,
+    0x3c202000, 0x003c2222, 0x221c0000, 0x001c021e,
+    0x023c0000, 0x0002021e, 0x223c0000, 0x1c203c22,
+    0x1e020200, 0x00222222, 0x0c000800, 0x003e0808,
+    0x20002000, 0x1c222020, 0x0a120200, 0x00320a06,
+    0x08080c00, 0x003e0808, 0x2a160000, 0x002a2a2a,
+    0x320e0000, 0x00222222, 0x221c0000, 0x001c2222,
+    0x221e0000, 0x02021e22, 0x223c0000, 0x70203c22,
+    0x261a0000, 0x00020202, 0x021c0000, 0x001e201c,
+    0x043e0400, 0x00180404, 0x22220000, 0x001c2222,
+    0x22220000, 0x00081422, 0x22220000, 0x00142a2a,
+    0x14220000, 0x00221408, 0x22220000, 0x1c203c22,
+    0x103e0000, 0x003e0408, 0x041c0000, 0x001c0406,
+    0x08080000, 0x00080808, 0x101c0000, 0x001c1030,
+    0x1a2c0000, 0x00000000, 0x14080000, 0x003e2222,
+    0x00000000, 0x00000000, 0x08080000, 0x00080008,
+    0x14140000, 0x00000000, 0x3e140000, 0x00143e14,
+    0x041c0800, 0x081c101c, 0x10040000, 0x00100408,
+    0x14080000, 0x00381438, 0x08080000, 0x00000000,
+    0x04080000, 0x00080404, 0x10080000, 0x00081010,
+    0x08140000, 0x0014081c, 0x08000000, 0x0000081c,
+    0x00000000, 0x04080000, 0x00000000, 0x0000001c,
+    0x00000000, 0x00080000, 0x18100000, 0x00040c08,
+    0x141c0000, 0x001c1414, 0x0c080000, 0x001c0808,
+    0x101c0000, 0x001c041c, 0x101c0000, 0x001c1018,
+    0x14140000, 0x0010101c, 0x041c0000, 0x001c101c,
+    0x041c0000, 0x001c141c, 0x101c0000, 0x00040408,
+    0x141c0000, 0x001c141c, 0x141c0000, 0x0010101c,
+    0x08000000, 0x00080000, 0x08000000, 0x04080000,
+    0x08100000, 0x00100804, 0x1c000000, 0x00001c00,
+    0x08040000, 0x00040810, 0x101c0000, 0x00080008,
+    0x141c0000, 0x001c0414, 0x141c0000, 0x00141c14,
+    0x141c0000, 0x001c140c, 0x041c0000, 0x001c0404,
+    0x140c0000, 0x000c1414, 0x041c0000, 0x001c040c,
+    0x041c0000, 0x0004040c, 0x041c0000, 0x001c1414,
+    0x14140000, 0x0014141c, 0x081c0000, 0x001c0808,
+    0x10100000, 0x001c1410, 0x14140000, 0x0014140c,
+    0x04040000, 0x001c0404, 0x1c140000, 0x00141414,
+    0x140c0000, 0x00141414, 0x140c0000, 0x001c1414,
+    0x141c0000, 0x0004041c, 0x141c0000, 0x00101c14,
+    0x141c0000, 0x0014140c, 0x04180000, 0x000c101c,
+    0x081c0000, 0x00080808, 0x14140000, 0x001c1414,
+    0x14140000, 0x00181414, 0x14140000, 0x00141c14,
+    0x14140000, 0x00141408, 0x14140000, 0x0008081c,
+    0x101c0000, 0x001c0408, 0x041c0000, 0x001c0404,
+    0x0c040000, 0x00101808, 0x101c0000, 0x001c1010,
+    0x14080000, 0x00000000, 0x00000000, 0x001c0000,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+    0x3232323e, 0x003e3232, 0x18181818, 0x00181818,
+    0x3e30303e, 0x003e0606, 0x3c30303e, 0x003e3030,
+    0x32323232, 0x0030303e, 0x3e06063e, 0x003e3030,
+    0x3e06063e, 0x003e2626, 0x1830303e, 0x0006060c,
+    0x3e32323e, 0x003e3232, 0x3e32323e, 0x003e3030,
+    0x00000000, 0x00060600, 0x18362600, 0x0032360c,
+    0x32323e1c, 0x0032323e, 0x1e26261e, 0x001e2626,
+    0x0606063c, 0x003c0606, 0x3232321e, 0x001e3232,
+    0x1e06063e, 0x003e0606, 0x1e06063e, 0x00060606,
+    0x3606063c, 0x001c2626, 0x3e323232, 0x00323232,
+    0x0c0c0c1e, 0x001e0c0c, 0x30303030, 0x003e3232,
+    0x0e162626, 0x00262616, 0x06060606, 0x003e0606,
+    0x6a7e7662, 0x00626262, 0x323a3e36, 0x00323232,
+    0x3232321c, 0x001c3232, 0x3e32323e, 0x00020202,
+    0x3232321c, 0x002c1a3a, 0x1e26261e, 0x00262626,
+    0x1e06063c, 0x001e3030, 0x1818187e, 0x00181818,
+    0x32323232, 0x003e3232, 0x32323232, 0x00183432,
+    0x32323232, 0x00143a32, 0x0c323232, 0x00323232,
+    0x32323232, 0x001e203c, 0x1c30303e, 0x003e0606,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+    0x00000000, 0x00000000, 0x00000000, 0x00000000,
+};
diff --git a/src/text/text.h b/src/text/text.h
index 0bcf090..ab525d4 100644
--- a/src/text/text.h
+++ b/src/text/text.h
@@ -2,10 +2,8 @@
 #define TEXT_H
 
 #include "posprintf.h"
-#include "renderer.h"
-
-#include "font.h"
 
+typedef void (*TxtDrawc)(char c, size_t x, size_t y, u8 clr);
 typedef struct TextEngine {
     // Cursor for tiled text mode The X and Y positions correspond to the tile
     // X and Y starting from the top left of the screen. For a 240x160 screen,
@@ -13,14 +11,10 @@ typedef struct TextEngine {
     size_t cursor_x;
     size_t cursor_y;
 
-    // Memory location of font tile data and tilemap. Likely located on the
-    // VRAM.
-    u32 *font_data;
-    u16 *font_tilemap;
-
-    // The font map for tiled text. Writing the character stored in this
-    // position on the tilemap will show a character on the screen.
-    u16 font_map[256];
+    u8 buffer[30 * 20];
+    u8 spacing;
+    u8 color;
+    TxtDrawc drawc;
 } TextEngine;
 
 static TextEngine text_engine = {0};
@@ -28,20 +22,11 @@ static TextEngine text_engine = {0};
 // Initializes the text engine.
 static inline
 void
-txt_init(u32 *font_data, u16 *font_tilemap, u16 font_offset) {
-    // Load font data into VRAM.
-    unpack_tiles(&font, font_data, 256);
-
-    // Initialize the font map translation table. That way we can write
-    // a character on the text background layer with:
-    //     FONT_TILEMAP[tile_x + 32 * tile_y] = font_map['A'];
-    for (size_t i = 0; i < 256; ++i) {
-        text_engine.font_map[i] = font_offset + i;
-    }
-
+txt_init(TxtDrawc drawc) {
     // Initialize remaining variables.
-    text_engine.font_data = font_data;
-    text_engine.font_tilemap = font_tilemap;
+    text_engine.spacing = 8;
+    text_engine.color = 1;
+    text_engine.drawc = drawc;
 }
 
 // Writes a message to the tile text background.
@@ -59,7 +44,7 @@ txt_puts(char *msg) {
         } else {
             int x = text_engine.cursor_x;
             int y = text_engine.cursor_y;
-            text_engine.font_tilemap[x + 32 * y] = text_engine.font_map[(u16)c];
+            text_engine.buffer[x + 30 * y] = c;
             text_engine.cursor_x += 1;
             if (text_engine.cursor_x >= 30) {
                 text_engine.cursor_x = 0;
@@ -76,18 +61,15 @@ txt_puts(char *msg) {
 static inline
 void
 txt_clear_line(void) {
-    for (size_t i = 0; i < 30; ++i) {
-        int x = text_engine.cursor_x;
-        int y = text_engine.cursor_y;
-        text_engine.font_tilemap[x + 32 * y] = text_engine.font_map[0];
-    }
+    dma_fill(text_engine.buffer, 0, sizeof(text_engine.buffer), 3);
     text_engine.cursor_x = 0;
+    text_engine.cursor_y = 0;
 }
 
 // Clears the screen on the tile text mode.
 static inline
 void
-txt_clear_screen(void) {
+txt_clear(void) {
     for (size_t j = 0; j < 20; ++j) {
         text_engine.cursor_y = j;
         txt_clear_line();
@@ -104,30 +86,49 @@ txt_position(size_t tile_x, size_t tile_y) {
     text_engine.cursor_y = tile_y;
 }
 
+static inline
+void
+txt_color(u8 clr) {
+    text_engine.color = clr;
+}
+
+static inline
+void
+txt_spacing(u8 spacing) {
+    text_engine.spacing = spacing;
+}
+
+// Renders the contents of the scrollback buffer to the screen.
+void
+txt_render(void) {
+    for (size_t y = 0; y < 20; y++) {
+        for (size_t x = 0; x < 30; x++) {
+            size_t pos = x + y * 30;
+            if (text_engine.buffer[pos] == 0) {
+                continue;
+            }
+            text_engine.drawc(
+                    text_engine.buffer[pos],
+                    x * text_engine.spacing,
+                    y * text_engine.spacing,
+                    text_engine.color);
+        }
+    }
+}
+
 // Draws a message where the first character's top-left corner begins at the
 // given x and y position. The character spacing can be adjusted, but beware of
 // color merging issues.
 static inline
 void
-txt_draws(char *msg, size_t x, size_t y, size_t spacing, u8 clr) {
+txt_draws(char *msg, size_t x, size_t y, u8 clr) {
+    size_t i = 0;
     while (*msg) {
         char c = *msg++;
-        Tile *tile = FONT_DATA;
-        tile += c;
-        draw_tile(x, y, tile, clr, true);
-        x += spacing;
+        text_engine.drawc(c, x + i++ * text_engine.spacing, y, clr);
     }
 }
 
-static inline
-void
-txt_drawc(char c, size_t x, size_t y, size_t spacing, u8 clr) {
-    Tile *tile = FONT_DATA;
-    tile += c;
-    draw_tile(x, y, tile, clr, true);
-    x += spacing;
-}
-
 // Print text to the screen with formatting.
 #define txt_printf(msg, ...) \
         { \
@@ -138,31 +139,11 @@ txt_drawc(char c, size_t x, size_t y, size_t spacing, u8 clr) {
 
 // Draws text to the screen with formatting starting on the x and y position and
 // with custom character spacing.
-#define txt_drawf(msg, x, y, s, c, ...) \
-        { \
-                char buf[256] = {0}; \
-                posprintf(buf, msg, ##__VA_ARGS__); \
-                txt_draws(buf, x, y, s, c); \
-        }
-
-// Small font is located after the initial ASCII characters, and only supports
-// lowercase characters.
-// NOTE: Slow, we should do this with a LUT.
-#define txt_drawf_small(msg, x, y, s, c, ...) \
+#define txt_drawf(msg, x, y, c, ...) \
         { \
                 char buf[256] = {0}; \
                 posprintf(buf, msg, ##__VA_ARGS__); \
-                for (size_t i = 0; i < 256; i++) { \
-                    if (buf[i] == 0) { \
-                        break; \
-                    } \
-                    if (buf[i] < 'a') { \
-                buf[i] += 16 * 6; \
-                    } else { \
-                buf[i] += 16 * 4; \
-                    }\
-                } \
-                txt_draws(buf, x, y, s, c); \
+                txt_draws(buf, x, y, c); \
         }
 
 #endif // TEXT_H