agx_dcp_compression.txt

DEV 00:
001948000 ... 00194c000 -> 0000010fc5fc8000 00360000
001d68000 ... 001d6c000 -> 0000010fbec44000 00298000


01500070000 -> 10fc1230000 -> 0x9034000
                width = 3456
                height = 2234
                base = 0x00000000
                offset = 0x00000000
                stride = 0x00036000
                size = 0x01E08000
                tile_size = 1024
                tile_w = 16
                tile_h = 16

                tile_w = 16
                tile_h = 16
                meta_offset = 30965760  0x1d88000
                data_offset = 0
                tile_meta_bytes = 8
                tiles_w = 216
                tiles_h = 140
                unk1 = 0
                compression_type = 3
                unk3 = 0
                pad = None
                tile_bytes = 1024
                row_stride = 221184
                pad2 = None

01501e78000 -> 10fc3038000 -> 0xae3c000 (+0x1e08000)
                width = 3456
                height = 2234
                base = 0x01E08000
                offset = 0x01E08000
                stride = 0x0000D800
                size = 0x007E2000
                tile_size = 256
                tile_w = 16
                tile_h = 16

                tile_w = 16
                tile_h = 16
                meta_offset = 39231488  0x256a000
                data_offset = 31490048  0x1e08000
                tile_meta_bytes = 8
                tiles_w = 216
                tiles_h = 140
                unk1 = 0
                compression_type = 3
                unk3 = 0
                pad = None
                tile_bytes = 256
                row_stride = 55296
                pad2 = None


PBE
  Dimension: 2D
  Layout: Compressed
  Channels: unknown 4C (XXX)
  Type: XR
  Swizzle R: B
  Swizzle G: G
  Swizzle B: R
  Swizzle A: A
  Width: 3456
  Height: 2234
  Unk 52: false
  Rotate 90: false
  Flip vertical: false
  Samples: 2
  Unk mipmapped: false
  Compressed 1: true
  Buffer: 0x1500070000
  Level: 0
  Levels: 1
  Layers: 216 <- tiles_w
  Page-aligned layers: false
  sRGB: false
  Extended: true
  Stride: 0xd70
  Acceleration buffer: 0xd71501e78000
                         ^^ ntiles_w - 1
  Depth (linear): 1
  Layer stride (linear): 0x1501e780

Texture
  Dimension: 2D
  Layout: Compressed
  Channels: unknown 4C (XXX)
  Type: XR
  Swizzle R: B
  Swizzle G: G
  Swizzle B: R
  Swizzle A: A
  Width: 3456
  Height: 2234
  First level: 0
  Last level: 0
  Samples: 2
  Address: 0x1500070000
  Unk mipmapped: false
  Compressed 1: true
  Null: false
  Compression: 0x0
  sRGB: false
  sRGB 2-channel: false
  Stride: 0x200d70
  Depth: 216
  Page-aligned layers: false
  Extended: true
  Acceleration buffer: 0xd71501e78000
  Depth (linear): 1
  Layer stride (linear): 0x1501e780


Store shader:
   0: 7e0000098000         mov              r0l, u0l
   6: b1810080004a00000800 TODO.unkB1
  10: 8800                 stop

Load shader:
   0: f2051004             get_sr           r1.cache, sr80 (thread_position_in_grid.x)
   4: f2011104             get_sr           r0.cache, sr81 (thread_position_in_grid.y)
   8: be850a242c00         convert          u32_to_f, r1.cache, r1.discard, rte
   e: be810a042c00         convert          u32_to_f, r0.cache, r0.discard, rte
  14: aa89c2020200         fadd32           r2.cache, r1.discard, 0.5
  1a: aa85c0020200         fadd32           r1.cache, r0.discard, 0.5
  20: ba81804128880100     fmadd32          r0.cache, u0, r2.cache, u4
  28: ba8982412c8a0100     fmadd32          r2.cache, u1, r2.discard, u5
  30: 3a81842128c00200     fmadd32          r0, u2, r1.cache, r0.discard
  38: 3a8586212cc40200     fmadd32          r1, u3, r1.discard, r2.discard
  40: 3180400000620f0020000000 texture_sample   0, 0b1100, 0b0, none, xyzw, 0b01, r0l_r0h_r1l_r1h, None, ts0, ss0, tex_2d, r0_r1.discard, auto_lod, 0
  4c: 3800                 wait             0
  4e: 09000001f0fc8003     st_tile          r0l_r0h_r1l_r1h, i16, 0, xyzw, 0, 255, 0
  56: 8800                 stop