2014-07-08 03:11:14 -04:00
# include "EmberCLPch.h"
# include "FinalAccumOpenCLKernelCreator.h"
namespace EmberCLns
{
/// <summary>
/// Constructor that creates all kernel strings.
/// The caller will access these strings through the accessor functions.
/// </summary>
2015-08-10 23:10:23 -04:00
FinalAccumOpenCLKernelCreator : : FinalAccumOpenCLKernelCreator ( bool doublePrecision )
2014-07-08 03:11:14 -04:00
{
2015-08-10 23:10:23 -04:00
m_DoublePrecision = doublePrecision ;
2017-07-22 16:43:35 -04:00
m_GammaCorrectionWithoutAlphaCalcKernel = CreateGammaCorrectionKernelString ( ) ;
m_FinalAccumEarlyClipWithoutAlphaCalcWithAlphaAccumKernel = CreateFinalAccumKernelString ( true ) ;
m_FinalAccumLateClipWithoutAlphaCalcWithAlphaAccumKernel = CreateFinalAccumKernelString ( false ) ;
2014-07-08 03:11:14 -04:00
}
/// <summary>
/// Kernel source and entry point properties, getters only.
/// </summary>
2015-08-12 21:51:07 -04:00
const string & FinalAccumOpenCLKernelCreator : : FinalAccumEarlyClipWithoutAlphaCalcWithAlphaAccumKernel ( ) const { return m_FinalAccumEarlyClipWithoutAlphaCalcWithAlphaAccumKernel ; }
const string & FinalAccumOpenCLKernelCreator : : FinalAccumEarlyClipWithoutAlphaCalcWithAlphaAccumEntryPoint ( ) const { return m_FinalAccumEarlyClipWithoutAlphaCalcWithAlphaAccumEntryPoint ; }
const string & FinalAccumOpenCLKernelCreator : : FinalAccumLateClipWithoutAlphaCalcWithAlphaAccumKernel ( ) const { return m_FinalAccumLateClipWithoutAlphaCalcWithAlphaAccumKernel ; }
const string & FinalAccumOpenCLKernelCreator : : FinalAccumLateClipWithoutAlphaCalcWithAlphaAccumEntryPoint ( ) const { return m_FinalAccumLateClipWithoutAlphaCalcWithAlphaAccumEntryPoint ; }
2014-07-08 03:11:14 -04:00
2017-07-22 16:43:35 -04:00
const string & FinalAccumOpenCLKernelCreator : : GammaCorrectionEntryPoint ( ) const { return m_GammaCorrectionWithoutAlphaCalcEntryPoint ; }
const string & FinalAccumOpenCLKernelCreator : : GammaCorrectionKernel ( ) const { return m_GammaCorrectionWithoutAlphaCalcKernel ; }
2014-07-08 03:11:14 -04:00
/// <summary>
/// Get the final accumulation entry point.
/// </summary>
/// <param name="earlyClip">True if early clip is desired, else false.</param>
/// <returns>The name of the final accumulation entry point kernel function</returns>
2017-07-22 16:43:35 -04:00
const string & FinalAccumOpenCLKernelCreator : : FinalAccumEntryPoint ( bool earlyClip ) const
2014-07-08 03:11:14 -04:00
{
if ( earlyClip )
2017-07-22 16:43:35 -04:00
return FinalAccumEarlyClipWithoutAlphaCalcWithAlphaAccumEntryPoint ( ) ;
2014-07-08 03:11:14 -04:00
else
2017-07-22 16:43:35 -04:00
return FinalAccumLateClipWithoutAlphaCalcWithAlphaAccumEntryPoint ( ) ;
2014-07-08 03:11:14 -04:00
}
/// <summary>
/// Get the final accumulation kernel string.
/// </summary>
/// <param name="earlyClip">True if early clip is desired, else false.</param>
/// <returns>The final accumulation kernel string</returns>
2017-07-22 16:43:35 -04:00
const string & FinalAccumOpenCLKernelCreator : : FinalAccumKernel ( bool earlyClip ) const
2014-07-08 03:11:14 -04:00
{
if ( earlyClip )
2017-07-22 16:43:35 -04:00
return FinalAccumEarlyClipWithoutAlphaCalcWithAlphaAccumKernel ( ) ;
2014-07-08 03:11:14 -04:00
else
2017-07-22 16:43:35 -04:00
return FinalAccumLateClipWithoutAlphaCalcWithAlphaAccumKernel ( ) ;
2014-07-08 03:11:14 -04:00
}
/// <summary>
/// Create the final accumulation kernel string
/// </summary>
/// <param name="earlyClip">True if early clip is desired, else false.</param>
/// <returns>The final accumulation kernel string</returns>
2017-07-22 16:43:35 -04:00
string FinalAccumOpenCLKernelCreator : : CreateFinalAccumKernelString ( bool earlyClip )
2014-07-08 03:11:14 -04:00
{
ostringstream os ;
os < <
2016-02-24 00:01:02 -05:00
ConstantDefinesString ( m_DoublePrecision ) < <
UnionCLStructString < <
RgbToHsvFunctionString < <
HsvToRgbFunctionString < <
CalcAlphaFunctionString < <
CurveAdjustFunctionString < <
SpatialFilterCLStructString ;
2014-07-08 03:11:14 -04:00
if ( earlyClip )
{
2017-07-22 16:43:35 -04:00
os < < " __kernel void " < < m_FinalAccumEarlyClipWithoutAlphaCalcWithAlphaAccumEntryPoint < < " ( \n " ;
2014-07-08 03:11:14 -04:00
}
else
{
os < <
2016-02-24 00:01:02 -05:00
CreateCalcNewRgbFunctionString ( false ) < <
2017-07-22 16:43:35 -04:00
CreateGammaCorrectionFunctionString ( false , true ) < <
" __kernel void " < < m_FinalAccumLateClipWithoutAlphaCalcWithAlphaAccumEntryPoint < < " ( \n " ;
2014-07-08 03:11:14 -04:00
}
os < <
2016-02-24 00:01:02 -05:00
" const __global real4reals_bucket* accumulator, \n "
" __write_only image2d_t pixels, \n "
" __constant SpatialFilterCL* spatialFilter, \n "
" __constant real_bucket_t* filterCoefs, \n "
2017-07-22 16:43:35 -04:00
" __global real4reals_bucket* csa, \n "
" const uint doCurves \n "
2016-02-24 00:01:02 -05:00
" \t ) \n "
" { \n "
" \n "
" if ((GLOBAL_ID_Y >= spatialFilter->m_FinalRasH) || (GLOBAL_ID_X >= spatialFilter->m_FinalRasW)) \n "
" return; \n "
" \n "
" uint accumX = spatialFilter->m_DensityFilterOffset + (GLOBAL_ID_X * spatialFilter->m_Supersample); \n "
" uint accumY = spatialFilter->m_DensityFilterOffset + (GLOBAL_ID_Y * spatialFilter->m_Supersample); \n "
2016-03-28 21:49:10 -04:00
" uint clampedFilterH = min((uint)spatialFilter->m_FilterWidth, spatialFilter->m_SuperRasH - accumY); "
" uint clampedFilterW = min((uint)spatialFilter->m_FilterWidth, spatialFilter->m_SuperRasW - accumX); "
2016-02-24 00:01:02 -05:00
" int2 finalCoord; \n "
" finalCoord.x = GLOBAL_ID_X; \n "
" finalCoord.y = (int)((spatialFilter->m_YAxisUp == 1) ? ((spatialFilter->m_FinalRasH - GLOBAL_ID_Y) - 1) : GLOBAL_ID_Y); \n "
" float4floats finalColor; \n "
" int ii, jj; \n "
" uint filterKRowIndex; \n "
" const __global real4reals_bucket* accumBucket; \n "
" real4reals_bucket newBucket; \n "
" newBucket.m_Real4 = 0; \n "
" \n "
2016-03-28 21:49:10 -04:00
" for (jj = 0; jj < clampedFilterH; jj++) \n "
2016-02-24 00:01:02 -05:00
" { \n "
2016-03-28 21:49:10 -04:00
" filterKRowIndex = jj * spatialFilter->m_FilterWidth; \n " //Use the full, non-clamped width to get the filter value.
2016-02-24 00:01:02 -05:00
" \n "
2016-03-28 21:49:10 -04:00
" for (ii = 0; ii < clampedFilterW; ii++) \n "
2016-02-24 00:01:02 -05:00
" { \n "
2016-03-28 21:49:10 -04:00
" real_bucket_t k = filterCoefs[filterKRowIndex + ii]; \n "
2016-02-24 00:01:02 -05:00
" \n "
2016-03-28 21:49:10 -04:00
" accumBucket = accumulator + ((accumY + jj) * spatialFilter->m_SuperRasW) + (accumX + ii); \n "
2016-02-24 00:01:02 -05:00
" newBucket.m_Real4 += (k * accumBucket->m_Real4); \n "
" } \n "
" } \n "
" \n " ;
2014-07-08 03:11:14 -04:00
2015-03-21 18:27:37 -04:00
if ( earlyClip ) //If early clip, simply assign values directly to the temp float4 since they've been gamma corrected already, then write it straight to the output image below.
2014-07-08 03:11:14 -04:00
{
os < <
2016-02-24 00:01:02 -05:00
" finalColor.m_Float4.x = (float)newBucket.m_Real4.x; \n " //CPU side clamps, skip here because write_imagef() does the clamping for us.
" finalColor.m_Float4.y = (float)newBucket.m_Real4.y; \n "
2017-07-22 16:43:35 -04:00
" finalColor.m_Float4.z = (float)newBucket.m_Real4.z; \n "
" finalColor.m_Float4.w = (float)newBucket.m_Real4.w; \n " ;
2014-07-08 03:11:14 -04:00
}
else
{
2016-03-28 21:49:10 -04:00
//Late clip, so must gamma correct from the temp newBucket to temp finalColor float4.
2015-08-10 23:10:23 -04:00
if ( m_DoublePrecision )
2014-07-08 03:11:14 -04:00
{
os < <
2016-02-24 00:01:02 -05:00
" real4reals_bucket realFinal; \n "
" \n "
2017-07-22 16:43:35 -04:00
" GammaCorrectionFloats(&newBucket, &(spatialFilter->m_Background[0]), spatialFilter->m_Gamma, spatialFilter->m_LinRange, spatialFilter->m_Vibrancy, spatialFilter->m_HighlightPower, &(realFinal.m_Reals[0])); \n "
2016-02-24 00:01:02 -05:00
" finalColor.m_Float4.x = (float)realFinal.m_Real4.x; \n "
" finalColor.m_Float4.y = (float)realFinal.m_Real4.y; \n "
" finalColor.m_Float4.z = (float)realFinal.m_Real4.z; \n "
" finalColor.m_Float4.w = (float)realFinal.m_Real4.w; \n "
;
2014-07-08 03:11:14 -04:00
}
else
{
os < <
2017-07-22 16:43:35 -04:00
" GammaCorrectionFloats(&newBucket, &(spatialFilter->m_Background[0]), spatialFilter->m_Gamma, spatialFilter->m_LinRange, spatialFilter->m_Vibrancy, spatialFilter->m_HighlightPower, &(finalColor.m_Floats[0])); \n " ;
2014-07-08 03:11:14 -04:00
}
}
os < <
2016-02-24 00:01:02 -05:00
" \n "
" if (doCurves) \n "
" { \n "
" CurveAdjust(csa, &(finalColor.m_Floats[0]), 1); \n "
" CurveAdjust(csa, &(finalColor.m_Floats[1]), 2); \n "
" CurveAdjust(csa, &(finalColor.m_Floats[2]), 3); \n "
" } \n "
" \n "
" write_imagef(pixels, finalCoord, finalColor.m_Float4); \n " //Use write_imagef instead of write_imageui because only the former works when sharing with an OpenGL texture.
" barrier(CLK_GLOBAL_MEM_FENCE); \n " //Required, or else page tearing will occur during interactive rendering.
" } \n "
;
2014-07-08 03:11:14 -04:00
return os . str ( ) ;
}
/// <summary>
/// Creates the gamma correction function string.
/// This is not a full kernel, just a function that is used in the kernels.
/// </summary>
/// <param name="globalBucket">True if writing to a global buffer (early clip), else false (late clip).</param>
/// <param name="finalOut">True if writing to global buffer (late clip), else false (early clip).</param>
/// <returns>The gamma correction function string</returns>
2017-07-22 16:43:35 -04:00
string FinalAccumOpenCLKernelCreator : : CreateGammaCorrectionFunctionString ( bool globalBucket , bool finalOut )
2014-07-08 03:11:14 -04:00
{
ostringstream os ;
string dataType ;
string unionMember ;
2015-08-10 23:10:23 -04:00
dataType = " real_bucket_t " ;
2014-07-08 03:11:14 -04:00
//Use real_t for all cases, early clip and final accum.
2017-07-22 16:43:35 -04:00
os < < " void GammaCorrectionFloats( " < < ( globalBucket ? " __global " : " " ) < < " real4reals_bucket* bucket, __constant real_bucket_t* background, real_bucket_t g, real_bucket_t linRange, real_bucket_t vibrancy, real_bucket_t highlightPower, " < < ( finalOut ? " " : " __global " ) < < " real_bucket_t* correctedChannels) \n " ;
os < < " { \n "
< < " real_bucket_t alpha, ls, tmp, a; \n "
< < " real4reals_bucket newRgb; \n "
< < " \n "
< < " if (bucket->m_Reals[3] <= 0) \n "
< < " { \n "
< < " alpha = 0; \n "
< < " ls = 0; \n "
< < " } \n "
< < " else \n "
< < " { \n "
< < " tmp = bucket->m_Reals[3]; \n "
< < " alpha = CalcAlpha(tmp, g, linRange); \n "
< < " ls = vibrancy * alpha / tmp; \n "
< < " alpha = clamp(alpha, (real_bucket_t)0.0, (real_bucket_t)1.0); \n "
< < " } \n "
< < " \n "
< < " CalcNewRgb(bucket, ls, highlightPower, &newRgb); \n "
< < " \n "
< < " for (uint rgbi = 0; rgbi < 3; rgbi++) \n "
< < " { \n "
< < " a = newRgb.m_Reals[rgbi] + ((1.0 - vibrancy) * pow(fabs(bucket->m_Reals[rgbi]), g)); \n "
< < " a += ((1.0 - alpha) * background[rgbi]); \n "
< < " correctedChannels[rgbi] = ( " < < dataType < < " )clamp(a, (real_bucket_t)0.0, (real_bucket_t)1.0); \n "
< < " } \n "
< < " \n "
< < " correctedChannels[3] = ( " < < dataType < < " )alpha; \n "
< < " } \n "
< < " \n " ;
2014-07-08 03:11:14 -04:00
return os . str ( ) ;
}
/// <summary>
/// OpenCL equivalent of Palette::CalcNewRgb().
/// </summary>
/// <param name="globalBucket">True if writing the corrected value to a global buffer (early clip), else false (late clip).</param>
/// <returns>The CalcNewRgb function string</returns>
2015-08-10 23:10:23 -04:00
string FinalAccumOpenCLKernelCreator : : CreateCalcNewRgbFunctionString ( bool globalBucket )
2014-07-08 03:11:14 -04:00
{
ostringstream os ;
os < <
2016-02-24 00:01:02 -05:00
" static void CalcNewRgb( " < < ( globalBucket ? " __global " : " " ) < < " real4reals_bucket* oldRgb, real_bucket_t ls, real_bucket_t highPow, real4reals_bucket* newRgb) \n "
" { \n "
" int rgbi; \n "
2017-07-05 02:08:06 -04:00
" real_bucket_t lsratio; \n "
2016-02-24 00:01:02 -05:00
" real4reals_bucket newHsv; \n "
2017-07-05 02:08:06 -04:00
" real_bucket_t maxa, maxc, newls; \n "
2016-02-24 00:01:02 -05:00
" real_bucket_t adjhlp; \n "
" \n "
" if (ls == 0 || (oldRgb->m_Real4.x == 0 && oldRgb->m_Real4.y == 0 && oldRgb->m_Real4.z == 0)) \n " //Can't do a vector compare to zero.
" { \n "
" newRgb->m_Real4 = 0; \n "
" return; \n "
" } \n "
" \n "
//Identify the most saturated channel.
" maxc = max(max(oldRgb->m_Reals[0], oldRgb->m_Reals[1]), oldRgb->m_Reals[2]); \n "
" maxa = ls * maxc; \n "
2017-07-05 02:08:06 -04:00
" newls = 1 / maxc; \n "
2016-02-24 00:01:02 -05:00
" \n "
//If a channel is saturated and highlight power is non-negative
//modify the color to prevent hue shift.
2017-07-05 02:08:06 -04:00
" if (maxa > 1 && highPow >= 0) \n "
2016-02-24 00:01:02 -05:00
" { \n "
" lsratio = pow(newls / ls, highPow); \n "
" \n "
//Calculate the max-value color (ranged 0 - 1).
" for (rgbi = 0; rgbi < 3; rgbi++) \n "
2017-07-05 02:08:06 -04:00
" newRgb->m_Reals[rgbi] = newls * oldRgb->m_Reals[rgbi]; \n "
2016-02-24 00:01:02 -05:00
" \n "
//Reduce saturation by the lsratio.
" RgbToHsv(&(newRgb->m_Real4), &(newHsv.m_Real4)); \n "
" newHsv.m_Real4.y *= lsratio; \n "
" HsvToRgb(&(newHsv.m_Real4), &(newRgb->m_Real4)); \n "
" } \n "
" else \n "
" { \n "
" adjhlp = -highPow; \n "
" \n "
" if (adjhlp > 1) \n "
" adjhlp = 1; \n "
" \n "
2017-07-05 02:08:06 -04:00
" if (maxa <= 1) \n "
2016-02-24 00:01:02 -05:00
" adjhlp = 1; \n "
" \n "
//Calculate the max-value color (ranged 0 - 1) interpolated with the old behavior.
" for (rgbi = 0; rgbi < 3; rgbi++) \n " //Unrolling, caching and vectorizing makes no difference.
" newRgb->m_Reals[rgbi] = ((1.0 - adjhlp) * newls + adjhlp * ls) * oldRgb->m_Reals[rgbi]; \n "
" } \n "
" } \n "
" \n " ;
2014-07-08 03:11:14 -04:00
return os . str ( ) ;
}
/// <summary>
/// Create the gamma correction kernel string used for early clipping.
/// </summary>
/// <returns>The gamma correction kernel string used for early clipping</returns>
2017-07-22 16:43:35 -04:00
string FinalAccumOpenCLKernelCreator : : CreateGammaCorrectionKernelString ( )
2014-07-08 03:11:14 -04:00
{
ostringstream os ;
string dataType ;
os < <
2016-02-24 00:01:02 -05:00
ConstantDefinesString ( m_DoublePrecision ) < <
UnionCLStructString < <
RgbToHsvFunctionString < <
HsvToRgbFunctionString < <
CalcAlphaFunctionString < <
CreateCalcNewRgbFunctionString ( true ) < <
SpatialFilterCLStructString < <
2017-07-22 16:43:35 -04:00
CreateGammaCorrectionFunctionString ( true , false ) ; //Will only be used with float in this case, early clip. Will always alpha accum.
os < < " __kernel void " < < m_GammaCorrectionWithoutAlphaCalcEntryPoint < < " ( \n " < <
2016-02-24 00:01:02 -05:00
" __global real4reals_bucket* accumulator, \n "
" __constant SpatialFilterCL* spatialFilter \n "
" ) \n "
" { \n "
" int testGutter = 0; \n "
" \n "
" if (GLOBAL_ID_Y >= (spatialFilter->m_SuperRasH - testGutter) || GLOBAL_ID_X >= (spatialFilter->m_SuperRasW - testGutter)) \n "
" return; \n "
" \n "
" uint superIndex = (GLOBAL_ID_Y * spatialFilter->m_SuperRasW) + GLOBAL_ID_X; \n "
" __global real4reals_bucket* bucket = accumulator + superIndex; \n "
2017-07-22 16:43:35 -04:00
" GammaCorrectionFloats(bucket, &(spatialFilter->m_Background[0]), spatialFilter->m_Gamma, spatialFilter->m_LinRange, spatialFilter->m_Vibrancy, spatialFilter->m_HighlightPower, &(bucket->m_Reals[0])); \n "
2016-02-24 00:01:02 -05:00
" } \n "
;
2014-07-08 03:11:14 -04:00
return os . str ( ) ;
}
2014-12-05 21:30:46 -05:00
}