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

experiment

Browse Source
This commit is contained in:
chudov
2010-11-12 05:44:39 +00:00
parent c3b0c1ae9c
commit dd9b86e4a5
4 changed files with 331 additions and 812 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
CUETools.Codecs.FLACCL/CUETools.Codecs.FLACCL.csproj
View File

@@ -65,9 +65,6 @@
</ProjectReference>
</ItemGroup>
<ItemGroup>
<Content Include="flaccpu.cl">
<CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
</Content>
<None Include="flac.cu">
</None>
<Content Include="flac.cl">

153
CUETools.Codecs.FLACCL/FLACCLWriter.cs
View File

@@ -826,16 +826,34 @@ 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);
if (task.riceOnGPU)
{
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
@@ -851,6 +869,7 @@ namespace CUETools.Codecs.FLACCL
j += cnt;
}
}
}
unsafe void
output_subframe_constant(FlacFrame frame, FlacSubframeInfo sub)
@@ -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,11 +2875,44 @@ namespace CUETools.Codecs.FLACCL
groupSize,
channels * frameCount);
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 (!writer._settings.MappedMemory)
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);
}

198
CUETools.Codecs.FLACCL/flac.cl
View File

@@ -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.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.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 == 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

759
CUETools.Codecs.FLACCL/flaccpu.cl
View File

@@ -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