Merge branch 'master' of https://bitbucket.org/mfeemster/fractorium

Builds/QtCreator/defaults.pri

@@ -217,8 +217,8 @@ win32 {
 
 macx {
         QMAKE_MAC_SDK = macosx10.14
-        QMAKE_MACOSX_DEPLOYMENT_TARGET = 10.14
+        QMAKE_MACOSX_DEPLOYMENT_TARGET = 10.10
-        QMAKE_CXXFLAGS += -mmacosx-version-min=10.9 -arch x86_64
+        QMAKE_CXXFLAGS += -mmacosx-version-min=10.10 -arch x86_64
 	QMAKE_CXXFLAGS += -stdlib=libc++
 }
 

Source/EmberCL/EmberCLPch.h

@@ -53,3 +53,7 @@ using namespace std;
 using namespace EmberNs;
 //#define TEST_CL 1
 //#define TEST_CL_BUFFERS 1
+
+#ifdef __APPLE__
+    #define KNL_USE_GLOBAL_CONSEC
+#endif

Source/EmberCL/IterOpenCLKernelCreator.cpp

@@ -359,7 +359,13 @@ string IterOpenCLKernelCreator<T>::CreateIterKernelString(const Ember<T>& ember,
 	os <<
 	   "	uint histSize,\n"
 	   "	__read_only image2d_t palette,\n"
-	   "	__global Point* points\n"
+       "	__global Point* points"
+#ifndef KNL_USE_GLOBAL_CONSEC
+       "\n"
+#else
+       ",\n"
+       "	__global uchar* consec\n"
+#endif
 	   "\t)\n"
 	   "{\n"
 	   "	bool fuse, ok;\n"
@@ -368,7 +374,9 @@ string IterOpenCLKernelCreator<T>::CreateIterKernelString(const Ember<T>& ember,
 	   "	uint blockStartIndex = BLOCK_START_INDEX_IN_GRID_2D;\n"
 	   "	uint blockStartThreadIndex = blockStartIndex + threadIndex;\n"
 	   "	uint i, itersToDo;\n"
+#ifndef KNL_USE_GLOBAL_CONSEC
 	   "	uint consec = 0;\n"
+#endif
 	   //"	int badvals = 0;\n"
 	   "	uint histIndex;\n"
 	   "	real_t p00, p01;\n"
@@ -454,36 +462,42 @@ string IterOpenCLKernelCreator<T>::CreateIterKernelString(const Ember<T>& ember,
 	   "\n"
 	   "	for (i = 0; i < itersToDo; i++)\n"
 	   "	{\n"
-	   "		consec = 0;\n"
+#ifndef KNL_USE_GLOBAL_CONSEC
-	   "\n"
+       "		consec = 0;\n"
+#else
+       "		consec[blockStartThreadIndex] = 0;\n"
+#endif
+       "\n";
+
+    if (ember.XformCount() > 1)
+    {
+        //If xaos is present, the a hybrid of the cuburn method is used.
+        //This makes each thread in a row pick the same offset into a distribution, using xfsel.
+        //However, the distribution the offset is in, is determined by firstPoint.m_LastXfUsed.
+        if (ember.XaosPresent())
+        {
+            os <<
+#ifdef STRAIGHT_RAND
+               "		secondPoint.m_LastXfUsed = xformDistributions[(MwcNext(&mwc) & " << CHOOSE_XFORM_GRAIN_M1 << "u) + (" << CHOOSE_XFORM_GRAIN << "u * (firstPoint.m_LastXfUsed + 1u))];\n\n";
+#else
+               "		secondPoint.m_LastXfUsed = xformDistributions[xfsel[THREAD_ID_Y] + (" << CHOOSE_XFORM_GRAIN << "u * (firstPoint.m_LastXfUsed + 1u))];\n\n";//Partial cuburn hybrid.
+#endif
+        }
+        else
+        {
+            os <<
+#ifdef STRAIGHT_RAND
+               "		secondPoint.m_LastXfUsed = xformDistributions[MwcNext(&mwc) & " << CHOOSE_XFORM_GRAIN_M1 << "u];\n\n";//For testing, using straight rand flam4/fractron style instead of cuburn.
+#else
+               "		secondPoint.m_LastXfUsed = xformDistributions[xfsel[THREAD_ID_Y]];\n\n";
+#endif
+        }
+    }
+
+    os <<
 	   "		do\n"
 	   "		{\n";
 
-	if (ember.XformCount() > 1)
-	{
-		//If xaos is present, the a hybrid of the cuburn method is used.
-		//This makes each thread in a row pick the same offset into a distribution, using xfsel.
-		//However, the distribution the offset is in, is determined by firstPoint.m_LastXfUsed.
-		if (ember.XaosPresent())
-		{
-			os <<
-#ifdef STRAIGHT_RAND
-			   "			secondPoint.m_LastXfUsed = xformDistributions[(MwcNext(&mwc) & " << CHOOSE_XFORM_GRAIN_M1 << "u) + (" << CHOOSE_XFORM_GRAIN << "u * (firstPoint.m_LastXfUsed + 1u))];\n\n";
-#else
-			   "			secondPoint.m_LastXfUsed = xformDistributions[xfsel[THREAD_ID_Y] + (" << CHOOSE_XFORM_GRAIN << "u * (firstPoint.m_LastXfUsed + 1u))];\n\n";//Partial cuburn hybrid.
-#endif
-		}
-		else
-		{
-			os <<
-#ifdef STRAIGHT_RAND
-			   "			secondPoint.m_LastXfUsed = xformDistributions[MwcNext(&mwc) & " << CHOOSE_XFORM_GRAIN_M1 << "u];\n\n";//For testing, using straight rand flam4/fractron style instead of cuburn.
-#else
-			   "			secondPoint.m_LastXfUsed = xformDistributions[xfsel[THREAD_ID_Y]];\n\n";
-#endif
-		}
-	}
-
 	for (i = 0; i < ember.XformCount(); i++)
 	{
 		if (ember.XformCount() > 1)
@@ -525,11 +539,19 @@ string IterOpenCLKernelCreator<T>::CreateIterKernelString(const Ember<T>& ember,
 	   "				firstPoint.m_Y = MwcNextFRange(&mwc, -ember->m_RandPointRange, ember->m_RandPointRange);\n"
 	   "				firstPoint.m_Z = 0.0;\n"
 	   "				firstPoint.m_ColorX = secondPoint.m_ColorX;\n"
+#ifndef KNL_USE_GLOBAL_CONSEC
 	   "				consec++;\n"
+#else
+       "				consec[blockStartThreadIndex]++;\n"
+#endif
 	   //"				badvals++;\n"
 	   "			}\n"
 	   "		}\n"
+#ifndef KNL_USE_GLOBAL_CONSEC
 	   "		while (!ok && consec < 5);\n"
+#else
+       "		while (!ok && consec[blockStartThreadIndex] < 5);\n"
+#endif
 	   "\n"
 	   "		if (!ok)\n"
 	   "		{\n"
@@ -771,7 +793,8 @@ string IterOpenCLKernelCreator<T>::CreateIterKernelString(const Ember<T>& ember,
 
 #endif
 	   os <<
-	   "	barrier(CLK_GLOBAL_MEM_FENCE);\n"
+       "	barrier(CLK_GLOBAL_MEM_FENCE);\n"
+       //"	printf(\"Global ID0: %d Global ID1: %d WorkDim: %d ThreadIndex: %d\\n\", get_global_id(0), get_global_id(1), get_work_dim(), blockStartThreadIndex);\n"
 	   "}\n";
 	return os.str();
 }

Source/EmberCL/RendererCL.cpp

@@ -727,7 +727,9 @@ bool RendererCL<T, bucketT>::Alloc(bool histOnly)
 		if (b && !(b = device->m_Wrapper.AddBuffer(m_HistBufferName, size)))                                           { ErrorStr(loc, "Failed to set histogram buffer", device.get()); break; }//Histogram. Will memset to zero later.
 
 		if (b && !(b = device->m_Wrapper.AddBuffer(m_PointsBufferName, IterGridKernelCount() * sizeof(PointCL<T>))))   { ErrorStr(loc, "Failed to set points buffer", device.get()); break; }//Points between iter calls.
-
+#ifdef KNL_USE_GLOBAL_CONSEC
+        if (b && !(b = device->m_Wrapper.AddBuffer(m_ConsecBufferName, IterGridKernelCount() * sizeof(cl_uchar))))                     { ErrorStr(loc, "Failed to set consec buffer", device.get()); break; }//Global sequence.
+#endif
 		if (m_VarStates.size())
 			if (b && !(b = device->m_Wrapper.AddBuffer(m_VarStateBufferName, SizeOf(m_VarStates))))                        { ErrorStr(loc, "Failed to set variation state buffer", device.get()); break; }//Points between iter calls.
 
@@ -1085,8 +1087,8 @@ bool RendererCL<T, bucketT>::RunIter(size_t iterCount, size_t temporalSample, si
 			//Similar to what's done in the base class.
 			//The number of iters per thread must be adjusted if they've requested less iters than is normally ran in a grid (256 * 256 * 64 * 2 = 32,768).
 			uint iterCountPerKernel = std::min<uint>(uint(adjustedIterCountPerKernel), uint(ceil(double(itersRemaining) / IterGridKernelCount())));
-			size_t iterCountThisLaunch = iterCountPerKernel * IterGridKernelWidth() * IterGridKernelHeight();
+            size_t iterCountThisLaunch = iterCountPerKernel * IterGridKernelWidth() * IterGridKernelHeight();
-			//cout << "itersRemaining " << itersRemaining << ", iterCountPerKernel " << iterCountPerKernel << ", iterCountThisLaunch " << iterCountThisLaunch << "\n";
+            //cout << "itersRemaining " << itersRemaining << ", iterCountPerKernel " << iterCountPerKernel << ", iterCountThisLaunch " << iterCountThisLaunch << "\n";
 
 			if (b && !(b = wrapper.SetArg	   (kernelIndex, argIndex++, iterCountPerKernel)))        { ErrorStr(loc, "Setting iter count argument failed", m_Devices[dev].get()); }//Number of iters for each thread to run.
 
@@ -1116,7 +1118,9 @@ bool RendererCL<T, bucketT>::RunIter(size_t iterCount, size_t temporalSample, si
 			if (b && !(b = wrapper.SetImageArg (kernelIndex, argIndex++, false, "Palette")))          { ErrorStr(loc, "Setting palette argument failed", m_Devices[dev].get()); }//Palette.
 
 			if (b && !(b = wrapper.SetBufferArg(kernelIndex, argIndex++, m_PointsBufferName)))        { ErrorStr(loc, "Setting points buffer argument failed", m_Devices[dev].get()); }//Random start points.
-
+#ifdef KNL_USE_GLOBAL_CONSEC
+            if (b && !(b = wrapper.SetBufferArg(kernelIndex, argIndex++, m_ConsecBufferName)))        { ErrorStr(loc, "Setting consec buffer argument failed", m_Devices[dev].get()); }//Global sequence.
+#endif
 			if (b && !(b = wrapper.RunKernel(kernelIndex,
 											 IterGridKernelWidth(),//Total grid dims.
 											 IterGridKernelHeight(),

Source/EmberCL/RendererCL.h

@@ -239,6 +239,9 @@ private:
 	string m_AccumBufferName = "Accum";
 	string m_FinalImageName = "Final";
 	string m_PointsBufferName = "Points";
+#ifdef KNL_USE_GLOBAL_CONSEC
+    string m_ConsecBufferName = "Consec";
+#endif
 	string m_VarStateBufferName = "VarState";
 
 	//Kernels.

Source/Fractorium/FractoriumPalette.cpp

@@ -85,6 +85,7 @@ size_t FractoriumEmberController<T>::InitPaletteList(const QString& s)
 		}
 	}
 
+    m_Fractorium->ui.PaletteFilenameCombo->model()->sort(0);
 	return m_PaletteList->Size();
 }
 

Source/Fractorium/PaletteEditor/PaletteEditor.cpp

@@ -58,6 +58,8 @@ PaletteEditor::PaletteEditor(QWidget* p) :
 		ui->PaletteFilenameCombo->addItem(info.fileName());
 	}
 
+    ui->PaletteFilenameCombo->model()->sort(0);
+
 	if (ui->PaletteFilenameCombo->count() > 0)
 		m_CurrentPaletteFilePath = ui->PaletteFilenameCombo->itemText(0).toStdString();
 }