🏡 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, 2020 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 "harness/compat.h"

// Bug: Missing in spec: atomic_intptr_t is always supported if device is 32-bits.
// Bug: Missing in spec: CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE

#define FLUSH fflush(stdout)

#define MAX_STR 16*1024

#define ALIGNMENT 128

#define OPTIONS "-cl-std=CL2.0"

// NUM_ROUNDS must be at least 1.
// It determines how many sets of random data we push through the global
// variables.
#define NUM_ROUNDS 1

// This is a shared property of the writer and reader kernels.
#define NUM_TESTED_VALUES 5

// TODO: pointer-to-half (and its vectors)
// TODO: union of...

#include <algorithm>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <string>
#include <vector>
#include <cassert>
#include <sys/types.h>
#include <sys/stat.h>
#include "harness/typeWrappers.h"
#include "harness/errorHelpers.h"
#include "harness/mt19937.h"
#include "procs.h"


////////////////////
// Device capabilities
static int l_has_double = 0;
static int l_has_half = 0;
static int l_64bit_device = 0;
static int l_has_int64_atomics = 0;
static int l_has_intptr_atomics = 0;
static int l_has_cles_int64 = 0;

static int l_host_is_big_endian = 1;

static size_t l_max_global_id0 = 0;
static cl_bool l_linker_available = false;

#define check_error(errCode,msg,...) ((errCode != CL_SUCCESS) ? (log_error("ERROR: " msg "! (%s:%d)\n", ## __VA_ARGS__, __FILE__, __LINE__), 1) : 0)

////////////////////
// Info about types we can use for program scope variables.


class TypeInfo {

public:
    TypeInfo() :
        name(""),
        m_buf_elem_type(""),
        m_is_vecbase(false),
        m_is_atomic(false),
        m_is_like_size_t(false),
        m_is_bool(false),
        m_elem_type(0), m_num_elem(0),
        m_size(0),
        m_value_size(0)
        {}
    TypeInfo(const char* name_arg) :
        name(name_arg),
        m_buf_elem_type(name_arg),
        m_is_vecbase(false),
        m_is_atomic(false),
        m_is_like_size_t(false),
        m_is_bool(false),
        m_elem_type(0), m_num_elem(0),
        m_size(0),
        m_value_size(0)
        { }

    // Vectors
    TypeInfo( TypeInfo* elem_type, int num_elem ) :
        m_is_vecbase(false),
        m_is_atomic(false),
        m_is_like_size_t(false),
        m_is_bool(false),
        m_elem_type(elem_type),
        m_num_elem(num_elem)
    {
        char the_name[10]; // long enough for longest vector type name "double16"
        snprintf(the_name,sizeof(the_name),"%s%d",elem_type->get_name_c_str(),m_num_elem);
        this->name = std::string(the_name);
        this->m_buf_elem_type = std::string(the_name);
        this->m_value_size = num_elem * elem_type->get_size();
        if ( m_num_elem == 3 ) {
            this->m_size = 4 * elem_type->get_size();
        } else {
            this->m_size = num_elem * elem_type->get_size();
        }
    }
    const std::string& get_name(void) const { return name; }
    const char* get_name_c_str(void) const { return name.c_str(); }
    TypeInfo& set_vecbase(void) { this->m_is_vecbase = true; return *this; }
    TypeInfo& set_atomic(void) { this->m_is_atomic = true; return *this; }
    TypeInfo& set_like_size_t(void) {
        this->m_is_like_size_t = true;
        this->set_size( l_64bit_device ? 8 : 4 );
        this->m_buf_elem_type = l_64bit_device ? "ulong" : "uint";
        return *this;
    }
    TypeInfo& set_bool(void) { this->m_is_bool = true; return *this; }
    TypeInfo& set_size(size_t n) { this->m_value_size = this->m_size = n; return *this; }
    TypeInfo& set_buf_elem_type( const char* name ) { this->m_buf_elem_type = std::string(name); return *this; }

    const TypeInfo* elem_type(void) const { return m_elem_type; }
    int num_elem(void) const { return m_num_elem; }

    bool is_vecbase(void) const {return m_is_vecbase;}
    bool is_atomic(void) const {return m_is_atomic;}
    bool is_atomic_64bit(void) const {return m_is_atomic && m_size == 8;}
    bool is_like_size_t(void) const {return m_is_like_size_t;}
    bool is_bool(void) const {return m_is_bool;}
    size_t get_size(void) const {return m_size;}
    size_t get_value_size(void) const {return m_value_size;}

    // When passing values of this type to a kernel, what buffer type
    // should be used?
    const char* get_buf_elem_type(void) const { return m_buf_elem_type.c_str(); }

    std::string as_string(const cl_uchar* value_ptr) const {
        // This method would be shorter if I had a real handle to element
        // vector type.
        if ( this->is_bool() ) {
            std::string result( name );
            result += "<";
            result += (*value_ptr ? "true" : "false");
            result += ", ";
            char buf[10];
            sprintf(buf,"%02x",*value_ptr);
            result += buf;
            result += ">";
            return result;
        } else if ( this->num_elem() ) {
            std::string result( name );
            result += "<";
            for ( unsigned ielem = 0 ; ielem < this->num_elem() ; ielem++ ) {
                char buf[MAX_STR];
                if ( ielem ) result += ", ";
                for ( unsigned ibyte = 0; ibyte < this->m_elem_type->get_size() ; ibyte++ ) {
                    sprintf(buf + 2*ibyte,"%02x", value_ptr[ ielem * this->m_elem_type->get_size() + ibyte ] );
                }
                result += buf;
            }
            result += ">";
            return result;
        } else {
            std::string result( name );
            result += "<";
            char buf[MAX_STR];
            for ( unsigned ibyte = 0; ibyte < this->get_size() ; ibyte++ ) {
                sprintf(buf + 2*ibyte,"%02x", value_ptr[ ibyte ] );
            }
            result += buf;
            result += ">";
            return result;
        }
    }

    // Initialize the given buffer to a constant value initialized as if it
    // were from the INIT_VAR macro below.
    // Only needs to support values 0 and 1.
    void init( cl_uchar* buf, cl_uchar val) const {
        if ( this->num_elem() ) {
            for ( unsigned ielem = 0 ; ielem < this->num_elem() ; ielem++ ) {
                // Delegate!
                this->init_elem( buf + ielem * this->get_value_size()/this->num_elem(), val );
            }
        } else {
            init_elem( buf, val );
        }
    }

private:
    void init_elem( cl_uchar* buf, cl_uchar val ) const {
        size_t elem_size = this->num_elem() ? this->get_value_size()/this->num_elem() : this->get_size();
        memset(buf,0,elem_size);
        if ( val ) {
            if ( strstr( name.c_str(), "float" ) ) {
                *(float*)buf = (float)val;
                return;
            }
            if ( strstr( name.c_str(), "double" ) ) {
                *(double*)buf = (double)val;
                return;
            }
            if ( this->is_bool() ) { *buf = (bool)val; return; }

            // Write a single character value to the correct spot,
            // depending on host endianness.
            if ( l_host_is_big_endian ) *(buf + elem_size-1) = (cl_uchar)val;
            else *buf = (cl_uchar)val;
        }
    }
public:

    void dump(FILE* fp) const {
        fprintf(fp,"Type %s : <%d,%d,%s> ", name.c_str(),
                (int)m_size,
                (int)m_value_size,
                m_buf_elem_type.c_str() );
        if ( this->m_elem_type ) fprintf(fp, " vec(%s,%d)", this->m_elem_type->get_name_c_str(), this->num_elem() );
        if ( this->m_is_vecbase ) fprintf(fp, " vecbase");
        if ( this->m_is_bool ) fprintf(fp, " bool");
        if ( this->m_is_like_size_t ) fprintf(fp, " like-size_t");
        if ( this->m_is_atomic ) fprintf(fp, " atomic");
        fprintf(fp,"\n");
        fflush(fp);
    }

private:
    std::string name;
    TypeInfo* m_elem_type;
    int m_num_elem;
    bool m_is_vecbase;
    bool m_is_atomic;
    bool m_is_like_size_t;
    bool m_is_bool;
    size_t m_size; // Number of bytes of storage occupied by this type.
    size_t m_value_size; // Number of bytes of value significant for this type. Differs for vec3.

    // When passing values of this type to a kernel, what buffer type
    // should be used?
    // For most types, it's just itself.
    // Use a std::string so I don't have to make a copy constructor.
    std::string m_buf_elem_type;
};


#define NUM_SCALAR_TYPES (8+2) // signed and unsigned integral types, float and double
#define NUM_VECTOR_SIZES (5)   // 2,3,4,8,16
#define NUM_PLAIN_TYPES \
      5 /*boolean and size_t family */  \
    + NUM_SCALAR_TYPES \
    + NUM_SCALAR_TYPES*NUM_VECTOR_SIZES \
    + 10 /* atomic types */

// Need room for plain, array, pointer, struct
#define MAX_TYPES (4*NUM_PLAIN_TYPES)

static TypeInfo type_info[MAX_TYPES];
static int num_type_info = 0; // Number of valid entries in type_info[]




// A helper class to form kernel source arguments for clCreateProgramWithSource.
class StringTable {
public:
    StringTable() : m_c_strs(NULL), m_lengths(NULL), m_frozen(false), m_strings() {}
    ~StringTable() { release_frozen(); }

    void add(std::string s) { release_frozen(); m_strings.push_back(s); }

    const size_t num_str() { freeze(); return m_strings.size(); }
    const char** strs() { freeze(); return m_c_strs; }
    const size_t* lengths() { freeze(); return m_lengths; }

private:
    void freeze(void) {
        if ( !m_frozen ) {
            release_frozen();

            m_c_strs = (const char**) malloc(sizeof(const char*) * m_strings.size());
            m_lengths = (size_t*) malloc(sizeof(size_t) * m_strings.size());
            assert( m_c_strs );
            assert( m_lengths );

            for ( size_t i = 0; i < m_strings.size() ; i++ ) {
                m_c_strs[i] = m_strings[i].c_str();
                m_lengths[i] = strlen(m_c_strs[i]);
            }

            m_frozen = true;
        }
    }
    void release_frozen(void) {
        if ( m_c_strs ) { free(m_c_strs); m_c_strs = 0; }
        if ( m_lengths ) { free(m_lengths); m_lengths = 0; }
        m_frozen = false;
    }

    typedef std::vector<std::string> strlist_t;
    strlist_t m_strings;
    const char** m_c_strs;
    size_t* m_lengths;
    bool m_frozen;
};


////////////////////
// File scope function declarations

static void l_load_abilities(cl_device_id device);
static const char* l_get_fp64_pragma(void);
static const char* l_get_cles_int64_pragma(void);
static int l_build_type_table(cl_device_id device);

static int l_get_device_info(cl_device_id device, size_t* max_size_ret, size_t* pref_size_ret);

static void l_set_randomly( cl_uchar* buf, size_t buf_size, RandomSeed& rand_state );
static int l_compare( const cl_uchar* expected, const cl_uchar* received, unsigned num_values, const TypeInfo&ti );
static int l_copy( cl_uchar* dest, unsigned dest_idx, const cl_uchar* src, unsigned src_idx, const TypeInfo&ti );

static std::string conversion_functions(const TypeInfo& ti);
static std::string global_decls(const TypeInfo& ti, bool with_init);
static std::string global_check_function(const TypeInfo& ti);
static std::string writer_function(const TypeInfo& ti);
static std::string reader_function(const TypeInfo& ti);

static int l_write_read( cl_device_id device, cl_context context, cl_command_queue queue );
static int l_write_read_for_type( cl_device_id device, cl_context context, cl_command_queue queue, const TypeInfo& ti, RandomSeed& rand_state );

static int l_init_write_read( cl_device_id device, cl_context context, cl_command_queue queue );
static int l_init_write_read_for_type( cl_device_id device, cl_context context, cl_command_queue queue, const TypeInfo& ti, RandomSeed& rand_state );

static int l_capacity( cl_device_id device, cl_context context, cl_command_queue queue, size_t max_size );
static int l_user_type( cl_device_id device, cl_context context, cl_command_queue queue, size_t max_size, bool separate_compilation );



////////////////////
// File scope function definitions

static cl_int print_build_log(cl_program program, cl_uint num_devices, cl_device_id *device_list, cl_uint count, const char **strings, const size_t *lengths, const char* options)
{
    cl_uint i;
    cl_int error;
    BufferOwningPtr<cl_device_id> devices;

    if(num_devices == 0 || device_list == NULL)
    {
        error = clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, sizeof(num_devices), &num_devices, NULL);
        test_error(error, "clGetProgramInfo CL_PROGRAM_NUM_DEVICES failed");

        device_list = (cl_device_id*)malloc(sizeof(cl_device_id)*num_devices);
        devices.reset(device_list);

        memset(device_list, 0, sizeof(cl_device_id) * num_devices);

        error = clGetProgramInfo(program, CL_PROGRAM_DEVICES, sizeof(cl_device_id) * num_devices, device_list, NULL);
        test_error(error, "clGetProgramInfo CL_PROGRAM_DEVICES failed");
    }

    cl_uint z;
    bool sourcePrinted = false;

    for(z = 0; z < num_devices; z++)
    {
        char deviceName[4096] = "";
        error = clGetDeviceInfo(device_list[z], CL_DEVICE_NAME, sizeof(deviceName), deviceName, NULL);
        check_error(error, "Device \"%d\" failed to return a name. clGetDeviceInfo CL_DEVICE_NAME failed", z);

        cl_build_status buildStatus;
        error = clGetProgramBuildInfo(program, device_list[z], CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL);
        check_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_STATUS failed");

        if(buildStatus != CL_BUILD_SUCCESS)
        {
            if(!sourcePrinted)
            {
                log_error("Build options: %s\n", options);
                if(count && strings)
                {
                    log_error("Original source is: ------------\n");
                    for(i = 0; i < count; i++) log_error("%s", strings[i]);
                }
                sourcePrinted = true;
            }

            char statusString[64] = "";
            if (buildStatus == (cl_build_status)CL_BUILD_SUCCESS)
              sprintf(statusString, "CL_BUILD_SUCCESS");
            else if (buildStatus == (cl_build_status)CL_BUILD_NONE)
              sprintf(statusString, "CL_BUILD_NONE");
            else if (buildStatus == (cl_build_status)CL_BUILD_ERROR)
              sprintf(statusString, "CL_BUILD_ERROR");
            else if (buildStatus == (cl_build_status)CL_BUILD_IN_PROGRESS)
              sprintf(statusString, "CL_BUILD_IN_PROGRESS");
            else
              sprintf(statusString, "UNKNOWN (%d)", buildStatus);

            log_error("Build not successful for device \"%s\", status: %s\n", deviceName, statusString);

            size_t paramSize = 0;
            error = clGetProgramBuildInfo(program, device_list[z], CL_PROGRAM_BUILD_LOG, 0, NULL, &paramSize);
            if(check_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_LOG failed")) break;

            std::string log;
            log.resize(paramSize/sizeof(char));

            error = clGetProgramBuildInfo(program, device_list[z], CL_PROGRAM_BUILD_LOG, paramSize, &log[0], NULL);
            if(check_error(error, "Device %d (%s) failed to return a build log", z, deviceName)) break;
            if(log[0] == 0) log_error("clGetProgramBuildInfo returned an empty log.\n");
            else
            {
                log_error("Build log:\n", deviceName);
                log_error("%s\n", log.c_str());
            }
        }
    }
    return error;
}

static void l_load_abilities(cl_device_id device)
{
    l_has_half       = is_extension_available(device,"cl_khr_fp16");
    l_has_double     = is_extension_available(device,"cl_khr_fp64");
    l_has_cles_int64 = is_extension_available(device,"cles_khr_int64");

    l_has_int64_atomics
    =  is_extension_available(device,"cl_khr_int64_base_atomics")
    && is_extension_available(device,"cl_khr_int64_extended_atomics");

    {
        int status = CL_SUCCESS;
        cl_uint addr_bits = 32;
        status = clGetDeviceInfo(device,CL_DEVICE_ADDRESS_BITS,sizeof(addr_bits),&addr_bits,0);
        l_64bit_device = ( status == CL_SUCCESS && addr_bits == 64 );
    }

    // 32-bit devices always have intptr atomics.
    l_has_intptr_atomics = !l_64bit_device || l_has_int64_atomics;

    union { char c[4]; int i; } probe;
    probe.i = 1;
    l_host_is_big_endian = !probe.c[0];

    // Determine max global id.
    {
        int status = CL_SUCCESS;
        cl_uint max_dim = 0;
        status = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,sizeof(max_dim),&max_dim,0);
        assert( status == CL_SUCCESS );
        assert( max_dim > 0 );
        size_t max_id[3];
        max_id[0] = 0;
    status = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_ITEM_SIZES,max_dim*sizeof(size_t),&max_id[0],0);
        assert( status == CL_SUCCESS );
        l_max_global_id0 = max_id[0];
    }

    { // Is separate compilation supported?
        int status = CL_SUCCESS;
        l_linker_available = false;
        status = clGetDeviceInfo(device,CL_DEVICE_LINKER_AVAILABLE,sizeof(l_linker_available),&l_linker_available,0);
        assert( status == CL_SUCCESS );
    }
}


static const char* l_get_fp64_pragma(void)
{
    return l_has_double ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" : "";
}

static const char* l_get_cles_int64_pragma(void)
{
    return l_has_cles_int64 ? "#pragma OPENCL EXTENSION cles_khr_int64 : enable\n" : "";
}

static const char* l_get_int64_atomic_pragma(void)
{
    return "#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable\n"
           "#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable\n";
}

static int l_build_type_table(cl_device_id device)
{
    int status = CL_SUCCESS;
    size_t iscalar = 0;
    size_t ivecsize = 0;
    int vecsizes[] = { 2, 3, 4, 8, 16 };
    const char* vecbase[] = {
        "uchar", "char",
        "ushort", "short",
        "uint", "int",
        "ulong", "long",
        "float",
        "double"
    };
    int vecbase_size[] = {
        1, 1,
        2, 2,
        4, 4,
        8, 8,
        4,
        8
    };
    const char* like_size_t[] = {
        "intptr_t",
        "uintptr_t",
        "size_t",
        "ptrdiff_t"
    };
    const char* atomics[] = {
        "atomic_int", "atomic_uint",
        "atomic_long", "atomic_ulong",
        "atomic_float",
        "atomic_double",
    };
    int atomics_size[] = {
        4, 4,
        8, 8,
        4,
        8
    };
    const char* intptr_atomics[] = {
        "atomic_intptr_t",
        "atomic_uintptr_t",
        "atomic_size_t",
        "atomic_ptrdiff_t"
    };

    l_load_abilities(device);
    num_type_info = 0;

    // Boolean.
    type_info[ num_type_info++ ] = TypeInfo( "bool" ).set_bool().set_size(1).set_buf_elem_type("uchar");

    // Vector types, and the related scalar element types.
    for ( iscalar=0; iscalar < sizeof(vecbase)/sizeof(vecbase[0]) ; ++iscalar ) {
        if ( !gHasLong && strstr(vecbase[iscalar],"long") ) continue;
        if ( !l_has_double && strstr(vecbase[iscalar],"double") ) continue;

        // Scalar
        TypeInfo* elem_type = type_info + num_type_info++;
        *elem_type = TypeInfo( vecbase[iscalar] ).set_vecbase().set_size( vecbase_size[iscalar] );

        // Vector
        for ( ivecsize=0; ivecsize < sizeof(vecsizes)/sizeof(vecsizes[0]) ; ivecsize++ ) {
            type_info[ num_type_info++ ] = TypeInfo( elem_type, vecsizes[ivecsize] );
        }
    }

    // Size_t-like types
    for ( iscalar=0; iscalar < sizeof(like_size_t)/sizeof(like_size_t[0]) ; ++iscalar ) {
        type_info[ num_type_info++ ] = TypeInfo( like_size_t[iscalar] ).set_like_size_t();
    }

    // Atomic types.
    for ( iscalar=0; iscalar < sizeof(atomics)/sizeof(atomics[0]) ; ++iscalar ) {
        if ( !l_has_int64_atomics && strstr(atomics[iscalar],"long") ) continue;
        if ( !(l_has_int64_atomics && l_has_double) && strstr(atomics[iscalar],"double") ) continue;

        // The +7 is used to skip over the "atomic_" prefix.
        const char* buf_type = atomics[iscalar] + 7;
        type_info[ num_type_info++ ] = TypeInfo( atomics[iscalar] ).set_atomic().set_size( atomics_size[iscalar] ).set_buf_elem_type( buf_type );
    }
    if ( l_has_intptr_atomics ) {
        for ( iscalar=0; iscalar < sizeof(intptr_atomics)/sizeof(intptr_atomics[0]) ; ++iscalar ) {
            type_info[ num_type_info++ ] = TypeInfo( intptr_atomics[iscalar] ).set_atomic().set_like_size_t();
        }
    }

    assert( num_type_info <= MAX_TYPES ); // or increase MAX_TYPES

#if 0
    for ( size_t i = 0 ; i < num_type_info ; i++ ) {
        type_info[ i ].dump(stdout);
    }
    exit(0);
#endif

    return status;
}

static const TypeInfo& l_find_type( const char* name )
{
    auto itr =
        std::find_if(type_info, type_info + num_type_info,
                     [name](TypeInfo& ti) { return ti.get_name() == name; });
    assert(itr != type_info + num_type_info);
    return *itr;
}



// Populate return parameters for max program variable size, preferred program variable size.

static int l_get_device_info(cl_device_id device, size_t* max_size_ret, size_t* pref_size_ret)
{
    int err = CL_SUCCESS;
    size_t return_size = 0;

    err = clGetDeviceInfo(device, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE, sizeof(*max_size_ret), max_size_ret, &return_size);
    if ( err != CL_SUCCESS ) {
        log_error("Error: Failed to get device info for CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE\n");
        return err;
    }
    if ( return_size != sizeof(size_t) ) {
        log_error("Error: Invalid size %d returned for CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE\n", (int)return_size );
        return 1;
    }
    if ( return_size != sizeof(size_t) ) {
        log_error("Error: Invalid size %d returned for CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE\n", (int)return_size );
        return 1;
    }

    return_size = 0;
    err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE, sizeof(*pref_size_ret), pref_size_ret, &return_size);
    if ( err != CL_SUCCESS ) {
        log_error("Error: Failed to get device info for CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE: %d\n",err);
        return err;
    }
    if ( return_size != sizeof(size_t) ) {
        log_error("Error: Invalid size %d returned for CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE\n", (int)return_size );
        return 1;
    }

    return CL_SUCCESS;
}


static void l_set_randomly( cl_uchar* buf, size_t buf_size, RandomSeed& rand_state )
{
    assert( 0 == (buf_size % sizeof(cl_uint) ) );
    for ( size_t i = 0; i < buf_size ; i += sizeof(cl_uint) ) {
        *( (cl_uint*)(buf + i) ) = genrand_int32( rand_state );
    }
#if 0
    for ( size_t i = 0; i < buf_size ; i++ ) {
        printf("%02x",buf[i]);
    }
    printf("\n");
#endif
}

// Return num_value values of the given type.
// Returns CL_SUCCESS if they compared as equal.
static int l_compare( const char* test_name, const cl_uchar* expected, const cl_uchar* received, size_t num_values, const TypeInfo&ti )
{
    // Compare only the valid returned bytes.
    for ( unsigned value_idx = 0; value_idx < num_values; value_idx++ ) {
        const cl_uchar* expv = expected + value_idx * ti.get_size();
        const cl_uchar* gotv = received + value_idx * ti.get_size();
        if ( memcmp( expv, gotv, ti.get_value_size() ) ) {
            std::string exp_str = ti.as_string( expv );
            std::string got_str = ti.as_string( gotv );
            log_error("Error: %s test for type %s, at index %d: Expected %s got %s\n",
                    test_name,
                    ti.get_name_c_str(), value_idx,
                    exp_str.c_str(),
                    got_str.c_str() );
            return 1;
        }
    }
    return CL_SUCCESS;
}

// Copy a target value from src[idx] to dest[idx]
static int l_copy( cl_uchar* dest, unsigned dest_idx, const cl_uchar* src, unsigned src_idx, const TypeInfo&ti )
{
    cl_uchar* raw_dest      = dest + dest_idx * ti.get_size();
    const cl_uchar* raw_src =  src +  src_idx * ti.get_size();
    memcpy( raw_dest, raw_src,  ti.get_value_size() );

    return 0;
}


static std::string conversion_functions(const TypeInfo& ti)
{
    std::string result;
    static char buf[MAX_STR];
    int num_printed = 0;
    // The atomic types just use the base type.
    if ( ti.is_atomic() || 0 == strcmp( ti.get_buf_elem_type(), ti.get_name_c_str() ) ) {
        // The type is represented in a buffer by itself.
        num_printed = snprintf(buf,MAX_STR,
                "%s from_buf(%s a) { return a; }\n"
                "%s to_buf(%s a) { return a; }\n",
                ti.get_buf_elem_type(), ti.get_buf_elem_type(),
                ti.get_buf_elem_type(), ti.get_buf_elem_type() );
    } else {
        // Just use C-style cast.
        num_printed = snprintf(buf,MAX_STR,
                "%s from_buf(%s a) { return (%s)a; }\n"
                "%s to_buf(%s a) { return (%s)a; }\n",
                ti.get_name_c_str(), ti.get_buf_elem_type(), ti.get_name_c_str(),
                ti.get_buf_elem_type(), ti.get_name_c_str(), ti.get_buf_elem_type() );
    }
    // Add initializations.
    if ( ti.is_atomic() ) {
        num_printed += snprintf( buf + num_printed, MAX_STR-num_printed,
                "#define INIT_VAR(a) ATOMIC_VAR_INIT(a)\n" );
    } else {
        // This cast works even if the target type is a vector type.
        num_printed += snprintf( buf + num_printed, MAX_STR-num_printed,
                "#define INIT_VAR(a) ((%s)(a))\n", ti.get_name_c_str());
    }
    assert( num_printed < MAX_STR ); // or increase MAX_STR
    result = buf;
    return result;
}

static std::string global_decls(const TypeInfo& ti, bool with_init )
{
    const char* tn = ti.get_name_c_str();
    const char* vol = (ti.is_atomic() ? " volatile " : " ");
    static char decls[MAX_STR];
    int num_printed = 0;
    if ( with_init ) {
        const char *decls_template_with_init =
            "%s %s var = INIT_VAR(0);\n"
            "global %s %s g_var = INIT_VAR(1);\n"
            "%s %s a_var[2] = { INIT_VAR(1), INIT_VAR(1) };\n"
            "volatile global %s %s* p_var = &a_var[1];\n\n";
        num_printed = snprintf(decls,sizeof(decls),decls_template_with_init,
                vol,tn,vol,tn,vol,tn,vol,tn);
    } else {
        const char *decls_template_no_init =
            "%s %s var;\n"
            "global %s %s g_var;\n"
            "%s %s a_var[2];\n"
            "global %s %s* p_var;\n\n";
        num_printed = snprintf(decls,sizeof(decls),decls_template_no_init,
             vol,tn,vol,tn,vol,tn,vol,tn);
    }
    assert( num_printed < sizeof(decls) );
    return std::string(decls);
}

// Return the source code for the "global_check" function for the given type.
// This function checks that all program-scope variables have appropriate
// initial values when no explicit initializer is used. If all tests pass the
// kernel writes a non-zero value to its output argument, otherwise it writes
// zero.
static std::string global_check_function(const TypeInfo& ti)
{
    const std::string type_name = ti.get_buf_elem_type();

    // all() should only be used on vector inputs. For scalar comparison, the
    // result of the equality operator can be used as a bool value.
    const bool is_scalar = ti.num_elem() == 0; // 0 is used to represent scalar types, not 1.
    const std::string is_equality_true = is_scalar ? "" : "all";

    std::string code = "kernel void global_check(global int* out) {\n";
    code += "  const " + type_name + " zero = ((" + type_name + ")0);\n";
    code += "  bool status = true;\n";
    if (ti.is_atomic()) {
        code += "  status &= " + is_equality_true + "(atomic_load(&var) == zero);\n";
        code += "  status &= " + is_equality_true + "(atomic_load(&g_var) == zero);\n";
        code += "  status &= " + is_equality_true + "(atomic_load(&a_var[0]) == zero);\n";
        code += "  status &= " + is_equality_true + "(atomic_load(&a_var[1]) == zero);\n";
    } else {
        code += "  status &= " + is_equality_true + "(var == zero);\n";
        code += "  status &= " + is_equality_true + "(g_var == zero);\n";
        code += "  status &= " + is_equality_true + "(a_var[0] == zero);\n";
        code += "  status &= " + is_equality_true + "(a_var[1] == zero);\n";
    }
    code += "  status &= (p_var == NULL);\n";
    code += "  *out = status ? 1 : 0;\n";
    code += "}\n\n";

    return code;
}

// Return the source text for the writer function for the given type.
// For types that can't be passed as pointer-to-type as a kernel argument,
// use a substitute base type of the same size.
static std::string writer_function(const TypeInfo& ti)
{
    static char writer_src[MAX_STR];
    int num_printed = 0;
    if ( !ti.is_atomic() ) {
        const char* writer_template_normal =
            "kernel void writer( global %s* src, uint idx ) {\n"
            "  var = from_buf(src[0]);\n"
            "  g_var = from_buf(src[1]);\n"
            "  a_var[0] = from_buf(src[2]);\n"
            "  a_var[1] = from_buf(src[3]);\n"
            "  p_var = a_var + idx;\n"
            "}\n\n";
        num_printed = snprintf(writer_src,sizeof(writer_src),writer_template_normal,ti.get_buf_elem_type());
    } else {
        const char* writer_template_atomic =
            "kernel void writer( global %s* src, uint idx ) {\n"
            "  atomic_store( &var, from_buf(src[0]) );\n"
            "  atomic_store( &g_var, from_buf(src[1]) );\n"
            "  atomic_store( &a_var[0], from_buf(src[2]) );\n"
            "  atomic_store( &a_var[1], from_buf(src[3]) );\n"
            "  p_var = a_var + idx;\n"
            "}\n\n";
        num_printed = snprintf(writer_src,sizeof(writer_src),writer_template_atomic,ti.get_buf_elem_type());
    }
    assert( num_printed < sizeof(writer_src) );
    std::string result = writer_src;
    return result;
}


// Return source text for teh reader function for the given type.
// For types that can't be passed as pointer-to-type as a kernel argument,
// use a substitute base type of the same size.
static std::string reader_function(const TypeInfo& ti)
{
    static char reader_src[MAX_STR];
    int num_printed = 0;
    if ( !ti.is_atomic() ) {
        const char* reader_template_normal =
            "kernel void reader( global %s* dest, %s ptr_write_val ) {\n"
            "  *p_var = from_buf(ptr_write_val);\n"
            "  dest[0] = to_buf(var);\n"
            "  dest[1] = to_buf(g_var);\n"
            "  dest[2] = to_buf(a_var[0]);\n"
            "  dest[3] = to_buf(a_var[1]);\n"
            "}\n\n";
        num_printed = snprintf(reader_src,sizeof(reader_src),reader_template_normal,ti.get_buf_elem_type(),ti.get_buf_elem_type());
    } else {
        const char* reader_template_atomic =
            "kernel void reader( global %s* dest, %s ptr_write_val ) {\n"
            "  atomic_store( p_var, from_buf(ptr_write_val) );\n"
            "  dest[0] = to_buf( atomic_load( &var ) );\n"
            "  dest[1] = to_buf( atomic_load( &g_var ) );\n"
            "  dest[2] = to_buf( atomic_load( &a_var[0] ) );\n"
            "  dest[3] = to_buf( atomic_load( &a_var[1] ) );\n"
            "}\n\n";
        num_printed = snprintf(reader_src,sizeof(reader_src),reader_template_atomic,ti.get_buf_elem_type(),ti.get_buf_elem_type());
    }
    assert( num_printed < sizeof(reader_src) );
    std::string result = reader_src;
    return result;
}

// Check that all globals where appropriately default-initialized.
static int check_global_initialization(cl_context context, cl_program program, cl_command_queue queue)
{
    int status = CL_SUCCESS;

    // Create a buffer on device to store a unique integer.
    cl_int is_init_valid = 0;
    clMemWrapper buffer(clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(is_init_valid), &is_init_valid, &status));
    test_error_ret(status, "Failed to allocate buffer", status);

    // Create, setup and invoke kernel.
    clKernelWrapper global_check(clCreateKernel(program, "global_check", &status));
    test_error_ret(status, "Failed to create global_check kernel", status);
    status = clSetKernelArg(global_check, 0, sizeof(cl_mem), &buffer);
    test_error_ret(status, "Failed to set up argument for the global_check kernel", status);
    const cl_uint work_dim = 1;
    const size_t global_work_offset[] = { 0 };
    const size_t global_work_size[] = { 1 };
    status = clEnqueueNDRangeKernel(queue, global_check, work_dim, global_work_offset, global_work_size, nullptr, 0, nullptr, nullptr);
    test_error_ret(status, "Failed to run global_check kernel", status);
    status = clFinish(queue);
    test_error_ret(status, "clFinish() failed", status);

    // Read back the memory buffer from the device.
    status = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, sizeof(is_init_valid), &is_init_valid, 0, nullptr, nullptr);
    test_error_ret(status, "Failed to read buffer from device", status);
    if (is_init_valid == 0) {
        log_error("Unexpected default values were detected");
        return 1;
    }

    return CL_SUCCESS;
}

// Check write-then-read.
static int l_write_read( cl_device_id device, cl_context context, cl_command_queue queue )
{
    int status = CL_SUCCESS;
    int itype;

    RandomSeed rand_state( gRandomSeed );

    for ( itype = 0; itype < num_type_info ; itype++ ) {
        status = status | l_write_read_for_type(device,context,queue,type_info[itype], rand_state );
        FLUSH;
    }

    return status;
}

static int l_write_read_for_type( cl_device_id device, cl_context context, cl_command_queue queue, const TypeInfo& ti, RandomSeed& rand_state )
{
    int err = CL_SUCCESS;
    std::string type_name( ti.get_name() );
    const char* tn = type_name.c_str();
    log_info("  %s ",tn);

    StringTable ksrc;
    ksrc.add( l_get_fp64_pragma() );
    ksrc.add( l_get_cles_int64_pragma() );
    if (ti.is_atomic_64bit())
      ksrc.add( l_get_int64_atomic_pragma() );
    ksrc.add( conversion_functions(ti) );
    ksrc.add( global_decls(ti,false) );
    ksrc.add( global_check_function(ti) );
    ksrc.add( writer_function(ti) );
    ksrc.add( reader_function(ti) );

    int status = CL_SUCCESS;
    clProgramWrapper program;
    clKernelWrapper writer;

    status = create_single_kernel_helper_with_build_options(context, &program, &writer, ksrc.num_str(), ksrc.strs(), "writer", OPTIONS);
    test_error_ret(status,"Failed to create program for read-after-write test",status);

    clKernelWrapper reader( clCreateKernel( program, "reader", &status ) );
    test_error_ret(status,"Failed to create reader kernel for read-after-write test",status);

    // Check size query.
    size_t used_bytes = 0;
    status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
    test_error_ret(status,"Failed to query global variable total size",status);
    size_t expected_used_bytes =
        (NUM_TESTED_VALUES-1)*ti.get_size() // Two regular variables and an array of 2 elements.
        + ( l_64bit_device ? 8 : 4 ); // The pointer
    if ( used_bytes < expected_used_bytes ) {
        log_error("Error program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)expected_used_bytes, (unsigned long long)used_bytes );
        err |= 1;
    }

    err |= check_global_initialization(context, program, queue);

    // We need to create 5 random values of the given type,
    // and read 4 of them back.
    const size_t write_data_size = NUM_TESTED_VALUES * sizeof(cl_ulong16);
    const size_t read_data_size = (NUM_TESTED_VALUES - 1) * sizeof(cl_ulong16);
    cl_uchar* write_data = (cl_uchar*)align_malloc(write_data_size, ALIGNMENT);
    cl_uchar* read_data = (cl_uchar*)align_malloc(read_data_size, ALIGNMENT);

    clMemWrapper write_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, write_data_size, write_data, &status ) );
    test_error_ret(status,"Failed to allocate write buffer",status);
    clMemWrapper read_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, read_data_size, read_data, &status ) );
    test_error_ret(status,"Failed to allocate read buffer",status);

    status = clSetKernelArg(writer,0,sizeof(cl_mem),&write_mem); test_error_ret(status,"set arg",status);
    status = clSetKernelArg(reader,0,sizeof(cl_mem),&read_mem); test_error_ret(status,"set arg",status);

    // Boolean random data needs to be massaged a bit more.
    const int num_rounds = ti.is_bool() ? (1 << NUM_TESTED_VALUES ) : NUM_ROUNDS;
    unsigned bool_iter = 0;

    for ( int iround = 0; iround < num_rounds ; iround++ ) {
        for ( cl_uint iptr_idx = 0; iptr_idx < 2 ; iptr_idx++ ) { // Index into array, to write via pointer
            // Generate new random data to push through.
            // Generate 5 * 128 bytes all the time, even though the test for many types use less than all that.

            cl_uchar *write_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, write_mem, CL_TRUE, CL_MAP_WRITE, 0, write_data_size, 0, 0, 0, 0);

            if ( ti.is_bool() ) {
                // For boolean, random data cast to bool isn't very random.
                // So use the bottom bit of bool_value_iter to get true
                // diversity.
                for ( unsigned value_idx = 0; value_idx < NUM_TESTED_VALUES ; value_idx++ ) {
                    write_data[value_idx] = (1<<value_idx) & bool_iter;
                    //printf(" %s", (write_data[value_idx] ? "true" : "false" ));
                }
                bool_iter++;
            } else {
                l_set_randomly( write_data, write_data_size, rand_state );
            }
            status = clSetKernelArg(writer,1,sizeof(cl_uint),&iptr_idx); test_error_ret(status,"set arg",status);

            // The value to write via the pointer should be taken from the
            // 5th typed slot of the write_data.
            status = clSetKernelArg(reader,1,ti.get_size(),write_data + (NUM_TESTED_VALUES-1)*ti.get_size()); test_error_ret(status,"set arg",status);

            // Determine the expected values.
            cl_uchar expected[read_data_size];
            memset( expected, -1, sizeof(expected) );
            l_copy( expected, 0, write_data, 0, ti );
            l_copy( expected, 1, write_data, 1, ti );
            l_copy( expected, 2, write_data, 2, ti );
            l_copy( expected, 3, write_data, 3, ti );
            // But we need to take into account the value from the pointer write.
            // The 2 represents where the "a" array values begin in our read-back.
            l_copy( expected, 2 + iptr_idx, write_data, 4, ti );

            clEnqueueUnmapMemObject(queue, write_mem, write_ptr, 0, 0, 0);

            if ( ti.is_bool() ) {
                // Collapse down to one bit.
                for ( unsigned i = 0; i <  NUM_TESTED_VALUES-1 ; i++ ) expected[i] = (bool)expected[i];
            }

            cl_uchar *read_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, read_mem, CL_TRUE, CL_MAP_READ, 0, read_data_size, 0, 0, 0, 0);
            memset(read_data, -1, read_data_size);
            clEnqueueUnmapMemObject(queue, read_mem, read_ptr, 0, 0, 0);

            // Now run the kernel
            const size_t one = 1;
            status = clEnqueueNDRangeKernel(queue,writer,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue writer",status);
            status = clEnqueueNDRangeKernel(queue,reader,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue reader",status);
            status = clFinish(queue); test_error_ret(status,"finish",status);

            read_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, read_mem, CL_TRUE, CL_MAP_READ, 0, read_data_size, 0, 0, 0, 0);

            if ( ti.is_bool() ) {
                // Collapse down to one bit.
                for ( unsigned i = 0; i <  NUM_TESTED_VALUES-1 ; i++ ) read_data[i] = (bool)read_data[i];
            }

            // Compare only the valid returned bytes.
            int compare_result = l_compare( "read-after-write", expected, read_data, NUM_TESTED_VALUES-1, ti );
            // log_info("Compared %d values each of size %llu. Result %d\n", NUM_TESTED_VALUES-1, (unsigned long long)ti.get_value_size(), compare_result );
            err |= compare_result;

            clEnqueueUnmapMemObject(queue, read_mem, read_ptr, 0, 0, 0);

            if ( err ) break;
        }
    }

    if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
    align_free(write_data);
    align_free(read_data);
    return err;
}


// Check initialization, then, read, then write, then read.
static int l_init_write_read( cl_device_id device, cl_context context, cl_command_queue queue )
{
    int status = CL_SUCCESS;
    int itype;

    RandomSeed rand_state( gRandomSeed );

    for ( itype = 0; itype < num_type_info ; itype++ ) {
        status = status | l_init_write_read_for_type(device,context,queue,type_info[itype], rand_state );
    }
    return status;
}
static int l_init_write_read_for_type( cl_device_id device, cl_context context, cl_command_queue queue, const TypeInfo& ti, RandomSeed& rand_state )
{
    int err = CL_SUCCESS;
    std::string type_name( ti.get_name() );
    const char* tn = type_name.c_str();
    log_info("  %s ",tn);

    StringTable ksrc;
    ksrc.add( l_get_fp64_pragma() );
    ksrc.add( l_get_cles_int64_pragma() );
    if (ti.is_atomic_64bit())
      ksrc.add( l_get_int64_atomic_pragma() );
    ksrc.add( conversion_functions(ti) );
    ksrc.add( global_decls(ti,true) );
    ksrc.add( writer_function(ti) );
    ksrc.add( reader_function(ti) );

    int status = CL_SUCCESS;
    clProgramWrapper program;
    clKernelWrapper writer;

    status = create_single_kernel_helper_with_build_options(context, &program, &writer, ksrc.num_str(), ksrc.strs(), "writer", OPTIONS);
    test_error_ret(status,"Failed to create program for init-read-after-write test",status);

    clKernelWrapper reader( clCreateKernel( program, "reader", &status ) );
    test_error_ret(status,"Failed to create reader kernel for init-read-after-write test",status);

    // Check size query.
    size_t used_bytes = 0;
    status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
    test_error_ret(status,"Failed to query global variable total size",status);
    size_t expected_used_bytes =
        (NUM_TESTED_VALUES-1)*ti.get_size() // Two regular variables and an array of 2 elements.
        + ( l_64bit_device ? 8 : 4 ); // The pointer
    if ( used_bytes < expected_used_bytes ) {
        log_error("Error: program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)expected_used_bytes, (unsigned long long)used_bytes );
        err |= 1;
    }

    // We need to create 5 random values of the given type,
    // and read 4 of them back.
    const size_t write_data_size = NUM_TESTED_VALUES * sizeof(cl_ulong16);
    const size_t read_data_size = (NUM_TESTED_VALUES-1) * sizeof(cl_ulong16);

    cl_uchar* write_data = (cl_uchar*)align_malloc(write_data_size, ALIGNMENT);
    cl_uchar* read_data = (cl_uchar*)align_malloc(read_data_size, ALIGNMENT);
    clMemWrapper write_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, write_data_size, write_data, &status ) );
    test_error_ret(status,"Failed to allocate write buffer",status);
    clMemWrapper read_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, read_data_size, read_data, &status ) );
    test_error_ret(status,"Failed to allocate read buffer",status);

    status = clSetKernelArg(writer,0,sizeof(cl_mem),&write_mem); test_error_ret(status,"set arg",status);
    status = clSetKernelArg(reader,0,sizeof(cl_mem),&read_mem); test_error_ret(status,"set arg",status);

    // Boolean random data needs to be massaged a bit more.
    const int num_rounds = ti.is_bool() ? (1 << NUM_TESTED_VALUES ) : NUM_ROUNDS;
    unsigned bool_iter = 0;

    // We need to count iterations.  We do something *different on the
    // first iteration, to ensure we actually pick up the initialized
    // values.
    unsigned iteration = 0;

    for ( int iround = 0; iround < num_rounds ; iround++ ) {
        for ( cl_uint iptr_idx = 0; iptr_idx < 2 ; iptr_idx++ ) { // Index into array, to write via pointer
            // Generate new random data to push through.
            // Generate 5 * 128 bytes all the time, even though the test for many types use less than all that.

            cl_uchar *write_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, write_mem, CL_TRUE, CL_MAP_WRITE, 0, write_data_size, 0, 0, 0, 0);

            if ( ti.is_bool() ) {
                // For boolean, random data cast to bool isn't very random.
                // So use the bottom bit of bool_value_iter to get true
                // diversity.
                for ( unsigned value_idx = 0; value_idx < NUM_TESTED_VALUES ; value_idx++ ) {
                    write_data[value_idx] = (1<<value_idx) & bool_iter;
                    //printf(" %s", (write_data[value_idx] ? "true" : "false" ));
                }
                bool_iter++;
            } else {
                l_set_randomly( write_data, write_data_size, rand_state );
            }
            status = clSetKernelArg(writer,1,sizeof(cl_uint),&iptr_idx); test_error_ret(status,"set arg",status);

            if ( !iteration ) {
                // On first iteration, the value we write via the last arg
                // to the "reader" function is 0.
                // It's way easier to code the test this way.
                ti.init( write_data + (NUM_TESTED_VALUES-1)*ti.get_size(), 0 );
            }

            // The value to write via the pointer should be taken from the
            // 5th typed slot of the write_data.
            status = clSetKernelArg(reader,1,ti.get_size(),write_data + (NUM_TESTED_VALUES-1)*ti.get_size()); test_error_ret(status,"set arg",status);

            // Determine the expected values.
            cl_uchar expected[read_data_size];
            memset( expected, -1, sizeof(expected) );
            if ( iteration ) {
                l_copy( expected, 0, write_data, 0, ti );
                l_copy( expected, 1, write_data, 1, ti );
                l_copy( expected, 2, write_data, 2, ti );
                l_copy( expected, 3, write_data, 3, ti );
                // But we need to take into account the value from the pointer write.
                // The 2 represents where the "a" array values begin in our read-back.
                // But we need to take into account the value from the pointer write.
                l_copy( expected, 2 + iptr_idx, write_data, 4, ti );
            } else {
                // On first iteration, expect these initialized values!
                // See the decls_template_with_init above.
                ti.init( expected, 0 );
                ti.init( expected + ti.get_size(), 1 );
                ti.init( expected + 2*ti.get_size(), 1 );
                // Emulate the effect of the write via the pointer.
                // The value is 0, not 1 (see above).
                // The pointer is always initialized to the second element
                // of the array. So it goes into slot 3 of the "expected" array.
                ti.init( expected + 3*ti.get_size(), 0 );
            }

            if ( ti.is_bool() ) {
                // Collapse down to one bit.
                for ( unsigned i = 0; i <  NUM_TESTED_VALUES-1 ; i++ ) expected[i] = (bool)expected[i];
            }

            clEnqueueUnmapMemObject(queue, write_mem, write_ptr, 0, 0, 0);

            cl_uchar *read_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, read_mem, CL_TRUE, CL_MAP_READ, 0, read_data_size, 0, 0, 0, 0);
            memset( read_data, -1, read_data_size );
            clEnqueueUnmapMemObject(queue, read_mem, read_ptr, 0, 0, 0);

            // Now run the kernel
            const size_t one = 1;
            if ( iteration ) {
                status = clEnqueueNDRangeKernel(queue,writer,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue writer",status);
            } else {
                // On first iteration, we should be picking up the
                // initialized value. So don't enqueue the writer.
            }
            status = clEnqueueNDRangeKernel(queue,reader,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue reader",status);
            status = clFinish(queue); test_error_ret(status,"finish",status);

            read_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, read_mem, CL_TRUE, CL_MAP_READ, 0, read_data_size, 0, 0, 0, 0);

            if ( ti.is_bool() ) {
                // Collapse down to one bit.
                for ( unsigned i = 0; i <  NUM_TESTED_VALUES-1 ; i++ ) read_data[i] = (bool)read_data[i];
            }

            // Compare only the valid returned bytes.
            //log_info(" Round %d ptr_idx %u\n", iround, iptr_idx );
            int compare_result = l_compare( "init-write-read", expected, read_data, NUM_TESTED_VALUES-1, ti );
            //log_info("Compared %d values each of size %llu. Result %d\n", NUM_TESTED_VALUES-1, (unsigned long long)ti.get_value_size(), compare_result );
            err |= compare_result;

            clEnqueueUnmapMemObject(queue, read_mem, read_ptr, 0, 0, 0);

            if ( err ) break;

            iteration++;
        }
    }

    if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
    align_free(write_data);
    align_free(read_data);

    return err;
}


// Check that we can make at least one variable with size
// max_size which is returned from the device info property : CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE.
static int l_capacity( cl_device_id device, cl_context context, cl_command_queue queue, size_t max_size )
{
    int err = CL_SUCCESS;
    // Just test one type.
    const TypeInfo ti( l_find_type("uchar") );
    log_info(" l_capacity...");

    const char prog_src_template[] =
#if defined(_WIN32)
        "uchar var[%Iu];\n\n"
#else
        "uchar var[%zu];\n\n"
#endif
        "kernel void get_max_size( global ulong* size_ret ) {\n"
#if defined(_WIN32)
        "  *size_ret = (ulong)%Iu;\n"
#else
        "  *size_ret = (ulong)%zu;\n"
#endif
        "}\n\n"
        "kernel void writer( global uchar* src ) {\n"
        "  var[get_global_id(0)] = src[get_global_linear_id()];\n"
        "}\n\n"
        "kernel void reader( global uchar* dest ) {\n"
        "  dest[get_global_linear_id()] = var[get_global_id(0)];\n"
        "}\n\n";
    char prog_src[MAX_STR];
    int num_printed = snprintf(prog_src,sizeof(prog_src),prog_src_template,max_size, max_size);
    assert( num_printed < MAX_STR ); // or increase MAX_STR

    StringTable ksrc;
    ksrc.add( prog_src );

    int status = CL_SUCCESS;
    clProgramWrapper program;
    clKernelWrapper get_max_size;

    status = create_single_kernel_helper_with_build_options(context, &program, &get_max_size, ksrc.num_str(), ksrc.strs(), "get_max_size", OPTIONS);
    test_error_ret(status,"Failed to create program for capacity test",status);

    // Check size query.
    size_t used_bytes = 0;
    status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
    test_error_ret(status,"Failed to query global variable total size",status);
    if ( used_bytes < max_size ) {
        log_error("Error: program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)max_size, (unsigned long long)used_bytes );
        err |= 1;
    }

    // Prepare to execute
    clKernelWrapper writer( clCreateKernel( program, "writer", &status ) );
    test_error_ret(status,"Failed to create writer kernel for capacity test",status);
    clKernelWrapper reader( clCreateKernel( program, "reader", &status ) );
    test_error_ret(status,"Failed to create reader kernel for capacity test",status);

    cl_ulong max_size_ret = 0;
    const size_t arr_size = 10*1024*1024;
    cl_uchar* buffer = (cl_uchar*) align_malloc( arr_size, ALIGNMENT );

    if ( !buffer ) { log_error("Failed to allocate buffer\n"); return 1; }

    clMemWrapper max_size_ret_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, sizeof(max_size_ret), &max_size_ret, &status ) );
    test_error_ret(status,"Failed to allocate size query buffer",status);
    clMemWrapper buffer_mem( clCreateBuffer( context, CL_MEM_READ_WRITE, arr_size, 0, &status ) );
    test_error_ret(status,"Failed to allocate write buffer",status);

    status = clSetKernelArg(get_max_size,0,sizeof(cl_mem),&max_size_ret_mem); test_error_ret(status,"set arg",status);
    status = clSetKernelArg(writer,0,sizeof(cl_mem),&buffer_mem); test_error_ret(status,"set arg",status);
    status = clSetKernelArg(reader,0,sizeof(cl_mem),&buffer_mem); test_error_ret(status,"set arg",status);

    // Check the macro value of CL_DEVICE_MAX_GLOBAL_VARIABLE
    const size_t one = 1;
    status = clEnqueueNDRangeKernel(queue,get_max_size,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue size query",status);
    status = clFinish(queue); test_error_ret(status,"finish",status);

    cl_uchar *max_size_ret_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, max_size_ret_mem, CL_TRUE, CL_MAP_READ, 0, sizeof(max_size_ret), 0, 0, 0, 0);
    if ( max_size_ret != max_size ) {
        log_error("Error: preprocessor definition for CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE is %llu and does not match device query value %llu\n",
                (unsigned long long) max_size_ret,
                (unsigned long long) max_size );
        err |= 1;
    }
    clEnqueueUnmapMemObject(queue, max_size_ret_mem, max_size_ret_ptr, 0, 0, 0);

    RandomSeed rand_state_write( gRandomSeed );
    for ( size_t offset = 0; offset < max_size ; offset += arr_size ) {
        size_t curr_size = (max_size - offset) < arr_size ? (max_size - offset) : arr_size;
        l_set_randomly( buffer, curr_size, rand_state_write );
        status = clEnqueueWriteBuffer (queue, buffer_mem, CL_TRUE, 0, curr_size, buffer, 0, 0, 0);test_error_ret(status,"populate buffer_mem object",status);
        status = clEnqueueNDRangeKernel(queue,writer,1,&offset,&curr_size,0,0,0,0); test_error_ret(status,"enqueue writer",status);
    status = clFinish(queue); test_error_ret(status,"finish",status);
    }

    RandomSeed rand_state_read( gRandomSeed );
    for ( size_t offset = 0; offset < max_size ; offset += arr_size ) {
        size_t curr_size = (max_size - offset) < arr_size ? (max_size - offset) : arr_size;
        status = clEnqueueNDRangeKernel(queue,reader,1,&offset,&curr_size,0,0,0,0); test_error_ret(status,"enqueue reader",status);
        cl_uchar* read_mem_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, buffer_mem, CL_TRUE, CL_MAP_READ, 0, curr_size, 0, 0, 0, &status);test_error_ret(status,"map read data",status);
        l_set_randomly( buffer, curr_size, rand_state_read );
        err |= l_compare( "capacity", buffer, read_mem_ptr, curr_size, ti );
        clEnqueueUnmapMemObject(queue, buffer_mem, read_mem_ptr, 0, 0, 0);
    }

    if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
    align_free(buffer);

    return err;
}


// Check operation on a user type.
static int l_user_type( cl_device_id device, cl_context context, cl_command_queue queue, bool separate_compile )
{
    int err = CL_SUCCESS;
    // Just test one type.
    const TypeInfo ti( l_find_type("uchar") );
    log_info(" l_user_type %s...", separate_compile ? "separate compilation" : "single source compilation" );

    if ( separate_compile && ! l_linker_available ) {
        log_info("Separate compilation is not supported. Skipping test\n");
        return err;
    }

    const char type_src[] =
        "typedef struct { uchar c; uint i; } my_struct_t;\n\n";
    const char def_src[] =
        "my_struct_t var = { 'a', 42 };\n\n";
    const char decl_src[] =
        "extern my_struct_t var;\n\n";

    // Don't use a host struct. We can't guarantee that the host
    // compiler has the same structure layout as the device compiler.
    const char writer_src[] =
        "kernel void writer( uchar c, uint i ) {\n"
        "  var.c = c;\n"
        "  var.i = i;\n"
        "}\n\n";
    const char reader_src[] =
        "kernel void reader( global uchar* C, global uint* I ) {\n"
        "  *C = var.c;\n"
        "  *I = var.i;\n"
        "}\n\n";

    clProgramWrapper program;

    if ( separate_compile ) {
        // Separate compilation flow.
        StringTable wksrc;
        wksrc.add( type_src );
        wksrc.add( def_src );
        wksrc.add( writer_src );

        StringTable rksrc;
        rksrc.add( type_src );
        rksrc.add( decl_src );
        rksrc.add( reader_src );

        int status = CL_SUCCESS;
        clProgramWrapper writer_program( clCreateProgramWithSource( context, wksrc.num_str(), wksrc.strs(), wksrc.lengths(), &status ) );
        test_error_ret(status,"Failed to create writer program for user type test",status);

        status = clCompileProgram( writer_program, 1, &device, OPTIONS, 0, 0, 0, 0, 0 );
        if(check_error(status, "Failed to compile writer program for user type test (%s)", IGetErrorString(status)))
        {
            print_build_log(writer_program, 1, &device, wksrc.num_str(), wksrc.strs(), wksrc.lengths(), OPTIONS);
            return status;
        }

        clProgramWrapper reader_program( clCreateProgramWithSource( context, rksrc.num_str(), rksrc.strs(), rksrc.lengths(), &status ) );
        test_error_ret(status,"Failed to create reader program for user type test",status);

        status = clCompileProgram( reader_program, 1, &device, OPTIONS, 0, 0, 0, 0, 0 );
        if(check_error(status, "Failed to compile reader program for user type test (%s)", IGetErrorString(status)))
        {
            print_build_log(reader_program, 1, &device, rksrc.num_str(), rksrc.strs(), rksrc.lengths(), OPTIONS);
            return status;
        }

        cl_program progs[2];
        progs[0] = writer_program;
        progs[1] = reader_program;

        program = clLinkProgram( context, 1, &device, "", 2, progs, 0, 0, &status );
        if(check_error(status, "Failed to link program for user type test (%s)", IGetErrorString(status)))
        {
            print_build_log(program, 1, &device, 0, NULL, NULL, "");
            return status;
        }
    } else {
        // Single compilation flow.
        StringTable ksrc;
        ksrc.add( type_src );
        ksrc.add( def_src );
        ksrc.add( writer_src );
        ksrc.add( reader_src );

        int status = CL_SUCCESS;

        status = create_single_kernel_helper_create_program(context, &program, ksrc.num_str(), ksrc.strs(), OPTIONS);
        if(check_error(status, "Failed to build program for user type test (%s)", IGetErrorString(status)))
        {
            print_build_log(program, 1, &device, ksrc.num_str(), ksrc.strs(), ksrc.lengths(), OPTIONS);
            return status;
        }

        status = clBuildProgram(program, 1, &device, OPTIONS, 0, 0);
        if(check_error(status, "Failed to compile program for user type test (%s)", IGetErrorString(status)))
        {
            print_build_log(program, 1, &device, ksrc.num_str(), ksrc.strs(), ksrc.lengths(), OPTIONS);
            return status;
        }
    }


    // Check size query.
    size_t used_bytes = 0;
    int status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
    test_error_ret(status,"Failed to query global variable total size",status);
    size_t expected_size = sizeof(cl_uchar) + sizeof(cl_uint);
    if ( used_bytes < expected_size ) {
        log_error("Error: program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)expected_size, (unsigned long long)used_bytes );
        err |= 1;
    }

    // Prepare to execute
    clKernelWrapper writer( clCreateKernel( program, "writer", &status ) );
    test_error_ret(status,"Failed to create writer kernel for user type test",status);
    clKernelWrapper reader( clCreateKernel( program, "reader", &status ) );
    test_error_ret(status,"Failed to create reader kernel for user type test",status);

    // Set up data.
    cl_uchar* uchar_data = (cl_uchar*)align_malloc(sizeof(cl_uchar), ALIGNMENT);
    cl_uint* uint_data = (cl_uint*)align_malloc(sizeof(cl_uint), ALIGNMENT);

    clMemWrapper uchar_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, sizeof(cl_uchar), uchar_data, &status ) );
    test_error_ret(status,"Failed to allocate uchar buffer",status);
    clMemWrapper uint_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, sizeof(cl_uint), uint_data, &status ) );
    test_error_ret(status,"Failed to allocate uint buffer",status);

    status = clSetKernelArg(reader,0,sizeof(cl_mem),&uchar_mem); test_error_ret(status,"set arg",status);
    status = clSetKernelArg(reader,1,sizeof(cl_mem),&uint_mem); test_error_ret(status,"set arg",status);

    cl_uchar expected_uchar = 'a';
    cl_uint expected_uint = 42;
    for ( unsigned iter = 0; iter < 5 ; iter++ ) { // Must go around at least twice
        // Read back data
        *uchar_data = -1;
        *uint_data = -1;
        const size_t one = 1;
        status = clEnqueueNDRangeKernel(queue,reader,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue reader",status);
        status = clFinish(queue); test_error_ret(status,"finish",status);

        cl_uchar *uint_data_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, uint_mem, CL_TRUE, CL_MAP_READ, 0, sizeof(cl_uint), 0, 0, 0, 0);
        cl_uchar *uchar_data_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, uchar_mem, CL_TRUE, CL_MAP_READ, 0, sizeof(cl_uchar), 0, 0, 0, 0);

        if ( expected_uchar != *uchar_data || expected_uint != *uint_data ) {
            log_error("FAILED: Iteration %d Got (0x%2x,%d) but expected (0x%2x,%d)\n",
                    iter, (int)*uchar_data, *uint_data, (int)expected_uchar, expected_uint );
            err |= 1;
        }

        clEnqueueUnmapMemObject(queue, uint_mem, uint_data_ptr, 0, 0, 0);
        clEnqueueUnmapMemObject(queue, uchar_mem, uchar_data_ptr, 0, 0, 0);

        // Mutate the data.
        expected_uchar++;
        expected_uint++;

        // Write the new values into persistent store.
        *uchar_data = expected_uchar;
        *uint_data = expected_uint;
        status = clSetKernelArg(writer,0,sizeof(cl_uchar),uchar_data); test_error_ret(status,"set arg",status);
        status = clSetKernelArg(writer,1,sizeof(cl_uint),uint_data); test_error_ret(status,"set arg",status);
        status = clEnqueueNDRangeKernel(queue,writer,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue writer",status);
        status = clFinish(queue); test_error_ret(status,"finish",status);
    }

    if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
    align_free(uchar_data);
    align_free(uint_data);
    return err;
}

// Determines whether its valid to skip this test based on the driver version
// and the features it optionally supports.
// Whether the test should be skipped is writen into the out paramter skip.
// The check returns an error code for the clDeviceInfo query.
static cl_int should_skip(cl_device_id device, cl_bool& skip)
{
    // Assume we can't skip to begin with.
    skip = CL_FALSE;

    // Progvar tests are already skipped for OpenCL < 2.0, so here we only need
    // to test for 3.0 since that is when program scope global variables become
    // optional.
    if (get_device_cl_version(device) >= Version(3, 0))
    {
        size_t max_global_variable_size{};
        test_error(clGetDeviceInfo(device, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE,
                                   sizeof(max_global_variable_size),
                                   &max_global_variable_size, nullptr),
                   "clGetDeviceInfo failed");
        skip = (max_global_variable_size != 0) ? CL_FALSE : CL_TRUE;
    }
    return CL_SUCCESS;
}

////////////////////
// Global functions


// Test support for variables at program scope. Miscellaneous
int test_progvar_prog_scope_misc(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
    cl_bool skip{ CL_FALSE };
    auto error = should_skip(device, skip);
    if (CL_SUCCESS != error)
    {
        return TEST_FAIL;
    }
    if (skip)
    {
        log_info("Skipping progvar_prog_scope_misc since it is optionally not "
                 "supported on this device\n");
        return TEST_SKIPPED_ITSELF;
    }
    size_t max_size = 0;
    size_t pref_size = 0;

    cl_int err = CL_SUCCESS;

    err = l_get_device_info( device, &max_size, &pref_size );
    err |= l_build_type_table( device );

    err |= l_capacity( device, context, queue, max_size );
    err |= l_user_type( device, context, queue, false );
    err |= l_user_type( device, context, queue, true );

    return err;
}


// Test support for variables at program scope. Unitialized data
int test_progvar_prog_scope_uninit(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
    cl_bool skip{ CL_FALSE };
    auto error = should_skip(device, skip);
    if (CL_SUCCESS != error)
    {
        return TEST_FAIL;
    }
    if (skip)
    {
        log_info(
            "Skipping progvar_prog_scope_uninit since it is optionally not "
            "supported on this device\n");
        return TEST_SKIPPED_ITSELF;
    }
    size_t max_size = 0;
    size_t pref_size = 0;

    cl_int err = CL_SUCCESS;

    err = l_get_device_info( device, &max_size, &pref_size );
    err |= l_build_type_table( device );

    err |= l_write_read( device, context, queue );

    return err;
}

// Test support for variables at program scope. Initialized data.
int test_progvar_prog_scope_init(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
    cl_bool skip{ CL_FALSE };
    auto error = should_skip(device, skip);
    if (CL_SUCCESS != error)
    {
        return TEST_FAIL;
    }
    if (skip)
    {
        log_info("Skipping progvar_prog_scope_init since it is optionally not "
                 "supported on this device\n");
        return TEST_SKIPPED_ITSELF;
    }
    size_t max_size = 0;
    size_t pref_size = 0;

    cl_int err = CL_SUCCESS;

    err = l_get_device_info( device, &max_size, &pref_size );
    err |= l_build_type_table( device );

    err |= l_init_write_read( device, context, queue );

    return err;
}


// A simple test for support of static variables inside a kernel.
int test_progvar_func_scope(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
    cl_bool skip{ CL_FALSE };
    auto error = should_skip(device, skip);
    if (CL_SUCCESS != error)
    {
        return TEST_FAIL;
    }
    if (skip)
    {
        log_info("Skipping progvar_func_scope since it is optionally not "
                 "supported on this device\n");
        return TEST_SKIPPED_ITSELF;
    }
    size_t max_size = 0;
    size_t pref_size = 0;

    cl_int err = CL_SUCCESS;

    // Deliberately have two variables with the same name but in different
    // scopes.
    // Also, use a large initialized structure in both cases.
    const char prog_src[] =
        "typedef struct { char c; int16 i; } mystruct_t;\n"
        "kernel void test_bump( global int* value, int which ) {\n"
        "  if ( which ) {\n"
        // Explicit address space.
        // Last element set to 0
        "     static global mystruct_t persistent = {'a',(int16)(0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,0) };\n"
        "     *value = persistent.i.sf++;\n"
        "  } else {\n"
        // Implicitly global
        // Last element set to 100
        "     static mystruct_t persistent = {'b',(int16)(0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,100) };\n"
        "     *value = persistent.i.sf++;\n"
        "  }\n"
        "}\n";

    StringTable ksrc;
    ksrc.add( prog_src );

    int status = CL_SUCCESS;
    clProgramWrapper program;
    clKernelWrapper test_bump;

    status = create_single_kernel_helper_with_build_options(context, &program, &test_bump, ksrc.num_str(), ksrc.strs(), "test_bump", OPTIONS);
    test_error_ret(status, "Failed to create program for function static variable test", status);

    // Check size query.
    size_t used_bytes = 0;
    status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
    test_error_ret(status,"Failed to query global variable total size",status);
    size_t expected_size = 2 * sizeof(cl_int); // Two ints.
    if ( used_bytes < expected_size ) {
        log_error("Error: program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)expected_size, (unsigned long long)used_bytes );
        err |= 1;
    }

    // Prepare the data.
    cl_int counter_value = 0;
    clMemWrapper counter_value_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, sizeof(counter_value), &counter_value, &status ) );
    test_error_ret(status,"Failed to allocate counter query buffer",status);

    status = clSetKernelArg(test_bump,0,sizeof(cl_mem),&counter_value_mem); test_error_ret(status,"set arg",status);

    // Go a few rounds, alternating between the two counters in the kernel.

    // Same as initial values in kernel.
    // But "true" which increments the 0-based counter, and "false" which
    // increments the 100-based counter.
    cl_int expected_counter[2] = { 100, 0 };

    const size_t one = 1;
    for ( int iround = 0; iround < 5 ; iround++ ) { // Must go at least twice around
        for ( int iwhich = 0; iwhich < 2 ; iwhich++ ) { // Cover both counters
            status = clSetKernelArg(test_bump,1,sizeof(iwhich),&iwhich); test_error_ret(status,"set arg",status);
            status = clEnqueueNDRangeKernel(queue,test_bump,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue test_bump",status);
            status = clFinish(queue); test_error_ret(status,"finish",status);

            cl_uchar *counter_value_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, counter_value_mem, CL_TRUE, CL_MAP_READ, 0, sizeof(counter_value), 0, 0, 0, 0);

            if ( counter_value != expected_counter[iwhich] ) {
                log_error("Error: Round %d on counter %d: Expected %d but got %d\n",
                        iround, iwhich, expected_counter[iwhich], counter_value );
                err |= 1;
            }
            expected_counter[iwhich]++; // Emulate behaviour of the kernel.

            clEnqueueUnmapMemObject(queue, counter_value_mem, counter_value_ptr, 0, 0, 0);
        }
    }

    if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }

    return err;
}