mirror of
https://bitbucket.org/mfeemster/fractorium.git
synced 2025-01-21 05:00:06 -05:00
Merged in mmastriani/fractorium_michel (pull request #50)
macOS improvements
This commit is contained in:
Michel Mastriani
Matt Feemster
commit
3f7816412f
4
Source/EmberCL/EmberCLPch.h
Normal file → Executable file
4
Source/EmberCL/EmberCLPch.h
Normal file → Executable file
@ -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
|
||||
|
||||
81
Source/EmberCL/IterOpenCLKernelCreator.cpp
Normal file → Executable file
81
Source/EmberCL/IterOpenCLKernelCreator.cpp
Normal file → Executable file
@ -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"
|
||||
"\n"
|
||||
#ifndef KNL_USE_GLOBAL_CONSEC
|
||||
" 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();
|
||||
}
|
||||
|
||||
12
Source/EmberCL/RendererCL.cpp
Normal file → Executable file
12
Source/EmberCL/RendererCL.cpp
Normal file → Executable file
@ -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();
|
||||
//cout << "itersRemaining " << itersRemaining << ", iterCountPerKernel " << iterCountPerKernel << ", iterCountThisLaunch " << iterCountThisLaunch << "\n";
|
||||
size_t iterCountThisLaunch = iterCountPerKernel * IterGridKernelWidth() * IterGridKernelHeight();
|
||||
//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(),
|
||||
|
||||
3
Source/EmberCL/RendererCL.h
Normal file → Executable file
3
Source/EmberCL/RendererCL.h
Normal file → Executable file
@ -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.
|
||||
|
||||
Loading…
Reference in New Issue
Block a user