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