summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--kernels/compiler_box_blur_image.cl8
-rw-r--r--utests/compiler_box_blur_image.cpp2
2 files changed, 6 insertions, 4 deletions
diff --git a/kernels/compiler_box_blur_image.cl b/kernels/compiler_box_blur_image.cl
index 42f463b2..69fa2293 100644
--- a/kernels/compiler_box_blur_image.cl
+++ b/kernels/compiler_box_blur_image.cl
@@ -6,13 +6,15 @@ __kernel void compiler_box_blur_image(__read_only image2d_t src,
CLK_FILTER_NEAREST;
const int2 coord = (int2)(get_global_id(0), get_global_id(1));
int2 offset;
- float4 sum = 0;
+ uint4 sum = 0;
for (offset.y = -1; offset.y <= 1; offset.y++) {
for (offset.x = -1; offset.x <= 1; offset.x++) {
- sum += read_imagef(src, sampler, coord + offset);
+ sum += read_imageui(src, sampler, coord + offset);
}
}
- write_imagef(dst, coord, (1.0f/9.0f)*sum);
+ uint4 result = sum / 9;
+
+ write_imageui(dst, coord, result);
}
diff --git a/utests/compiler_box_blur_image.cpp b/utests/compiler_box_blur_image.cpp
index d94a97c9..711f2b38 100644
--- a/utests/compiler_box_blur_image.cpp
+++ b/utests/compiler_box_blur_image.cpp
@@ -15,7 +15,7 @@ static void compiler_box_blur_image()
src = cl_read_bmp("lenna128x128.bmp", &w, &h);
format.image_channel_order = CL_RGBA;
- format.image_channel_data_type = CL_UNORM_INT8;
+ format.image_channel_data_type = CL_UNSIGNED_INT8;
desc.image_type = CL_MEM_OBJECT_IMAGE2D;
desc.image_width = w;
desc.image_height = h;