macOS improvements

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,10 +462,12 @@ string IterOpenCLKernelCreator<T>::CreateIterKernelString(const Ember<T>& ember,
 	   "\n"
 	   "	for (i = 0; i < itersToDo; i++)\n"
 	   "	{\n"
+#ifndef KNL_USE_GLOBAL_CONSEC
        "		consec = 0;\n"
-	   "\n"
+#else
-	   "		do\n"
+       "		consec[blockStartThreadIndex] = 0;\n"
-	   "		{\n";
+#endif
+       "\n";
 
     if (ember.XformCount() > 1)
     {
@@ -484,6 +494,10 @@ string IterOpenCLKernelCreator<T>::CreateIterKernelString(const Ember<T>& ember,
         }
     }
 
+    os <<
+	   "		do\n"
+	   "		{\n";
+
 	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"
@@ -772,6 +794,7 @@ string IterOpenCLKernelCreator<T>::CreateIterKernelString(const Ember<T>& ember,
 #endif
 	   os <<
        "	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.
 
@@ -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.