[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [59169] branches/soc-2013-vse/source/ blender/sequencer/kernel/kernel_working.cl: Adding VSE image kernels.
Alexander Kuznetsov
kuzsasha at gmail.com
Fri Aug 16 00:29:56 CEST 2013
Revision: 59169
http://projects.blender.org/scm/viewvc.php?view=rev&root=bf-blender&revision=59169
Author: alexk
Date: 2013-08-15 22:29:56 +0000 (Thu, 15 Aug 2013)
Log Message:
-----------
Adding VSE image kernels. Those are already converted to support the position format to blending of differently size images. Therefore, we use GETPOINT and STOREPOINT in order to have abstruction. The rest will be fixed up when there won't be major changes
Added Paths:
-----------
branches/soc-2013-vse/source/blender/sequencer/kernel/kernel_working.cl
Added: branches/soc-2013-vse/source/blender/sequencer/kernel/kernel_working.cl
===================================================================
--- branches/soc-2013-vse/source/blender/sequencer/kernel/kernel_working.cl (rev 0)
+++ branches/soc-2013-vse/source/blender/sequencer/kernel/kernel_working.cl 2013-08-15 22:29:56 UTC (rev 59169)
@@ -0,0 +1,396 @@
+
+
+#define GETPOINT int i = get_global_id(0); \
+ int j = get_global_id(1); \
+ size_t pos = (j*sizex+i)*4;\
+ vst r = in[pos];\
+ vst g = in[pos+1];\
+ vst b = in[pos+2];\
+ vst a = in[pos+3];
+
+
+
+
+#define STOREPOINT \
+ out[pos] = r;\
+ out[pos+1] = g;\
+ out[pos+2] = b;\
+ out[pos+3] = a;
+
+
+
+
+__kernel void kBrightnessContrast
+(__global vst *in, __global vst *out,
+ const unsigned int sizex,
+ const unsigned int sizey,
+ const float multiply,
+ const float add
+)
+{
+ GETPOINT
+
+ r = r*multiply + add;
+ g = g*multiply + add;
+ b = b*multiply + add;
+
+ STOREPOINT
+
+}
+
+
+
+
+
+
+__kernel void kConv4FromChar(
+ __global uchar4 *in,
+ __global vst4 *out,
+ const unsigned int lsize
+
+)
+{
+ size_t pos = get_global_id(0);
+
+ if(pos < lsize)
+ out[pos] = (1.0f/255)*convert_vst4(in[pos]);
+}
+
+
+#define GETPOINT1 \
+ int i = get_global_id(0); \
+ int j = get_global_id(1); \
+ vst dr, dg, db, da;\
+ vst dr1, dg1, db1, da1;\
+ size_t pos;\
+ int pos1[2] = {i-in1i.p.posx, j-in1i.p.posy};\
+ if(pos1[0] < 0 || pos1[1] < 0 ||\
+ pos1[0] >= in1i.s.sizex || pos1[1] >= in1i.s.sizey)\
+ dr1 = dg1 = db1 = da1 = 0;\
+ else {\
+ pos = (pos1[1]*in1i.s.sizex+pos1[0])*4;\
+ dr1 = in1[pos];\
+ dg1 = in1[pos+1];\
+ db1 = in1[pos+2];\
+ da1 = in1[pos+3];\
+ }
+
+#define GETPOINT2 \
+ int pos2[2] = {i-in2i.p.posx, j-in2i.p.posy};\
+ vst dr2, dg2, db2, da2;\
+ if(pos2[0] < 0 || pos2[1] < 0 ||\
+ pos2[0] >= in2i.s.sizex || pos2[1] >= in2i.s.sizey)\
+ dr2 = dg2 = db2 = da2 = 0;\
+ else {\
+ pos = (pos2[1]*in2i.s.sizex+pos2[0])*4;\
+ dr2 = in2[pos];\
+ dg2 = in2[pos+1];\
+ db2 = in2[pos+2];\
+ da2 = in2[pos+3];\
+ }
+
+#define STOREPOINT3 \
+ pos = (j*sizex+i)*4;\
+ out[pos] = dr;\
+ out[pos+1] = dg;\
+ out[pos+2] = db;\
+ out[pos+3] = da;
+
+
+
+
+
+__kernel void kAlphaOver(
+ __global vst *in1,
+ const imgbuf in1i,
+
+ __global vst *in2,
+ const imgbuf in2i,
+
+ __global vst *out,
+ const unsigned int sizex,
+ const unsigned int sizey,
+ const float fac
+)
+{
+ GETPOINT1
+ GETPOINT2
+
+ float mfac = 1.0f - (fac*da1);
+
+ if(fac <= 0.0f) {
+ dr = dr2;
+ db = db2;
+ dg = dg2;
+ da = da2;
+ }
+ else if(mfac <= 0.0f) {
+ dr = dr1;
+ dg = dg1;
+ db = db1;
+ da = da1;
+ } else
+ {
+ dr = fac*dr1 + mfac*dr2;
+ dg = fac*dg1 + mfac*dg2;
+ db = fac*db1 + mfac*db2;
+ da = fac*da1 + mfac*da2;
+ }
+
+ STOREPOINT3
+
+}
+
+
+
+
+__kernel void kAdd(
+ __global vst *in1,
+ const imgbuf in1i,
+
+ __global vst *in2,
+ const imgbuf in2i,
+
+ __global vst *out,
+ const unsigned int sizex,
+ const unsigned int sizey,
+ const float fac
+)
+{
+ GETPOINT1
+ GETPOINT2
+
+
+ float m = 1.0f - (da1 * (1.0f - fac));
+
+ dr = dr1 + m * dr2;
+ dg = dg1 + m * dg2;
+ db = db1 + m * db2;
+ da = da1;
+
+ STOREPOINT3
+
+}
+
+
+__kernel void kSubtract(
+ __global vst *in1,
+ const imgbuf in1i,
+
+ __global vst *in2,
+ const imgbuf in2i,
+
+ __global vst *out,
+ const unsigned int sizex,
+ const unsigned int sizey,
+ const float fac
+)
+{
+ GETPOINT1
+ GETPOINT2
+
+
+ float m = 1.0f - (da1 * (1.0f - fac));
+
+ dr = dr1 - m * dr2;
+ dg = dg1 - m * dg2;
+ db = db1 - m * db2;
+ da = da1;
+
+ STOREPOINT3
+
+}
+
+
+
+__kernel void kCross(
+ __global vst *in1,
+ const imgbuf in1i,
+
+ __global vst *in2,
+ const imgbuf in2i,
+
+ __global vst *out,
+ const unsigned int sizex,
+ const unsigned int sizey,
+ const float fac
+)
+{
+ GETPOINT1
+ GETPOINT2
+
+/*
+ float fac1 = 1.0f - fac);
+
+ dr = dr1 * fac1 + dr2 ;
+ dg = dg1 * fac1 + dg2;
+ db = db1 * fac1 + db2;
+ da = da1 * fac1 ;
+ */
+ STOREPOINT3
+
+}
+
+
+
+#define GETCOLOR(in, x, y)\
+ ((0 <= x && 0 <= y && \
+ x < in ## size.sizex &&\
+y < in ## size.sizey)\
+? in[y*in ## size.sizex + x] : 0)
+
+
+__kernel void kSampleNearestNeighbor(
+ __global vst4 *in,
+ const bufsize insize,
+
+ __global vst4 *out,
+ const bufsize outsize,
+
+ const matrix_affine tran
+)
+{
+ size_t pos_i = get_global_id(0);
+ size_t pos_j = get_global_id(1);
+
+ float fpos_i = tran.row1[0]*pos_i + tran.row1[1]*pos_j + tran.row1[2];
+ float fpos_j = tran.row2[0]*pos_i + tran.row2[1]*pos_j + tran.row2[2];
+
+ int in_pos_i = floor(fpos_i);
+ int in_pos_j = floor(fpos_j);
+
+
+
+ if(pos_i < outsize.sizex &&
+ pos_j < outsize.sizey)
+ {
+ out[pos_i + pos_j*outsize.sizex] = GETCOLOR(in, in_pos_i, in_pos_j);
+ }
+
+
+
+
+}
+
+
+
+
+
+
+__kernel void kSampleBilinear(
+ __global vst4 *in,
+ const bufsize insize,
+
+ __global vst4 *out,
+ const bufsize outsize,
+
+ const matrix_affine tran
+
+)
+{
+ int pos_i = get_global_id(0);
+ int pos_j = get_global_id(1);
+
+ float fpos_i = tran.row1[0]*pos_i + tran.row1[1]*pos_j + tran.row1[2];
+ float fpos_j = tran.row2[0]*pos_i + tran.row2[1]*pos_j + tran.row2[2];
+
+ int in_pos_i_f = floor(fpos_i);
+ int in_pos_i_c = in_pos_i_f+1;
+ int in_pos_j_f = floor(fpos_j);
+ int in_pos_j_c = in_pos_j_f+1;
+
+ vst4 row1 = GETCOLOR(in, in_pos_i_f, in_pos_j_f);
+ vst4 row2 = GETCOLOR(in, in_pos_i_f, in_pos_j_c);
+ vst4 row3 = GETCOLOR(in, in_pos_i_c, in_pos_j_f);
+ vst4 row4 = GETCOLOR(in, in_pos_i_c, in_pos_j_c);
+
+ float a = fpos_i - in_pos_i_f;
+ float b = fpos_j - in_pos_j_f;
+
+ float a_b, ma_b, a_mb, ma_mb;
+ a_b = a * b; ma_b = (1.0f - a) * b; a_mb = a * (1.0f - b); ma_mb = (1.0f - a) * (1.0f - b);
+
+ if(pos_i < outsize.sizex &&
+ pos_j < outsize.sizey)
+ {
+ out[pos_i + pos_j*outsize.sizex] = ma_mb * row1 + a_mb * row3 + ma_b * row2 + a_b * row4;
+ }
+}
+
+
+
+
+float hf_BicubicR(float x)
+{
+ float p1, p2, p3, p4;
+ p1 = fmax(x + 2.0f, 0.0f);
+ p2 = fmax(x + 1.0f, 0.0f);
+ p3 = fmax(x, 0.0f);
+ p4 = fmax(x - 1.0f, 0.0f);
+
+
+ return (1.0f/6.0f)*(p1*p1*p1 - 4.0f*(p2*p2*p2+p4*p4*p4) + 6.0f*(p3*p3*p3)) ;
+}
+
+
+
+
+__kernel void kSampleBicubic(
+ __global vst4 *in,
+ const bufsize insize,
+
+ __global vst4 *out,
+ const bufsize outsize,
+
+ const matrix_affine tran
+ )
+{
+ int pos_i = get_global_id(0);
+ int pos_j = get_global_id(1);
+
+ float fpos_i = tran.row1[0]*pos_i + tran.row1[1]*pos_j + tran.row1[2];
+ float fpos_j = tran.row2[0]*pos_i + tran.row2[1]*pos_j + tran.row2[2];
+
+ int in_pos_i_f = floor(fpos_i);
+ int in_pos_j_f = floor(fpos_j);
+
+ float a = fpos_i - in_pos_i_f;
+ float b = fpos_j - in_pos_j_f;
+
+ float rdy[4] = {hf_BicubicR(b - (-1)),
+ hf_BicubicR(b - (0)),
+ hf_BicubicR(b - (1)),
+ hf_BicubicR(b - (2))
+ };
+
+ vst4 dout = 0;
+
+ int n, m;
+ for(n = -1; n <= 2; n++){
+ int x1 = in_pos_i_f + n ;
+ //x1 = clamp(x1, 0, (int)insize.sizex-1);
+ float rdx = hf_BicubicR(n-a);
+ for(m = -1; m <= 2; m++){
+ vst4 data;
+
+ int y1 = in_pos_j_f + m;
+ //y1 = clamp(y1, 0, (int)insize.sizey-1);
+ float rval = rdx * rdy[m+1];
+
+ dout += rval*GETCOLOR(in, x1, y1);
+
+
+ }
+
+ }
+
+
+
+ if(pos_i < outsize.sizex &&
+ pos_j < outsize.sizey)
+ {
+ out[pos_i + pos_j*outsize.sizex] = dout;
+ }
+}
+
+
More information about the Bf-blender-cvs
mailing list