optimizations

CUETools.Codecs.FLACCL/FLACCLWriter.cs

@@ -444,6 +444,20 @@ namespace CUETools.Codecs.FLACCL
 			}
 		}
 
+		public int OrdersPerChannel
+		{
+			get
+			{
+				return eparams.orders_per_channel;
+			}
+			set
+			{
+				if (value < 0 || value > 32)
+					throw new Exception("invalid OrdersPerWindow " + value.ToString());
+				eparams.orders_per_channel = value;
+			}
+		}
+
 		public int MinLPCOrder
 		{
 			get
@@ -998,6 +1012,10 @@ namespace CUETools.Codecs.FLACCL
 			task.nResidualTasksPerChannel = task.nWindowFunctions * task.nTasksPerWindow + (eparams.do_constant ? 1 : 0) + Math.Max(0, 1 + eparams.max_fixed_order - eparams.min_fixed_order);
             if (task.nResidualTasksPerChannel > 32)
                 throw new Exception("too many tasks");
+			if (channels == 2 && channelsCount == 4)
+				task.nEstimateTasksPerChannel = Math.Min(eparams.orders_per_channel, task.nResidualTasksPerChannel);
+			else
+				task.nEstimateTasksPerChannel = task.nResidualTasksPerChannel;
 
 			//if (task.nResidualTasksPerChannel >= 4)
 			//    task.nResidualTasksPerChannel = (task.nResidualTasksPerChannel + 7) & ~7;
@@ -1005,6 +1023,19 @@ namespace CUETools.Codecs.FLACCL
 			{
 				for (int ch = 0; ch < channelsCount; ch++)
 				{
+					int *selectedTasks = (int*)task.clSelectedTasksPtr;
+					for (int j = 0; j < task.nEstimateTasksPerChannel; j++)
+					{
+						int k = j;
+						if (j < task.nWindowFunctions * task.nTasksPerWindow && task.nWindowFunctions > 1)
+						{
+							k = (j % task.nWindowFunctions) * task.nTasksPerWindow
+								+ (j / task.nWindowFunctions);
+						}
+						selectedTasks[(iFrame * channelsCount + ch) * task.nEstimateTasksPerChannel + j] =
+							(iFrame * channelsCount + ch) * task.nResidualTasksPerChannel + k;
+					}
+
 					for (int iWindow = 0; iWindow < task.nWindowFunctions; iWindow++)
 					{
 						// LPC tasks
@@ -1018,6 +1049,9 @@ 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++;
 						}
 					}
@@ -1031,6 +1065,9 @@ 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;
 						task.ResidualTasks[task.nResidualTasks].shift = 0;
 						task.ResidualTasks[task.nResidualTasks].coefs[0] = 1;
@@ -1047,6 +1084,9 @@ 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;
 						switch (order)
 						{
@@ -1094,6 +1134,8 @@ namespace CUETools.Codecs.FLACCL
 
             if (!_settings.MappedMemory)
 			    task.openCLCQ.EnqueueWriteBuffer(task.clResidualTasks, true, 0, sizeof(FLACCLSubframeTask) * task.nResidualTasks, task.clResidualTasksPtr);
+			if (!_settings.MappedMemory)
+				task.openCLCQ.EnqueueWriteBuffer(task.clSelectedTasks, true, 0, sizeof(int) * (nFrames * channelsCount * task.nEstimateTasksPerChannel), task.clSelectedTasksPtr);		
 		}
 
 		unsafe void encode_residual(FLACCLTask task)
@@ -1536,7 +1578,7 @@ namespace CUETools.Codecs.FLACCL
 				openCLContext = OCLMan.Context;
 				try
 				{
-					openCLProgram = OCLMan.CompileFile(_settings.DeviceType == OpenCLDeviceType.CPU ? "flaccpu.cl" : "flac.cl");
+					openCLProgram = OCLMan.CompileFile("flac.cl");
 				}
 				catch (OpenCLBuildException ex)
 				{
@@ -2012,6 +2054,8 @@ namespace CUETools.Codecs.FLACCL
 
 		public int orders_per_window;
 
+		public int orders_per_channel;
+
 		// minimum fixed prediction order
 		// set by user prior to calling flake_encode_init
 		// if set to less than 0, it is chosen based on compression.
@@ -2080,6 +2124,7 @@ namespace CUETools.Codecs.FLACCL
 			variable_block_size = 0;
 			lpc_min_precision_search = 0;
 			lpc_max_precision_search = 0;
+			orders_per_channel = 32;
 			do_seektable = true;
 			do_wasted = true;
 			do_constant = true;
@@ -2120,48 +2165,52 @@ namespace CUETools.Codecs.FLACCL
 					max_partition_order = 4;
 					break;
 				case 3:
-					window_function = WindowFunction.Bartlett;
 					do_constant = false;
 					min_fixed_order = 2;
 					max_fixed_order = 2;
 					orders_per_window = 1;
+					orders_per_channel = 1;
 					max_prediction_order = 8;
-					max_partition_order = 4;
 					break;
 				case 4:
 					do_constant = false;
 					min_fixed_order = 2;
 					max_fixed_order = 2;
-					orders_per_window = 1;
+					orders_per_window = 2;
+					orders_per_channel = 2;
 					max_prediction_order = 8;
-					max_partition_order = 4;
 					break;
 				case 5:
 					do_constant = false;
 					min_fixed_order = 2;
 					max_fixed_order = 2;
-					orders_per_window = 2;
+					orders_per_window = 4;
+					orders_per_channel = 4;
 					max_prediction_order = 8;
 					break;
 				case 6:
 					do_constant = false;
 					min_fixed_order = 2;
 					max_fixed_order = 2;
-					orders_per_window = 1;
+					orders_per_window = 2;
+					orders_per_channel = 2;
 					break;
 				case 7:
 					do_constant = false;
 					min_fixed_order = 2;
 					max_fixed_order = 2;
-					orders_per_window = 3;
+					orders_per_window = 4;
+					orders_per_channel = 4;
 					break;
 				case 8:
-					orders_per_window = 12;
+					orders_per_window = 8;
+					orders_per_channel = 8;
 					break;
 				case 9:
 					min_fixed_order = 2;
 					max_fixed_order = 2;
-					orders_per_window = 3;
+					orders_per_window = 4;
+					orders_per_channel = 4;
 					max_prediction_order = 32;
 					break;
 				case 10:
@@ -2201,7 +2250,8 @@ namespace CUETools.Codecs.FLACCL
 		public int wbits;
 		public int abits;
 		public int porder;
-		public fixed int reserved[2];
+		public int ignore;
+		public int reserved;
 		public fixed int coefs[32];
 	};
 
@@ -2220,8 +2270,6 @@ namespace CUETools.Codecs.FLACCL
         public Kernel clSelectStereoTasks;
 		public Kernel clEstimateResidual;
 		public Kernel clChooseBestMethod;
-		public Kernel clCopyBestMethod;
-		public Kernel clCopyBestMethodStereo;
 		public Kernel clEncodeResidual;
 		public Kernel clCalcPartition;
 		public Kernel clCalcPartition16;
@@ -2236,6 +2284,9 @@ namespace CUETools.Codecs.FLACCL
 		public Mem clRiceParams;
 		public Mem clBestRiceParams;
 		public Mem clAutocorOutput;
+		public Mem clSelectedTasks;
+		public Mem clSelectedTasksSecondEstimate;
+		public Mem clSelectedTasksBestMethod;
 		public Mem clResidualTasks;
 		public Mem clBestResidualTasks;
 		public Mem clWindowFunctions;
@@ -2246,6 +2297,7 @@ namespace CUETools.Codecs.FLACCL
         public Mem clResidualTasksPinned;
         public Mem clBestResidualTasksPinned;
         public Mem clWindowFunctionsPinned;
+		public Mem clSelectedTasksPinned;
 
         public IntPtr clSamplesBytesPtr;
         public IntPtr clResidualPtr;
@@ -2253,6 +2305,7 @@ namespace CUETools.Codecs.FLACCL
         public IntPtr clResidualTasksPtr;
         public IntPtr clBestResidualTasksPtr;
         public IntPtr clWindowFunctionsPtr;
+		public IntPtr clSelectedTasksPtr;
 
         public int[] samplesBuffer;
 		public byte[] outputBuffer;
@@ -2267,6 +2320,7 @@ namespace CUETools.Codecs.FLACCL
 		public int samplesBufferLen;
 		public int nResidualTasks = 0;
 		public int nResidualTasksPerChannel = 0;
+		public int nEstimateTasksPerChannel = 0;
 		public int nTasksPerWindow = 0;
 		public int nWindowFunctions = 0;
 		public int max_porder = 0;
@@ -2306,6 +2360,7 @@ namespace CUETools.Codecs.FLACCL
             int lpcDataLen = autocorLen * 32;
             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;
 
             if (!writer._settings.MappedMemory)
             {
@@ -2315,6 +2370,7 @@ namespace CUETools.Codecs.FLACCL
                 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);
 
                 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);
@@ -2322,6 +2378,7 @@ namespace CUETools.Codecs.FLACCL
                 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);
 
                 clSamplesBytesPtr = openCLCQ.EnqueueMapBuffer(clSamplesBytesPinned, true, MapFlags.WRITE, 0, samplesBufferLen / 2);
                 clResidualPtr = openCLCQ.EnqueueMapBuffer(clResidualPinned, true, MapFlags.WRITE, 0, samplesBufferLen);
@@ -2329,7 +2386,8 @@ namespace CUETools.Codecs.FLACCL
                 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);
+			}
             else
             {
                 clSamplesBytes = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE | MemFlags.ALLOC_HOST_PTR, (uint)samplesBufferLen / 2);
@@ -2338,6 +2396,7 @@ namespace CUETools.Codecs.FLACCL
                 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);
 
                 clSamplesBytesPtr = openCLCQ.EnqueueMapBuffer(clSamplesBytes, true, MapFlags.WRITE, 0, samplesBufferLen / 2);
                 clResidualPtr = openCLCQ.EnqueueMapBuffer(clResidual, true, MapFlags.WRITE, 0, samplesBufferLen);
@@ -2345,6 +2404,7 @@ namespace CUETools.Codecs.FLACCL
                 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);
 
                 //clSamplesBytesPtr = clSamplesBytes.HostPtr;
                 //clResidualPtr = clResidual.HostPtr;
@@ -2357,7 +2417,9 @@ namespace CUETools.Codecs.FLACCL
             clSamples = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, samplesBufferLen);
             clLPCData = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, lpcDataLen);
             clAutocorOutput = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, autocorLen);
-            if (writer._settings.GPUOnly)
+			clSelectedTasksSecondEstimate = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, selectedLen);
+			clSelectedTasksBestMethod = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, selectedLen);
+			if (writer._settings.GPUOnly)
             {
                 clPartitions = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, partitionsLen);
                 clRiceParams = openCLProgram.Context.CreateBuffer(MemFlags.READ_WRITE, riceParamsLen);
@@ -2376,8 +2438,6 @@ namespace CUETools.Codecs.FLACCL
             clSelectStereoTasks = openCLProgram.CreateKernel("clSelectStereoTasks");
 			clEstimateResidual = openCLProgram.CreateKernel("clEstimateResidual");
 			clChooseBestMethod = openCLProgram.CreateKernel("clChooseBestMethod");
-			clCopyBestMethod = openCLProgram.CreateKernel("clCopyBestMethod");
-			clCopyBestMethodStereo = openCLProgram.CreateKernel("clCopyBestMethodStereo");
 			if (writer._settings.GPUOnly)
 			{
 				clEncodeResidual = openCLProgram.CreateKernel("clEncodeResidual");
@@ -2426,8 +2486,6 @@ namespace CUETools.Codecs.FLACCL
             clSelectStereoTasks.Dispose();
 			clEstimateResidual.Dispose();
 			clChooseBestMethod.Dispose();
-			clCopyBestMethod.Dispose();
-			clCopyBestMethodStereo.Dispose();
 			if (writer._settings.GPUOnly)
 			{
 				clEncodeResidual.Dispose();
@@ -2461,6 +2519,9 @@ namespace CUETools.Codecs.FLACCL
                 if (clWindowFunctionsPtr != IntPtr.Zero)
                     openCLCQ.EnqueueUnmapMemObject(clWindowFunctionsPinned, clWindowFunctionsPtr);
                 clWindowFunctionsPtr = IntPtr.Zero;
+				if (clSelectedTasksPtr != IntPtr.Zero)
+					openCLCQ.EnqueueUnmapMemObject(clSelectedTasksPinned, clSelectedTasksPtr);
+				clSelectedTasksPtr = IntPtr.Zero;
 
                 clSamplesBytesPinned.Dispose();
                 clResidualPinned.Dispose();
@@ -2468,6 +2529,7 @@ namespace CUETools.Codecs.FLACCL
                 clResidualTasksPinned.Dispose();
                 clBestResidualTasksPinned.Dispose();
                 clWindowFunctionsPinned.Dispose();
+				clSelectedTasksPinned.Dispose();
             }
             else
             {
@@ -2477,6 +2539,7 @@ namespace CUETools.Codecs.FLACCL
                 openCLCQ.EnqueueUnmapMemObject(clResidualTasks, clResidualTasksPtr);
                 openCLCQ.EnqueueUnmapMemObject(clBestResidualTasks, clBestResidualTasksPtr);
                 openCLCQ.EnqueueUnmapMemObject(clWindowFunctions, clWindowFunctionsPtr);
+				openCLCQ.EnqueueUnmapMemObject(clSelectedTasks, clSelectedTasksPtr);
             }
 
 			clSamples.Dispose();
@@ -2484,9 +2547,12 @@ namespace CUETools.Codecs.FLACCL
 			clLPCData.Dispose();
 			clResidual.Dispose();
 			clAutocorOutput.Dispose();
+			clSelectedTasksSecondEstimate.Dispose();
+			clSelectedTasksBestMethod.Dispose();
 			clResidualTasks.Dispose();
 			clBestResidualTasks.Dispose();
 			clWindowFunctions.Dispose();
+			clSelectedTasks.Dispose();
 
 			openCLCQ.Dispose();
 
@@ -2527,9 +2593,10 @@ namespace CUETools.Codecs.FLACCL
 			clChannelDecorr.SetArgs(
 				clSamples,
 				clSamplesBytes,
-				FLACCLWriter.MAX_BLOCKSIZE/4);
+				FLACCLWriter.MAX_BLOCKSIZE / 4);
 
-            openCLCQ.EnqueueNDRangeKernel(clChannelDecorr, 0, frameSize * frameCount / 4);
+			openCLCQ.EnqueueNDRangeKernel(clChannelDecorr, 0, FLACCLWriter.MAX_BLOCKSIZE / 4);
+			//openCLCQ.EnqueueNDRangeKernel(clChannelDecorr, 0, (frameSize * frameCount + 3) / 4);
 
             if (eparams.do_wasted)
 			{
@@ -2587,44 +2654,74 @@ namespace CUETools.Codecs.FLACCL
 				nWindowFunctions,
 				channelsCount * frameCount);
 
-            clEstimateResidual.SetArgs(
+			if (channels == 2 && channelsCount == 4)
-				clSamples,
-				clResidualTasks);
-
-			openCLCQ.EnqueueNDRangeKernel(
-				clEstimateResidual,
-				groupSize,
-				nResidualTasksPerChannel * channelsCount * frameCount);
-
-            clChooseBestMethod.SetArgs(
-				clResidualTasks,
-				nResidualTasksPerChannel);
-
-			openCLCQ.EnqueueNDRangeKernel(
-				clChooseBestMethod,
-				Math.Min(groupSize, 32), channelsCount * frameCount);
-
-            if (channels == 2 && channelsCount == 4)
 			{
-				clCopyBestMethodStereo.SetArgs(
+				clEstimateResidual.SetArgs(
+					clSamples,
+					clSelectedTasks,
+					clResidualTasks);
+
+				openCLCQ.EnqueueNDRangeKernel(
+					clEstimateResidual,
+					groupSize,
+					nEstimateTasksPerChannel * channelsCount * frameCount); // 1 per channel, 4 channels
+
+				clSelectStereoTasks.SetArgs(
+					clResidualTasks,
+					clSelectedTasks,
+					clSelectedTasksSecondEstimate,
+					clSelectedTasksBestMethod,
+					nResidualTasksPerChannel,
+					nEstimateTasksPerChannel);
+
+				openCLCQ.EnqueueNDRangeKernel(
+					clSelectStereoTasks,
+					0, frameCount);
+
+				if (nEstimateTasksPerChannel < nResidualTasksPerChannel)
+				{
+					clEstimateResidual.SetArgs(
+						clSamples,
+						clSelectedTasksSecondEstimate,
+						clResidualTasks);
+
+					openCLCQ.EnqueueNDRangeKernel(
+						clEstimateResidual,
+						groupSize,
+						(nResidualTasksPerChannel - nEstimateTasksPerChannel) * channels * frameCount);
+				}
+
+				clChooseBestMethod.SetArgs(
 					clBestResidualTasks,
 					clResidualTasks,
+					clSelectedTasksBestMethod,
 					nResidualTasksPerChannel);
 
 				openCLCQ.EnqueueNDRangeKernel(
-					clCopyBestMethodStereo,
+					clChooseBestMethod,
-					Math.Min(groupSize, 64), frameCount);
+					0, channels * frameCount);
 			}
 			else
 			{
-				clCopyBestMethod.SetArgs(
+				clEstimateResidual.SetArgs(
+					clSamples,
+					clSelectedTasks,
+					clResidualTasks);
+
+				openCLCQ.EnqueueNDRangeKernel(
+					clEstimateResidual,
+					groupSize,
+					nResidualTasksPerChannel * channelsCount * frameCount);
+
+				clChooseBestMethod.SetArgs(
 					clBestResidualTasks,
 					clResidualTasks,
+					clSelectedTasks,
 					nResidualTasksPerChannel);
 
 				openCLCQ.EnqueueNDRangeKernel(
-					clCopyBestMethod,
+					clChooseBestMethod,
-					Math.Min(groupSize, 64), channels * frameCount);
+					0, channels * frameCount);
 			}
 
 			if (writer._settings.GPUOnly)
@@ -2726,6 +2823,7 @@ namespace CUETools.Codecs.FLACCL
         }
 	}
 
+#if LKJLKJLJK
     public static class OpenCLExtensions
     {
         public static void SetArgs(this Kernel kernel, params object[] args)
@@ -2756,4 +2854,5 @@ namespace CUETools.Codecs.FLACCL
             queue.EnqueueNDRangeKernel(kernel, 2, null, new long[] { localSizeX * globalSizeX, localSizeY * globalSizeY }, new long[] { localSizeX, localSizeY });
         }
     }
+#endif
 }

CUETools.Codecs.FLACCL/flac.cl

@@ -20,12 +20,14 @@
 #ifndef _FLACCL_KERNEL_H_
 #define _FLACCL_KERNEL_H_
 
-#if defined(__Cedar__) || defined(__Redwood__) || defined(__Juniper__) || defined(__Cypress__)
+#if defined(__Cedar__) || defined(__Redwood__) || defined(__Juniper__) || defined(__Cypress__) || defined(__CPU__)
 #define AMD
 #ifdef DEBUG
 #pragma OPENCL EXTENSION cl_amd_printf : enable
 #endif
-//#pragma OPENCL EXTENSION cl_amd_fp64 : enable
+#ifdef __CPU__
+#pragma OPENCL EXTENSION cl_amd_fp64 : enable
+#endif
 #define iclamp(a,b,c) clamp(a,b,c)
 #else
 #define iclamp(a,b,c) max(b,min(a,c))
@@ -58,7 +60,8 @@ typedef struct
     int wbits;
     int abits;
     int porder;
-    int reserved[2];
+    int ignore;
+    int reserved;
 } FLACCLSubframeData;
 
 typedef struct
@@ -67,36 +70,54 @@ typedef struct
     int coefs[32]; // fixme: should be short?
 } FLACCLSubframeTask;
 
+__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_F * 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_F * n / Np);
+}
+
 __kernel void clStereoDecorr(
-    __global int *samples,
+    __global int4 *samples,
-    __global short2 *src,
+    __global int4 *src,
     int offset
 )
 {
     int pos = get_global_id(0);
-    if (pos < offset)
+    int4 s = src[pos];
-    {
+    int4 x = (s << 16) >> 16;
-	short2 s = src[pos];
+    int4 y = s >> 16;
-	samples[pos] = s.x;
+    samples[pos] = x;
-	samples[1 * offset + pos] = s.y;
+    samples[1 * offset + pos] = y;
-	samples[2 * offset + pos] = (s.x + s.y) >> 1;
+    samples[2 * offset + pos] = (x + y) >> 1;
-	samples[3 * offset + pos] = s.x - s.y;
+    samples[3 * offset + pos] = x - y;
-    }
 }
 
 __kernel void clChannelDecorr2(
-    __global int *samples,
+    __global int4 *samples,
-    __global short2 *src,
+    __global int4 *src,
     int offset
 )
 {
     int pos = get_global_id(0);
-    if (pos < offset)
+    int4 s = src[pos];
-    {
+    samples[pos] = (s << 16) >> 16;
-	short2 s = src[pos];
+    samples[offset + pos] = s >> 16;
-	samples[pos] = s.x;
-	samples[1 * offset + pos] = s.y;
-    }
 }
 
 //__kernel void clChannelDecorr(
@@ -113,6 +134,32 @@ __kernel void clChannelDecorr2(
 #define __ffs(a) (32 - clz(a & (-a)))
 //#define __ffs(a) (33 - clz(~a & (a - 1)))
 
+#ifdef __CPU__
+__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.size = ptask[i].data.obits * ptask[i].data.blocksize;
+    }
+}
+#else
 __kernel __attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
 void clFindWastedBits(
     __global FLACCLSubframeTask *tasks,
@@ -154,11 +201,88 @@ void clFindWastedBits(
     w = max(0,__ffs(wbits[0]) - 1);
     a = 32 - clz(abits[0]) - w;
     if (tid < tasksPerChannel)
-	tasks[get_group_id(0) * tasksPerChannel + tid].data.wbits = w;
+    {
-    if (tid < tasksPerChannel)
+	int i = get_group_id(0) * tasksPerChannel + tid;
-	tasks[get_group_id(0) * tasksPerChannel + tid].data.abits = a;
+	tasks[i].data.wbits = w;
+	tasks[i].data.abits = a;
+	//tasks[i].data.size = tasks[i].data.obits * tasks[i].data.blocksize;
+    }
 }
+#endif
 
+#ifdef __CPU__
+#define TEMPBLOCK 128
+#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)
+}
+#else
 // get_num_groups(0) == number of tasks
 // get_num_groups(1) == number of windows
 __kernel __attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
@@ -181,7 +305,16 @@ void clComputeAutocor(
     int bs = task.blocksize;
     int windowOffs = get_group_id(1) * bs;
 
-    data[tid] = tid < bs ? samples[task.samplesOffs + tid] * window[windowOffs + tid] : 0.0f;
+ //   if (tid < GROUP_SIZE / 4)
+ //   {
+	//float4 dd = 0.0f;
+	//if (tid * 4 < bs)
+	//    dd = vload4(tid, window + windowOffs) * convert_float4(vload4(tid, samples + task.samplesOffs));
+	//vstore4(dd, tid, &data[0]);
+ //   }
+    data[tid] = 0.0f;
+    // This simpler code doesn't work somehow!!!
+    //data[tid] = tid < bs ? samples[task.samplesOffs + tid] * window[windowOffs + tid] : 0.0f;
 
     const int THREADS_FOR_ORDERS = MAX_ORDER < 8 ? 8 : MAX_ORDER < 16 ? 16 : MAX_ORDER < 32 ? 32 : 64;
     float corr = 0.0f;
@@ -189,24 +322,24 @@ void clComputeAutocor(
     for (int pos = 0; pos < bs; pos += GROUP_SIZE)
     {
 	// fetch samples
-	float nextData = pos + tid + GROUP_SIZE < bs ? samples[task.samplesOffs + pos + tid + GROUP_SIZE] * window[windowOffs + pos + tid + GROUP_SIZE] : 0.0f;
+	float nextData = pos + tid < bs ? samples[task.samplesOffs + pos + tid] * window[windowOffs + pos + tid] : 0.0f;
 	data[tid + GROUP_SIZE] = nextData;
 	barrier(CLK_LOCAL_MEM_FENCE);
 
-#ifdef XXXAMD
+	int lag = tid & (THREADS_FOR_ORDERS - 1);
-	__local float * dptr = &data[tid & ~(THREADS_FOR_ORDERS - 1)];
+	int tid1 = tid + GROUP_SIZE - lag;
+#ifdef AMD
 	float4 res = 0.0f;
 	for (int i = 0; i < THREADS_FOR_ORDERS / 4; i++)
-	    res += vload4(i, dptr) * vload4(i, &data[tid]);
+	    res += vload4(i, &data[tid1 - lag]) * vload4(i, &data[tid1]);
 	corr += res.x + res.y + res.w + res.z;
 #else
-	int tid1 = tid & ~(THREADS_FOR_ORDERS - 1);
 	float res = 0.0f;
 	for (int i = 0; i < THREADS_FOR_ORDERS; i++)
-	    res += data[tid1 + i] * data[tid + i];
+	    res += data[tid1 - lag + i] * data[tid1 + i];
 	corr += res;
 #endif
-	if (THREADS_FOR_ORDERS > 8 && (pos & (GROUP_SIZE * 7)) == 0)
+	if ((pos & (GROUP_SIZE * 15)) == 0)
 	{
 	    corr1 += corr;
 	    corr = 0.0f;
@@ -228,7 +361,68 @@ void clComputeAutocor(
     if (tid <= MAX_ORDER)
 	output[(get_group_id(0) * get_num_groups(1) + get_group_id(1)) * (MAX_ORDER + 1) + tid] = data[tid];
 }
+#endif
 
+#ifdef __CPU__
+__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];
+}
+#else
 __kernel __attribute__((reqd_work_group_size(32, 1, 1)))
 void clComputeLPC(
     __global float *autoc,
@@ -311,7 +505,85 @@ void clComputeLPC(
     if (get_local_id(0) < MAX_ORDER)
 	lpcs[lpcOffs + MAX_ORDER * 32 + get_local_id(0)] = shared.error[get_local_id(0)];
 }
+#endif
 
+#ifdef __CPU__
+__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;
+    }
+}
+#else
 __kernel __attribute__((reqd_work_group_size(32, 1, 1)))
 void clQuantizeLPC(
     __global FLACCLSubframeTask *tasks,
@@ -443,10 +715,82 @@ void clQuantizeLPC(
 	    tasks[taskNo].coefs[tid] = coef;
     }
 }
+#endif
 
+#ifdef __CPU__
+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__(( vec_type_hint (int4)))*/ __attribute__((reqd_work_group_size(1, 1, 1)))
+void clEstimateResidual(
+    __global int*samples,
+    __global int*selectedTasks,
+    __global FLACCLSubframeTask *tasks
+    )
+{
+    int selectedTask = selectedTasks[get_group_id(0)];
+    FLACCLSubframeTask task = tasks[selectedTask];
+    int ro = task.data.residualOrder;
+    int bs = task.data.blocksize;
+#define EPO 6
+    int len[1 << EPO]; // blocksize / 64!!!!
+
+    __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[selectedTask].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);
+}
+#else
 __kernel /*__attribute__(( vec_type_hint (int4)))*/ __attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
 void clEstimateResidual(
     __global int*samples,
+    __global int*selectedTasks,
     __global FLACCLSubframeTask *tasks
     )
 {
@@ -454,11 +798,16 @@ void clEstimateResidual(
     __local FLACCLSubframeTask task;
     __local int psum[64];
     __local float fcoef[32];
+    __local int selectedTask;
+    
+    if (get_local_id(0) == 0)
+	selectedTask = selectedTasks[get_group_id(0)];
+    barrier(CLK_LOCAL_MEM_FENCE);
 
     const int tid = get_local_id(0);
     if (tid < sizeof(task)/sizeof(int))
-	((__local int*)&task)[tid] = ((__global int*)(&tasks[get_group_id(0)]))[tid];
+	((__local int*)&task)[tid] = ((__global int*)(&tasks[selectedTask]))[tid];
-    barrier(CLK_GLOBAL_MEM_FENCE);
+    barrier(CLK_LOCAL_MEM_FENCE);
 
     int ro = task.data.residualOrder;
     int bs = task.data.blocksize;
@@ -571,106 +920,101 @@ void clEstimateResidual(
 		task.data.type == LPC ? task.data.residualOrder * obits + 4 + 5 + task.data.residualOrder * task.data.cbits + 6 + (4 * 1/2)/* << porder */ + pl :
 		task.data.type == Constant ? obits * select(1, task.data.blocksize, pl != task.data.blocksize - task.data.residualOrder) : 
 		obits * task.data.blocksize);
-	tasks[get_group_id(0)].data.size = len;
+	tasks[selectedTask].data.size = len;
+    }
+}
+#endif
+
+__kernel
+void clSelectStereoTasks(
+    __global FLACCLSubframeTask *tasks,
+    __global int*selectedTasks,
+    __global int*selectedTasksSecondEstimate,
+    __global int*selectedTasksBestMethod,
+    int taskCount,
+    int selectedCount
+    )
+{
+    int best_size[4];
+    for (int ch = 0; ch < 4; ch++)
+    {
+	int first_no = selectedTasks[(get_global_id(0) * 4 + ch) * selectedCount];
+	int best_len = tasks[first_no].data.size;
+	for (int i = 1; i < selectedCount; i++)
+	{
+	    int task_no = selectedTasks[(get_global_id(0) * 4 + ch) * selectedCount + i];
+	    int task_len = tasks[task_no].data.size;
+	    best_len = min(task_len, best_len);
+	}
+	best_size[ch] = best_len;
+    }
+        
+    int bitsBest = best_size[2] + best_size[3]; // MidSide
+    int chMask = 2 | (3 << 2);
+    int bits = best_size[3] + best_size[1];
+    chMask = select(chMask, 3 | (1 << 2), bits < bitsBest); // RightSide
+    bitsBest = min(bits, bitsBest);
+    bits = best_size[0] + best_size[3];
+    chMask = select(chMask, 0 | (3 << 2), bits < bitsBest); // LeftSide
+    bitsBest = min(bits, bitsBest);
+    bits = best_size[0] + best_size[1];
+    chMask = select(chMask, 0 | (1 << 2), bits < bitsBest); // LeftRight
+    bitsBest = min(bits, bitsBest);
+    for (int ich = 0; ich < 2; ich++)
+    {
+	int ch = select(chMask & 3, chMask >> 2, ich > 0);
+	int roffs = tasks[(get_global_id(0) * 4 + ich) * taskCount].data.samplesOffs;
+	int nonSelectedNo = 0;
+	for (int i = 0; i < taskCount; i++)
+	{
+	    int no = (get_global_id(0) * 4 + ch) * taskCount + i;
+	    selectedTasksBestMethod[(get_global_id(0) * 2 + ich) * taskCount + i] = no;
+	    tasks[no].data.residualOffs = roffs;
+	    int selectedFound = 0;
+	    for(int selectedNo = 0; selectedNo < selectedCount; selectedNo++)
+		selectedFound |= (selectedTasks[(get_global_id(0) * 4 + ch) * selectedCount + selectedNo] == no);
+	    if (!selectedFound)
+		selectedTasksSecondEstimate[(get_global_id(0) * 2 + ich) * (taskCount - selectedCount) + nonSelectedNo++] = no;
+	}
     }
 }
 
-__kernel __attribute__((reqd_work_group_size(32, 1, 1)))
+__kernel
 void clChooseBestMethod(
+    __global FLACCLSubframeTask *tasks_out,
     __global FLACCLSubframeTask *tasks,
+    __global int*selectedTasks,
     int taskCount
     )
 {
-    int best_length = 0x7fffffff;
+    int best_no = selectedTasks[get_global_id(0) * taskCount];
-    int best_index = 0;
+    int best_len = tasks[best_no].data.size;
-    const int tid = get_local_id(0);
+    for (int i = 1; i < taskCount; i++)
-    
-    for (int taskNo = 0; taskNo < taskCount; taskNo++)
     {
-	if (tid == 0)
+	int task_no = selectedTasks[get_global_id(0) * taskCount + i];
-	{
+	int task_len = tasks[task_no].data.size;
-	    int len = tasks[taskNo + taskCount * get_group_id(0)].data.size;
+	best_no = select(best_no, task_no, task_len < best_len);
-	    if (len < best_length)
+	best_len = min(best_len, task_len);
-	    {
-		best_length = len;
-		best_index = taskNo;
-	    }
-	}
-	barrier(CLK_LOCAL_MEM_FENCE);
     }
-
+    tasks_out[get_global_id(0)] = tasks[best_no];
-    if (tid == 0)
-	tasks[taskCount * get_group_id(0)].data.best_index = taskCount * get_group_id(0) + best_index;
 }
 
-__kernel __attribute__((reqd_work_group_size(64, 1, 1)))
+#ifdef __CPU__
-void clCopyBestMethod(
+// get_group_id(0) == task index
-    __global FLACCLSubframeTask *tasks_out,
+__kernel __attribute__((reqd_work_group_size(1, 1, 1)))
-    __global FLACCLSubframeTask *tasks,
+void clEncodeResidual(
-    int count
+    __global int *residual,
+    __global int *samples,
+    __global FLACCLSubframeTask *tasks
     )
 {
-    __local int best_index;
+    FLACCLSubframeTask task = tasks[get_group_id(0)];
-    if (get_local_id(0) == 0)
+    int bs = task.data.blocksize;
-	best_index = tasks[count * get_group_id(0)].data.best_index;
+    int ro = task.data.residualOrder;
-    barrier(CLK_LOCAL_MEM_FENCE);
+    __global int *data = &samples[task.data.samplesOffs];
-    if (get_local_id(0) < sizeof(FLACCLSubframeTask)/sizeof(int))
+    SWITCH_N(residual[task.data.residualOffs + pos] = t);
-	((__global int*)(tasks_out + get_group_id(0)))[get_local_id(0)] = ((__global int*)(tasks + best_index))[get_local_id(0)];
 }
-
+#else
-__kernel __attribute__((reqd_work_group_size(64, 1, 1)))
-void clCopyBestMethodStereo(
-    __global FLACCLSubframeTask *tasks_out,
-    __global FLACCLSubframeTask *tasks,
-    int count
-    )
-{
-    __local struct {
-	int best_index[4];
-	int best_size[4];
-	int lr_index[2];
-    } shared;
-    if (get_local_id(0) < 4)
-	shared.best_index[get_local_id(0)] = tasks[count * (get_group_id(0) * 4 + get_local_id(0))].data.best_index;
-    barrier(CLK_LOCAL_MEM_FENCE);
-    if (get_local_id(0) < 4)
-	shared.best_size[get_local_id(0)] = tasks[shared.best_index[get_local_id(0)]].data.size;
-    barrier(CLK_LOCAL_MEM_FENCE);
-    if (get_local_id(0) == 0)
-    {
-	int bitsBest = shared.best_size[2] + shared.best_size[3]; // MidSide
-	shared.lr_index[0] = shared.best_index[2];
-	shared.lr_index[1] = shared.best_index[3];
-	if (bitsBest > shared.best_size[3] + shared.best_size[1]) // RightSide
-	{
-	    bitsBest = shared.best_size[3] + shared.best_size[1];
-	    shared.lr_index[0] = shared.best_index[3];
-	    shared.lr_index[1] = shared.best_index[1];
-	}
-	if (bitsBest > shared.best_size[0] + shared.best_size[3]) // LeftSide
-	{
-	    bitsBest = shared.best_size[0] + shared.best_size[3];
-	    shared.lr_index[0] = shared.best_index[0];
-	    shared.lr_index[1] = shared.best_index[3];
-	}
-	if (bitsBest > shared.best_size[0] + shared.best_size[1]) // LeftRight
-	{
-	    bitsBest = shared.best_size[0] + shared.best_size[1];
-	    shared.lr_index[0] = shared.best_index[0];
-	    shared.lr_index[1] = shared.best_index[1];
-	}
-    }
-    barrier(CLK_LOCAL_MEM_FENCE);
-    if (get_local_id(0) < sizeof(FLACCLSubframeTask)/sizeof(int))
-	((__global int*)(tasks_out + 2 * get_group_id(0)))[get_local_id(0)] = ((__global int*)(tasks + shared.lr_index[0]))[get_local_id(0)];
-    if (get_local_id(0) == 0)
-	tasks_out[2 * get_group_id(0)].data.residualOffs = tasks[shared.best_index[0]].data.residualOffs;
-    if (get_local_id(0) < sizeof(FLACCLSubframeTask)/sizeof(int))
-	((__global int*)(tasks_out + 2 * get_group_id(0) + 1))[get_local_id(0)] = ((__global int*)(tasks + shared.lr_index[1]))[get_local_id(0)];
-    if (get_local_id(0) == 0)
-	tasks_out[2 * get_group_id(0) + 1].data.residualOffs = tasks[shared.best_index[1]].data.residualOffs;
-}
-
 // get_group_id(0) == task index
 __kernel __attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
 void clEncodeResidual(
@@ -737,7 +1081,38 @@ void clEncodeResidual(
 	data[tid] = nextData;
     }
 }
+#endif
 
+#ifdef __CPU__
+__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;
+    }
+}
+#else
 // get_group_id(0) == partition index / (GROUP_SIZE / 16)
 // get_group_id(1) == task index
 __kernel __attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
@@ -794,7 +1169,31 @@ void clCalcPartition(
 	}
     }
 }
+#endif
 
+#ifdef __CPU__
+// 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));
+}
+#else
 // get_group_id(0) == task index
 __kernel __attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
 void clCalcPartition16(
@@ -881,7 +1280,33 @@ void clCalcPartition16(
 	    partition_lengths[lpos] = min(0x7fffff, s) + (16 - select(0, ro, offs < 16)) * (k + 1);
     }    
 }
+#endif
 
+#ifdef __CPU__
+// 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;
+    }
+}
+#else
 // Sums partition lengths for a certain k == get_group_id(0)
 // Requires 128 threads
 // get_group_id(0) == k
@@ -914,7 +1339,42 @@ void clSumPartition(
     if (get_local_size(0) + get_local_id(0) < (1 << max_porder))
 	partition_lengths[pos + (1 << max_porder) + get_local_size(0) + get_local_id(0)] = data[get_local_size(0) + get_local_id(0)];
 }
+#endif
 
+#ifdef __CPU__
+// 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;
+    }
+}
+#else
 // Finds optimal rice parameter for each partition.
 // get_group_id(0) == task index
 __kernel __attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
@@ -943,7 +1403,69 @@ void clFindRiceParameter(
 	rice_parameters[(get_group_id(0) << (max_porder + 2)) + (1 << (max_porder + 1)) + offs] = best_l;
     }
 }
+#endif
 
+#ifdef __CPU__
+// 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];
+}
+#else
 // get_group_id(0) == task index
 __kernel __attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
 void clFindPartitionOrder(
@@ -998,3 +1520,4 @@ void clFindPartitionOrder(
     // FIXME: should be bytes?
 }
 #endif
+#endif