diff options
author | Zoltan Gilian <zoltan.gilian@gmail.com> | 2015-08-13 18:28:24 +0200 |
---|---|---|
committer | Tom Stellard <thomas.stellard@amd.com> | 2016-03-03 20:15:58 +0000 |
commit | 13236d29b47859a2e46c1839045f20f1bddc137a (patch) | |
tree | a28eddefdf36150db83b85cc62ffa2526207f72f /tests/cl/program | |
parent | 1274800d6418266c28e744712a0714b1ffb44560 (diff) |
cl: add image reading tests
Reviewed-by: Tom Stellard <thomas.stellard@amd.com>
Diffstat (limited to 'tests/cl/program')
-rw-r--r-- | tests/cl/program/execute/image-read-2d.cl | 86 | ||||
-rw-r--r-- | tests/cl/program/execute/sampler.cl | 65 |
2 files changed, 151 insertions, 0 deletions
diff --git a/tests/cl/program/execute/image-read-2d.cl b/tests/cl/program/execute/image-read-2d.cl new file mode 100644 index 000000000..3d8128854 --- /dev/null +++ b/tests/cl/program/execute/image-read-2d.cl @@ -0,0 +1,86 @@ +/*! +[config] +name: 2D image reading + +dimensions: 1 +global_size: 1 0 0 +local_size: 1 0 0 + +[test] +name: read float from CL_FLOAT CL_RGBA image. +kernel_name: readf + +arg_in: 1 image float4 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 \ + 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 \ + 0.0 0.0 0.0 0.0 0.1 0.2 0.3 0.4 \ + image_type 2d \ + image_width 2 \ + image_height 3 \ + image_channel_order RGBA \ + image_channel_data_type FLOAT +arg_in: 2 sampler normalized_coords 0 \ + addressing_mode NONE \ + filter_mode NEAREST +arg_in: 3 int2 1 2 +arg_in: 4 float2 1.5 2.5 +arg_out: 0 buffer float4[2] 0.1 0.2 0.3 0.4 \ + 0.1 0.2 0.3 0.4 + +[test] +name: read signed integer from CL_SIGNED_INT8 CL_RGBA image. +kernel_name: readi + +arg_in: 1 image char4 0 0 0 0 0 0 0 0 \ + 0 0 0 0 0 0 0 0 \ + 0 0 0 0 -128 -1 0 127 \ + image_type 2d \ + image_width 2 \ + image_height 3 \ + image_channel_order RGBA \ + image_channel_data_type SIGNED_INT8 +arg_in: 2 sampler normalized_coords 0 \ + addressing_mode NONE \ + filter_mode NEAREST +arg_in: 3 int2 1 2 +arg_in: 4 float2 1.5 2.5 +arg_out: 0 buffer int4[2] -128 -1 0 127 \ + -128 -1 0 127 + +[test] +name: read unsigned integer from CL_UNSIGNED_INT8 CL_RGBA image. +kernel_name: readui + +arg_in: 1 image uchar4 0 0 0 0 0 0 0 0 \ + 0 0 0 0 0 0 0 0 \ + 0 0 0 0 0 1 127 255 \ + image_type 2d \ + image_width 2 \ + image_height 3 \ + image_channel_order RGBA \ + image_channel_data_type UNSIGNED_INT8 +arg_in: 2 sampler normalized_coords 0 \ + addressing_mode NONE \ + filter_mode NEAREST +arg_in: 3 int2 1 2 +arg_in: 4 float2 1.5 2.5 +arg_out: 0 buffer uint4[2] 0 1 127 255 \ + 0 1 127 255 +!*/ + +kernel void readf(global float4 *out, read_only image2d_t img, sampler_t s, + int2 coords_i, float2 coords_f) { + out[0] = read_imagef(img, s, coords_i); + out[1] = read_imagef(img, s, coords_f); +} + +kernel void readi(global int4 *out, read_only image2d_t img, sampler_t s, + int2 coords_i, float2 coords_f) { + out[0] = read_imagei(img, s, coords_i); + out[1] = read_imagei(img, s, coords_f); +} + +kernel void readui(global uint4 *out, read_only image2d_t img, sampler_t s, + int2 coords_i, float2 coords_f) { + out[0] = read_imageui(img, s, coords_i); + out[1] = read_imageui(img, s, coords_f); +} diff --git a/tests/cl/program/execute/sampler.cl b/tests/cl/program/execute/sampler.cl new file mode 100644 index 000000000..45960b778 --- /dev/null +++ b/tests/cl/program/execute/sampler.cl @@ -0,0 +1,65 @@ +/*! +[config] +name: Sampler tests + +dimensions: 1 +global_size: 4 0 0 +local_size: 1 0 0 + +[test] +name: read from image using linear filtering and unnormalized coords +kernel_name: readf_f_x + +arg_in: 1 image float 0.0 1.0 \ + 2.0 3.0 \ + image_type 2d \ + image_width 2 \ + image_height 2 \ + image_channel_order INTENSITY \ + image_channel_data_type FLOAT +arg_in: 2 sampler normalized_coords 0 \ + addressing_mode NONE \ + filter_mode LINEAR +arg_in: 3 buffer float2[4] 0.5 0.5 1.0 0.5 1.0 1.0 0.5 1.0 +arg_out: 0 buffer float[4] 0.0 0.5 1.5 1.0 + +[test] +name: read from image using linear filtering and normalized coords +kernel_name: readf_f_x + +arg_in: 1 image float 0.0 1.0 \ + 2.0 3.0 \ + image_type 2d \ + image_width 2 \ + image_height 2 \ + image_channel_order INTENSITY \ + image_channel_data_type FLOAT +arg_in: 2 sampler normalized_coords 1 \ + addressing_mode NONE \ + filter_mode LINEAR +arg_in: 3 buffer float2[4] 0.25 0.25 0.5 0.25 0.5 0.5 0.25 0.5 +arg_out: 0 buffer float[4] 0.0 0.5 1.5 1.0 + +[test] +name: read from image using clamp_to_edge addressing mode +kernel_name: readf_f_x + +arg_in: 1 image float 0.0 1.0 \ + 2.0 3.0 \ + image_type 2d \ + image_width 2 \ + image_height 2 \ + image_channel_order INTENSITY \ + image_channel_data_type FLOAT +arg_in: 2 sampler normalized_coords 0 \ + addressing_mode CLAMP_TO_EDGE \ + filter_mode NEAREST +arg_in: 3 buffer float2[4] -1.0 -1.0 3.0 -1.0 3.0 3.0 -1.0 3.0 +arg_out: 0 buffer float[4] 0.0 1.0 3.0 2.0 +!*/ + +kernel void readf_f_x(global float *out, read_only image2d_t img, sampler_t s, + global float2 *coords) { + int i = get_global_id(0); + out[i] = read_imagef(img, s, coords[i]).x; +} |