🏡 index : cold-air-networks/mirror-opencl-cts.git

📄 .appveyor.yml
📄 .clang-format
📄 .gitattributes
📁 .github
📁 workflows
📄 presubmit.yml
📄 CMakeLists.txt
📄 CODE_OF_CONDUCT.md
📄 LICENSE.txt
📄 README.md
📄 check-format.sh
📄 presubmit.sh
📄 readme-spir-v-binaries.txt
📁 test_common
📄 CMakeLists.txt
📁 autotest
📄 autotest.hpp
📄 test_suite.hpp
📄 config.hpp
📁 gl
📄 gl_headers.h
📄 helpers.cpp
📄 helpers.h
📄 setup.h
📄 setup_osx.cpp
📄 setup_win32.cpp
📄 setup_x11.cpp
📁 gles
📄 gl_headers.h
📄 helpers.cpp
📄 helpers.h
📄 setup.h
📁 harness
📄 ThreadPool.cpp
📄 ThreadPool.h
📄 alloc.h
📄 clImageHelper.h
📄 cl_offline_compiler-interface.txt
📄 compat.h
📄 conversions.cpp
📄 conversions.h
📄 crc32.cpp
📄 crc32.h
📄 deviceInfo.cpp
📄 deviceInfo.h
📄 errorHelpers.cpp
📄 errorHelpers.h
📄 featureHelpers.cpp
📄 featureHelpers.h
📄 fpcontrol.h
📄 genericThread.cpp
📄 genericThread.h
📄 imageHelpers.cpp
📄 imageHelpers.h
📄 kernelHelpers.cpp
📄 kernelHelpers.h
📄 mingw_compat.c
📄 mingw_compat.h
📄 msvc9.c
📄 mt19937.cpp
📄 mt19937.h
📄 os_helpers.cpp
📄 os_helpers.h
📄 parseParameters.cpp
📄 parseParameters.h
📄 propertyHelpers.cpp
📄 propertyHelpers.h
📄 ref_counting.h
📄 rounding_mode.cpp
📄 rounding_mode.h
📄 testHarness.cpp
📄 testHarness.h
📄 test_mt19937.c
📄 threadTesting.cpp
📄 threadTesting.h
📄 typeWrappers.cpp
📄 typeWrappers.h
📁 miniz
📄 CMakeLists.txt
📄 miniz.c
📄 miniz.h
📁 test_conformance
📄 CMakeCommon.txt
📄 CMakeLists.txt
📁 SVM
📄 CMakeLists.txt
📄 common.h
📄 main.cpp
📄 test_allocate_shared_buffer.cpp
📄 test_byte_granularity.cpp
📄 test_cross_buffer_pointers.cpp
📄 test_enqueue_api.cpp
📄 test_fine_grain_memory_consistency.cpp
📄 test_fine_grain_sync_buffers.cpp
📄 test_migrate.cpp
📄 test_pointer_passing.cpp
📄 test_set_kernel_exec_info_svm_ptrs.cpp
📄 test_shared_address_space_coarse_grain.cpp
📄 test_shared_address_space_fine_grain.cpp
📄 test_shared_address_space_fine_grain_buffers.cpp
📄 test_shared_sub_buffers.cpp
📁 allocations
📄 CMakeLists.txt
📄 allocation_execute.cpp
📄 allocation_execute.h
📄 allocation_fill.cpp
📄 allocation_fill.h
📄 allocation_functions.cpp
📄 allocation_functions.h
📄 allocation_utils.cpp
📄 allocation_utils.h
📄 main.cpp
📄 testBase.h
📁 api
📄 CMakeLists.txt
📄 main.cpp
📄 negative_platform.cpp
📄 procs.h
📄 testBase.h
📄 test_api_consistency.cpp
📄 test_api_min_max.cpp
📄 test_binary.cpp
📄 test_bool.cpp
📄 test_clone_kernel.cpp
📄 test_context_destructor_callback.cpp
📄 test_create_context_from_type.cpp
📄 test_create_kernels.cpp
📄 test_device_min_data_type_align_size_alignment.cpp
📄 test_kernel_arg_changes.cpp
📄 test_kernel_arg_info.cpp
📄 test_kernel_arg_info_compatibility.cpp
📄 test_kernel_arg_multi_setup.cpp
📄 test_kernel_private_memory_size.cpp
📄 test_kernels.cpp
📄 test_mem_object_info.cpp
📄 test_mem_object_properties_queries.cpp
📄 test_mem_objects.cpp
📄 test_min_image_formats.cpp
📄 test_native_kernel.cpp
📄 test_null_buffer_arg.cpp
📄 test_pipe_properties_queries.cpp
📄 test_platform.cpp
📄 test_queries.cpp
📄 test_queue.cpp
📄 test_queue_hint.cpp
📄 test_queue_properties.cpp
📄 test_queue_properties_queries.cpp
📄 test_retain.cpp
📄 test_retain_program.cpp
📄 test_sub_group_dispatch.cpp
📄 test_zero_sized_enqueue.cpp
📁 atomics
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 testBase.h
📄 test_atomics.cpp
📄 test_indexed_cases.cpp
📁 basic
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 run_array
📄 run_array_image_copy
📄 run_image
📄 run_multi_read_image
📄 test_arraycopy.cpp
📄 test_arrayimagecopy.cpp
📄 test_arrayimagecopy3d.cpp
📄 test_arrayreadwrite.cpp
📄 test_astype.cpp
📄 test_async_copy.cpp
📄 test_async_copy2D.cpp
📄 test_async_copy3D.cpp
📄 test_async_copy_fence.cpp
📄 test_async_strided_copy.cpp
📄 test_barrier.cpp
📄 test_basic_parameter_types.cpp
📄 test_bufferreadwriterect.cpp
📄 test_constant.cpp
📄 test_constant_source.cpp
📄 test_createkernelsinprogram.cpp
📄 test_enqueue_map.cpp
📄 test_enqueued_local_size.cpp
📄 test_explicit_s2v.cpp
📄 test_float2int.cpp
📄 test_fpmath_float.cpp
📄 test_fpmath_float2.cpp
📄 test_fpmath_float4.cpp
📄 test_get_linear_ids.cpp
📄 test_global_linear_id.cpp
📄 test_global_work_offsets.cpp
📄 test_hiloeo.cpp
📄 test_hostptr.cpp
📄 test_if.cpp
📄 test_imagecopy.cpp
📄 test_imagecopy3d.cpp
📄 test_imagedim.cpp
📄 test_imagenpot.cpp
📄 test_imagerandomcopy.cpp
📄 test_imagereadwrite.cpp
📄 test_imagereadwrite3d.cpp
📄 test_int2float.cpp
📄 test_intmath.cpp
📄 test_kernel_call_kernel_function.cpp
📄 test_kernel_memory_alignment.cpp
📄 test_local.cpp
📄 test_local_kernel_scope.cpp
📄 test_local_linear_id.cpp
📄 test_loop.cpp
📄 test_multireadimagemultifmt.cpp
📄 test_multireadimageonefmt.cpp
📄 test_numeric_constants.cpp
📄 test_pointercast.cpp
📄 test_preprocessors.cpp
📄 test_progvar.cpp
📄 test_queue_priority.cpp
📄 test_readimage.cpp
📄 test_readimage3d.cpp
📄 test_readimage3d_fp32.cpp
📄 test_readimage3d_int16.cpp
📄 test_readimage_fp32.cpp
📄 test_readimage_int16.cpp
📄 test_rw_image_access_qualifier.cpp
📄 test_simple_image_pitch.cpp
📄 test_sizeof.cpp
📄 test_vec_type_hint.cpp
📄 test_vector_creation.cpp
📄 test_vector_swizzle.cpp
📄 test_vloadstore.cpp
📄 test_wg_barrier.cpp
📄 test_work_item_functions.cpp
📄 test_writeimage.cpp
📄 test_writeimage_fp32.cpp
📄 test_writeimage_int16.cpp
📁 buffers
📄 CMakeLists.txt
📄 array_info.cpp
📄 main.cpp
📄 procs.h
📄 test_buffer_copy.cpp
📄 test_buffer_fill.cpp
📄 test_buffer_map.cpp
📄 test_buffer_mem.cpp
📄 test_buffer_migrate.cpp
📄 test_buffer_read.cpp
📄 test_buffer_write.cpp
📄 test_image_migrate.cpp
📄 test_sub_buffers.cpp
📁 c11_atomics
📄 CMakeLists.txt
📄 common.cpp
📄 common.h
📄 host_atomics.cpp
📄 host_atomics.h
📄 main.cpp
📄 test_atomics.cpp
📁 clcpp
📄 CMakeLists.txt
📁 address_spaces
📄 CMakeLists.txt
📄 common.hpp
📄 main.cpp
📄 test_pointer_types.hpp
📄 test_storage_types.hpp
📁 api
📄 CMakeLists.txt
📄 main.cpp
📄 test_ctors.hpp
📄 test_ctors_dtors.hpp
📄 test_dtors.hpp
📄 test_spec_consts.hpp
📁 atomics
📄 CMakeLists.txt
📄 atomic_fetch.hpp
📄 main.cpp
📁 attributes
📄 CMakeLists.txt
📄 main.cpp
📄 test_ivdep.hpp
📄 test_max_size.hpp
📄 test_required_num_sub_groups.hpp
📄 common.hpp
📁 common_funcs
📄 CMakeLists.txt
📄 common_funcs.hpp
📄 main.cpp
📁 convert
📄 CMakeLists.txt
📄 convert_cast.hpp
📄 main.cpp
📁 device_queue
📄 CMakeLists.txt
📄 main.cpp
📄 test_enqueue.hpp
📄 funcs_test_utils.hpp
📁 geometric_funcs
📄 CMakeLists.txt
📄 fast_geometric_funcs.hpp
📄 geometric_funcs.hpp
📄 main.cpp
📁 images
📄 CMakeLists.txt
📄 common.hpp
📄 main.cpp
📄 test_read.hpp
📄 test_sample.hpp
📄 test_write.hpp
📁 integer_funcs
📄 24bit_funcs.hpp
📄 CMakeLists.txt
📄 bitwise_funcs.hpp
📄 common.hpp
📄 main.cpp
📄 numeric_funcs.hpp
📁 math_funcs
📄 CMakeLists.txt
📄 common.hpp
📄 comparison_funcs.hpp
📄 exponential_funcs.hpp
📄 floating_point_funcs.hpp
📄 half_math_funcs.hpp
📄 logarithmic_funcs.hpp
📄 main.cpp
📄 other_funcs.hpp
📄 power_funcs.hpp
📄 reference.hpp
📄 trigonometric_funcs.hpp
📁 pipes
📄 CMakeLists.txt
📄 main.cpp
📄 test_pipes.hpp
📁 program_scope_ctors_dtors
📄 CMakeLists.txt
📄 common.hpp
📄 main.cpp
📄 test_ctors_dtors.hpp
📁 reinterpret
📄 CMakeLists.txt
📄 as_type.hpp
📄 main.cpp
📁 relational_funcs
📄 CMakeLists.txt
📄 common.hpp
📄 comparison_funcs.hpp
📄 main.cpp
📄 select_funcs.hpp
📄 test_funcs.hpp
📁 spec_constants
📄 CMakeLists.txt
📄 common.hpp
📄 main.cpp
📄 test_spec_consts_attributes.hpp
📄 test_spec_consts_if.hpp
📄 test_spec_consts_init_vars.hpp
📄 spirv10_2016.04.27.7z
📁 subgroups
📄 CMakeLists.txt
📄 common.hpp
📄 main.cpp
📄 test_sg_all.hpp
📄 test_sg_any.hpp
📄 test_sg_broadcast.hpp
📄 test_sg_reduce.hpp
📄 test_sg_scan_exclusive.hpp
📄 test_sg_scan_inclusive.hpp
📁 synchronization
📄 CMakeLists.txt
📄 main.cpp
📁 named_barrier
📄 common.hpp
📄 test_named_barrier.hpp
📄 test_spec_example.hpp
📄 test_sub_group_barrier.hpp
📄 test_work_group_barrier.hpp
📁 utils_common
📄 errors.hpp
📄 is_vector_type.hpp
📄 kernel_helpers.hpp
📄 make_vector_type.hpp
📄 scalar_type.hpp
📄 string.hpp
📄 type_name.hpp
📄 type_supported.hpp
📄 vector_size.hpp
📁 utils_test
📄 binary.hpp
📄 compare.hpp
📁 detail
📄 base_func_type.hpp
📄 vec_helpers.hpp
📄 generate_inputs.hpp
📄 ternary.hpp
📄 unary.hpp
📁 vload_vstore
📄 CMakeLists.txt
📄 common.hpp
📄 half_utils.hpp
📄 main.cpp
📄 vload_funcs.hpp
📄 vstore_funcs.hpp
📁 workgroups
📄 CMakeLists.txt
📄 common.hpp
📄 main.cpp
📄 test_wg_all.hpp
📄 test_wg_any.hpp
📄 test_wg_broadcast.hpp
📄 test_wg_reduce.hpp
📄 test_wg_scan_exclusive.hpp
📄 test_wg_scan_inclusive.hpp
📁 workitems
📄 CMakeLists.txt
📄 main.cpp
📄 test_workitems.hpp
📁 commonfns
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 test_binary_fn.cpp
📄 test_clamp.cpp
📄 test_degrees.cpp
📄 test_fmax.cpp
📄 test_fmaxf.cpp
📄 test_fmin.cpp
📄 test_fminf.cpp
📄 test_max.cpp
📄 test_maxf.cpp
📄 test_min.cpp
📄 test_minf.cpp
📄 test_mix.cpp
📄 test_radians.cpp
📄 test_sign.cpp
📄 test_smoothstep.cpp
📄 test_smoothstepf.cpp
📄 test_step.cpp
📄 test_stepf.cpp
📁 compiler
📄 CMakeLists.txt
📁 includeTestDirectory
📄 testIncludeFile.h
📄 main.cpp
📄 procs.h
📁 secondIncludeTestDirectory
📄 testIncludeFile.h
📄 testBase.h
📄 test_async_build.cpp
📄 test_build_helpers.cpp
📄 test_build_options.cpp
📄 test_compile.cpp
📄 test_compiler_defines_for_extensions.cpp
📄 test_feature_macro.cpp
📄 test_image_macro.cpp
📄 test_opencl_c_versions.cpp
📄 test_pragma_unroll.cpp
📄 test_preprocessor.cpp
📄 test_unload_platform_compiler.cpp
📄 test_unload_platform_compiler_resources.hpp
📄 write_kernel.spvasm32
📄 write_kernel.spvasm64
📁 computeinfo
📄 CMakeLists.txt
📄 conforming_version.cpp
📄 device_uuid.cpp
📄 extended_versioning.cpp
📄 main.cpp
📁 contractions
📄 CMakeLists.txt
📄 contractions.cpp
📁 conversions
📄 CMakeLists.txt
📄 Sleep.cpp
📄 Sleep.h
📄 basic_test_conversions.cpp
📄 basic_test_conversions.h
📄 fplib.cpp
📄 fplib.h
📄 run_batch
📄 test_conversions.cpp
📁 d3d10
📄 CMakeLists.txt
📄 buffer.cpp
📄 harness.cpp
📄 harness.h
📄 main.cpp
📄 misc.cpp
📄 texture2d.cpp
📄 texture3d.cpp
📁 d3d11
📄 CMakeLists.txt
📄 buffer.cpp
📄 harness.cpp
📄 harness.h
📄 main.cpp
📄 misc.cpp
📄 texture2d.cpp
📄 texture3d.cpp
📁 device_execution
📄 CMakeLists.txt
📄 device_info.cpp
📄 device_queue.cpp
📄 enqueue_block.cpp
📄 enqueue_flags.cpp
📄 enqueue_multi_queue.cpp
📄 enqueue_ndrange.cpp
📄 enqueue_profiling.cpp
📄 enqueue_wg_size.cpp
📄 execute_block.cpp
📄 host_multi_queue.cpp
📄 host_queue_order.cpp
📄 main.cpp
📄 nested_blocks.cpp
📄 procs.h
📄 utils.cpp
📄 utils.h
📁 device_partition
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 testBase.h
📄 test_device_partition.cpp
📁 device_timer
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 test_device_timer.cpp
📁 events
📄 CMakeLists.txt
📄 action_classes.cpp
📄 action_classes.h
📄 main.cpp
📄 procs.h
📄 testBase.h
📄 test_callbacks.cpp
📄 test_event_dependencies.cpp
📄 test_events.cpp
📄 test_userevents.cpp
📄 test_userevents_multithreaded.cpp
📄 test_waitlists.cpp
📁 extensions
📄 CMakeLists.txt
📁 cl_ext_cxx_for_opencl
📄 CMakeLists.txt
📄 cxx_for_opencl_ext.cpp
📄 cxx_for_opencl_ver.cpp
📄 main.cpp
📄 procs.h
📁 cl_khr_dx9_media_sharing
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 test_create_context.cpp
📄 test_functions_api.cpp
📄 test_functions_kernel.cpp
📄 test_get_device_ids.cpp
📄 test_interop_sync.cpp
📄 test_memory_access.cpp
📄 test_other_data_types.cpp
📄 utils.cpp
📄 utils.h
📄 wrappers.cpp
📄 wrappers.h
📄 generate_spirv_offline.py
📁 generic_address_space
📄 CMakeLists.txt
📄 advanced_tests.cpp
📄 base.h
📄 basic_tests.cpp
📄 main.cpp
📄 stress_tests.cpp
📁 geometrics
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 testBase.h
📄 test_geometrics.cpp
📄 test_geometrics_double.cpp
📁 gl
📄 CMakeLists.txt
📁 GLEW
📁 GL
📄 eglew.h
📄 glew.h
📄 glxew.h
📄 wglew.h
📄 glew.c
📄 common.h
📄 helpers.cpp
📄 main.cpp
📄 procs.h
📄 testBase.h
📄 test_buffers.cpp
📄 test_fence_sync.cpp
📄 test_image_methods.cpp
📄 test_images_1D.cpp
📄 test_images_1Darray.cpp
📄 test_images_2D.cpp
📄 test_images_2Darray.cpp
📄 test_images_3D.cpp
📄 test_images_depth.cpp
📄 test_images_getinfo_common.cpp
📄 test_images_multisample.cpp
📄 test_images_read_common.cpp
📄 test_images_write_common.cpp
📄 test_renderbuffer.cpp
📄 test_renderbuffer_info.cpp
📁 gles
📄 CMakeLists.txt
📄 helpers.cpp
📄 main.cpp
📄 procs.h
📄 setup_egl.cpp
📄 testBase.h
📄 test_buffers.cpp
📄 test_fence_sync.cpp
📄 test_images_2D.cpp
📄 test_images_2D_info.cpp
📄 test_images_3D.cpp
📄 test_images_3D_info.cpp
📄 test_renderbuffer.cpp
📄 test_renderbuffer_info.cpp
📁 half
📄 CMakeLists.txt
📄 Test_roundTrip.cpp
📄 Test_vLoadHalf.cpp
📄 Test_vStoreHalf.cpp
📄 cl_utils.cpp
📄 cl_utils.h
📄 main.cpp
📄 test_config.h
📄 tests.h
📁 images
📄 CMakeLists.txt
📁 clCopyImage
📄 CMakeLists.txt
📄 main.cpp
📄 test_copy_1D.cpp
📄 test_copy_1D_array.cpp
📄 test_copy_2D.cpp
📄 test_copy_2D_2D_array.cpp
📄 test_copy_2D_3D.cpp
📄 test_copy_2D_array.cpp
📄 test_copy_3D.cpp
📄 test_copy_3D_2D_array.cpp
📄 test_copy_generic.cpp
📄 test_loops.cpp
📁 clFillImage
📄 CMakeLists.txt
📄 main.cpp
📄 test_fill_1D.cpp
📄 test_fill_1D_array.cpp
📄 test_fill_2D.cpp
📄 test_fill_2D_array.cpp
📄 test_fill_3D.cpp
📄 test_fill_generic.cpp
📄 test_loops.cpp
📁 clGetInfo
📄 CMakeLists.txt
📄 main.cpp
📄 test_1D.cpp
📄 test_1D_2D_array.cpp
📄 test_2D.cpp
📄 test_3D.cpp
📄 test_loops.cpp
📁 clReadWriteImage
📄 CMakeLists.txt
📄 main.cpp
📄 test_loops.cpp
📄 test_read_1D.cpp
📄 test_read_1D_array.cpp
📄 test_read_2D.cpp
📄 test_read_2D_array.cpp
📄 test_read_3D.cpp
📄 common.cpp
📄 common.h
📁 kernel_image_methods
📄 CMakeLists.txt
📄 main.cpp
📄 test_1D.cpp
📄 test_1D_array.cpp
📄 test_2D.cpp
📄 test_2D_array.cpp
📄 test_3D.cpp
📄 test_loops.cpp
📁 kernel_read_write
📄 CMakeLists.txt
📄 main.cpp
📄 test_common.cpp
📄 test_common.h
📄 test_iterations.cpp
📄 test_loops.cpp
📄 test_read_1D.cpp
📄 test_read_1D_array.cpp
📄 test_read_2D_array.cpp
📄 test_read_3D.cpp
📄 test_write_1D.cpp
📄 test_write_1D_array.cpp
📄 test_write_2D_array.cpp
📄 test_write_3D.cpp
📄 test_write_image.cpp
📁 samplerlessReads
📄 CMakeLists.txt
📄 main.cpp
📄 test_iterations.cpp
📄 test_loops.cpp
📄 test_read_1D.cpp
📄 test_read_1D_array.cpp
📄 test_read_1D_buffer.cpp
📄 test_read_2D_array.cpp
📄 test_read_3D.cpp
📄 testBase.h
📁 integer_ops
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 testBase.h
📄 test_abs.cpp
📄 test_absdiff.cpp
📄 test_add_sat.cpp
📄 test_int_basic_ops.cpp
📄 test_integers.cpp
📄 test_intmad24.cpp
📄 test_intmul24.cpp
📄 test_popcount.cpp
📄 test_sub_sat.cpp
📄 test_unary_ops.cpp
📄 test_upsample.cpp
📄 verification_and_generation_functions.cpp
📁 math_brute_force
📄 CMakeLists.txt
📄 README.txt
📄 binary_double.cpp
📄 binary_float.cpp
📄 binary_i_double.cpp
📄 binary_i_float.cpp
📄 binary_operator_double.cpp
📄 binary_operator_float.cpp
📄 binary_two_results_i_double.cpp
📄 binary_two_results_i_float.cpp
📄 function_list.cpp
📄 function_list.h
📄 i_unary_double.cpp
📄 i_unary_float.cpp
📄 macro_binary_double.cpp
📄 macro_binary_float.cpp
📄 macro_unary_double.cpp
📄 macro_unary_float.cpp
📄 mad_double.cpp
📄 mad_float.cpp
📄 main.cpp
📄 reference_math.cpp
📄 reference_math.h
📄 run_math_brute_force_in_parallel.py
📄 sleep.cpp
📄 sleep.h
📄 ternary_double.cpp
📄 ternary_float.cpp
📄 test_functions.h
📄 unary_double.cpp
📄 unary_float.cpp
📄 unary_two_results_double.cpp
📄 unary_two_results_float.cpp
📄 unary_two_results_i_double.cpp
📄 unary_two_results_i_float.cpp
📄 unary_u_double.cpp
📄 unary_u_float.cpp
📄 utility.cpp
📄 utility.h
📁 mem_host_flags
📄 CMakeLists.txt
📄 C_host_memory_block.h
📄 checker.h
📄 checker_image_mem_host_no_access.hpp
📄 checker_image_mem_host_read_only.hpp
📄 checker_image_mem_host_write_only.hpp
📄 checker_mem_host_no_access.hpp
📄 checker_mem_host_read_only.hpp
📄 checker_mem_host_write_only.hpp
📄 main.cpp
📄 mem_host_buffer.cpp
📄 mem_host_image.cpp
📄 procs.h
📄 testBase.h
📁 multiple_device_context
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 testBase.h
📄 test_multiple_contexts.cpp
📄 test_multiple_devices.cpp
📁 non_uniform_work_group
📄 CMakeLists.txt
📄 TestNonUniformWorkGroup.cpp
📄 TestNonUniformWorkGroup.h
📄 main.cpp
📄 procs.h
📄 test_advanced_2d.cpp
📄 test_advanced_3d.cpp
📄 test_advanced_other.cpp
📄 test_basic.cpp
📄 tools.cpp
📄 tools.h
📄 opencl_conformance_tests_conversions.csv
📄 opencl_conformance_tests_d3d.csv
📄 opencl_conformance_tests_full.csv
📄 opencl_conformance_tests_full_binary.csv
📄 opencl_conformance_tests_full_no_math_or_conversions.csv
📄 opencl_conformance_tests_full_spirv.csv
📄 opencl_conformance_tests_math.csv
📄 opencl_conformance_tests_quick.csv
📁 pipes
📄 CMakeLists.txt
📄 kernels.h
📄 main.cpp
📄 procs.h
📄 test_pipe_info.cpp
📄 test_pipe_limits.cpp
📄 test_pipe_query_functions.cpp
📄 test_pipe_read_write.cpp
📄 test_pipe_readwrite_errors.cpp
📄 test_pipe_subgroups.cpp
📁 printf
📄 CMakeLists.txt
📄 test_printf.cpp
📄 test_printf.h
📄 util_printf.cpp
📁 profiling
📄 CMakeLists.txt
📄 copy.cpp
📄 execute.cpp
📄 execute_multipass.cpp
📄 main.cpp
📄 procs.h
📄 readArray.cpp
📄 readImage.cpp
📄 writeArray.cpp
📄 writeImage.cpp
📁 relationals
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 testBase.h
📄 test_comparisons_double.cpp
📄 test_comparisons_float.cpp
📄 test_relationals.cpp
📄 test_shuffles.cpp
📄 run_conformance.py
📁 select
📄 CMakeLists.txt
📄 test_select.cpp
📄 test_select.h
📄 util_select.cpp
📁 spir
📄 CMakeLists.txt
📄 api.zip
📄 atomics.zip
📄 basic.zip
📄 binary_type.zip
📄 commonfns.zip
📄 compile_and_link.zip
📄 conversions.zip
📄 datagen.cpp
📄 datagen.h
📄 enum_values.zip
📄 exceptions.h
📄 geometrics.zip
📄 half.zip
📄 images_kernel_read_write.zip
📄 images_samplerlessRead.zip
📄 integer_ops.zip
📄 kernel_attributes.zip
📄 kernel_image_methods.zip
📄 kernelargs.cpp
📄 kernelargs.h
📄 khr.csv
📄 main.cpp
📄 math_brute_force.zip
📄 printf.zip
📄 profiling.zip
📄 relationals.zip
📄 run_build_test.cpp
📄 run_build_test.h
📄 run_services.cpp
📄 run_services.h
📄 sampler_enumeration.zip
📄 select.zip
📄 typeinfo.h
📄 vec_align.zip
📄 vec_step.zip
📁 spirv_new
📄 CMakeLists.txt
📄 README.txt
📄 assemble_spirv.py
📄 main.cpp
📄 procs.h
📁 spirv_asm
📄 atomic_dec_global.spvasm32
📄 atomic_dec_global.spvasm64
📄 atomic_inc_global.spvasm32
📄 atomic_inc_global.spvasm64
📄 branch_conditional.spvasm32
📄 branch_conditional.spvasm64
📄 branch_conditional_weighted.spvasm32
📄 branch_conditional_weighted.spvasm64
📄 branch_simple.spvasm32
📄 branch_simple.spvasm64
📄 composite_construct_int4.spvasm32
📄 composite_construct_int4.spvasm64
📄 composite_construct_struct.spvasm32
📄 composite_construct_struct.spvasm64
📄 constant_char_simple.spvasm32
📄 constant_char_simple.spvasm64
📄 constant_double_simple.spvasm32
📄 constant_double_simple.spvasm64
📄 constant_false_simple.spvasm32
📄 constant_false_simple.spvasm64
📄 constant_float_simple.spvasm32
📄 constant_float_simple.spvasm64
📄 constant_half_simple.spvasm32
📄 constant_half_simple.spvasm64
📄 constant_int3_simple.spvasm32
📄 constant_int3_simple.spvasm64
📄 constant_int4_simple.spvasm32
📄 constant_int4_simple.spvasm64
📄 constant_int_simple.spvasm32
📄 constant_int_simple.spvasm64
📄 constant_long_simple.spvasm32
📄 constant_long_simple.spvasm64
📄 constant_short_simple.spvasm32
📄 constant_short_simple.spvasm64
📄 constant_struct_int_char_simple.spvasm32
📄 constant_struct_int_char_simple.spvasm64
📄 constant_true_simple.spvasm64
📄 constant_uchar_simple.spvasm32
📄 constant_uchar_simple.spvasm64
📄 constant_uint_simple.spvasm32
📄 constant_uint_simple.spvasm64
📄 constant_ulong_simple.spvasm32
📄 constant_ulong_simple.spvasm64
📄 constant_ushort_simple.spvasm32
📄 constant_ushort_simple.spvasm64
📄 copy_char_simple.spvasm32
📄 copy_char_simple.spvasm64
📄 copy_double_simple.spvasm32
📄 copy_double_simple.spvasm64
📄 copy_float_simple.spvasm32
📄 copy_float_simple.spvasm64
📄 copy_half_simple.spvasm32
📄 copy_half_simple.spvasm64
📄 copy_int3_simple.spvasm32
📄 copy_int3_simple.spvasm64
📄 copy_int4_simple.spvasm32
📄 copy_int4_simple.spvasm64
📄 copy_int_simple.spvasm32
📄 copy_int_simple.spvasm64
📄 copy_long_simple.spvasm32
📄 copy_long_simple.spvasm64
📄 copy_short_simple.spvasm32
📄 copy_short_simple.spvasm64
📄 copy_struct_int_char_simple.spvasm32
📄 copy_struct_int_char_simple.spvasm64
📄 copy_struct_int_float_simple.spvasm32
📄 copy_struct_int_float_simple.spvasm64
📄 copy_uint_simple.spvasm64
📄 copy_ulong_simple.spvasm32
📄 copy_ulong_simple.spvasm64
📄 copy_ushort_simple.spvasm32
📄 copy_ushort_simple.spvasm64
📄 decorate_aliased.spvasm32
📄 decorate_aliased.spvasm64
📄 decorate_alignment.spvasm32
📄 decorate_alignment.spvasm64
📄 decorate_constant.spvasm32
📄 decorate_constant.spvasm64
📄 decorate_constant_fail.spvasm32
📄 decorate_constant_fail.spvasm64
📄 decorate_cpacked.spvasm32
📄 decorate_cpacked.spvasm64
📄 decorate_restrict.spvasm32
📄 decorate_restrict.spvasm64
📄 decorate_rounding_rte_double_long.spvasm32
📄 decorate_rounding_rte_double_long.spvasm64
📄 decorate_rounding_rte_float_int.spvasm32
📄 decorate_rounding_rte_float_int.spvasm64
📄 decorate_rounding_rtn_double_long.spvasm32
📄 decorate_rounding_rtn_double_long.spvasm64
📄 decorate_rounding_rtn_float_int.spvasm32
📄 decorate_rounding_rtn_float_int.spvasm64
📄 decorate_rounding_rtp_double_long.spvasm32
📄 decorate_rounding_rtp_double_long.spvasm64
📄 decorate_rounding_rtp_float_int.spvasm32
📄 decorate_rounding_rtp_float_int.spvasm64
📄 decorate_rounding_rtz_double_long.spvasm32
📄 decorate_rounding_rtz_double_long.spvasm64
📄 decorate_saturated_conversion_int.spvasm64
📄 decorate_saturated_conversion_short.spvasm32
📄 decorate_saturated_conversion_short.spvasm64
📄 decorate_saturated_conversion_uchar.spvasm32
📄 decorate_saturated_conversion_uchar.spvasm64
📄 decorate_saturated_conversion_uint.spvasm32
📄 decorate_saturated_conversion_uint.spvasm64
📄 decorate_saturated_conversion_ushort.spvasm32
📄 decorate_saturated_conversion_ushort.spvasm64
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_int.spvasm32
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_int.spvasm64
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_uint.spvasm32
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_uint.spvasm64
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_int.spvasm32
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_int.spvasm64
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_uint.spvasm32
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_uint.spvasm64
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fnegate_int.spvasm32
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fnegate_int.spvasm64
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_int.spvasm32
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_int.spvasm64
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_uint.spvasm32
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_uint.spvasm64
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_int.spvasm32
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_int.spvasm64
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_uint.spvasm32
📄 ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_uint.spvasm64
📄 fadd_double.spvasm32
📄 fadd_double.spvasm64
📄 fadd_double2.spvasm32
📄 fadd_double2.spvasm64
📄 fadd_half.spvasm64
📄 fdiv_double.spvasm32
📄 fdiv_double.spvasm64
📄 fdiv_double2.spvasm32
📄 fdiv_double2.spvasm64
📄 fdiv_float.spvasm32
📄 fdiv_float.spvasm64
📄 fdiv_float4.spvasm32
📄 fdiv_float4.spvasm64
📄 fdiv_half.spvasm32
📄 fdiv_half.spvasm64
📄 fmod_double.spvasm32
📄 fmod_double.spvasm64
📄 fmod_double2.spvasm32
📄 fmod_double2.spvasm64
📄 fmod_float.spvasm32
📄 fmod_float.spvasm64
📄 fmod_float4.spvasm32
📄 fmod_float4.spvasm64
📄 fmod_half.spvasm32
📄 fmod_half.spvasm64
📄 fmul_double.spvasm32
📄 fmul_double.spvasm64
📄 fmul_double2.spvasm32
📄 fmul_double2.spvasm64
📄 fmul_float.spvasm32
📄 frem_double.spvasm32
📄 frem_double.spvasm64
📄 frem_double2.spvasm32
📄 frem_double2.spvasm64
📄 frem_float.spvasm32
📄 frem_float.spvasm64
📄 frem_float4.spvasm32
📄 frem_float4.spvasm64
📄 frem_half.spvasm32
📄 frem_half.spvasm64
📄 fsub_double.spvasm32
📄 fsub_double.spvasm64
📄 fsub_double2.spvasm32
📄 fsub_double2.spvasm64
📄 fsub_float.spvasm32
📄 fsub_float.spvasm64
📄 fsub_float4.spvasm32
📄 fsub_float4.spvasm64
📄 fsub_half.spvasm32
📄 fsub_half.spvasm64
📄 label_simple.spvasm32
📄 label_simple.spvasm64
📄 lifetime_simple.spvasm32
📄 lifetime_simple.spvasm64
📄 linkage_export.spvasm32
📄 linkage_export.spvasm64
📄 linkage_import.spvasm32
📄 linkage_import.spvasm64
📄 loop_merge_branch_conditional_dont_unroll.spvasm32
📄 loop_merge_branch_conditional_dont_unroll.spvasm64
📄 loop_merge_branch_conditional_none.spvasm32
📄 loop_merge_branch_none.spvasm32
📄 loop_merge_branch_none.spvasm64
📄 loop_merge_branch_unroll.spvasm32
📄 loop_merge_branch_unroll.spvasm64
📄 op_function_const.spvasm32
📄 op_function_const.spvasm64
📄 op_function_inline.spvasm32
📄 op_function_inline.spvasm64
📄 op_function_noinline.spvasm32
📄 op_function_noinline.spvasm64
📄 op_function_none.spvasm32
📄 op_function_none.spvasm64
📄 op_function_pure.spvasm32
📄 op_function_pure.spvasm64
📄 op_function_pure_ptr.spvasm32
📄 op_function_pure_ptr.spvasm64
📄 op_neg_double.spvasm32
📄 op_neg_double.spvasm64
📄 op_neg_float.spvasm32
📄 op_neg_float.spvasm64
📄 op_neg_float4.spvasm32
📄 op_neg_float4.spvasm64
📄 op_neg_int.spvasm32
📄 op_neg_int.spvasm64
📄 op_neg_int4.spvasm32
📄 op_neg_int4.spvasm64
📄 op_neg_long.spvasm32
📄 op_neg_long.spvasm64
📄 op_neg_short.spvasm32
📄 op_neg_short.spvasm64
📄 op_not_int.spvasm32
📄 op_not_short.spvasm32
📄 op_not_short.spvasm64
📄 op_spec_constant_double_simple.spvasm32
📄 op_spec_constant_double_simple.spvasm64
📄 op_spec_constant_false_simple.spvasm32
📄 op_spec_constant_false_simple.spvasm64
📄 op_spec_constant_float_simple.spvasm32
📄 op_spec_constant_float_simple.spvasm64
📄 op_spec_constant_half_simple.spvasm32
📄 op_spec_constant_half_simple.spvasm64
📄 op_spec_constant_true_simple.spvasm32
📄 op_spec_constant_true_simple.spvasm64
📄 op_spec_constant_uchar_simple.spvasm32
📄 op_spec_constant_uchar_simple.spvasm64
📄 op_spec_constant_uint_simple.spvasm32
📄 op_spec_constant_uint_simple.spvasm64
📄 op_spec_constant_ulong_simple.spvasm32
📄 op_spec_constant_ulong_simple.spvasm64
📄 op_spec_constant_ushort_simple.spvasm32
📄 op_spec_constant_ushort_simple.spvasm64
📄 opaque.spvasm32
📄 opaque.spvasm64
📄 phi_2.spvasm32
📄 phi_2.spvasm64
📄 phi_3.spvasm32
📄 phi_3.spvasm64
📄 phi_4.spvasm32
📄 phi_4.spvasm64
📄 select_if_dont_flatten.spvasm32
📄 select_if_dont_flatten.spvasm64
📄 select_if_flatten.spvasm32
📄 select_switch_flatten.spvasm32
📄 select_switch_flatten.spvasm64
📄 select_switch_none.spvasm32
📄 select_switch_none.spvasm64
📄 undef_char_simple.spvasm32
📄 undef_char_simple.spvasm64
📄 undef_double_simple.spvasm32
📄 undef_double_simple.spvasm64
📄 undef_false_simple.spvasm32
📄 undef_false_simple.spvasm64
📄 undef_float_simple.spvasm32
📄 undef_float_simple.spvasm64
📄 undef_half_simple.spvasm32
📄 undef_half_simple.spvasm64
📄 undef_int3_simple.spvasm32
📄 undef_int3_simple.spvasm64
📄 undef_int4_simple.spvasm32
📄 undef_int4_simple.spvasm64
📄 undef_int_simple.spvasm32
📄 undef_int_simple.spvasm64
📄 undef_long_simple.spvasm32
📄 undef_long_simple.spvasm64
📄 undef_short_simple.spvasm32
📄 undef_short_simple.spvasm64
📄 undef_struct_int_char_simple.spvasm32
📄 undef_struct_int_char_simple.spvasm64
📄 undef_struct_int_float_simple.spvasm32
📄 undef_struct_int_float_simple.spvasm64
📄 undef_struct_struct_simple.spvasm32
📄 undef_struct_struct_simple.spvasm64
📄 undef_true_simple.spvasm32
📄 undef_ulong_simple.spvasm32
📄 undef_ulong_simple.spvasm64
📄 undef_ushort_simple.spvasm32
📄 undef_ushort_simple.spvasm64
📄 unreachable_simple.spvasm32
📄 unreachable_simple.spvasm64
📄 vector_char16_extract.spvasm32
📄 vector_char16_extract.spvasm64
📄 vector_char16_insert.spvasm32
📄 vector_char16_insert.spvasm64
📄 vector_double2_extract.spvasm32
📄 vector_double2_extract.spvasm64
📄 vector_double2_insert.spvasm32
📄 vector_double2_insert.spvasm64
📄 vector_float4_extract.spvasm32
📄 vector_float4_extract.spvasm64
📄 vector_float4_insert.spvasm32
📄 vector_float4_insert.spvasm64
📄 vector_int4_extract.spvasm32
📄 vector_int4_extract.spvasm64
📄 vector_int4_insert.spvasm32
📄 vector_int4_insert.spvasm64
📄 vector_long2_extract.spvasm32
📄 vector_long2_extract.spvasm64
📄 vector_long2_insert.spvasm32
📄 vector_long2_insert.spvasm64
📄 vector_times_scalar_double.spvasm32
📄 vector_times_scalar_double.spvasm64
📄 vector_times_scalar_float.spvasm32
📄 vector_times_scalar_float.spvasm64
📄 testBase.h
📄 test_cl_khr_spirv_no_integer_wrap_decoration.cpp
📄 test_decorate.cpp
📄 test_get_program_il.cpp
📄 test_linkage.cpp
📄 test_op_atomic.cpp
📄 test_op_branch.cpp
📄 test_op_branch_conditional.cpp
📄 test_op_composite_construct.cpp
📄 test_op_constant.cpp
📄 test_op_copy_object.cpp
📄 test_op_fmath.cpp
📄 test_op_function.cpp
📄 test_op_lifetime.cpp
📄 test_op_loop_merge.cpp
📄 test_op_negate.cpp
📄 test_op_opaque.cpp
📄 test_op_phi.cpp
📄 test_op_selection_merge.cpp
📄 test_op_spec_constant.cpp
📄 test_op_undef.cpp
📄 test_op_vector_extract.cpp
📄 test_op_vector_insert.cpp
📄 test_op_vector_times_scalar.cpp
📄 types.hpp
📁 subgroups
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 subgroup_common_kernels.cpp
📄 subgroup_common_kernels.h
📄 subgroup_common_templates.h
📄 subhelpers.h
📄 test_barrier.cpp
📄 test_ifp.cpp
📄 test_queries.cpp
📄 test_subgroup.cpp
📄 test_subgroup_ballot.cpp
📄 test_subgroup_clustered_reduce.cpp
📄 test_subgroup_extended_types.cpp
📄 test_subgroup_non_uniform_arithmetic.cpp
📄 test_subgroup_non_uniform_vote.cpp
📄 test_subgroup_shuffle.cpp
📄 test_subgroup_shuffle_relative.cpp
📄 test_workitem.cpp
📄 submission_details_template.txt
📁 thread_dimensions
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 test_thread_dimensions.cpp
📁 vectors
📄 CMakeLists.txt
📄 defines.h
📄 globals.cpp
📄 main.cpp
📄 procs.h
📄 structs.cpp
📄 structs.h
📄 testBase.h
📄 test_step.cpp
📄 test_vec_align.cpp
📄 type_replacer.cpp
📄 type_replacer.h
📁 workgroups
📄 CMakeLists.txt
📄 main.cpp
📄 procs.h
📄 testBase.h
📄 test_wg_all.cpp
📄 test_wg_any.cpp
📄 test_wg_broadcast.cpp
📄 test_wg_reduce.cpp
📄 test_wg_reduce_max.cpp
📄 test_wg_reduce_min.cpp
📄 test_wg_scan_exclusive_add.cpp
📄 test_wg_scan_exclusive_max.cpp
📄 test_wg_scan_exclusive_min.cpp
📄 test_wg_scan_inclusive_add.cpp
📄 test_wg_scan_inclusive_max.cpp
📄 test_wg_scan_inclusive_min.cpp 
//
// Copyright (c) 2017 The Khronos Group Inc.
// 
// 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
//
//    http://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.
//
#include "testBase.h"
#include "common.h"
#include <limits.h>

#if defined( __APPLE__ )
    #include <OpenGL/glu.h>
#else
    #include <GL/glu.h>
    #include <CL/cl_gl.h>
#endif

#pragma mark -
#pragma mark Write test kernels

static const char *kernelpattern_image_write_1D =
"__kernel void sample_test( __global %s4 *source, write_only image1d_t dest )\n"
"{\n"
"    uint index = get_global_id(0);\n"
"    %s4 value = source[index];\n"
"    write_image%s( dest, index, %s(value));\n"
"}\n";

static const char *kernelpattern_image_write_1D_half =
"__kernel void sample_test( __global half4 *source, write_only image1d_t dest )\n"
"{\n"
"    uint index = get_global_id(0);\n"
"    write_imagef( dest, index, vload_half4(index, (__global half *)source));\n"
"}\n";

static const char *kernelpattern_image_write_1D_buffer =
"__kernel void sample_test( __global %s4 *source, write_only image1d_buffer_t dest )\n"
"{\n"
"    uint index = get_global_id(0);\n"
"    %s4 value = source[index];\n"
"    write_image%s( dest, index, %s(value));\n"
"}\n";

static const char *kernelpattern_image_write_1D_buffer_half =
"__kernel void sample_test( __global half4 *source, write_only image1d_buffer_t dest )\n"
"{\n"
"    uint index = get_global_id(0);\n"
"    write_imagef( dest, index, vload_half4(index, (__global half *)source));\n"
"}\n";

static const char *kernelpattern_image_write_2D =
"__kernel void sample_test( __global %s4 *source, write_only image2d_t dest )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    uint index = tidY * get_image_width( dest ) + tidX;\n"
"    %s4 value = source[index];\n"
"    write_image%s( dest, (int2)( tidX, tidY ), %s(value));\n"
"}\n";

static const char *kernelpattern_image_write_2D_half =
"__kernel void sample_test( __global half4 *source, write_only image2d_t dest )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    uint index = tidY * get_image_width( dest ) + tidX;\n"
"    write_imagef( dest, (int2)( tidX, tidY ), vload_half4(index, (__global half *)source));\n"
"}\n";

static const char *kernelpattern_image_write_1Darray =
"__kernel void sample_test( __global %s4 *source, write_only image1d_array_t dest )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    uint index = tidY * get_image_width( dest ) + tidX;\n"
"    %s4 value = source[index];\n"
"    write_image%s( dest, (int2)( tidX, tidY ), %s(value));\n"
"}\n";

static const char *kernelpattern_image_write_1Darray_half =
"__kernel void sample_test( __global half4 *source, write_only image1d_array_t dest )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    uint index = tidY * get_image_width( dest ) + tidX;\n"
"    write_imagef( dest, (int2)( tidX, tidY ), vload_half4(index, (__global half *)source));\n"
"}\n";

static const char *kernelpattern_image_write_3D =
"#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable\n"
"__kernel void sample_test( __global %s4 *source, write_only image3d_t dest )\n"
"{\n"
"    int  tidX   = get_global_id(0);\n"
"    int  tidY   = get_global_id(1);\n"
"    int  tidZ   = get_global_id(2);\n"
"    int  width  = get_image_width( dest );\n"
"    int  height = get_image_height( dest );\n"
"    int  index = tidZ * width * height + tidY * width + tidX;\n"
"    %s4 value = source[index];\n"
"    write_image%s( dest, (int4)( tidX, tidY, tidZ, 0 ), %s(value));\n"
"}\n";

static const char *kernelpattern_image_write_3D_half =
"#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable\n"
"__kernel void sample_test( __global half4 *source, write_only image3d_t dest )\n"
"{\n"
"    int  tidX   = get_global_id(0);\n"
"    int  tidY   = get_global_id(1);\n"
"    int  tidZ   = get_global_id(2);\n"
"    int  width  = get_image_width( dest );\n"
"    int  height = get_image_height( dest );\n"
"    int  index = tidZ * width * height + tidY * width + tidX;\n"
"    write_imagef( dest, (int4)( tidX, tidY, tidZ, 0 ), vload_half4(index, (__global half *)source));\n"
"}\n";

static const char *kernelpattern_image_write_2Darray =
"__kernel void sample_test( __global %s4 *source, write_only image2d_array_t dest )\n"
"{\n"
"    int  tidX   = get_global_id(0);\n"
"    int  tidY   = get_global_id(1);\n"
"    int  tidZ   = get_global_id(2);\n"
"    int  width  = get_image_width( dest );\n"
"    int  height = get_image_height( dest );\n"
"    int  index = tidZ * width * height + tidY * width + tidX;\n"
"    %s4 value = source[index];\n"
"    write_image%s( dest, (int4)( tidX, tidY, tidZ, 0 ), %s(value));\n"
"}\n";

static const char *kernelpattern_image_write_2Darray_half =
"__kernel void sample_test( __global half4 *source, write_only image2d_array_t dest )\n"
"{\n"
"    int  tidX   = get_global_id(0);\n"
"    int  tidY   = get_global_id(1);\n"
"    int  tidZ   = get_global_id(2);\n"
"    int  width  = get_image_width( dest );\n"
"    int  height = get_image_height( dest );\n"
"    int  index = tidZ * width * height + tidY * width + tidX;\n"
"    write_imagef( dest, (int4)( tidX, tidY, tidZ, 0 ), vload_half4(index, (__global half *)source));\n"
"}\n";

#ifdef GL_VERSION_3_2

static const char * kernelpattern_image_write_2D_depth =
"__kernel void sample_test( __global %s *source, write_only image2d_depth_t dest )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    uint index = tidY * get_image_width( dest ) + tidX;\n"
"    float value = source[index];\n"
"    write_imagef( dest, (int2)( tidX, tidY ), value);\n"
"}\n";

static const char * kernelpattern_image_write_2D_array_depth =
"__kernel void sample_test( __global %s *source, write_only image2d_array_depth_t dest )\n"
"{\n"
"    int  tidX   = get_global_id(0);\n"
"    int  tidY   = get_global_id(1);\n"
"    int  tidZ   = get_global_id(2);\n"
"    int  width  = get_image_width( dest );\n"
"    int  height = get_image_height( dest );\n"
"    int  index = tidZ * width * height + tidY * width + tidX;\n"
"    %s value = source[index];\n"
"    write_image%s( dest, (int4)( tidX, tidY, tidZ, 0 ), %s(value));\n"
"}\n";


#endif

#pragma mark -
#pragma mark Utility functions

static const char* get_appropriate_write_kernel(GLenum target,
  ExplicitType type, cl_channel_order channel_order)
{
  switch (get_base_gl_target(target)) {
    case GL_TEXTURE_1D:

      if (type == kHalf)
        return kernelpattern_image_write_1D_half;
      else
        return kernelpattern_image_write_1D;
      break;
    case GL_TEXTURE_BUFFER:
       if (type == kHalf)
        return kernelpattern_image_write_1D_buffer_half;
      else
        return kernelpattern_image_write_1D_buffer;
      break;
    case GL_TEXTURE_1D_ARRAY:
      if (type == kHalf)
        return kernelpattern_image_write_1Darray_half;
      else
        return kernelpattern_image_write_1Darray;
      break;
    case GL_COLOR_ATTACHMENT0:
    case GL_RENDERBUFFER:
    case GL_TEXTURE_RECTANGLE_EXT:
    case GL_TEXTURE_2D:
    case GL_TEXTURE_CUBE_MAP:
#ifdef GL_VERSION_3_2
      if (channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
        return kernelpattern_image_write_2D_depth;
#endif
      if (type == kHalf)
        return kernelpattern_image_write_2D_half;
      else
        return kernelpattern_image_write_2D;
      break;

    case GL_TEXTURE_2D_ARRAY:
#ifdef GL_VERSION_3_2
      if (channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
        return kernelpattern_image_write_2D_array_depth;
#endif
      if (type == kHalf)
        return kernelpattern_image_write_2Darray_half;
      else
        return kernelpattern_image_write_2Darray;
      break;

    case GL_TEXTURE_3D:
      if (type == kHalf)
        return kernelpattern_image_write_3D_half;
      else
        return kernelpattern_image_write_3D;
      break;

    default:
      log_error("Unsupported GL tex target (%s) passed to write test: "
        "%s (%s):%d", GetGLTargetName(target), __FUNCTION__,
        __FILE__, __LINE__);
      return NULL;
  }
}

void set_dimensions_by_target(GLenum target, size_t *dims, size_t sizes[3],
  size_t width, size_t height, size_t depth)
{
  switch (get_base_gl_target(target)) {
    case GL_TEXTURE_1D:
      sizes[0] = width;
      *dims = 1;
      break;

    case GL_TEXTURE_BUFFER:
      sizes[0] = width;
      *dims = 1;
      break;

    case GL_TEXTURE_1D_ARRAY:
      sizes[0] = width;
      sizes[1] = height;
      *dims = 2;
      break;

    case GL_COLOR_ATTACHMENT0:
    case GL_RENDERBUFFER:
    case GL_TEXTURE_RECTANGLE_EXT:
    case GL_TEXTURE_2D:
    case GL_TEXTURE_CUBE_MAP:

      sizes[0] = width;
      sizes[1] = height;
      *dims = 2;
      break;

    case GL_TEXTURE_2D_ARRAY:
      sizes[0] = width;
      sizes[1] = height;
      sizes[2] = depth;
      *dims = 3;
      break;

    case GL_TEXTURE_3D:
      sizes[0] = width;
      sizes[1] = height;
      sizes[2] = depth;
      *dims = 3;
      break;

    default:
      log_error("Unsupported GL tex target (%s) passed to write test: "
        "%s (%s):%d", GetGLTargetName(target), __FUNCTION__,
        __FILE__, __LINE__);
  }
}

int test_cl_image_write( cl_context context, cl_command_queue queue,
  GLenum target, cl_mem clImage, size_t width, size_t height, size_t depth,
  cl_image_format *outFormat, ExplicitType *outType, void **outSourceBuffer,
  MTdata d, bool supports_half )
{
  size_t global_dims, global_sizes[3];
  clProgramWrapper program;
  clKernelWrapper kernel;
  clMemWrapper inStream;
  char* programPtr;
  int error;
  char kernelSource[2048];

  // What CL format did we get from the texture?

  error = clGetImageInfo(clImage, CL_IMAGE_FORMAT, sizeof(cl_image_format),
    outFormat, NULL);
  test_error(error, "Unable to get the CL image format");

  // Create the kernel source.  The target and the data type will influence
  // which particular kernel we choose.

  *outType = get_write_kernel_type( outFormat );
  size_t channelSize = get_explicit_type_size(*outType);

  const char* appropriateKernel = get_appropriate_write_kernel(target,
    *outType, outFormat->image_channel_order);
  if (*outType == kHalf && !supports_half) {
    log_info("cl_khr_fp16 isn't supported. Skip this test.\n");
    return 0;
  }

  const char* suffix = get_kernel_suffix( outFormat );
  const char* convert = get_write_conversion( outFormat, *outType );

  sprintf(kernelSource, appropriateKernel, get_explicit_type_name( *outType ),
    get_explicit_type_name( *outType ), suffix, convert);

  programPtr = kernelSource;
  if( create_single_kernel_helper_with_build_options( context, &program, &kernel, 1,
    (const char **)&programPtr, "sample_test", "" ) )
  {
      return -1;
  }

  // Create an appropriately-sized output buffer.

  // Check to see if the output buffer will fit on the device
  size_t bytes = channelSize * 4 * width * height * depth;
  cl_ulong alloc_size = 0;

  cl_device_id device = NULL;
  error = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(device), &device, NULL);
  test_error( error, "Unable to query command queue for device" );

  error = clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(alloc_size), &alloc_size, NULL);
  test_error( error, "Unable to device for max mem alloc size" );

  if (bytes > alloc_size) {
    log_info("  Skipping: Buffer size (%lu) is greater than CL_DEVICE_MAX_MEM_ALLOC_SIZE (%lu)\n", bytes, alloc_size);
    *outSourceBuffer = NULL;
    return 0;
  }

  *outSourceBuffer = CreateRandomData(*outType, width * height * depth * 4, d);

  inStream = clCreateBuffer( context, CL_MEM_COPY_HOST_PTR,
    channelSize * 4 * width * height * depth, *outSourceBuffer, &error );
  test_error( error, "Unable to create output buffer" );

  clSamplerWrapper sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_NONE, CL_FILTER_NEAREST, &error);
  test_error( error, "Unable to create sampler" );

  error = clSetKernelArg( kernel, 0, sizeof( inStream ), &inStream );
  test_error( error, "Unable to set kernel arguments" );

  error = clSetKernelArg( kernel, 1, sizeof( clImage ), &clImage );
  test_error( error, "Unable to set kernel arguments" );

  // Flush and Acquire.

  glFinish();

  error = (*clEnqueueAcquireGLObjects_ptr)( queue, 1, &clImage, 0, NULL, NULL);
  test_error( error, "Unable to acquire GL obejcts");

  // Execute ( letting OpenCL choose the local size )

  // Setup the global dimensions and sizes based on the target type.
  set_dimensions_by_target(target, &global_dims, global_sizes,
    width, height, depth);

  error = clEnqueueNDRangeKernel( queue, kernel, global_dims, NULL,
    global_sizes, NULL, 0, NULL, NULL );
  test_error( error, "Unable to execute test kernel" );

  clEventWrapper event;
  error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &clImage, 0, NULL, &event );
  test_error(error, "clEnqueueReleaseGLObjects failed");

  error = clWaitForEvents( 1, &event );
  test_error(error, "clWaitForEvents failed");

  return 0;
}

static int test_image_write( cl_context context, cl_command_queue queue,
  GLenum glTarget, GLuint glTexture, size_t width, size_t height, size_t depth,
  cl_image_format *outFormat, ExplicitType *outType, void **outSourceBuffer,
  MTdata d, bool supports_half )
{
  int error;

  // Create a CL image from the supplied GL texture
  clMemWrapper image = (*clCreateFromGLTexture_ptr)( context, CL_MEM_WRITE_ONLY,
    glTarget, 0, glTexture, &error );

  if ( error != CL_SUCCESS ) {
    print_error( error, "Unable to create CL image from GL texture" );
    GLint fmt;
    glGetTexLevelParameteriv( glTarget, 0, GL_TEXTURE_INTERNAL_FORMAT, &fmt );
    log_error( "    Supplied GL texture was base format %s and internal "
      "format %s\n", GetGLBaseFormatName( fmt ), GetGLFormatName( fmt ) );
    return error;
  }

  return test_cl_image_write( context, queue, glTarget, image,
    width, height, depth, outFormat, outType, outSourceBuffer, d, supports_half );
}

int supportsHalf(cl_context context, bool* supports_half)
{
  int error;
  size_t  size;
  cl_uint numDev;

  error = clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES, sizeof(cl_uint), &numDev, NULL);
  test_error(error, "clGetContextInfo for CL_CONTEXT_NUM_DEVICES failed");

  cl_device_id* devices = new cl_device_id[numDev];
  error = clGetContextInfo(context, CL_CONTEXT_DEVICES, numDev * sizeof(cl_device_id), devices, NULL);
  test_error(error, "clGetContextInfo for CL_CONTEXT_DEVICES failed");

  *supports_half = is_extension_available(devices[0], "cl_khr_fp16");
  delete [] devices;

  return error;
}

int supportsMsaa(cl_context context, bool* supports_msaa)
{
  int error;
  size_t  size;
  cl_uint numDev;

  error = clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES, sizeof(cl_uint), &numDev, NULL);
  test_error(error, "clGetContextInfo for CL_CONTEXT_NUM_DEVICES failed");

  cl_device_id* devices = new cl_device_id[numDev];
  error = clGetContextInfo(context, CL_CONTEXT_DEVICES, numDev * sizeof(cl_device_id), devices, NULL);
  test_error(error, "clGetContextInfo for CL_CONTEXT_DEVICES failed");

  *supports_msaa = is_extension_available(devices[0], "cl_khr_gl_msaa_sharing");
  delete [] devices;

  return error;
}

int supportsDepth(cl_context context, bool* supports_depth)
{
  int error;
  size_t  size;
  cl_uint numDev;

  error = clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES, sizeof(cl_uint), &numDev, NULL);
  test_error(error, "clGetContextInfo for CL_CONTEXT_NUM_DEVICES failed");

  cl_device_id* devices = new cl_device_id[numDev];
  error = clGetContextInfo(context, CL_CONTEXT_DEVICES, numDev * sizeof(cl_device_id), devices, NULL);
  test_error(error, "clGetContextInfo for CL_CONTEXT_DEVICES failed");

  *supports_depth = is_extension_available(devices[0], "cl_khr_gl_depth_images");
  delete [] devices;

  return error;
}

static int test_image_format_write( cl_context context, cl_command_queue queue,
  size_t width, size_t height, size_t depth, GLenum target, GLenum format,
  GLenum internalFormat,  GLenum glType, ExplicitType type, MTdata d )
{
  int error;
  int samples = 8;
  // If we're testing a half float format, then we need to determine the
  // rounding mode of this machine.  Punt if we fail to do so.

  if( type == kHalf )
    if( DetectFloatToHalfRoundingMode(queue) )
      return 1;

  // Create an appropriate GL texture or renderbuffer, given the target.

  glTextureWrapper glTexture;
  glBufferWrapper glBuf;
  glFramebufferWrapper glFramebuffer;
  glRenderbufferWrapper glRenderbuffer;
  switch (get_base_gl_target(target)) {
    case GL_TEXTURE_1D:
      CreateGLTexture1D( width, target, format, internalFormat, glType,
        type, &glTexture, &error, false, d );
      break;
    case GL_TEXTURE_BUFFER:
      CreateGLTextureBuffer( width, target, format, internalFormat, glType,
        type, &glTexture, &glBuf, &error, false, d );
      break;
    case GL_TEXTURE_1D_ARRAY:
      CreateGLTexture1DArray( width, height, target, format, internalFormat,
        glType, type, &glTexture, &error, false, d );
      break;
    case GL_TEXTURE_RECTANGLE_EXT:
    case GL_TEXTURE_2D:
    case GL_TEXTURE_CUBE_MAP:
      CreateGLTexture2D( width, height, target, format, internalFormat, glType,
        type, &glTexture, &error, false, d );
      break;
    case GL_COLOR_ATTACHMENT0:
    case GL_RENDERBUFFER:
      CreateGLRenderbuffer(width, height, target, format, internalFormat,
        glType, type, &glFramebuffer, &glRenderbuffer, &error, d, false);
    case GL_TEXTURE_2D_ARRAY:
      CreateGLTexture2DArray( width, height, depth, target, format,
        internalFormat, glType, type, &glTexture, &error, false, d );
      break;
    case GL_TEXTURE_3D:
      CreateGLTexture3D( width, height, depth, target, format,
        internalFormat, glType, type, &glTexture, &error, d, false );
      break;

    default:
      log_error("Unsupported GL tex target (%s) passed to write test: "
        "%s (%s):%d", GetGLTargetName(target), __FUNCTION__,
        __FILE__, __LINE__);
  }

  // If there was a problem during creation, make sure it isn't a known
  // cause, and then complain.
  if ( error == -2 ) {
    log_info("OpenGL texture couldn't be created, because a texture is too big. Skipping test.\n");
    return 0;
  }

  if ( error != 0 ) {
    if ((format == GL_RGBA_INTEGER_EXT) && (!CheckGLIntegerExtensionSupport())){
      log_info("OpenGL version does not support GL_RGBA_INTEGER_EXT. "
        "Skipping test.\n");
      return 0;
    } else {
      return error;
    }
  }

  // Run and get the results
  cl_image_format clFormat;
  ExplicitType sourceType;
  ExplicitType validationType;
  void *outSourceBuffer = NULL;

  GLenum globj = glTexture;
  if (target == GL_RENDERBUFFER || target == GL_COLOR_ATTACHMENT0) {
    globj = glRenderbuffer;
  }

  bool supports_half = false;
  error = supportsHalf(context, &supports_half);
  if( error != 0 )
    return error;

  error = test_image_write( context, queue, target, globj, width, height,
    depth, &clFormat, &sourceType, (void **)&outSourceBuffer, d, supports_half );

  if( error != 0 || ((sourceType == kHalf ) && !supports_half)) {
    if (outSourceBuffer)
      free(outSourceBuffer);
    return error;
  }

  if (!outSourceBuffer)
    return 0;

  // If actual source type was half, convert to float for validation.

  if ( sourceType == kHalf )
    validationType = kFloat;
  else
    validationType = sourceType;

  BufferOwningPtr<char> validationSource;

  if ( clFormat.image_channel_data_type == CL_UNORM_INT_101010 )
  {
    validationSource.reset( outSourceBuffer );
  }
  else
  {
    validationSource.reset( convert_to_expected( outSourceBuffer,
      width * height * depth, sourceType, validationType, get_channel_order_channel_count(clFormat.image_channel_order) ) );
    free(outSourceBuffer);
  }

  log_info( "- Write for %s [%4ld x %4ld x %4ld] : GL Texture : %s : %s : %s =>"
    " CL Image : %s : %s \n",
    GetGLTargetName(target),
    width, height, depth,
    GetGLFormatName( format ),
    GetGLFormatName( internalFormat ),
    GetGLTypeName( glType),
    GetChannelOrderName( clFormat.image_channel_order ),
    GetChannelTypeName( clFormat.image_channel_data_type ));

  // Read the results from the GL texture.

  ExplicitType readType = type;
  BufferOwningPtr<char> glResults( ReadGLTexture(
    target, glTexture, glBuf, width, format,
    internalFormat, glType, readType, /* unused */ 1, 1 ) );
  if( glResults == NULL )
    return -1;

  // We have to convert our input buffer to the returned type, so we can validate.
  BufferOwningPtr<char> convertedGLResults;
  if ( clFormat.image_channel_data_type != CL_UNORM_INT_101010 )
  {
    convertedGLResults.reset( convert_to_expected(
      glResults, width * height * depth, readType, validationType, get_channel_order_channel_count(clFormat.image_channel_order), glType ));
  }

  // Validate.

  int valid = 0;
  if (convertedGLResults) {
    if( sourceType == kFloat || sourceType == kHalf )
    {
      if ( clFormat.image_channel_data_type == CL_UNORM_INT_101010 )
      {
        valid = validate_float_results_rgb_101010( validationSource, glResults, width, height, depth, 1 );
      }
      else
      {
        valid = validate_float_results( validationSource, convertedGLResults,
          width, height, depth, 1, get_channel_order_channel_count(clFormat.image_channel_order) );
      }
    }
    else
    {
      valid = validate_integer_results( validationSource, convertedGLResults,
        width, height, depth, 1, get_explicit_type_size( readType ) );
    }
  }

  return valid;
}

#pragma mark -
#pragma mark Write test common entry point

// This is the main loop for all of the write tests.  It iterates over the
// given formats & targets, testing a variety of sizes against each
// combination.

int test_images_write_common(cl_device_id device, cl_context context,
  cl_command_queue queue, struct format* formats, size_t nformats,
  GLenum *targets, size_t ntargets, sizevec_t* sizes, size_t nsizes )
{
  int err = 0;
  int error = 0;
  RandomSeed seed(gRandomSeed);

  // First, ensure this device supports images.

  if (checkForImageSupport(device)) {
    log_info("Device does not support images.  Skipping test.\n");
    return 0;
  }

  // Get the value of CL_DEVICE_MAX_MEM_ALLOC_SIZE
  cl_ulong max_individual_allocation_size = 0;
  err = clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
                        sizeof(max_individual_allocation_size),
                        &max_individual_allocation_size, NULL);
  if (err) {
    log_error("ERROR: clGetDeviceInfo failed for CL_DEVICE_MAX_MEM_ALLOC_SIZE.\n");
    error++;
    return error;
  }

  size_t total_allocation_size;
  size_t fidx, tidx, sidx;

  for ( fidx = 0; fidx < nformats; fidx++ ) {
    for ( tidx = 0; tidx < ntargets; tidx++ ) {

      // Texture buffer only takes an internal format, so the level data passed
      // by the test and used for verification must match the internal format
      if ((targets[tidx] == GL_TEXTURE_BUFFER) && (GetGLFormat(formats[ fidx ].internal) != formats[fidx].formattype))
        continue;

      if ( formats[ fidx ].datatype == GL_UNSIGNED_INT_2_10_10_10_REV )
      {
        // Check if the RGB 101010 format is supported
        if ( is_rgb_101010_supported( context, targets[ tidx ] ) == 0 )
          continue; // skip
      }

      if (formats[ fidx ].datatype == GL_UNSIGNED_INT_24_8)
      {
        //check if a implementation supports writing to the depth stencil formats
        cl_image_format imageFormat = { CL_DEPTH_STENCIL, CL_UNORM_INT24 };
        if (!is_image_format_supported(context, CL_MEM_WRITE_ONLY, (targets[tidx] == GL_TEXTURE_2D || targets[tidx] == GL_TEXTURE_RECTANGLE) ? CL_MEM_OBJECT_IMAGE2D: CL_MEM_OBJECT_IMAGE2D_ARRAY, &imageFormat))
          continue;
      }

      if (formats[ fidx ].datatype == GL_FLOAT_32_UNSIGNED_INT_24_8_REV)
      {
        //check if a implementation supports writing to the depth stencil formats
        cl_image_format imageFormat = { CL_DEPTH_STENCIL, CL_FLOAT};
        if (!is_image_format_supported(context, CL_MEM_WRITE_ONLY, (targets[tidx] == GL_TEXTURE_2D || targets[tidx] == GL_TEXTURE_RECTANGLE) ? CL_MEM_OBJECT_IMAGE2D: CL_MEM_OBJECT_IMAGE2D_ARRAY, &imageFormat))
          continue;
      }

      if (targets[tidx] != GL_TEXTURE_BUFFER)
        log_info( "Testing image write for GL format %s : %s : %s : %s\n",
                 GetGLTargetName( targets[ tidx ] ),
                 GetGLFormatName( formats[ fidx ].internal ),
                 GetGLBaseFormatName( formats[ fidx ].formattype ),
                 GetGLTypeName( formats[ fidx ].datatype ) );
      else
        log_info( "Testing image write for GL format %s : %s\n",
                 GetGLTargetName( targets[ tidx ] ),
                 GetGLFormatName( formats[ fidx ].internal ));


      for (sidx = 0; sidx < nsizes; sidx++) {

        // All tested formats are 4-channel formats
        total_allocation_size =
           sizes[sidx].width * sizes[sidx].height * sizes[sidx].depth *
           4 * get_explicit_type_size( formats[ fidx ].type );

        if (total_allocation_size > max_individual_allocation_size) {
          log_info( "The requested allocation size (%gMB) is larger than the "
                    "maximum individual allocation size (%gMB)\n",
                    total_allocation_size/(1024.0*1024.0),
                    max_individual_allocation_size/(1024.0*1024.0));
          log_info( "Skipping write test for %s : %s : %s : %s "
                    " and size (%ld, %ld, %ld)\n",
                    GetGLTargetName( targets[ tidx ] ),
                    GetGLFormatName( formats[ fidx ].internal ),
                    GetGLBaseFormatName( formats[ fidx ].formattype ),
                    GetGLTypeName( formats[ fidx ].datatype ),
                    sizes[sidx].width,
                    sizes[sidx].height,
                    sizes[sidx].depth);
          continue;
        }
#ifdef GL_VERSION_3_2
        if (get_base_gl_target(targets[ tidx ]) == GL_TEXTURE_2D_MULTISAMPLE ||
            get_base_gl_target(targets[ tidx ]) == GL_TEXTURE_2D_MULTISAMPLE_ARRAY)
        {
            bool supports_msaa;
            int errorInGetInfo = supportsMsaa(context, &supports_msaa);
            if (errorInGetInfo != 0) return errorInGetInfo;
            if (!supports_msaa) return 0;
        }
        if (formats[ fidx ].formattype == GL_DEPTH_COMPONENT ||
            formats[ fidx ].formattype == GL_DEPTH_STENCIL)
        {
            bool supports_depth;
            int errorInGetInfo = supportsDepth(context, &supports_depth);
            if (errorInGetInfo != 0) return errorInGetInfo;
            if (!supports_depth) return 0;
        }
#endif

        if( test_image_format_write( context, queue,
                                     sizes[sidx].width,
                                     sizes[sidx].height,
                                     sizes[sidx].depth,
                                     targets[ tidx ],
                                     formats[ fidx ].formattype,
                                     formats[ fidx ].internal,
                                     formats[ fidx ].datatype,
                                     formats[ fidx ].type, seed ) )
        {
          log_error( "ERROR: Image write test failed for %s : %s : %s : %s "
            " and size (%ld, %ld, %ld)\n\n",
            GetGLTargetName( targets[ tidx ] ),
            GetGLFormatName( formats[ fidx ].internal ),
            GetGLBaseFormatName( formats[ fidx ].formattype ),
            GetGLTypeName( formats[ fidx ].datatype ),
            sizes[sidx].width,
            sizes[sidx].height,
            sizes[sidx].depth);

          error++;
          break;    // Skip other sizes for this combination
        }
      }

      // If we passed all sizes (check versus size loop count):

      if (sidx == nsizes) {
        log_info( "passed: Image write for GL format  %s : %s : %s : %s\n\n",
          GetGLTargetName( targets[ tidx ] ),
          GetGLFormatName( formats[ fidx ].internal ),
          GetGLBaseFormatName( formats[ fidx ].formattype ),
          GetGLTypeName( formats[ fidx ].datatype ) );
      }
    }
  }

  return error;
}