[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [60296] branches/soc-2013-vse/source/ blender/sequencer/kernel: Updating OpenCL kernels for VSE effects.
Alexander Kuznetsov
kuzsasha at gmail.com
Sun Sep 22 03:37:57 CEST 2013
Revision: 60296
http://projects.blender.org/scm/viewvc.php?view=rev&root=bf-blender&revision=60296
Author: alexk
Date: 2013-09-22 01:37:55 +0000 (Sun, 22 Sep 2013)
Log Message:
-----------
Updating OpenCL kernels for VSE effects. Also adding stuctures and specific definition for CPU part in order to share data structures.
Modified Paths:
--------------
branches/soc-2013-vse/source/blender/sequencer/kernel/include_type_float.cl
branches/soc-2013-vse/source/blender/sequencer/kernel/kernel_working.cl
Added Paths:
-----------
branches/soc-2013-vse/source/blender/sequencer/kernel/common_stuctures.cl
branches/soc-2013-vse/source/blender/sequencer/kernel/common_stuctures.h
Added: branches/soc-2013-vse/source/blender/sequencer/kernel/common_stuctures.cl
===================================================================
--- branches/soc-2013-vse/source/blender/sequencer/kernel/common_stuctures.cl (rev 0)
+++ branches/soc-2013-vse/source/blender/sequencer/kernel/common_stuctures.cl 2013-09-22 01:37:55 UTC (rev 60296)
@@ -0,0 +1,42 @@
+#ifndef __COMMON_STRUCTURES_CL__
+#define __COMMON_STRUCTURES_CL__
+
+
+
+typedef struct bufsize
+{
+ unsigned int sizex;
+ unsigned int sizey;
+} bufsize;
+
+typedef struct bufpos
+{
+ unsigned int posx;
+ unsigned int posy;
+} bufpos;
+
+
+typedef struct imgbuf
+{
+bufpos p;
+bufsize s;
+} imgbuf;
+
+
+
+
+typedef struct matrix_affine
+{
+ float row1[3];
+ float row2[3];
+} matrix_affine;
+
+
+typedef struct color_balance_data
+{
+ float4 lift;
+ float4 gamma;
+ float4 gain;
+} color_balance_data;
+
+#endif /* __COMMON_STRUCTURES_CL__*/
Added: branches/soc-2013-vse/source/blender/sequencer/kernel/common_stuctures.h
===================================================================
--- branches/soc-2013-vse/source/blender/sequencer/kernel/common_stuctures.h (rev 0)
+++ branches/soc-2013-vse/source/blender/sequencer/kernel/common_stuctures.h 2013-09-22 01:37:55 UTC (rev 60296)
@@ -0,0 +1,9 @@
+#ifndef __COMMON_STRUCTURES_H__
+#define __COMMON_STRUCTURES_H__
+
+typedef struct float4 {float s[4];} float4;
+
+#include "common_stuctures.cl"
+
+
+#endif /*__COMMON_STRUCTURES_H__*/
Modified: branches/soc-2013-vse/source/blender/sequencer/kernel/include_type_float.cl
===================================================================
--- branches/soc-2013-vse/source/blender/sequencer/kernel/include_type_float.cl 2013-09-22 00:52:08 UTC (rev 60295)
+++ branches/soc-2013-vse/source/blender/sequencer/kernel/include_type_float.cl 2013-09-22 01:37:55 UTC (rev 60296)
@@ -6,3 +6,7 @@
typedef float8 vst8;
typedef float16 vst16;
+#define convert_vst4(a) convert_float4(a)
+#define convert_vst(a) convert_float(a)
+
+
Modified: branches/soc-2013-vse/source/blender/sequencer/kernel/kernel_working.cl
===================================================================
--- branches/soc-2013-vse/source/blender/sequencer/kernel/kernel_working.cl 2013-09-22 00:52:08 UTC (rev 60295)
+++ branches/soc-2013-vse/source/blender/sequencer/kernel/kernel_working.cl 2013-09-22 01:37:55 UTC (rev 60296)
@@ -1,5 +1,4 @@
-
#define GETPOINT int i = get_global_id(0); \
int j = get_global_id(1); \
size_t pos = (j*sizex+i)*4;\
@@ -18,8 +17,31 @@
out[pos+3] = a;
+#define GETPOINT4 \
+ int pos_i = get_global_id(0); \
+ int pos_j = get_global_id(1); \
+ vst4 dv;\
+ size_t pos;\
+ bool legal_pixel;\
+ if(pos_i < in1s.sizex && pos_j < in1s.sizey){\
+ pos = (pos_j*in1s.sizex+pos_i);\
+ dv = in[pos];\
+ legal_pixel = true;\
+ } else {\
+ dv = 0;\
+ legal_pixel = false;\
+ };
+
+
+
+
+#define STOREPOINT4 \
+ if(legal_pixel) out[pos] = dv;
+
+
+
__kernel void kBrightnessContrast
(__global vst *in, __global vst *out,
const unsigned int sizex,
@@ -38,16 +60,35 @@
}
+__kernel void kColorBalance
+(__global vst4 *in, __global vst4 *out,
+ const bufsize in1s,
+ const color_balance_data cb,
+ const float mul
+)
+{
+ GETPOINT4
+ dv = (((dv - 1.0f) * cb.lift) + 1.0f) * cb.gain;
+// dv = cb.gain*((float4)dv);
+ dv = max(dv, 0.0f);
+ dv = pow(dv, cb.gamma);
+ dv.s012 = dv.s012*mul;
+ STOREPOINT4
+
+}
+
+
+
+
__kernel void kConv4FromChar(
__global uchar4 *in,
__global vst4 *out,
const unsigned int lsize
-
)
{
size_t pos = get_global_id(0);
@@ -56,85 +97,112 @@
out[pos] = (1.0f/255)*convert_vst4(in[pos]);
}
+__kernel void kConvFromChar(
+ __global uchar *in,
+ __global vst *out,
+ const unsigned int lsize
+)
+{
+ size_t pos = get_global_id(0);
+ if(pos < lsize)
+ out[pos] = (1.0f/255)*convert_vst(in[pos]);
+}
+
+
+
+
+
+__kernel void kConv4FromFloat(
+ __global float4 *in,
+ __global vst4 *out,
+ const unsigned int lsize
+)
+{
+ size_t pos = get_global_id(0);
+
+ if(pos < lsize)
+ out[pos] = (1.0f/1.0f)*convert_vst4(in[pos]);
+}
+
+__kernel void kConvFromFloat(
+ __global float *in,
+ __global vst *out,
+ const unsigned int lsize
+)
+{
+ size_t pos = get_global_id(0);
+
+ if(pos < lsize)
+ out[pos] = (1.0f/1.0f)*convert_vst(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;\
+ vst4 dv;\
+ vst4 dv1;\
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;\
+ dv1 = 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];\
+ pos = (pos1[1]*in1i.s.sizex+pos1[0]);\
+ dv1 = in1[pos];\
}
#define GETPOINT2 \
int pos2[2] = {i-in2i.p.posx, j-in2i.p.posy};\
- vst dr2, dg2, db2, da2;\
+ vst4 dv2;\
if(pos2[0] < 0 || pos2[1] < 0 ||\
pos2[0] >= in2i.s.sizex || pos2[1] >= in2i.s.sizey)\
- dr2 = dg2 = db2 = da2 = 0;\
+ dv2 = 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];\
+ pos = (pos2[1]*in2i.s.sizex+pos2[0]);\
+ dv2 = in2[pos];\
}
#define STOREPOINT3 \
- pos = (j*sizex+i)*4;\
- out[pos] = dr;\
- out[pos+1] = dg;\
- out[pos+2] = db;\
- out[pos+3] = da;
+ if(i < outs.sizex && j < outs.sizey){\
+ pos = (j*outs.sizex+i);\
+ out[pos] = dv;\
+ }
__kernel void kAlphaOver(
- __global vst *in1,
+ __global vst4 *in1,
const imgbuf in1i,
- __global vst *in2,
+ __global vst4 *in2,
const imgbuf in2i,
- __global vst *out,
- const unsigned int sizex,
- const unsigned int sizey,
+ __global vst4 *out,
+ const bufsize outs,
const float fac
)
{
GETPOINT1
GETPOINT2
- float mfac = 1.0f - (fac*da1);
+ float mfac = 1.0f - (fac*dv1.s3);
if(fac <= 0.0f) {
- dr = dr2;
- db = db2;
- dg = dg2;
- da = da2;
+ dv = dv2;
}
else if(mfac <= 0.0f) {
- dr = dr1;
- dg = dg1;
- db = db1;
- da = da1;
+ dv = dv1;
} else
{
- dr = fac*dr1 + mfac*dr2;
- dg = fac*dg1 + mfac*dg2;
- db = fac*db1 + mfac*db2;
- da = fac*da1 + mfac*da2;
+ dv = fac*dv1 + mfac*dv2;
}
STOREPOINT3
@@ -144,16 +212,50 @@
+__kernel void kAlphaUnder(
+ __global vst4 *in1,
+ const imgbuf in1i,
+
+ __global vst4 *in2,
+ const imgbuf in2i,
+
+ __global vst4 *out,
+ const bufsize outs,
+ const float fac2
+)
+{
+ GETPOINT1
+ GETPOINT2
+
+ float mfac = dv2.s3;
+ float fac = fac2 * (1.0f - mfac);
+
+ if(fac <= 0.0f) {
+ dv = dv2;
+ }
+ else if(mfac <= 0.0f) {
+ dv = dv1;
+ } else
+ {
+ dv = fac*dv1 + mfac*dv2;
+ }
+
+ STOREPOINT3
+
+}
+
+
+
+
__kernel void kAdd(
- __global vst *in1,
+ __global vst4 *in1,
const imgbuf in1i,
- __global vst *in2,
+ __global vst4 *in2,
const imgbuf in2i,
- __global vst *out,
- const unsigned int sizex,
- const unsigned int sizey,
+ __global vst4 *out,
+ const bufsize outs,
const float fac
)
{
@@ -161,12 +263,10 @@
GETPOINT2
- float m = 1.0f - (da1 * (1.0f - fac));
+ float m = 1.0f - (dv1.s3 * (1.0f - fac));
- dr = dr1 + m * dr2;
- dg = dg1 + m * dg2;
- db = db1 + m * db2;
- da = da1;
+ dv.s012 = dv1.s012 + m * dv2.s012;
+ dv.s3 = dv1.s3;
STOREPOINT3
@@ -174,15 +274,14 @@
__kernel void kSubtract(
- __global vst *in1,
+ __global vst4 *in1,
const imgbuf in1i,
- __global vst *in2,
+ __global vst4 *in2,
const imgbuf in2i,
- __global vst *out,
- const unsigned int sizex,
- const unsigned int sizey,
+ __global vst4 *out,
+ const bufsize outs,
const float fac
)
{
@@ -190,43 +289,56 @@
GETPOINT2
- float m = 1.0f - (da1 * (1.0f - fac));
+ float m = 1.0f - (dv1.s3 * (1.0f - fac));
- dr = dr1 - m * dr2;
- dg = dg1 - m * dg2;
- db = db1 - m * db2;
- da = da1;
+ dv.s012 = dv1.s012 - m * dv2.s012;
+ dv.s3 = dv1.s3;
STOREPOINT3
+}
+
+__kernel void kMultiply(
+ __global vst4 *in1,
+ const imgbuf in1i,
+
+ __global vst4 *in2,
+ const imgbuf in2i,
+
+ __global vst4 *out,
+ const bufsize outs,
+ const float fac
+)
+{
+ GETPOINT1
+ GETPOINT2
+ dv = dv1 + fac*dv1*(dv2 - 1.0f);
+
+ STOREPOINT3
}
__kernel void kCross(
- __global vst *in1,
+ __global vst4 *in1,
const imgbuf in1i,
- __global vst *in2,
+ __global vst4 *in2,
const imgbuf in2i,
- __global vst *out,
- const unsigned int sizex,
- const unsigned int sizey,
+ __global vst4 *out,
+ const bufsize outs,
const float fac
)
{
GETPOINT1
GETPOINT2
-/*
- float fac1 = 1.0f - fac);
+ float lfac = 1.0f - fac;
+
- dr = dr1 * fac1 + dr2 ;
- dg = dg1 * fac1 + dg2;
- db = db1 * fac1 + db2;
- da = da1 * fac1 ;
- */
+ dv = lfac*dv1 + fac*dv2;
+
STOREPOINT3
}
@@ -397,7 +509,6 @@
-
int get_id_from_value(float value, float offset, float stepsize)
{
@@ -817,5 +928,4 @@
-}
-
+}
\ No newline at end of file
More information about the Bf-blender-cvs
mailing list