Update profiling macros and m0 renderer

19 files changed, 1704 insertions, 2062 deletions

diff --git a/Makefile b/Makefile
index 98fbfd1..8db7b02 100644
--- a/Makefile
+++ b/Makefile
@@ -17,8 +17,8 @@ BUILD_DIR   := build
 SRC_MAIN    := $(SRC_DIR)/main.c
 SRC_BDGBA   := $(wildcard $(SRC_DIR)/gba/*.s)
 SRC_BDGBA   += $(wildcard $(SRC_DIR)/gba/*.c)
-SRC_BDTEXT  := $(wildcard $(SRC_DIR)/text/*.s)
-SRC_BDTEXT  += $(wildcard $(SRC_DIR)/text/*.c)
+SRC_BDTEXT  := $(wildcard $(SRC_DIR)/gba/text/*.s)
+SRC_BDTEXT  += $(wildcard $(SRC_DIR)/gba/text/*.c)
 SRC         := $(SRC_MAIN) $(SRC_BDGBA) $(SRC_BDTEXT)
 WATCH_SRC   := $(shell find $(SRC_DIR) -name *.c -or -name *.s -or -name *.h)
 INC_DIRS    := $(shell find $(SRC_DIR) -type d)
diff --git a/src/gba/draw.s b/src/gba/draw.s
new file mode 100644
index 0000000..a5705ee
--- /dev/null
+++ b/src/gba/draw.s
@@ -0,0 +1,183 @@
+.file "draw.s"
+.section .iwram, "ax", %progbits
+.arm
+.align
+
+.global decode_1bpp_mask
+decode_1bpp_mask:
+    .word   0x00000000
+    .word   0xff000000
+    .word   0x00ff0000
+    .word   0xffff0000
+    .word   0x0000ff00
+    .word   0xff00ff00
+    .word   0x00ffff00
+    .word   0xffffff00
+    .word   0x000000ff
+    .word   0xff0000ff
+    .word   0x00ff00ff
+    .word   0xffff00ff
+    .word   0x0000ffff
+    .word   0xff00ffff
+    .word   0x00ffffff
+    .word   0xffffffff
+
+.global decode_1bpp_mask_flip_x
+decode_1bpp_mask_flip_x:
+    .word   0x00000000
+    .word   0x000000ff
+    .word   0x0000ff00
+    .word   0x0000ffff
+    .word   0x00ff0000
+    .word   0x00ff00ff
+    .word   0x00ffff00
+    .word   0x00ffffff
+    .word   0xff000000
+    .word   0xff0000ff
+    .word   0xff00ff00
+    .word   0xff00ffff
+    .word   0xffff0000
+    .word   0xffff00ff
+    .word   0xffffff00
+    .word   0xffffffff
+
+color_mask:
+    .word   0x01010101
+
+@ Arguments
+@
+@     r0, r1: x, y
+@     r2: *sprite
+@     r3: clr
+@     sp0: flip_x (r4)
+@     sp1: flip_y (r5)
+@     sp2: framebuffer (r10)
+@
+.global draw_icn_asm
+draw_icn_asm:
+    @ Save fp and lr so that we can use it in our function. The fp becomes
+    @ a frame of reference to be able to access the stack passed arguments.
+    @ Right now lr can be used as a temporary volatile value to do operations,
+    @ we will restore both of these before exiting.
+    push    {fp, lr}
+    mov     fp, sp
+    add     fp, #8
+
+    @ Make sure we don't clobber these registers.
+    push    {r4-r10}
+
+    @ Get the u8 color into the proper u32 format
+    ldr     r4, color_mask
+    mul     r3, r3, r4
+
+    @ Fetch arguments 4-6.
+    ldmfd   fp, {r4-r6}
+
+    @ Prepare backbuffer position
+    @ fp: framebuffer position for the initial tile
+    mov     lr, #240
+    mul     lr, r1, lr
+    add     r6, lr
+    mov     lr, r0, lsr #3
+    add     r6, lr, lsl #3
+    mov     fp, r6
+
+    @ r10: decoding table
+    adr     r10, decode_1bpp_mask
+
+    @ TODO: Ignoring flip_x/flip_y for now
+    @ TODO: Once we calculate the number of rows, r0 and r1 are free to use,
+    @       assuming n_rows is 8 for now for simplicity. Using r0/r1 for shift_left/right
+
+    @ start_col
+    and     lr, r0, #7
+    cmp     lr, #3
+    bls     .draw_icn_asm_shift_setup_b
+.draw_icn_asm_shift_setup_a:
+    @ dst++;
+    @ shift_left = 8 * (start_col - 4);
+    @ shift_right = 8 * (8 - start_col);
+    add     fp, #4
+    sub     r0, lr, #4
+    lsl     r0, #3
+    mov     r4, #8
+    sub     r1, r4, lr
+    lsl     r1, #3
+    b       .draw_icn_asm_loop_start
+
+    @ r0: shift_left / r1: shift_right
+.draw_icn_asm_shift_setup_b:
+    @ size_t shift_left = start_col * 8;
+    @ size_t shift_right = (4 - start_col) * 8;
+    mov     r0, lr, lsl #3
+    mov     r1, #4
+    sub     r1, r1, lr
+    lsl     r1, #3
+
+.draw_icn_asm_loop_start:
+    @ lr: loop counter
+    mov     lr, #8
+.Ldraw_icn_asm_loop:
+    @ Decode current sprite mask
+    @ u32 msk0 = decode_1bpp_mask[*sprite >> 4];
+    @ u32 msk1 = decode_1bpp_mask[*sprite & 0xf];
+    @ r4: msk0 (left)
+    @ r6: msk1 (right)
+    ldrb    r4, [r2]
+    mov     r6, r4
+    lsr     r4, #4
+    and     r6, #0xf
+    ldr     r4, [r10, r4, lsl #2] @ mks0
+    ldr     r6, [r10, r6, lsl #2] @ mks1
+
+    @ u32 msk2 = (msk0 >> shift_right) | (msk1 << shift_left);
+    @ r5: msk2 (middle)
+    mov     r5, r4, lsr r1
+    orr     r5, r6, lsl r0
+
+    @ msk0 <<= shift_left;
+    @ msk1 >>= shift_right;
+    lsl     r4, r0
+    lsr     r6, r1
+
+    @ dst[0] = (dst[0] & ~msk0) | clr0;
+    @ dst[1] = (dst[1] & ~msk2) | clr2;
+    @ dst[2] = (dst[2] & ~msk1) | clr1;
+    @ Load existing row data.
+    ldm     fp, {r7-r9}
+    @ dst[0]
+    bic     r7, r4
+    and     r4, r3, r4
+    orr     r7, r4
+    @ dst[1]
+    bic     r8, r5
+    and     r5, r3, r5
+    orr     r8, r5
+    @ dst[2]
+    bic     r9, r6
+    and     r6, r3, r6
+    orr     r9, r6
+
+    @ Store all three tile rows at once
+    stm   fp, {r7-r9}
+
+    @ Next loop iteration
+    add     fp, fp, #240
+    subs    lr, lr, #1
+    add     r2, #1
+    bhi     .Ldraw_icn_asm_loop
+
+    @ DEBUG: Write color in framebuffer
+    @ str     r3, [fp]
+    @ add     fp, #4
+    @ str     r3, [fp]
+
+    @ DEBUG: return value.
+    @ mov     r0, r4
+    @ mov     r0, r6
+    @ mov     r0, r5
+
+.draw_icn_asm_end:
+    pop     {r4-r10}
+    pop     {fp, lr}
+    bx      lr
diff --git a/src/gba/gba.h b/src/gba/gba.h
index 27a6a9a..045f2f4 100644
--- a/src/gba/gba.h
+++ b/src/gba/gba.h
@@ -289,12 +289,24 @@ profile_measure(void) {
 static u16 key_curr = 0;
 static u16 key_prev = 0;
 
-static inline
-void
-poll_keys(void) {
-    key_prev = key_curr;
-    key_curr = ~KEY_INPUTS & KEY_MASK;
-}
+// Stores number of frames since a keay was pressed.
+typedef struct Controller {
+    int key_up;
+    int key_down;
+    int key_left;
+    int key_right;
+    int key_select;
+    int key_start;
+    int key_b;
+    int key_a;
+    int key_l;
+    int key_r;
+} Controller;
+
+static Controller ctrl = {0};
+
+#define RETRIG_OFFSET 16
+#define RETRIG_FRAMES 2
 
 // Returns true if the given key has been pressed at time of calling and was not
 // pressed since the previous call. For example, if a key is being held, this
@@ -327,40 +339,110 @@ key_hold(u32 key) {
     return key_curr & key_prev & key;
 }
 
-// Back/unpack bits.
 static inline
-u32
-unpack_1bb(u8 hex) {
-    const u32 conversion_u32[16] = {
-        0x00000000, 0x00000001, 0x00000010, 0x00000011,
-        0x00000100, 0x00000101, 0x00000110, 0x00000111,
-        0x00001000, 0x00001001, 0x00001010, 0x00001011,
-        0x00001100, 0x00001101, 0x00001110, 0x00001111,
-    };
-    u8 low = hex & 0xF;
-    u8 high = (hex >> 4) & 0xF;
-    return (conversion_u32[high] << 16) | conversion_u32[low];
+bool
+_key_retrig(int key, int offset, int frames) {
+    if (key_tap(key)) {
+        return true;
+    }
+    switch (key) {
+        case KEY_L: {
+            if (key_hold(key)) {
+                if (ctrl.key_l < offset) { return false; }
+                if (ctrl.key_l % frames == 0) { return true; }
+            }
+        } break;
+        case KEY_R: {
+            if (key_hold(key)) {
+                if (ctrl.key_r < offset) { return false; }
+                if (ctrl.key_r % frames == 0) { return true; }
+            }
+        } break;
+        case KEY_A: {
+            if (key_hold(key)) {
+                if (ctrl.key_a < offset) { return false; }
+                if (ctrl.key_a % frames == 0) { return true; }
+            }
+        } break;
+        case KEY_B: {
+            if (key_hold(key)) {
+                if (ctrl.key_b < offset) { return false; }
+                if (ctrl.key_b % frames == 0) { return true; }
+            }
+        } break;
+        case KEY_SELECT: {
+            if (key_hold(key)) {
+                if (ctrl.key_select < offset) { return false; }
+                if (ctrl.key_select % frames == 0) { return true; }
+            }
+        } break;
+        case KEY_START: {
+            if (key_hold(key)) {
+                if (ctrl.key_start < offset) { return false; }
+                if (ctrl.key_start % frames == 0) { return true; }
+            }
+        } break;
+        case KEY_UP: {
+            if (key_hold(key)) {
+                if (ctrl.key_up < offset) { return false; }
+                if (ctrl.key_up % frames == 0) { return true; }
+            }
+        } break;
+        case KEY_DOWN: {
+            if (key_hold(key)) {
+                if (ctrl.key_down < offset) { return false; }
+                if (ctrl.key_down % frames == 0) { return true; }
+            }
+        } break;
+        case KEY_LEFT: {
+            if (key_hold(key)) {
+                if (ctrl.key_left < offset) { return false; }
+                if (ctrl.key_left % frames == 0) { return true; }
+            }
+        } break;
+        case KEY_RIGHT: {
+            if (key_hold(key)) {
+                if (ctrl.key_right < offset) { return false; }
+                if (ctrl.key_right % frames == 0) { return true; }
+            }
+        } break;
+    }
+    return false;
+}
+
+static inline
+bool
+key_retrig(int key) {
+    return _key_retrig(key, RETRIG_OFFSET, RETRIG_FRAMES);
 }
 
-// Unpack N tiles packed at 1bpp.
 static inline
 void
-unpack_tiles(const u32 *src, u32 *dst, size_t n_tiles) {
-    const u32 *target_src = src + n_tiles * 2;
-    while (src != target_src) {
-        *dst++ = unpack_1bb((*src >> 24) & 0xFF);
-        *dst++ = unpack_1bb((*src >> 16) & 0xFF);
-        *dst++ = unpack_1bb((*src >> 8) & 0xFF);
-        *dst++ = unpack_1bb(*src & 0xFF);
-        src++;
-    }
+update_controller(void) {
+    if (key_pressed(KEY_UP)) { ctrl.key_up++; } else if (key_released(KEY_UP)) { ctrl.key_up = 0; }
+    if (key_pressed(KEY_DOWN)) { ctrl.key_down++; } else if (key_released(KEY_DOWN)) { ctrl.key_down = 0; }
+    if (key_pressed(KEY_LEFT)) { ctrl.key_left++; } else if (key_released(KEY_LEFT)) { ctrl.key_left = 0; }
+    if (key_pressed(KEY_RIGHT)) { ctrl.key_right++; } else if (key_released(KEY_RIGHT)) { ctrl.key_right = 0; }
+    if (key_pressed(KEY_L)) { ctrl.key_l++; } else if (key_released(KEY_L)) { ctrl.key_l = 0; }
+    if (key_pressed(KEY_R)) { ctrl.key_r++; } else if (key_released(KEY_R)) { ctrl.key_r = 0; }
+    if (key_pressed(KEY_A)) { ctrl.key_a++; } else if (key_released(KEY_A)) { ctrl.key_a = 0; }
+    if (key_pressed(KEY_B)) { ctrl.key_b++; } else if (key_released(KEY_B)) { ctrl.key_b = 0; }
+    if (key_pressed(KEY_SELECT)) { ctrl.key_select++; } else if (key_released(KEY_SELECT)) { ctrl.key_select = 0; }
+    if (key_pressed(KEY_START)) { ctrl.key_start++; } else if (key_released(KEY_START)) { ctrl.key_start = 0; }
+}
+
+static inline
+void
+poll_keys(void) {
+    key_prev = key_curr;
+    key_curr = ~KEY_INPUTS & KEY_MASK;
+    update_controller();
 }
 
 //
 // Direct Memory Access (DMA)
 //
 
-
 // Source, destination, and control registers.
 #define DMA_SRC(N)  *((vu32*) 0x040000B0 + (N) * 12)
 #define DMA_DST(N)  *((vu32*) 0x040000B4 + (N) * 12)
@@ -482,6 +564,24 @@ int bios_vblank_wait();
 int bios_div(int num, int denom);
 
 //
+// SIO Link Cable
+//
+
+#define SIO_MODE *((vu16*)(MEM_IO + 0x0134))
+#define SIO_CNT  *((vu16*)(MEM_IO + 0x0128))
+
+#define SIO_MODE_GP     (2 << 14)
+#define SIO_SC(X)       ((X) << 0)
+#define SIO_SD(X)       ((X) << 1)
+#define SIO_SI(X)       ((X) << 2)
+#define SIO_SO(X)       ((X) << 3)
+#define SIO_SC_OUT(X)   ((X) << 4)
+#define SIO_SD_OUT(X)   ((X) << 5)
+#define SIO_SI_OUT(X)   ((X) << 6)
+#define SIO_SO_OUT(X)   ((X) << 7)
+#define SIO_IRQ_ENABLE  (1 << 8)
+
+//
 // Sound.
 //
 
@@ -522,10 +622,35 @@ typedef enum {
 } SoundChannel;
 
 inline u16
-sound_volume(SoundChannel channels, u8 volume) {
+sound_volume(SoundChannel channels, u8 volume, u8 pan) {
     volume = volume & 0x7;
     channels = channels & 0xF;
-    return volume | (volume << 0x4) | (channels << 0x8) | (channels << 0xC);
+    u16 right = volume | (channels << 0x8);
+    u16 left = (volume << 0x4) | (channels << 0xC);
+    if (pan == 1) {
+        return left;
+    }
+    if (pan == 2) {
+        return right;
+    }
+    return left | right;
+}
+
+inline u16 dmg_stereo_vol(u8 vol) {
+    return vol | (vol << 0x4);
+}
+
+inline u16
+channel_vol(SoundChannel channel, s8 pan) {
+    u16 left = (channel << 0xc);
+    u16 right = (channel << 0x8);
+    if (pan == -1) {
+        return left;
+    }
+    if (pan == +1) {
+        return right;
+    }
+    return left | right;
 }
 
 // Sound Direct Sound master bits.
@@ -534,12 +659,12 @@ sound_volume(SoundChannel channels, u8 volume) {
 #define SOUND_DMG100          0x2
 #define SOUND_DSOUND_RATIO_A  (1 << 0x2)
 #define SOUND_DSOUND_RATIO_B  (1 << 0x3)
-#define SOUND_DSOUND_LEFT_A   (1 << 0x8)
-#define SOUND_DSOUND_RIGHT_A  (1 << 0x9)
+#define SOUND_DSOUND_RIGHT_A  (1 << 0x8)
+#define SOUND_DSOUND_LEFT_A   (1 << 0x9)
 #define SOUND_DSOUND_TIMER_A  (1 << 0xA)
 #define SOUND_DSOUND_RESET_A  (1 << 0xB)
-#define SOUND_DSOUND_LEFT_B   (1 << 0xC)
-#define SOUND_DSOUND_RIGHT_B  (1 << 0xD)
+#define SOUND_DSOUND_RIGHT_B  (1 << 0xC)
+#define SOUND_DSOUND_LEFT_B   (1 << 0xD)
 #define SOUND_DSOUND_TIMER_B  (1 << 0xE)
 #define SOUND_DSOUND_RESET_B  (1 << 0xF)
 
@@ -661,6 +786,7 @@ wait_vsync(void) {
 // General utility macros.
 #define MIN(A, B) ((A) <= (B) ? (A) : (B))
 #define MAX(A, B) ((A) >= (B) ? (A) : (B))
+#define ABS(A)    (((A) ^ ((A) >> (sizeof(A) * 8 - 1))) - ((A) >> (sizeof(A) * 8 - 1)))
 #define CLAMP(X, MIN, MAX) ((X) <= (MIN) ? (MIN) : (X) > (MAX) ? (MAX): (X))
 #define LEN(ARR) (sizeof(ARR) / sizeof((ARR)[0]))
 
@@ -688,10 +814,23 @@ memcpy32(u32 *dst, const u32 *src, u32 size) {
     }
 }
 
+static inline
+void
+memset32(u32 *dst, const u32 data, u32 size) {
+    for (size_t i = 0; i < size / 4; i++) {
+        dst[i] = data;
+    }
+}
+
+// Optimized ARMASM versions of memcpy32 and memset32.
+extern void copy32(u32 *dst, u32 *src, u32 chunks);
+extern void set32(u32 *dst, u32 data, u32 chunks);
+
 //
 // Compiler hints.
 //
 
 #define UNROLL_LOOPS __attribute__((optimize("unroll-loops")))
+#define INLINE __attribute__((always_inline)) inline
 
 #endif // GBA_H
diff --git a/src/gba/profiling.h b/src/gba/profiling.h
new file mode 100755
index 0000000..5fc8d83
--- /dev/null
+++ b/src/gba/profiling.h
@@ -0,0 +1,134 @@
+//
+// Profiling macros.
+//
+
+// NOTE: Profiling uses the last two timers to count cycles, and thus can't be
+//       used for measuring applications that
+
+
+static u32 frame_time = 0;
+
+#ifndef PROF_ENABLE
+#define PROF_ENABLE 0
+#endif
+
+#if PROF_ENABLE > 0
+
+#ifndef PROF_RESET_MINMAX
+#define PROF_RESET_MINMAX false
+#endif
+
+// Maximum number of profiling to monitor.
+typedef enum ProfType {
+    PROF_INPUT,
+    PROF_UPDATE,
+    PROF_RENDER,
+    PROF_FLIP,
+    PROF_FILL,
+    PROF_NUM,
+} ProfType;
+
+char *prof_type_str[PROF_NUM] = {
+    "INPUT    ",
+    "UPDATE   ",
+    "RENDER   ",
+    "FLIPBUF  ",
+    "SCRFILL  ",
+};
+
+u32 prof_frame_time  = 0;
+u32 prof_frame_count = 0;
+u32 prof_frame_avg  = 0;
+u32 prof_frame_time_max  = 0;
+u32 prof_times[PROF_NUM] = {0};
+u32 prof_count[PROF_NUM] = {0};
+u32 prof_avg[PROF_NUM] = {0};
+u32 prof_max[PROF_NUM] = {0};
+u32 prof_min[PROF_NUM] = {0};
+
+bool prof_reset_minmax = PROF_RESET_MINMAX;
+bool prof_show = true;
+
+#define PROF_INIT() do { \
+    for (size_t i = 0; i < PROF_NUM; i++) { \
+        prof_min[i] = -1; \
+    } \
+} while(0);
+
+#define PROF(func, idx) do { \
+    u32 time_before = profile_measure(); \
+    (func); \
+    u32 time_after = profile_measure(); \
+    u32 time_current = time_after - time_before; \
+    prof_times[idx] += time_current; \
+    prof_count[idx]++; \
+    prof_max[idx] = MAX(time_current, prof_max[idx]);\
+    prof_min[idx] = MIN(time_current, prof_min[idx]);\
+} while(0);
+
+#define FRAME_START() do { \
+    profile_start();\
+} while(0)
+
+
+#define FRAME_END() do { \
+        prof_frame_count++;\
+        frame_time = profile_measure();\
+        prof_frame_time_max = MAX(prof_frame_time_max, frame_time);\
+        prof_frame_time += profile_stop();\
+        if (prof_show) { \
+            draw_filled_rect(0, 0, SCREEN_WIDTH - 1, 8 * (PROF_NUM + 1), 2); \
+            u32 fps = (u64)280896 * 60 / (prof_frame_avg + 1); \
+            if (prof_frame_avg == 0) { \
+                fps = 0; \
+            } \
+            txt_drawf_small("FRAME TIME/FPS: %.9l/%.2l", 0, 0, COL_FG, \
+                    prof_frame_avg, fps);\
+            txt_drawf_small("MAX: %.9l/%l", 8 * 19, 0, COL_FG, \
+                    prof_frame_time_max, 280896);\
+            for (size_t idx = 0; idx < PROF_NUM; idx++) { \
+                txt_drawf_small("%s %.9l (%.9l %.9l) %08x:%08x", 0, 8 * (idx + 1), COL_FG, \
+                        prof_type_str[idx], \
+                        prof_avg[idx], \
+                        prof_min[idx], \
+                        prof_max[idx], \
+                        prof_avg[idx], \
+                        prof_max[idx]);\
+            }; \
+            draw_filled_rect(0, SCREEN_HEIGHT - 9, 58, SCREEN_HEIGHT - 1, 2); \
+            txt_drawf_small("CPU USAGE: %.3l", 0, SCREEN_HEIGHT - 9, COL_FG, \
+                    (u64)prof_frame_avg * 100 / 280896);\
+        } \
+        if (prof_frame_count >= PROF_ENABLE) { \
+            for (size_t idx = 0; idx < PROF_NUM; idx++) { \
+                prof_avg[idx] = prof_times[idx] / prof_frame_count; \
+                if (prof_reset_minmax) { \
+                    prof_min[idx] = -1; \
+                    prof_max[idx] = 0; \
+                } \
+                prof_times[idx] = 0; \
+                prof_count[idx] = 0; \
+            }; \
+            prof_frame_avg = prof_frame_time / prof_frame_count; \
+            prof_frame_count = 0; \
+            prof_frame_time = 0; \
+        } \
+    } while(0)
+
+#define PROF_SHOW() do { \
+    prof_show ^= 1; \
+} while(0)
+
+#else
+
+// No profiling.
+#define PROF_INIT()
+#define PROF(F,VAR) do {F;} while(0)
+#define FRAME_START() do { \
+    profile_start();\
+} while(0)
+#define FRAME_END() do { \
+    frame_time = profile_stop();\
+} while(0)
+#define PROF_SHOW()
+#endif
diff --git a/src/gba/renderer/renderer.c b/src/gba/renderer/renderer.c
new file mode 100644
index 0000000..fec5b6f
--- /dev/null
+++ b/src/gba/renderer/renderer.c
@@ -0,0 +1 @@
+#include "renderer_m0.c"
diff --git a/src/renderer.h b/src/gba/renderer/renderer.h
index e6637ef..8b37496 100644
--- a/src/renderer.h
+++ b/src/gba/renderer/renderer.h
@@ -19,14 +19,8 @@ void draw_filled_rect(size_t x0, size_t y0, size_t x1, size_t y1, u8 clr);
 // 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);
-
-// 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);
+// Draws a 1bpp sprite in the given color.
+void draw_1bpp(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y);
 
 // Copies data and performs page flipping if needed.
 // To be called exactly once at the beginning of the VBlank.
diff --git a/src/renderer_m0.c b/src/gba/renderer/renderer_m0.c
index 0145f64..6eb45c5 100644
--- a/src/renderer_m0.c
+++ b/src/gba/renderer/renderer_m0.c
@@ -13,8 +13,9 @@
 //
 
 #define SUBPIXEL_LINES 1
-#define DEC_BIG_LUT 1
-#define FLIP_TYPE 3
+#define FLIP_TYPE 1
+#define DISABLE_BOUNDCHECK_SCREEN 0
+#define NO_DMA 1
 
 // Front/back buffers for double buffering.
 #define BUF_0 ((u32*)(MEM_VRAM))
@@ -38,7 +39,7 @@ static u32 dirty_tiles[21] = {0};
 // 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
+#if DISABLE_BOUNDCHECK_SCREEN > 0
 #define BOUNDCHECK_SCREEN(X,Y)
 #else
 #define BOUNDCHECK_SCREEN(X,Y) if ((X) >= SCREEN_WIDTH || (Y) >= SCREEN_HEIGHT) return;
@@ -50,6 +51,24 @@ static u32 dirty_tiles[21] = {0};
 // Swap A and B values to make sure A <= B.
 #define MAYBE_SWAP(A,B) if ((A) > (B)) { SWAP(A,B); }
 
+// Color macros.
+#define COL_BG 0
+#define COL_FG 1
+#define COL_00 2
+#define COL_01 3
+#define COL_02 4
+#define COL_03 5
+#define COL_04 6
+#define COL_05 7
+#define COL_06 8
+#define COL_07 9
+#define COL_08 10
+#define COL_09 11
+#define COL_10 12
+#define COL_11 13
+#define COL_12 14
+#define COL_13 15
+
 //
 // Basic primitives.
 //
@@ -66,7 +85,11 @@ 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.
+#if NO_DMA == 0
     dma_fill(backbuf, 0x11111111 * clr, KB(20) - 32, 3);
+#else
+    set32(backbuf, 0x11111111 * clr, (KB(20) / 4) - 8);
+#endif
     redraw();
 }
 
@@ -85,12 +108,13 @@ draw_pixel(size_t x, size_t y, u8 clr) {
     // Update backbuffer.
     size_t shift = start_col * sizeof(u32);
     u32 mask = 0xF << shift;
-    u32 row = clr << shift;
-    *dst = (*dst & ~mask) | row;
+    u32 color = clr << shift;
+    *dst = (*dst & ~mask) | color;
     dirty_tiles[tile_y] |= 1 << tile_x;
 }
 
 IWRAM_CODE
+UNROLL_LOOPS
 static inline
 void
 draw_hline(size_t x0, size_t x1, size_t y0, u8 clr) {
@@ -117,21 +141,21 @@ draw_hline(size_t x0, size_t x1, size_t y0, u8 clr) {
         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;
+        u32 color = (0x11111111 * clr) & mask;
+        *dst = (*dst & ~mask) | color;
     } 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);
+        u32 color = 0x11111111 * clr;
+        *dst = (*dst & ~(mask << shift_left)) | (color << shift_left);
         dst += 8;
         for (size_t i = 1; i < dtx; i++) {
             dirty |= (1 << (tile_x0 + i));
-            *dst = row;
+            *dst = color;
             dst += 8;
         }
-        *dst = (*dst & ~(mask >> shift_right)) | (row >> shift_right);
+        *dst = (*dst & ~(mask >> shift_right)) | (color >> shift_right);
     }
     dirty_tiles[tile_y] |= dirty;
 }
@@ -156,26 +180,26 @@ draw_vline(size_t x0, size_t y0, size_t y1, u8 clr) {
 
     u32 *dst = &backbuf[start_row0 + (tile_x + tile_y * 32) * 8];
     u32 mask = 0x0000000F << shift_left;
-    u32 row = (0x11111111 * clr) & mask;
+    u32 color = clr << shift_left;
     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;
+            dst[0] = (dst[0] & ~mask) | color;
         }
     } else {
         for (size_t i = 0; i < (8 - start_row0); i++, dst++) {
-            dst[0] = (dst[0] & ~mask) | row;
+            dst[0] = (dst[0] & ~mask) | color;
         }
         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[0] = (dst[0] & ~mask) | color;
             }
             dst += 8 * 31;
         }
         for (size_t i = 0; i <= start_row1; i++, dst++) {
-            dst[0] = (dst[0] & ~mask) | row;
+            dst[0] = (dst[0] & ~mask) | color;
         }
     }
     dirty_tiles[tile_y0] |= dirty;
@@ -313,8 +337,8 @@ draw_filled_rect(size_t x0, size_t y0, size_t x1, size_t y1, u8 clr) {
     MAYBE_SWAP(x0, x1);
     MAYBE_SWAP(y0, y1);
 
-    // Special condition. If the screen is to be completely filled, use the DMA
-    // instead.
+    // Special condition. If the screen is to be completely filled, use the
+    // full clearing functions instead.
     if (x0 == 0 && x1 >= (SCREEN_WIDTH - 1) && y0 == 0 && y1 >= (SCREEN_HEIGHT - 1)) {
         screen_fill(clr);
         return;
@@ -326,161 +350,137 @@ draw_filled_rect(size_t x0, size_t y0, size_t x1, size_t y1, u8 clr) {
 }
 
 //
-// Sprites (chr/icn).
+// Sprites (1bpp).
 //
 
-#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
+// NOTE: If IWRAM is too full, these could be made `const` at the cost of some
+//       performance when decoding 1bpp sprites.
+static u32 lut_1bpp_mask[256] = {
+    0x00000000, 0xf0000000, 0x0f000000, 0xff000000, 0x00f00000,
+    0xf0f00000, 0x0ff00000, 0xfff00000, 0x000f0000, 0xf00f0000,
+    0x0f0f0000, 0xff0f0000, 0x00ff0000, 0xf0ff0000, 0x0fff0000,
+    0xffff0000, 0x0000f000, 0xf000f000, 0x0f00f000, 0xff00f000,
+    0x00f0f000, 0xf0f0f000, 0x0ff0f000, 0xfff0f000, 0x000ff000,
+    0xf00ff000, 0x0f0ff000, 0xff0ff000, 0x00fff000, 0xf0fff000,
+    0x0ffff000, 0xfffff000, 0x00000f00, 0xf0000f00, 0x0f000f00,
+    0xff000f00, 0x00f00f00, 0xf0f00f00, 0x0ff00f00, 0xfff00f00,
+    0x000f0f00, 0xf00f0f00, 0x0f0f0f00, 0xff0f0f00, 0x00ff0f00,
+    0xf0ff0f00, 0x0fff0f00, 0xffff0f00, 0x0000ff00, 0xf000ff00,
+    0x0f00ff00, 0xff00ff00, 0x00f0ff00, 0xf0f0ff00, 0x0ff0ff00,
+    0xfff0ff00, 0x000fff00, 0xf00fff00, 0x0f0fff00, 0xff0fff00,
+    0x00ffff00, 0xf0ffff00, 0x0fffff00, 0xffffff00, 0x000000f0,
+    0xf00000f0, 0x0f0000f0, 0xff0000f0, 0x00f000f0, 0xf0f000f0,
+    0x0ff000f0, 0xfff000f0, 0x000f00f0, 0xf00f00f0, 0x0f0f00f0,
+    0xff0f00f0, 0x00ff00f0, 0xf0ff00f0, 0x0fff00f0, 0xffff00f0,
+    0x0000f0f0, 0xf000f0f0, 0x0f00f0f0, 0xff00f0f0, 0x00f0f0f0,
+    0xf0f0f0f0, 0x0ff0f0f0, 0xfff0f0f0, 0x000ff0f0, 0xf00ff0f0,
+    0x0f0ff0f0, 0xff0ff0f0, 0x00fff0f0, 0xf0fff0f0, 0x0ffff0f0,
+    0xfffff0f0, 0x00000ff0, 0xf0000ff0, 0x0f000ff0, 0xff000ff0,
+    0x00f00ff0, 0xf0f00ff0, 0x0ff00ff0, 0xfff00ff0, 0x000f0ff0,
+    0xf00f0ff0, 0x0f0f0ff0, 0xff0f0ff0, 0x00ff0ff0, 0xf0ff0ff0,
+    0x0fff0ff0, 0xffff0ff0, 0x0000fff0, 0xf000fff0, 0x0f00fff0,
+    0xff00fff0, 0x00f0fff0, 0xf0f0fff0, 0x0ff0fff0, 0xfff0fff0,
+    0x000ffff0, 0xf00ffff0, 0x0f0ffff0, 0xff0ffff0, 0x00fffff0,
+    0xf0fffff0, 0x0ffffff0, 0xfffffff0, 0x0000000f, 0xf000000f,
+    0x0f00000f, 0xff00000f, 0x00f0000f, 0xf0f0000f, 0x0ff0000f,
+    0xfff0000f, 0x000f000f, 0xf00f000f, 0x0f0f000f, 0xff0f000f,
+    0x00ff000f, 0xf0ff000f, 0x0fff000f, 0xffff000f, 0x0000f00f,
+    0xf000f00f, 0x0f00f00f, 0xff00f00f, 0x00f0f00f, 0xf0f0f00f,
+    0x0ff0f00f, 0xfff0f00f, 0x000ff00f, 0xf00ff00f, 0x0f0ff00f,
+    0xff0ff00f, 0x00fff00f, 0xf0fff00f, 0x0ffff00f, 0xfffff00f,
+    0x00000f0f, 0xf0000f0f, 0x0f000f0f, 0xff000f0f, 0x00f00f0f,
+    0xf0f00f0f, 0x0ff00f0f, 0xfff00f0f, 0x000f0f0f, 0xf00f0f0f,
+    0x0f0f0f0f, 0xff0f0f0f, 0x00ff0f0f, 0xf0ff0f0f, 0x0fff0f0f,
+    0xffff0f0f, 0x0000ff0f, 0xf000ff0f, 0x0f00ff0f, 0xff00ff0f,
+    0x00f0ff0f, 0xf0f0ff0f, 0x0ff0ff0f, 0xfff0ff0f, 0x000fff0f,
+    0xf00fff0f, 0x0f0fff0f, 0xff0fff0f, 0x00ffff0f, 0xf0ffff0f,
+    0x0fffff0f, 0xffffff0f, 0x000000ff, 0xf00000ff, 0x0f0000ff,
+    0xff0000ff, 0x00f000ff, 0xf0f000ff, 0x0ff000ff, 0xfff000ff,
+    0x000f00ff, 0xf00f00ff, 0x0f0f00ff, 0xff0f00ff, 0x00ff00ff,
+    0xf0ff00ff, 0x0fff00ff, 0xffff00ff, 0x0000f0ff, 0xf000f0ff,
+    0x0f00f0ff, 0xff00f0ff, 0x00f0f0ff, 0xf0f0f0ff, 0x0ff0f0ff,
+    0xfff0f0ff, 0x000ff0ff, 0xf00ff0ff, 0x0f0ff0ff, 0xff0ff0ff,
+    0x00fff0ff, 0xf0fff0ff, 0x0ffff0ff, 0xfffff0ff, 0x00000fff,
+    0xf0000fff, 0x0f000fff, 0xff000fff, 0x00f00fff, 0xf0f00fff,
+    0x0ff00fff, 0xfff00fff, 0x000f0fff, 0xf00f0fff, 0x0f0f0fff,
+    0xff0f0fff, 0x00ff0fff, 0xf0ff0fff, 0x0fff0fff, 0xffff0fff,
+    0x0000ffff, 0xf000ffff, 0x0f00ffff, 0xff00ffff, 0x00f0ffff,
+    0xf0f0ffff, 0x0ff0ffff, 0xfff0ffff, 0x000fffff, 0xf00fffff,
+    0x0f0fffff, 0xff0fffff, 0x00ffffff, 0xf0ffffff, 0x0fffffff,
+    0xffffffff
 };
 
-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
+static u32 lut_1bpp_mask_flip_x[256] = {
+    0x00000000, 0x0000000f, 0x000000f0, 0x000000ff, 0x00000f00,
+    0x00000f0f, 0x00000ff0, 0x00000fff, 0x0000f000, 0x0000f00f,
+    0x0000f0f0, 0x0000f0ff, 0x0000ff00, 0x0000ff0f, 0x0000fff0,
+    0x0000ffff, 0x000f0000, 0x000f000f, 0x000f00f0, 0x000f00ff,
+    0x000f0f00, 0x000f0f0f, 0x000f0ff0, 0x000f0fff, 0x000ff000,
+    0x000ff00f, 0x000ff0f0, 0x000ff0ff, 0x000fff00, 0x000fff0f,
+    0x000ffff0, 0x000fffff, 0x00f00000, 0x00f0000f, 0x00f000f0,
+    0x00f000ff, 0x00f00f00, 0x00f00f0f, 0x00f00ff0, 0x00f00fff,
+    0x00f0f000, 0x00f0f00f, 0x00f0f0f0, 0x00f0f0ff, 0x00f0ff00,
+    0x00f0ff0f, 0x00f0fff0, 0x00f0ffff, 0x00ff0000, 0x00ff000f,
+    0x00ff00f0, 0x00ff00ff, 0x00ff0f00, 0x00ff0f0f, 0x00ff0ff0,
+    0x00ff0fff, 0x00fff000, 0x00fff00f, 0x00fff0f0, 0x00fff0ff,
+    0x00ffff00, 0x00ffff0f, 0x00fffff0, 0x00ffffff, 0x0f000000,
+    0x0f00000f, 0x0f0000f0, 0x0f0000ff, 0x0f000f00, 0x0f000f0f,
+    0x0f000ff0, 0x0f000fff, 0x0f00f000, 0x0f00f00f, 0x0f00f0f0,
+    0x0f00f0ff, 0x0f00ff00, 0x0f00ff0f, 0x0f00fff0, 0x0f00ffff,
+    0x0f0f0000, 0x0f0f000f, 0x0f0f00f0, 0x0f0f00ff, 0x0f0f0f00,
+    0x0f0f0f0f, 0x0f0f0ff0, 0x0f0f0fff, 0x0f0ff000, 0x0f0ff00f,
+    0x0f0ff0f0, 0x0f0ff0ff, 0x0f0fff00, 0x0f0fff0f, 0x0f0ffff0,
+    0x0f0fffff, 0x0ff00000, 0x0ff0000f, 0x0ff000f0, 0x0ff000ff,
+    0x0ff00f00, 0x0ff00f0f, 0x0ff00ff0, 0x0ff00fff, 0x0ff0f000,
+    0x0ff0f00f, 0x0ff0f0f0, 0x0ff0f0ff, 0x0ff0ff00, 0x0ff0ff0f,
+    0x0ff0fff0, 0x0ff0ffff, 0x0fff0000, 0x0fff000f, 0x0fff00f0,
+    0x0fff00ff, 0x0fff0f00, 0x0fff0f0f, 0x0fff0ff0, 0x0fff0fff,
+    0x0ffff000, 0x0ffff00f, 0x0ffff0f0, 0x0ffff0ff, 0x0fffff00,
+    0x0fffff0f, 0x0ffffff0, 0x0fffffff, 0xf0000000, 0xf000000f,
+    0xf00000f0, 0xf00000ff, 0xf0000f00, 0xf0000f0f, 0xf0000ff0,
+    0xf0000fff, 0xf000f000, 0xf000f00f, 0xf000f0f0, 0xf000f0ff,
+    0xf000ff00, 0xf000ff0f, 0xf000fff0, 0xf000ffff, 0xf00f0000,
+    0xf00f000f, 0xf00f00f0, 0xf00f00ff, 0xf00f0f00, 0xf00f0f0f,
+    0xf00f0ff0, 0xf00f0fff, 0xf00ff000, 0xf00ff00f, 0xf00ff0f0,
+    0xf00ff0ff, 0xf00fff00, 0xf00fff0f, 0xf00ffff0, 0xf00fffff,
+    0xf0f00000, 0xf0f0000f, 0xf0f000f0, 0xf0f000ff, 0xf0f00f00,
+    0xf0f00f0f, 0xf0f00ff0, 0xf0f00fff, 0xf0f0f000, 0xf0f0f00f,
+    0xf0f0f0f0, 0xf0f0f0ff, 0xf0f0ff00, 0xf0f0ff0f, 0xf0f0fff0,
+    0xf0f0ffff, 0xf0ff0000, 0xf0ff000f, 0xf0ff00f0, 0xf0ff00ff,
+    0xf0ff0f00, 0xf0ff0f0f, 0xf0ff0ff0, 0xf0ff0fff, 0xf0fff000,
+    0xf0fff00f, 0xf0fff0f0, 0xf0fff0ff, 0xf0ffff00, 0xf0ffff0f,
+    0xf0fffff0, 0xf0ffffff, 0xff000000, 0xff00000f, 0xff0000f0,
+    0xff0000ff, 0xff000f00, 0xff000f0f, 0xff000ff0, 0xff000fff,
+    0xff00f000, 0xff00f00f, 0xff00f0f0, 0xff00f0ff, 0xff00ff00,
+    0xff00ff0f, 0xff00fff0, 0xff00ffff, 0xff0f0000, 0xff0f000f,
+    0xff0f00f0, 0xff0f00ff, 0xff0f0f00, 0xff0f0f0f, 0xff0f0ff0,
+    0xff0f0fff, 0xff0ff000, 0xff0ff00f, 0xff0ff0f0, 0xff0ff0ff,
+    0xff0fff00, 0xff0fff0f, 0xff0ffff0, 0xff0fffff, 0xfff00000,
+    0xfff0000f, 0xfff000f0, 0xfff000ff, 0xfff00f00, 0xfff00f0f,
+    0xfff00ff0, 0xfff00fff, 0xfff0f000, 0xfff0f00f, 0xfff0f0f0,
+    0xfff0f0ff, 0xfff0ff00, 0xfff0ff0f, 0xfff0fff0, 0xfff0ffff,
+    0xffff0000, 0xffff000f, 0xffff00f0, 0xffff00ff, 0xffff0f00,
+    0xffff0f0f, 0xffff0ff0, 0xffff0fff, 0xfffff000, 0xfffff00f,
+    0xfffff0f0, 0xfffff0ff, 0xffffff00, 0xffffff0f, 0xfffffff0,
+    0xffffffff
 };
 
-IWRAM_CODE
-static inline
+// Create a mask for zero sprite values in each nibble.
+// For example: 0x12305008 -> 0xFFF0F00F
+INLINE
 u32
-decode_1bpp(u8 row, u8 flip_x) {
-    if (flip_x) {
-        return dec_byte_flip_x[row];
-    }
-    return dec_byte[row];
+create_zero_mask(u32 x) {
+    x |= x >> 2;
+    x |= x >> 1;
+    x &= 0x11111111;
+    return x * 0xf;
 }
-#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) {
+draw_sprite(size_t x, size_t y, u32 *sprite, u8 clear) {
+    // Copy a 4bpp sprite into memory. Color 0 is the transparency color.
     BOUNDCHECK_SCREEN(x, y);
     size_t tile_x0 = x / 8;
     size_t tile_x1 = (x + 7) / 8;
@@ -491,77 +491,71 @@ draw_chr(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y) {
     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;
-            }
+    size_t n_rows = 8;
+    if (y + 8 > SCREEN_HEIGHT) {
+        n_rows = 8 - ((y + 8) - SCREEN_HEIGHT);
+    }
+
+    size_t n0 = MIN(8 - start_row, n_rows);
+    if (clear) {
+        for(size_t v = 0; v < n0; v++, dst++) {
+            u32 row = sprite[v];
+
+            u32 mask = create_zero_mask(row);
+            u32 msk0 = mask << shift_left;
+            u32 msk1 = mask >> shift_right;
+
+            dst[0] = (dst[0] & ~msk0);
+            dst[8] = (dst[8] & ~msk1);
+        }
+        dst += (32 - 1) * 8;
+        for(size_t v = n0; v < n_rows; v++, dst++) {
+            u32 row = sprite[v];
+
+            u32 mask = create_zero_mask(row);
+            u32 msk0 = mask << shift_left;
+            u32 msk1 = mask >> shift_right;
+
+            dst[0] = (dst[0] & ~msk0);
+            dst[8] = (dst[8] & ~msk1);
         }
     } 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;
-            }
+        for(size_t v = 0; v < n0; v++, dst++) {
+            u32 row = sprite[v];
+
+            u32 mask = create_zero_mask(row);
+            u32 msk0 = mask << shift_left;
+            u32 msk1 = mask >> shift_right;
+            u32 clr0 = row << shift_left;
+            u32 clr1 = row >> shift_right;
+
+            dst[0] = (dst[0] & ~msk0) | clr0;
+            dst[8] = (dst[8] & ~msk1) | clr1;
+        }
+        dst += (32 - 1) * 8;
+        for(size_t v = n0; v < n_rows; v++, dst++) {
+            u32 row = sprite[v];
+
+            u32 mask = create_zero_mask(row);
+            u32 msk0 = mask << shift_left;
+            u32 msk1 = mask >> shift_right;
+            u32 clr0 = row << shift_left;
+            u32 clr1 = row >> shift_right;
+
+            dst[0] = (dst[0] & ~msk0) | clr0;
+            dst[8] = (dst[8] & ~msk1) | clr1;
         }
     }
     dirty_tiles[tile_y] |= dirty;
+    if (start_row != 0) {
+        dirty_tiles[tile_y + 1] |= dirty;
+    }
 }
 
 IWRAM_CODE
 UNROLL_LOOPS
 void
-draw_icn(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y) {
+draw_1bpp(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;
@@ -572,47 +566,63 @@ draw_icn(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y) {
     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
+    u32 color = clr * 0x11111111;
+    u32 *lut = flip_x ? lut_1bpp_mask_flip_x : lut_1bpp_mask;
+    size_t n_rows = 8;
+    if (y + 8 > SCREEN_HEIGHT) {
+        n_rows = 8 - ((y + 8) - SCREEN_HEIGHT);
+    }
+    size_t n0 = MIN(8 - start_row, n_rows);
     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;
-            }
+        for(size_t v = 0; v < n0; v++, dst++) {
+            u32 mask = lut[*sprite];
+            u32 msk0 = mask << shift_left;
+            u32 msk1 = mask >> shift_right;
+            u32 clr0 = msk0 & color;
+            u32 clr1 = msk1 & color;
+            dst[0] = (dst[0] & ~msk0) | clr0;
+            dst[8] = (dst[8] & ~msk1) | clr1;
+            sprite++;
+        }
+        dst += (32 - 1) * 8;
+        for(size_t v = n0; v < n_rows; v++, dst++) {
+            u32 mask = lut[*sprite];
+            u32 msk0 = mask << shift_left;
+            u32 msk1 = mask >> shift_right;
+            u32 clr0 = msk0 & color;
+            u32 clr1 = msk1 & color;
+            dst[0] = (dst[0] & ~msk0) | clr0;
+            dst[8] = (dst[8] & ~msk1) | clr1;
+            sprite++;
         }
     } 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;
-            }
+        sprite += 7;
+        for(size_t v = 0; v < n0; v++, dst++) {
+            u32 mask = lut[*sprite];
+            u32 msk0 = mask << shift_left;
+            u32 msk1 = mask >> shift_right;
+            u32 clr0 = msk0 & color;
+            u32 clr1 = msk1 & color;
+            dst[0] = (dst[0] & ~msk0) | clr0;
+            dst[8] = (dst[8] & ~msk1) | clr1;
+            sprite--;
+        }
+        dst += (32 - 1) * 8;
+        for(size_t v = n0; v < n_rows; v++, dst++) {
+            u32 mask = lut[*sprite];
+            u32 msk0 = mask << shift_left;
+            u32 msk1 = mask >> shift_right;
+            u32 clr0 = msk0 & color;
+            u32 clr1 = msk1 & color;
+            dst[0] = (dst[0] & ~msk0) | clr0;
+            dst[8] = (dst[8] & ~msk1) | clr1;
+            sprite--;
         }
     }
     dirty_tiles[tile_y] |= dirty;
+    if (start_row != 0) {
+        dirty_tiles[tile_y + 1] |= dirty;
+    }
 }
 
 //
@@ -620,9 +630,11 @@ draw_icn(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y) {
 //
 
 IWRAM_CODE
+UNROLL_LOOPS
 void
 flip_buffer(void) {
-// Mode 0: double buffering without dirty tiles.
+// Mode 0: double buffering without dirty tiles. Use this when we are clearing
+// the screen every single frame.
 #if FLIP_TYPE == 0
     if (backbuf == BUF_0) {
         backbuf = BUF_1;
@@ -635,7 +647,7 @@ flip_buffer(void) {
     }
 
 // Mode 1: single buffer, copy the dirty lines from backbuffer (BUF_1) to
-// frontbuffer (BUF_0) using the DMA.
+// frontbuffer (BUF_0).
 #elif FLIP_TYPE == 1
     u32 *front = BUF_0;
     u32 *back = BUF_1;
@@ -646,7 +658,11 @@ flip_buffer(void) {
             continue;
         }
         u32 offset = j * 32 * 8;
+#if NO_DMA == 0
         dma_copy(front + offset, back + offset, (30 * 8 * 4), 3);
+#else
+        copy32(front + offset, back + offset, (30 * 8));
+#endif
         dirty_tiles[j] = 0;
     }
 
@@ -673,7 +689,7 @@ flip_buffer(void) {
     }
 
 // Mode 3: Double buffering with dirty line, copying the dirty lines if needed
-// after flipping buffers with the DMA.
+// after flipping buffers.
 #elif FLIP_TYPE == 3
     bool should_flip = false;
     for (size_t j = 0; j < 20; ++j) {
@@ -701,7 +717,11 @@ flip_buffer(void) {
             continue;
         }
         u32 offset = j * 32 * 8;
+#if NO_DMA == 0
         dma_copy(backbuf + offset, frontbuf + offset, (30 * 8 * 4), 3);
+#else
+        copy32(backbuf + offset, frontbuf + offset, (30 * 8));
+#endif
         dirty_tiles[j] = 0;
     }
 
@@ -746,6 +766,22 @@ flip_buffer(void) {
 #endif
 }
 
+IWRAM_CODE
+UNROLL_LOOPS
+void
+decode_1bpp(u32 *dst, u8 *src, u8 clr, u8 flip_x, u32 n_tiles) {
+    u32 color = 0x11111111 * clr;
+    if (!flip_x) {
+        for (size_t i = 0; i < n_tiles * 8; i++) {
+            *dst++ = lut_1bpp_mask[*src++] & color;
+        }
+    } else {
+        for (size_t i = 0; i < n_tiles * 8; i++) {
+            *dst++ = lut_1bpp_mask_flip_x[*src++] & color;
+        }
+    }
+}
+
 //
 // Text rendering.
 //
@@ -756,9 +792,28 @@ flip_buffer(void) {
 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);
+    draw_1bpp(x, y, tile + 8 * c, clr, 1, 0);
 }
 
+void
+txt_drawc_small(char c, size_t x, size_t y, u8 clr) {
+    u8 *tile = font_icn;
+    c = c < 'a' ? c + 16 * 6 : c + 16 * 4;
+    draw_1bpp(x, y, tile + 8 * c, clr, 1, 0);
+}
+
+#define txt_drawf_small(msg, x, y, clr, ...) \
+        { \
+                char buf[256] = {0}; \
+                posprintf(buf, msg, ##__VA_ARGS__); \
+                u8 tmp = text_engine.spacing;\
+        txt_spacing(4);\
+        text_engine.drawc = txt_drawc_small;\
+                txt_draws(buf, x, y, clr); \
+        txt_spacing(tmp);\
+        text_engine.drawc = txt_drawc;\
+        }
+
 //
 // Initialization.
 //
@@ -769,7 +824,11 @@ renderer_init(void) {
     DISP_CTRL = DISP_MODE_0 | DISP_BG_0 | DISP_BG_1;
 
     // Clear VRAM.
+#if NO_DMA == 0
     dma_fill((u32*)MEM_VRAM, 0, KB(96), 3);
+#else
+    set32((u32*)MEM_VRAM, 0, KB(96)/4);
+#endif
 
     // Initialize backgrounds.
     BG_CTRL(0) = BG_CHARBLOCK(CB_0) | BG_SCREENBLOCK(SB_0) | BG_PRIORITY(0) | BG_SIZE(1);
@@ -796,15 +855,22 @@ renderer_init(void) {
     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_PURPLE;
-    PAL_BUFFER_BG[6] = COLOR_YELLOW;
-    PAL_BUFFER_BG[7] = COLOR_GREEN;
-    PAL_BUFFER_BG[8] = COLOR_GREY;
+    PAL_BUFFER_BG[0x0] = COLOR_BLACK;
+    PAL_BUFFER_BG[0x1] = COLOR_WHITE;
+    PAL_BUFFER_BG[0x2] = COLOR_RED;
+    PAL_BUFFER_BG[0x3] = COLOR_BLUE;
+    PAL_BUFFER_BG[0x4] = COLOR_CYAN;
+    PAL_BUFFER_BG[0x5] = COLOR_PURPLE;
+    PAL_BUFFER_BG[0x6] = COLOR_YELLOW;
+    PAL_BUFFER_BG[0x7] = COLOR_GREEN;
+    PAL_BUFFER_BG[0x8] = COLOR_GREY;
+    PAL_BUFFER_BG[0x9] = COLOR_WHITE;
+    PAL_BUFFER_BG[0xa] = COLOR_RED;
+    PAL_BUFFER_BG[0xb] = COLOR_BLUE;
+    PAL_BUFFER_BG[0xc] = COLOR_CYAN;
+    PAL_BUFFER_BG[0xd] = COLOR_PURPLE;
+    PAL_BUFFER_BG[0xe] = COLOR_YELLOW;
+    PAL_BUFFER_BG[0xf] = COLOR_GREEN;
 
     // Initialize text engine.
     txt_init(txt_drawc);
diff --git a/src/renderer_m3.c b/src/gba/renderer/renderer_m3.c
index 709c982..709c982 100644
--- a/src/renderer_m3.c
+++ b/src/gba/renderer/renderer_m3.c
diff --git a/src/gba/renderer/renderer_m4.c b/src/gba/renderer/renderer_m4.c
new file mode 100644
index 0000000..722899b
--- /dev/null
+++ b/src/gba/renderer/renderer_m4.c
@@ -0,0 +1,475 @@
+#include "renderer.h"
+#include "text.h"
+
+//
+// Parameters.
+//
+
+#define SUBPIXEL_LINES 0
+#define DEC_BIG_LUT 0
+#define DISABLE_BOUNDCHECK_SCREEN 0
+
+// Front/back buffers for double buffering.
+#define BUF_0 ((u32*)(MEM_VRAM))
+#define BUF_1 ((u32*)(MEM_VRAM ^ 0x0A000))
+
+// Pointer to the backbuffer.
+static u16 *backbuf = BUF_1;
+
+// 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 optimiztions but not sure at this moment.
+#if DISABLE_BOUNDCHECK_SCREEN > 0
+#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.
+//
+
+IWRAM_CODE
+void screen_fill(u8 clr)  {
+    set32(backbuf, 0x01010101 * clr, KB(75) / 2 / 4);
+}
+
+IWRAM_CODE
+void
+draw_pixel(size_t x, size_t y, u8 clr) {
+    BOUNDCHECK_SCREEN(x, y);
+    u16 *dst = &backbuf[(x + y * SCREEN_WIDTH) / 2];
+    if(x & 1) {
+        *dst = (*dst &  0xF) | (clr << 8);
+    } else {
+        *dst = (*dst & ~0xF) | clr;
+    }
+}
+
+IWRAM_CODE
+UNROLL_LOOPS
+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 start_col = x0 % 8;
+    size_t end_col = x1 % 8;
+    // 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 dx = tile_x1 - tile_x0;
+    u64 *dst = &backbuf[(tile_x0 * 8 + y0 * SCREEN_WIDTH) / 2];
+    if (dx < 1) {
+        u64 row_mask = 0xFFFFFFFFFFFFFFFF;
+        row_mask >>= (7 - end_col - dx) * 8;
+        row_mask &= 0xFFFFFFFFFFFFFFFF << start_col * 8;
+        u64 row = (0x0101010101010101 * clr) & row_mask;
+        *dst = (*dst & ~row_mask) | row;
+    } else {
+        size_t shift_left = start_col * 8;
+        size_t shift_right = (7 - end_col) * 8;
+        u64 row_mask = 0xFFFFFFFFFFFFFFFF;
+        u64 row = 0x0101010101010101 * clr;
+        dst[0] = (dst[0] & ~(row_mask << shift_left)) | row << shift_left;
+        if (dx != 1) {
+            set32(&dst[1], 0x01010101 * clr, (dx - 1) * 8 / 4);
+        }
+        dst[dx] = dst[dx] & ~(row_mask >> shift_right);
+        dst[dx] |= row >> shift_right;
+    }
+}
+
+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_x0 = x0 / 8;
+    size_t start_col = x0 % 8;
+    u16 *dst = &backbuf[(start_col + tile_x0 * 8 + y0 * SCREEN_WIDTH) / 2];
+    if(x0 & 1) {
+        for (size_t i = 0; i <= y1 - y0; i++, dst += SCREEN_WIDTH / 2) {
+            *dst = (*dst &  0xF) | (clr << 8);
+        }
+    } else {
+        for (size_t i = 0; i <= y1 - y0; i++, dst += SCREEN_WIDTH / 2) {
+            *dst = (*dst & ~0xF) | clr;
+        }
+    }
+}
+
+IWRAM_CODE
+UNROLL_LOOPS
+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 {
+#if SUBPIXEL_LINES == 1
+        // 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;
+        int dxf = (dx << fp_bit);
+        int dyf = (dy << fp_bit);
+
+        if ((dx >= dy && x0 > x1) || (dx < dy && y0 > y1)) {
+            SWAP(x0, x1);
+            SWAP(y0, y1);
+        }
+
+        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 ? -SCREEN_WIDTH : SCREEN_WIDTH;
+
+        u16 *dst = NULL;
+        uintptr_t addr = ((uintptr_t)backbuf + y0 * SCREEN_WIDTH + x0);
+        u32 mask = x0 & 1 ? ~0xFF : 0xFF;
+        u32 color = (clr & 0xFF) | ((clr & 0xFF) << 8);
+        if (dx >= dy) {
+            int distance = (frac_y - fp_one) * dx - (frac_x - fp_half) * dy;
+            int remaining = dx;
+            while (distance <= 0  && remaining > 0) {
+                dst = (u16*)(addr - (mask >> 31));
+                *dst = (*dst & ~mask) | (color & mask);
+                distance += 2 * dyf;
+                addr += x_step;
+                remaining--;
+                mask = ~mask;
+            }
+            distance -= 2 * dxf;
+            addr += y_step;
+
+            while (remaining >= 0) {
+                dst = (u16*)(addr - (mask >> 31));
+                *dst = (*dst & ~mask) | (color & mask);
+                if (distance >= 0) {
+                    distance -= 2 * dxf;
+                    addr += y_step;
+                }
+                distance += 2 * dyf;
+                addr += x_step;
+                mask = ~mask;
+                remaining--;
+            }
+        } else {
+            int distance = (frac_x - fp_one) * dy - (frac_y - fp_half) * dx;
+            int remaining = dy;
+            while (distance <= 0  && remaining > 0) {
+                dst = (u16*)(addr - (mask >> 31));
+                *dst = (*dst & ~mask) | (color & mask);
+                distance += 2 * dxf;
+                addr += y_step;
+                remaining--;
+            }
+            distance -= 2 * dyf;
+            addr += x_step;
+            mask = ~mask;
+
+            while (remaining >= 0) {
+                dst = (u16*)(addr - (mask >> 31));
+                *dst = (*dst & ~mask) | (color & mask);
+                if (distance >= 0) {
+                    distance -= 2 * dyf;
+                    addr += x_step;
+                    mask = ~mask;
+                }
+                distance += 2 * dxf;
+                addr += y_step;
+                remaining--;
+            }
+        }
+#else
+        // Diagonal line.
+        int     dx = x0 > x1 ?       x0 - x1 : x1 - x0;
+        int     dy = y0 > y1 ?       y0 - y1 : y1 - y0;
+        int x_step = x0 > x1 ?            -1 : 1;
+        int y_step = y0 > y1 ? -SCREEN_WIDTH : SCREEN_WIDTH;
+
+        u16 *dst = NULL;
+        uintptr_t addr = ((uintptr_t)backbuf + y0 * SCREEN_WIDTH + x0);
+        u32 mask = x0 & 1 ? ~0xFF : 0xFF;
+        u32 color = (clr & 0xFF) | ((clr & 0xFF) << 8);
+        if (dx >= dy) {
+            int diff = 2 * dy - dx;
+            for (int i = 0; i < dx + 1; i++) {
+                dst = (u16*)(addr - (mask >> 31));
+                *dst = (*dst & ~mask) | (color & mask);
+                if (diff >= 0) {
+                    diff -= 2 * dx;
+                    addr += y_step;
+                }
+                diff += 2 * dy;
+                addr += x_step;
+                mask = ~mask;
+            }
+        } else {
+            int diff = 2 * dx - dy;
+            for (int i = 0; i < dy + 1; i++) {
+                dst = (u16*)(addr - (mask >> 31));
+                *dst = (*dst & ~mask) | (color & mask);
+                if (diff >= 0) {
+                    diff -= 2 * dy;
+                    addr += x_step;
+                    mask = ~mask;
+                }
+                diff += 2 * dx;
+                addr += 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
+    // filling function instead.
+    if (x0 == 0 && x1 >= (SCREEN_WIDTH - 1) && y0 == 0 && y1 >= (SCREEN_HEIGHT - 1)) {
+        screen_fill(clr);
+        return;
+    }
+
+    // Drawline implementation.
+    for (size_t y = y0; y <= y1; y++) {
+        draw_hline(x0, x1, y, clr);
+    }
+}
+
+//
+// Sprites (chr/icn).
+//
+
+extern u32 decode_1bpp_mask[16];
+extern u32 decode_1bpp_mask_flip_x[16];
+
+IWRAM_CODE
+UNROLL_LOOPS
+void
+draw_1bpp(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y) {
+    BOUNDCHECK_SCREEN(x, y);
+    size_t tile_x = x / 8;
+    size_t start_col = x % 8;
+
+    u32 *dst = &backbuf[(y * 30 + tile_x) * 8 / 2];
+    size_t shift_left = start_col * 8;
+    size_t shift_right = (4 - start_col) * 8;
+    if (start_col > 3) {
+        dst++;
+        shift_left = 8 * (start_col - 4);
+        shift_right = 8 * (8 - start_col);
+    }
+
+    size_t n_rows = 8;
+    if (y + 8 > SCREEN_HEIGHT) {
+        n_rows = 8 - ((y + 8) - SCREEN_HEIGHT);
+    }
+
+    u32 color = clr * 0x01010101;
+    if (!flip_y) {
+        if (!flip_x) {
+            for(size_t v = 0; v < n_rows; v++) {
+                // Decode colors.
+                u32 msk0 = decode_1bpp_mask[*sprite >> 4];
+                u32 msk1 = decode_1bpp_mask[*sprite & 0xf];
+                u32 msk2 = (msk0 >> shift_right) | (msk1 << shift_left);
+                msk0 <<= shift_left;
+                msk1 >>= shift_right;
+                u32 clr0 = msk0 & color;
+                u32 clr1 = msk1 & color;
+                u32 clr2 = msk2 & color;
+
+                // Overlay current colors on top of existing framebuffer.
+                dst[0] = (dst[0] & ~msk0) | clr0;
+                dst[1] = (dst[1] & ~msk2) | clr2;
+                dst[2] = (dst[2] & ~msk1) | clr1;
+
+                // Next row.
+                dst += 60;
+                sprite++;
+            }
+        } else {
+            for(size_t v = 0; v < n_rows; v++) {
+                // Decode colors.
+                u32 msk0 = decode_1bpp_mask_flip_x[*sprite & 0xf];
+                u32 msk1 = decode_1bpp_mask_flip_x[*sprite >> 4];
+                u32 msk2 = (msk0 >> shift_right) | (msk1 << shift_left);
+                msk0 <<= shift_left;
+                msk1 >>= shift_right;
+                u32 clr0 = msk0 & color;
+                u32 clr1 = msk1 & color;
+                u32 clr2 = msk2 & color;
+
+                // Overlay current colors on top of existing framebuffer.
+                dst[0] = (dst[0] & ~msk0) | clr0;
+                dst[1] = (dst[1] & ~msk2) | clr2;
+                dst[2] = (dst[2] & ~msk1) | clr1;
+
+                // Next row.
+                dst += 60;
+                sprite++;
+            }
+        }
+    } else {
+        sprite += 7;
+        if (!flip_x) {
+            for(size_t v = 0; v < n_rows; v++) {
+                // Decode colors.
+                u32 msk0 = decode_1bpp_mask[*sprite >> 4];
+                u32 msk1 = decode_1bpp_mask[*sprite & 0xf];
+                u32 msk2 = (msk0 >> shift_right) | (msk1 << shift_left);
+                msk0 <<= shift_left;
+                msk1 >>= shift_right;
+                u32 clr0 = msk0 & color;
+                u32 clr1 = msk1 & color;
+                u32 clr2 = msk2 & color;
+
+                // Overlay current colors on top of existing framebuffer.
+                dst[0] = (dst[0] & ~msk0) | clr0;
+                dst[1] = (dst[1] & ~msk2) | clr2;
+                dst[2] = (dst[2] & ~msk1) | clr1;
+
+                // Next row.
+                dst += 60;
+                sprite--;
+            }
+        } else {
+            for(size_t v = 0; v < 8; v++) {
+                // Decode colors.
+                u32 msk0 = decode_1bpp_mask_flip_x[*sprite & 0xf];
+                u32 msk1 = decode_1bpp_mask_flip_x[*sprite >> 4];
+                u32 msk2 = (msk0 >> shift_right) | (msk1 << shift_left);
+                msk0 <<= shift_left;
+                msk1 >>= shift_right;
+                u32 clr0 = msk0 & color;
+                u32 clr1 = msk1 & color;
+                u32 clr2 = msk2 & color;
+
+                // Overlay current colors on top of existing framebuffer.
+                dst[0] = (dst[0] & ~msk0) | clr0;
+                dst[1] = (dst[1] & ~msk2) | clr2;
+                dst[2] = (dst[2] & ~msk1) | clr1;
+
+                // Next row.
+                dst += 60;
+                sprite--;
+            }
+        }
+    }
+}
+
+//
+// Flipping buffers/copying memory.
+//
+
+IWRAM_CODE
+void
+flip_buffer(void) {
+    backbuf = (u16*)((u32)backbuf ^ 0x0A000);
+    DISP_CTRL ^= DISP_PAGE;
+}
+
+//
+// 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_1bpp(x, y, tile + 8 * c, clr, 1, 0);
+}
+
+void
+txt_drawc_small(char c, size_t x, size_t y, u8 clr) {
+    u8 *tile = font_icn;
+    c = c < 'a' ? c + 16 * 6 : c + 16 * 4;
+    draw_1bpp(x, y, tile + 8 * c, clr, 1, 0);
+}
+
+#define txt_drawf_small(msg, x, y, clr, ...) \
+        { \
+                char buf[256] = {0}; \
+                posprintf(buf, msg, ##__VA_ARGS__); \
+                u8 tmp = text_engine.spacing;\
+        txt_spacing(4);\
+        text_engine.drawc = txt_drawc_small;\
+                txt_draws(buf, x, y, clr); \
+        txt_spacing(tmp);\
+        text_engine.drawc = txt_drawc;\
+        }
+
+//
+// Initialization.
+//
+
+void
+renderer_init(void) {
+    // Initialize display mode and bg palette.
+    DISP_CTRL = DISP_MODE_4 | DISP_BG_2;
+
+    // Clear VRAM.
+    set32((u16*)MEM_VRAM, 0, KB(96)/4);
+
+    // 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_PURPLE;
+    PAL_BUFFER_BG[6] = COLOR_YELLOW;
+    PAL_BUFFER_BG[7] = COLOR_GREEN;
+    PAL_BUFFER_BG[8] = COLOR_GREY;
+
+    // Initialize text engine.
+    txt_init(txt_drawc);
+}
diff --git a/src/gba/text/font.h b/src/gba/text/font.h
new file mode 100644
index 0000000..70a5e28
--- /dev/null
+++ b/src/gba/text/font.h
@@ -0,0 +1,130 @@
+static const 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/posprintf.h b/src/gba/text/posprintf.h
index a560db5..a560db5 100644
--- a/src/text/posprintf.h
+++ b/src/gba/text/posprintf.h
diff --git a/src/text/posprintf.s b/src/gba/text/posprintf.s
index f3ef1e6..f3ef1e6 100644
--- a/src/text/posprintf.s
+++ b/src/gba/text/posprintf.s
diff --git a/src/text/text.h b/src/gba/text/text.h
index 24573f8..87db439 100644
--- a/src/text/text.h
+++ b/src/gba/text/text.h
@@ -92,6 +92,12 @@ 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) {
@@ -119,7 +125,9 @@ txt_draws(char *msg, size_t x, size_t y, u8 clr) {
     size_t i = 0;
     while (*msg) {
         char c = *msg++;
-        text_engine.drawc(c, x + i++ * text_engine.spacing, y, clr);
+        if (c != 0) {
+            text_engine.drawc(c, x + i++ * text_engine.spacing, y, clr);
+        }
     }
 }
 
diff --git a/src/gba/utils.s b/src/gba/utils.s
new file mode 100755
index 0000000..d70d7ef
--- /dev/null
+++ b/src/gba/utils.s
@@ -0,0 +1,82 @@
+.file "utils.s"
+.section .iwram, "ax", %progbits
+.arm
+.align
+
+@ Efficient memcpy32 function (borrowed from TONC). It uses a two step
+@ approach. It tries to copy 8 u32 chunks at a time with the ldm and stm
+@ instructions and then copy the remainder if there are less than 8 chunks
+@ left.
+@
+@   r0: destination address
+@   r1: source address
+@   r2: number of 32bit chunks to copy
+@
+.global copy32
+copy32:
+    cmp     r2, #0
+    beq     .copy32_end
+
+    and     r12, r2, #7     @ r12 = r2 % 8
+    movs    r2, r2, lsr #3  @ r2  = r2 / 8
+    beq     .Lcopy32_residual
+
+    @ Copy 8 32B chunks at a time
+    push    {r4-r10}
+.Lcopy32_chunks:
+    ldmia   r1!, {r3-r10}
+    stmia   r0!, {r3-r10}
+    subs    r2, r2, #1
+    bhi     .Lcopy32_chunks
+    pop     {r4-r10}
+
+    @ Copy residual 32B chunks (0-7)
+.Lcopy32_residual:
+    subs    r12, r12, #1
+    ldrhs   r3, [r1], #4
+    strhs   r3, [r0], #4
+    bhi     .Lcopy32_residual
+
+.copy32_end:
+    bx      lr
+
+
+@ Efficient memset32 function (borrowed from TONC). It uses a two step
+@ approach. Uses the same stmia approach from memcpy32 but, no need for ldmia
+@
+@   r0: destination address
+@   r1: u32 value to set
+@   r2: number of 32bit chunks to set
+@
+.global set32
+set32:
+    cmp     r2, #0
+    beq     .set32_end
+
+    and     r12, r2, #7     @ r12 = r2 % 8
+    movs    r2, r2, lsr #3  @ r2  = r2 / 8
+    beq     .Lset32_residual
+
+    @ Set 8 32B chunks at a time
+    push    {r4-r9}
+    mov     r3, r1
+    mov     r4, r1
+    mov     r5, r1
+    mov     r6, r1
+    mov     r7, r1
+    mov     r8, r1
+    mov     r9, r1
+.Lset32_chunks:
+    stmia   r0!, {r1, r3-r9}
+    subs    r2, r2, #1
+    bhi     .Lset32_chunks
+    pop     {r4-r9}
+
+    @ Set residual 32B chunks (0-7)
+.Lset32_residual:
+    subs    r12, r12, #1
+    strhs   r1, [r0], #4
+    bhi     .Lset32_residual
+
+.set32_end:
+    bx      lr
diff --git a/src/main.c b/src/main.c
index 0a78367..7ca8370 100644
--- a/src/main.c
+++ b/src/main.c
@@ -17,7 +17,7 @@ WITH REGARD TO THIS SOFTWARE.
 // Config parameters.
 //
 
-#define PROF_ENABLE 1
+#define PROF_ENABLE 30
 #include "profiling.c"
 
 void
@@ -42,20 +42,6 @@ test_fill_rect(void) {
 }
 
 void
-test_chr(void) {
-    u8 tile[16] = {
-        0xf8, 0xf8, 0xf8, 0xf8, 0xf8, 0x00, 0x00, 0x00,
-        0x00, 0x00, 0x3e, 0x3e, 0x3e, 0x3e, 0x3e, 0x00
-    };
-
-    for (size_t y = 10; y < 20; y++) {
-        for (size_t x = 20; x < 30; x++) {
-            draw_chr(8 + x * 8, 2 + y * 8, tile, 0, 0, 0);
-        }
-    }
-}
-
-void
 test_icn(void) {
     u8 tile[8] = {
         0xf8, 0xf8, 0xf8, 0xf8 | 0x3e, 0xf8 | 0x3e, 0x3e, 0x3e, 0x00,
@@ -63,7 +49,7 @@ test_icn(void) {
 
     for (size_t y = 10; y < 20; y++) {
         for (size_t x = 10; x < 20; x++) {
-            draw_icn(8 + x * 8, 2 + y * 8, tile, 7, 0, 1);
+            draw_1bpp(8 + x * 8, 2 + y * 8, tile, 7, 0, 1);
         }
     }
 }
@@ -96,13 +82,16 @@ test_moving_line(void) {
         } else {
             should_move = 1;
         }
+        if (key_hold(KEY_SELECT) && key_tap(KEY_START)) {
+            PROF_SHOW();
+        }
         bios_vblank_wait();
         FRAME_START();
-        PROF(flip_buffer(), flip_cycles);
+        PROF(flip_buffer(), PROF_FLIP);
 #if FLIP_TYPE == 0
-        PROF(screen_fill(3), clear_cycles);
+        PROF(screen_fill(3), PROF_FILL);
 #endif
-        PROF(draw_line(x, y, 239 - x, 159 - y, 3), line_cycles);
+        PROF(draw_line(x, y, 239 - x, 159 - y, 3), PROF_LINE);
         x += inc_x * should_move;
         y += inc_y * should_move;
         if (x == 239 && inc_x == 1) {
@@ -115,9 +104,8 @@ test_moving_line(void) {
             inc_x = 1;
             inc_y = 0;
         }
-        PROF(draw_line(x, y, 239 - x, 159 - y, 2), line_cycles);
+        PROF(draw_line(x, y, 239 - x, 159 - y, 2), PROF_LINE);
         FRAME_END();
-        PROF_SHOW();
     }
 }
 
@@ -128,17 +116,18 @@ test_all_static(void) {
         if (key_tap(KEY_A)) {
             break;
         }
+        if (key_hold(KEY_SELECT) && key_tap(KEY_START)) {
+            PROF_SHOW();
+        }
         bios_vblank_wait();
         FRAME_START();
-        PROF(flip_buffer(), flip_cycles);
-        PROF(test_clear(), clear_cycles);
-        PROF(test_lines(), line_cycles);
-        PROF(test_rect(), rect_cycles);
-        PROF(test_fill_rect(), fill_rect_cycles);
-        PROF(test_chr(), chr_cycles);
-        PROF(test_icn(), icn_cycles);
+        PROF(flip_buffer(), PROF_FLIP);
+        PROF(test_clear(), PROF_FILL);
+        PROF(test_lines(), PROF_LINE);
+        PROF(test_rect(), PROF_RECT);
+        PROF(test_fill_rect(), PROF_FRECT);
+        PROF(test_icn(), PROF_ICN);
         FRAME_END();
-        PROF_SHOW();
     }
 }
 
@@ -169,6 +158,9 @@ test_growing_rects(void) {
         if (key_tap(KEY_A)) {
             break;
         }
+        if (key_hold(KEY_SELECT) && key_tap(KEY_START)) {
+            PROF_SHOW();
+        }
         if (key_pressed(KEY_B)) {
             if (key_pressed(KEY_LEFT)) {
                 rects[0].x0 = CLAMP(rects[0].x0 - 1, 0, SCREEN_WIDTH - 1);
@@ -200,22 +192,21 @@ test_growing_rects(void) {
         }
         bios_vblank_wait();
         FRAME_START();
-        PROF(flip_buffer(), flip_cycles);
-        PROF(screen_fill(0), clear_cycles);
+        PROF(flip_buffer(), PROF_FLIP);
+        PROF(screen_fill(0), PROF_FILL);
         PROF(draw_filled_rect(
                 rects[0].x0,
                 rects[0].y0,
                 rects[0].x1,
                 rects[0].y1,
-                2), rect_cycles);
+                2), PROF_RECT);
         PROF(draw_rect(
                 rects[1].x0,
                 rects[1].y0,
                 rects[1].x1,
                 rects[1].y1,
-                3), fill_rect_cycles);
+                3), PROF_FRECT);
         FRAME_END();
-        PROF_SHOW();
     }
 }
 
@@ -224,10 +215,6 @@ test_sprites_bounce(void) {
     u8 sprite_icn[8] = {
         0xf8, 0xf8, 0xf8, 0xf8 | 0x3e, 0xf8 | 0x3e, 0x3e, 0x3e, 0x00,
     };
-    u8 sprite_chr[16] = {
-        0xf8, 0xf8, 0xf8, 0xf8, 0xf8, 0x00, 0x00, 0x00,
-        0x00, 0x00, 0x3e, 0x3e, 0x3e, 0x3e, 0x3e, 0x00,
-    };
     typedef struct Sprite {
         int x;
         int y;
@@ -260,36 +247,29 @@ test_sprites_bounce(void) {
         } else {
             should_move = 1;
         }
+        if (key_hold(KEY_SELECT) && key_tap(KEY_START)) {
+            PROF_SHOW();
+        }
         bios_vblank_wait();
         FRAME_START();
-        PROF(flip_buffer(), flip_cycles);
+        PROF(flip_buffer(), PROF_FLIP);
 #if FLIP_TYPE == 0
-        PROF(screen_fill(8), clear_cycles);
+        PROF(screen_fill(8), PROF_FILL);
 #endif
-        for (size_t i = 0; i < LEN(sprites) * should_move; i++) {
+        for (size_t i = 0; i < LEN(sprites); i++) {
             Sprite *s = &sprites[i];
-            PROF(draw_icn(
+            PROF(draw_1bpp(
                     s->x, s->y,
                     sprite_icn,
                     8,
-                    s->flip_x, s->flip_y), icn_cycles);
-            PROF(draw_chr(
-                    (240 - s->x - 8), (160 - s->y - 8),
-                    sprite_chr,
-                    8,
-                    s->flip_x, s->flip_y), chr_cycles);
-            s->x += s->inc_x;
-            s->y += s->inc_y;
-            PROF(draw_icn(
+                    s->flip_x, s->flip_y), PROF_ICN);
+            s->x += s->inc_x * should_move;
+            s->y += s->inc_y * should_move;
+            PROF(draw_1bpp(
                     s->x, s->y,
                     sprite_icn,
                     s->clr,
-                    s->flip_x, s->flip_y), icn_cycles);
-            PROF(draw_chr(
-                    (240 - s->x - 8), (160 - s->y - 8),
-                    sprite_chr,
-                    s->clr,
-                    s->flip_x, s->flip_y), chr_cycles);
+                    s->flip_x, s->flip_y), PROF_ICN);
             if (s->x >= (240 - 8) && s->inc_x > 0) {
                 s->inc_x *= -1;
             } else if (s->x <= 0 && s->inc_x < 0){
@@ -302,7 +282,6 @@ test_sprites_bounce(void) {
             }
         }
         FRAME_END();
-        PROF_SHOW();
     }
 }
 
@@ -313,27 +292,29 @@ test_text_rendering(void) {
         if (key_tap(KEY_A)) {
             break;
         }
+        if (key_hold(KEY_SELECT) && key_tap(KEY_START)) {
+            PROF_SHOW();
+        }
         bios_vblank_wait();
         FRAME_START();
-        PROF(flip_buffer(), flip_cycles);
-        PROF(screen_fill(0), clear_cycles);
+        PROF(flip_buffer(), PROF_FLIP);
+        PROF(screen_fill(0), PROF_FILL);
         txt_color(2);
-        PROF(txt_drawf("The strongest bulwark of", 4, 8 * 2, 3), txt_drawf_cycles);
-        PROF(txt_drawf("authority is uniformity;", 2, 8 * 3, 3), txt_drawf_cycles);
-        PROF(txt_drawf("the least divergence from it", 8, 8 * 4, 3), txt_drawf_cycles);
-        PROF(txt_drawf("it's the greatest crime", 6, 8 * 5, 3), txt_drawf_cycles);
-        PROF(txt_drawf("- Emma Goldman", 100, 8 * 6 + 3, 3), txt_drawf_cycles);
+        PROF(txt_drawf("The strongest bulwark of", 4, 8 * 2, 3), PROF_DRAWTXT);
+        PROF(txt_drawf("authority is uniformity;", 2, 8 * 3, 3), PROF_DRAWTXT);
+        PROF(txt_drawf("the least divergence from it", 8, 8 * 4, 3), PROF_DRAWTXT);
+        PROF(txt_drawf("it's the greatest crime", 6, 8 * 5, 3), PROF_DRAWTXT);
+        PROF(txt_drawf("- Emma Goldman", 100, 8 * 6 + 3, 3), PROF_DRAWTXT);
         txt_position(0, 10);
         PROF(txt_printf("The only way to deal with an\n"
                    "unfree world is to become\n"
                    "so absolutely free,\n"
                    "that your very existence\n"
-                   "is an act of rebellion.\n"), txt_printf_cycles);
-        PROF(txt_printf("\n              - Albert Camus\n"), txt_printf_cycles);
-        PROF(txt_render(), txt_render_cycles);
-        PROF(txt_clear(), txt_clear_cycles);
+                   "is an act of rebellion.\n"), PROF_PRINTXT);
+        PROF(txt_printf("\n              - Albert Camus\n"), PROF_PRINTXT);
+        PROF(txt_render(), PROF_TXTREND);
+        PROF(txt_clear(), PROF_TXTCLEAR);
         FRAME_END();
-        PROF_SHOW();
     }
 }
 
@@ -348,6 +329,7 @@ int main(void) {
     irq_init();
     irs_set(IRQ_VBLANK, irs_stub);
 
+    PROF_INIT();
     while (true) {
         test_sprites_bounce();
         test_text_rendering();
diff --git a/src/profiling.c b/src/profiling.c
index 90215c8..e1f4feb 100644
--- a/src/profiling.c
+++ b/src/profiling.c
@@ -2,199 +2,143 @@
 // Profiling macros.
 //
 
+// NOTE: Profiling uses the last two timers to count cycles, and thus can't be
+//       used for measuring applications that
+
+
+static u32 frame_time = 0;
+
 #ifndef PROF_ENABLE
 #define PROF_ENABLE 0
 #endif
 
-#if PROF_ENABLE > 0 && PROF_ENABLE < 3
+#if PROF_ENABLE > 0
 
-#ifndef PROF_N_FRAMES
-#define PROF_N_FRAMES 30
+#ifndef PROF_RESET_MINMAX
+#define PROF_RESET_MINMAX false
 #endif
 
-// Profile method 1: Average per N frames.
-#if PROF_ENABLE == 1
-#define TEXT_ENABLE 1
-#define PROF(F,VAR) \
-    do { \
-        u32 __tmp_prof = profile_measure();\
-        F;\
-        (VAR) += profile_measure() - __tmp_prof;\
-    } while (0)
-
-// Profile method 2: Maximum in N frames.
-#elif PROF_ENABLE == 2
-#define TEXT_ENABLE 1
-#define PROF(F,VAR) \
-    do { \
-        u32 __tmp_prof = profile_measure();\
-        (F);\
-        (VAR) = MAX(profile_measure() - __tmp_prof, (VAR));\
-    } while (0)
-#endif
+// Maximum number of profiling to monitor.
+typedef enum ProfType {
+    PROF_FLIP,
+    PROF_FILL,
+    PROF_LINE,
+    PROF_RECT,
+    PROF_FRECT,
+    PROF_ICN,
+    PROF_DRAWTXT,
+    PROF_PRINTXT,
+    PROF_TXTREND,
+    PROF_TXTCLEAR,
+    PROF_NUM,
+} ProfType;
 
-#ifndef PROF_SHOW_X
-#define PROF_SHOW_X 0
-#endif
-#ifndef PROF_SHOW_Y
-#define PROF_SHOW_Y 0
-#endif
+char *prof_type_str[PROF_NUM] = {
+    "FLIPBUF  ",
+    "SCRFILL  ",
+    "LINE     ",
+    "RECT     ",
+    "FRECT    ",
+    "ICN      ",
+    "DRAWTXT  ",
+    "PRINTXT  ",
+    "TXTREND  ",
+    "TXTCLEAR ",
+};
 
-static bool profile_show = true;
-static bool profile_bg_show = true;
-
-#define PROF_SHOW() \
-    do { \
-        if (key_tap(KEY_START)) {\
-            profile_show ^= 1;\
-        }\
-        if (key_tap(KEY_SELECT)) {\
-            profile_bg_show ^= 1;\
-        }\
-        if (profile_show) {\
-            txt_color(1);\
-            txt_position((PROF_SHOW_X), (PROF_SHOW_Y));\
-            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("TOTAL   %.8lu\n", avg_frame_cycles);\
-            txt_render();\
-        }\
-        if (profile_bg_show) {\
-            u32 frame_time =\
-                FP_DIV(\
-                    FP_NUM(avg_frame_cycles + 1, 2),\
-                    FP_NUM(2809, 2),\
-                    2) * 166;\
-            u32 fps =\
-                FP_DIV(\
-                    FP_NUM(280896 * 60, 2),\
-                    FP_NUM(avg_frame_cycles + 1, 2),\
-                    2);\
-            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 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;
-
-#if PROF_ENABLE == 1
-#define FRAME_START()\
-    do { \
-        if (prof_frame_counter == PROF_N_FRAMES) {\
-            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;\
-            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;\
-            prof_frame_counter    = 0;\
-        }\
-        profile_start();\
-    } while (0)
-#elif PROF_ENABLE == 2
-#define FRAME_START()\
-    do { \
-        if (prof_frame_counter == PROF_N_FRAMES) {\
-            avg_frame_cycles      = frame_cycles;\
-            avg_flip_cycles       = flip_cycles;\
-            avg_clear_cycles      = clear_cycles;\
-            avg_line_cycles       = line_cycles;\
-            avg_rect_cycles       = rect_cycles;\
-            avg_fill_rect_cycles  = fill_rect_cycles;\
-            avg_chr_cycles        = chr_cycles;\
-            avg_icn_cycles        = icn_cycles;\
-            avg_txt_drawf_cycles  = txt_drawf_cycles;\
-            avg_txt_printf_cycles = txt_printf_cycles;\
-            avg_txt_render_cycles = txt_render_cycles;\
-            avg_txt_clear_cycles  = txt_clear_cycles;\
-            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;\
-            prof_frame_counter    = 0;\
-        }\
-        profile_start();\
-    } while (0)
-#endif
+u32 prof_frame_time  = 0;
+u32 prof_frame_count = 0;
+u32 prof_frame_avg  = 0;
+u32 prof_frame_time_max  = 0;
+u32 prof_times[PROF_NUM] = {0};
+u32 prof_count[PROF_NUM] = {0};
+u32 prof_avg[PROF_NUM] = {0};
+u32 prof_max[PROF_NUM] = {0};
+u32 prof_min[PROF_NUM] = {0};
+
+bool prof_reset_minmax = PROF_RESET_MINMAX;
+bool prof_show = true;
+
+#define PROF_INIT() do { \
+    for (size_t i = 0; i < PROF_NUM; i++) { \
+        prof_min[i] = -1; \
+    } \
+} while(0);
+
+#define PROF(func, idx) do { \
+    u32 time_before = profile_measure(); \
+    (func); \
+    u32 time_after = profile_measure(); \
+    u32 time_current = time_after - time_before; \
+    prof_times[idx] += time_current; \
+    prof_count[idx]++; \
+    prof_max[idx] = MAX(time_current, prof_max[idx]);\
+    prof_min[idx] = MIN(time_current, prof_min[idx]);\
+} while(0);
+
+#define FRAME_START() do { \
+    profile_start();\
+} while(0)
+
+
+#define FRAME_END() do { \
+        prof_frame_count++;\
+        frame_time = profile_measure();\
+        prof_frame_time_max = MAX(prof_frame_time_max, frame_time);\
+        prof_frame_time += profile_stop();\
+        if (prof_show) { \
+            draw_filled_rect(0, 0, SCREEN_WIDTH - 1, 8 * (PROF_NUM + 1), 2); \
+            u32 fps = (u64)280896 * 60 / (prof_frame_avg + 1); \
+            if (prof_frame_avg == 0) { \
+                fps = 0; \
+            } \
+            txt_drawf_small("FRAME TIME/FPS: %.9l/%.2l", 0, 0, COL_FG, \
+                    prof_frame_avg, fps);\
+            txt_drawf_small("MAX: %.9l/%l", 8 * 19, 0, COL_FG, \
+                    prof_frame_time_max, 280896);\
+            for (size_t idx = 0; idx < PROF_NUM; idx++) { \
+                txt_drawf_small("%s %.9l (%.9l %.9l) %08x:%08x", 0, 8 * (idx + 1), COL_FG, \
+                        prof_type_str[idx], \
+                        prof_avg[idx], \
+                        prof_min[idx], \
+                        prof_max[idx], \
+                        prof_avg[idx], \
+                        prof_max[idx]);\
+            }; \
+            draw_filled_rect(0, SCREEN_HEIGHT - 9, 58, SCREEN_HEIGHT - 1, 2); \
+            txt_drawf_small("CPU USAGE: %.3l", 0, SCREEN_HEIGHT - 9, COL_FG, \
+                    (u64)prof_frame_avg * 100 / 280896);\
+        } \
+        if (prof_frame_count >= PROF_ENABLE) { \
+            for (size_t idx = 0; idx < PROF_NUM; idx++) { \
+                prof_avg[idx] = prof_times[idx] / prof_frame_count; \
+                if (prof_reset_minmax) { \
+                    prof_min[idx] = -1; \
+                    prof_max[idx] = 0; \
+                } \
+                prof_times[idx] = 0; \
+                prof_count[idx] = 0; \
+            }; \
+            prof_frame_avg = prof_frame_time / prof_frame_count; \
+            prof_frame_count = 0; \
+            prof_frame_time = 0; \
+        } \
+    } while(0)
 
-#define FRAME_END() \
-    do { \
-        prof_frame_counter++;\
-        frame_cycles += profile_stop();\
-    } while (0)
+#define PROF_SHOW() do { \
+    prof_show ^= 1; \
+} while(0)
 
 #else
 
 // No profiling.
-#define PROF(F,VAR) (F)
+#define PROF_INIT()
+#define PROF(F,VAR) do {F;} while(0)
+#define FRAME_START() do { \
+    profile_start();\
+} while(0)
+#define FRAME_END() do { \
+    frame_time = profile_stop();\
+} while(0)
 #define PROF_SHOW()
-#define FRAME_START()
-#define FRAME_END()
 #endif
diff --git a/src/renderer.c b/src/renderer.c
deleted file mode 100644
index 5ef9cab..0000000
--- a/src/renderer.c
+++ /dev/null
@@ -1,653 +0,0 @@
-//
-// 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.
-//
-// In addition to the frontbuffer (displayed on background 0), a tiled text
-// layer is displayed on background 1, which can be used for application
-// development or for debug information.
-//
-// These two layers occupy the first and second background charblocks, leaving
-// the remaining two available for other background layers. There are 14KB of
-// sprite memory available, since the backbuffer is located at the end of the
-// VRAM, but if more space is needed it can be moved to the end of the BG
-// charblocks instead.
-//
-
-#include "renderer.h"
-#include "text.h"
-
-// The frontbuffs,eer 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)))
-
-// Keep track of which tiles need to be copied to the frontbuffer.
-static u32 dirty_tiles[21] = {0};
-
-// 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 optimiztions 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); }
-
-IWRAM_CODE
-void screen_fill(u8 clr)  {
-// #if 0
-//     u32 *dst = backbuf;
-//     for(int i = 0; i < KB(75) / 8; i++) {
-//         *dst++ = 0x01010101 * clr;
-//     }
-// #else
-//     dma_fill(backbuf, 0x01010101 * clr, KB(75) / 2, 3);
-//     screen_updated = true;
-// #endif
-}
-
-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;
-    size_t pos = start_row + (tile_x + tile_y * 32) * 8;
-
-    // Update backbuffer.
-    size_t shift = start_col * sizeof(u32);
-    BACKBUF[pos] = (BACKBUF[pos] & ~(0xF << shift)) | clr << shift;
-
-    // Mark tile as dirty.
-    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) {
-    // TODO
-}
-
-IWRAM_CODE
-UNROLL_LOOPS
-static inline
-void
-draw_vline(size_t x0, size_t y0, size_t y1, u8 clr) {
-    // TODO
-}
-
-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 {
-        // // Diagonal line.
-        // int     dx = x0 > x1 ?       x0 - x1 : x1 - x0;
-        // int     dy = y0 > y1 ?       y0 - y1 : y1 - y0;
-        // int x_step = x0 > x1 ?            -1 : 1;
-        // int y_step = y0 > y1 ? -SCREEN_WIDTH : SCREEN_WIDTH;
-
-        // u16 *dst = NULL;
-        // uintptr_t addr = ((uintptr_t)backbuf + y0 * SCREEN_WIDTH + x0);
-        // u32 mask = x0 & 1 ? ~0xFF : 0xFF;
-        // u32 color = (clr & 0xFF) | ((clr & 0xFF) << 8);
-        // if (dx >= dy) {
-        //     int diff = 2 * dy - dx;
-        //     for (int i = 0; i < dx + 1; i++) {
-        //         dst = (u16 *)(addr - (mask >> 31));
-        //         *dst = (*dst & ~mask) | (color & mask);
-        //         if (diff >= 0) {
-        //             diff -= 2 * dx;
-        //             addr += y_step;
-        //         }
-        //         diff += 2 * dy;
-        //         addr += x_step;
-        //         mask = ~mask;
-        //     }
-        // } else {
-        //     int diff = 2 * dx - dy;
-        //     for (int i = 0; i < dy + 1; i++) {
-        //         dst = (u16 *)(addr - (mask >> 31));
-        //         *dst = (*dst & ~mask) | (color & mask);
-        //         if (diff >= 0) {
-        //             diff -= 2 * dy;
-        //             addr += x_step;
-        //             mask = ~mask;
-        //         }
-        //         diff += 2 * dx;
-        //         addr += y_step;
-        //     }
-        // }
-    }
-
-    // // Find row positions for the given x/y coordinates.
-    // size_t tile_x0 = x0 / 8;
-    // size_t tile_y0 = y0 / 8;
-    // size_t tile_x1 = x1 / 8;
-    // size_t tile_y1 = y1 / 8;
-    // size_t start_col0 = x0 % 8;
-    // size_t start_col1 = x1 % 8;
-    // size_t start_row0 = y0 % 8;
-    // size_t start_row1 = y1 % 8;
-
-    // // Get a pointer to the backbuffer and the tile row.
-    // u32 *backbuffer = &BACKBUF[start_row0 + (tile_x0 + tile_y0 * 32) * 8];
-
-    // if (y0 == y1) {
-    //     // 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.
-    //     size_t dx = tile_x1 - tile_x0;
-    //     if (dx < 1) {
-    //         u32 row_mask = 0xFFFFFFFF;
-    //         row_mask >>= (7 - start_col1 - dx) * 4;
-    //         row_mask &= 0xFFFFFFFF << start_col0 * 4;
-    //         u32 row = (0x11111111 * clr) & row_mask;
-    //         backbuffer[0] = (backbuffer[0] & ~row_mask) | row;
-    //         dirty_tiles[tile_y0] |= 1 << tile_x0;
-    //     } else {
-    //         size_t shift_left = start_col0 * 4;
-    //         size_t shift_right = (7 - start_col1) * 4;
-    //         u32 row_mask = 0xFFFFFFFF;
-    //         u32 row = 0x11111111 * clr;
-    //         backbuffer[0] = backbuffer[0] & ~(row_mask << shift_left);
-    //         backbuffer[0] |= row << shift_left;
-    //         dirty_tiles[tile_y0] |= 1 << tile_x0;
-    //         for (size_t i = 1; i < dx; i++) {
-    //             backbuffer[i * 8] = row;
-    //             dirty_tiles[tile_y0] |= 1 << (tile_x0 + i);
-    //         }
-    //         backbuffer[dx * 8] = backbuffer[dx * 8] & ~(row_mask >> shift_right);
-    //         backbuffer[dx * 8] |= row >> shift_right;
-    //         dirty_tiles[tile_y0] |= 1 << (tile_x0 + dx);
-    //     }
-    // } else if (x0 == x1) {
-    //     // Vertical line. The cases are analogous to the horizontal ones.
-    //     size_t dy = tile_y1 - tile_y0;
-    //     u32 row_mask = 0xF << start_col0 * 4;
-    //     u32 row_left = (0x11111111 * clr) & row_mask;
-    //     if (dy < 1) {
-    //         for (size_t i = 0; i <= y1 - y0; i++, backbuffer++) {
-    //             backbuffer[0] = (backbuffer[0] & ~row_mask) | row_left;
-    //         }
-    //     } else {
-    //         for (size_t i = 0; i < (8 - start_row0); i++, backbuffer++) {
-    //             backbuffer[0] = (backbuffer[0] & ~row_mask) | row_left;
-    //         }
-    //         dirty_tiles[tile_y0] |= 1 << tile_x0;
-    //         backbuffer += 8 * 31;
-    //         for (size_t j = 1; j < dy; j++) {
-    //             for (size_t i = 0; i < 8; i++, backbuffer++) {
-    //                 backbuffer[0] = (backbuffer[0] & ~row_mask) | row_left;
-    //             }
-    //             backbuffer += 8 * 31;
-    //             dirty_tiles[tile_y0 + j] |= 1 << tile_x0;
-    //         }
-    //         for (size_t i = 0; i <= start_row1; i++, backbuffer++) {
-    //             backbuffer[0] = (backbuffer[0] & ~row_mask) | row_left;
-    //         }
-    //         dirty_tiles[tile_y1] |= 1 << tile_x0;
-    //     }
-    // } else {
-    //     // Diagonal line.
-    //     int dx = x0 > x1 ? x0 - x1 : x1 - x0;
-    //     int dy = y0 > y1 ? y1 - y0 : y0 - y1;
-    //     int x_step = x0 < x1 ? 1 : -1;
-    //     int y_step = y0 < y1 ? 1 : -1;
-    //     int err = dx + dy;
-    //     while (!(x0 == x1 && y0 == y1)) {
-    //         draw_pixel(x0, y0, clr);
-    //         int diff = 2 * err;
-    //         if (diff >= dy) {
-    //             err += dy;
-    //             x0 += x_step;
-    //         }
-    //         if (diff <= dx) {
-    //             err += dx;
-    //             y0 += y_step;
-    //         }
-    //     }
-    // }
-}
-
-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);
-
-    // TODO: check if this is better.
-//     BOUNDCHECK_SCREEN(x0, y0);
-//     BOUNDCHECK_SCREEN(x1, y1);
-
-//     // Find row positions for the given x/y coordinates.
-//     size_t tile_x0 = x0 / 8;
-//     size_t tile_y0 = y0 / 8;
-//     size_t tile_x1 = x1 / 8;
-//     size_t tile_y1 = y1 / 8;
-//     size_t start_col0 = x0 % 8;
-//     size_t start_col1 = x1 % 8;
-//     size_t start_row0 = y0 % 8;
-//     size_t start_row1 = y1 % 8;
-
-//     // Get a pointer to the backbuffer and the tile row.
-//     u32 *buf_top = &BACKBUF[start_row0 + (tile_x0 + tile_y0 * 32) * 8];
-//     u32 *buf_bot = &BACKBUF[start_row1 + (tile_x0 + tile_y1 * 32) * 8];
-
-//     size_t dx = tile_x1 - tile_x0;
-//     size_t dy = tile_y1 - tile_y0;
-
-//     memset(buf_top, 3, x1 - x0);
-//     memset(buf_bot, 3, x1 - x0);
-
-//     if (dx < 1) {
-//         dirty_tiles[tile_y0] |= 1 << tile_x0;
-//         dirty_tiles[tile_y1] |= 1 << tile_x0;
-//     } else {
-//         dirty_tiles[tile_y0] |= 1 << tile_x0;
-//         dirty_tiles[tile_y1] |= 1 << tile_x0;
-//         for (size_t i = 1; i < dx; i++) {
-//             dirty_tiles[tile_y0] |= 1 << (tile_x0 + i);
-//             dirty_tiles[tile_y1] |= 1 << (tile_x0 + i);
-//         }
-//         dirty_tiles[tile_y0] |= 1 << (tile_x0 + dx);
-//         dirty_tiles[tile_y1] |= 1 << (tile_x0 + dx);
-//     }
-//     if (dy < 1) {
-//     } else {
-//         for (size_t j = 1; j < dy; j++) {
-//             dirty_tiles[tile_y0 + j] |= 1 << tile_x0;
-//             dirty_tiles[tile_y0 + j] |= 1 << (tile_x0 + dx);
-//         }
-//     }
-}
-
-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;
-    }
-
-    // Drawline implementation.
-    for (size_t y = y0; y <= y1; y++) {
-        // NOTE: Unclear why here draw_hline is faster than draw_line.
-        draw_hline(x0, x1, y, clr);
-    }
-
-    // TODO: check if this is better.
-    // BOUNDCHECK_SCREEN(x0, y0);
-    // BOUNDCHECK_SCREEN(x1, y1);
-
-    // size_t dx = x1 - x0;
-    // size_t dy = y1 - y0;
-    // u8 *buf = &BACKBUF[0];
-    // memset(buf, 0x11 * clr, 16);
-    //for (size_t j = 0; j < 1; j++) {
-    //    // for (size_t i = 0; i < dx; i++) {
-    //    //     buf[i + j * 16] = clr;
-    //    // }
-    //    //
-    //    // BACKBUF[j + 0] = 0x11111111 * clr;
-    //    // BACKBUF[j + 1] = 0x11111111 * clr;
-    //    // BACKBUF[j + 2] = 0x11111111 * clr;
-    //    // BACKBUF[j + 3] = 0x11111111 * clr;
-    //    // BACKBUF[j + 4] = 0x11111111 * clr;
-    //    // BACKBUF[j + 5] = 0x11111111 * clr;
-    //    // BACKBUF[j + 6] = 0x11111111 * clr;
-    //    // BACKBUF[j + 7] = 0x11111111 * clr;
-
-    //    buf[j + 0] = 0x1 * clr;
-    //    buf[j + 1] = 0x1 * clr;
-    //    buf[j + 2] = 0x1 * clr;
-    //    buf[j + 3] = 0x1 * clr;
-    //    // buf[j + 4] = 0x1111 * clr;
-    //    // buf[j + 5] = 0x1111 * clr;
-    //    // buf[j + 6] = 0x1111 * clr;
-    //    // buf[j + 7] = 0x1111 * clr;
-    //}
-    // u8 *buf = &BACKBUF[0];
-    // buf[8 * 16 + 0] = clr;
-    // buf[8 * 16 + 1] = clr;
-    // buf[8 * 16 + 2] = clr;
-    // buf[8 * 16 + 3] = clr;
-    // buf[8 * 16 + 4] = clr;
-    // buf[8 * 16 + 5] = clr;
-    // buf[8 * 16 + 6] = clr;
-    // buf[8 * 16 + 7] = clr;
-    // for (size_t j = 0; j < dy; j++) {
-    //     for (size_t i = 0; i < dx; i++) {
-    //         buf[i + j * 16] = clr;
-    //     }
-    // }
-    // size_t n_rect = MIN(dx, dy);
-    // n_rect = n_rect / 2 + 1;
-    // for (size_t i = 0; i < n_rect; i++) {
-    //     draw_rect(x0 + i, y0 + i, x1 - i, y1 - i, clr);
-    // }
-}
-
-// IWRAM_CODE
-// void
-// draw_tile(size_t x, size_t y, Tile *tile, 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;
-
-//     // Get a pointer to the backbuffer and the tile row.
-//     size_t pos = start_row + (tile_x + tile_y * 32) * 8;
-//     u32 *backbuffer = &BACKBUF[pos];
-//     u32 *row = tile;
-
-//     // This will blend all colors weirdly if using tiles that contain colors
-//     // higher than 1.
-//     size_t shift_left = start_col * 4;
-//     size_t shift_right = (8 - start_col) * 4;
-//     // u32 row_mask_left = merge ? 0 : 0xFFFFFFFF << shift_left;
-//     // u32 row_mask_right = merge ? 0 : 0xFFFFFFFF >> shift_right;
-
-//     // Draw the tiles. There are 4 possible cases:
-//     //     1. The tile is exactly at the tile boundary.
-//     //     2. The tile spans 2 tiles horizontally.
-//     //     3. The tile spans 2 tiles vertically.
-//     //     4. The tile spans 4 tiles.
-//     if (start_col == 0 && start_row == 0) {
-//         for (size_t i = 0; i < (8 - start_row); i++, backbuffer++) {
-//             BOUNDCHECK_SCREEN(x, y + i);
-//             backbuffer[0] = (backbuffer[0] & ~row_mask_left) | row[i] * clr;
-//         }
-//         dirty_tiles[tile_y] |= 1 << tile_x;
-//     } else if (start_row == 0) {
-//         for (size_t i = 0; i < 8; i++, backbuffer++) {
-//             BOUNDCHECK_SCREEN(x, y + i);
-//             backbuffer[0] = (backbuffer[0] & ~row_mask_left) | (row[i] * clr << shift_left);
-//             backbuffer[8] = (backbuffer[8] & ~row_mask_right) | (row[i] * clr >> shift_right);
-//         }
-//         dirty_tiles[tile_y] |= 1 << tile_x;
-//         dirty_tiles[tile_y] |= 1 << (tile_x + 1);
-//     } else if (start_col == 0) {
-//         for (size_t i = 0; i < (8 - start_row); i++, backbuffer++) {
-//             BOUNDCHECK_SCREEN(x, y + i);
-//             backbuffer[0] = (backbuffer[0] & ~row_mask_left) | row[i] * clr;
-//         }
-//         backbuffer += 8 * 31;
-//         for (size_t i = (8 - start_row); i < 8; i++, backbuffer++) {
-//             BOUNDCHECK_SCREEN(x, y + i);
-//             backbuffer[0] = (backbuffer[0] & ~row_mask_left) | row[i] * clr;
-//         }
-//         dirty_tiles[tile_y] |= 1 << tile_x;
-//         dirty_tiles[tile_y + 1] |= 1 << tile_x;
-//     } else {
-//         for (size_t i = 0; i < (8 - start_row); i++, backbuffer++) {
-//             BOUNDCHECK_SCREEN(x, y + i);
-//             backbuffer[0] = (backbuffer[0] & ~row_mask_left) | (row[i] * clr << shift_left);
-//             backbuffer[8] = (backbuffer[8] & ~row_mask_right) | (row[i] * clr >> shift_right);
-//         }
-//         backbuffer += 8 * 31;
-//         for (size_t i = (8 - start_row); i < 8; i++, backbuffer++) {
-//             BOUNDCHECK_SCREEN(x, y + i);
-//             backbuffer[0] = (backbuffer[0] & ~row_mask_left) | (row[i] * clr << shift_left);
-//             backbuffer[8] = (backbuffer[8] & ~row_mask_right) | (row[i] * clr >> shift_right);
-//         }
-//         dirty_tiles[tile_y] |= 1 << tile_x;
-//         dirty_tiles[tile_y] |= 1 << (tile_x + 1);
-//         dirty_tiles[tile_y + 1] |= 1 << tile_x;
-//         dirty_tiles[tile_y + 1] |= 1 << (tile_x + 1);
-//     }
-// }
-
-// void
-// clear_screen(void) {
-//     dma_fill(FRONTBUF, 0, KB(20), 3);
-// }
-
-IWRAM_CODE
-void
-flip_buffer(void) {
-    // TODO: Copying all tiles for now. Study if it's better to use dirty_tiles
-    // or dirty_lines.
-    // Copy dirty tiles from the backbuffer to the frontbuffer.
-    Tile *dst = FRONTBUF;
-    Tile *src = BACKBUF;
-    for (size_t j = 0; j < 20; ++j) {
-        // if (dirty_tiles[j] == 0) {
-        //     continue;
-        // }
-        for (size_t i = 0, k = 1; i < 30; ++i, k <<= 1) {
-            // if (dirty_tiles[j] & k) {
-                dst[i + j * 32] = src[i + j * 32];
-            // }
-        }
-        // dirty_tiles[j] = 0;
-    }
-}
-
-static u32 dec_nibble[] = {
-    0x00000000, 0x01000000, 0x00010000, 0x01010000,
-    0x00000100, 0x01000100, 0x00010100, 0x01010100,
-    0x00000001, 0x01000001, 0x00010001, 0x01010001,
-    0x00000101, 0x01000101, 0x00010101, 0x01010101,
-};
-
-static u32 dec_nibble_flip_x[] = {
-    0x00000000, 0x00000001, 0x00000100, 0x00000101,
-    0x00010000, 0x00010001, 0x00010100, 0x00010101,
-    0x01000000, 0x01000001, 0x01000100, 0x01000101,
-    0x01010000, 0x01010001, 0x01010100, 0x01010101,
-};
-
-IWRAM_CODE
-static inline
-u64
-decode_1bpp(u8 row, u8 flip_x) {
-    if (flip_x) {
-        u32 *lut = dec_nibble_flip_x;
-        return (u64)lut[(row >> 4) & 0xF] << 32 | (u64)lut[(row >> 0) & 0xF];
-    }
-    u32 *lut = dec_nibble;
-    return (u64)lut[(row >> 0) & 0xF] << 32 | (u64)lut[(row >> 4) & 0xF];
-}
-
-IWRAM_CODE
-static inline
-void
-draw_2bpp_row(size_t x, size_t y, u8 a, u8 b, u8 flip_x) {
-    // BOUNDCHECK_SCREEN(x, y);
-
-    // size_t tile_x = x / 8;
-    // size_t start_col = x % 8;
-    // size_t shift_left = start_col * 8;
-    // size_t shift_right = (8 - start_col) * 8;
-
-    // u64 *dst = &backbuf[(y * 30 + tile_x) * 8 / 2];
-    // if (start_col == 0) {
-    //     u64 clr_a = decode_1bpp(a, flip_x);
-    //     u64 clr_b = decode_1bpp(b, flip_x);
-    //     u64 mask_a = (clr_a * 0xF);
-    //     u64 mask_b = (clr_b * 0xF);
-    //     u64 mask = (mask_a | mask_b);
-    //     u64 color = clr_a + (clr_b << 1);
-    //     dst[0] = (dst[0] & ~mask) | color;
-    // } else {
-    //     u64 clr_a = decode_1bpp(a, flip_x);
-    //     u64 clr_b = decode_1bpp(b, flip_x);
-    //     u64 mask_a = (clr_a * 0xF);
-    //     u64 mask_b = (clr_b * 0xF);
-    //     u64 mask = (mask_a | mask_b);
-    //     u64 color = clr_a + (clr_b << 1);
-    //     dst[0] = (dst[0] & ~(mask << shift_left))  | (color << shift_left);
-    //     if ((x + 7) > (SCREEN_WIDTH)) {
-    //         return;
-    //     }
-    //     dst[1] = (dst[1] & ~(mask >> shift_right)) | (color >> shift_right);
-    // }
-
-    // TODO: different blend modes?
-}
-
-IWRAM_CODE
-static inline
-void
-draw_1bpp_row(size_t x, size_t y, u8 a, u8 clr, u8 flip_x) {
-    // BOUNDCHECK_SCREEN(x, y);
-
-    // size_t tile_x = x / 8;
-    // size_t start_col = x % 8;
-    // size_t shift_left = start_col * 8;
-    // size_t shift_right = (8 - start_col) * 8;
-
-    // u64 *dst = &backbuf[(y * 30 + tile_x) * 8 / 2];
-    // if (start_col == 0) {
-    //     u64 color = decode_1bpp(a, flip_x);
-    //     u64 mask = (color * 0xF);
-    //     color *= clr;
-    //     dst[0] = (dst[0] & ~mask) | color;
-    // } else {
-    //     u64 color = decode_1bpp(a, flip_x);
-    //     u64 mask = (color * 0xF);
-    //     color *= clr;
-    //     dst[0] = (dst[0] & ~(mask << shift_left))  | (color << shift_left);
-    //     if ((x + 7) > (SCREEN_WIDTH)) {
-    //         return;
-    //     }
-    //     dst[1] = (dst[1] & ~(mask >> shift_right)) | (color >> shift_right);
-    // }
-
-    // TODO: different blend modes?
-}
-
-IWRAM_CODE
-void
-draw_chr(size_t x, size_t y, u8 *sprite, u8 flip_x, u8 flip_y) {
-    // BOUNDCHECK_SCREEN(x, y);
-    // if (!flip_y) {
-    //     for(size_t v = 0; v < 8; v++) {
-    //         if ((y + v) >= SCREEN_HEIGHT) break;
-    //         u8 ch1 = sprite[v + 0];
-    //         u8 ch2 = sprite[v + 8];
-    //         draw_2bpp_row(x, y + v, ch1, ch2, flip_x);
-    //     }
-    // } else {
-    //     for(size_t v = 0; v < 8; v++) {
-    //         if ((y + v) >= SCREEN_HEIGHT) break;
-    //         u8 ch1 = sprite[(7 - v) + 0];
-    //         u8 ch2 = sprite[(7 - v) + 8];
-    //         draw_2bpp_row(x, y + v, ch1, ch2, flip_x);
-    //     }
-    // }
-}
-
-IWRAM_CODE
-void
-draw_icn(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y) {
-    // BOUNDCHECK_SCREEN(x, y);
-    // if (!flip_y) {
-    //     for(size_t v = 0; v < 8; v++) {
-    //         if ((y + v) >= SCREEN_HEIGHT) break;
-    //         u8 ch1 = sprite[v];
-    //         draw_1bpp_row(x, y + v, ch1, clr, flip_x);
-    //     }
-    // } else {
-    //     for(size_t v = 0; v < 8; v++) {
-    //         if ((y + v) >= SCREEN_HEIGHT) break;
-    //         u8 ch1 = sprite[(7 - v)];
-    //         draw_1bpp_row(x, y + v, ch1, clr, flip_x);
-    //     }
-    // }
-}
-
-#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);
-}
-
-void
-renderer_init(void) {
-    // Initialize display mode and bg palette.
-    DISP_CTRL = DISP_MODE_3 | DISP_BG_2;
-
-    // Clear VRAM.
-    dma_fill(MEM_VRAM, 0, KB(96), 3);
-
-    // Initialize display mode and bg palette.
-    DISP_CTRL = DISP_MODE_0 | DISP_BG_0 | DISP_BG_1 | DISP_OBJ;
-
-    // Initialize backgrounds.
-    BG_CTRL(0) = BG_CHARBLOCK(0) | BG_SCREENBLOCK(FRONTBUF_SB) | BG_PRIORITY(1);
-    // BG_CTRL(1) = BG_CHARBLOCK(1) | BG_SCREENBLOCK(FONT_SB) | BG_PRIORITY(0);
-
-    // Initialize background memory map for frontbuffer.
-    for (size_t i = 0; i < 32 * 20; ++i) {
-        FRONTBUF_TILEMAP[i] = i;
-    }
-
-    // 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_PURPLE;
-    PAL_BUFFER_BG[6] = COLOR_YELLOW;
-    PAL_BUFFER_BG[7] = COLOR_GREEN;
-    PAL_BUFFER_BG[8] = COLOR_GREY;
-
-    // Initialize text engine.
-    txt_init(txt_drawc);
-}
diff --git a/src/renderer_m4.c b/src/renderer_m4.c
deleted file mode 100644
index 657bb73..0000000
--- a/src/renderer_m4.c
+++ /dev/null
@@ -1,582 +0,0 @@
-#include "renderer.h"
-#include "text.h"
-
-//
-// Parameters.
-//
-
-#define SUBPIXEL_LINES 0
-#define DEC_BIG_LUT 1
-
-// Front/back buffers for double buffering.
-#define BUF_0 ((u32*)(MEM_VRAM))
-#define BUF_1 ((u32*)(MEM_VRAM ^ 0x0A000))
-
-// Pointer to the backbuffer.
-static u16 *backbuf = BUF_1;
-
-// 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 optimiztions 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.
-//
-
-IWRAM_CODE
-void screen_fill(u8 clr)  {
-    dma_fill(backbuf, 0x01010101 * clr, KB(75) / 2, 3);
-}
-
-IWRAM_CODE
-void
-draw_pixel(size_t x, size_t y, u8 clr) {
-    BOUNDCHECK_SCREEN(x, y);
-    u16 *dst = &backbuf[(x + y * SCREEN_WIDTH) / 2];
-    if(x & 1) {
-        *dst = (*dst &  0xF) | (clr << 8);
-    } else {
-        *dst = (*dst & ~0xF) | clr;
-    }
-}
-
-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 start_col = x0 % 8;
-    size_t end_col = x1 % 8;
-    // 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 dx = tile_x1 - tile_x0;
-    u64 *dst = &backbuf[(tile_x0 * 8 + y0 * SCREEN_WIDTH) / 2];
-    if (dx < 1) {
-        u64 row_mask = 0xFFFFFFFFFFFFFFFF;
-        row_mask >>= (7 - end_col - dx) * 8;
-        row_mask &= 0xFFFFFFFFFFFFFFFF << start_col * 8;
-        u64 row = (0x0101010101010101 * clr) & row_mask;
-        *dst = (*dst & ~row_mask) | row;
-    } else {
-        size_t shift_left = start_col * 8;
-        size_t shift_right = (7 - end_col) * 8;
-        u64 row_mask = 0xFFFFFFFFFFFFFFFF;
-        u64 row = 0x0101010101010101 * clr;
-        dst[0] = (dst[0] & ~(row_mask << shift_left)) | row << shift_left;
-        if (dx != 1) {
-            dma_fill(&dst[1], 0x01010101 * clr, (dx - 1) * 8, 3);
-        }
-        dst[dx] = dst[dx] & ~(row_mask >> shift_right);
-        dst[dx] |= row >> shift_right;
-    }
-}
-
-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_x0 = x0 / 8;
-    size_t start_col = x0 % 8;
-    u16 *dst = &backbuf[(start_col + tile_x0 * 8 + y0 * SCREEN_WIDTH) / 2];
-    if(x0 & 1) {
-        for (size_t i = 0; i <= y1 - y0; i++, dst += SCREEN_WIDTH / 2) {
-            *dst = (*dst &  0xF) | (clr << 8);
-        }
-    } else {
-        for (size_t i = 0; i <= y1 - y0; i++, dst += SCREEN_WIDTH / 2) {
-            *dst = (*dst & ~0xF) | clr;
-        }
-    }
-}
-
-IWRAM_CODE
-UNROLL_LOOPS
-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 {
-#if SUBPIXEL_LINES == 1
-        // 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;
-        int dxf = (dx << fp_bit);
-        int dyf = (dy << fp_bit);
-
-        if ((dx >= dy && x0 > x1) || (dx < dy && y0 > y1)) {
-            SWAP(x0, x1);
-            SWAP(y0, y1);
-        }
-
-        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 ? -SCREEN_WIDTH : SCREEN_WIDTH;
-
-        u16 *dst = NULL;
-        uintptr_t addr = ((uintptr_t)backbuf + y0 * SCREEN_WIDTH + x0);
-        u32 mask = x0 & 1 ? ~0xFF : 0xFF;
-        u32 color = (clr & 0xFF) | ((clr & 0xFF) << 8);
-        if (dx >= dy) {
-            int distance = (frac_y - fp_one) * dx - (frac_x - fp_half) * dy;
-            int remaining = dx;
-            while (distance <= 0  && remaining > 0) {
-                dst = (u16*)(addr - (mask >> 31));
-                *dst = (*dst & ~mask) | (color & mask);
-                distance += 2 * dyf;
-                addr += x_step;
-                remaining--;
-                mask = ~mask;
-            }
-            distance -= 2 * dxf;
-            addr += y_step;
-
-            while (remaining >= 0) {
-                dst = (u16*)(addr - (mask >> 31));
-                *dst = (*dst & ~mask) | (color & mask);
-                if (distance >= 0) {
-                    distance -= 2 * dxf;
-                    addr += y_step;
-                }
-                distance += 2 * dyf;
-                addr += x_step;
-                mask = ~mask;
-                remaining--;
-            }
-        } else {
-            int distance = (frac_x - fp_one) * dy - (frac_y - fp_half) * dx;
-            int remaining = dy;
-            while (distance <= 0  && remaining > 0) {
-                dst = (u16*)(addr - (mask >> 31));
-                *dst = (*dst & ~mask) | (color & mask);
-                distance += 2 * dxf;
-                addr += y_step;
-                remaining--;
-            }
-            distance -= 2 * dyf;
-            addr += x_step;
-            mask = ~mask;
-
-            while (remaining >= 0) {
-                dst = (u16*)(addr - (mask >> 31));
-                *dst = (*dst & ~mask) | (color & mask);
-                if (distance >= 0) {
-                    distance -= 2 * dyf;
-                    addr += x_step;
-                    mask = ~mask;
-                }
-                distance += 2 * dxf;
-                addr += y_step;
-                remaining--;
-            }
-        }
-#else
-        // Diagonal line.
-        int     dx = x0 > x1 ?       x0 - x1 : x1 - x0;
-        int     dy = y0 > y1 ?       y0 - y1 : y1 - y0;
-        int x_step = x0 > x1 ?            -1 : 1;
-        int y_step = y0 > y1 ? -SCREEN_WIDTH : SCREEN_WIDTH;
-
-        u16 *dst = NULL;
-        uintptr_t addr = ((uintptr_t)backbuf + y0 * SCREEN_WIDTH + x0);
-        u32 mask = x0 & 1 ? ~0xFF : 0xFF;
-        u32 color = (clr & 0xFF) | ((clr & 0xFF) << 8);
-        if (dx >= dy) {
-            int diff = 2 * dy - dx;
-            for (int i = 0; i < dx + 1; i++) {
-                dst = (u16*)(addr - (mask >> 31));
-                *dst = (*dst & ~mask) | (color & mask);
-                if (diff >= 0) {
-                    diff -= 2 * dx;
-                    addr += y_step;
-                }
-                diff += 2 * dy;
-                addr += x_step;
-                mask = ~mask;
-            }
-        } else {
-            int diff = 2 * dx - dy;
-            for (int i = 0; i < dy + 1; i++) {
-                dst = (u16*)(addr - (mask >> 31));
-                *dst = (*dst & ~mask) | (color & mask);
-                if (diff >= 0) {
-                    diff -= 2 * dy;
-                    addr += x_step;
-                    mask = ~mask;
-                }
-                diff += 2 * dx;
-                addr += 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;
-    }
-
-    // Drawline implementation.
-    for (size_t y = y0; y <= y1; y++) {
-        draw_hline(x0, x1, y, clr);
-    }
-}
-
-//
-// Sprites (chr/icn).
-//
-
-#if DEC_BIG_LUT == 1
-
-static u64 dec_byte_flip_x[256] = {
-    0x0000000000000000, 0x0000000000000001, 0x0000000000000100, 0x0000000000000101, 0x0000000000010000,
-    0x0000000000010001, 0x0000000000010100, 0x0000000000010101, 0x0000000001000000, 0x0000000001000001,
-    0x0000000001000100, 0x0000000001000101, 0x0000000001010000, 0x0000000001010001, 0x0000000001010100,
-    0x0000000001010101, 0x0000000100000000, 0x0000000100000001, 0x0000000100000100, 0x0000000100000101,
-    0x0000000100010000, 0x0000000100010001, 0x0000000100010100, 0x0000000100010101, 0x0000000101000000,
-    0x0000000101000001, 0x0000000101000100, 0x0000000101000101, 0x0000000101010000, 0x0000000101010001,
-    0x0000000101010100, 0x0000000101010101, 0x0000010000000000, 0x0000010000000001, 0x0000010000000100,
-    0x0000010000000101, 0x0000010000010000, 0x0000010000010001, 0x0000010000010100, 0x0000010000010101,
-    0x0000010001000000, 0x0000010001000001, 0x0000010001000100, 0x0000010001000101, 0x0000010001010000,
-    0x0000010001010001, 0x0000010001010100, 0x0000010001010101, 0x0000010100000000, 0x0000010100000001,
-    0x0000010100000100, 0x0000010100000101, 0x0000010100010000, 0x0000010100010001, 0x0000010100010100,
-    0x0000010100010101, 0x0000010101000000, 0x0000010101000001, 0x0000010101000100, 0x0000010101000101,
-    0x0000010101010000, 0x0000010101010001, 0x0000010101010100, 0x0000010101010101, 0x0001000000000000,
-    0x0001000000000001, 0x0001000000000100, 0x0001000000000101, 0x0001000000010000, 0x0001000000010001,
-    0x0001000000010100, 0x0001000000010101, 0x0001000001000000, 0x0001000001000001, 0x0001000001000100,
-    0x0001000001000101, 0x0001000001010000, 0x0001000001010001, 0x0001000001010100, 0x0001000001010101,
-    0x0001000100000000, 0x0001000100000001, 0x0001000100000100, 0x0001000100000101, 0x0001000100010000,
-    0x0001000100010001, 0x0001000100010100, 0x0001000100010101, 0x0001000101000000, 0x0001000101000001,
-    0x0001000101000100, 0x0001000101000101, 0x0001000101010000, 0x0001000101010001, 0x0001000101010100,
-    0x0001000101010101, 0x0001010000000000, 0x0001010000000001, 0x0001010000000100, 0x0001010000000101,
-    0x0001010000010000, 0x0001010000010001, 0x0001010000010100, 0x0001010000010101, 0x0001010001000000,
-    0x0001010001000001, 0x0001010001000100, 0x0001010001000101, 0x0001010001010000, 0x0001010001010001,
-    0x0001010001010100, 0x0001010001010101, 0x0001010100000000, 0x0001010100000001, 0x0001010100000100,
-    0x0001010100000101, 0x0001010100010000, 0x0001010100010001, 0x0001010100010100, 0x0001010100010101,
-    0x0001010101000000, 0x0001010101000001, 0x0001010101000100, 0x0001010101000101, 0x0001010101010000,
-    0x0001010101010001, 0x0001010101010100, 0x0001010101010101, 0x0100000000000000, 0x0100000000000001,
-    0x0100000000000100, 0x0100000000000101, 0x0100000000010000, 0x0100000000010001, 0x0100000000010100,
-    0x0100000000010101, 0x0100000001000000, 0x0100000001000001, 0x0100000001000100, 0x0100000001000101,
-    0x0100000001010000, 0x0100000001010001, 0x0100000001010100, 0x0100000001010101, 0x0100000100000000,
-    0x0100000100000001, 0x0100000100000100, 0x0100000100000101, 0x0100000100010000, 0x0100000100010001,
-    0x0100000100010100, 0x0100000100010101, 0x0100000101000000, 0x0100000101000001, 0x0100000101000100,
-    0x0100000101000101, 0x0100000101010000, 0x0100000101010001, 0x0100000101010100, 0x0100000101010101,
-    0x0100010000000000, 0x0100010000000001, 0x0100010000000100, 0x0100010000000101, 0x0100010000010000,
-    0x0100010000010001, 0x0100010000010100, 0x0100010000010101, 0x0100010001000000, 0x0100010001000001,
-    0x0100010001000100, 0x0100010001000101, 0x0100010001010000, 0x0100010001010001, 0x0100010001010100,
-    0x0100010001010101, 0x0100010100000000, 0x0100010100000001, 0x0100010100000100, 0x0100010100000101,
-    0x0100010100010000, 0x0100010100010001, 0x0100010100010100, 0x0100010100010101, 0x0100010101000000,
-    0x0100010101000001, 0x0100010101000100, 0x0100010101000101, 0x0100010101010000, 0x0100010101010001,
-    0x0100010101010100, 0x0100010101010101, 0x0101000000000000, 0x0101000000000001, 0x0101000000000100,
-    0x0101000000000101, 0x0101000000010000, 0x0101000000010001, 0x0101000000010100, 0x0101000000010101,
-    0x0101000001000000, 0x0101000001000001, 0x0101000001000100, 0x0101000001000101, 0x0101000001010000,
-    0x0101000001010001, 0x0101000001010100, 0x0101000001010101, 0x0101000100000000, 0x0101000100000001,
-    0x0101000100000100, 0x0101000100000101, 0x0101000100010000, 0x0101000100010001, 0x0101000100010100,
-    0x0101000100010101, 0x0101000101000000, 0x0101000101000001, 0x0101000101000100, 0x0101000101000101,
-    0x0101000101010000, 0x0101000101010001, 0x0101000101010100, 0x0101000101010101, 0x0101010000000000,
-    0x0101010000000001, 0x0101010000000100, 0x0101010000000101, 0x0101010000010000, 0x0101010000010001,
-    0x0101010000010100, 0x0101010000010101, 0x0101010001000000, 0x0101010001000001, 0x0101010001000100,
-    0x0101010001000101, 0x0101010001010000, 0x0101010001010001, 0x0101010001010100, 0x0101010001010101,
-    0x0101010100000000, 0x0101010100000001, 0x0101010100000100, 0x0101010100000101, 0x0101010100010000,
-    0x0101010100010001, 0x0101010100010100, 0x0101010100010101, 0x0101010101000000, 0x0101010101000001,
-    0x0101010101000100, 0x0101010101000101, 0x0101010101010000, 0x0101010101010001, 0x0101010101010100,
-    0x0101010101010101
-};
-
-static u64 dec_byte[256] = {
-    0x0000000000000000, 0x0100000000000000, 0x0001000000000000, 0x0101000000000000, 0x0000010000000000,
-    0x0100010000000000, 0x0001010000000000, 0x0101010000000000, 0x0000000100000000, 0x0100000100000000,
-    0x0001000100000000, 0x0101000100000000, 0x0000010100000000, 0x0100010100000000, 0x0001010100000000,
-    0x0101010100000000, 0x0000000001000000, 0x0100000001000000, 0x0001000001000000, 0x0101000001000000,
-    0x0000010001000000, 0x0100010001000000, 0x0001010001000000, 0x0101010001000000, 0x0000000101000000,
-    0x0100000101000000, 0x0001000101000000, 0x0101000101000000, 0x0000010101000000, 0x0100010101000000,
-    0x0001010101000000, 0x0101010101000000, 0x0000000000010000, 0x0100000000010000, 0x0001000000010000,
-    0x0101000000010000, 0x0000010000010000, 0x0100010000010000, 0x0001010000010000, 0x0101010000010000,
-    0x0000000100010000, 0x0100000100010000, 0x0001000100010000, 0x0101000100010000, 0x0000010100010000,
-    0x0100010100010000, 0x0001010100010000, 0x0101010100010000, 0x0000000001010000, 0x0100000001010000,
-    0x0001000001010000, 0x0101000001010000, 0x0000010001010000, 0x0100010001010000, 0x0001010001010000,
-    0x0101010001010000, 0x0000000101010000, 0x0100000101010000, 0x0001000101010000, 0x0101000101010000,
-    0x0000010101010000, 0x0100010101010000, 0x0001010101010000, 0x0101010101010000, 0x0000000000000100,
-    0x0100000000000100, 0x0001000000000100, 0x0101000000000100, 0x0000010000000100, 0x0100010000000100,
-    0x0001010000000100, 0x0101010000000100, 0x0000000100000100, 0x0100000100000100, 0x0001000100000100,
-    0x0101000100000100, 0x0000010100000100, 0x0100010100000100, 0x0001010100000100, 0x0101010100000100,
-    0x0000000001000100, 0x0100000001000100, 0x0001000001000100, 0x0101000001000100, 0x0000010001000100,
-    0x0100010001000100, 0x0001010001000100, 0x0101010001000100, 0x0000000101000100, 0x0100000101000100,
-    0x0001000101000100, 0x0101000101000100, 0x0000010101000100, 0x0100010101000100, 0x0001010101000100,
-    0x0101010101000100, 0x0000000000010100, 0x0100000000010100, 0x0001000000010100, 0x0101000000010100,
-    0x0000010000010100, 0x0100010000010100, 0x0001010000010100, 0x0101010000010100, 0x0000000100010100,
-    0x0100000100010100, 0x0001000100010100, 0x0101000100010100, 0x0000010100010100, 0x0100010100010100,
-    0x0001010100010100, 0x0101010100010100, 0x0000000001010100, 0x0100000001010100, 0x0001000001010100,
-    0x0101000001010100, 0x0000010001010100, 0x0100010001010100, 0x0001010001010100, 0x0101010001010100,
-    0x0000000101010100, 0x0100000101010100, 0x0001000101010100, 0x0101000101010100, 0x0000010101010100,
-    0x0100010101010100, 0x0001010101010100, 0x0101010101010100, 0x0000000000000001, 0x0100000000000001,
-    0x0001000000000001, 0x0101000000000001, 0x0000010000000001, 0x0100010000000001, 0x0001010000000001,
-    0x0101010000000001, 0x0000000100000001, 0x0100000100000001, 0x0001000100000001, 0x0101000100000001,
-    0x0000010100000001, 0x0100010100000001, 0x0001010100000001, 0x0101010100000001, 0x0000000001000001,
-    0x0100000001000001, 0x0001000001000001, 0x0101000001000001, 0x0000010001000001, 0x0100010001000001,
-    0x0001010001000001, 0x0101010001000001, 0x0000000101000001, 0x0100000101000001, 0x0001000101000001,
-    0x0101000101000001, 0x0000010101000001, 0x0100010101000001, 0x0001010101000001, 0x0101010101000001,
-    0x0000000000010001, 0x0100000000010001, 0x0001000000010001, 0x0101000000010001, 0x0000010000010001,
-    0x0100010000010001, 0x0001010000010001, 0x0101010000010001, 0x0000000100010001, 0x0100000100010001,
-    0x0001000100010001, 0x0101000100010001, 0x0000010100010001, 0x0100010100010001, 0x0001010100010001,
-    0x0101010100010001, 0x0000000001010001, 0x0100000001010001, 0x0001000001010001, 0x0101000001010001,
-    0x0000010001010001, 0x0100010001010001, 0x0001010001010001, 0x0101010001010001, 0x0000000101010001,
-    0x0100000101010001, 0x0001000101010001, 0x0101000101010001, 0x0000010101010001, 0x0100010101010001,
-    0x0001010101010001, 0x0101010101010001, 0x0000000000000101, 0x0100000000000101, 0x0001000000000101,
-    0x0101000000000101, 0x0000010000000101, 0x0100010000000101, 0x0001010000000101, 0x0101010000000101,
-    0x0000000100000101, 0x0100000100000101, 0x0001000100000101, 0x0101000100000101, 0x0000010100000101,
-    0x0100010100000101, 0x0001010100000101, 0x0101010100000101, 0x0000000001000101, 0x0100000001000101,
-    0x0001000001000101, 0x0101000001000101, 0x0000010001000101, 0x0100010001000101, 0x0001010001000101,
-    0x0101010001000101, 0x0000000101000101, 0x0100000101000101, 0x0001000101000101, 0x0101000101000101,
-    0x0000010101000101, 0x0100010101000101, 0x0001010101000101, 0x0101010101000101, 0x0000000000010101,
-    0x0100000000010101, 0x0001000000010101, 0x0101000000010101, 0x0000010000010101, 0x0100010000010101,
-    0x0001010000010101, 0x0101010000010101, 0x0000000100010101, 0x0100000100010101, 0x0001000100010101,
-    0x0101000100010101, 0x0000010100010101, 0x0100010100010101, 0x0001010100010101, 0x0101010100010101,
-    0x0000000001010101, 0x0100000001010101, 0x0001000001010101, 0x0101000001010101, 0x0000010001010101,
-    0x0100010001010101, 0x0001010001010101, 0x0101010001010101, 0x0000000101010101, 0x0100000101010101,
-    0x0001000101010101, 0x0101000101010101, 0x0000010101010101, 0x0100010101010101, 0x0001010101010101,
-    0x0101010101010101
-};
-
-IWRAM_CODE
-static inline
-u64
-decode_1bpp(u8 row, u8 flip_x) {
-    if (flip_x) {
-        return dec_byte_flip_x[row];
-    }
-    return dec_byte[row];
-}
-
-#else
-
-static u32 dec_nibble[] = {
-    0x00000000, 0x01000000, 0x00010000, 0x01010000,
-    0x00000100, 0x01000100, 0x00010100, 0x01010100,
-    0x00000001, 0x01000001, 0x00010001, 0x01010001,
-    0x00000101, 0x01000101, 0x00010101, 0x01010101,
-};
-
-static u32 dec_nibble_flip_x[] = {
-    0x00000000, 0x00000001, 0x00000100, 0x00000101,
-    0x00010000, 0x00010001, 0x00010100, 0x00010101,
-    0x01000000, 0x01000001, 0x01000100, 0x01000101,
-    0x01010000, 0x01010001, 0x01010100, 0x01010101,
-};
-
-IWRAM_CODE
-static inline
-u64
-decode_1bpp(u8 row, u8 flip_x) {
-    if (flip_x) {
-        u32 *lut = dec_nibble_flip_x;
-        return (u64)lut[(row >> 4) & 0xF] << 32 | (u64)lut[(row >> 0) & 0xF];
-    }
-    u32 *lut = dec_nibble;
-    return (u64)lut[(row >> 0) & 0xF] << 32 | (u64)lut[(row >> 4) & 0xF];
-}
-
-#endif
-
-IWRAM_CODE
-static inline
-void
-draw_2bpp_row(size_t x, size_t y, u8 a, u8 b, u8 clr, u8 flip_x) {
-    BOUNDCHECK_SCREEN(x, y);
-
-    size_t tile_x = x / 8;
-    size_t start_col = x % 8;
-    size_t shift_left = start_col * 8;
-    size_t shift_right = (8 - start_col) * 8;
-
-    u64 *dst = &backbuf[(y * 30 + tile_x) * 8 / 2];
-#if DEC_BIG_LUT
-    u64 *lut = flip_x ? dec_byte_flip_x : dec_byte;
-    u64 clr_a = lut[a];
-    u64 clr_b = lut[b];
-#else
-    u64 clr_a = decode_1bpp(a, flip_x);
-    u64 clr_b = decode_1bpp(b, flip_x);
-#endif
-    u64 mask_a = (clr_a * 0xF);
-    u64 mask_b = (clr_b * 0xF);
-    u64 mask = (mask_a | mask_b);
-    u64 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[1] = (dst[1] & ~(mask >> shift_right)) | (color >> shift_right);
-}
-
-IWRAM_CODE
-static inline
-void
-draw_1bpp_row(size_t x, size_t y, u8 a, u8 clr, u8 flip_x) {
-    BOUNDCHECK_SCREEN(x, y);
-
-    size_t tile_x = x / 8;
-    size_t start_col = x % 8;
-    size_t shift_left = start_col * 8;
-    size_t shift_right = (8 - start_col) * 8;
-
-    u64 *dst = &backbuf[(y * 30 + tile_x) * 8 / 2];
-    u64 color = decode_1bpp(a, flip_x);
-    u64 mask = (color * 0xF);
-    color *= clr;
-    dst[0] = (dst[0] & ~(mask << shift_left))  | (color << shift_left);
-    dst[1] = (dst[1] & ~(mask >> shift_right)) | (color >> shift_right);
-}
-
-IWRAM_CODE
-void
-draw_chr(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y) {
-    BOUNDCHECK_SCREEN(x, y);
-    if (!flip_y) {
-        for(size_t v = 0; v < 8; v++) {
-            if ((y + v) >= SCREEN_HEIGHT) break;
-            u8 ch1 = sprite[v + 0];
-            u8 ch2 = sprite[v + 8];
-            draw_2bpp_row(x, y + v, ch1, ch2, clr, flip_x);
-        }
-    } else {
-        for(size_t v = 0; v < 8; v++) {
-            if ((y + v) >= SCREEN_HEIGHT) break;
-            u8 ch1 = sprite[(7 - v) + 0];
-            u8 ch2 = sprite[(7 - v) + 8];
-            draw_2bpp_row(x, y + v, ch1, ch2, clr, flip_x);
-        }
-    }
-}
-
-IWRAM_CODE
-void
-draw_icn(size_t x, size_t y, u8 *sprite, u8 clr, u8 flip_x, u8 flip_y) {
-    BOUNDCHECK_SCREEN(x, y);
-    if (!flip_y) {
-        for(size_t v = 0; v < 8; v++) {
-            if ((y + v) >= SCREEN_HEIGHT) break;
-            u8 ch1 = sprite[v];
-            draw_1bpp_row(x, y + v, ch1, clr, flip_x);
-        }
-    } else {
-        for(size_t v = 0; v < 8; v++) {
-            if ((y + v) >= SCREEN_HEIGHT) break;
-            u8 ch1 = sprite[(7 - v)];
-            draw_1bpp_row(x, y + v, ch1, clr, flip_x);
-        }
-    }
-}
-
-//
-// Flipping buffers/copying memory.
-//
-
-IWRAM_CODE
-void
-flip_buffer(void) {
-    backbuf = (u16*)((u32)backbuf ^ 0x0A000);
-    DISP_CTRL ^= DISP_PAGE;
-}
-
-//
-// 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_4 | DISP_BG_2;
-
-    // Clear VRAM.
-    dma_fill((u16*)MEM_VRAM, 0, KB(96), 3);
-
-    // 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_PURPLE;
-    PAL_BUFFER_BG[6] = COLOR_YELLOW;
-    PAL_BUFFER_BG[7] = COLOR_GREEN;
-    PAL_BUFFER_BG[8] = COLOR_GREY;
-
-    // Initialize text engine.
-    txt_init(txt_drawc);
-}
diff --git a/src/text/font.h b/src/text/font.h
deleted file mode 100644
index e4b089c..0000000
--- a/src/text/font.h
+++ /dev/null
@@ -1,261 +0,0 @@
-static const u32 font[] = {
-    0x00000000, 0x00000000, 0x00002400, 0x423c0000,
-    0x00002400, 0x3c420000, 0x0000363e, 0x3e1c0800,
-    0x00081c3e, 0x3e1c0800, 0x001c1c3e, 0x363e081c,
-    0x00081c3e, 0x3e3e081c, 0x00000018, 0x18000000,
-    0x7e7e7e66, 0x667e7e7e, 0x00001824, 0x24180000,
-    0x7e7e665a, 0x5a667e7e, 0x00081c3e, 0x081c221c,
-    0x001c221c, 0x08083e08, 0x00183828, 0x08080c0c,
-    0x003c2424, 0x24343606, 0x00082208, 0x1c082208,
-    0x040c1c3c, 0x1c0c0400, 0x2030383c, 0x38302000,
-    0x081c3e08, 0x083e1c08, 0x00141414, 0x14001400,
-    0x003c2a2a, 0x2c282828, 0x0038043c, 0x423c201e,
-    0x00000000, 0x7e000000, 0x081c3e08, 0x3e1c083e,
-    0x00081c3e, 0x08080800, 0x00080808, 0x3e1c0800,
-    0x00001030, 0x7e301000, 0x0000080c, 0x7e0c0800,
-    0x00000000, 0x0002023e, 0x00001436, 0x7f361400,
-    0x0008081c, 0x1c3e3e00, 0x003e3e1c, 0x1c080800,
-    0x00000000, 0x00000000, 0x00080808, 0x08000800,
-    0x00141414, 0x00000000, 0x0000143e, 0x143e1400,
-    0x00081c02, 0x1c201e08, 0x00002616, 0x08343200,
-    0x00081408, 0x34122c00, 0x00080808, 0x00000000,
-    0x00000804, 0x04040800, 0x00000810, 0x10100800,
-    0x00001408, 0x3e081400, 0x00000808, 0x3e080800,
-    0x00000000, 0x000c0804, 0x00000000, 0x3e000000,
-    0x00000000, 0x000c0c00, 0x00203018, 0x0c060200,
-    0x00001c22, 0x2a221c00, 0x00000c0a, 0x08083e00,
-    0x00001c20, 0x1c023e00, 0x00001c22, 0x18221c00,
-    0x00001212, 0x123e1000, 0x00003e02, 0x1e201e00,
-    0x00001c02, 0x1e221c00, 0x00003e20, 0x10080400,
-    0x00001c22, 0x1c221c00, 0x00001c22, 0x3c201c00,
-    0x00000c0c, 0x000c0c00, 0x00000c0c, 0x000c0804,
-    0x00001008, 0x04081000, 0x0000003e, 0x003e0000,
-    0x00000408, 0x10080400, 0x001c2210, 0x08000800,
-    0x00001c2a, 0x3a021c00, 0x00001c22, 0x223e2200,
-    0x00001e22, 0x1e221e00, 0x00001c22, 0x02221c00,
-    0x00001e22, 0x22221e00, 0x00003e02, 0x1e023e00,
-    0x00003e02, 0x1e020200, 0x00001c02, 0x32221c00,
-    0x00002222, 0x3e222200, 0x00003e08, 0x08083e00,
-    0x00002020, 0x20221c00, 0x00002212, 0x0e122200,
-    0x00000202, 0x02023e00, 0x00002236, 0x2a222200,
-    0x0000262a, 0x32222200, 0x00001c22, 0x22221c00,
-    0x00001e22, 0x221e0200, 0x00001c22, 0x22122c00,
-    0x00001e22, 0x221e2200, 0x00001c02, 0x1c201e00,
-    0x00003e08, 0x08080800, 0x00002222, 0x22221c00,
-    0x00002222, 0x22140800, 0x0000222a, 0x2a2a1400,
-    0x00002214, 0x08142200, 0x00002222, 0x14080800,
-    0x00003e10, 0x08043e00, 0x00001c04, 0x04041c00,
-    0x0002060c, 0x18302000, 0x00001c10, 0x10101c00,
-    0x00081422, 0x00000000, 0x00000000, 0x00003e00,
-    0x00040810, 0x00000000, 0x00001c20, 0x3c223c00,
-    0x0002021e, 0x22221e00, 0x00001c02, 0x02021c00,
-    0x0020203c, 0x22223c00, 0x00001c22, 0x1e021c00,
-    0x00003c02, 0x1e020200, 0x00003c22, 0x223c201c,
-    0x0002021e, 0x22222200, 0x0008000c, 0x08083e00,
-    0x00200020, 0x2020221c, 0x0002120a, 0x060a3200,
-    0x000c0808, 0x08083e00, 0x0000162a, 0x2a2a2a00,
-    0x00000e32, 0x22222200, 0x00001c22, 0x22221c00,
-    0x00001e22, 0x221e0202, 0x00003c22, 0x223c2070,
-    0x00001a26, 0x02020200, 0x00001c02, 0x1c201e00,
-    0x00043e04, 0x04041800, 0x00002222, 0x22221c00,
-    0x00002222, 0x22140800, 0x00002222, 0x2a2a1400,
-    0x00002214, 0x08142200, 0x00002222, 0x223c201c,
-    0x00003e10, 0x08043e00, 0x00001c04, 0x06041c00,
-    0x00000808, 0x08080800, 0x00001c10, 0x30101c00,
-    0x00002c1a, 0x00000000, 0x00000814, 0x22223e00,
-    0x00000000, 0x00000000, 0x00000808, 0x08000800,
-    0x00001414, 0x00000000, 0x0000143e, 0x143e1400,
-    0x00081c04, 0x1c101c08, 0x00000410, 0x08041000,
-    0x00000814, 0x38143800, 0x00000808, 0x00000000,
-    0x00000804, 0x04040800, 0x00000810, 0x10100800,
-    0x00001408, 0x1c081400, 0x00000008, 0x1c080000,
-    0x00000000, 0x00000804, 0x00000000, 0x1c000000,
-    0x00000000, 0x00000800, 0x00001018, 0x080c0400,
-    0x00001c14, 0x14141c00, 0x0000080c, 0x08081c00,
-    0x00001c10, 0x1c041c00, 0x00001c10, 0x18101c00,
-    0x00001414, 0x1c101000, 0x00001c04, 0x1c101c00,
-    0x00001c04, 0x1c141c00, 0x00001c10, 0x08040400,
-    0x00001c14, 0x1c141c00, 0x00001c14, 0x1c101000,
-    0x00000008, 0x00000800, 0x00000008, 0x00000804,
-    0x00001008, 0x04081000, 0x0000001c, 0x001c0000,
-    0x00000408, 0x10080400, 0x00001c10, 0x08000800,
-    0x00001c14, 0x14041c00, 0x00001c14, 0x141c1400,
-    0x00001c14, 0x0c141c00, 0x00001c04, 0x04041c00,
-    0x00000c14, 0x14140c00, 0x00001c04, 0x0c041c00,
-    0x00001c04, 0x0c040400, 0x00001c04, 0x14141c00,
-    0x00001414, 0x1c141400, 0x00001c08, 0x08081c00,
-    0x00001010, 0x10141c00, 0x00001414, 0x0c141400,
-    0x00000404, 0x04041c00, 0x0000141c, 0x14141400,
-    0x00000c14, 0x14141400, 0x00000c14, 0x14141c00,
-    0x00001c14, 0x1c040400, 0x00001c14, 0x141c1000,
-    0x00001c14, 0x0c141400, 0x00001804, 0x1c100c00,
-    0x00001c08, 0x08080800, 0x00001414, 0x14141c00,
-    0x00001414, 0x14141800, 0x00001414, 0x141c1400,
-    0x00001414, 0x08141400, 0x00001414, 0x1c080800,
-    0x00001c10, 0x08041c00, 0x00001c04, 0x04041c00,
-    0x0000040c, 0x08181000, 0x00001c10, 0x10101c00,
-    0x00000814, 0x00000000, 0x00000000, 0x00001c00,
-    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,
-    0x3e323232, 0x32323e00, 0x18181818, 0x18181800,
-    0x3e30303e, 0x06063e00, 0x3e30303c, 0x30303e00,
-    0x32323232, 0x3e303000, 0x3e06063e, 0x30303e00,
-    0x3e06063e, 0x26263e00, 0x3e303018, 0x0c060600,
-    0x3e32323e, 0x32323e00, 0x3e32323e, 0x30303e00,
-    0x00000000, 0x00060600, 0x00263618, 0x0c363200,
-    0x1c3e3232, 0x3e323200, 0x1e26261e, 0x26261e00,
-    0x3c060606, 0x06063c00, 0x1e323232, 0x32321e00,
-    0x3e06061e, 0x06063e00, 0x3e06061e, 0x06060600,
-    0x3c060636, 0x26261c00, 0x3232323e, 0x32323200,
-    0x1e0c0c0c, 0x0c0c1e00, 0x30303030, 0x32323e00,
-    0x2626160e, 0x16262600, 0x06060606, 0x06063e00,
-    0x62767e6a, 0x62626200, 0x363e3a32, 0x32323200,
-    0x1c323232, 0x32321c00, 0x3e32323e, 0x02020200,
-    0x1c323232, 0x3a1a2c00, 0x1e26261e, 0x26262600,
-    0x3c06061e, 0x30301e00, 0x7e181818, 0x18181800,
-    0x32323232, 0x32323e00, 0x32323232, 0x32341800,
-    0x32323232, 0x323a1400, 0x3232320c, 0x32323200,
-    0x32323232, 0x3c201e00, 0x3e30301c, 0x06063e00,
-    0x00000000, 0x00000000, 0x00000000, 0x00000000,
-    0x00000000, 0x00000000, 0x00000000, 0x00000000,
-    0x00000000, 0x00000000, 0x00000000, 0x00000000,
-    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,
-};