claunia/cuetools.net
1
0
Fork 0
mirror of https://github.com/claunia/cuetools.net.git synced 2025-12-16 18:14:25 +00:00
Code Issues Packages Projects Releases Wiki Activity

optimizations

Browse Source
This commit is contained in:
chudov
2009-09-24 21:34:34 +00:00
parent d8957f6a3f
commit 693c2fbf41
3 changed files with 156 additions and 163 deletions
Download Patch File Download Diff File Expand all files Collapse all files

14
CUETools.FlaCuda/FlaCudaWriter.cs
View File

@@ -1099,17 +1099,14 @@ namespace CUETools.Codecs.FlaCuda
unsafe void compute_autocorellation(int blocksize, int channelsCount, int max_order, int nFrames, FlaCudaTask task) unsafe void compute_autocorellation(int blocksize, int channelsCount, int max_order, int nFrames, FlaCudaTask task)
{ {
int autocorThreads = 256; if (blocksize <= 4)
int partSize = 2 * autocorThreads - max_order; return;
partSize &= 0xffffff0;
int partSize = 256 + 128;// (2 * 256 - max_order) & ~31;
int partCount = (blocksize + partSize - 1) / partSize; int partCount = (blocksize + partSize - 1) / partSize;
if (partCount > maxAutocorParts) if (partCount > maxAutocorParts)
throw new Exception("internal error"); throw new Exception("internal error");
if (blocksize <= 4)
return;
cuda.SetParameter(task.cudaStereoDecorr, 0, (uint)task.cudaSamples.Pointer); cuda.SetParameter(task.cudaStereoDecorr, 0, (uint)task.cudaSamples.Pointer);
cuda.SetParameter(task.cudaStereoDecorr, sizeof(uint), (uint)MAX_BLOCKSIZE); cuda.SetParameter(task.cudaStereoDecorr, sizeof(uint), (uint)MAX_BLOCKSIZE);
cuda.SetParameterSize(task.cudaStereoDecorr, sizeof(uint) * 2U); cuda.SetParameterSize(task.cudaStereoDecorr, sizeof(uint) * 2U);
@@ -1129,8 +1126,8 @@ namespace CUETools.Codecs.FlaCuda
cuda.SetParameter(task.cudaComputeAutocor, sizeof(uint) * 4, (uint)max_order); cuda.SetParameter(task.cudaComputeAutocor, sizeof(uint) * 4, (uint)max_order);
cuda.SetParameter(task.cudaComputeAutocor, sizeof(uint) * 4 + sizeof(uint), (uint)blocksize); cuda.SetParameter(task.cudaComputeAutocor, sizeof(uint) * 4 + sizeof(uint), (uint)blocksize);
cuda.SetParameter(task.cudaComputeAutocor, sizeof(uint) * 4 + sizeof(uint) * 2, (uint)partSize); cuda.SetParameter(task.cudaComputeAutocor, sizeof(uint) * 4 + sizeof(uint) * 2, (uint)partSize);
cuda.SetParameterSize(task.cudaComputeAutocor, (uint)(sizeof(uint) * 4) + sizeof(uint) * 3); cuda.SetParameterSize(task.cudaComputeAutocor, sizeof(uint) * 7U);
cuda.SetFunctionBlockShape(task.cudaComputeAutocor, autocorThreads, 1, 1); cuda.SetFunctionBlockShape(task.cudaComputeAutocor, 32, 8, 1);
cuda.SetParameter(task.cudaComputeLPC, 0, (uint)task.cudaResidualTasks.Pointer); cuda.SetParameter(task.cudaComputeLPC, 0, (uint)task.cudaResidualTasks.Pointer);
cuda.SetParameter(task.cudaComputeLPC, sizeof(uint), (uint)task.cudaAutocorOutput.Pointer); cuda.SetParameter(task.cudaComputeLPC, sizeof(uint), (uint)task.cudaAutocorOutput.Pointer);
@@ -1768,6 +1765,7 @@ namespace CUETools.Codecs.FlaCuda
public int windowOffs; public int windowOffs;
public int residualOffs; public int residualOffs;
public int blocksize; public int blocksize;
public fixed int reserved[12];
}; };
unsafe struct encodeResidualTaskStruct unsafe struct encodeResidualTaskStruct

50
CUETools.FlaCuda/flacuda.cu
View File

@@ -26,6 +26,7 @@ typedef struct
int windowOffs; int windowOffs;
int residualOffs; int residualOffs;
int blocksize; int blocksize;
int reserved[12];
} computeAutocorTaskStruct; } computeAutocorTaskStruct;
typedef enum typedef enum
@@ -115,46 +116,41 @@ extern "C" __global__ void cudaComputeAutocor(
__shared__ struct { __shared__ struct {
float data[512]; float data[512];
volatile float product[256]; volatile float product[256];
volatile float sum[33];
computeAutocorTaskStruct task; computeAutocorTaskStruct task;
} shared; } shared;
const int tid = threadIdx.x; const int tid = threadIdx.x + (threadIdx.y * 32);
const int tid2 = threadIdx.x + 256;
// fetch task data // fetch task data
if (tid < sizeof(shared.task) / sizeof(int)) if (tid < sizeof(shared.task) / sizeof(int))
((int*)&shared.task)[tid] = ((int*)(tasks + blockIdx.y))[tid]; ((int*)&shared.task)[tid] = ((int*)(tasks + blockIdx.y))[tid];
__syncthreads(); __syncthreads();
const int pos = blockIdx.x * partSize;
const int productLen = min(frameSize - pos - max_order, partSize);
const int dataLen = productLen + max_order;
// fetch samples // fetch samples
shared.data[tid] = tid < dataLen ? samples[shared.task.samplesOffs + pos + tid] * window[shared.task.windowOffs + pos + tid]: 0.0f; {
shared.data[tid2] = tid2 < dataLen ? samples[shared.task.samplesOffs + pos + tid2] * window[shared.task.windowOffs + pos + tid2]: 0.0f; const int pos = blockIdx.x * partSize;
const int dataLen = min(frameSize - pos, partSize + max_order);
shared.data[tid] = tid < dataLen ? samples[shared.task.samplesOffs + pos + tid] * window[shared.task.windowOffs + pos + tid]: 0.0f;
shared.data[tid + 256] = tid + 256 < dataLen ? samples[shared.task.samplesOffs + pos + tid + 256] * window[shared.task.windowOffs + pos + tid + 256]: 0.0f;
}
__syncthreads(); __syncthreads();
for (int lag = 0; lag <= max_order; lag++) for (int lag = threadIdx.y; lag <= max_order; lag += 8)
{ {
shared.product[tid] = (tid < productLen) * shared.data[tid] * shared.data[tid + lag] + const int productLen = min(frameSize - blockIdx.x * partSize - lag, partSize);
+ (tid2 < productLen) * shared.data[tid2] * shared.data[tid2 + lag]; shared.product[tid] = 0.0;
__syncthreads(); for (int ptr = threadIdx.x; ptr < productLen + threadIdx.x; ptr += 128)
shared.product[tid] += ((ptr < productLen) * shared.data[ptr] * shared.data[ptr + lag]
+ (ptr + 32 < productLen) * shared.data[ptr + 32] * shared.data[ptr + 32 + lag])
+ ((ptr + 64 < productLen) * shared.data[ptr + 64] * shared.data[ptr + 64 + lag]
+ (ptr + 96 < productLen) * shared.data[ptr + 96] * shared.data[ptr + 96 + lag]);
// product sum: reduction in shared mem // product sum: reduction in shared mem
//if (tid < 256) shared.product[tid] += shared.product[tid + 256]; __syncthreads(); //shared.product[tid] += shared.product[tid + 16];
if (tid < 128) shared.product[tid] += shared.product[tid + 128]; __syncthreads(); shared.product[tid] = (shared.product[tid] + shared.product[tid + 16]) + (shared.product[tid + 8] + shared.product[tid + 24]);
if (tid < 64) shared.product[tid] += shared.product[tid + 64]; __syncthreads(); shared.product[tid] = (shared.product[tid] + shared.product[tid + 4]) + (shared.product[tid + 2] + shared.product[tid + 6]);
if (tid < 32) shared.product[tid] += shared.product[tid + 32]; __syncthreads(); // return results
shared.product[tid] += shared.product[tid + 16]; if (threadIdx.x == 0)
shared.product[tid] += shared.product[tid + 8]; output[(blockIdx.x + blockIdx.y * gridDim.x) * (max_order + 1) + lag] = shared.product[tid] + shared.product[tid + 1];
shared.product[tid] += shared.product[tid + 4];
shared.product[tid] += shared.product[tid + 2];
if (tid == 0) shared.sum[lag] = shared.product[0] + shared.product[1];
__syncthreads();
} }
// return results
if (tid <= max_order)
output[(blockIdx.x + blockIdx.y * gridDim.x) * (max_order + 1) + tid] = shared.sum[tid];
} }
extern "C" __global__ void cudaComputeLPC( extern "C" __global__ void cudaComputeLPC(

255
CUETools.FlaCuda/flacuda.cubin
View File

@@ -4,93 +4,92 @@ modname {cubin}
code { code {
name = cudaComputeAutocor name = cudaComputeAutocor
lmem = 0 lmem = 0
smem = 3264 smem = 3180
reg = 10 reg = 10
bar = 1 bar = 1
const { const {
segname = const segname = const
segnum = 1 segnum = 1
offset = 0 offset = 0
bytes = 16 bytes = 8
mem { mem {
0x00000003 0x0000007f 0x0000003f 0x0000001f 0x000003ff 0x0000000f
} }
} }
bincode { bincode {
0xa0000009 0x04000780 0x308005fd 0x644107c8 0xd0800205 0x00400780 0xa0000219 0x04000780
0xa000b003 0x00000000 0x3002040d 0xc4100780 0xa0000011 0x04000780 0x30050c01 0xc4100780
0x1000b003 0x00000280 0xa0004e01 0x04200780 0x20000815 0x04000780 0x30810bfd 0x644107c8
0x30040001 0xc4100780 0x2100ee00 0x20008600 0xa000f003 0x00000000 0x30020a0d 0xc4100780
0x1000f003 0x00000280 0xa0004e01 0x04200780
0x30060001 0xc4100780 0x2100ee00 0x20008600
0xd00e0001 0x80c00780 0x00000605 0xc0000780 0xd00e0001 0x80c00780 0x00000605 0xc0000780
0x04065801 0xe4200780 0xf0000001 0xe0000002 0x04061601 0xe4200780 0xf0000001 0xe0000002
0x861ffe03 0x00000000 0xa0004c05 0x04200780 0x861ffe03 0x00000000 0xa0004c01 0x04200780
0x1000d401 0x0423c780 0x40020211 0x00000780 0x1000d405 0x0423c780 0x40000609 0x00000780
0x30100811 0xc4100780 0x60020001 0x00010780 0x30100409 0xc4100780 0x60000421 0x00008780
0x2140f210 0x3104f010 0x3004d411 0xac200780 0x1100f400 0x2100f000 0x2040d209 0x04220780
0x2000d015 0x04210780 0x30020bfd 0x6c00c7c8 0x3000041d 0xac000780 0x30050ffd 0x6c00c7c8
0xa0024003 0x00000000 0x10023003 0x00000280 0xa0028003 0x00000000 0x10027003 0x00000280
0xd0196005 0x20000780 0x2500e018 0x2500e21c 0xd0185805 0x20000780 0x2508e004 0x2508e200
0x20068418 0x2007841c 0x30020c19 0xc4100780 0x20018a04 0x20008a00 0x30020205 0xc4100780
0x30020e1d 0xc4100780 0x2106ea18 0x2107ec20 0x30020001 0xc4100780 0x2101ea04 0x2100ec00
0xd00e0c1d 0x80c00780 0xd00e1019 0x80c00780 0xd00e0205 0x80c00780 0xd00e0001 0x80c00780
0xa0000e1d 0x44014780 0xc0060e19 0x00000780 0xa0000205 0x44014780 0xc0000201 0x00000780
0x10024003 0x00000780 0x1000f819 0x0403c780 0x10028003 0x00000780 0x1000f801 0x0403c780
0x00000605 0xc0000782 0x04001601 0xe4218780 0x00000605 0xc0000782 0x04001601 0xe4200780
0x20008419 0x00000013 0x30050dfd 0x6c0187c8 0x20008a01 0x00000013 0x30000ffd 0x6c00c7c8
0xa0039003 0x00000000 0x10038003 0x00000280 0xa003c003 0x00000000 0x1003b003 0x00000280
0xd0196005 0x20000780 0x2500e014 0x2500e21c 0xd0185805 0x20000780 0x2508e000 0x2508e204
0x20058400 0x20078414 0x20008001 0x00000013 0x20008a00 0x20018a04 0x30020001 0xc4100780
0x20008a15 0x00000013 0x30020001 0xc4100780 0x30020205 0xc4100780 0x2100ea00 0x2101ec04
0x30020a1d 0xc4100780 0x2000ca01 0x04200780 0x20008001 0x00000043 0xd00e0001 0x80c00780
0xd00e0015 0x80c00780 0x2000cc01 0x0421c780 0x20008205 0x00000043 0xd00e0205 0x80c00780
0xd00e0001 0x80c00780 0xa0000a15 0x44014780 0xa0000001 0x44014780 0xc0010001 0x00000780
0xc0000a01 0x00000780 0x10039003 0x00000780 0x1003c003 0x00000780 0x1000f801 0x0403c780
0x1000f801 0x0403c780 0x00020c05 0xc0000782 0x00000605 0xc0000782 0x04021601 0xe4200780
0x04001601 0xe4200780 0x861ffe03 0x00000000 0x861ffe03 0x00000000 0x3006d1fd 0x6c2047c8
0x307cd1fd 0x6c2047c8 0x10085003 0x00000280 0x10000c15 0x0403c780 0x30000003 0x00000280
0x300209fd 0x6c00c7e8 0x30040dfd 0x6c0187f8 0x307c09fd 0x640087c8 0x20400401 0x04014780
0x308105fd 0x6c40c7c8 0x00000019 0x20000780 0x3000d419 0xa4200780 0x20000805 0x04018780
0x2101f011 0x00000003 0x1000f815 0x0403c780 0x00000605 0xc0000780 0x300403fd 0x6400c7d8
0x308205fd 0x6c40c7c8 0x0000001d 0x20000780 0xa0072003 0x00000000 0x10000825 0x0403c780
0x308305fd 0x6c40c7c8 0x00000021 0x20000780 0x04041601 0xe43f0780 0x10072003 0x00001280
0x307c05fd 0x6c0087d8 0x20000a25 0x04008780 0x200b8801 0x00000003 0x00020005 0xc0000780
0x20009201 0x00000013 0x00020009 0xc0000780 0x20001201 0x04014780 0x00020009 0xc0000780
0x1800d601 0x0423c780 0x0002120d 0xc0000780 0x20209201 0x00000003 0x300601fd 0x6c0187d8
0xc400d625 0x00200780 0x00000609 0xc0000780 0x20009201 0x00000007 0x300601fd 0x6c0187e8
0x1c00d601 0x0423c780 0x1000f825 0x0403f280 0xd8035811 0x20000780 0x20209201 0x00000007
0xe800d601 0x00224780 0x10001201 0x0403e280 0xd403000d 0x20000780 0x1000c01d 0x0423c784
0x08041601 0xe4200780 0x861ffe03 0x00000000 0x300601fd 0x6c0187f8 0xd8025811 0x20000780
0x00000c01 0xa00007c0 0x00000609 0xc0000680 0xcc07c001 0x00200780 0xd801580d 0x20000780
0xd8145811 0x20000680 0xd810580d 0x20000680 0x1000c021 0x0423c784 0xd4020011 0x20000780
0x1000c001 0x0423c684 0xbc00c001 0x00200680 0x1000f801 0x0403f280 0x1c00c01d 0x0423c780
0x08041601 0xe4200680 0x861ffe03 0x00000000 0xd401000d 0x20000780 0xe008c021 0x00200784
0x00000e01 0xa00007c0 0x00000609 0xc0000680 0xcc07c01d 0x00200780 0x10000021 0x0403e280
0xd8125811 0x20000680 0xd810580d 0x20000680 0x1800d601 0x0423c780 0x1000f81d 0x0403d280
0x1000c001 0x0423c684 0xbc00c001 0x00200680 0xe400c001 0x0021c780 0x300613fd 0x6c0187d8
0x08041601 0xe4200680 0x861ffe03 0x00000000 0xb0001001 0x00000780 0xb0000e01 0x00021280
0x00001001 0xa00007c0 0x00000609 0xc0000680
0xd8115811 0x20000680 0xd810580d 0x20000680
0x1000c001 0x0423c684 0xbc00c001 0x00200680
0x08041601 0xe4200680 0x861ffe03 0x00000000
0x00000609 0xc0000780 0xd810580d 0x20000780 0x00000609 0xc0000780 0xd810580d 0x20000780
0x1c00e001 0x0423c780 0xbc00c001 0x00200780 0xbc00c001 0x00200780 0x08041601 0xe4200780
0x08041601 0xe4200780 0x1d00f000 0xbd006000 0x20009225 0x0000000b 0x300113fd 0x640047d8
0x08041601 0xe4200780 0x1d00e800 0xbd006000 0xd4040005 0x20000780 0x1004e003 0x00001280
0x08041601 0xe4200780 0x1d00e400 0xbd006000 0x00000605 0xc0000782 0xd4105809 0x20000780
0x08041601 0xe4200780 0xa0080003 0x00000000 0x1800e005 0x0423c780 0x1800f001 0x0423c780
0x10080003 0x00001100 0xd010580d 0x20000780 0xb9016004 0xb9007000 0xb0000201 0x00000780
0x1c00c201 0x0423c780 0x00020a09 0xc0000780 0x04041601 0xe4200780 0x1900e804 0x1900ec00
0xbc00c001 0x00200780 0x08061601 0xe4200780 0xb9016004 0xb9006400 0xb0000201 0x00000780
0xf0000001 0xe0000002 0x861ffe03 0x00000000 0x04041601 0xe4200780 0xa008d003 0x00000000
0x20018a15 0x00000003 0x30040bfd 0x6c0147c8 0x1008d003 0x00000100 0x10004e31 0x0023c780
0x10049003 0x00000280 0x3002d1fd 0x6c2047c8 0x2101f005 0x00000003 0xa0004c01 0x04200780
0x30000003 0x00000280 0x10004e01 0x0023c780 0x600c4801 0x00200780 0x40010419 0x00000780
0x60004805 0x00204780 0x2101f001 0x00000003 0x60000619 0x00018780 0x30100c19 0xc4100780
0x40030011 0x00000780 0x60020211 0x00010780 0x00000605 0xc0000780 0x60000401 0x00018780
0x30100811 0xc4100780 0x60020001 0x00010780 0xd4105805 0x20000780 0x20058004 0x1500e200
0x00000605 0xc0000780 0x20000001 0x04008780 0x30020205 0xc4100780 0xb5006000 0x2101e804
0xd4185805 0x20000780 0x30020005 0xc4100780 0xd00e0201 0xa0c00780 0xf0000001 0xe0000002
0x1500e000 0x2101e804 0xd00e0201 0xa0c00781 0x20088a15 0x00000003 0x3005d1fd 0x6c2187d8
0x10043003 0x00001280 0xf0000001 0xe0000001
} }
} }
code { code {
@@ -585,25 +584,25 @@ code {
code { code {
name = cudaComputeLPC name = cudaComputeLPC
lmem = 0 lmem = 0
smem = 1212 smem = 1260
reg = 10 reg = 10
bar = 1 bar = 1
const { const {
segname = const segname = const
segnum = 1 segnum = 1
offset = 0 offset = 0
bytes = 56 bytes = 52
mem { mem {
0x00000003 0x0000001f 0x0000003f 0x00000040 0x0000000f 0x0000001f 0x0000003f 0x00000040
0x00000001 0x00000020 0x7e800000 0x00000008 0x00000001 0x00000020 0x7e800000 0x00000008
0x0000000c 0x0000000f 0xfffff000 0x00000fff 0x0000000c 0xfffff000 0x00000fff 0x3e800000
0x3e800000 0x0000009e 0x0000009e
} }
} }
bincode { bincode {
0xa0000009 0x04000780 0x308005fd 0x644107c8 0xa0000009 0x04000780 0x308005fd 0x644107c8
0xa000b003 0x00000000 0x1000b003 0x00000280 0xa000b003 0x00000000 0x1000b003 0x00000280
0xa0004e01 0x04200780 0x30040001 0xc4100780 0xa0004e01 0x04200780 0x30060001 0xc4100780
0x30020405 0xc4100780 0x2100ec00 0x20008200 0x30020405 0xc4100780 0x2100ec00 0x20008200
0xd00e0001 0x80c00780 0x00020405 0xc0000780 0xd00e0001 0x80c00780 0x00020405 0xc0000780
0x04001201 0xe4200780 0x307ccffd 0x6c2047ca 0x04001201 0xe4200780 0x307ccffd 0x6c2047ca
@@ -621,75 +620,75 @@ code {
0x60000401 0x00018780 0x20000001 0x04014780 0x60000401 0x00018780 0x20000001 0x04014780
0x30020001 0xc4100780 0x2000ca01 0x04200780 0x30020001 0xc4100780 0x2000ca01 0x04200780
0xd00e0001 0x80c00780 0x10029003 0x00000780 0xd00e0001 0x80c00780 0x10029003 0x00000780
0x1000f801 0x0403c780 0x04015c01 0xe4200782 0x1000f801 0x0403c780 0x04017401 0xe4200782
0x861ffe03 0x00000000 0x307c09fd 0x6c0087e8 0x861ffe03 0x00000000 0x307c09fd 0x6c0087e8
0xd407700d 0x20002500 0xd4057009 0x20002500 0xd407d00d 0x20002500 0xd405d009 0x20002500
0x1c00c001 0x0423e500 0xb800c001 0x00202500 0x1c00c001 0x0423e500 0xb800c001 0x00202500
0x04015c01 0xe4202500 0x861ffe03 0x00000000 0x04017401 0xe4202500 0x861ffe03 0x00000000
0xa004c003 0x00000000 0x1004c003 0x00000100 0xa004c003 0x00000000 0x1004c003 0x00000100
0x308507fd 0x6440c7e8 0x1003a003 0x00002280 0x308507fd 0x6440c7e8 0x1003a003 0x00002280
0xd406700d 0x20000780 0xd4057009 0x20000780 0xd406d00d 0x20000780 0xd405d009 0x20000780
0x1d00e000 0xb9006000 0x04015c01 0xe4200780 0x1d00e000 0xb9006000 0x04017401 0xe4200780
0xd4057009 0x20000780 0x1800e001 0x0423c780 0xd405d009 0x20000780 0x1800e001 0x0423c780
0xb800c001 0x00200780 0x04015c01 0xe4200780 0xb800c001 0x00200780 0x04017401 0xe4200780
0x1900f000 0xb9006000 0x04015c01 0xe4200780 0x1900f000 0xb9006000 0x04017401 0xe4200780
0x1900e800 0xb9006000 0x04015c01 0xe4200780 0x1900e800 0xb9006000 0x04017401 0xe4200780
0x1900e400 0xb9006000 0x04015c01 0xe4200780 0x1900e400 0xb9006000 0x04017401 0xe4200780
0x1900e200 0xb9006000 0x307c05fd 0x6c0147e8 0x1900e200 0xb9006000 0x307c05fd 0x6c0147e8
0x04015c01 0xe4200780 0x1004c003 0x00002280 0x04017401 0xe4200780 0x1004c003 0x00002280
0xd005700d 0x20000780 0x00020a09 0xc0000780 0xd005d00d 0x20000780 0x00020a09 0xc0000780
0x1c00c001 0x0423c780 0x08009a01 0xe4200780 0x1c00c001 0x0423c780 0x0800b201 0xe4200780
0xf0000001 0xe0000002 0x20018a15 0x00000003 0xf0000001 0xe0000002 0x20018a15 0x00000003
0x300503fd 0x6c0147e8 0x10018003 0x00002280 0x300503fd 0x6c0147e8 0x10018003 0x00002280
0x10052003 0x00000780 0x308105fd 0x6c40c7c8 0x10052003 0x00000780 0x308105fd 0x6c40c7c8
0x30000003 0x00000100 0x00020405 0xc0000780 0x30000003 0x00000100 0x00020405 0xc0000780
0xd4027009 0x20000780 0x1800c001 0x0423c780 0xd402d009 0x20000780 0x1800c001 0x0423c780
0x0400dc01 0xe4200780 0x1800c001 0x0423c780 0x0400f401 0xe4200780 0x1800c001 0x0423c780
0x04011c01 0xe4200780 0xd0026809 0x20000780 0x04013401 0xe4200780 0xd002c809 0x20000780
0x04001a01 0xe43f0780 0x307ccffd 0x6c20c7c8 0x04003201 0xe43f0780 0x307ccffd 0x6c20c7c8
0x3002040d 0xc4100780 0x1800c001 0x0423c780 0x3002040d 0xc4100780 0x1800c001 0x0423c780
0x30000003 0x00000280 0x307c05fd 0x6c0087c8 0x30000003 0x00000280 0x307c05fd 0x6c0087c8
0x213fee11 0x0fffffff 0x1000f815 0x0403c780 0x213fee11 0x0fffffff 0x1000f815 0x0403c780
0xd0047005 0x20000780 0xb08601fd 0x605107d8 0xd004d005 0x20000780 0xb08601fd 0x605107d8
0x10000005 0x0403c780 0xa400c019 0xe4204780 0x10000005 0x0403c780 0xa400c019 0xe4204780
0xc08c0c19 0x00401680 0xc08c0205 0x00401680 0xc08b0c19 0x00401680 0xc08b0205 0x00401680
0x90000204 0xc0010c04 0xd0047005 0x20000780 0x90000204 0xc0010c04 0xd004d005 0x20000780
0xc401c019 0x0020c780 0xb0060000 0x20458818 0xc401c019 0x0020c780 0xb0060000 0x20458818
0x300605fd 0x6c0187d8 0xa0077003 0x00000000 0x300605fd 0x6c0187d8 0xa0077003 0x00000000
0x10077003 0x00001280 0x00000605 0xc0000780 0x10077003 0x00001280 0x00000605 0xc0000780
0xd403700d 0x20000780 0xd4047809 0x20000780 0xd403d00d 0x20000780 0xd404d809 0x20000780
0xcc01c019 0x0020c780 0xc801c01d 0x0020c780 0xcc01c019 0x0020c780 0xc801c01d 0x0020c780
0xb9066018 0xbd07601c 0x04011c01 0xe4218780 0xb9066018 0xbd07601c 0x04013401 0xe4218780
0x0400dc01 0xe421c780 0x20400a19 0x04008782 0x0400f401 0xe421c780 0x20400a19 0x04008782
0x00020c05 0xc0000780 0x30020bfd 0x6c00c7d8 0x00020c05 0xc0000780 0x30020bfd 0x6c00c7d8
0xc401d819 0x0020c780 0x1000f819 0x0403d280 0xc401f019 0x0020c780 0x1000f819 0x0403d280
0x30020bfd 0x6c0147d8 0xb0000c05 0x00004780 0x30020bfd 0x6c0147d8 0xb0000c05 0x00004780
0x10000c05 0x0403d280 0x00000609 0xc0000780 0x10000c05 0x0403d280 0x00000609 0xc0000780
0xb800da05 0x00204780 0x08001a01 0xe4204780 0xb800f205 0x00204780 0x08003201 0xe4204780
0xa800da05 0xc4304780 0xc0000205 0x04700003 0xa800f205 0xc4304780 0xc0000205 0x04700003
0xa0000205 0x8c0047d0 0x2000d619 0x04214780 0xa0000205 0x8c0047d0 0x2000d619 0x04214780
0xa0000205 0x44065680 0x30170205 0xec101680 0xa0000205 0x44065680 0x30170205 0xec101680
0x31000205 0x04435680 0x10000a05 0x2440d100 0x31000205 0x04431680 0x10000a05 0x2440d100
0x30870bfd 0x6c4107d8 0x100d801d 0x00000003 0x30870bfd 0x6c4107d8 0x100d801d 0x00000003
0x1000101d 0x2440d280 0x20000e05 0x04004780 0x1000101d 0x2440d280 0x20000e05 0x04004780
0x30020a1d 0x6c0187e0 0xd0840e1d 0x04400780 0x30020a1d 0x6c0187e0 0xd0840e1d 0x04400780
0x30218205 0x00000003 0x40070205 0x00018780 0x30218205 0x00000003 0x40070205 0x00018780
0x00000609 0xc0000780 0x08005a01 0xe4204780 0x00000609 0xc0000780 0x08007201 0xe4204780
0xd801680d 0x20000780 0x1c00e005 0x0423c780 0xd801c80d 0x20000780 0x1c00e005 0x0423c780
0x3c01c005 0x8c200780 0x08005a01 0xe4204780 0x3c01c005 0x8c200780 0x08007201 0xe4204780
0x1c00d005 0x0423c780 0x3c01c005 0x8c200780 0x1c00d005 0x0423c780 0x3c01c005 0x8c200780
0x08005a01 0xe4204780 0x1c00c805 0x0423c780 0x08007201 0xe4204780 0x1c00c805 0x0423c780
0x3c01c005 0x8c200780 0x08005a01 0xe4204780 0x3c01c005 0x8c200780 0x08007201 0xe4204780
0x1c00c405 0x0423c780 0x3c01c005 0x8c200780 0x1c00c405 0x0423c780 0x3c01c005 0x8c200780
0x08005a01 0xe4204780 0x1c00c205 0x0423c780 0x08007201 0xe4204780 0x1c00c205 0x0423c780
0x3c01c005 0x8c200780 0x08005a01 0xe4204780 0x3c01c005 0x8c200780 0x08007201 0xe4204780
0xd0016809 0x20000780 0x390fe005 0x00000003 0xd001c809 0x20000780 0x390fe005 0x00000003
0x30890205 0xac400780 0x1001801d 0x00000003 0x30800205 0xac400780 0x1001801d 0x00000003
0x307c0205 0x8c000780 0x30010e1d 0xc4000780 0x307c0205 0x8c000780 0x30010e1d 0xc4000780
0xa0000e21 0x44014780 0x103f801d 0x000001ff 0xa0000e21 0x44014780 0x103f801d 0x000001ff
0xc408da21 0x00200780 0x1000161d 0x2440d280 0xc408f221 0x00200780 0x1000141d 0x2440d280
0xa0001021 0xac004780 0x30080e21 0xac000780 0xa0001021 0xac004780 0x30080e21 0xac000780
0x1000801d 0x0ffffe03 0x1000141d 0x2440d280 0x1000801d 0x0ffffe03 0x1000121d 0x2440d280
0x30080e1d 0x8c000780 0xa00bb003 0x00000000 0x30080e1d 0x8c000780 0xa00bb003 0x00000000
0x100bb003 0x00002100 0x30070c21 0xc4100780 0x100bb003 0x00002100 0x30070c21 0xc4100780
0x30060c25 0xc4100780 0x20099020 0x2108e820 0x30060c25 0xc4100780 0x20099020 0x2108e820
@@ -699,24 +698,24 @@ code {
0x20001021 0x04024680 0x2000c821 0x04220680 0x20001021 0x04024680 0x2000c821 0x04220680
0x21001021 0x0441c680 0xd00e1005 0xa0c00680 0x21001021 0x0441c680 0xd00e1005 0xa0c00680
0x307c0ffd 0x6c0087d8 0xa0000e05 0x44065500 0x307c0ffd 0x6c0087d8 0xa0000e05 0x44065500
0x30170205 0xec101500 0x31000205 0x04435500 0x30170205 0xec101500 0x31000205 0x04431500
0x10000a05 0x2440d280 0xd007001d 0x0402c780 0x10000a05 0x2440d280 0xd007001d 0x0402c780
0x307c0ffd 0x6c0087d8 0xa0000e1d 0x44065500 0x307c0ffd 0x6c0087d8 0xa0000e1d 0x44065500
0x30170e1d 0xec101500 0x31000e1d 0x04435500 0x30170e1d 0xec101500 0x31000e1d 0x04431500
0x10000a1d 0x2440d280 0x30070205 0x8c000780 0x10000a1d 0x2440d280 0x30070205 0x8c000780
0x00000605 0xc0000780 0x30218205 0x00000003 0x00000605 0xc0000780 0x30218205 0x00000003
0x04005a01 0xe4204780 0xd4016809 0x20000780 0x04007201 0xe4204780 0xd401c809 0x20000780
0x1800e005 0x0423c780 0x3801c005 0x8c200780 0x1800e005 0x0423c780 0x3801c005 0x8c200780
0x04005a01 0xe4204780 0x1800d005 0x0423c780 0x04007201 0xe4204780 0x1800d005 0x0423c780
0x3801c005 0x8c200780 0x04005a01 0xe4204780 0x3801c005 0x8c200780 0x04007201 0xe4204780
0x1800c805 0x0423c780 0x3801c005 0x8c200780 0x1800c805 0x0423c780 0x3801c005 0x8c200780
0x04005a01 0xe4204780 0x1800c405 0x0423c780 0x04007201 0xe4204780 0x1800c405 0x0423c780
0x3801c005 0x8c200780 0x04005a01 0xe4204780 0x3801c005 0x8c200780 0x04007201 0xe4204780
0x1800c205 0x0423c780 0x3801c005 0x8c200780 0x1800c205 0x0423c780 0x3801c005 0x8c200780
0x04005a01 0xe4204780 0xa00ea003 0x00000000 0x04007201 0xe4204780 0xa00ea003 0x00000000
0x100ea003 0x00000100 0x30070c05 0xc4100780 0x100ea003 0x00000100 0x30070c05 0xc4100780
0x30060c19 0xc4100780 0x20000205 0x04018780 0x30060c19 0xc4100780 0x20000205 0x04018780
0xd0016805 0x20000780 0x2101e818 0x1500e004 0xd001c805 0x20000780 0x2101e818 0x1500e004
0x200c8c19 0x00000003 0xd00e0c05 0xa0c00780 0x200c8c19 0x00000003 0xd00e0c05 0xa0c00780
0xf0000001 0xe0000002 0x20018a15 0x00000003 0xf0000001 0xe0000002 0x20018a15 0x00000003
0x3005cffd 0x6c2147d8 0x10062003 0x00001280 0x3005cffd 0x6c2147d8 0x10062003 0x00001280