mirror of
https://github.com/dolphin-emu/dolphin.git
synced 2025-07-29 09:09:52 -06:00
Hg:
enable newline normalization get revision number via `hg svn info` for svnrev.h ignore incremental/generated binary files (windows/VS at least) leave a comment if some files need native eol set in svnprops git-svn-id: https://dolphin-emu.googlecode.com/svn/trunk@5637 8ced0084-cf51-0410-be5f-012b33b47a6e
This commit is contained in:
Shawn Hoffman
@ -1,182 +1,182 @@
|
||||
[Default]
|
||||
IRLeft = 266
|
||||
IRTop = 215
|
||||
IRWidth = 486
|
||||
IRHeight = 490
|
||||
[RSPE] # Wii Sports (NTSC)
|
||||
IRLeft = 266
|
||||
IRTop = 215
|
||||
IRWidth = 486
|
||||
IRHeight = 490
|
||||
[RSPP] # Wii Sports (PAL)
|
||||
IRLeft = 266
|
||||
IRTop = 215
|
||||
IRWidth = 486
|
||||
IRHeight = 490
|
||||
[RMGE] # Mario Galaxy (NTSC)
|
||||
IRLeft = 255
|
||||
IRTop = 278
|
||||
IRWidth = 452
|
||||
IRHeight = 456
|
||||
[RMGP] # Mario Galaxy (PAL)
|
||||
IRLeft = 255
|
||||
IRTop = 278
|
||||
IRWidth = 452
|
||||
IRHeight = 456
|
||||
[RMCE] # Mario Kart Wii (NTSC)
|
||||
IRLeft = 253
|
||||
IRTop = 272
|
||||
IRWidth = 454
|
||||
IRHeight = 455
|
||||
[RMCP] # Mario Kart Wii (PAL)
|
||||
IRLeft = 254
|
||||
IRTop = 278
|
||||
IRWidth = 451
|
||||
IRHeight = 448
|
||||
[R7PE] # Punch Out (NTSC)
|
||||
IRLeft = 265
|
||||
IRTop = 289
|
||||
IRWidth = 408
|
||||
IRHeight = 416
|
||||
[R7PP] # Punch Out (PAL)
|
||||
IRLeft = 265
|
||||
IRTop = 289
|
||||
IRWidth = 408
|
||||
IRHeight = 416
|
||||
[RZDE] # Zelda - Twilight Princess (NTSC)
|
||||
IRLeft = 233
|
||||
IRTop = 181
|
||||
IRWidth = 559
|
||||
IRHeight = 409
|
||||
[RZDP] # Zelda - Twilight Princess (PAL)
|
||||
IRLeft = 233
|
||||
IRTop = 181
|
||||
IRWidth = 559
|
||||
IRHeight = 409
|
||||
[RM8E] # Mario Part 8 (NTSC)
|
||||
IRLeft = 277
|
||||
IRTop = 273
|
||||
IRWidth = 460
|
||||
IRHeight = 394
|
||||
[RM8P] # Mario Part 8 (PAL)
|
||||
IRLeft = 277
|
||||
IRTop = 273
|
||||
IRWidth = 460
|
||||
IRHeight = 394
|
||||
[R8PE] # Super Paper Mario (NTSC)
|
||||
IRLeft = 399
|
||||
IRTop = 373
|
||||
IRWidth = 227
|
||||
IRHeight = 228
|
||||
[R8PP] # Super Paper Mario (PAL)
|
||||
IRLeft = 399
|
||||
IRTop = 373
|
||||
IRWidth = 227
|
||||
IRHeight = 228
|
||||
[R4QP] # Mario Strikers (NTSC)
|
||||
IRLeft = 200
|
||||
IRTop = 54
|
||||
IRWidth = 615
|
||||
IRHeight = 657
|
||||
[R4QP] # Mario Strikers (PAL)
|
||||
IRLeft = 200
|
||||
IRTop = 54
|
||||
IRWidth = 615
|
||||
IRHeight = 657
|
||||
[RBUE] # Resident Evil - The Umbrella Chronicles (NTSC)
|
||||
IRLeft = 335
|
||||
IRTop = 351
|
||||
IRWidth = 357
|
||||
IRHeight = 273
|
||||
[RBUP] # Resident Evil - The Umbrella Chronicles (PAL)
|
||||
IRLeft = 335
|
||||
IRTop = 351
|
||||
IRWidth = 357
|
||||
IRHeight = 273
|
||||
[RB4E] # Resident Evil 4 (NTSC)
|
||||
IRLeft = 286
|
||||
IRTop = 256
|
||||
IRWidth = 450
|
||||
IRHeight = 455
|
||||
[RB4P] # Resident Evil 4 (PAL)
|
||||
IRLeft = 286
|
||||
IRTop = 256
|
||||
IRWidth = 450
|
||||
IRHeight = 455
|
||||
[R3IJ] # Metroid Prime - Wii De Asobu (JAP)
|
||||
IRLeft = 228
|
||||
IRTop = 112
|
||||
IRWidth = 486
|
||||
IRHeight = 577
|
||||
[RM3E] # Metroid Prime 3 (NTSC)
|
||||
IRLeft = 258
|
||||
IRTop = 84
|
||||
IRWidth = 489
|
||||
IRHeight = 613
|
||||
[RM3P] # Metroid Prime 3 (PAL)
|
||||
IRLeft = 258
|
||||
IRTop = 84
|
||||
IRWidth = 489
|
||||
IRHeight = 613
|
||||
[RSUP] # Sports Party (NTSC)
|
||||
IRLeft = 391
|
||||
IRTop = 377
|
||||
IRWidth = 241
|
||||
IRHeight = 225
|
||||
[RSUP] # Sports Party (PAL)
|
||||
IRLeft = 391
|
||||
IRTop = 377
|
||||
IRWidth = 241
|
||||
IRHeight = 225
|
||||
[RDZE] # Disaster - Day of Crisis (NTSC)
|
||||
IRLeft = 253
|
||||
IRTop = 276
|
||||
IRWidth = 453
|
||||
IRHeight = 421
|
||||
[RDZP] # Disaster - Day of Crisis (PAL)
|
||||
IRLeft = 253
|
||||
IRTop = 276
|
||||
IRWidth = 453
|
||||
IRHeight = 421
|
||||
[R4QE] # Mario Strikers (NTSC)
|
||||
IRLeft = 266
|
||||
IRTop = 215
|
||||
IRWidth = 486
|
||||
IRHeight = 490
|
||||
[R4QP] # Mario Strikers (PAL)
|
||||
IRLeft = 266
|
||||
IRTop = 215
|
||||
IRWidth = 486
|
||||
IRHeight = 490
|
||||
[RPBE] # Pokemon Battle Revolution (NTSC)
|
||||
IRLeft = 287
|
||||
IRTop = 261
|
||||
IRWidth = 451
|
||||
IRHeight = 449
|
||||
[RPBP] # Pokemon Battle Revolution (PAL)
|
||||
IRLeft = 287
|
||||
IRTop = 261
|
||||
IRWidth = 451
|
||||
IRHeight = 449
|
||||
[R2GJ] # Fragile: Sayonara Tsuki no Haikyo (JAP)
|
||||
IRLeft = 254
|
||||
IRTop = 280
|
||||
IRWidth = 451
|
||||
IRHeight = 453
|
||||
[RFNE] # Wii Fit (NTSC)
|
||||
[RFNP] # Wii Fit (PAL)
|
||||
[RSBE] # Super Smash Bros. Brawl (NTSC)
|
||||
[RSBP] # Super Smash Bros. Brawl (PAL)
|
||||
[R3TE] # Top Spin 3 (NTSC)
|
||||
[RSBP] # Top Spin 3 (PAL)
|
||||
[RLBE] # Lego Batman (NTSC)
|
||||
[RLBP] # Lego Batman (PAL)
|
||||
##########################################################################
|
||||
# These games don't use the IR Pointer at any time
|
||||
##########################################################################
|
||||
[RSRE] # Sonic and the Secret Rings (NTSC)
|
||||
[RSRP] # Sonic and the Secret Rings (PAL)
|
||||
[RWLE] # Wario Land - Shake It (NTSC)
|
||||
[RWLP] # Wario Land - Shake It (PAL)
|
||||
[RTNE] # Tenchu: Shadow Assassins (NTSC)
|
||||
[RTNE] # Tenchu: Shadow Assassins (PAL)
|
||||
[Default]
|
||||
IRLeft = 266
|
||||
IRTop = 215
|
||||
IRWidth = 486
|
||||
IRHeight = 490
|
||||
[RSPE] # Wii Sports (NTSC)
|
||||
IRLeft = 266
|
||||
IRTop = 215
|
||||
IRWidth = 486
|
||||
IRHeight = 490
|
||||
[RSPP] # Wii Sports (PAL)
|
||||
IRLeft = 266
|
||||
IRTop = 215
|
||||
IRWidth = 486
|
||||
IRHeight = 490
|
||||
[RMGE] # Mario Galaxy (NTSC)
|
||||
IRLeft = 255
|
||||
IRTop = 278
|
||||
IRWidth = 452
|
||||
IRHeight = 456
|
||||
[RMGP] # Mario Galaxy (PAL)
|
||||
IRLeft = 255
|
||||
IRTop = 278
|
||||
IRWidth = 452
|
||||
IRHeight = 456
|
||||
[RMCE] # Mario Kart Wii (NTSC)
|
||||
IRLeft = 253
|
||||
IRTop = 272
|
||||
IRWidth = 454
|
||||
IRHeight = 455
|
||||
[RMCP] # Mario Kart Wii (PAL)
|
||||
IRLeft = 254
|
||||
IRTop = 278
|
||||
IRWidth = 451
|
||||
IRHeight = 448
|
||||
[R7PE] # Punch Out (NTSC)
|
||||
IRLeft = 265
|
||||
IRTop = 289
|
||||
IRWidth = 408
|
||||
IRHeight = 416
|
||||
[R7PP] # Punch Out (PAL)
|
||||
IRLeft = 265
|
||||
IRTop = 289
|
||||
IRWidth = 408
|
||||
IRHeight = 416
|
||||
[RZDE] # Zelda - Twilight Princess (NTSC)
|
||||
IRLeft = 233
|
||||
IRTop = 181
|
||||
IRWidth = 559
|
||||
IRHeight = 409
|
||||
[RZDP] # Zelda - Twilight Princess (PAL)
|
||||
IRLeft = 233
|
||||
IRTop = 181
|
||||
IRWidth = 559
|
||||
IRHeight = 409
|
||||
[RM8E] # Mario Part 8 (NTSC)
|
||||
IRLeft = 277
|
||||
IRTop = 273
|
||||
IRWidth = 460
|
||||
IRHeight = 394
|
||||
[RM8P] # Mario Part 8 (PAL)
|
||||
IRLeft = 277
|
||||
IRTop = 273
|
||||
IRWidth = 460
|
||||
IRHeight = 394
|
||||
[R8PE] # Super Paper Mario (NTSC)
|
||||
IRLeft = 399
|
||||
IRTop = 373
|
||||
IRWidth = 227
|
||||
IRHeight = 228
|
||||
[R8PP] # Super Paper Mario (PAL)
|
||||
IRLeft = 399
|
||||
IRTop = 373
|
||||
IRWidth = 227
|
||||
IRHeight = 228
|
||||
[R4QP] # Mario Strikers (NTSC)
|
||||
IRLeft = 200
|
||||
IRTop = 54
|
||||
IRWidth = 615
|
||||
IRHeight = 657
|
||||
[R4QP] # Mario Strikers (PAL)
|
||||
IRLeft = 200
|
||||
IRTop = 54
|
||||
IRWidth = 615
|
||||
IRHeight = 657
|
||||
[RBUE] # Resident Evil - The Umbrella Chronicles (NTSC)
|
||||
IRLeft = 335
|
||||
IRTop = 351
|
||||
IRWidth = 357
|
||||
IRHeight = 273
|
||||
[RBUP] # Resident Evil - The Umbrella Chronicles (PAL)
|
||||
IRLeft = 335
|
||||
IRTop = 351
|
||||
IRWidth = 357
|
||||
IRHeight = 273
|
||||
[RB4E] # Resident Evil 4 (NTSC)
|
||||
IRLeft = 286
|
||||
IRTop = 256
|
||||
IRWidth = 450
|
||||
IRHeight = 455
|
||||
[RB4P] # Resident Evil 4 (PAL)
|
||||
IRLeft = 286
|
||||
IRTop = 256
|
||||
IRWidth = 450
|
||||
IRHeight = 455
|
||||
[R3IJ] # Metroid Prime - Wii De Asobu (JAP)
|
||||
IRLeft = 228
|
||||
IRTop = 112
|
||||
IRWidth = 486
|
||||
IRHeight = 577
|
||||
[RM3E] # Metroid Prime 3 (NTSC)
|
||||
IRLeft = 258
|
||||
IRTop = 84
|
||||
IRWidth = 489
|
||||
IRHeight = 613
|
||||
[RM3P] # Metroid Prime 3 (PAL)
|
||||
IRLeft = 258
|
||||
IRTop = 84
|
||||
IRWidth = 489
|
||||
IRHeight = 613
|
||||
[RSUP] # Sports Party (NTSC)
|
||||
IRLeft = 391
|
||||
IRTop = 377
|
||||
IRWidth = 241
|
||||
IRHeight = 225
|
||||
[RSUP] # Sports Party (PAL)
|
||||
IRLeft = 391
|
||||
IRTop = 377
|
||||
IRWidth = 241
|
||||
IRHeight = 225
|
||||
[RDZE] # Disaster - Day of Crisis (NTSC)
|
||||
IRLeft = 253
|
||||
IRTop = 276
|
||||
IRWidth = 453
|
||||
IRHeight = 421
|
||||
[RDZP] # Disaster - Day of Crisis (PAL)
|
||||
IRLeft = 253
|
||||
IRTop = 276
|
||||
IRWidth = 453
|
||||
IRHeight = 421
|
||||
[R4QE] # Mario Strikers (NTSC)
|
||||
IRLeft = 266
|
||||
IRTop = 215
|
||||
IRWidth = 486
|
||||
IRHeight = 490
|
||||
[R4QP] # Mario Strikers (PAL)
|
||||
IRLeft = 266
|
||||
IRTop = 215
|
||||
IRWidth = 486
|
||||
IRHeight = 490
|
||||
[RPBE] # Pokemon Battle Revolution (NTSC)
|
||||
IRLeft = 287
|
||||
IRTop = 261
|
||||
IRWidth = 451
|
||||
IRHeight = 449
|
||||
[RPBP] # Pokemon Battle Revolution (PAL)
|
||||
IRLeft = 287
|
||||
IRTop = 261
|
||||
IRWidth = 451
|
||||
IRHeight = 449
|
||||
[R2GJ] # Fragile: Sayonara Tsuki no Haikyo (JAP)
|
||||
IRLeft = 254
|
||||
IRTop = 280
|
||||
IRWidth = 451
|
||||
IRHeight = 453
|
||||
[RFNE] # Wii Fit (NTSC)
|
||||
[RFNP] # Wii Fit (PAL)
|
||||
[RSBE] # Super Smash Bros. Brawl (NTSC)
|
||||
[RSBP] # Super Smash Bros. Brawl (PAL)
|
||||
[R3TE] # Top Spin 3 (NTSC)
|
||||
[RSBP] # Top Spin 3 (PAL)
|
||||
[RLBE] # Lego Batman (NTSC)
|
||||
[RLBP] # Lego Batman (PAL)
|
||||
##########################################################################
|
||||
# These games don't use the IR Pointer at any time
|
||||
##########################################################################
|
||||
[RSRE] # Sonic and the Secret Rings (NTSC)
|
||||
[RSRP] # Sonic and the Secret Rings (PAL)
|
||||
[RWLE] # Wario Land - Shake It (NTSC)
|
||||
[RWLP] # Wario Land - Shake It (PAL)
|
||||
[RTNE] # Tenchu: Shadow Assassins (NTSC)
|
||||
[RTNE] # Tenchu: Shadow Assassins (PAL)
|
||||
|
||||
@ -1,41 +1,41 @@
|
||||
.text
|
||||
80003488 000000b8 80003488 0 __fill_mem
|
||||
80003540 00000050 80003540 0 memcpy
|
||||
8006cff8 0000004c 8006cff8 0 .LoadQuantizers
|
||||
8006d044 0000002c 8006d044 0 .kill_infinites_helper
|
||||
8006d070 00000018 8006d070 0 .kill_infinites
|
||||
8006d088 0000002c 8006d088 0 .rsqrt
|
||||
8006d0b4 00000034 8006d0b4 0 .sqrt_internal_fz
|
||||
8006d0e8 00000030 8006d0e8 0 .rsqrt_internal_fz
|
||||
8006d118 00000070 8006d118 0 .sqrt_fz
|
||||
8006d188 00000030 8006d188 0 .wrapping_once_fp_lookup
|
||||
8006d1b8 00000064 8006d1b8 0 .weird2
|
||||
8006d1c4 00000058 8006d1c4 0 .into_weird2
|
||||
8006d21c 00000030 8006d21c 0 .lookup_some_float_in_table_with_neg_wrap
|
||||
8006d24c 00000184 8006d24c 0 .atan2
|
||||
8006d3d0 0000009c 8006d3d0 0 .asin_fz
|
||||
8006d46c 000000c8 8006d46c 0 .acos_fz
|
||||
8006d534 00000070 8006d534 0 .evil_vec_cosine
|
||||
8006d5f0 00000078 8006d5f0 0 .evil_vec_setlength
|
||||
8006d668 00000094 8006d668 0 .evil_vec_something
|
||||
8006d6fc 0000005c 8006d6fc 0 .func
|
||||
8006d784 0000002c 8006d784 0 .load_strange_matrix1
|
||||
8006d7b0 0000002c 8006d7b0 0 .load_strange_matrix2
|
||||
8006d7f4 0000003c 8006d7f4 0 .some_strange_destination
|
||||
8006db30 00000044 8006db30 0 .push_matrix_3x3?
|
||||
8006db74 00000038 8006db74 0 .write_top_3x3_matrix
|
||||
8006dbe4 0000003c 8006dbe4 0 .read_current_3x3_matrix
|
||||
8006dc20 00000014 8006dc20 0 .pop_matrix_stack
|
||||
8006e424 00000074 8006e424 0 .weird_param_in_p1_p2
|
||||
8006e978 000001d4 8006e978 0 zz_006e978_
|
||||
8006eb4c 000001c0 8006eb4c 0 zz_006eb4c_
|
||||
8006f6a8 000000cc 8006f6a8 0 .z_last_skum_function
|
||||
800798f0 000000ec 800798f0 0 __div2u
|
||||
800799dc 00000138 800799dc 0 __div2i
|
||||
80079b14 000000e0 80079b14 0 __mod2u
|
||||
80079bf8 0000010c 80079bf8 0 __mod2i
|
||||
80079d04 00000024 80079d04 0 __shl2i
|
||||
80079d28 00000024 80079d28 0 __shr2u
|
||||
80079d4c 00000028 80079d4c 0 __shr2i
|
||||
8008596c 00000310 8008596c 0 big_matrix_trickery
|
||||
80088538 00000020 80088538 0 zz_0088538_
|
||||
.text
|
||||
80003488 000000b8 80003488 0 __fill_mem
|
||||
80003540 00000050 80003540 0 memcpy
|
||||
8006cff8 0000004c 8006cff8 0 .LoadQuantizers
|
||||
8006d044 0000002c 8006d044 0 .kill_infinites_helper
|
||||
8006d070 00000018 8006d070 0 .kill_infinites
|
||||
8006d088 0000002c 8006d088 0 .rsqrt
|
||||
8006d0b4 00000034 8006d0b4 0 .sqrt_internal_fz
|
||||
8006d0e8 00000030 8006d0e8 0 .rsqrt_internal_fz
|
||||
8006d118 00000070 8006d118 0 .sqrt_fz
|
||||
8006d188 00000030 8006d188 0 .wrapping_once_fp_lookup
|
||||
8006d1b8 00000064 8006d1b8 0 .weird2
|
||||
8006d1c4 00000058 8006d1c4 0 .into_weird2
|
||||
8006d21c 00000030 8006d21c 0 .lookup_some_float_in_table_with_neg_wrap
|
||||
8006d24c 00000184 8006d24c 0 .atan2
|
||||
8006d3d0 0000009c 8006d3d0 0 .asin_fz
|
||||
8006d46c 000000c8 8006d46c 0 .acos_fz
|
||||
8006d534 00000070 8006d534 0 .evil_vec_cosine
|
||||
8006d5f0 00000078 8006d5f0 0 .evil_vec_setlength
|
||||
8006d668 00000094 8006d668 0 .evil_vec_something
|
||||
8006d6fc 0000005c 8006d6fc 0 .func
|
||||
8006d784 0000002c 8006d784 0 .load_strange_matrix1
|
||||
8006d7b0 0000002c 8006d7b0 0 .load_strange_matrix2
|
||||
8006d7f4 0000003c 8006d7f4 0 .some_strange_destination
|
||||
8006db30 00000044 8006db30 0 .push_matrix_3x3?
|
||||
8006db74 00000038 8006db74 0 .write_top_3x3_matrix
|
||||
8006dbe4 0000003c 8006dbe4 0 .read_current_3x3_matrix
|
||||
8006dc20 00000014 8006dc20 0 .pop_matrix_stack
|
||||
8006e424 00000074 8006e424 0 .weird_param_in_p1_p2
|
||||
8006e978 000001d4 8006e978 0 zz_006e978_
|
||||
8006eb4c 000001c0 8006eb4c 0 zz_006eb4c_
|
||||
8006f6a8 000000cc 8006f6a8 0 .z_last_skum_function
|
||||
800798f0 000000ec 800798f0 0 __div2u
|
||||
800799dc 00000138 800799dc 0 __div2i
|
||||
80079b14 000000e0 80079b14 0 __mod2u
|
||||
80079bf8 0000010c 80079bf8 0 __mod2i
|
||||
80079d04 00000024 80079d04 0 __shl2i
|
||||
80079d28 00000024 80079d28 0 __shr2u
|
||||
80079d4c 00000028 80079d4c 0 __shr2i
|
||||
8008596c 00000310 8008596c 0 big_matrix_trickery
|
||||
80088538 00000020 80088538 0 zz_0088538_
|
||||
|
||||
@ -1,34 +1,34 @@
|
||||
.text
|
||||
800031f0 0000001c 800031f0 0 load_sp_rtoc
|
||||
80007034 0000004c 80007034 0 .LoadQuantizers
|
||||
80007080 00000030 80007080 0 .LoadInfinitiesEtc
|
||||
800070b0 00000038 800070b0 0 .rsqrt
|
||||
800070ec 00000040 800070ec 0 .sqrt_internal_needs_cr1
|
||||
8000712c 00000040 8000712c 0 .rsqrt_internal_needs_cr1
|
||||
800071e0 00000030 800071e0 0 .wrapping_once_fp_lookup
|
||||
80007210 00000064 80007210 0 .weird2
|
||||
80007274 00000030 80007274 0 .lookup_some_float_in_table_with_neg_wrap
|
||||
800072a4 00000180 800072a4 0 .atan2
|
||||
80007424 000000b8 80007424 0 .calls_sqrt
|
||||
800074dc 0000005c 800074dc 0 .func
|
||||
80007538 0000002c 80007538 0 .load_strange_matrix1
|
||||
80007564 0000002c 80007564 0 .load_strange_matrix3
|
||||
80007590 0000002c 80007590 0 .load_strange_matrix2
|
||||
80007834 00000044 80007834 0 .push_matrix_3x3?
|
||||
80007878 00000038 80007878 0 .read_top_3x3matrix
|
||||
800078b0 00000038 800078b0 0 .write_top_3x3_matrix
|
||||
800078e8 0000003c 800078e8 0 .read_current_3x3_matrix
|
||||
80007924 00000014 80007924 0 .pop_matrix_stack
|
||||
80007a50 00000170 80007a50 0 .mult_matrix?
|
||||
80007ecc 000000bc 80007ecc 0 .weird_vector_op_status_in_cr2
|
||||
80007f88 00000074 80007f88 0 .weird_param_in_p1_p2
|
||||
800080fc 00000078 800080fc 0 .evil_normalize
|
||||
80008174 00000078 80008174 0 .evil_vec_setlength
|
||||
800081ec 00000070 800081ec 0 .evil_vec_cosine
|
||||
80008538 000000f0 80008538 0 .calls_evil1
|
||||
8000875c 00000088 8000875c 0 .another_caller
|
||||
800087e4 0000008c 800087e4 0 .another_caller2
|
||||
80008d30 000001b4 80008d30 0 .another_caller3
|
||||
80036544 000001b4 80036544 0 .fctiwi_weird2
|
||||
8003dd1c 00000110 8003dd1c 0 .fctwi_weird
|
||||
80043b48 000005bc 80043b48 0 .fctwi_weird3
|
||||
.text
|
||||
800031f0 0000001c 800031f0 0 load_sp_rtoc
|
||||
80007034 0000004c 80007034 0 .LoadQuantizers
|
||||
80007080 00000030 80007080 0 .LoadInfinitiesEtc
|
||||
800070b0 00000038 800070b0 0 .rsqrt
|
||||
800070ec 00000040 800070ec 0 .sqrt_internal_needs_cr1
|
||||
8000712c 00000040 8000712c 0 .rsqrt_internal_needs_cr1
|
||||
800071e0 00000030 800071e0 0 .wrapping_once_fp_lookup
|
||||
80007210 00000064 80007210 0 .weird2
|
||||
80007274 00000030 80007274 0 .lookup_some_float_in_table_with_neg_wrap
|
||||
800072a4 00000180 800072a4 0 .atan2
|
||||
80007424 000000b8 80007424 0 .calls_sqrt
|
||||
800074dc 0000005c 800074dc 0 .func
|
||||
80007538 0000002c 80007538 0 .load_strange_matrix1
|
||||
80007564 0000002c 80007564 0 .load_strange_matrix3
|
||||
80007590 0000002c 80007590 0 .load_strange_matrix2
|
||||
80007834 00000044 80007834 0 .push_matrix_3x3?
|
||||
80007878 00000038 80007878 0 .read_top_3x3matrix
|
||||
800078b0 00000038 800078b0 0 .write_top_3x3_matrix
|
||||
800078e8 0000003c 800078e8 0 .read_current_3x3_matrix
|
||||
80007924 00000014 80007924 0 .pop_matrix_stack
|
||||
80007a50 00000170 80007a50 0 .mult_matrix?
|
||||
80007ecc 000000bc 80007ecc 0 .weird_vector_op_status_in_cr2
|
||||
80007f88 00000074 80007f88 0 .weird_param_in_p1_p2
|
||||
800080fc 00000078 800080fc 0 .evil_normalize
|
||||
80008174 00000078 80008174 0 .evil_vec_setlength
|
||||
800081ec 00000070 800081ec 0 .evil_vec_cosine
|
||||
80008538 000000f0 80008538 0 .calls_evil1
|
||||
8000875c 00000088 8000875c 0 .another_caller
|
||||
800087e4 0000008c 800087e4 0 .another_caller2
|
||||
80008d30 000001b4 80008d30 0 .another_caller3
|
||||
80036544 000001b4 80036544 0 .fctiwi_weird2
|
||||
8003dd1c 00000110 8003dd1c 0 .fctwi_weird
|
||||
80043b48 000005bc 80043b48 0 .fctwi_weird3
|
||||
|
||||
@ -1,224 +1,224 @@
|
||||
// Copyright (C) 2003 Dolphin Project.
|
||||
|
||||
// This program is free software: you can redistribute it and/or modify
|
||||
// it under the terms of the GNU General Public License as published by
|
||||
// the Free Software Foundation, version 2.0.
|
||||
|
||||
// This program is distributed in the hope that it will be useful,
|
||||
// but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
// GNU General Public License 2.0 for more details.
|
||||
|
||||
// A copy of the GPL 2.0 should have been included with the program.
|
||||
// If not, see http://www.gnu.org/licenses/
|
||||
|
||||
// Official SVN repository and contact information can be found at
|
||||
// http://code.google.com/p/dolphin-emu/
|
||||
|
||||
kernel void DecodeI4(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 8, y = get_global_id(1) * 8;
|
||||
int srcOffset = x + y * width / 8;
|
||||
for (int iy = 0; iy < 8; iy++)
|
||||
{
|
||||
uchar4 val = vload4(srcOffset, src);
|
||||
uchar8 res;
|
||||
res.even = (val >> (uchar4)4) & (uchar4)0x0F;
|
||||
res.odd = val & (uchar4)0x0F;
|
||||
res |= res << (uchar8)4;
|
||||
vstore8(res, 0, dst + ((y + iy)*width + x));
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeI8(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 8, y = get_global_id(1) * 4;
|
||||
int srcOffset = ((x * 4) + (y * width)) / 8;
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
vstore8(vload8(srcOffset, src),
|
||||
0, dst + ((y + iy)*width + x));
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeIA8(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 4, y = get_global_id(1) * 4;
|
||||
int srcOffset = ((x * 4) + (y * width)) / 4;
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
uchar8 val = vload8(srcOffset, src);
|
||||
uchar8 res;
|
||||
res.odd = val.even;
|
||||
res.even = val.odd;
|
||||
vstore8(res, 0, dst + ((y + iy)*width + x) * 2);
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeIA4(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 8, y = get_global_id(1) * 4;
|
||||
int srcOffset = ((x * 4) + (y * width)) / 8;
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
uchar8 val = vload8(srcOffset, src);
|
||||
uchar16 res;
|
||||
res.odd = (val >> (uchar8)4) & (uchar8)0x0F;
|
||||
res.even = val & (uchar8)0x0F;
|
||||
res |= res << (uchar16)4;
|
||||
vstore16(res, 0, dst + ((y + iy)*width + x) * 2);
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeRGBA8(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 4, y = get_global_id(1) * 4;
|
||||
int srcOffset = (x * 2) + (y * width) / 2;
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
uchar8 ar = vload8(srcOffset, src);
|
||||
uchar8 gb = vload8(srcOffset + 4, src);
|
||||
uchar16 res;
|
||||
res.even.even = gb.odd;
|
||||
res.even.odd = ar.odd;
|
||||
res.odd.even = gb.even;
|
||||
res.odd.odd = ar.even;
|
||||
vstore16(res, 0, dst + ((y + iy)*width + x) * 4);
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeRGB565(global ushort *dst,
|
||||
const global ushort *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 4, y = get_global_id(1) * 4;
|
||||
int srcOffset = x + (y * width) / 4;
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
ushort4 val = vload4(srcOffset, src);
|
||||
val = (val >> (ushort4)8) | (val << (ushort4)8);
|
||||
vstore4(val, 0, dst + ((y + iy)*width + x));
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeRGB5A3(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 4, y = get_global_id(1) * 4;
|
||||
int srcOffset = x + (y * width) / 4;
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
ushort8 val = convert_ushort8(vload8(srcOffset, src));
|
||||
ushort4 vs = val.odd | (ushort4)(val.even << (ushort4)8);
|
||||
|
||||
uchar16 resNoAlpha;
|
||||
resNoAlpha.s26AE = convert_uchar4(vs >> (ushort4)7); // R
|
||||
resNoAlpha.s159D = convert_uchar4(vs >> (ushort4)2); // G
|
||||
resNoAlpha.s048C = convert_uchar4(vs << (ushort4)3); // B
|
||||
resNoAlpha &= (uchar16)0xF8;
|
||||
resNoAlpha |= (uchar16)(resNoAlpha >> (uchar16)5) & (uchar16)3; // 5 -> 8
|
||||
resNoAlpha.s37BF = (uchar4)(0xFF);
|
||||
|
||||
uchar16 resAlpha;
|
||||
resAlpha.s26AE = convert_uchar4(vs >> (ushort4)8); // R
|
||||
resAlpha.s159D = convert_uchar4(vs >> (ushort4)4); // G
|
||||
resAlpha.s048C = convert_uchar4(vs); // B
|
||||
resAlpha &= (uchar16)0x0F;
|
||||
resAlpha |= (resAlpha << (uchar16)4);
|
||||
resAlpha.s37BF = convert_uchar4(vs >> (ushort4)7) & (uchar4)0xE0;
|
||||
resAlpha.s37BF |= ((resAlpha.s37BF >> (uchar4)3) & (uchar4)0x1C)
|
||||
| ((resAlpha.s37BF >> (uchar4)6) & (uchar4)0x3);
|
||||
uchar16 choice = (uchar16)((uchar4)(vs.s0 >> 8),
|
||||
(uchar4)(vs.s1 >> 8),
|
||||
(uchar4)(vs.s2 >> 8),
|
||||
(uchar4)(vs.s3 >> 8));
|
||||
uchar16 res;
|
||||
res = select(resAlpha, resNoAlpha, choice);
|
||||
vstore16(res, 0, dst + ((y + iy) * width + x) * 4);
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
uint4 unpack2bits(uchar b)
|
||||
{
|
||||
return (uint4)(b >> 6,
|
||||
(b >> 4) & 3,
|
||||
(b >> 2) & 3,
|
||||
b & 3);
|
||||
}
|
||||
|
||||
kernel void decodeCMPRBlock(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 4, y = get_global_id(1) * 4;
|
||||
uchar8 val = vload8(0, src);
|
||||
ushort2 color565 = (ushort2)((val.s1 & 0xFF) | (val.s0 << 8), (val.s3 & 0xFF) | (val.s2 << 8));
|
||||
uchar8 color32 = convert_uchar8((ushort8)
|
||||
(((color565 << (ushort2)3) & (ushort2)0xF8) | ((color565 >> (ushort2)2) & (ushort2)0x7),
|
||||
((color565 >> (ushort2)3) & (ushort2)0xFC) | ((color565 >> (ushort2)9) & (ushort2)0x3),
|
||||
((color565 >> (ushort2)8) & (ushort2)0xF8) | ((color565 >> (ushort2)13) & (ushort2)0x7),
|
||||
0xFF, 0xFF));
|
||||
uint4 colors;
|
||||
uint4 colorNoAlpha;
|
||||
uchar4 frac = convert_uchar4((((convert_ushort4(color32.even) & (ushort4)0xFF) - (convert_ushort4(color32.odd) & (ushort4)0xFF)) * (ushort4)3) / (ushort4)8);
|
||||
colorNoAlpha = convert_uint4(color32.odd + frac);
|
||||
colorNoAlpha = (colorNoAlpha << (uint4)8) | convert_uint4(color32.even - frac);
|
||||
colorNoAlpha = (colorNoAlpha << (uint4)8) | convert_uint4(color32.odd);
|
||||
colorNoAlpha = (colorNoAlpha << (uint4)8) | convert_uint4(color32.even);
|
||||
|
||||
uint4 colorAlpha;
|
||||
uchar4 midpoint = convert_uchar4((convert_ushort4(color32.odd) + convert_ushort4(color32.even) + (ushort4)1) / (ushort4)2);
|
||||
midpoint.s3 = 0xFF;
|
||||
colorAlpha = convert_uint4((uchar4)(0, 0, 0, 0));
|
||||
colorAlpha = (colorAlpha << (uint4)8) | convert_uint4(midpoint);
|
||||
colorAlpha = (colorAlpha << (uint4)8) | convert_uint4(color32.odd);
|
||||
colorAlpha = (colorAlpha << (uint4)8) | convert_uint4(color32.even);
|
||||
|
||||
colors = color565.s0 > color565.s1 ? colorNoAlpha : colorAlpha;
|
||||
|
||||
uint16 colorsFull = (uint16)(colors, colors, colors, colors);
|
||||
|
||||
uint4 shift0 = unpack2bits(val.s4);
|
||||
uint4 shift1 = unpack2bits(val.s5);
|
||||
uint4 shift2 = unpack2bits(val.s6);
|
||||
uint4 shift3 = unpack2bits(val.s7);
|
||||
uint16 shifts = (uint16)((uint4)(shift3.s0), (uint4)(shift3.s1), (uint4)(shift3.s2), (uint4)(shift3.s3));
|
||||
shifts = (shifts << (uint16)8) | (uint16)((uint4)(shift2.s0), (uint4)(shift2.s1), (uint4)(shift2.s2), (uint4)(shift2.s3));
|
||||
shifts = (shifts << (uint16)8) | (uint16)((uint4)(shift1.s0), (uint4)(shift1.s1), (uint4)(shift1.s2), (uint4)(shift1.s3));
|
||||
shifts = (shifts << (uint16)8) | (uint16)((uint4)(shift0.s0), (uint4)(shift0.s1), (uint4)(shift0.s2), (uint4)(shift0.s3)) << (uint16)3;
|
||||
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
uchar16 res;
|
||||
res = convert_uchar16(colorsFull >> (shifts & (uint16)0xFF)) >> (uchar16)8;
|
||||
vstore16(res, 0, dst);
|
||||
dst += width * 4;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeCMPR(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 8, y = get_global_id(1) * 8;
|
||||
|
||||
src += x * 4 + (y * width) / 2;
|
||||
|
||||
decodeCMPRBlock(dst + (y * width + x) * 4, src, width);
|
||||
src += 8;
|
||||
decodeCMPRBlock(dst + (y * width + x + 4) * 4, src, width);
|
||||
src += 8;
|
||||
decodeCMPRBlock(dst + ((y + 4) * width + x) * 4, src, width);
|
||||
src += 8;
|
||||
decodeCMPRBlock(dst + ((y + 4) * width + x + 4) * 4, src, width);
|
||||
|
||||
// Copyright (C) 2003 Dolphin Project.
|
||||
|
||||
// This program is free software: you can redistribute it and/or modify
|
||||
// it under the terms of the GNU General Public License as published by
|
||||
// the Free Software Foundation, version 2.0.
|
||||
|
||||
// This program is distributed in the hope that it will be useful,
|
||||
// but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
// GNU General Public License 2.0 for more details.
|
||||
|
||||
// A copy of the GPL 2.0 should have been included with the program.
|
||||
// If not, see http://www.gnu.org/licenses/
|
||||
|
||||
// Official SVN repository and contact information can be found at
|
||||
// http://code.google.com/p/dolphin-emu/
|
||||
|
||||
kernel void DecodeI4(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 8, y = get_global_id(1) * 8;
|
||||
int srcOffset = x + y * width / 8;
|
||||
for (int iy = 0; iy < 8; iy++)
|
||||
{
|
||||
uchar4 val = vload4(srcOffset, src);
|
||||
uchar8 res;
|
||||
res.even = (val >> (uchar4)4) & (uchar4)0x0F;
|
||||
res.odd = val & (uchar4)0x0F;
|
||||
res |= res << (uchar8)4;
|
||||
vstore8(res, 0, dst + ((y + iy)*width + x));
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeI8(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 8, y = get_global_id(1) * 4;
|
||||
int srcOffset = ((x * 4) + (y * width)) / 8;
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
vstore8(vload8(srcOffset, src),
|
||||
0, dst + ((y + iy)*width + x));
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeIA8(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 4, y = get_global_id(1) * 4;
|
||||
int srcOffset = ((x * 4) + (y * width)) / 4;
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
uchar8 val = vload8(srcOffset, src);
|
||||
uchar8 res;
|
||||
res.odd = val.even;
|
||||
res.even = val.odd;
|
||||
vstore8(res, 0, dst + ((y + iy)*width + x) * 2);
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeIA4(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 8, y = get_global_id(1) * 4;
|
||||
int srcOffset = ((x * 4) + (y * width)) / 8;
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
uchar8 val = vload8(srcOffset, src);
|
||||
uchar16 res;
|
||||
res.odd = (val >> (uchar8)4) & (uchar8)0x0F;
|
||||
res.even = val & (uchar8)0x0F;
|
||||
res |= res << (uchar16)4;
|
||||
vstore16(res, 0, dst + ((y + iy)*width + x) * 2);
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeRGBA8(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 4, y = get_global_id(1) * 4;
|
||||
int srcOffset = (x * 2) + (y * width) / 2;
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
uchar8 ar = vload8(srcOffset, src);
|
||||
uchar8 gb = vload8(srcOffset + 4, src);
|
||||
uchar16 res;
|
||||
res.even.even = gb.odd;
|
||||
res.even.odd = ar.odd;
|
||||
res.odd.even = gb.even;
|
||||
res.odd.odd = ar.even;
|
||||
vstore16(res, 0, dst + ((y + iy)*width + x) * 4);
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeRGB565(global ushort *dst,
|
||||
const global ushort *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 4, y = get_global_id(1) * 4;
|
||||
int srcOffset = x + (y * width) / 4;
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
ushort4 val = vload4(srcOffset, src);
|
||||
val = (val >> (ushort4)8) | (val << (ushort4)8);
|
||||
vstore4(val, 0, dst + ((y + iy)*width + x));
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeRGB5A3(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 4, y = get_global_id(1) * 4;
|
||||
int srcOffset = x + (y * width) / 4;
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
ushort8 val = convert_ushort8(vload8(srcOffset, src));
|
||||
ushort4 vs = val.odd | (ushort4)(val.even << (ushort4)8);
|
||||
|
||||
uchar16 resNoAlpha;
|
||||
resNoAlpha.s26AE = convert_uchar4(vs >> (ushort4)7); // R
|
||||
resNoAlpha.s159D = convert_uchar4(vs >> (ushort4)2); // G
|
||||
resNoAlpha.s048C = convert_uchar4(vs << (ushort4)3); // B
|
||||
resNoAlpha &= (uchar16)0xF8;
|
||||
resNoAlpha |= (uchar16)(resNoAlpha >> (uchar16)5) & (uchar16)3; // 5 -> 8
|
||||
resNoAlpha.s37BF = (uchar4)(0xFF);
|
||||
|
||||
uchar16 resAlpha;
|
||||
resAlpha.s26AE = convert_uchar4(vs >> (ushort4)8); // R
|
||||
resAlpha.s159D = convert_uchar4(vs >> (ushort4)4); // G
|
||||
resAlpha.s048C = convert_uchar4(vs); // B
|
||||
resAlpha &= (uchar16)0x0F;
|
||||
resAlpha |= (resAlpha << (uchar16)4);
|
||||
resAlpha.s37BF = convert_uchar4(vs >> (ushort4)7) & (uchar4)0xE0;
|
||||
resAlpha.s37BF |= ((resAlpha.s37BF >> (uchar4)3) & (uchar4)0x1C)
|
||||
| ((resAlpha.s37BF >> (uchar4)6) & (uchar4)0x3);
|
||||
uchar16 choice = (uchar16)((uchar4)(vs.s0 >> 8),
|
||||
(uchar4)(vs.s1 >> 8),
|
||||
(uchar4)(vs.s2 >> 8),
|
||||
(uchar4)(vs.s3 >> 8));
|
||||
uchar16 res;
|
||||
res = select(resAlpha, resNoAlpha, choice);
|
||||
vstore16(res, 0, dst + ((y + iy) * width + x) * 4);
|
||||
srcOffset++;
|
||||
}
|
||||
}
|
||||
|
||||
uint4 unpack2bits(uchar b)
|
||||
{
|
||||
return (uint4)(b >> 6,
|
||||
(b >> 4) & 3,
|
||||
(b >> 2) & 3,
|
||||
b & 3);
|
||||
}
|
||||
|
||||
kernel void decodeCMPRBlock(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 4, y = get_global_id(1) * 4;
|
||||
uchar8 val = vload8(0, src);
|
||||
ushort2 color565 = (ushort2)((val.s1 & 0xFF) | (val.s0 << 8), (val.s3 & 0xFF) | (val.s2 << 8));
|
||||
uchar8 color32 = convert_uchar8((ushort8)
|
||||
(((color565 << (ushort2)3) & (ushort2)0xF8) | ((color565 >> (ushort2)2) & (ushort2)0x7),
|
||||
((color565 >> (ushort2)3) & (ushort2)0xFC) | ((color565 >> (ushort2)9) & (ushort2)0x3),
|
||||
((color565 >> (ushort2)8) & (ushort2)0xF8) | ((color565 >> (ushort2)13) & (ushort2)0x7),
|
||||
0xFF, 0xFF));
|
||||
uint4 colors;
|
||||
uint4 colorNoAlpha;
|
||||
uchar4 frac = convert_uchar4((((convert_ushort4(color32.even) & (ushort4)0xFF) - (convert_ushort4(color32.odd) & (ushort4)0xFF)) * (ushort4)3) / (ushort4)8);
|
||||
colorNoAlpha = convert_uint4(color32.odd + frac);
|
||||
colorNoAlpha = (colorNoAlpha << (uint4)8) | convert_uint4(color32.even - frac);
|
||||
colorNoAlpha = (colorNoAlpha << (uint4)8) | convert_uint4(color32.odd);
|
||||
colorNoAlpha = (colorNoAlpha << (uint4)8) | convert_uint4(color32.even);
|
||||
|
||||
uint4 colorAlpha;
|
||||
uchar4 midpoint = convert_uchar4((convert_ushort4(color32.odd) + convert_ushort4(color32.even) + (ushort4)1) / (ushort4)2);
|
||||
midpoint.s3 = 0xFF;
|
||||
colorAlpha = convert_uint4((uchar4)(0, 0, 0, 0));
|
||||
colorAlpha = (colorAlpha << (uint4)8) | convert_uint4(midpoint);
|
||||
colorAlpha = (colorAlpha << (uint4)8) | convert_uint4(color32.odd);
|
||||
colorAlpha = (colorAlpha << (uint4)8) | convert_uint4(color32.even);
|
||||
|
||||
colors = color565.s0 > color565.s1 ? colorNoAlpha : colorAlpha;
|
||||
|
||||
uint16 colorsFull = (uint16)(colors, colors, colors, colors);
|
||||
|
||||
uint4 shift0 = unpack2bits(val.s4);
|
||||
uint4 shift1 = unpack2bits(val.s5);
|
||||
uint4 shift2 = unpack2bits(val.s6);
|
||||
uint4 shift3 = unpack2bits(val.s7);
|
||||
uint16 shifts = (uint16)((uint4)(shift3.s0), (uint4)(shift3.s1), (uint4)(shift3.s2), (uint4)(shift3.s3));
|
||||
shifts = (shifts << (uint16)8) | (uint16)((uint4)(shift2.s0), (uint4)(shift2.s1), (uint4)(shift2.s2), (uint4)(shift2.s3));
|
||||
shifts = (shifts << (uint16)8) | (uint16)((uint4)(shift1.s0), (uint4)(shift1.s1), (uint4)(shift1.s2), (uint4)(shift1.s3));
|
||||
shifts = (shifts << (uint16)8) | (uint16)((uint4)(shift0.s0), (uint4)(shift0.s1), (uint4)(shift0.s2), (uint4)(shift0.s3)) << (uint16)3;
|
||||
|
||||
for (int iy = 0; iy < 4; iy++)
|
||||
{
|
||||
uchar16 res;
|
||||
res = convert_uchar16(colorsFull >> (shifts & (uint16)0xFF)) >> (uchar16)8;
|
||||
vstore16(res, 0, dst);
|
||||
dst += width * 4;
|
||||
}
|
||||
}
|
||||
|
||||
kernel void DecodeCMPR(global uchar *dst,
|
||||
const global uchar *src, int width)
|
||||
{
|
||||
int x = get_global_id(0) * 8, y = get_global_id(1) * 8;
|
||||
|
||||
src += x * 4 + (y * width) / 2;
|
||||
|
||||
decodeCMPRBlock(dst + (y * width + x) * 4, src, width);
|
||||
src += 8;
|
||||
decodeCMPRBlock(dst + (y * width + x + 4) * 4, src, width);
|
||||
src += 8;
|
||||
decodeCMPRBlock(dst + ((y + 4) * width + x) * 4, src, width);
|
||||
src += 8;
|
||||
decodeCMPRBlock(dst + ((y + 4) * width + x + 4) * 4, src, width);
|
||||
|
||||
}
|
||||
Reference in New Issue
Block a user