Provide an example on multiscreen scrolling

Signed-off-by: Miquel Sabaté Solà <mikisabate@gmail.com>

12 files changed, 1702 insertions, 99 deletions

diff --git a/scroll/include/all.s b/scroll/include/all.s
new file mode 100644
index 0000000..e9f059d
--- /dev/null
+++ b/scroll/include/all.s
@@ -0,0 +1,18 @@
+;;; Instead of figuring out the includes for all the examples, simply include
+;;; this one and run with it.
+
+.include "apu.s"
+.include "oam.s"
+.include "ppu.s"
+.include "../../shared/asm.s"
+.include "../../shared/joypad.s"
+
+.include "reset.s"
+.include "palettes.s"
+.include "metatile.s"
+.include "globals.s"
+.include "collision.s"
+.include "buffer.s"
+.include "background.s"
+.include "player.s"
+.include "driver.s"
diff --git a/scroll/include/background.s b/scroll/include/background.s
new file mode 100644
index 0000000..f392792
--- /dev/null
+++ b/scroll/include/background.s
@@ -0,0 +1,498 @@
+;; Reset the background cycle so a new column can be loaded even if the cycle
+;; was stopped at the end of the loading of the previous screen.
+.macro RESET_BACKGROUND_CYCLE
+    lda #$FF
+    sta Background::zp_background_cycle
+.endmacro
+
+;; All functions and related variables for keeping up with the background
+;; elements. This is by far the most extensive and complex piece of code from
+;; the scrolling examples, and it's the core of the render engine for them.
+;;
+;; Documentation has been added for it as much as possible, so follow along to
+;; better grasp how to use and grok this piece of code.
+.scope Background
+    ;; Current value for the scroll on the X axis.
+    zp_scroll = $95
+
+    ;; Instead of attempting to load a column every time the player movement
+    ;; requires it, stall things a bit (as much as defined in
+    ;; `BACKGROUND_CYCLE_MAX`). This allows for loading columns at a slower
+    ;; pace. This cycle is diregarded if you set it to
+    ;; `DISABLED_BACGROUND_CYCLE`, which it just means something like "no more
+    ;; columns to load".
+    zp_background_cycle = $96
+
+    ;; See `zp_background_cycle`.
+    DISABLED_BACKGROUND_CYCLE = $FE
+    BACKGROUND_CYCLE_MAX = $02
+
+    ;; Current column (or the next to be loaded by functions like
+    ;; `load_column`).
+    zp_cur_column = $75
+
+    ;; Current row (or the next to be loaded by functions like `load_column`).
+    zp_cur_row = $76
+
+    .proc init
+        ;; Load the first level.
+        ldx #$00
+        JAL load_level_x
+        ;; NOTE: returned value ignored as the level *must* have loaded at this
+        ;; point.
+    .endproc
+
+    ;; Resets all the variables which touch in any way the loading of columns,
+    ;; levels, etc. This is already being called by `load_level_x`, so outside
+    ;; of this usage you shouldn't be calling it.
+    .proc reset
+        RESET_BACKGROUND_CYCLE
+
+        lda #$00
+        sta zp_scroll
+        sta Buffer::zp_vram_idx
+        sta zp_cur_row
+        sta zp_cur_column
+
+        rts
+    .endproc
+
+    ;; Load the level indexed by the `x` register. This function requires that
+    ;; the PPU is disabled. If this is not the case, then it sets PPU::zp_mask
+    ;; in a way so to disable the PPU on the next `nmi` call, and finally sets
+    ;; `a` to 0 (i.e. level could not be loaded). Otherwise you can expect the
+    ;; `a` register to contain 1 after calling this function.
+    ;;
+    ;; Moreover, the screen pointer is going to be left at the start of the
+    ;; second screen, so future `load_column` calls can rely on this pointer
+    ;; being set properly.
+    .proc load_level_x
+        ;; Is the PPU set to be disabled?
+        lda PPU::zp_mask
+        and #%00111100
+        beq @load_level
+
+        ;; The PPU is not disabled, set the mask so the PPU is disabled.
+        lda #$00
+        sta PPU::zp_mask
+
+        ;; Reset the nametable being used.
+        lda #%11111100
+        and PPU::zp_control
+        sta PPU::zp_control
+
+        ;; Enable the `ppu` flag so these changes are not ignored.
+        lda Globals::zp_flags
+        ora #%01000000
+        and #%11011111
+        sta Globals::zp_flags
+        lda #0
+        rts
+
+    @load_level:
+        ;; Start by resetting any variables that touch on the loading of columns
+        ;; in any way.
+        jsr Background::reset
+
+        ;; Reset all relevant flags.
+        lda #%11000011
+        and Globals::zp_flags
+        sta Globals::zp_flags
+
+        ;; Initialize the screen pointer by using the argument as passed through
+        ;; the `x` register.
+        lda levels_lo, x
+        sta Metatile::zp_screen_ptr
+        lda levels_hi, x
+        sta Metatile::zp_screen_ptr + 1
+
+    @loop:
+        jsr load_column
+        ;; NOTE: returned value ignored.
+
+        ;; Increase the column being used and loop if we are not at the end of
+        ;; the screen yet.
+        lda #0
+        sta zp_cur_row
+        inc zp_cur_column
+        lda zp_cur_column
+        cmp #$10
+        bne @loop
+
+        ;; Advance the screen pointer so it points to the second screen.
+        ADVANCE_SCREEN_PTR(1)
+
+        ;; Enable back the PPU on the next NMI call.
+        lda #%00011110
+        sta PPU::zp_mask
+
+        ;; And set the `ppu`, `level` and `nametable` flags.
+        lda #%01100100
+        ora Globals::zp_flags
+        sta Globals::zp_flags
+
+        ;; Reset things again just in case.
+        jsr Background::reset
+
+        lda #1
+        rts
+    .endproc
+
+    ;; Translates the given metatile coordinate byte as given on the `a`
+    ;; register and sets the first four argument variables from it. A "metatile
+    ;; coordinate byte" is the first byte of a metatile reference on the screen
+    ;; definition. Hence, the high nibble contains the Y metatile coordinate,
+    ;; and the low nibble contains the X metatile coordinate.
+    ;;
+    ;; After calling this function you will have the following:
+    ;;
+    ;;  - zp_arg0: Y metatile coordinate (i.e. the high nibble as given on the
+    ;;    `a` register shifted 4 times right).
+    ;;  - zp_arg1: X metatile coordinate (i.e. the low nibble as given on the
+    ;;    `a` register and the high nibble zero'ed out).
+    ;;  - zp_arg2: low byte for the corresponding PPU base address.
+    ;;  - zp_arg3: high byte for the corresponding PPU base address.
+    .proc translate_mt_coordinates_to_arguments
+        ;; Save the X metatile coordinate. This is also saved into a temporary
+        ;; variable so it can be messed up when figuring out the PPU address.
+        tay
+        and #%00001111
+        sta Globals::zp_arg1
+        sta Globals::zp_tmp0
+
+        ;; Save the Y metatile coordinate.
+        tya
+        and #%11110000
+        lsr
+        lsr
+        lsr
+        lsr
+        sta Globals::zp_arg0
+
+        ;; Now that the metatile coordinates have been isolated, let's go for
+        ;; the PPU address.
+
+        ;; The initial value for the high byte of the PPU address depends on the
+        ;; nametable being used. For this we can take advantage of the fact that
+        ;; this information is on bit 2 of the global flags. If we mask this
+        ;; out, and add to it the base value of `$20`, we get either `$20` or
+        ;; `$24`, which is the base high byte for either nametable.
+        lda #%00000100
+        and Globals::zp_flags
+        clc
+        adc #$20
+        sta Globals::zp_arg3
+
+        ;; Every 4 rows the high byte of the PPU address increases by one. This
+        ;; means that with the Y metatile position, divide it by 4 (shift right
+        ;; twice), and you can add that to the high byte directly.
+        lda Globals::zp_arg0
+        lsr
+        lsr
+        clc
+        adc Globals::zp_arg3
+        sta Globals::zp_arg3
+
+        ;; After the high byte has been sorted out, notice that for the low byte
+        ;; the base can only be the four values as defined in `ppu_offsets`,
+        ;; as they repeat every four times.
+        lda Globals::zp_arg0
+        and #%00000011
+        tax
+        lda ppu_offsets, x
+
+        ;; To that base we need to apply the X coordinate. Since each metatile
+        ;; takes two tiles on the PPU, we can simply take the value on the X and
+        ;; shift it left once. And finally add the base low byte from the row
+        ;; with this shifted X value.
+        asl Globals::zp_tmp0
+        clc
+        adc Globals::zp_tmp0
+        sta Globals::zp_arg2
+
+        rts
+    ppu_offsets:
+        .byte $00, $40, $80, $C0
+    .endproc
+
+    ;; Load the column as indexed by `Background::zp_cur_column` and
+    ;; `Background:zp_cur_row`. This means that this function can be called at
+    ;; any time, and so rendering columns in multiple iterations is indeed
+    ;; possible. This is important to keep a small VRAM buffer.
+    ;;
+    ;; Returns 1 if the column has been completely loaded, 0 if there are rows
+    ;; still to be loaded for the current column after calling this function.
+    ;;
+    ;; This function also sets the `ppu` and the `column` flags as expected in
+    ;; each scenario, but it will *not* call functions such as
+    ;; `prepare_next_column`, as that is to be done by the caller if necessary.
+    .proc load_column
+    @start:
+        ;; First, compute the offset of the current metatile definition. That
+        ;; is, what's left from the current row until the Y position of the next
+        ;; metatile to push. First things first, check whether we are at the end
+        ;; of the screen definition.
+        ldy #0
+        lda (Metatile::zp_screen_ptr), y
+        cmp #$FF
+        bne @fetch_offset
+
+        ;; We are actually at the end of the screen definition. Then the offset
+        ;; is simply the current row until the end. If that turns out to be
+        ;; zero, then we are done.
+        lda #$0F
+        sec
+        sbc zp_cur_row
+        beq @done
+        sta Globals::zp_idx
+        jmp @default_loop
+
+    @fetch_offset:
+        ;; Otherwise, if we weren't at the end of the screen definition, fetch
+        ;; the offset now. Note that if the offset is zero, then we can simply
+        ;; go and push the metatile definition.
+        jsr find_offset
+        beq @push
+
+        ;; This is an inner loop, which pushes the default metatile
+        ;; `Globals::zp_idx` - 1 times.
+        sta Globals::zp_idx
+    @default_loop:
+        ;; Make up a fake metatile position at the currently evaluated pair of
+        ;; row/column. Afterwards just push it. Doing all this dance is
+        ;; certainly not too expensive, but at the same time it could be easily
+        ;; optimized. This is not done out of simplicity.
+        lda zp_cur_row
+        asl
+        asl
+        asl
+        asl
+        ora zp_cur_column
+        jsr translate_mt_coordinates_to_arguments
+        lda #$00
+        jsr Buffer::push_metatile
+        beq @end
+
+        ;; The metatile was successfully pushed, now increase the row and check
+        ;; if we have to push another default metatile.
+        inc zp_cur_row
+        dec Globals::zp_idx
+        bne @default_loop
+
+        ;; We are done pushing default metatiles. Check if we are done with the
+        ;; column entirely.
+        lda zp_cur_row
+        cmp #$0F
+        beq @done
+
+        ;; Push the metatile pointed by the screen pointer.
+    @push:
+        ;; Pick up the current metatile position byte and set the arguments as
+        ;; expected by `Buffer::push_metatile`,
+        ;; `Collision::set_background_collision` and related functions.
+        ldy #0
+        lda (Metatile::zp_screen_ptr), y
+        jsr translate_mt_coordinates_to_arguments
+
+        ;; And now call `push_metatile` with the proper metatile definition.
+        ldy #1
+        lda (Metatile::zp_screen_ptr), y
+        jsr Buffer::push_metatile
+        beq @end
+
+        ;; The two bytes for the metatile have been consumed, now advance the
+        ;; screen pointer two bytes to point to the next metatile for the
+        ;; screen.
+        ADVANCE_SCREEN_PTR(2)
+
+        ;; Increase the row being evaluated and go for another iteration if this
+        ;; wasn't the last row for the column.
+        inc zp_cur_row
+        lda zp_cur_row
+        cmp #$10
+        bne @start
+
+    @done:
+        ;; Set the `ppu` flag and unset the `column` one. That is, a column is
+        ;; now fully buffered and the PPU needs to update things.
+        lda Globals::zp_flags
+        ora #%01000000
+        and #%11101111
+        sta Globals::zp_flags
+
+        lda #1
+        rts
+    @end:
+        ;; Set the `ppu` and the `column` flags. Even if the column is not fully
+        ;; buffered yet, we can already consume the VRAM buffer and empty it.
+        ;; Setting the `column` flag will ask code elsewhere to beware of a
+        ;; pending column to be loaded and act on it (e.g. by calling
+        ;; `load_column` again to finish the task).
+        lda Globals::zp_flags
+        ora #%01010000
+        sta Globals::zp_flags
+
+        lda #0
+        rts
+    .endproc
+
+    ;; Load a new column if the engine detects that it was needed. That is, it
+    ;; will ask to load a column unless a column is already being rendered, or
+    ;; the background cycle has not been consumed yet (or disabled entirely).
+    .proc load_column_if_needed
+        ;; If the `column` or the `level` flags are set, then we can skip the
+        ;; loading altogether.
+        lda #%00011000
+        and Globals::zp_flags
+        bne @end
+
+        ;; Is the background cycle disabled? If so just go to the end.
+        lda Background::zp_background_cycle
+        cmp #DISABLED_BACKGROUND_CYCLE
+        beq @end
+
+        ;; Instead of loading a column every time it can be loaded, we add a bit
+        ;; of a hiccup (i.e. "background cycle") so an entire level is not
+        ;; loaded right away after the player moves one bit.
+        inc Background::zp_background_cycle
+        beq @do_load
+        lda Background::zp_background_cycle
+        cmp #BACKGROUND_CYCLE_MAX
+        bne @end
+        lda #$00
+        sta Background::zp_background_cycle
+
+    @do_load:
+        ;; Load the column and increase the current column if it was fully
+        ;; buffered. Otherwise it will be picked up by the `update` function.
+        jsr load_column
+        beq @end
+        JAL prepare_next_column
+
+    @end:
+        rts
+    .endproc
+
+    ;; Load the next background column. You want to call this function on the
+    ;; outside whenever a new column might be needed. This function will
+    ;; eventually call `load_column_if_needed`, and it will also handle
+    ;; overflows on the scroll values, changing the nametable being used, etc.
+    ;;
+    ;; Because of the above, this is the most suitable function to call whenever
+    ;; you want to load a new column from the outside code, even if you are
+    ;; unsure whether it's really needed.
+    .proc load_next_background
+        ;; Was there an overflow on the scroll's value?
+        lda Background::zp_scroll
+        bne @load_column_if_needed
+
+        ;; Reset the background cycle so the next screen is picked up.
+        RESET_BACKGROUND_CYCLE
+
+        ;; Flip the nametable being used so the scroll makes sense for the
+        ;; current context. After that, skip the loading of the column because
+        ;; otherwise it would start loading on the currently displayed nametable
+        ;; and that would make for weird glitches.
+        lda PPU::zp_control
+        eor #%00000001
+        sta PPU::zp_control
+        jmp @update_ppu_and_quit
+
+    @load_column_if_needed:
+        jsr Background::load_column_if_needed
+
+    @update_ppu_and_quit:
+        ;; Regardless of whether the loading of the column happened or not, the
+        ;; scroll register has to change. Hence, set the `ppu` flag so at least
+        ;; the PPU::SCROLL register is updated.
+        lda Globals::zp_flags
+        ora #%01000000
+        sta Globals::zp_flags
+
+        rts
+    .endproc
+
+    ;; Reset all the variables in preparation for the loading of a new column.
+    ;; This function assumes that a column has already been fully buffered; and
+    ;; it will also handle the case where it was the last column of the
+    ;; screen/level by poking the right flags.
+    .proc prepare_next_column
+        ;; The row index always has to be set to zero.
+        lda #0
+        sta zp_cur_row
+
+        ;; Is this the last column? If not, just quit.
+        lda zp_cur_column
+        cmp #$0F
+        bne @end
+
+        ;; Flip the nametable to be used.
+        lda Globals::zp_flags
+        eor #%00000100
+        sta Globals::zp_flags
+
+        ;; Mark the background cycle as over. This will make functions like
+        ;; `load_column_if_needed` to refrain from loading a new column
+        ;; altogether.
+        lda #DISABLED_BACKGROUND_CYCLE
+        sta zp_background_cycle
+
+        ;; Reset the column to use. Note that the `@end` section will increase
+        ;; the column to be used.
+        lda #$FF
+        sta zp_cur_column
+
+        ;; If we were at the last column, we were out of metatiles for the
+        ;; current screen. Hence, move to the next screen if possible.
+        ADVANCE_SCREEN_PTR(1)
+
+        ;; Do we have another "$FF" byte? If so then we are at the end of the
+        ;; level entirely.
+        ldy #0
+        lda (Metatile::zp_screen_ptr), y
+        cmp #$FF
+        bne @end
+
+        ;; Mark the end of the level.
+        lda #%00001000
+        ora Globals::zp_flags
+        sta Globals::zp_flags
+
+    @end:
+        ;; Increase the current column and quit.
+        inc zp_cur_column
+        rts
+    .endproc
+
+    ;;;
+    ;; Find the distance between the currently evaluated row and the one pointed
+    ;; by the given argument. The argument is a byte where the Y value is
+    ;; located on the high nibble, as expected from a metatile definition.
+    ;;
+    ;; NOTE: the argument and the returned value are on the `a` register.
+    .proc find_offset
+        ;; Are we at the right column?
+        sta Globals::zp_arg0
+        and #%00001111
+        cmp zp_cur_column
+        bne @to_the_end
+
+        ;; We are: go until the current Y index.
+        lda Globals::zp_arg0
+        lsr
+        lsr
+        lsr
+        lsr
+        jmp @compute
+
+    @to_the_end:
+        ;; We aren't: go until the last row.
+        lda #$0F
+
+    @compute:
+        ;; Subtract the current row with the one we are trying to reach.
+        sec
+        sbc zp_cur_row
+        rts
+    .endproc
+.endscope
diff --git a/scroll/include/buffer.s b/scroll/include/buffer.s
new file mode 100644
index 0000000..fd08c92
--- /dev/null
+++ b/scroll/include/buffer.s
@@ -0,0 +1,258 @@
+;;;
+;; VRAM buffer
+;;
+;; As a general rule, you do *not* modify the PPU memory outside of `nmi` code:
+;; this can incur into graphical glitches because of writing into PPU memory
+;; while the PPU is consuming it. Instead, you append all writes into a VRAM
+;; buffer so the changes are finally committed during VBlank.
+;;
+;; The format of the VRAM buffer being used here is quite straight-forward
+;; and needs three bytes per update: two bytes for the PPU address, and
+;; another for the data to be passed into the PPU. This means that for each
+;; metatile we need 3 bytes per tile update * 4 tiles per metatile = 12
+;; bytes per update on a metatile. This is important when considering the
+;; capacity for this buffer. One idea could be to allow a whole column of
+;; metatiles to fit this buffer, which would require 12 bytes per metatile *
+;; 15 rows = 180 bytes. This does not look like a lot at first, but bear in
+;; mind that flushing this buffer has to be done during VBlank, which has a
+;; very tight time frame. Testing this number I noticed that even with 160
+;; bytes for this buffer (the suggested size on the NesDev wiki:
+;; https://www.nesdev.org/wiki/Sample_RAM_map) the NMI code would be
+;; struggling to fit the time frame.
+;;
+;; Thus, by poking with different values, I ended up with 96 bytes for this
+;; buffer, which allows for updates on 8 metatiles (a little more than half
+;; a screen's column). This looks quite reasonable with all the tests I have
+;; done.
+;;
+;; As a final note, on a real game you'd want to be more "aggressive" than this,
+;; as the format on the VRAM buffer directly influences how fast the code during
+;; VBlank can go. Hence, this buffer might not be suitable for your case.
+.scope Buffer
+    ;; The index of the next element to be evaluated on the VRAM buffer. If this
+    ;; index reaches `VRAM_BUFFER_CAP` then we are sure that we are done
+    ;; flushing updates. It's up for the user to notice when the buffer is full
+    ;; and should stop sending in more updates: the user should stop things as
+    ;; they are and continue on the next frame by picking up this index again.
+    zp_vram_idx = $19
+
+    ;; The VRAM buffer lives at a somewhat dangerous RAM address: $0100. This
+    ;; can be dangerous because we set the stack to start at $01FF and grows in
+    ;; decreasing addresses. Considering that the buffer is capped at $60, it
+    ;; means that if there are no buffer overflows, the memory layout for the
+    ;; $01xx RAM space is:
+    ;;
+    ;;   - $0100-$0159: VRAM buffer (96 bytes).
+    ;;   - $0160-$01FF: Stack (160 bytes).
+    ;;
+    ;; Crossing fingers there are no leaks on the VRAM buffer, we have 160 bytes
+    ;; for the stack, which realistically should be more than enough for any
+    ;; game out there. That is, unless you are abusing the stack a la
+    ;; Battletoads.
+    m_vram_buffer = $0100
+
+    ;; As explained above, the maximum capacity for the VRAM buffer, which
+    ;; amounts to 96 bytes (allowing for updates on 8 metatiles during VBlank).
+    VRAM_BUFFER_CAP = $60
+
+    ;; A tile in the VRAM buffer spans 3 bytes: 2 for the address, 1 for the
+    ;; sprite identifier. Since a metatile is made up of 4 tiles, then a
+    ;; metatile "bucket" is 3 * 4 = 12 bytes long.
+    METATILE_SIZE = 3 * 4
+
+    ;; Push the given metatile into the graphics and the collision map. How the
+    ;; graphics are updated depends on whether the PPU is enabled or not. If
+    ;; it's enabled then all updates are pushed into the VRAM buffer (given that
+    ;; enough space is available on this buffer, see more on the return value).
+    ;; Otherwise it pushes data into the PPU directly so everything is good
+    ;; after enabling the PPU back. This function expects the following
+    ;; *parameters*:
+    ;;
+    ;;   - `a` register: the metatile definition as described in `metatile.s`.
+    ;;     More specifically, the second byte on a metatile reference on the
+    ;;     screen. Hence, something like `$81` for metatile with ID = 1 and with
+    ;;     collision.
+    ;;   - `zp_arg0`: the Y metatile coordinate.
+    ;;   - `zp_arg1`: the X metatile coordinate.
+    ;;   - `zp_arg2`: low byte for the corresponding PPU base address.
+    ;;   - `zp_arg3`: high byte for the corresponding PPU base address.
+    ;;
+    ;; For the memory arguments see:
+    ;; `Background::translate_mt_coordinates_to_arguments`, as calling this
+    ;; function will set these values properly.
+    ;;
+    ;; *Returns* 1 if the push was successful, 0 otherwise. A push can only
+    ;; *not* be successful if the VRAM buffer doesn't have enough space for the
+    ;; push.
+    ;;
+    ;; NOTE: the VRAM size check also happens even if the PPU is disabled. That
+    ;; is, whenever you disable the PPU, make sure to mark the VRAM buffer as
+    ;; free as well.
+    .proc push_metatile
+        ;; Do we actually have room for this metatile? If not then just return
+        ;; early.
+        ldx zp_vram_idx
+        cpx #(VRAM_BUFFER_CAP - METATILE_SIZE)
+        bne @has_enough_space
+        lda #0
+        rts
+
+    @has_enough_space:
+        ;; Set the proper collision for the given metatile definition, while
+        ;; also preserving the value from the `a` register.
+        pha
+        jsr Collision::set_background_collision
+        pla
+
+        ;; Set the `y` register to contain the index on the `metatiles` table
+        ;; from `metatile.s`. This index will then be incremented on each
+        ;; iteration to index each tile on a metatile.
+        A_TO_METATILE_INDEX
+        tay
+
+        ;; Is the PPU enabled?
+        lda #%00011000
+        and PPU::zp_mask
+        bne @buffered_push
+
+        ;; The PPU is disabled, we can push directly graphics into it.
+        ldx #$00
+    @direct_push_loop:
+        bit PPU::STATUS
+
+        ;; Write the high byte of the PPU address being used.
+        lda Globals::zp_arg3
+        sta PPU::ADDRESS
+
+        ;; The low byte for the PPU address is the base as given from the
+        ;; parameter plus an offset that we have pre-computed (nothing too
+        ;; misterious, just check what we need to add for each tile of the
+        ;; metatile).
+        lda ppu_offsets, x
+        clc
+        adc Globals::zp_arg2
+        sta PPU::ADDRESS
+
+        ;; Load the tile id for this part of the metatile and increment the `y`
+        ;; register in preparation for the next iteration.
+        lda (Metatile::zp_metatile_ptr), y
+        sta PPU::DATA
+        iny
+
+        inx
+        cpx #4
+        bne @direct_push_loop
+        jmp @end
+
+    @buffered_push:
+        ;; Set on `Globals::zp_tmp0` the index on the metatile table, and on
+        ;; `Globals::zp_tmp1` the loop index.
+        sty Globals::zp_tmp0
+        lda #$00
+        sta Globals::zp_tmp1
+        ldx zp_vram_idx
+
+    @buffered_push_loop:
+        ;; The first element being pushed is the high byte for the PPU address
+        ;; to be used, which can be pushed directly.
+        lda Globals::zp_arg3
+        sta m_vram_buffer, x
+        inx
+
+        ;; The low byte needs some addition as in `@direct_push_loop`, but
+        ;; taking into account that the index is handled in memory instead of
+        ;; registers.
+        ldy Globals::zp_tmp1
+        lda ppu_offsets, y
+        clc
+        adc Globals::zp_arg2
+        sta m_vram_buffer, x
+        inx
+
+        ;; Load on `y` the index on the metatile table and increase it for the
+        ;; next iteration.
+        ldy Globals::zp_tmp0
+        inc Globals::zp_tmp0
+
+        ;; And push the tile id for this part.
+        lda (Metatile::zp_metatile_ptr), y
+        sta m_vram_buffer, x
+        inx
+
+        ;; Should we continue?
+        inc Globals::zp_tmp1
+        lda Globals::zp_tmp1
+        cmp #4
+        bne @buffered_push_loop
+
+        ;; And save the current index for the VRAM buffer.
+        stx zp_vram_idx
+
+    @end:
+        ;; End of the push, set the proper return value.
+        lda #1
+        rts
+    ppu_offsets:
+        .byte $00, $01, $20, $21
+    .endproc
+
+    ;; Flush the current contents of the VRAM buffer into the PPU. This function
+    ;; will also set the "ppu" flag if data was actually pushed.
+    ;;
+    ;; NOTE (NMI): this function should *only* be called during VBlank.
+    .proc flush_vram_buffer
+        ldx #$FF
+
+    @loop:
+        inx
+        cpx zp_vram_idx
+        beq @after_loop
+
+        bit PPU::STATUS
+
+        lda m_vram_buffer, x
+        sta PPU::ADDRESS
+
+        inx
+        lda m_vram_buffer, x
+        sta PPU::ADDRESS
+
+        inx
+        lda m_vram_buffer, x
+        sta PPU::DATA
+
+        jmp @loop
+
+    @after_loop:
+        ;; If no loop cycle was performed then, in theory, there is no data to
+        ;; be buffered into PPU. If this was not the case, and things had to be
+        ;; updated anyways (e.g. the scroll moved regardless of this), then it's
+        ;; up for the caller to decide.
+        lda zp_vram_idx
+        beq @end
+
+        ;; Buffer a write into the PPU control register, so the scroll is also
+        ;; left at a known state after touching the PPU address register.
+        lda #%01000000
+        ora Globals::zp_flags
+        sta Globals::zp_flags
+
+        ;; If there were some bytes actually consumed by this loop, ensure that
+        ;; it's reset.
+        lda #0
+        sta zp_vram_idx
+
+    @end:
+        rts
+    .endproc
+.endscope
+
+;; Flush the VRAM buffer if there are pending updates.
+;;
+;; NOTE (NMI): this macro should *only* be used during VBlank.
+.macro FLUSH_PENDING_VRAM_BUFFER
+    lda Buffer::zp_vram_idx
+    beq :+
+    jsr Buffer::flush_vram_buffer
+:
+.endmacro
diff --git a/scroll/include/collision.s b/scroll/include/collision.s
new file mode 100644
index 0000000..37debc5
--- /dev/null
+++ b/scroll/include/collision.s
@@ -0,0 +1,234 @@
+;;;
+;; Collisions for the background are handled by building up an array starting at
+;; RAM address $0700. This array contains a list of values to be considered when
+;; checking up background collision, and it's expected to be built when loading
+;; a screen or parts of it (e.g. `load_column`).
+;;
+;; In summary, each row of metatiles is mapped by two bytes. There are 16 tiles
+;; in a row (see `metatile.s`), so each bit from these two bytes map whether
+;; that metatile has a collision block or not. More specifically, the first bit
+;; from the first byte maps the first metatile on that row, the second bit from
+;; the first byte the second metatile on that row, and so on. Considering that
+;; there are 15 rows of tiles * 2 bytes for each row, it means that we can get a
+;; whole screen mapped with 30 bytes. As explained on the `metatile.s` file and
+;; taking the example of Super Mario Bros., things could have been further
+;; compressed by disregarding some of the rows which are anyways unaccessible by
+;; the player (e.g. by the HUD), but I have tried to keep things generic.
+;;
+;; Last but not least, we reserve space for two screens: one per nametable.
+;; Thus, the memory usage looks like this:
+;;
+;;   - $0700-$071D: collision map for the screen on nametable 0.
+;;   - $071E-$073B: collision map for the screen on nametable 1.
+.scope Collision
+    ;; The high byte value for the collision map (which lives at $0700 in RAM).
+    COLLISION_MAP_HI = $07
+
+    ;; Starting index for the collision map of the second screen.
+    COLLISION_MAP_SECOND = $1E
+
+    ;; Shadow references to Player::zp_screen_{x,y} and Background::zp_scroll.
+    ;; Not awesome, but cc65 doesn't allow to re-open scopes or similar
+    ;; code-sharing scenarios.
+    zp_player_screen_x   = $30
+    zp_player_screen_y   = $31
+    zp_background_scroll = $95
+
+    ;; Set/unset the proper bit on the collision map for the given metatile
+    ;; position and collision bit.
+    ;;
+    ;; The `a` register has to contain the collision information as given on a
+    ;; metatile reference for a screen. Hence, something like `$81` informs that
+    ;; metatile with id 1 is a solid block, while `$01` would mean the same but
+    ;; unsetting collision for that block.
+    ;;
+    ;; In memory you have to pass `Globals::zp_arg0` and `Globals::zp_arg1`,
+    ;; which contain the Y and X metatile coordinates as passed around in
+    ;; `Buffer::push_metatile`. This function won't touch other parameters, so
+    ;; future code can rely on `Globals::zp_arg2` and the rest to be preserved
+    ;; (i.e. only `Globals::zp_arg0` and `Globals::zp_arg1` will be affected).
+    .proc set_background_collision
+        ;; Preserve the collision byte for later.
+        sta Globals::zp_tmp0
+
+        ;; Select the screen to be used and save it into the `y` register, as
+        ;; expected by the `get_background_collision_y` function.
+        lda Globals::zp_flags
+        and #%00000100
+        lsr
+        lsr
+        tay
+
+        ;; Get the bitmap address for the Y and X coordinates.
+        jsr get_background_collision_y
+        tax
+
+        ;; Should we set or unset the collision bit?
+        lda Globals::zp_tmp0
+        bmi @set_collision
+
+        ;; We have to unset it: flip the mask, unset the given bit and save the
+        ;; result.
+        txa
+        eor #$FF
+        and (Globals::zp_arg0), y
+        sta (Globals::zp_arg0), y
+        rts
+
+    @set_collision:
+        ;; We have to set it: apply the given mask and save the result.
+        txa
+        ora (Globals::zp_arg0), y
+        sta (Globals::zp_arg0), y
+        rts
+    .endproc
+
+    ;; Get the collision value for the Y and X positions of a metatile. These
+    ;; two coordinates are given in memory arguments in this exact order, while
+    ;; the `y` register selects which map to pick from (i.e. the screen).
+    ;;
+    ;; Returns two bytes in memory which correspond to the 16-bit pointer to the
+    ;; row section that contains this metatile (remember that a row of metatiles
+    ;; in screen is split in two). This pointer is complemented by the fact that
+    ;; the `y` register will be set to the proper value. Hence, after calling
+    ;; this function you can get the collision bitmap for the row section with:
+    ;;
+    ;;    lda (Globals::zp_arg0), y
+    ;;
+    ;; The `a` register will contain the mask which pin points the exact
+    ;; location on the X coordinate. This is useful for collision detection.
+    ;; Hence, combining with the above you can do:
+    ;;
+    ;;    jsr get_background_collision_y
+    ;;    and (Globals::zp_arg0), y    ; Use the mask on `a` against the bitmap.
+    ;;    bne @collision               ; Collision!
+    ;;
+    .proc get_background_collision_y
+        ;; Each row is made up of two bytes worth of bitmaps. Hence, the Y
+        ;; metatile position has to be multiplied by two to get to the proper
+        ;; row.
+        lda Globals::zp_arg0
+        asl
+
+        ;; Depending on the nametable being mapped, we have to pick one map or
+        ;; the other. Since they are contiguous, this is a matter of simply
+        ;; adding the base address for the second map if that's the case.
+        cpy #1
+        bne @store_base_pointer
+        clc
+        adc #COLLISION_MAP_SECOND
+    @store_base_pointer:
+        sta Globals::zp_arg0
+
+        ;; Now it's time to figure things out given the value for the X
+        ;; coordinate. First of all, set on the `y` register which bitmap to
+        ;; pick for the row of metatiles (remember that a row is split in two
+        ;; bitmaps). Moreover, update the coordinate if it's on the second
+        ;; bitmap.
+        lda Globals::zp_arg1
+        ldy #0
+        cmp #8
+        bcc @eval_x
+        sec
+        sbc #8
+        iny
+
+        ;; With that, now we have to shift as many times as the X coordinate to
+        ;; get the mask for the bitmap we are pointing at.
+    @eval_x:
+        tax
+        lda #1
+        cpx #0
+        beq @end
+    @loop:
+        asl
+        dex
+        bne @loop
+
+    @end:
+        ;; The only thing missing is to set the high byte of the end 16-bit
+        ;; pointer.
+        ldx #COLLISION_MAP_HI
+        stx Globals::zp_arg1
+
+        ;; NOTE: Returned value is left untouched into the `a` register,
+        ;; computed in the previous loop.
+
+        rts
+    .endproc
+
+    ;; Translates the player's screen coordinates into metatile ones.
+    ;;
+    ;; This function takes two memory arguments, which are offsets to the Y and
+    ;; X screen coordinates respectively. The screen coordinates are taken
+    ;; directly from `Player::zp_screen_{y,x}`, and it also handles
+    ;; `Background::zp_scroll`.
+    ;;
+    ;; In turn this function will store in `zp_arg0` the Y metatile coordinate,
+    ;; and in `zp_arg1` the X one. The `y` register will also point if we are in
+    ;; nametable 0 or 1.
+    .proc screen_to_mt_coordinates
+        ;; Add up the player's screen coordinates with the given Y offset.
+        ;; Transforming that result into a metatile index is basically a matter
+        ;; of shifting right 4 times, since metatiles are 16x16 pixels.
+        lda Globals::zp_arg0
+        clc
+        adc Collision::zp_player_screen_y
+        lsr
+        lsr
+        lsr
+        lsr
+        sta Globals::zp_arg0
+
+        ;; The `y` register is generally the nametable where the focus is on.
+        ;; This is handled on the PPU control register, and shadowed by a
+        ;; `PPU::zp_control` variable. This will have to be adjusted depending
+        ;; on the current scroll. See more below.
+        lda #%00000001
+        and PPU::zp_control
+        tay
+
+        ;; Before figuring out the X metatile coordinates, compute the metatile
+        ;; coordinates if we were only to consider the scroll. This gives us
+        ;; information on from which metatile column is the current nametable
+        ;; visible. It will be used to correct the `y` register whenever columns
+        ;; from the next nametable start appearing on the player's area.
+        lda Collision::zp_background_scroll
+        lsr
+        lsr
+        lsr
+        lsr
+        sta Globals::zp_tmp0
+
+        ;; The coordinates for the X axis is similar to the Y one but we also
+        ;; need to take into account scrolling. Note that overflow is permitted
+        ;; as this will be handled by the correction code below (i.e. if the X
+        ;; metatile is beyond what can be seen by the current nametable, focus
+        ;; on the next one).
+        lda Globals::zp_arg1
+        clc
+        adc Collision::zp_player_screen_x
+        clc
+        adc Collision::zp_background_scroll
+        lsr
+        lsr
+        lsr
+        lsr
+        sta Globals::zp_arg1
+
+        ;; Is the X metatile coordinate actually smaller than the scrolling
+        ;; view?
+        cmp Globals::zp_tmp0
+        bcs @end
+
+        ;; Yes! Then we are on the next nametable. Amend the value on the `y`
+        ;; register and make sure it's either 1 or 0.
+        iny
+        tya
+        and #%00000001
+        tay
+
+    @end:
+        rts
+    .endproc
+.endscope
diff --git a/scroll/include/driver.s b/scroll/include/driver.s
new file mode 100644
index 0000000..abb3f5e
--- /dev/null
+++ b/scroll/include/driver.s
@@ -0,0 +1,98 @@
+;; The runtime of the game.
+.scope Driver
+    ;; Index of the current level.
+    zp_level = $77
+
+    ;; Whenever "Select" is pressed, it's used to count how many cycles have to
+    ;; pass before allowing another level switch.
+    zp_counter = $97
+
+    ;; Value that will be set for the `zp_counter` whenever the "Select" button
+    ;; is pressed for the first time.
+    SELECT_COUNTER_VALUE = $20
+
+    ;; Initialize global variables and other variables which have an impact on
+    ;; the runtime of the game.
+    .proc init
+        lda #0
+        sta Globals::zp_arg0
+        sta Globals::zp_arg1
+        sta Globals::zp_arg2
+        sta Globals::zp_arg3
+        sta Globals::zp_arg4
+        sta Globals::zp_tmp0
+        sta Globals::zp_tmp1
+        sta Globals::zp_tmp2
+        sta Globals::zp_tmp3
+        sta Globals::zp_idx
+        sta Globals::zp_flags
+        sta zp_level
+        sta zp_counter
+
+        rts
+    .endproc
+
+    ;; Function to be called at each iteration of the game loop. It loads any
+    ;; pending background columns and checks for level selection.
+    .proc update
+        ;; Are we actually loading a new level?
+        lda #%00100000
+        and Globals::zp_flags
+        beq @new_level
+
+        ;; Is the user actually allowed to hit the 'Select' button?
+        lda zp_counter
+        beq @check_select
+        dec zp_counter
+        jmp @check_column
+
+    @check_select:
+        ;; If `select` is pressed, then go for a new level. Otherwise check
+        ;; whether there is a pending column to be loaded.
+        lda #Joypad::BUTTON_SELECT
+        and Joypad::m_buttons1
+        beq @check_column
+
+        ;; Set the counter.
+        lda #SELECT_COUNTER_VALUE
+        sta zp_counter
+
+        ;; 'Select' was pressed and we weren't loading another level. Hence,
+        ;; increase the level index and load it.
+        inc zp_level
+    @new_level:
+        ;; Check the level index.
+        ldx zp_level
+        cpx #$02
+        bne :+
+        ldx #$00
+    :
+        ;; And load the computed level.
+        stx zp_level
+        jsr Background::load_level_x
+        beq :+
+
+        ;; Level could be rendered, enable back the PPU.
+        lda PPU::zp_mask
+        sta PPU::MASK
+
+        ;; Set the player back to its position.
+        RESET_PLAYER_POSITION
+    :
+        rts
+
+    @check_column:
+        ;; Do nothing if the `column` flag is not set.
+        lda #%00010000
+        and Globals::zp_flags
+        beq @end
+
+        ;; Continue loading the pending column.
+        jsr Background::load_column
+        beq @end
+        JAL Background::prepare_next_column
+
+    @end:
+        rts
+    .endproc
+.endscope
diff --git a/scroll/include/globals.s b/scroll/include/globals.s
index d941736..87a93ea 100644
--- a/scroll/include/globals.s
+++ b/scroll/include/globals.s
@@ -1,43 +1,39 @@
-;; Resets the background buffer to a zero-sized array.
-.macro CLEAR_BACKGROUND_BUFFER
-    lda #$FF
-    sta Globals::m_background_buffer
-.endmacro
-
-;; TODO: macro PUSH_TO_BACKGROUND_BUFFER_X
-
-;; Global variables used throughout the scrolling examples.
+;; Global variables used throughout the code base.
 .scope Globals
     ;;;
-    ;; Temporary values.
-
-    m_tmp_1  = $90
-    m_tmp_2  = $91
-    m_tmp_3  = $92
-
-    ;; Current value for the scroll on the X axis.
-    m_scroll = $93
-
-    m_background_idx = $19
+    ;; Argument values as defined in https://github.com/mssola/style.nes. Note
+    ;; that these variables can also be used as temporary variables.
+    zp_arg0 = $90
+    zp_arg1 = $91
+    zp_arg2 = $92
+    zp_arg3 = $93
+    zp_arg4 = $94
 
-    ;; TODO
-    m_collisions = $20
-
-    ;; TODO
-    m_background_buffer = $40
-
-    ;; Initialize global variables.
-    .proc init
-        lda #0
-        sta m_tmp_1
-        sta m_tmp_2
-        sta m_tmp_3
-        sta m_scroll
-        sta m_background_idx
-
-        CLEAR_BACKGROUND_BUFFER
-        sta m_collisions
+    ;;;
+    ;; Random values that can be used inside of functions for temporary values
+    ;; so `zp_argX` variables are not overwritten as often.
+    zp_tmp0 = $9A
+    zp_tmp1 = $9B
+    zp_tmp2 = $9C
+    zp_tmp3 = $9D
 
-        rts
-    .endproc
+    ;;;
+    ;; Reserve a byte of memory for preserving indices on memory. This is needed
+    ;; whenever the `x` and `y` registers might not be reliable because of
+    ;; underlying `jsr` calls that might tamper with their values. Sometimes
+    ;; saving the value in memory is enough instead of playing with the stack.
+    zp_idx = $9E
+
+    ;; Flags that manage the state of the game.
+    ;;
+    ;; | Bit | Short name | Meaning when set                                            |
+    ;; |-----+------------+-------------------------------------------------------------|
+    ;; |   7 | render     | Game logic is over, block main code until NMI code is over. |
+    ;; |   6 | ppu        | PPU registers (2000 and scroll) have to be updated.         |
+    ;; |   5 | level      | The current level is active.                                |
+    ;; |   4 | column     | The current column is still being buffered into the PPU.    |
+    ;; |   3 | end        | We are at the end of the level                              |
+    ;; |   2 | nametable  | Next nametable to be used                                   |
+    ;; | 1-0 | -          | Unused                                                      |
+    zp_flags = $20
 .endscope
diff --git a/scroll/include/metatile.s b/scroll/include/metatile.s
new file mode 100644
index 0000000..01dd0e8
--- /dev/null
+++ b/scroll/include/metatile.s
@@ -0,0 +1,209 @@
+;;;
+;; This file contains definitions for metatiles and screens.
+;;
+;; *Metatiles* are a way to group multiple related background tiles into a
+;; single entity. For the programs using this file, a metatile is a square of 4
+;; tiles: that is, a block of 16x16 pixels. Screens are going to use metatiles
+;; as building blocks, never tiles in on themselves. This means that rendering
+;; and collision checking is done at the metatile level, never at the tile
+;; level. This simplifies memory consumption a lot (e.g. smaller collision maps,
+;; smaller screen definitions, etc.). As how things are defined here, a metatile
+;; is a list of four consecutive bytes, which each represent the index on the
+;; pattern table for the top-left, top-right, bottom-left and bottom-right tiles
+;; respectively.
+;;
+;; Because of the above, a *screen* is laid out as a grid of 16x15 metatiles.
+;; This means that we can set the metatile position with a single byte. In our
+;; case, the high nibble will represent the "y coordinate" and the low one the
+;; "x coordinate". For example, if we have a byte like so "$12", then it means
+;; that it's the metatile at "1" on the vertical axis of the metatile grid, and
+;; "2" on the horizontal one. In the rest of the code we would also say that $01
+;; are its Y "metatile coordinates" and $02 its X "metatile coordinates". We
+;; could have gone a step further and do like Super Mario Bros. which disregards
+;; some of the metatile rows because they are not entirely visible anyways on a
+;; CRT screen and because of the HUD on top. This amount of compression is not
+;; needed here. That being said, there is one important gotcha: tiles have to be
+;; sorted, both on the X and the Y axis. This makes the background loader
+;; snappier.
+;;
+;; Other than that, we use another byte to store properties for each metatile
+;; that we are placing. Since this is quite simple, we just reserve bit 7 to set
+;; whether the metatile is to be considered for collisions or not, and the rest
+;; encodes the metatile index. The metatile index simply identifies which of the
+;; metatiles in `metatiles` we are referencing. Thus, a value of "$81" means
+;; that we want the metatile with index 1 on `metatiles` and that it has to be
+;; considered for collision checking.
+;;
+;; All in all, a screen here is a bunch of pairs of bytes, each pair encoding a
+;; metatile for the screen. A screen definition stops whenever there is the byte
+;; $FF.
+;;
+;; A *level* consists of one or more screens, each of them enclosed by the
+;; termination byte $FF. A level is finished whenever we find the termination
+;; $FF after another $FF one. Levels are indexed by the `levels_lo` and
+;; `levels_hi` lists, which contain the low byte and the high byte respectively
+;; of the address for each level. Splitting a 16-bit pointer into two separate
+;; lists is quite convenient due to the architecture of the 6502 and indexing
+;; operations. The level list is also closed by a '$FF' byte, which would result
+;; in an impossible $FFFF pointer.
+;;
+;; One missing (and obvious) feature is allowing to also specify different
+;; values for the attribute table. But I thought that this would add more
+;; complexity to a code that was already too complex for the sake of giving away
+;; an example.
+
+;; Advance the screen pointer by the given `increment`.
+.macro ADVANCE_SCREEN_PTR increment
+    lda Metatile::zp_screen_ptr
+    clc
+    adc #increment
+    sta Metatile::zp_screen_ptr
+    lda #0
+    adc Metatile::zp_screen_ptr + 1
+    sta Metatile::zp_screen_ptr + 1
+.endmacro
+
+;; Transform the current value of the `a` register to a proper index for the
+;; `metatiles` table. Note that the current value on the `a` register is assumed
+;; to be a metatile definition (that is, the second byte of a metatile on a
+;; screen). The end result is also left on the `a` register.
+.macro A_TO_METATILE_INDEX
+    ;; Each metatile definition on `metatiles` is 4 bytes long. Hence, the index
+    ;; on that table is simply the given index multiplied by four. Or more
+    ;; simply, shifted left twice. Moreover, shifting left at least once removes
+    ;; the most significant bit, which was the collision bit that we needed to
+    ;; discard anyways. All in all, this operation for now is simply shifting
+    ;; the current value twice.
+    asl
+    asl
+.endmacro
+
+;; Holds variables which are useful to manipulate metatiles and how they are
+;; laid out on screens.
+.scope Metatile
+    ;; The "screen pointer". A 16-bit pointer which points to the current
+    ;; metatile definition for the current screen.
+    zp_screen_ptr  = $70
+    zp_screen_ptr1 = $71
+
+    ;; The "metatile pointer". A 16-bit pointer which points to the base table
+    ;; of metatile definitions.
+    zp_metatile_ptr  = $72
+    zp_metatile_ptr2 = $73
+
+    ;; Initialize the pointers to handle metatiles on the game.
+    .proc init
+        lda #<metatiles
+        sta zp_metatile_ptr
+        lda #>metatiles
+        sta zp_metatile_ptr + 1
+
+        ;; NOTE: the screen pointer is supposed to be initialized when loading a
+        ;; new level.
+
+        rts
+    .endproc
+.endscope
+
+;; List of metatiles.
+metatiles:
+    ;; Default metatile: transparent.
+    .byte $00, $00, $00, $00
+
+    ;; Super Mario Bros. block.
+    .byte $05, $07, $06, $08
+
+    ;; Super Mario Bros. ground.
+    .byte $02, $01, $03, $04
+
+;; Low bytes of the address for each level.
+levels_lo:
+    .byte <level1, <level2
+
+    ;; End of levels.
+    .byte $FF
+
+;; High bytes of the address for each level.
+levels_hi:
+    .byte >level1, >level2
+
+    ;; End of levels.
+    .byte $FF
+
+;;;
+;; List of levels.
+
+level1:
+    ;;;
+    ;; Screen 1.
+    .byte $10, $81
+    .byte $40, $81
+    .byte $50, $81
+    .byte $80, $81
+    .byte $A0, $81
+    .byte $94, $01
+    .byte $65, $81
+    .byte $75, $81
+    .byte $37, $81
+    .byte $38, $81
+    .byte $39, $81
+    .byte $3A, $81
+    .byte $3B, $81
+    .byte $3C, $81
+    .byte $4C, $81
+    .byte $3D, $81
+    .byte $4F, $81
+    .byte $FF
+
+    ;;;
+    ;; Screen 2.
+    .byte $22, $81
+    .byte $A4, $81
+    .byte $C6, $81
+    .byte $FF
+
+    ;;;
+    ;; Screen 3.
+    .byte $43, $81
+    .byte $24, $81
+    .byte $A6, $81
+    .byte $FF
+
+    ;;;
+    ;; Screen 4.
+    .byte $75, $81
+    .byte $37, $81
+    .byte $38, $81
+    .byte $39, $81
+    .byte $3A, $81
+    .byte $FF
+
+    ;;;
+    ;; End of level.
+    .byte $FF
+
+level2:
+    ;;;
+    ;; Screen 1.
+    .byte $14, $81
+    .byte $25, $81
+    .byte $36, $81
+    .byte $FF
+
+    ;;;
+    ;; Screen 2.
+    .byte $47, $81
+    .byte $58, $81
+    .byte $69, $81
+    .byte $FF
+
+    ;;;
+    ;; Screen 3.
+    .byte $7A, $81
+    .byte $8B, $81
+    .byte $9C, $81
+    .byte $FF
+
+    ;;;
+    ;; End of level.
+    .byte $FF
diff --git a/scroll/include/oam.s b/scroll/include/oam.s
index ecd1b7d..5680a93 100644
--- a/scroll/include/oam.s
+++ b/scroll/include/oam.s
@@ -1,5 +1,3 @@
-.segment "CODE"
-
 .scope OAM
     ADDR = $2003
     DMA  = $4014
diff --git a/scroll/include/palettes.s b/scroll/include/palettes.s
index 9403b7b..1b9f20c 100644
--- a/scroll/include/palettes.s
+++ b/scroll/include/palettes.s
@@ -3,8 +3,18 @@
     DEFAULT_COLOR = $11
 
     ;; Copies all the palettes for our game into the proper PPU address.
+    ;;
+    ;; NOTE: as explained in `metatile.s`, the engine for the scrolling examples
+    ;; lack the ability to update the attributes for each tile. This is
+    ;; embarrasing, but adding support for it would make these examples more
+    ;; complex than they need to be. Hence we just reproduce the same palette
+    ;; all over and avoid glitches on real hardware or emulators with randomized
+    ;; memory.
     .proc init
-        PPU_ADDR $3F00
+        lda #$3F
+        sta PPU::ADDRESS
+        lda #$00
+        sta PPU::ADDRESS
 
         ldx #0
     @load_palettes_loop:
@@ -17,14 +27,14 @@
     palettes:
         ;; Background
         .byte DEFAULT_COLOR, $36, $17, $0F
-        .byte DEFAULT_COLOR, $00, $00, $00
-        .byte DEFAULT_COLOR, $00, $00, $00
-        .byte DEFAULT_COLOR, $00, $00, $00
+        .byte DEFAULT_COLOR, $36, $17, $0F
+        .byte DEFAULT_COLOR, $36, $17, $0F
+        .byte DEFAULT_COLOR, $36, $17, $0F
 
         ;; Foreground
         .byte DEFAULT_COLOR, $28, $0F, $30
-        .byte DEFAULT_COLOR, $00, $00, $00
-        .byte DEFAULT_COLOR, $00, $00, $00
-        .byte DEFAULT_COLOR, $00, $00, $00
+        .byte DEFAULT_COLOR, $28, $0F, $30
+        .byte DEFAULT_COLOR, $28, $0F, $30
+        .byte DEFAULT_COLOR, $28, $0F, $30
     .endproc
 .endscope
diff --git a/scroll/include/player.s b/scroll/include/player.s
index 95d710d..239cec6 100644
--- a/scroll/include/player.s
+++ b/scroll/include/player.s
@@ -1,87 +1,322 @@
+;; Set the player to its initial coordinates.
+.macro RESET_PLAYER_POSITION
+    lda #40
+    sta Player::zp_screen_x
+    sta Player::zp_screen_y
+.endmacro
+
 ;; This is the player code that gets re-used on all scrolling examples. This
 ;; will show a meta-sprite made up of four sprites that make up diskun. You can
 ;; move this player with the arrows, and the movement will be pretty basic (1
-;; pixel each time). Furthermore, the sprite handles collision with background
-;; elements (TODO), which is shared across all the scrolling examples.
+;; pixel at a time). Furthermore, the sprite handles collision with background
+;; elements, which is shared across all the scrolling examples.
 ;;
-;; The sprite is stored in OAM memory $200-$20F, and the $30 and $31 memory
-;; addresses are used for keeping up with the screen coordinates.
+;; The sprite is stored in OAM memory $204-$213 (hence sprite0 is left out so
+;; other examples can use it), and the $30 and $31 memory addresses are used for
+;; keeping up with the screen coordinates.
 .scope Player
-    m_screen_x = $30
-    m_screen_y = $31
+    ;; The minimum value that the player is allowed to have on the Y axis.
+    MIN_SCREEN_Y = 8
+
+    ;; The maximum value of X that is allowed before we start updating the
+    ;; scroll value.
+    MAX_SCREEN_X = 144
+
+    ;; The height for the player when it comes to detecting the edge of the
+    ;; screen or background collisions.
+    PLAYER_HEIGHT = 25
+
+    ;; The width for the player when it comes to detecting the edge of the
+    ;; screen or background collisions.
+    PLAYER_WIDTH = 16
+
+    ;; Indeces for the collisions table as defined on `check_bg_collision_x` so
+    ;; the setting of the `x` register is less magical.
+    TOP_COLLISION_INDEX   = (0 << 2)
+    RIGHT_COLLISION_INDEX = (1 << 2)
+    DOWN_COLLISION_INDEX  = (2 << 2)
+    LEFT_COLLISION_INDEX  = (3 << 2)
 
+    ;; Coordinates on the X axis.
+    ;; NOTE: shadowed in collision.s. If you change it here, change it there as
+    ;; well.
+    zp_screen_x = $30
+
+    ;; Coordinates on the Y axis.
+    ;; NOTE: shadowed in collision.s. If you change it here, change it there as
+    ;; well.
+    zp_screen_y = $31
+
+    ;; Initialize the sprite for the player.
     .proc init
-        lda #40
-        sta m_screen_x
-        sta m_screen_y
+        RESET_PLAYER_POSITION
 
         lda #$01
-        sta $201
-        lda #$02
         sta $205
-        lda #$11
+        lda #$02
         sta $209
-        lda #$12
+        lda #$11
         sta $20D
+        lda #$12
+        sta $211
 
         lda #0
-        sta $202
         sta $206
         sta $20A
         sta $20E
+        sta $212
 
         rts
     .endproc
 
+    ;; The update for the player is simply about checking for the d-pad. Note
+    ;; that if pressing right, then the scroll value might also move, which
+    ;; might mean to load the next background column until the end of the
+    ;; screen.
     .proc update
-        jsr update_coordinates
-        jsr update_sprites
-        rts
-    .endproc
-
-    .proc update_coordinates
+        ;; Is the player requesting to go up?
         lda #Joypad::BUTTON_UP
         and Joypad::m_buttons1
         beq @check_down
-        dec m_screen_y
+
+        ;; If we are already at the top disregard this button press and check
+        ;; for the left button.
+        lda #MIN_SCREEN_Y
+        cmp zp_screen_y
+        beq @check_left
+
+        ;; If, for whatever reason, we are below the minimum value, then reset
+        ;; the value to this minimum Y value and jump to check the left button.
+        bcc @update_y_up
+        sta zp_screen_y
+        jmp @check_left
+
+    @update_y_up:
+        ;; Try to go up and check for a collision with the background.
+        dec zp_screen_y
+        dec zp_screen_y
+        ldx #TOP_COLLISION_INDEX
+        jsr check_bg_collision_x
+        beq @check_left
+
+        ;; A collision was detected. Then get back to the old value of Y and
+        ;; move on.
+        inc zp_screen_y
+        inc zp_screen_y
         jmp @check_left
+
     @check_down:
+        ;; Is the player requesting to go down?
         lda #Joypad::BUTTON_DOWN
         and Joypad::m_buttons1
         beq @check_left
-        inc m_screen_y
+
+        ;; We have to move down unless we are already at the very bottom.
+        lda zp_screen_y
+        cmp #(240 - PLAYER_HEIGHT)
+        bcc @update_y_down
+        lda #(240 - PLAYER_HEIGHT)
+        sta zp_screen_y
+        jmp @check_left
+    @update_y_down:
+        ;; Try to go down as requested, but check for a collision with the
+        ;; background.
+        inc zp_screen_y
+        inc zp_screen_y
+        ldx #DOWN_COLLISION_INDEX
+        jsr check_bg_collision_x
+        beq @check_left
+
+        ;; There was a collision with the background, restore back the value on
+        ;; Y and move on.
+        dec zp_screen_y
+        dec zp_screen_y
+
     @check_left:
+        ;; Is the player requesting to go left?
         lda #Joypad::BUTTON_LEFT
         and Joypad::m_buttons1
         beq @check_right
-        dec m_screen_x
+
+        ;; We have to move left unless we are already at the leftmost edge.
+        lda zp_screen_x
+        bne :+
+        rts
+    :
+        ;; We are not at the leftmost edge, try to go left while also checking
+        ;; for a background collision.
+        dec zp_screen_x
+        dec zp_screen_x
+        ldx #LEFT_COLLISION_INDEX
+        jsr check_bg_collision_x
+        bne :+
         rts
+    :
+        ;; There was a collision, restore back the value on X and quit.
+        inc zp_screen_x
+        inc zp_screen_x
+        rts
+
     @check_right:
+        ;; Last check! Is the player requesting to go right?
         lda #Joypad::BUTTON_RIGHT
         and Joypad::m_buttons1
+        bne @check_level_end
+        rts
+
+    @check_level_end:
+        ;; Are we at the last screen? If so consume the rest of the missing
+        ;; scroll. Whenever that is done (i.e. there are no more screens and the
+        ;; scroll sits at a zero value), then we can no longer load more columns
+        ;; or update the scroll: stick to updating the X position.
+        lda #%00001000
+        and Globals::zp_flags
+        beq @check_max_screen
+        lda Background::zp_scroll
+        beq @update_x
+
+    @check_max_screen:
+        ;; Is the X position already at the MAX_SCREEN_X? If so, then we don't
+        ;; move the character: we move the scroll instead.
+        lda zp_screen_x
+        cmp #Player::MAX_SCREEN_X
+        bcc @update_x
+
+        ;; Moving will then be a matter of increasing the scroll value.
+        inc Background::zp_scroll
+        inc Background::zp_scroll
+        ldx #RIGHT_COLLISION_INDEX
+        jsr check_bg_collision_x
+        bne @collision_on_scroll
+
+        ;; If there are no collisions in the scrolling scenario, then we might
+        ;; need to load the next background column in advance. Whether that is
+        ;; possible or something that we want to do is handled automatically by
+        ;; the `Background::load_next_background` function. Hence, call this
+        ;; function and quit.
+        JAL Background::load_next_background
+
+    @collision_on_scroll:
+        ;; Collision was given. Hence, restore the value for the scroll and
+        ;; quit.
+        dec Background::zp_scroll
+        dec Background::zp_scroll
+        rts
+
+    @update_x:
+        ;; Do not overwrap the screen.
+        lda zp_screen_x
+        cmp #(256 - PLAYER_WIDTH)
         beq @end
-        inc m_screen_x
-    ;; TODO: move scroll
+
+        ;; We are not about to overwrap the screen, let's update the value on
+        ;; the X axis and check whether a collision would happen.
+        inc zp_screen_x
+        inc zp_screen_x
+        ldx #RIGHT_COLLISION_INDEX
+        jsr check_bg_collision_x
+        beq @end
+
+        ;; There was a collision, roll back the value for the X axis.
+        dec zp_screen_x
+        dec zp_screen_x
+
     @end:
         rts
     .endproc
 
-    .proc update_sprites
-        lda m_screen_x
-        sta $203
-        sta $20B
-        clc
-        adc #8
+    ;; Check collision between the player and any background element. Note that
+    ;; collision is not tested on all edges for the player, but you have to
+    ;; provide a "collision index", which refers to which edges to test. See the
+    ;; "*_COLLISION_INDEX" constants above and the `edges` data included inside
+    ;; of this function. This collision index has to be provided on the `x`
+    ;; register.
+    ;;
+    ;; Returns a non-zero value on the `a` register if a collision was detected,
+    ;; zero otherwise.
+    .proc check_bg_collision_x
+        ;; We have to check for a collision on both the left and the right
+        ;; edges, which should follow the same code. Unrolling it is simply
+        ;; faster and more clear than setting up a loop.
+        .repeat 2, I
+            ;; Save the offset on the Y coordinate.
+            lda edges, x
+            sta Globals::zp_arg0
+
+            ;; Save the offset on the X coordinate.
+            inx
+            lda edges, x
+            sta Globals::zp_arg1
+
+            ;; Save the edge index in preparation for the following calls.
+            .if I = 0
+                inx
+                stx Globals::zp_idx
+            .endif
+
+            ;; Transform screen coordinates into metatile ones and fetch the
+            ;; collision bitmap for it with its mask.
+            jsr Collision::screen_to_mt_coordinates
+            jsr Collision::get_background_collision_y
+
+            ;; Restore back the edge index.
+            .if I = 0
+                ldx Globals::zp_idx
+            .endif
+
+            ;; Checking for a collision is as simple as performing an AND with the
+            ;; bitmap at the `a` register and the byte that can be addressed with
+            ;; the given pointer. If the result is non-zero, then we found a match.
+            and (Globals::zp_arg0), y
+            .if I = 0
+                ;; If I = 1 we are falling through as expected anyways, so this
+                ;; instruction is not needed.
+                bne @collision
+            .endif
+        .endrepeat
+
+        ;; The last instruction before ending up here will already have the `a`
+        ;; register lined up to the proper return value. If a collision was
+        ;; found (in either left/right case), `bne @collision` is good, and
+        ;; hence `a` has an expected non-zero value. The same is true when not
+        ;; in a collision.
+    @collision:
+        rts
+
+    edges:
+        ;; Up: top-left (y, x), top-right (y, x)
+        .byte $02, $02, $02, $0F
+
+        ;; Right: top-right (y, x), bottom-right (y, x)
+        .byte $02, $0F, $0F, $0F
+
+        ;; Down: bottom-left (y, x), bottom-right (y, x)
+        .byte $0F, $02, $0F, $0F
+
+        ;; Left: top-left (y, x), bottom-left (y, x)
+        .byte $02, $01, $0F, $01
+    .endproc
+
+    ;; Update the memory on the PPU with the current values for the screen
+    ;; coordinates of the player.
+    ;;
+    ;; NOTE: only call this function on NMI Code.
+    .proc update_sprite
+        lda zp_screen_x
         sta $207
         sta $20F
+        clc
+        adc #8
+        sta $20B
+        sta $213
 
-        lda m_screen_y
-        sta $200
+        lda zp_screen_y
         sta $204
+        sta $208
         clc
         adc #8
-        sta $208
         sta $20C
+        sta $210
 
         rts
     .endproc
diff --git a/scroll/include/ppu.s b/scroll/include/ppu.s
index 23fd978..34410a0 100644
--- a/scroll/include/ppu.s
+++ b/scroll/include/ppu.s
@@ -1,5 +1,3 @@
-.segment "CODE"
-
 .scope PPU
     CONTROL = $2000
     MASK    = $2001
@@ -7,21 +5,10 @@
     SCROLL  = $2005
     ADDRESS = $2006
     DATA    = $2007
-.endscope
 
-.macro PPU_ADDR address
-    lda #.HIBYTE(address)
-    sta PPU::ADDRESS
-    lda #.LOBYTE(address)
-    sta PPU::ADDRESS
-.endmacro
+    ;; Variable that shadows the value on PPU::CONTROL.
+    zp_control = $80
 
-;; WRITE_PPU_DATA is a macro that will write into PPU::ADDRESS the given address
-;; and into PPU::DATA the given byte value.
-.macro WRITE_PPU_DATA address, value
-    bit PPU::STATUS
-
-    PPU_ADDR address
-    lda #value
-    sta PPU::DATA
-.endmacro
+    ;; Variable that shadows the value on PPU::MASK.
+    zp_mask = $81
+.endscope
diff --git a/scroll/include/reset.s b/scroll/include/reset.s
new file mode 100644
index 0000000..2402795
--- /dev/null
+++ b/scroll/include/reset.s
@@ -0,0 +1,62 @@
+;; Check `basics/sprite.s` for a deeper look on the logic below.
+.proc reset
+    sei
+    cld
+
+    ldx #$40
+    stx APU::FRAME_COUNTER
+
+    ldx #$FF
+    txs
+
+    inx
+    stx PPU::CONTROL
+    stx PPU::MASK
+    stx APU::DMC
+
+@vblankwait1:
+    bit PPU::STATUS
+    bpl @vblankwait1
+
+    ldx #0
+    lda #0
+@ram_reset_loop:
+    sta $000, x
+    sta $100, x
+    sta $300, x
+    sta $400, x
+    sta $500, x
+    sta $600, x
+    sta $700, x
+    inx
+    bne @ram_reset_loop
+
+    lda #$EF
+@sprite_reset_loop:
+    sta $200, x
+    inx
+    bne @sprite_reset_loop
+
+    lda #$00
+    sta OAM::ADDR
+    lda #$02
+    sta OAM::DMA
+
+@vblankwait2:
+    bit PPU::STATUS
+    bpl @vblankwait2
+
+    lda #$3F
+    sta PPU::ADDRESS
+    lda #$00
+    sta PPU::ADDRESS
+
+    lda #$0F
+    ldx #$20
+@palettes_reset_loop:
+    sta PPU::DATA
+    dex
+    bne @palettes_reset_loop
+
+    jmp main
+.endproc