[gegl] Make box-blur.cl match the committed box-blur.cl.h
- From: Daniel Sabo <daniels src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] Make box-blur.cl match the committed box-blur.cl.h
- Date: Sun, 16 Feb 2014 01:37:02 +0000 (UTC)
commit 1ad4dacaa7c03634213fee4c4081d1d781aef5aa
Author: Daniel Sabo <DanielSabo gmail com>
Date: Sat Feb 15 17:04:27 2014 -0800
Make box-blur.cl match the committed box-blur.cl.h
opencl/box-blur.cl | 100 ++++++++++++++++++++++++++--------------------------
1 files changed, 50 insertions(+), 50 deletions(-)
---
diff --git a/opencl/box-blur.cl b/opencl/box-blur.cl
index 5011fa3..7ed2475 100644
--- a/opencl/box-blur.cl
+++ b/opencl/box-blur.cl
@@ -45,7 +45,7 @@ __kernel void kernel_blur_ver (__global const float4 *aux,
}
__kernel
-__attibute__((reqd_work_group_size(256,1,1)))
+__attribute__((reqd_work_group_size(256,1,1)))
void kernel_box_blur_fast(const __global float4 *in,
__global float4 *out,
__local float4 *column_sum,
@@ -54,60 +54,60 @@ void kernel_box_blur_fast(const __global float4 *in,
const int radius,
const int size)
{
- const int local_id0 = get_local_id(0);
- const int twice_radius = 2 * radius;
- const int in_width = twice_radius + width;
- const int in_height = twice_radius + height;
- const float4 area = (float4)( (twice_radius+1) * (twice_radius+1) );
- int column_index_start,column_index_end;
- int y = get_global_id(1) * size;
- const int out_x = get_group_id(0)
+ const int local_id0 = get_local_id(0);
+ const int twice_radius = 2 * radius;
+ const int in_width = twice_radius + width;
+ const int in_height = twice_radius + height;
+ const float4 area = (float4)( (twice_radius+1) * (twice_radius+1) );
+ int column_index_start,column_index_end;
+ int y = get_global_id(1) * size;
+ const int out_x = get_group_id(0)
* ( get_local_size(0) - twice_radius ) + local_id0 - radius;
- const int in_x = out_x + radius;
- int tmp_size = size;
- int tmp_index = 0;
- float4 tmp_sum = (float4)0.0f;
- float4 total_sum = (float4)0.0f;
- if( in_x < in_width )
- {
- column_index_start = y;
- column_index_end = y + twice_radius;
- for( int i=0; i<twice_radius+1; ++i )
- tmp_sum+=in[(y+i)*in_width+in_x];
- column_sum[local_id0] = tmp_sum;
- }
+ const int in_x = out_x + radius;
+ int tmp_size = size;
+ int tmp_index = 0;
+ float4 tmp_sum = (float4)0.0f;
+ float4 total_sum = (float4)0.0f;
+ if( in_x < in_width )
+ {
+ column_index_start = y;
+ column_index_end = y + twice_radius;
+ for( int i=0; i<twice_radius+1; ++i )
+ tmp_sum+=in[(y+i)*in_width+in_x];
+ column_sum[local_id0] = tmp_sum;
+ }
- barrier( CLK_LOCAL_MEM_FENCE );
+ barrier( CLK_LOCAL_MEM_FENCE );
- while(1)
- {
- if( out_x < width )
- {
- if( local_id0 >= radius
+ while(1)
+ {
+ if( out_x < width )
+ {
+ if( local_id0 >= radius
&& local_id0 < get_local_size(0) - radius )
- {
- total_sum = (float4)0.0f;
- for( int i=0; i<twice_radius+1; ++i )
- total_sum += column_sum[local_id0-radius+i];
- out[y*width+out_x] = total_sum/area;
- }
- }
- if( --tmp_size ==0 || y == height - 1 )
- break;
+ {
+ total_sum = (float4)0.0f;
+ for( int i=0; i<twice_radius+1; ++i )
+ total_sum += column_sum[local_id0-radius+i];
+ out[y*width+out_x] = total_sum/area;
+ }
+ }
+ if( --tmp_size ==0 || y == height - 1 )
+ break;
- barrier( CLK_LOCAL_MEM_FENCE );
+ barrier( CLK_LOCAL_MEM_FENCE );
- ++y;
- if( in_x < in_width )
- {
- tmp_sum = column_sum[local_id0];
- tmp_sum -= in[(column_index_start)*in_width+in_x];
- tmp_sum += in[(column_index_end+1)*in_width+in_x];
- ++column_index_start;
- ++column_index_end;
- column_sum[local_id0] = tmp_sum;
- }
+ ++y;
+ if( in_x < in_width )
+ {
+ tmp_sum = column_sum[local_id0];
+ tmp_sum -= in[(column_index_start)*in_width+in_x];
+ tmp_sum += in[(column_index_end+1)*in_width+in_x];
+ ++column_index_start;
+ ++column_index_end;
+ column_sum[local_id0] = tmp_sum;
+ }
- barrier( CLK_LOCAL_MEM_FENCE );
- }
+ barrier( CLK_LOCAL_MEM_FENCE );
+ }
}
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]