diff --git a/CUETools.Codecs.FLACCL/CUETools.Codecs.FLACCL.csproj b/CUETools.Codecs.FLACCL/CUETools.Codecs.FLACCL.csproj index ff9d027..8e31c17 100644 --- a/CUETools.Codecs.FLACCL/CUETools.Codecs.FLACCL.csproj +++ b/CUETools.Codecs.FLACCL/CUETools.Codecs.FLACCL.csproj @@ -65,9 +65,6 @@ - - PreserveNewest - diff --git a/CUETools.Codecs.FLACCL/FLACCLWriter.cs b/CUETools.Codecs.FLACCL/FLACCLWriter.cs index 205bbea..1fe461b 100644 --- a/CUETools.Codecs.FLACCL/FLACCLWriter.cs +++ b/CUETools.Codecs.FLACCL/FLACCLWriter.cs @@ -826,29 +826,48 @@ namespace CUETools.Codecs.FLACCL return size; } - unsafe void output_residual(FlacFrame frame, FlacSubframeInfo sub) + unsafe void output_residual(FLACCLTask task, FlacSubframeInfo sub, int offs0, int index) { + FlacFrame frame = task.frame; + // rice-encoded block frame.writer.writebits(2, 0); - // partition order int porder = sub.best.rc.porder; - int psize = frame.blocksize >> porder; //assert(porder >= 0); frame.writer.writebits(4, porder); - int res_cnt = psize - sub.best.order; - // residual - int j = sub.best.order; - fixed (byte* fixbuf = frame.writer.Buffer) - for (int p = 0; p < (1 << porder); p++) + if (task.riceOnGPU) { - int k = sub.best.rc.rparams[p]; - frame.writer.writebits(4, k); - if (p == 1) res_cnt = psize; - int cnt = Math.Min(res_cnt, frame.blocksize - j); - frame.writer.write_rice_block_signed(fixbuf, k, sub.best.residual + j, cnt); - j += cnt; + if (task.BestResidualTasks[index].size != (int)sub.best.size) + throw new Exception("Encoding offset mismatch"); + if (task.BestResidualTasks[index].headerLen != offs0 + 6) + throw new Exception("Encoding offset mismatch"); + if (task.BestResidualTasks[index].encodingOffset != frame.writer.BitLength) + throw new Exception("Encoding offset mismatch"); + int len = task.BestResidualTasks[index].size - task.BestResidualTasks[index].headerLen; + // task.BestResidualTasks[index].headerLen + frame.writer.writeints(len, (byte*)task.clRiceOutputPtr); + if (task.BestResidualTasks[index].encodingOffset + len != frame.writer.BitLength) + throw new Exception("Encoding offset mismatch"); + } + else + { + int psize = frame.blocksize >> porder; + int res_cnt = psize - sub.best.order; + + // residual + int j = sub.best.order; + fixed (byte* fixbuf = frame.writer.Buffer) + for (int p = 0; p < (1 << porder); p++) + { + int k = sub.best.rc.rparams[p]; + frame.writer.writebits(4, k); + if (p == 1) res_cnt = psize; + int cnt = Math.Min(res_cnt, frame.blocksize - j); + frame.writer.write_rice_block_signed(fixbuf, k, sub.best.residual + j, cnt); + j += cnt; + } } } @@ -868,14 +887,16 @@ namespace CUETools.Codecs.FLACCL } unsafe void - output_subframe_fixed(FlacFrame frame, FlacSubframeInfo sub) + output_subframe_fixed(FLACCLTask task, FlacSubframeInfo sub, int index) { + FlacFrame frame = task.frame; + // warm-up samples for (int i = 0; i < sub.best.order; i++) frame.writer.writebits_signed(sub.obits, sub.samples[i]); // residual - output_residual(frame, sub); + output_residual(task, sub, sub.obits * sub.best.order, index); } unsafe uint @@ -904,8 +925,10 @@ namespace CUETools.Codecs.FLACCL } unsafe void - output_subframe_lpc(FlacFrame frame, FlacSubframeInfo sub) + output_subframe_lpc(FLACCLTask task, FlacSubframeInfo sub, int index) { + FlacFrame frame = task.frame; + // warm-up samples for (int i = 0; i < sub.best.order; i++) frame.writer.writebits_signed(sub.obits, sub.samples[i]); @@ -917,11 +940,12 @@ namespace CUETools.Codecs.FLACCL frame.writer.writebits_signed(sub.best.cbits, sub.best.coefs[i]); // residual - output_residual(frame, sub); + output_residual(task, sub, (sub.obits + sub.best.cbits) * sub.best.order + 9, index); } - unsafe void output_subframes(FlacFrame frame) + unsafe void output_subframes(FLACCLTask task, int iFrame) { + FlacFrame frame = task.frame; for (int ch = 0; ch < channels; ch++) { FlacSubframeInfo sub = frame.subframes[ch]; @@ -940,6 +964,8 @@ namespace CUETools.Codecs.FLACCL //if (frame_writer.Length >= frame_buffer.Length) // throw new Exception("buffer overflow"); + int index = ch + iFrame * channels; + // subframe switch (sub.best.type) { @@ -950,10 +976,10 @@ namespace CUETools.Codecs.FLACCL output_subframe_verbatim(frame, sub); break; case SubframeType.Fixed: - output_subframe_fixed(frame, sub); + output_subframe_fixed(task, sub, index); break; case SubframeType.LPC: - output_subframe_lpc(frame, sub); + output_subframe_lpc(task, sub, index); break; } //if (frame_writer.Length >= frame_buffer.Length) @@ -1049,7 +1075,6 @@ namespace CUETools.Codecs.FLACCL task.ResidualTasks[task.nResidualTasks].residualOrder = order + 1; task.ResidualTasks[task.nResidualTasks].samplesOffs = ch * FLACCLWriter.MAX_BLOCKSIZE + iFrame * blocksize; task.ResidualTasks[task.nResidualTasks].residualOffs = task.ResidualTasks[task.nResidualTasks].samplesOffs; - task.ResidualTasks[task.nResidualTasks].ignore = 0; task.ResidualTasks[task.nResidualTasks].wbits = 0; task.ResidualTasks[task.nResidualTasks].size = task.ResidualTasks[task.nResidualTasks].obits * blocksize; task.nResidualTasks++; @@ -1065,7 +1090,6 @@ namespace CUETools.Codecs.FLACCL task.ResidualTasks[task.nResidualTasks].blocksize = blocksize; task.ResidualTasks[task.nResidualTasks].samplesOffs = ch * FLACCLWriter.MAX_BLOCKSIZE + iFrame * blocksize; task.ResidualTasks[task.nResidualTasks].residualOffs = task.ResidualTasks[task.nResidualTasks].samplesOffs; - task.ResidualTasks[task.nResidualTasks].ignore = 0; task.ResidualTasks[task.nResidualTasks].wbits = 0; task.ResidualTasks[task.nResidualTasks].size = task.ResidualTasks[task.nResidualTasks].obits * blocksize; task.ResidualTasks[task.nResidualTasks].residualOrder = 1; @@ -1084,7 +1108,6 @@ namespace CUETools.Codecs.FLACCL task.ResidualTasks[task.nResidualTasks].residualOrder = order; task.ResidualTasks[task.nResidualTasks].samplesOffs = ch * FLACCLWriter.MAX_BLOCKSIZE + iFrame * blocksize; task.ResidualTasks[task.nResidualTasks].residualOffs = task.ResidualTasks[task.nResidualTasks].samplesOffs; - task.ResidualTasks[task.nResidualTasks].ignore = 0; task.ResidualTasks[task.nResidualTasks].wbits = 0; task.ResidualTasks[task.nResidualTasks].size = task.ResidualTasks[task.nResidualTasks].obits * blocksize; task.ResidualTasks[task.nResidualTasks].shift = 0; @@ -1173,7 +1196,7 @@ namespace CUETools.Codecs.FLACCL #if DEBUG // check size - if (_settings.GPUOnly) + if (_settings.GPUOnly && !task.riceOnGPU) { uint real_size = measure_subframe(task.frame, task.frame.subframes[ch]); if (real_size != task.frame.subframes[ch].best.size) @@ -1181,7 +1204,7 @@ namespace CUETools.Codecs.FLACCL } #endif - if ((csum << task.frame.subframes[ch].obits) >= 1UL << 32 || !_settings.GPUOnly) + if (((csum << task.frame.subframes[ch].obits) >= 1UL << 32 || !_settings.GPUOnly) && !task.riceOnGPU) { if (!unpacked) unpack_samples(task, task.frameSize); unpacked = true; if ((csum << task.frame.subframes[ch].obits) >= 1UL << 32) @@ -1250,9 +1273,12 @@ namespace CUETools.Codecs.FLACCL frame.subframes[ch].best.size = (uint)(frame.subframes[ch].obits * frame.blocksize); frame.subframes[ch].wbits = 0; + if (frame.blocksize <= Math.Max(4, eparams.max_prediction_order)) + continue; if (task.BestResidualTasks[index].size < 0) throw new Exception("internal error"); - if (frame.blocksize > Math.Max(4, eparams.max_prediction_order) && frame.subframes[ch].best.size > task.BestResidualTasks[index].size) + if (frame.subframes[ch].best.size > task.BestResidualTasks[index].size + && (SubframeType)task.BestResidualTasks[index].type != SubframeType.Verbatim) { frame.subframes[ch].best.type = (SubframeType)task.BestResidualTasks[index].type; frame.subframes[ch].best.size = (uint)task.BestResidualTasks[index].size; @@ -1261,11 +1287,12 @@ namespace CUETools.Codecs.FLACCL frame.subframes[ch].best.shift = task.BestResidualTasks[index].shift; frame.subframes[ch].obits -= task.BestResidualTasks[index].wbits; frame.subframes[ch].wbits = task.BestResidualTasks[index].wbits; - frame.subframes[ch].best.rc.porder = task.BestResidualTasks[index].porder; for (int i = 0; i < task.BestResidualTasks[index].residualOrder; i++) frame.subframes[ch].best.coefs[i] = task.BestResidualTasks[index].coefs[task.BestResidualTasks[index].residualOrder - 1 - i]; - if (_settings.GPUOnly && (frame.subframes[ch].best.type == SubframeType.Fixed || frame.subframes[ch].best.type == SubframeType.LPC)) + if (_settings.GPUOnly && !task.riceOnGPU && (frame.subframes[ch].best.type == SubframeType.Fixed || frame.subframes[ch].best.type == SubframeType.LPC)) + //if (_settings.GPUOnly && (frame.subframes[ch].best.type == SubframeType.Fixed || frame.subframes[ch].best.type == SubframeType.LPC)) { + frame.subframes[ch].best.rc.porder = task.BestResidualTasks[index].porder; int* riceParams = ((int*)task.clBestRiceParamsPtr) + (index << task.max_porder); fixed (int* dstParams = frame.subframes[ch].best.rc.rparams) AudioSamples.MemCpy(dstParams, riceParams, (1 << frame.subframes[ch].best.rc.porder)); @@ -1276,6 +1303,11 @@ namespace CUETools.Codecs.FLACCL throw new Exception("size reported incorrectly"); } } + else + { + if (task.riceOnGPU) + throw new Exception("size reported incorrectly"); + } } } @@ -1392,7 +1424,7 @@ namespace CUETools.Codecs.FLACCL task.frame.writer_offset = task.frame.writer.Length; output_frame_header(task.frame); - output_subframes(task.frame); + output_subframes(task, iFrame); output_frame_footer(task.frame); if (task.frame.writer.Length - task.frame.writer_offset >= max_frame_size) throw new Exception("buffer overflow"); @@ -1619,6 +1651,7 @@ namespace CUETools.Codecs.FLACCL _IO = new FileStream(_path, FileMode.Create, FileAccess.Write, FileShare.Read); int header_size = flake_encode_init(); _IO.Write(header, 0, header_size); + _totalSize += header_size; if (_IO.CanSeek) first_frame_offset = _IO.Position; @@ -2250,8 +2283,8 @@ namespace CUETools.Codecs.FLACCL public int wbits; public int abits; public int porder; - public int ignore; - public int reserved; + public int headerLen; + public int encodingOffset; public fixed int coefs[32]; }; @@ -2276,6 +2309,8 @@ namespace CUETools.Codecs.FLACCL public Kernel clSumPartition; public Kernel clFindRiceParameter; public Kernel clFindPartitionOrder; + public Kernel clCalcOutputOffsets; + public Kernel clRiceEncoding; public Mem clSamplesBytes; public Mem clSamples; public Mem clLPCData; @@ -2283,6 +2318,7 @@ namespace CUETools.Codecs.FLACCL public Mem clPartitions; public Mem clRiceParams; public Mem clBestRiceParams; + public Mem clRiceOutput; public Mem clAutocorOutput; public Mem clSelectedTasks; public Mem clSelectedTasksSecondEstimate; @@ -2298,6 +2334,7 @@ namespace CUETools.Codecs.FLACCL public Mem clBestResidualTasksPinned; public Mem clWindowFunctionsPinned; public Mem clSelectedTasksPinned; + public Mem clRiceOutputPinned; public IntPtr clSamplesBytesPtr; public IntPtr clResidualPtr; @@ -2306,6 +2343,7 @@ namespace CUETools.Codecs.FLACCL public IntPtr clBestResidualTasksPtr; public IntPtr clWindowFunctionsPtr; public IntPtr clSelectedTasksPtr; + public IntPtr clRiceOutputPtr; public int[] samplesBuffer; public byte[] outputBuffer; @@ -2317,7 +2355,6 @@ namespace CUETools.Codecs.FLACCL public FlacFrame frame; public int residualTasksLen; public int bestResidualTasksLen; - public int samplesBufferLen; public int nResidualTasks = 0; public int nResidualTasksPerChannel = 0; public int nEstimateTasksPerChannel = 0; @@ -2325,6 +2362,8 @@ namespace CUETools.Codecs.FLACCL public int nWindowFunctions = 0; public int max_porder = 0; + public bool riceOnGPU = false; + public FlakeReader verify; public Thread workThread = null; @@ -2353,7 +2392,8 @@ namespace CUETools.Codecs.FLACCL int MAX_ORDER = this.writer.eparams.max_prediction_order; residualTasksLen = sizeof(FLACCLSubframeTask) * 32 * channelsCount * FLACCLWriter.maxFrames; bestResidualTasksLen = sizeof(FLACCLSubframeTask) * channels * FLACCLWriter.maxFrames; - samplesBufferLen = sizeof(int) * FLACCLWriter.MAX_BLOCKSIZE * channelsCount; + int samplesBufferLen = sizeof(int) * FLACCLWriter.MAX_BLOCKSIZE * channelsCount; + int residualBufferLen = sizeof(int) * FLACCLWriter.MAX_BLOCKSIZE * channelsCount; // *channels! but need to adjust residualOffser int partitionsLen = sizeof(int) * (30 << 8) * channels * FLACCLWriter.maxFrames; int riceParamsLen = sizeof(int) * (4 << 8) * channels * FLACCLWriter.maxFrames; int autocorLen = sizeof(float) * (MAX_ORDER + 1) * lpc.MAX_LPC_WINDOWS * channelsCount * FLACCLWriter.maxFrames; @@ -2361,50 +2401,56 @@ namespace CUETools.Codecs.FLACCL int resOutLen = sizeof(int) * channelsCount * (lpc.MAX_LPC_WINDOWS * lpc.MAX_LPC_ORDER + 8) * FLACCLWriter.maxFrames; int wndLen = sizeof(float) * FLACCLWriter.MAX_BLOCKSIZE /** 2*/ * lpc.MAX_LPC_WINDOWS; int selectedLen = sizeof(int) * 32 * channelsCount * FLACCLWriter.maxFrames; + int riceLen = sizeof(int) * channelsCount * FLACCLWriter.MAX_BLOCKSIZE; if (!writer._settings.MappedMemory) { clSamplesBytes = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, samplesBufferLen / 2); - clResidual = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, samplesBufferLen); + clResidual = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, residualBufferLen); clBestRiceParams = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, riceParamsLen / 4); clResidualTasks = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, residualTasksLen); clBestResidualTasks = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, bestResidualTasksLen); clWindowFunctions = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, wndLen); clSelectedTasks = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, selectedLen); + clRiceOutput = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, riceLen); clSamplesBytesPinned = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, samplesBufferLen / 2); - clResidualPinned = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, samplesBufferLen); + clResidualPinned = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, residualBufferLen); clBestRiceParamsPinned = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, riceParamsLen / 4); clResidualTasksPinned = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, residualTasksLen); clBestResidualTasksPinned = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, bestResidualTasksLen); clWindowFunctionsPinned = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, wndLen); clSelectedTasksPinned = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, selectedLen); + clRiceOutputPinned = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, riceLen); clSamplesBytesPtr = openCLCQ.EnqueueMapBuffer(clSamplesBytesPinned, true, MapFlags.WRITE, 0, samplesBufferLen / 2); - clResidualPtr = openCLCQ.EnqueueMapBuffer(clResidualPinned, true, MapFlags.WRITE, 0, samplesBufferLen); + clResidualPtr = openCLCQ.EnqueueMapBuffer(clResidualPinned, true, MapFlags.WRITE, 0, residualBufferLen); clBestRiceParamsPtr = openCLCQ.EnqueueMapBuffer(clBestRiceParamsPinned, true, MapFlags.WRITE, 0, riceParamsLen / 4); clResidualTasksPtr = openCLCQ.EnqueueMapBuffer(clResidualTasksPinned, true, MapFlags.WRITE, 0, residualTasksLen); clBestResidualTasksPtr = openCLCQ.EnqueueMapBuffer(clBestResidualTasksPinned, true, MapFlags.WRITE, 0, bestResidualTasksLen); clWindowFunctionsPtr = openCLCQ.EnqueueMapBuffer(clWindowFunctionsPinned, true, MapFlags.WRITE, 0, wndLen); clSelectedTasksPtr = openCLCQ.EnqueueMapBuffer(clSelectedTasksPinned, true, MapFlags.WRITE, 0, selectedLen); + clRiceOutputPtr = openCLCQ.EnqueueMapBuffer(clRiceOutputPinned, true, MapFlags.WRITE, 0, riceLen); } else { clSamplesBytes = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, (uint)samplesBufferLen / 2); - clResidual = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, samplesBufferLen); + clResidual = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, residualBufferLen); clBestRiceParams = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, riceParamsLen / 4); clResidualTasks = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, residualTasksLen); clBestResidualTasks = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, bestResidualTasksLen); clWindowFunctions = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, wndLen); clSelectedTasks = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, selectedLen); + clRiceOutput = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, riceLen); clSamplesBytesPtr = openCLCQ.EnqueueMapBuffer(clSamplesBytes, true, MapFlags.WRITE, 0, samplesBufferLen / 2); - clResidualPtr = openCLCQ.EnqueueMapBuffer(clResidual, true, MapFlags.WRITE, 0, samplesBufferLen); + clResidualPtr = openCLCQ.EnqueueMapBuffer(clResidual, true, MapFlags.WRITE, 0, residualBufferLen); clBestRiceParamsPtr = openCLCQ.EnqueueMapBuffer(clBestRiceParams, true, MapFlags.WRITE, 0, riceParamsLen / 4); clResidualTasksPtr = openCLCQ.EnqueueMapBuffer(clResidualTasks, true, MapFlags.WRITE, 0, residualTasksLen); clBestResidualTasksPtr = openCLCQ.EnqueueMapBuffer(clBestResidualTasks, true, MapFlags.WRITE, 0, bestResidualTasksLen); clWindowFunctionsPtr = openCLCQ.EnqueueMapBuffer(clWindowFunctions, true, MapFlags.WRITE, 0, wndLen); clSelectedTasksPtr = openCLCQ.EnqueueMapBuffer(clSelectedTasks, true, MapFlags.WRITE, 0, selectedLen); + clRiceOutputPtr = openCLCQ.EnqueueMapBuffer(clRiceOutput, true, MapFlags.WRITE, 0, riceLen); //clSamplesBytesPtr = clSamplesBytes.HostPtr; //clResidualPtr = clResidual.HostPtr; @@ -2446,6 +2492,11 @@ namespace CUETools.Codecs.FLACCL clSumPartition = openCLProgram.CreateKernel("clSumPartition"); clFindRiceParameter = openCLProgram.CreateKernel("clFindRiceParameter"); clFindPartitionOrder = openCLProgram.CreateKernel("clFindPartitionOrder"); + if (riceOnGPU) + { + clCalcOutputOffsets = openCLProgram.CreateKernel("clCalcOutputOffsets"); + clRiceEncoding = openCLProgram.CreateKernel("clRiceEncoding"); + } } samplesBuffer = new int[FLACCLWriter.MAX_BLOCKSIZE * channelsCount]; @@ -2494,6 +2545,11 @@ namespace CUETools.Codecs.FLACCL clSumPartition.Dispose(); clFindRiceParameter.Dispose(); clFindPartitionOrder.Dispose(); + if (riceOnGPU) + { + clCalcOutputOffsets.Dispose(); + clRiceEncoding.Dispose(); + } clPartitions.Dispose(); clRiceParams.Dispose(); @@ -2522,6 +2578,9 @@ namespace CUETools.Codecs.FLACCL if (clSelectedTasksPtr != IntPtr.Zero) openCLCQ.EnqueueUnmapMemObject(clSelectedTasksPinned, clSelectedTasksPtr); clSelectedTasksPtr = IntPtr.Zero; + if (clRiceOutputPtr != IntPtr.Zero) + openCLCQ.EnqueueUnmapMemObject(clRiceOutputPinned, clRiceOutputPtr); + clRiceOutputPtr = IntPtr.Zero; clSamplesBytesPinned.Dispose(); clResidualPinned.Dispose(); @@ -2530,6 +2589,7 @@ namespace CUETools.Codecs.FLACCL clBestResidualTasksPinned.Dispose(); clWindowFunctionsPinned.Dispose(); clSelectedTasksPinned.Dispose(); + clRiceOutputPinned.Dispose(); } else { @@ -2540,6 +2600,7 @@ namespace CUETools.Codecs.FLACCL openCLCQ.EnqueueUnmapMemObject(clBestResidualTasks, clBestResidualTasksPtr); openCLCQ.EnqueueUnmapMemObject(clWindowFunctions, clWindowFunctionsPtr); openCLCQ.EnqueueUnmapMemObject(clSelectedTasks, clSelectedTasksPtr); + openCLCQ.EnqueueUnmapMemObject(clRiceOutput, clRiceOutputPtr); } clSamples.Dispose(); @@ -2553,6 +2614,7 @@ namespace CUETools.Codecs.FLACCL clBestResidualTasks.Dispose(); clWindowFunctions.Dispose(); clSelectedTasks.Dispose(); + clRiceOutput.Dispose(); openCLCQ.Dispose(); @@ -2813,10 +2875,43 @@ namespace CUETools.Codecs.FLACCL groupSize, channels * frameCount); - if (!writer._settings.MappedMemory) - openCLCQ.EnqueueReadBuffer(clBestRiceParams, false, 0, sizeof(int) * (1 << max_porder) * channels * frameCount, clBestRiceParamsPtr); - if (!writer._settings.MappedMemory) - openCLCQ.EnqueueReadBuffer(clResidual, false, 0, sizeof(int) * FLACCLWriter.MAX_BLOCKSIZE * channels, clResidualPtr); + if (riceOnGPU) + { + clCalcOutputOffsets.SetArgs( + clResidual, + clSamples, + clBestResidualTasks, + channels, + frameCount, + frameNumber); + + openCLCQ.EnqueueNDRangeKernel( + clCalcOutputOffsets, + groupSize, + 1); + + clRiceEncoding.SetArgs( + clResidual, + clSamples, + clBestRiceParams, + clBestResidualTasks, + clRiceOutput, + max_porder); + + openCLCQ.EnqueueNDRangeKernel( + clRiceEncoding, + groupSize, + channels * frameCount); + } + + if (!writer._settings.MappedMemory) + { + openCLCQ.EnqueueReadBuffer(clBestRiceParams, false, 0, sizeof(int) * (1 << max_porder) * channels * frameCount, clBestRiceParamsPtr); + if (riceOnGPU) + openCLCQ.EnqueueReadBuffer(clRiceOutput, false, 0, sizeof(int) * FLACCLWriter.MAX_BLOCKSIZE * channels, clRiceOutputPtr); + else + openCLCQ.EnqueueReadBuffer(clResidual, false, 0, sizeof(int) * FLACCLWriter.MAX_BLOCKSIZE * channels, clResidualPtr); + } } if (!writer._settings.MappedMemory) openCLCQ.EnqueueReadBuffer(clBestResidualTasks, false, 0, sizeof(FLACCLSubframeTask) * channels * frameCount, clBestResidualTasksPtr); diff --git a/CUETools.Codecs.FLACCL/flac.cl b/CUETools.Codecs.FLACCL/flac.cl index b6e2e73..0df9604 100644 --- a/CUETools.Codecs.FLACCL/flac.cl +++ b/CUETools.Codecs.FLACCL/flac.cl @@ -63,8 +63,8 @@ typedef struct int wbits; int abits; int porder; - int ignore; - int reserved; + int headerLen; + int encodingOffset; } FLACCLSubframeData; typedef struct @@ -1371,7 +1371,7 @@ void clFindRiceParameter( __global FLACCLSubframeTask* task = tasks + get_group_id(0); const int tid = get_local_id(0); int lim = (2 << max_porder) - 1; - int psize = task->data.blocksize >> max_porder; + //int psize = task->data.blocksize >> max_porder; int bs = task->data.blocksize; int ro = task->data.residualOrder; for (int offs = 0; offs < lim; offs ++) @@ -1474,10 +1474,18 @@ void clFindPartitionOrder( int obits = task->data.obits - task->data.wbits; task->data.porder = best_porder; - task->data.size = - task->data.type == Fixed ? task->data.residualOrder * obits + 6 + best_length : - task->data.type == LPC ? task->data.residualOrder * obits + 6 + best_length + 4 + 5 + task->data.residualOrder * task->data.cbits : - task->data.type == Constant ? obits : obits * task->data.blocksize; + task->data.headerLen = + task->data.type == Constant ? obits : + task->data.type == Verbatim ? obits * task->data.blocksize : + task->data.type == Fixed ? task->data.residualOrder * obits + 6 : + task->data.type == LPC ? task->data.residualOrder * obits + 6 + 4 + 5 + task->data.residualOrder * task->data.cbits : 0; + task->data.size = + task->data.headerLen + ((task->data.type == Fixed || task->data.type == LPC) ? best_length : 0); + if (task->data.size >= obits * task->data.blocksize) + { + task->data.headerLen = task->data.size = obits * task->data.blocksize; + task->data.type = Verbatim; + } for (int offs = 0; offs < (1 << best_porder); offs ++) best_rice_parameters[(get_group_id(0) << max_porder) + offs] = rice_parameters[pos - (2 << best_porder) + offs]; } @@ -1536,4 +1544,182 @@ void clFindPartitionOrder( // FIXME: should be bytes? } #endif + +#ifdef __CPU__ +typedef struct BitWriter_t +{ + __global int *buffer; + unsigned int bit_buf; + int bit_left; + int buf_ptr; +} BitWriter; + +inline void writebits(BitWriter *bw, int bits, int v) +{ + uint val = ((uint)v) & ((1 << bits) - 1); + if (bits < bw->bit_left) + { + bw->bit_buf = (bw->bit_buf << bits) | val; + bw->bit_left -= bits; + } + else + { +// if (bits >= 32) printf("\n\n\n\n-------------------------\n\n\n"); + unsigned int bb = (bw->bit_buf << bw->bit_left) | (val >> (bits - bw->bit_left)); + bw->buffer[bw->buf_ptr++] = (bb >> 24) | ((bb >> 8) & 0xff00) | ((bb << 8) & 0xff0000) | ((bb << 24) & 0xff000000); + bw->bit_left += (32 - bits); + bw->bit_buf = val; +// bw->bit_buf = val & ((1 << (32 - bw->bit_left)) - 1); + } +} + +inline void flush(BitWriter *bw) +{ + if (bw->bit_left < 32) + writebits(bw, bw->bit_left, 0); +} +#endif + +inline int len_utf8(int n) +{ + int bts = 31 - clz(n); + if (bts < 7) + return 8; + return 8 * ((bts + 4) / 5); +} + +// get_global_id(0) * channels == task index +__kernel +void clCalcOutputOffsets( + __global int *residual, + __global int *samples, + __global FLACCLSubframeTask *tasks, + int channels, + int frameCount, + int firstFrame + ) +{ + int offset = 0; + for (int iFrame = 0; iFrame < frameCount; iFrame++) + { + //printf("len_utf8(%d) == %d\n", firstFrame + iFrame, len_utf8(firstFrame + iFrame)); + offset += 15 + 1 + 4 + 4 + 4 + 3 + 1 + len_utf8(firstFrame + iFrame) + // + 8-16 // custom block size + // + 8-16 // custom sample rate + ; + int bs = tasks[iFrame * channels].data.blocksize; + //public static readonly int[] flac_blocksizes = new int[15] { 0, 192, 576, 1152, 2304, 4608, 0, 0, 256, 512, 1024, 2048, 4096, 8192, 16384 }; + if (bs != 4096 && bs != 4608) // TODO: check all other standard sizes + offset += select(8, 16, bs >= 256); + + // assert (offset % 8) == 0 + offset += 8; + for (int ch = 0; ch < channels; ch++) + { + __global FLACCLSubframeTask* task = tasks + iFrame * channels + ch; + offset += 8 + task->data.wbits; + task->data.encodingOffset = offset + task->data.headerLen; + offset += task->data.size; + } + offset = (offset + 7) & ~7; + offset += 16; + } +} + +// get_group_id(0) == task index +__kernel __attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1))) +void clRiceEncoding( + __global int *residual, + __global int *samples, + __global int* best_rice_parameters, + __global FLACCLSubframeTask *tasks, + __global int* output, + int max_porder + ) +{ +#ifdef __CPU__ + __global FLACCLSubframeTask* task = tasks + get_group_id(0); + if (task->data.type == Fixed || task->data.type == LPC) + { + int ro = task->data.residualOrder; + int bs = task->data.blocksize; + int porder = task->data.porder; + int psize = bs >> porder; + + BitWriter bw; + bw.buffer = output; + bw.buf_ptr = task->data.encodingOffset / 32; + bw.bit_left = 32 - (task->data.encodingOffset & 31); + bw.bit_buf = 0; + + //if (get_group_id(0) == 0) printf("%d\n", offs); + + int res_cnt = psize - ro; + // residual + int j = ro; + __global int * kptr = &best_rice_parameters[get_group_id(0) << max_porder]; + for (int p = 0; p < (1 << porder); p++) + { + int k = kptr[p]; + writebits(&bw, 4, k); + //if (get_group_id(0) == 0) printf("[%x] ", k); + //if (get_group_id(0) == 0) printf("(%x) ", bw.bit_buf); + if (p == 1) res_cnt = psize; + int cnt = min(res_cnt, bs - j); + for (int i = 0; i < cnt; i++) + { + int v = residual[task->data.residualOffs + j + i]; + v = (v << 1) ^ (v >> 31); + // write quotient in unary + int q = (v >> k) + 1; + int bits = k + q; + while (bits > 31) + { + int b = min(bits - 31, 31); + if (b < bw.bit_left) + { + bw.bit_buf <<= b; + bw.bit_left -= b; + } + else + { + unsigned int bb = bw.bit_buf << bw.bit_left; + bw.bit_buf = 0; + bw.bit_left += (32 - b); + bw.buffer[bw.buf_ptr++] = (bb >> 24) | ((bb >> 8) & 0xff00) | ((bb << 8) & 0xff0000) | ((bb << 24) & 0xff000000); + } + bits -= b; + } + unsigned int val = (unsigned int)((v & ((1 << k) - 1)) | (1 << k)); + if (bits < bw.bit_left) + { + bw.bit_buf = (bw.bit_buf << bits) | val; + bw.bit_left -= bits; + } + else + { + unsigned int bb = (bw.bit_buf << bw.bit_left) | (val >> (bits - bw.bit_left)); + bw.bit_buf = val; + bw.bit_left += (32 - bits); + bw.buffer[bw.buf_ptr++] = (bb >> 24) | ((bb >> 8) & 0xff00) | ((bb << 8) & 0xff0000) | ((bb << 24) & 0xff000000); + } + ////if (get_group_id(0) == 0) printf("%x ", v); + //writebits(&bw, (v >> k) + 1, 1); + ////if (get_group_id(0) == 0) printf("(%x) ", bw.bit_buf); + //writebits(&bw, k, v); + ////if (get_group_id(0) == 0) printf("(%x) ", bw.bit_buf); + } + j += cnt; + } + //if (bw.buf_ptr * 32 + 32 - bw.bit_left != task->data.encodingOffset - task->data.headerLen + task->data.size) + // printf("bit length mismatch: encodingOffset == %d, headerLen == %d, size == %d, so should be %d, but is %d\n", + // task->data.encodingOffset, task->data.headerLen, task->data.size, + // task->data.encodingOffset - task->data.headerLen + task->data.size, + // bw.buf_ptr * 32 + 32 - bw.bit_left + // ); + //if (get_group_id(0) == 0) printf("\n"); + flush(&bw); + } +#endif +} #endif diff --git a/CUETools.Codecs.FLACCL/flaccpu.cl b/CUETools.Codecs.FLACCL/flaccpu.cl deleted file mode 100644 index 676dcf1..0000000 --- a/CUETools.Codecs.FLACCL/flaccpu.cl +++ /dev/null @@ -1,759 +0,0 @@ -/** - * CUETools.FLACCL: FLAC audio encoder using OpenCL - * Copyright (c) 2010 Gregory S. Chudov - * - * This library is free software; you can redistribute it and/or - * modify it under the terms of the GNU Lesser General Public - * License as published by the Free Software Foundation; either - * version 2.1 of the License, or (at your option) any later version. - * - * This library 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 - * Lesser General Public License for more details. - * - * You should have received a copy of the GNU Lesser General Public - * License along with this library; if not, write to the Free Software - * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA - */ - -#ifndef _FLACCL_KERNEL_H_ -#define _FLACCL_KERNEL_H_ - -#ifdef DEBUG -#pragma OPENCL EXTENSION cl_amd_printf : enable -#endif - -#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable - -#pragma OPENCL EXTENSION cl_amd_fp64 : enable - -typedef enum -{ - Constant = 0, - Verbatim = 1, - Fixed = 8, - LPC = 32 -} SubframeType; - -typedef struct -{ - int residualOrder; // <= 32 - int samplesOffs; - int shift; - int cbits; - int size; - int type; - int obits; - int blocksize; - int best_index; - int channel; - int residualOffs; - int wbits; - int abits; - int porder; - int ignore; - int reserved; -} FLACCLSubframeData; - -typedef struct -{ - FLACCLSubframeData data; - int coefs[32]; // fixme: should be short? -} FLACCLSubframeTask; - -__kernel void clStereoDecorr( - __global int4 *samples, - __global int4 *src, - int offset -) -{ - int pos = get_global_id(0); - if (pos < offset) - { - int4 s = src[pos]; - int4 x = (s << 16) >> 16; - int4 y = s >> 16; - samples[pos] = x; - samples[1 * offset + pos] = y; - samples[2 * offset + pos] = (x + y) >> 1; - samples[3 * offset + pos] = x - y; - } -} - -__kernel void clWindowRectangle(__global float* window, int windowOffset) -{ - window[get_global_id(0)] = 1.0f; -} - -__kernel void clWindowFlattop(__global float* window, int windowOffset) -{ - float p = M_PI * get_global_id(0) / (get_global_size(0) - 1); - window[get_global_id(0)] = 1.0f - - 1.93f * cos(2 * p) - + 1.29f * cos(4 * p) - - 0.388f * cos(6 * p) - + 0.0322f * cos(8 * p); -} - -__kernel void clWindowTukey(__global float* window, int windowOffset, float p) -{ - int Np = (int)(p / 2.0f * get_global_size(0)) - 1; - int n = select(max(Np, get_global_id(0) - (get_global_size(0) - Np - 1) + Np), get_global_id(0), get_global_id(0) <= Np); - window[get_global_id(0)] = 0.5f - 0.5f * cos(M_PI * n / Np); -} - -__kernel void clChannelDecorr2( - __global int4 *samples, - __global int4 *src, - int offset -) -{ - int pos = get_global_id(0); - if (pos < offset) - { - int4 s = src[pos]; - int4 x = (s << 16) >> 16; - int4 y = s >> 16; - samples[pos] = x; - samples[1 * offset + pos] = y; - } -} - -//__kernel void clChannelDecorr( -// int *samples, -// short *src, -// int offset -//) -//{ -// int pos = get_global_id(0); -// if (pos < offset) -// samples[get_group_id(1) * offset + pos] = src[pos * get_num_groups(1) + get_group_id(1)]; -//} - -#define __ffs(a) (32 - clz(a & (-a))) -//#define __ffs(a) (33 - clz(~a & (a - 1))) - -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clFindWastedBits( - __global FLACCLSubframeTask *tasks, - __global int *samples, - int tasksPerChannel -) -{ - __global FLACCLSubframeTask* ptask = &tasks[get_group_id(0) * tasksPerChannel]; - int w = 0, a = 0; - for (int pos = 0; pos < ptask->data.blocksize; pos ++) - { - int smp = samples[ptask->data.samplesOffs + pos]; - w |= smp; - a |= smp ^ (smp >> 31); - } - w = max(0,__ffs(w) - 1); - a = 32 - clz(a) - w; - for (int i = 0; i < tasksPerChannel; i++) - { - ptask[i].data.wbits = w; - ptask[i].data.abits = a; - ptask[i].data.ignore = 0;//i != 0; - ptask[i].data.size = ptask[i].data.obits * ptask[i].data.blocksize; - } -} - -#define TEMPBLOCK 128 - -#if 0 -// get_num_groups(0) == number of tasks -// get_num_groups(1) == number of windows -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clComputeAutocor( - __global float *output, - __global const int *samples, - __global const float *window, - __global FLACCLSubframeTask *tasks, - const int taskCount // tasks per block -) -{ - FLACCLSubframeData task = tasks[get_group_id(0) * taskCount].data; - int len = task.blocksize; - int windowOffs = get_group_id(1) * len; - float data[TEMPBLOCK + MAX_ORDER + 3]; - double ac[MAX_ORDER + 4]; - - for (int i = 0; i <= MAX_ORDER; ++i) - ac[i] = 0.0; - - for (int pos = 0; pos < len; pos += TEMPBLOCK) - { - for (int tid = 0; tid < TEMPBLOCK + MAX_ORDER + 3; tid++) - data[tid] = tid < len - pos ? samples[task.samplesOffs + pos + tid] * window[windowOffs + pos + tid] : 0.0f; - - for (int i = 0; i <= MAX_ORDER; i += 4) - { - float4 temp = 0.0; - for (int j = 0; j < min(TEMPBLOCK, len - pos); j++) - temp += data[j] * vload4(0, &data[j + i]); - ac[i] += temp.x; - ac[i+1] += temp.y; - ac[i+2] += temp.z; - ac[i+3] += temp.w; - } - } - __global float * pout = &output[(get_group_id(0) * get_num_groups(1) + get_group_id(1)) * (MAX_ORDER + 1)]; - for (int i = 0; i <= MAX_ORDER; ++i) - pout[i] = ac[i]; -} -#else -#define STORE_AC(ro, val) if (ro <= MAX_ORDER) pout[ro] = val; -#define STORE_AC4(ro, val) STORE_AC(ro*4+0, val##ro.x) STORE_AC(ro*4+1, val##ro.y) STORE_AC(ro*4+2, val##ro.z) STORE_AC(ro*4+3, val##ro.w) - -// get_num_groups(0) == number of tasks -// get_num_groups(1) == number of windows -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clComputeAutocor( - __global float *output, - __global const int *samples, - __global const float *window, - __global FLACCLSubframeTask *tasks, - const int taskCount // tasks per block -) -{ - FLACCLSubframeData task = tasks[get_group_id(0) * taskCount].data; - int len = task.blocksize; - int windowOffs = get_group_id(1) * len; - float data[TEMPBLOCK + MAX_ORDER + 3]; - double4 ac0 = 0.0, ac1 = 0.0, ac2 = 0.0, ac3 = 0.0, ac4 = 0.0, ac5 = 0.0, ac6 = 0.0, ac7 = 0.0, ac8 = 0.0; - - for (int pos = 0; pos < len; pos += TEMPBLOCK) - { - for (int tid = 0; tid < TEMPBLOCK + MAX_ORDER + 3; tid++) - data[tid] = tid < len - pos ? samples[task.samplesOffs + pos + tid] * window[windowOffs + pos + tid] : 0.0f; - - for (int j = 0; j < TEMPBLOCK;) - { - float4 temp0 = 0.0f, temp1 = 0.0f, temp2 = 0.0f, temp3 = 0.0f, temp4 = 0.0f, temp5 = 0.0f, temp6 = 0.0f, temp7 = 0.0f, temp8 = 0.0f; - for (int k = 0; k < 32; k++) - { - float d0 = data[j]; - temp0 += d0 * vload4(0, &data[j]); - temp1 += d0 * vload4(1, &data[j]); -#if MAX_ORDER >= 8 - temp2 += d0 * vload4(2, &data[j]); -#if MAX_ORDER >= 12 - temp3 += d0 * vload4(3, &data[j]); -#if MAX_ORDER >= 16 - temp4 += d0 * vload4(4, &data[j]); - temp5 += d0 * vload4(5, &data[j]); - temp6 += d0 * vload4(6, &data[j]); - temp7 += d0 * vload4(7, &data[j]); - temp8 += d0 * vload4(8, &data[j]); -#endif -#endif -#endif - j++; - } - ac0 += convert_double4(temp0); - ac1 += convert_double4(temp1); - #if MAX_ORDER >= 8 - ac2 += convert_double4(temp2); - #if MAX_ORDER >= 12 - ac3 += convert_double4(temp3); - #if MAX_ORDER >= 16 - ac4 += convert_double4(temp4); - ac5 += convert_double4(temp5); - ac6 += convert_double4(temp6); - ac7 += convert_double4(temp7); - ac8 += convert_double4(temp8); - #endif - #endif - #endif - } - } - __global float * pout = &output[(get_group_id(0) * get_num_groups(1) + get_group_id(1)) * (MAX_ORDER + 1)]; - STORE_AC4(0, ac) STORE_AC4(1, ac) STORE_AC4(2, ac) STORE_AC4(3, ac) - STORE_AC4(4, ac) STORE_AC4(5, ac) STORE_AC4(6, ac) STORE_AC4(7, ac) - STORE_AC4(8, ac) -} -#endif - -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clComputeLPC( - __global float *pautoc, - __global float *lpcs, - int windowCount -) -{ - int lpcOffs = (get_group_id(0) + get_group_id(1) * windowCount) * (MAX_ORDER + 1) * 32; - int autocOffs = (get_group_id(0) + get_group_id(1) * get_num_groups(0)) * (MAX_ORDER + 1); - volatile double ldr[32]; - volatile double gen0[32]; - volatile double gen1[32]; - volatile double err[32]; - __global float* autoc = pautoc + autocOffs; - - for (int i = 0; i < MAX_ORDER; i++) - { - gen0[i] = gen1[i] = autoc[i + 1]; - ldr[i] = 0.0; - } - - // Compute LPC using Schur and Levinson-Durbin recursion - double error = autoc[0]; - for (int order = 0; order < MAX_ORDER; order++) - { - // Schur recursion - double reff = -gen1[0] / error; - //error += gen1[0] * reff; // Equivalent to error *= (1 - reff * reff); - error *= (1 - reff * reff); - - for (int j = 0; j < MAX_ORDER - 1 - order; j++) - { - gen1[j] = gen1[j + 1] + reff * gen0[j]; - gen0[j] = gen1[j + 1] * reff + gen0[j]; - } - - err[order] = error; - - // Levinson-Durbin recursion - - ldr[order] = reff; - for (int j = 0; j < order / 2; j++) - { - double tmp = ldr[j]; - ldr[j] += reff * ldr[order - 1 - j]; - ldr[order - 1 - j] += reff * tmp; - } - if (0 != (order & 1)) - ldr[order / 2] += ldr[order / 2] * reff; - - // Output coeffs - for (int j = 0; j <= order; j++) - lpcs[lpcOffs + order * 32 + j] = -ldr[order - j]; - } - // Output prediction error estimates - for (int j = 0; j < MAX_ORDER; j++) - lpcs[lpcOffs + MAX_ORDER * 32 + j] = err[j]; -} - -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clQuantizeLPC( - __global FLACCLSubframeTask *tasks, - __global float*lpcs, - int taskCount, // tasks per block - int taskCountLPC, // tasks per set of coeffs (<= 32) - int minprecision, - int precisions -) -{ - int bs = tasks[get_group_id(1) * taskCount].data.blocksize; - int abits = tasks[get_group_id(1) * taskCount].data.abits; - int lpcOffs = (get_group_id(0) + get_group_id(1) * get_num_groups(0)) * (MAX_ORDER + 1) * 32; - float error[MAX_ORDER]; - int best_orders[MAX_ORDER]; - - // Load prediction error estimates based on Akaike's Criteria - for (int tid = 0; tid < MAX_ORDER; tid++) - { - error[tid] = bs * log(lpcs[lpcOffs + MAX_ORDER * 32 + tid]) + tid * 4.12f * log(bs); - best_orders[tid] = tid; - } - - // Select best orders - for (int i = 0; i < MAX_ORDER && i < taskCountLPC; i++) - { - for (int j = i + 1; j < MAX_ORDER; j++) - { - if (error[best_orders[j]] < error[best_orders[i]]) - { - int tmp = best_orders[j]; - best_orders[j] = best_orders[i]; - best_orders[i] = tmp; - } - } - } - - // Quantization - for (int i = 0; i < taskCountLPC; i ++) - { - int order = best_orders[i >> precisions]; - int tmpi = 0; - for (int tid = 0; tid <= order; tid ++) - { - float lpc = lpcs[lpcOffs + order * 32 + tid]; - // get 15 bits of each coeff - int c = convert_int_rte(lpc * (1 << 15)); - // remove sign bits - tmpi |= c ^ (c >> 31); - } - // choose precision - //int cbits = max(3, min(10, 5 + (abits >> 1))); // - convert_int_rte(shared.PE[order - 1]) - int cbits = max(3, min(min(13 - minprecision + (i - ((i >> precisions) << precisions)) - (bs <= 2304) - (bs <= 1152) - (bs <= 576), abits), clz(order) + 1 - abits)); - // calculate shift based on precision and number of leading zeroes in coeffs - int shift = max(0,min(15, clz(tmpi) - 18 + cbits)); - - int taskNo = get_group_id(1) * taskCount + get_group_id(0) * taskCountLPC + i; - tmpi = 0; - for (int tid = 0; tid <= order; tid ++) - { - float lpc = lpcs[lpcOffs + order * 32 + tid]; - // quantize coeffs with given shift - int c = convert_int_rte(clamp(lpc * (1 << shift), -1 << (cbits - 1), 1 << (cbits - 1))); - // remove sign bits - tmpi |= c ^ (c >> 31); - tasks[taskNo].coefs[tid] = c; - } - // calculate actual number of bits (+1 for sign) - cbits = 1 + 32 - clz(tmpi); - // output shift, cbits, ro - tasks[taskNo].data.shift = shift; - tasks[taskNo].data.cbits = cbits; - tasks[taskNo].data.residualOrder = order + 1; - } -} - -inline int calc_residual(__global int *ptr, int * coefs, int ro) -{ - int sum = 0; - for (int i = 0; i < ro; i++) - sum += ptr[i] * coefs[i]; - return sum; -} - -#define ENCODE_N(cro,action) for (int pos = cro; pos < bs; pos ++) { \ - int t = (data[pos] - (calc_residual(data + pos - cro, task.coefs, cro) >> task.data.shift)) >> task.data.wbits; \ - action; \ - } -#define SWITCH_N(action) \ - switch (ro) \ - { \ - case 0: ENCODE_N(0, action) break; \ - case 1: ENCODE_N(1, action) break; \ - case 2: ENCODE_N(2, action) /*if (task.coefs[0] == -1 && task.coefs[1] == 2) ENCODE_N(2, 2 * ptr[1] - ptr[0], action) else*/ break; \ - case 3: ENCODE_N(3, action) break; \ - case 4: ENCODE_N(4, action) break; \ - case 5: ENCODE_N(5, action) break; \ - case 6: ENCODE_N(6, action) break; \ - case 7: ENCODE_N(7, action) break; \ - case 8: ENCODE_N(8, action) break; \ - case 9: ENCODE_N(9, action) break; \ - case 10: ENCODE_N(10, action) break; \ - case 11: ENCODE_N(11, action) break; \ - case 12: ENCODE_N(12, action) break; \ - default: ENCODE_N(ro, action) \ - } - -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clSelectStereoTasks( - __global FLACCLSubframeTask *tasks, - int count - ) -{ - for (int i = 0; i < count; i++) - { - __global FLACCLSubframeTask* ptask = tasks + count * get_group_id(0) + i; - ptask->data.ignore = i != 0; - ptask->data.size = ptask->data.obits * ptask->data.blocksize; - } -} - -__kernel /*__attribute__(( vec_type_hint (int4)))*/ __attribute__((reqd_work_group_size(1, 1, 1))) -void clEstimateResidual( - __global int*samples, - __global FLACCLSubframeTask *tasks - ) -{ - FLACCLSubframeTask task = tasks[get_group_id(0)]; - int ro = task.data.residualOrder; - int bs = task.data.blocksize; -#define EPO 6 - int len[1 << EPO]; // blocksize / 64!!!! - - if (task.data.ignore) - { - tasks[get_group_id(0)].data.size = task.data.obits * bs; - return; - } - - __global int *data = &samples[task.data.samplesOffs]; - // for (int i = ro; i < 32; i++) - //task.coefs[i] = 0; - for (int i = 0; i < 1 << EPO; i++) - len[i] = 0; - - SWITCH_N((t = clamp(t, -0x7fffff, 0x7fffff), len[pos >> (12 - EPO)] += (t << 1) ^ (t >> 31))) - - int total = 0; - for (int i = 0; i < 1 << EPO; i++) - { - int res = min(0x7fffff,len[i]); - int k = clamp(clz(1 << (12 - EPO)) - clz(res), 0, 14); // 27 - clz(res) == clz(16) - clz(res) == log2(res / 16) - total += (k << (12 - EPO)) + (res >> k); - } - int partLen = min(0x7ffffff, total) + (bs - ro); - int obits = task.data.obits - task.data.wbits; - tasks[get_group_id(0)].data.size = min(obits * bs, - task.data.type == Fixed ? ro * obits + 6 + (4 * 1/2) + partLen : - task.data.type == LPC ? ro * obits + 4 + 5 + ro * task.data.cbits + 6 + (4 * 1/2)/* << porder */ + partLen : - task.data.type == Constant ? obits * select(1, bs, partLen != bs - ro) : - obits * bs); -} - -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clChooseBestMethod( - __global FLACCLSubframeTask *tasks, - int taskCount - ) -{ - int best_length = 0x7fffff; - int best_no = 0; - for (int taskNo = 0; taskNo < taskCount; taskNo++) - { - int len = tasks[taskNo + taskCount * get_group_id(0)].data.size; - if (len < best_length) - { - best_length = len; - best_no = taskNo; - } - } - - tasks[taskCount * get_group_id(0)].data.best_index = taskCount * get_group_id(0) + best_no; -} - -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clCopyBestMethod( - __global FLACCLSubframeTask *tasks_out, - __global FLACCLSubframeTask *tasks, - int count - ) -{ - int best_index = tasks[count * get_group_id(0)].data.best_index; - tasks_out[get_group_id(0)] = tasks[best_index]; -} - -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clCopyBestMethodStereo( - __global FLACCLSubframeTask *tasks_out, - __global FLACCLSubframeTask *tasks, - int count - ) -{ - int best_index[4]; - int best_size[4]; - int lr_index[2]; - - for (int i = 0; i < 4; i++) - { - int best = tasks[count * (get_group_id(0) * 4 + i)].data.best_index; - best_index[i] = best; - best_size[i] = tasks[best].data.size; - } - - int bitsBest = best_size[2] + best_size[3]; // MidSide - lr_index[0] = best_index[2]; - lr_index[1] = best_index[3]; - if (bitsBest > best_size[3] + best_size[1]) // RightSide - { - bitsBest = best_size[3] + best_size[1]; - lr_index[0] = best_index[3]; - lr_index[1] = best_index[1]; - } - if (bitsBest > best_size[0] + best_size[3]) // LeftSide - { - bitsBest = best_size[0] + best_size[3]; - lr_index[0] = best_index[0]; - lr_index[1] = best_index[3]; - } - if (bitsBest > best_size[0] + best_size[1]) // LeftRight - { - bitsBest = best_size[0] + best_size[1]; - lr_index[0] = best_index[0]; - lr_index[1] = best_index[1]; - } - tasks_out[2 * get_group_id(0)] = tasks[lr_index[0]]; - tasks_out[2 * get_group_id(0)].data.residualOffs = tasks[best_index[0]].data.residualOffs; - tasks_out[2 * get_group_id(0) + 1] = tasks[lr_index[1]]; - tasks_out[2 * get_group_id(0) + 1].data.residualOffs = tasks[best_index[1]].data.residualOffs; -} - -// get_group_id(0) == task index -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clEncodeResidual( - __global int *residual, - __global int *samples, - __global FLACCLSubframeTask *tasks - ) -{ - FLACCLSubframeTask task = tasks[get_group_id(0)]; - int bs = task.data.blocksize; - int ro = task.data.residualOrder; - __global int *data = &samples[task.data.samplesOffs]; - SWITCH_N(residual[task.data.residualOffs + pos] = t); -} - -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clCalcPartition( - __global int *partition_lengths, - __global int *residual, - __global FLACCLSubframeTask *tasks, - int max_porder, // <= 8 - int psize // == task.blocksize >> max_porder? - ) -{ - FLACCLSubframeTask task = tasks[get_group_id(1)]; - int bs = task.data.blocksize; - int ro = task.data.residualOrder; - //int psize = bs >> max_porder; - __global int *pl = partition_lengths + (1 << (max_porder + 1)) * get_group_id(1); - - for (int p = 0; p < (1 << max_porder); p++) - pl[p] = 0; - - for (int pos = ro; pos < bs; pos ++) - { - int t = residual[task.data.residualOffs + pos]; - // overflow protection - t = clamp(t, -0x7fffff, 0x7fffff); - // convert to unsigned - t = (t << 1) ^ (t >> 31); - pl[pos / psize] += t; - } -} -// get_group_id(0) == task index -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clCalcPartition16( - __global int *partition_lengths, - __global int *residual, - __global int *samples, - __global FLACCLSubframeTask *tasks, - int max_porder // <= 8 - ) -{ - FLACCLSubframeTask task = tasks[get_global_id(0)]; - int bs = task.data.blocksize; - int ro = task.data.residualOrder; - __global int *data = &samples[task.data.samplesOffs]; - __global int *pl = partition_lengths + (1 << (max_porder + 1)) * get_global_id(0); - for (int p = 0; p < (1 << max_porder); p++) - pl[p] = 0; - //__global int *rptr = residual + task.data.residualOffs; - //SWITCH_N((rptr[pos] = t, pl[pos >> 4] += (t << 1) ^ (t >> 31))); - SWITCH_N((residual[task.data.residualOffs + pos] = t, t = clamp(t, -0x7fffff, 0x7fffff), t = (t << 1) ^ (t >> 31), pl[pos >> 4] += t)); -} - -// Sums partition lengths for a certain k == get_group_id(0) -// get_group_id(0) == k -// get_group_id(1) == task index -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clSumPartition( - __global int* partition_lengths, - int max_porder - ) -{ - if (get_group_id(0) != 0) // ignore k != 0 - return; - __global int * sums = partition_lengths + (1 << (max_porder + 1)) * get_group_id(1); - for (int i = max_porder - 1; i >= 0; i--) - { - for (int j = 0; j < (1 << i); j++) - { - sums[(2 << i) + j] = sums[2 * j] + sums[2 * j + 1]; - // if (get_group_id(1) == 0) - //printf("[%d][%d]: %d + %d == %d\n", i, j, sums[2 * j], sums[2 * j + 1], sums[2 * j] + sums[2 * j + 1]); - } - sums += 2 << i; - } -} - -// Finds optimal rice parameter for each partition. -// get_group_id(0) == task index -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clFindRiceParameter( - __global FLACCLSubframeTask *tasks, - __global int* rice_parameters, - __global int* partition_lengths, - int max_porder - ) -{ - __global FLACCLSubframeTask* task = tasks + get_group_id(0); - const int tid = get_local_id(0); - int lim = (2 << max_porder) - 1; - int psize = task->data.blocksize >> max_porder; - int bs = task->data.blocksize; - int ro = task->data.residualOrder; - for (int offs = 0; offs < lim; offs ++) - { - int pl = partition_lengths[(1 << (max_porder + 1)) * get_group_id(0) + offs]; - int porder = 31 - clz(lim - offs); - int ps = (bs >> porder) - select(0, ro, offs == lim + 1 - (2 << porder)); - //if (ps <= 0) - // printf("max_porder == %d, porder == %d, ro == %d\n", max_porder, porder, ro); - int k = clamp(31 - clz(pl / max(1, ps)), 0, 14); - int plk = ps * (k + 1) + (pl >> k); - - // output rice parameter - rice_parameters[(get_group_id(0) << (max_porder + 2)) + offs] = k; - // output length - rice_parameters[(get_group_id(0) << (max_porder + 2)) + (1 << (max_porder + 1)) + offs] = plk; - } -} - -// get_group_id(0) == task index -__kernel __attribute__((reqd_work_group_size(1, 1, 1))) -void clFindPartitionOrder( - __global int *residual, - __global int* best_rice_parameters, - __global FLACCLSubframeTask *tasks, - __global int* rice_parameters, - int max_porder - ) -{ - __global FLACCLSubframeTask* task = tasks + get_group_id(0); - int partlen[9]; - for (int p = 0; p < 9; p++) - partlen[p] = 0; - // fetch partition lengths - const int pos = (get_group_id(0) << (max_porder + 2)) + (2 << max_porder); - int lim = (2 << max_porder) - 1; - for (int offs = 0; offs < lim; offs ++) - { - int len = rice_parameters[pos + offs]; - int porder = 31 - clz(lim - offs); - partlen[porder] += len; - } - - int best_length = partlen[0] + 4; - int best_porder = 0; - for (int porder = 1; porder <= max_porder; porder++) - { - int length = (4 << porder) + partlen[porder]; - best_porder = select(best_porder, porder, length < best_length); - best_length = min(best_length, length); - } - - best_length = (4 << best_porder) + task->data.blocksize - task->data.residualOrder; - int best_psize = task->data.blocksize >> best_porder; - int start = task->data.residualOffs + task->data.residualOrder; - int fin = task->data.residualOffs + best_psize; - for (int p = 0; p < (1 << best_porder); p++) - { - int k = rice_parameters[pos - (2 << best_porder) + p]; - best_length += k * (fin - start); - for (int i = start; i < fin; i++) - { - int t = residual[i]; - best_length += ((t << 1) ^ (t >> 31)) >> k; - } - start = fin; - fin += best_psize; - } - - int obits = task->data.obits - task->data.wbits; - task->data.porder = best_porder; - task->data.size = - task->data.type == Fixed ? task->data.residualOrder * obits + 6 + best_length : - task->data.type == LPC ? task->data.residualOrder * obits + 6 + best_length + 4 + 5 + task->data.residualOrder * task->data.cbits : - task->data.type == Constant ? obits : obits * task->data.blocksize; - for (int offs = 0; offs < (1 << best_porder); offs ++) - best_rice_parameters[(get_group_id(0) << max_porder) + offs] = rice_parameters[pos - (2 << best_porder) + offs]; -} -#endif