From 044f29f46f49bf0ee53bcb5e50a1f378a0e476e7 Mon Sep 17 00:00:00 2001 From: alan-baker Date: Tue, 17 Dec 2019 20:28:37 -0500 Subject: Auto-generate OpenCL literal samplers (#759) * Parse literal sampler descriptor map entries * add them to the pipeline sampler info * Before generating engine pipelines, create the sampler objects for the literal samplers * Tests --- .../cases/opencl_read_image_literal_sampler.amber | 69 ++++++++++++++++++++++ tests/run_tests.py | 1 + 2 files changed, 70 insertions(+) create mode 100644 tests/cases/opencl_read_image_literal_sampler.amber (limited to 'tests') diff --git a/tests/cases/opencl_read_image_literal_sampler.amber b/tests/cases/opencl_read_image_literal_sampler.amber new file mode 100644 index 0000000..4c72cfe --- /dev/null +++ b/tests/cases/opencl_read_image_literal_sampler.amber @@ -0,0 +1,69 @@ +#!amber +# Copyright 2019 The Amber Authors. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +SHADER vertex vs PASSTHROUGH + +SHADER fragment fs GLSL +#version 430 +layout(location = 0) out vec4 color_out; +void main() { + color_out = vec4(1.0, 0.0, 0.0, 1.0); +} +END + +SHADER compute read_imagef OPENCL-C +const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST; +kernel void foo(read_only image2d_t image, global float4* out) { + int gid_x = get_global_id(0); + int gid_y = get_global_id(1); + int linear = 2 * gid_y + gid_x; + float2 coord = (float2)(gid_x, gid_y); + out[linear] = read_imagef(image, sampler, coord); +} +END + +IMAGE texture DATA_TYPE vec4 DIM_2D WIDTH 2 HEIGHT 2 FILL 0.0 +BUFFER out_buf DATA_TYPE vec4 DATA +2.0 2.0 2.0 2.0 +2.0 2.0 2.0 2.0 +2.0 2.0 2.0 2.0 +2.0 2.0 2.0 2.0 +END +SAMPLER sampler + +PIPELINE compute read_pipe + ATTACH read_imagef ENTRY_POINT foo + BIND BUFFER out_buf KERNEL ARG_NAME out + BIND BUFFER texture KERNEL ARG_NAME image +END + +PIPELINE graphics fill_red + ATTACH vs + ATTACH fs + FRAMEBUFFER_SIZE 2 2 + BIND BUFFER texture AS color LOCATION 0 +END + +CLEAR_COLOR fill_red 0 0 3 3 +CLEAR fill_red +RUN fill_red DRAW_RECT POS 0 0 SIZE 2 2 + +RUN read_pipe 2 2 1 + +EXPECT out_buf IDX 0 EQ 1.0 0.0 0.0 1.0 +EXPECT out_buf IDX 16 EQ 1.0 0.0 0.0 1.0 +EXPECT out_buf IDX 32 EQ 1.0 0.0 0.0 1.0 +EXPECT out_buf IDX 48 EQ 1.0 0.0 0.0 1.0 + diff --git a/tests/run_tests.py b/tests/run_tests.py index e514fd7..2484be1 100755 --- a/tests/run_tests.py +++ b/tests/run_tests.py @@ -90,6 +90,7 @@ OPENCL_CASES = [ "opencl_c_copy.amber", "opencl_read_and_write_image3d_rgba32i.amber", "opencl_read_image.amber", + "opencl_read_image_literal_sampler.amber", "opencl_set_arg.amber", "opencl_write_image.amber", ] -- cgit v1.2.3