summaryrefslogtreecommitdiff
path: root/kernels
diff options
context:
space:
mode:
authorChuanbo Weng <chuanbo.weng@intel.com>2015-06-18 16:29:39 +0800
committerYang Rong <rong.r.yang@intel.com>2015-06-19 10:43:01 +0800
commit05d343de8b44d24e6319cfb54cabaab1484969e6 (patch)
tree91a6a0042f10f4d5e4f2d5077fe25ffd694c51f0 /kernels
parent45ead9765dd1d7fc4b1fdd51aeb3f0047caadcf6 (diff)
Add example to show v4l2 buffer sharing with extension clGetMemObjectFdIntel.
This example captures yuy2 frame directly to cl buffer object by the way of dma, processed by OpenCL kernel, then convert to nv12 format and shown by libva. v2: Close cl buffer's fd by clCloseMemObjectFdIntel instead of close function. v3: Just use close function, no need of clCloseMemObjectFdIntel. v4: Some modifcation of examples/CMakeLists.txt after code rebase. Signed-off-by: Chuanbo Weng <chuanbo.weng@intel.com> Reviewed-by: Yuan, Feng <feng.yuan@intel.com>
Diffstat (limited to 'kernels')
-rw-r--r--kernels/runtime_yuy2_processing.cl15
1 files changed, 15 insertions, 0 deletions
diff --git a/kernels/runtime_yuy2_processing.cl b/kernels/runtime_yuy2_processing.cl
new file mode 100644
index 00000000..1478e656
--- /dev/null
+++ b/kernels/runtime_yuy2_processing.cl
@@ -0,0 +1,15 @@
+__kernel void
+runtime_yuy2_processing(__global uchar *src,
+ int image_height,
+ int image_pitch)
+{
+ int gx = get_global_id(0);
+ int gy = get_global_id(1);
+
+ int src_y = image_height / 2 + gy;
+ int mirror_y = image_height - src_y;
+
+ uchar4 mirror_val = *(__global uchar4*)(src + mirror_y*image_pitch + gx*4);
+ *(__global uchar4*)(src + src_y*image_pitch + gx*4) = mirror_val;
+
+}