[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [59342] branches/soc-2013-vse/source/ blender/sequencer/kernel/kernel_working.cl: Adding kernels for histograms, waveform, vectorscope.
Alexander Kuznetsov
kuzsasha at gmail.com
Wed Aug 21 05:39:21 CEST 2013
Revision: 59342
http://projects.blender.org/scm/viewvc.php?view=rev&root=bf-blender&revision=59342
Author: alexk
Date: 2013-08-21 03:39:20 +0000 (Wed, 21 Aug 2013)
Log Message:
-----------
Adding kernels for histograms, waveform, vectorscope. They are seperated into analysis and rendering parts. The rendering part might change to write directly to a texture (or make a GLSL shader do it).
Modified Paths:
--------------
branches/soc-2013-vse/source/blender/sequencer/kernel/kernel_working.cl
Modified: branches/soc-2013-vse/source/blender/sequencer/kernel/kernel_working.cl
===================================================================
--- branches/soc-2013-vse/source/blender/sequencer/kernel/kernel_working.cl 2013-08-21 02:30:19 UTC (rev 59341)
+++ branches/soc-2013-vse/source/blender/sequencer/kernel/kernel_working.cl 2013-08-21 03:39:20 UTC (rev 59342)
@@ -393,4 +393,429 @@
}
}
+
+
+
+
+
+int get_id_from_value(float value, float offset, float stepsize)
+{
+
+ return (int)((value-offset)/stepsize);
+
+}
+
+
+
+
+__kernel void kHistogram(
+ __global vst *in,
+ const bufsize insize,
+
+ const float zeropoint,
+ const float stepsize,
+
+ const unsigned int histsize,
+ __global unsigned int *globhist
+ )
+{
+
+ int pos_i = get_global_id(0);
+ int pos_j = get_global_id(1);
+
+
+ float dr = 0;
+ float da = 0;
+
+ if(pos_i < insize.sizex &&
+ pos_j < insize.sizey)
+ {
+
+ int histogram_p = 0;
+ dr = in[4*(pos_j*insize.sizex + pos_i)];
+ da = in[4*(pos_j*insize.sizex + pos_i)+3];
+ histogram_p = get_id_from_value(dr, zeropoint, stepsize);
+
+ if(histogram_p >= 0 && histogram_p < histsize)
+ {
+ atomic_add(&globhist[histogram_p], (unsigned int)255.0f*da);
+
+ }
+
+
+ }
+
+
+}
+
+
+__kernel void kHistogramRGB(
+ __global vst4 *in,
+ const bufsize insize,
+
+ const float zeropoint,
+ const float stepsize,
+
+ const unsigned int histsize,
+ __global unsigned int *globhistr,
+ __global unsigned int *globhistg,
+ __global unsigned int *globhistb
+ )
+{
+
+ int pos_i = get_global_id(0);
+ int pos_j = get_global_id(1);
+
+
+ vst4 dv = 0;
+
+ if(pos_i < insize.sizex &&
+ pos_j < insize.sizey)
+ {
+
+ int histogram_r = 0;
+ int histogram_g = 0;
+ int histogram_b = 0;
+
+ dv = in[(pos_j*insize.sizex + pos_i)];
+
+
+ histogram_r = get_id_from_value(dv.s0, zeropoint, stepsize);
+ histogram_g = get_id_from_value(dv.s1, zeropoint, stepsize);
+ histogram_b = get_id_from_value(dv.s2, zeropoint, stepsize);
+
+ if(histogram_r >= 0 && histogram_r < histsize)
+ {
+ atomic_add(&globhistr[histogram_r], (unsigned int)255.0f*dv.s3);
+ }
+
+ if(histogram_g >= 0 && histogram_g < histsize)
+ {
+ atomic_add(&globhistg[histogram_g], (unsigned int)255.0f*dv.s3);
+ }
+
+ if(histogram_b >= 0 && histogram_b < histsize)
+ {
+ atomic_add(&globhistb[histogram_b], (unsigned int)255.0f*dv.s3);
+ }
+
+
+ }
+
+
+}
+
+
+
+
+
+
+
+
+
+__kernel void kMaxArrayUINT(
+ __global unsigned int *indata,
+ __global unsigned int *outdata,
+ const unsigned int size
+)
+{
+ int pos_i = get_global_id(0);
+
+ unsigned int ta = pos_i*2 > size ? 0 : indata[pos_i*2];
+ unsigned int tb = pos_i*2+1 > size ? 0 : indata[pos_i*2+1];
+
+ outdata[pos_i] = ta > tb ? ta : tb;
+
+}
+
+
+__kernel void kSetZero(
+ __global unsigned int *outdata,
+ const unsigned int size
+)
+{
+ int pos_i = get_global_id(0);
+ if(pos_i < size)
+ outdata[pos_i] = 0;
+}
+
+
+
+
+/* warning, get_global_size must be == sizex and group size must have width of 1 */
+__kernel void kHistogramRender(
+ __global unsigned int *globhist,
+ const unsigned int histsize,
+
+ __global vst4 *out,
+ const bufsize outsize,
+
+ __global unsigned int *maxval
+)
+{
+ int pos_i = get_global_id(0);
+ int pos_j = get_global_id(1);
+
+ __local unsigned int cheight1;
+
+ if(get_local_id(0) == 0)
+ {
+ cheight1 = outsize.sizey*((float)globhist[pos_i]/maxval[0]);
+ }
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ if(pos_j < outsize.sizey)
+ {
+ if(pos_j < cheight1)
+ out[pos_j* outsize.sizex + pos_i] = 1;
+ else
+ out[pos_j* outsize.sizex + pos_i] = 0;
+
+ }
+
+
+
+}
+
+
+
+/* warning, get_global_size must be == sizex and group size must have width of 1 */
+__kernel void kHistogramRenderRGB(
+ __global unsigned int *globhistr,
+ __global unsigned int *globhistg,
+ __global unsigned int *globhistb,
+ const unsigned int histsize,
+
+ __global vst4 *out,
+ const bufsize outsize,
+
+ __global unsigned int *maxvalr,
+ __global unsigned int *maxvalg,
+ __global unsigned int *maxvalb
+)
+{
+ int pos_i = get_global_id(0);
+ int pos_j = get_global_id(1);
+
+ __local unsigned int cheight1;
+ __local unsigned int cheight2;
+ __local unsigned int cheight3;
+
+ if(get_local_id(0) == 0)
+ {
+ cheight1 = outsize.sizey*((float)globhistr[pos_i]/maxvalr[0]);
+ cheight2 = outsize.sizey*((float)globhistg[pos_i]/maxvalg[0]);
+ cheight3 = outsize.sizey*((float)globhistb[pos_i]/maxvalb[0]);
+ }
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ if(pos_j < outsize.sizey)
+ {
+ vst4 val = 0;
+
+ if(pos_j < cheight1)
+ val.s03 = 1;
+ if(pos_j < cheight2)
+ val.s13 = 1;
+ if(pos_j < cheight3)
+ val.s23 = 1;
+
+ out[pos_j* outsize.sizex + pos_i] = val;
+
+ }
+
+
+
+}
+
+float3 rgb_to_yuv(float3 rgb)
+{
+float3 yuv;
+ yuv.s0 = dot(rgb, (float3)(0.299f, 0.587f, 0.114f));
+ yuv.s1 = dot(rgb, (float3)(-0.147f, -0.289f, 0.436f));
+ yuv.s2 = dot(rgb, (float3)(0.615f, -0.515f, -0.100f));
+
+return yuv;
+}
+
+
+
+
+__kernel void kWaveform(
+ __global vst4 *in,
+ const bufsize insize,
+
+ const float zeropoint,
+ const float stepsize,
+
+ const unsigned int wavesize,
+
+ __global unsigned int *globalwave,
+ __local unsigned int *localwave
+ )
+{
+
+ int pos_i = get_global_id(0);
+ int pos_j = get_global_id(1);
+
+ size_t loc_i = get_local_id(0);
+ size_t loc_j = get_local_id(1);
+
+ vst4 dv = 0;
+
+ size_t local_size_y = get_global_size(1);
+
+ size_t copy_pos;
+ for(copy_pos = loc_j; copy_pos < wavesize; copy_pos+=local_size_y)
+ {
+ localwave[copy_pos] = 0;
+ }
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ if(pos_i < insize.sizex &&
+ pos_j < insize.sizey)
+ {
+
+ int wave_p = 0;
+ dv = in[pos_j*insize.sizex + pos_i];
+
+
+ wave_p = get_id_from_value(dv.s0, zeropoint, stepsize);
+
+ if(wave_p >= 0 && wave_p < wavesize)
+ {
+ atomic_add(&localwave[wave_p], (unsigned int)255.0f*dv.s3);
+ }
+
+
+ }
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+
+ for(copy_pos = loc_j; copy_pos < wavesize; copy_pos+=local_size_y)
+ {
+ globalwave[wavesize*(get_group_id(0)) + copy_pos] = localwave[copy_pos];
+ }
+
+
+}
+
+
+
+
+
+/* warning, get_global_size must be == sizex and group size must have width of 1 */
+__kernel void kWaveformRender(
+ __global unsigned int *globwave,
+ const unsigned int wavesize,
+
+ __global vst4 *out,
+ const bufsize outsize,
+
+ __global unsigned int *maxval
+)
+{
+ int pos_i = get_global_id(0);
+ int pos_j = get_global_id(1);
+
+
+ if(pos_j < outsize.sizey &&
+ pos_i < outsize.sizex
+ )
+ {
+ vst4 val = (float4)(0.0f, 0.0f, 0.0f, 1.0f);
+
+ val.s012 = ((float)globwave[pos_j+pos_i*wavesize]/maxval[0]);
+
+
+ out[pos_j* outsize.sizex + pos_i] = val;
+
+ }
+
+
+
+}
+
+
+
+__kernel void kVectorscope(
+ __global vst4 *in,
+ const bufsize insize,
+
+ const float zeropoint,
+ const float stepsize,
+ const unsigned int vectorsize,
+
+ __global unsigned int *globvector
+ )
+{
+
+ int pos_i = get_global_id(0);
+ int pos_j = get_global_id(1);
+
+
+ vst4 dv = 0;
+
+ if(pos_i < insize.sizex &&
+ pos_j < insize.sizey)
+ {
+
+ int vectorpos_x, vectorpos_y;
+
+ float3 df;
+ dv = in[pos_j*insize.sizex + pos_i];
+ df = dv.s012;
+ float3 yuv;
+ yuv = rgb_to_yuv(df);
+
+
+ vectorpos_x = get_id_from_value(yuv.s1, zeropoint, stepsize);
+ vectorpos_y = get_id_from_value(yuv.s2, zeropoint, stepsize);
+
+ if(vectorpos_x >= 0 && vectorpos_x < vectorsize &&
+ vectorpos_y >= 0 && vectorpos_y < vectorsize)
+ {
+ atomic_add(&globvector[vectorpos_y*vectorsize+vectorpos_x], (unsigned int)255.0f*dv.s3);
+ }
+
+
+ }
+
+
+}
+
+
+
+
+
+__kernel void kVectorscopeRender(
+ __global unsigned int *globvector,
+
+ __global vst4 *out,
+ const bufsize outsize,
+
+ __global unsigned int *maxval
+)
+{
+ int pos_i = get_global_id(0);
+ int pos_j = get_global_id(1);
+
+
+ if(pos_j < outsize.sizey &&
+ pos_i < outsize.sizex)
+ {
+ vst4 val = 1;
+ /* increase brightness to 4 */
+ val.s3 = 4.0f*((float)globvector[pos_j*outsize.sizey+pos_i]/maxval[0]);//
+
+ out[pos_j* outsize.sizex + pos_i] = val;
+ }
+
+
+
+}
More information about the Bf-blender-cvs
mailing list