diff --git a/CUETools.FlaCuda/FlaCudaWriter.cs b/CUETools.FlaCuda/FlaCudaWriter.cs index 5439834..706210c 100644 --- a/CUETools.FlaCuda/FlaCudaWriter.cs +++ b/CUETools.FlaCuda/FlaCudaWriter.cs @@ -1171,7 +1171,7 @@ namespace CUETools.Codecs.FlaCuda cuda.SetParameter(cudaComputeLPC, sizeof(uint) * 3, (uint)max_order); cuda.SetParameter(cudaComputeLPC, sizeof(uint) * 3 + sizeof(uint), (uint)partCount); cuda.SetParameterSize(cudaComputeLPC, (uint)(sizeof(uint) * 3) + sizeof(uint) * 2); - cuda.SetFunctionBlockShape(cudaComputeLPC, 64, 1, 1); + cuda.SetFunctionBlockShape(cudaComputeLPC, (partCount + 31) & ~31, 1, 1); // issue work to the GPU cuda.CopyHostToDeviceAsync(cudaSamples, samplesBufferPtr, (uint)(sizeof(int) * FlaCudaWriter.MAX_BLOCKSIZE * channelsCount), cudaStream); diff --git a/CUETools.FlaCuda/flacuda.cu b/CUETools.FlaCuda/flacuda.cu index 79017aa..fd166a1 100644 --- a/CUETools.FlaCuda/flacuda.cu +++ b/CUETools.FlaCuda/flacuda.cu @@ -110,6 +110,7 @@ extern "C" __global__ void cudaComputeLPC( volatile float autoc[33]; volatile float gen0[32]; volatile float gen1[32]; + volatile float parts[128]; //volatile float reff[32]; int cbits; } shared; @@ -119,15 +120,26 @@ extern "C" __global__ void cudaComputeLPC( if (tid < sizeof(shared.task) / sizeof(int)) ((int*)&shared.task)[tid] = ((int*)(tasks + blockIdx.y))[tid]; - // initialize autoc sums - if (tid <= max_order) - shared.autoc[tid] = 0.0f; - // add up parts - for (int part = 0; part < partCount; part++) - if (tid <= max_order) - shared.autoc[tid] += autoc[(blockIdx.y * partCount + part) * (max_order + 1) + tid]; - + for (int order = 0; order <= max_order; order++) + { + shared.parts[tid] = tid < partCount ? autoc[(blockIdx.y * partCount + tid) * (max_order + 1) + order] : 0; + __syncthreads(); + if (tid < 64 && blockDim.x > 64) shared.parts[tid] += shared.parts[tid + 64]; + __syncthreads(); + if (tid < 32) + { + if (blockDim.x > 32) shared.parts[tid] += shared.parts[tid + 32]; + shared.parts[tid] += shared.parts[tid + 16]; + shared.parts[tid] += shared.parts[tid + 8]; + shared.parts[tid] += shared.parts[tid + 4]; + shared.parts[tid] += shared.parts[tid + 2]; + shared.parts[tid] += shared.parts[tid + 1]; + if (tid == 0) + shared.autoc[order] = shared.parts[0]; + } + } + if (tid < 32) { shared.gen0[tid] = shared.autoc[tid+1]; @@ -195,7 +207,7 @@ extern "C" __global__ void cudaEstimateResidual( { __shared__ struct { int data[32*8]; - int residual[32*8]; + volatile int residual[32*8]; encodeResidualTaskStruct task[8]; } shared; const int tid = threadIdx.x + threadIdx.y * blockDim.x; @@ -207,11 +219,11 @@ extern "C" __global__ void cudaEstimateResidual( // fetch samples shared.data[tid] = tid < dataLen ? samples[shared.task[0].samplesOffs + pos + tid] : 0; - shared.residual[tid] = 0; const int residualLen = max(0,min(frameSize - pos - shared.task[threadIdx.y].residualOrder, partSize)) * (shared.task[threadIdx.y].residualOrder != 0); __syncthreads(); + shared.residual[tid] = 0; shared.task[threadIdx.y].coefs[threadIdx.x] = threadIdx.x < max_order ? tasks[blockIdx.y * blockDim.y + threadIdx.y].coefs[threadIdx.x] : 0; for (int i = threadIdx.x; i - threadIdx.x < residualLen; i += blockDim.x) // += 32 @@ -235,7 +247,6 @@ extern "C" __global__ void cudaEstimateResidual( // rice parameter search shared.residual[tid] = __mul24(threadIdx.x >= 15, 0x7fffff) + residualLen * (threadIdx.x + 1) + ((shared.residual[threadIdx.y * blockDim.x] - (residualLen >> 1)) >> threadIdx.x); - __syncthreads(); shared.residual[tid] = min(shared.residual[tid], shared.residual[tid + 8]); shared.residual[tid] = min(shared.residual[tid], shared.residual[tid + 4]); shared.residual[tid] = min(shared.residual[tid], shared.residual[tid + 2]); diff --git a/CUETools.FlaCuda/flacuda.cubin b/CUETools.FlaCuda/flacuda.cubin index 20f5e92..c6c1117 100644 --- a/CUETools.FlaCuda/flacuda.cubin +++ b/CUETools.FlaCuda/flacuda.cubin @@ -125,55 +125,56 @@ code { 0xa000420d 0x04200780 0x30100c19 0xc4100780 0x60080411 0x00018780 0x40060a05 0x00000780 0x30100219 0xc4100780 0x1000d205 0x0423c780 - 0x60060819 0x00018780 0x2101ee0c 0x2144f004 - 0x20000c1d 0x04000780 0x3003020d 0xac000780 - 0x30030ffd 0x6c0187c8 0xd010580d 0x20000780 - 0x2c00c00d 0x04210500 0x20000e0d 0x0400c500 - 0x3002060d 0xc4100500 0x2000ca0d 0x0420c500 - 0xd00e060d 0x80c00500 0x1000f80d 0x0403c280 - 0x00020e0d 0xc0000780 0x0c001401 0xe420c780 - 0xd4105011 0x20000780 0x0c021401 0xe43f0780 - 0x3000c005 0x04204784 0x307cc00d 0x6c208784 - 0x3001d205 0xac200780 0x307c0205 0x8c000780 - 0xd0030211 0x04020780 0xd4005005 0x20000780 - 0x861ffe03 0x00000000 0x3000cffd 0x6420c7c8 - 0xa0044003 0x00000000 0x10043003 0x00000280 - 0x10004409 0x0023c780 0x60024e05 0x00208780 - 0x3007020d 0xc4100780 0x30060205 0xc4100780 - 0x20018604 0x2101ec04 0x20000a05 0x04004780 - 0x20008205 0x00000007 0xd00e0205 0x80c00780 - 0x10044003 0x00000780 0x1000f805 0x0403c780 - 0x307c09fd 0x640087ca 0x08043401 0xe4204780 - 0xa0076003 0x00000000 0x10000005 0x0403c780 - 0x10073003 0x00000280 0x00020e0d 0xc0000780 - 0xdc085009 0x20000780 0x1800c015 0x0423c780 + 0x60060819 0x00018780 0x2101ee20 0x2144f004 + 0x20000c1d 0x04000780 0x30080221 0xac000780 + 0x30080ffd 0x6c0187c8 0xd010580d 0x20000780 + 0x2c00c011 0x04210500 0x20000e11 0x04010500 + 0x30020811 0xc4100500 0x2000ca11 0x04210500 + 0xd00e0811 0x80c00500 0x1000f811 0x0403c280 + 0x00020e0d 0xc0000780 0x0c001401 0xe4210780 + 0xd410500d 0x20000780 0x3c00c005 0x04204780 + 0x3c7cc011 0x6c208780 0x3001d205 0xac200780 + 0x307c0205 0x8c000780 0xd0040211 0x04020780 + 0xd4005005 0x20000780 0x861ffe03 0x00000000 + 0x3000cffd 0x6420c7c8 0xa0045003 0x00000000 + 0x00020e0d 0xc0000780 0x0c021401 0xe43f0780 + 0x10044003 0x00000280 0x10004409 0x0023c780 + 0x60024e05 0x00208780 0x30070221 0xc4100780 + 0x30060205 0xc4100780 0x20019004 0x2101ec04 + 0x20000a05 0x04004780 0x20008205 0x00000007 + 0xd00e0205 0x80c00780 0x10045003 0x00000780 + 0x1000f805 0x0403c780 0x307c09fd 0x640087ca + 0x08043401 0xe4204780 0xa0072003 0x00000000 + 0x10000005 0x0403c780 0x10072003 0x00000280 0xd4100009 0x20000780 0x387cc1fd 0x6c20c7c8 - 0xa0060003 0x00000000 0x1000f80d 0x0403c780 - 0x1000f825 0x0403c780 0x10060003 0x00000280 - 0xa005e003 0x00000000 0xd4000009 0x20000780 - 0x20000221 0x0400c780 0x0002100d 0xc0000780 + 0xa005e003 0x00000000 0x1000f815 0x0403c780 + 0x1000f825 0x0403c780 0x1005e003 0x00000280 + 0xa005c003 0x00000000 0xd4000009 0x20000780 + 0x20000221 0x04014780 0x0002100d 0xc0000780 0xd8108011 0x20000780 0x1000c021 0x0423c784 0x6c08d425 0x80224780 0xd410000d 0x20000780 - 0x2001860d 0x00000003 0x3c03c1fd 0x6c2147c8 - 0xd8000809 0x20000780 0x10054003 0x00000280 - 0xd4100009 0x20000782 0x1800c00d 0x0423c780 - 0x2000020d 0x0400c782 0x00020609 0xc0000780 - 0xd410100d 0x20000780 0x1c00c00d 0x0423c780 - 0x3003120d 0xec000780 0x2840d40d 0x0420c780 - 0x301f0621 0xec100780 0x3001060d 0xc4100780 - 0xd0031021 0x04008780 0x3001080d 0x6c010780 - 0xa000060d 0x2c014780 0x60080615 0x80014780 - 0xa000420d 0x04200780 0x20038204 0x2040820c - 0x300309fd 0x640107c8 0x00020e0d 0xc0000780 - 0x0c021401 0xe4214780 0x1004c003 0x00000280 - 0x10076003 0x00000780 0x00020e0d 0xc0000780 - 0xdc085009 0x20000780 0x1800c015 0x0423c780 - 0x00020e0d 0xc0000782 0xdc085809 0x20000780 - 0x2800de05 0x04214780 0x0c021401 0xe4204780 - 0x2800ce05 0x04204780 0x0c021401 0xe4204780 - 0x2800c605 0x04204780 0x0c021401 0xe4204780 - 0x2800c20d 0x04204780 0x20018005 0x00000003 - 0x0c021401 0xe420c780 0x2903e00c 0x40031014 + 0x20018a15 0x00000003 0x3c05c1fd 0x6c2147c8 + 0xd8000809 0x20000780 0x10052003 0x00000280 + 0xd4100009 0x20000782 0x1800c015 0x0423c780 + 0x20000215 0x04014782 0x00020a09 0xc0000780 + 0xd410100d 0x20000780 0x1c00c015 0x0423c780 + 0x30051215 0xec000780 0x2840d415 0x04214780 + 0x301f0a21 0xec100780 0x30010a15 0xc4100780 + 0xd0051021 0x04008780 0x00020e0d 0xc0000780 + 0xdc085009 0x20000780 0x30010825 0x6c010780 + 0x1800c015 0x0423c780 0xa0001225 0x2c014780 + 0x60081215 0x80014780 0x20000205 0x0400c780 + 0x0c021401 0xe4214780 0x20400215 0x04000780 + 0x300509fd 0x640107c8 0x1004a003 0x00000280 + 0x00020e0d 0xc0000782 0xdc085011 0x20000780 + 0x1000e005 0x0423c784 0x2000c005 0x04204784 + 0x0c021401 0xe4204780 0x1000d005 0x0423c784 + 0x2000c005 0x04204784 0x0c021401 0xe4204780 + 0x1000c805 0x0423c784 0x2000c005 0x04204784 + 0x0c021401 0xe4204780 0x1000c405 0x0423c784 + 0x2000c005 0x04204784 0x0c021401 0xe4204780 + 0x20018005 0x00000003 0x1000c20d 0x0423c784 + 0x2000c00d 0x0420c784 0x40031015 0x00000780 0x0c021401 0xe420c780 0x6002120d 0x00014780 0x30820021 0x64410780 0x00020c09 0xc0000780 0xd8085009 0x20000780 0x30100615 0xc4100780 @@ -181,20 +182,21 @@ code { 0x60021011 0x00014780 0x407f8c15 0x0007ffff 0x2943e004 0x2005880c 0x30000205 0xec000780 0x20000205 0x0400c780 0x0c021401 0xe4204780 - 0x861ffe03 0x00000000 0x00020e0d 0xc0000780 - 0xdc085009 0x20000780 0x1800d005 0x0423c780 - 0x3801c005 0xac200780 0x0c021401 0xe4204780 - 0x3801c805 0xac200780 0x0c021401 0xe4204780 - 0x3801c405 0xac200780 0x0c021401 0xe4204780 - 0x3801c205 0xac200780 0x307c01fd 0x640147c8 + 0x1000d005 0x0423c784 0x3001c005 0xac200784 + 0x0c021401 0xe4204780 0x1000c805 0x0423c784 + 0x3001c005 0xac200784 0x0c021401 0xe4204780 + 0x1000c405 0x0423c784 0x3001c005 0xac200784 + 0x0c021401 0xe4204780 0x1000c205 0x0423c784 + 0x3001c005 0xac200784 0x307c01fd 0x640147c8 0x0c021401 0xe4204780 0x30000003 0x00000280 0xd4100005 0x20000780 0x347cc1fd 0x6c2087c8 0x30000003 0x00000280 0x10004401 0x0023c780 - 0x60004e01 0x00208780 0x40014809 0x00200780 - 0x30100409 0xc4100780 0x60004801 0x00208780 + 0x60004e01 0x00208780 0x40014805 0x00200780 + 0x30100205 0xc4100780 0x60004801 0x00204780 0xa0004c11 0x04200780 0x20000001 0x04010780 - 0x30020001 0xc4100780 0x2000c801 0x04200780 - 0xd00e0005 0xa0c00781 + 0x00020e0d 0xc0000780 0xdc085005 0x20000780 + 0x30020005 0xc4100780 0x1500e000 0x2101e804 + 0xd00e0201 0xa0c00781 } } code { @@ -278,117 +280,138 @@ code { code { name = cudaComputeLPC lmem = 0 - smem = 700 + smem = 1212 reg = 10 - bar = 0 + bar = 1 const { segname = const segnum = 1 offset = 0 - bytes = 44 + bytes = 52 mem { - 0x00000003 0x0000001f 0x7e800000 0x00000001 - 0x0000000f 0x00001fff 0xffffe000 0x3e800000 - 0x00000020 0x0000009e 0x00000008 + 0x00000003 0x0000001f 0x0000003f 0x00000040 + 0x00000001 0x00000020 0x7e800000 0x0000000f + 0x00001fff 0xffffe000 0x3e800000 0x0000009e + 0x00000008 } } bincode { - 0xa000000d 0x04000780 0x308007fd 0x644107c8 + 0xa0000009 0x04000780 0x308005fd 0x644107c8 0xa000b003 0x00000000 0x1000b003 0x00000280 0xa0004e01 0x04200780 0x30040001 0xc4100780 - 0x30020605 0xc4100780 0x2100ec00 0x20008200 - 0xd00e0001 0x80c00780 0x00020605 0xc0000780 - 0x04001201 0xe4200780 0x3003ce01 0x6c2187d2 - 0xa00001fd 0x0c0147c8 0x00020605 0xc0001680 - 0x04009a01 0xe43f1680 0x307cd1fd 0x6c20c7d8 - 0x1002a003 0x00001280 0x1000f805 0x0403c780 - 0xa0026003 0x00000000 0x10026003 0x00000100 - 0x1000d001 0x0423c780 0x40014e09 0x00200780 - 0x30100409 0xc4100780 0x60004e09 0x00208780 - 0x2101ee01 0x00000003 0x20000409 0x04004780 - 0x40010811 0x00000780 0x60000a11 0x00010780 - 0x30100811 0xc4100780 0x60000801 0x00010780 - 0x20000601 0x04000780 0x30020001 0xc4100780 - 0x00020605 0xc0000780 0x2000ca01 0x04200780 - 0xd00e0001 0x80c00780 0xd4026809 0x20000780 - 0xb800c001 0x00200780 0x04009a01 0xe4200780 - 0xf0000001 0xe0000002 0x20018205 0x00000003 - 0x3001d1fd 0x6c2147d8 0x10012003 0x00001280 - 0x308107fd 0x6c4107c8 0x30000003 0x00000280 - 0x00020605 0xc0000780 0xd4027009 0x20000780 - 0x1800c001 0x0423c780 0x0400dc01 0xe4200780 - 0x1800c001 0x0423c780 0x04011c01 0xe4200780 - 0xd0026809 0x20000780 0x04001a01 0xe43f0780 - 0x307ccffd 0x6c20c7c8 0x30020611 0xc4100780 - 0x1800c001 0x0423c780 0x30000003 0x00000280 - 0x307c07fd 0x6c0087c8 0x213fee15 0x0fffffff - 0x1000f819 0x0403c780 0xd0047005 0x20000780 - 0xb08201fd 0x605107d8 0x10000005 0x0403c780 - 0xa400c009 0xe4204780 0xc0870409 0x00401680 - 0xc0870205 0x00401680 0x90000204 0xc001041c - 0x20400a05 0x04018780 0xd0047005 0x20000780 - 0x300107fd 0x6c0187d8 0xa0050003 0x00000000 - 0xe407c001 0x00200780 0x10050003 0x00001280 - 0x00000805 0xc0000780 0xd404780d 0x20000780 - 0xd4037009 0x20000780 0x1d00e004 0x1900e008 - 0xe807c005 0x00204780 0xec07c009 0x00208780 - 0x04011c01 0xe4204780 0x0400dc01 0xe4208780 - 0x20400c05 0x0400c782 0x00020205 0xc0000780 - 0x30030dfd 0x6c00c7d8 0xc407d809 0x00200780 - 0x1000f809 0x0403d280 0x30030dfd 0x6c0147d8 - 0xb0000405 0x0001c780 0x10000405 0x0403d280 - 0x00000809 0xc0000780 0xb800da05 0x00204780 - 0x08001a01 0xe4204780 0xa800da05 0xc4304780 - 0xc0000205 0x04700003 0xa0000205 0x8c0047d0 - 0x2000d609 0x04218780 0xa0000205 0x44065680 - 0x30170205 0xec101680 0x31000205 0x04425680 - 0x10001005 0x2440d100 0x30030c1d 0x6c0187d0 - 0x30148205 0x00000003 0xd0830e1d 0x04400780 - 0x40070205 0x00018780 0x00000809 0xc0000780 - 0x08005a01 0xe4204780 0xd801680d 0x20000780 - 0x1c00e005 0x0423c780 0x3c01c005 0x8c200780 - 0x08005a01 0xe4204780 0x1c00d005 0x0423c780 + 0x30020405 0xc4100780 0x2100ec00 0x20008200 + 0xd00e0001 0x80c00780 0x00020405 0xc0000780 + 0x04001201 0xe4200780 0x307ccffd 0x6c2047ca + 0x10051003 0x00000280 0xa000420d 0x04200780 + 0x30820401 0x6c40c780 0x30830605 0x64410780 + 0xd0840001 0x04400780 0xd0840205 0x04400780 + 0xd0010011 0x04000780 0x308105fd 0x6c40c7c8 + 0x3002d1fd 0x6c2107d8 0x00020405 0xc0000780 + 0x1000f815 0x0403c780 0x2101ee05 0x00000003 + 0xa0029003 0x00000000 0x10028003 0x00001100 + 0x1000d001 0x0423c780 0x40014e19 0x00200780 + 0x30100c19 0xc4100780 0x60004e01 0x00218780 + 0x20000001 0x04008780 0x40010419 0x00000780 + 0x60000619 0x00018780 0x30100c19 0xc4100780 + 0x60000401 0x00018780 0x20000001 0x04014780 + 0x30020001 0xc4100780 0x2000ca01 0x04200780 + 0xd00e0001 0x80c00780 0x10029003 0x00000780 + 0x1000f801 0x0403c780 0x04015c01 0xe4200782 + 0x861ffe03 0x00000000 0x307c09fd 0x6c0087e8 + 0xd407700d 0x20002500 0xd4057009 0x20002500 + 0x1c00c001 0x0423e500 0xb800c001 0x00202500 + 0x04015c01 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 + 0x1900e200 0xb9006000 0x307c05fd 0x6c0147e8 + 0x04015c01 0xe4200780 0x1004c003 0x00002280 + 0xd005700d 0x20000780 0x00020a09 0xc0000780 + 0x1c00c001 0x0423c780 0x08009a01 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 + 0x30020411 0xc4100780 0x1800c001 0x0423c780 + 0x30000003 0x00000280 0x307c05fd 0x6c0087c8 + 0x213fee15 0x0fffffff 0x1000f819 0x0403c780 + 0xd0047005 0x20000780 0xb08601fd 0x605107d8 + 0x10000005 0x0403c780 0xa400c00d 0xe4204780 + 0xc08a060d 0x00401680 0xc08a0205 0x00401680 + 0x90000204 0xc001061c 0x20400a05 0x04018780 + 0xd0047005 0x20000780 0x300105fd 0x6c0187d8 + 0xa0077003 0x00000000 0xe407c001 0x00200780 + 0x10077003 0x00001280 0x00000805 0xc0000780 + 0xd404780d 0x20000780 0xd4037009 0x20000780 + 0x1d00e004 0x1900e00c 0xe807c005 0x00204780 + 0xec07c00d 0x0020c780 0x04011c01 0xe4204780 + 0x0400dc01 0xe420c780 0x20400c05 0x04008782 + 0x00020205 0xc0000780 0x30020dfd 0x6c00c7d8 + 0xc407d80d 0x00200780 0x1000f80d 0x0403d280 + 0x30020dfd 0x6c0147d8 0xb0000605 0x0001c780 + 0x10000605 0x0403d280 0x00000809 0xc0000780 + 0xb800da05 0x00204780 0x08001a01 0xe4204780 + 0xa800da05 0xc4304780 0xc0000205 0x04700003 + 0xa0000205 0x8c0047d0 0x2000d60d 0x04218780 + 0xa0000205 0x44065680 0x30170205 0xec101680 + 0x31000205 0x0442d680 0x10000a05 0x2440d100 + 0x30020c1d 0x6c0187d0 0x30148205 0x00000003 + 0xd0840e1d 0x04400780 0x40070205 0x00018780 + 0x00000809 0xc0000780 0x08005a01 0xe4204780 + 0xd801680d 0x20000780 0x1c00e005 0x0423c780 0x3c01c005 0x8c200780 0x08005a01 0xe4204780 - 0x1c00c805 0x0423c780 0x3c01c005 0x8c200780 - 0x08005a01 0xe4204780 0x1c00c405 0x0423c780 + 0x1c00d005 0x0423c780 0x3c01c005 0x8c200780 + 0x08005a01 0xe4204780 0x1c00c805 0x0423c780 0x3c01c005 0x8c200780 0x08005a01 0xe4204780 - 0x1c00c205 0x0423c780 0x3c01c005 0x8c200780 - 0x08005a01 0xe4204780 0xd0016809 0x20000780 - 0x390fe005 0x00000003 0x30840205 0xac400780 - 0x1001801d 0x00000003 0x307c0205 0x8c000780 - 0x30010e1d 0xc4000780 0xa0000e1d 0x44014780 - 0xc407da1d 0x00200780 0xa0000e1d 0xac004780 - 0x30850e1d 0xac400780 0xa008c003 0x00000000 - 0x30860e1d 0x8c400780 0x1008c003 0x00001100 - 0x30070421 0xc4100780 0x30060425 0xc4100780 - 0x20099020 0x2108e820 0x20000821 0x04020780 - 0x20009021 0x00000007 0xd00e101d 0xa0c00780 - 0xf0000001 0xe0000002 0x30070421 0xc4100680 - 0x30060425 0xc4100680 0x20001021 0x04024680 - 0x2000c821 0x04220680 0x21001021 0x04428680 - 0xd00e1005 0xa0c00680 0x307c0ffd 0x6c0087d8 - 0xa0000e05 0x44065500 0x30170205 0xec101500 - 0x31000205 0x04425500 0x10001005 0x2440d280 - 0xd007001d 0x0402c780 0x307c0ffd 0x6c0087d8 - 0xa0000e1d 0x44065500 0x30170e1d 0xec101500 - 0x31000e1d 0x04425500 0x1000101d 0x2440d280 - 0x30070205 0x8c000780 0x00000805 0xc0000780 - 0x30218205 0x00000003 0x04005a01 0xe4204780 - 0xd4016809 0x20000780 0x1800e005 0x0423c780 + 0x1c00c405 0x0423c780 0x3c01c005 0x8c200780 + 0x08005a01 0xe4204780 0x1c00c205 0x0423c780 + 0x3c01c005 0x8c200780 0x08005a01 0xe4204780 + 0xd0016809 0x20000780 0x390fe005 0x00000003 + 0x30870205 0xac400780 0x1001801d 0x00000003 + 0x307c0205 0x8c000780 0x30010e1d 0xc4000780 + 0xa0000e1d 0x44014780 0xc407da1d 0x00200780 + 0xa0000e1d 0xac004780 0x30880e1d 0xac400780 + 0xa00b3003 0x00000000 0x30890e1d 0x8c400780 + 0x100b3003 0x00001100 0x30070621 0xc4100780 + 0x30060625 0xc4100780 0x20099020 0x2108e820 + 0x20000821 0x04020780 0x20009021 0x00000007 + 0xd00e101d 0xa0c00780 0xf0000001 0xe0000002 + 0x30070621 0xc4100680 0x30060625 0xc4100680 + 0x20001021 0x04024680 0x2000c821 0x04220680 + 0x21001021 0x04430680 0xd00e1005 0xa0c00680 + 0x307c0ffd 0x6c0087d8 0xa0000e05 0x44065500 + 0x30170205 0xec101500 0x31000205 0x0442d500 + 0x10000a05 0x2440d280 0xd007001d 0x0402c780 + 0x307c0ffd 0x6c0087d8 0xa0000e1d 0x44065500 + 0x30170e1d 0xec101500 0x31000e1d 0x0442d500 + 0x10000a1d 0x2440d280 0x30070205 0x8c000780 + 0x00000805 0xc0000780 0x30218205 0x00000003 + 0x04005a01 0xe4204780 0xd4016809 0x20000780 + 0x1800e005 0x0423c780 0x3801c005 0x8c200780 + 0x04005a01 0xe4204780 0x1800d005 0x0423c780 0x3801c005 0x8c200780 0x04005a01 0xe4204780 - 0x1800d005 0x0423c780 0x3801c005 0x8c200780 - 0x04005a01 0xe4204780 0x1800c805 0x0423c780 + 0x1800c805 0x0423c780 0x3801c005 0x8c200780 + 0x04005a01 0xe4204780 0x1800c405 0x0423c780 0x3801c005 0x8c200780 0x04005a01 0xe4204780 - 0x1800c405 0x0423c780 0x3801c005 0x8c200780 - 0x04005a01 0xe4204780 0x1800c205 0x0423c780 - 0x3801c005 0x8c200780 0x04005a01 0xe4204780 - 0xa00bb003 0x00000000 0x100bb003 0x00000100 - 0x30070405 0xc4100780 0x30060409 0xc4100780 - 0x20000205 0x04008780 0xd0016805 0x20000780 - 0x2101e808 0x1500e004 0x200c8409 0x00000003 - 0xd00e0405 0xa0c00780 0xf0000001 0xe0000002 - 0x20018c19 0x00000003 0x3006cffd 0x6c2147d8 - 0x1003b003 0x00001280 0xf0000001 0xe0000001 + 0x1800c205 0x0423c780 0x3801c005 0x8c200780 + 0x04005a01 0xe4204780 0xa00e2003 0x00000000 + 0x100e2003 0x00000100 0x30070605 0xc4100780 + 0x3006060d 0xc4100780 0x20000205 0x0400c780 + 0xd0016805 0x20000780 0x2101e80c 0x1500e004 + 0x200c860d 0x00000003 0xd00e0605 0xa0c00780 + 0xf0000001 0xe0000002 0x20018c19 0x00000003 + 0x3006cffd 0x6c2147d8 0x10062003 0x00001280 + 0xf0000001 0xe0000001 } } code {