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

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

#include "common.h"
#include "host_atomics.h"

#include <sstream>
#include <vector>

template<typename HostAtomicType, typename HostDataType>
class CBasicTestStore : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::OldValueCheck;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryScope;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  using CBasicTest<HostAtomicType, HostDataType>::CheckCapabilities;
  CBasicTestStore(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    OldValueCheck(false);
  }
  virtual cl_uint NumResults(cl_uint threadCount, cl_device_id deviceID)
  {
    return threadCount;
  }
  virtual int ExecuteSingleTest(cl_device_id deviceID, cl_context context, cl_command_queue queue)
  {
    if(MemoryOrder() == MEMORY_ORDER_ACQUIRE ||
      MemoryOrder() == MEMORY_ORDER_ACQ_REL)
      return 0; //skip test - not applicable

    if (CheckCapabilities(MemoryScope(), MemoryOrder()) == TEST_SKIPPED_ITSELF)
        return 0; // skip test - not applicable

    return CBasicTestMemOrderScope<HostAtomicType, HostDataType>::ExecuteSingleTest(deviceID, context, queue);
  }
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScopeStr();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      "  atomic_store"+postfix+"(&destMemory[tid], tid"+memoryOrderScope+");\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    host_atomic_store(&destMemory[tid], (HostDataType)tid, MemoryOrder());
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    expected = (HostDataType)whichDestValue;
    return true;
  }
};

int test_atomic_store_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestStore<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestStore<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestStore<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestStore<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  CBasicTestStore<HOST_ATOMIC_FLOAT, HOST_FLOAT> test_float(TYPE_ATOMIC_FLOAT, useSVM);
  EXECUTE_TEST(error, test_float.Execute(deviceID, context, queue, num_elements));
  CBasicTestStore<HOST_ATOMIC_DOUBLE, HOST_DOUBLE> test_double(TYPE_ATOMIC_DOUBLE, useSVM);
  EXECUTE_TEST(error, test_double.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestStore<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestStore<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestStore<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestStore<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestStore<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestStore<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestStore<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestStore<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_store(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_store_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_store(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_store_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestInit : public CBasicTest<HostAtomicType, HostDataType>
{
public:
  using CBasicTest<HostAtomicType, HostDataType>::OldValueCheck;
  CBasicTestInit(TExplicitAtomicType dataType, bool useSVM) : CBasicTest<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    OldValueCheck(false);
  }
  virtual cl_uint NumResults(cl_uint threadCount, cl_device_id deviceID)
  {
    return threadCount;
  }
  virtual std::string ProgramCore()
  {
    return
      "  atomic_init(&destMemory[tid], tid);\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    host_atomic_init(&destMemory[tid], (HostDataType)tid);
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    expected = (HostDataType)whichDestValue;
    return true;
  }
};

int test_atomic_init_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestInit<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestInit<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestInit<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestInit<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  CBasicTestInit<HOST_ATOMIC_FLOAT, HOST_FLOAT> test_float(TYPE_ATOMIC_FLOAT, useSVM);
  EXECUTE_TEST(error, test_float.Execute(deviceID, context, queue, num_elements));
  CBasicTestInit<HOST_ATOMIC_DOUBLE, HOST_DOUBLE> test_double(TYPE_ATOMIC_DOUBLE, useSVM);
  EXECUTE_TEST(error, test_double.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestInit<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestInit<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestInit<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestInit<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestInit<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestInit<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestInit<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestInit<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_init(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_init_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_init(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_init_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestLoad : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::OldValueCheck;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryScope;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryScopeStr;
  using CBasicTest<HostAtomicType, HostDataType>::CheckCapabilities;
  CBasicTestLoad(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    OldValueCheck(false);
  }
  virtual cl_uint NumResults(cl_uint threadCount, cl_device_id deviceID)
  {
    return threadCount;
  }
  virtual int ExecuteSingleTest(cl_device_id deviceID, cl_context context, cl_command_queue queue)
  {
    if(MemoryOrder() == MEMORY_ORDER_RELEASE ||
      MemoryOrder() == MEMORY_ORDER_ACQ_REL)
      return 0; //skip test - not applicable

    if (CheckCapabilities(MemoryScope(), MemoryOrder()) == TEST_SKIPPED_ITSELF)
        return 0; // skip test - not applicable

    return CBasicTestMemOrderScope<HostAtomicType, HostDataType>::ExecuteSingleTest(deviceID, context, queue);
  }
  virtual std::string ProgramCore()
  {
      // In the case this test is run with MEMORY_ORDER_ACQUIRE, the store
      // should be MEMORY_ORDER_RELEASE
      std::string memoryOrderScopeLoad = MemoryOrderScopeStr();
      std::string memoryOrderScopeStore =
          (MemoryOrder() == MEMORY_ORDER_ACQUIRE)
          ? (", memory_order_release" + MemoryScopeStr())
          : memoryOrderScopeLoad;
      std::string postfix(memoryOrderScopeLoad.empty() ? "" : "_explicit");
      return "  atomic_store" + postfix + "(&destMemory[tid], tid"
          + memoryOrderScopeStore
          + ");\n"
            "  oldValues[tid] = atomic_load"
          + postfix + "(&destMemory[tid]" + memoryOrderScopeLoad + ");\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    host_atomic_store(&destMemory[tid], (HostDataType)tid, MEMORY_ORDER_SEQ_CST);
    oldValues[tid] = host_atomic_load<HostAtomicType, HostDataType>(&destMemory[tid], MemoryOrder());
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    expected = (HostDataType)whichDestValue;
    return true;
  }
  virtual bool VerifyRefs(bool &correct, cl_uint threadCount, HostDataType *refValues, HostAtomicType *finalValues)
  {
    correct = true;
    for(cl_uint i = 0; i < threadCount; i++ )
    {
      if(refValues[i] != (HostDataType)i)
      {
        log_error("Invalid value for thread %u\n", (cl_uint)i);
        correct = false;
        return true;
      }
    }
    return true;
  }
};

int test_atomic_load_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestLoad<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestLoad<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestLoad<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestLoad<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  CBasicTestLoad<HOST_ATOMIC_FLOAT, HOST_FLOAT> test_float(TYPE_ATOMIC_FLOAT, useSVM);
  EXECUTE_TEST(error, test_float.Execute(deviceID, context, queue, num_elements));
  CBasicTestLoad<HOST_ATOMIC_DOUBLE, HOST_DOUBLE> test_double(TYPE_ATOMIC_DOUBLE, useSVM);
  EXECUTE_TEST(error, test_double.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestLoad<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestLoad<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestLoad<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestLoad<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestLoad<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestLoad<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestLoad<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestLoad<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_load(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_load_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_load(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_load_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestExchange : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::OldValueCheck;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::Iterations;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::IterationsStr;
  CBasicTestExchange(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    StartValue(123456);
  }
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScopeStr();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      "  oldValues[tid] = atomic_exchange"+postfix+"(&destMemory[0], tid"+memoryOrderScope+");\n"
      "  for(int i = 0; i < "+IterationsStr()+"; i++)\n"
      "    oldValues[tid] = atomic_exchange"+postfix+"(&destMemory[0], oldValues[tid]"+memoryOrderScope+");\n";
  }

  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    oldValues[tid] = host_atomic_exchange(&destMemory[0], (HostDataType)tid, MemoryOrder());
    for(int i = 0; i < Iterations(); i++)
      oldValues[tid] = host_atomic_exchange(&destMemory[0], oldValues[tid], MemoryOrder());
  }
  virtual bool VerifyRefs(bool &correct, cl_uint threadCount, HostDataType *refValues, HostAtomicType *finalValues)
  {
    OldValueCheck(Iterations()%2 == 0); //check is valid for even number of iterations only
    correct = true;
    /* We are expecting values from 0 to size-1 and initial value from atomic variable */
    /* These values must be distributed across refValues array and atomic variable finalVaue[0] */
    /* Any repeated value is treated as an error */
    std::vector<bool> tidFound(threadCount);
    bool startValueFound = false;
    cl_uint i;

    for(i = 0; i <= threadCount; i++)
    {
      cl_uint value;
      if(i == threadCount)
        value = (cl_uint)finalValues[0]; //additional value from atomic variable (last written)
      else
        value = (cl_uint)refValues[i];
      if(value == (cl_uint)StartValue())
      {
        // Special initial value
        if(startValueFound)
        {
          log_error("ERROR: Starting reference value (%u) occurred more thane once\n", (cl_uint)StartValue());
          correct = false;
          return true;
        }
        startValueFound = true;
        continue;
      }
      if(value >= threadCount)
      {
        log_error("ERROR: Reference value %u outside of valid range! (%u)\n", i, value);
        correct = false;
        return true;
      }
      if(tidFound[value])
      {
        log_error("ERROR: Value (%u) occurred more thane once\n", value);
        correct = false;
        return true;
      }
      tidFound[value] = true;
    }
    return true;
  }
};

int test_atomic_exchange_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestExchange<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestExchange<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestExchange<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestExchange<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  CBasicTestExchange<HOST_ATOMIC_FLOAT, HOST_FLOAT> test_float(TYPE_ATOMIC_FLOAT, useSVM);
  EXECUTE_TEST(error, test_float.Execute(deviceID, context, queue, num_elements));
  CBasicTestExchange<HOST_ATOMIC_DOUBLE, HOST_DOUBLE> test_double(TYPE_ATOMIC_DOUBLE, useSVM);
  EXECUTE_TEST(error, test_double.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestExchange<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestExchange<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestExchange<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestExchange<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestExchange<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestExchange<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestExchange<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestExchange<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_exchange(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_exchange_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_exchange(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_exchange_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestCompareStrong : public CBasicTestMemOrder2Scope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrder2Scope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrder2Scope<HostAtomicType, HostDataType>::OldValueCheck;
  using CBasicTestMemOrder2Scope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrder2Scope<HostAtomicType, HostDataType>::MemoryOrder2;
  using CBasicTestMemOrder2Scope<HostAtomicType, HostDataType>::MemoryOrderScope;
  using CBasicTestMemOrder2Scope<HostAtomicType, HostDataType>::MemoryScope;
  using CBasicTestMemOrder2Scope<HostAtomicType, HostDataType>::DataType;
  using CBasicTestMemOrder2Scope<HostAtomicType, HostDataType>::Iterations;
  using CBasicTestMemOrder2Scope<HostAtomicType, HostDataType>::IterationsStr;
  using CBasicTest<HostAtomicType, HostDataType>::CheckCapabilities;
  CBasicTestCompareStrong(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrder2Scope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    StartValue(123456);
    OldValueCheck(false);
  }
  virtual int ExecuteSingleTest(cl_device_id deviceID, cl_context context, cl_command_queue queue)
  {
    if(MemoryOrder2() == MEMORY_ORDER_RELEASE ||
      MemoryOrder2() == MEMORY_ORDER_ACQ_REL)
      return 0; // not allowed as 'failure' argument
    if((MemoryOrder() == MEMORY_ORDER_RELAXED && MemoryOrder2() != MEMORY_ORDER_RELAXED) ||
      (MemoryOrder() != MEMORY_ORDER_SEQ_CST && MemoryOrder2() == MEMORY_ORDER_SEQ_CST))
      return 0; // failure argument shall be no stronger than the success

    if (CheckCapabilities(MemoryScope(), MemoryOrder()) == TEST_SKIPPED_ITSELF)
        return 0; // skip test - not applicable

    if (CheckCapabilities(MemoryScope(), MemoryOrder2()) == TEST_SKIPPED_ITSELF)
        return 0; // skip test - not applicable

    return CBasicTestMemOrder2Scope<HostAtomicType, HostDataType>::ExecuteSingleTest(deviceID, context, queue);
  }
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScope();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      std::string("  ")+DataType().RegularTypeName()+" expected, previous;\n"
      "  int successCount = 0;\n"
      "  oldValues[tid] = tid;\n"
      "  expected = tid;  // force failure at the beginning\n"
      "  if(atomic_compare_exchange_strong"+postfix+"(&destMemory[0], &expected, oldValues[tid]"+memoryOrderScope+") || expected == tid)\n"
      "    oldValues[tid] = threadCount+1; //mark unexpected success with invalid value\n"
      "  else\n"
      "  {\n"
      "    for(int i = 0; i < "+IterationsStr()+" || successCount == 0; i++)\n"
      "    {\n"
      "      previous = expected;\n"
      "      if(atomic_compare_exchange_strong"+postfix+"(&destMemory[0], &expected, oldValues[tid]"+memoryOrderScope+"))\n"
      "      {\n"
      "        oldValues[tid] = expected;\n"
      "        successCount++;\n"
      "      }\n"
      "      else\n"
      "      {\n"
      "        if(previous == expected) // spurious failure - shouldn't occur for 'strong'\n"
      "        {\n"
      "          oldValues[tid] = threadCount; //mark fail with invalid value\n"
      "          break;\n"
      "        }\n"
      "      }\n"
      "    }\n"
      "  }\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    HostDataType expected = (HostDataType)StartValue(), previous;
    oldValues[tid] = (HostDataType)tid;
    for(int i = 0; i < Iterations(); i++)
    {
      previous = expected;
      if(host_atomic_compare_exchange(&destMemory[0], &expected, oldValues[tid], MemoryOrder(), MemoryOrder2()))
        oldValues[tid] = expected;
      else
      {
        if(previous == expected) // shouldn't occur for 'strong'
        {
          oldValues[tid] = threadCount; //mark fail with invalid value
        }
      }
    }
  }
  virtual bool VerifyRefs(bool &correct, cl_uint threadCount, HostDataType *refValues, HostAtomicType *finalValues)
  {
    correct = true;
    /* We are expecting values from 0 to size-1 and initial value from atomic variable */
    /* These values must be distributed across refValues array and atomic variable finalVaue[0] */
    /* Any repeated value is treated as an error */
    std::vector<bool> tidFound(threadCount);
    bool startValueFound = false;
    cl_uint i;

    for(i = 0; i <= threadCount; i++)
    {
      cl_uint value;
      if(i == threadCount)
        value = (cl_uint)finalValues[0]; //additional value from atomic variable (last written)
      else
        value = (cl_uint)refValues[i];
      if(value == (cl_uint)StartValue())
      {
        // Special initial value
        if(startValueFound)
        {
          log_error("ERROR: Starting reference value (%u) occurred more thane once\n", (cl_uint)StartValue());
          correct = false;
          return true;
        }
        startValueFound = true;
        continue;
      }
      if(value >= threadCount)
      {
        if(value == threadCount)
          log_error("ERROR: Spurious failure detected for atomic_compare_exchange_strong\n");
        log_error("ERROR: Reference value %u outside of valid range! (%u)\n", i, value);
        correct = false;
        return true;
      }
      if(tidFound[value])
      {
        log_error("ERROR: Value (%u) occurred more thane once\n", value);
        correct = false;
        return true;
      }
      tidFound[value] = true;
    }
    return true;
  }
};

int test_atomic_compare_exchange_strong_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestCompareStrong<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestCompareStrong<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestCompareStrong<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestCompareStrong<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestCompareStrong<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestCompareStrong<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestCompareStrong<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestCompareStrong<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestCompareStrong<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestCompareStrong<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestCompareStrong<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestCompareStrong<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_compare_exchange_strong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_compare_exchange_strong_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_compare_exchange_strong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_compare_exchange_strong_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestCompareWeak : public CBasicTestCompareStrong<HostAtomicType, HostDataType>
{
public:
  using CBasicTestCompareStrong<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestCompareStrong<HostAtomicType, HostDataType>::MemoryOrderScope;
  using CBasicTestCompareStrong<HostAtomicType, HostDataType>::DataType;
  using CBasicTestCompareStrong<HostAtomicType, HostDataType>::Iterations;
  using CBasicTestCompareStrong<HostAtomicType, HostDataType>::IterationsStr;
  CBasicTestCompareWeak(TExplicitAtomicType dataType, bool useSVM) : CBasicTestCompareStrong<HostAtomicType, HostDataType>(dataType, useSVM)
  {
  }
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScope();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      std::string("  ")+DataType().RegularTypeName()+" expected , previous;\n"
      "  int successCount = 0;\n"
      "  oldValues[tid] = tid;\n"
      "  expected = tid;  // force failure at the beginning\n"
      "  if(atomic_compare_exchange_weak"+postfix+"(&destMemory[0], &expected, oldValues[tid]"+memoryOrderScope+") || expected == tid)\n"
      "    oldValues[tid] = threadCount+1; //mark unexpected success with invalid value\n"
      "  else\n"
      "  {\n"
      "    for(int i = 0; i < "+IterationsStr()+" || successCount == 0; i++)\n"
      "    {\n"
      "      previous = expected;\n"
      "      if(atomic_compare_exchange_weak"+postfix+"(&destMemory[0], &expected, oldValues[tid]"+memoryOrderScope+"))\n"
      "      {\n"
      "        oldValues[tid] = expected;\n"
      "        successCount++;\n"
      "      }\n"
      "    }\n"
      "  }\n";
  }
};

int test_atomic_compare_exchange_weak_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestCompareWeak<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestCompareWeak<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestCompareWeak<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestCompareWeak<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestCompareWeak<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestCompareWeak<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestCompareWeak<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestCompareWeak<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestCompareWeak<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestCompareWeak<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestCompareWeak<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestCompareWeak<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_compare_exchange_weak(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_compare_exchange_weak_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_compare_exchange_weak(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_compare_exchange_weak_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestFetchAdd : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::DataType;
  CBasicTestFetchAdd(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
  }
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScopeStr();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      "  oldValues[tid] = atomic_fetch_add"+postfix+"(&destMemory[0], ("+DataType().AddSubOperandTypeName()+")tid + 3"+memoryOrderScope+");\n"+
      "  atomic_fetch_add"+postfix+"(&destMemory[0], ("+DataType().AddSubOperandTypeName()+")tid + 3"+memoryOrderScope+");\n"
      "  atomic_fetch_add"+postfix+"(&destMemory[0], ("+DataType().AddSubOperandTypeName()+")tid + 3"+memoryOrderScope+");\n"
      "  atomic_fetch_add"+postfix+"(&destMemory[0], (("+DataType().AddSubOperandTypeName()+")tid + 3) << (sizeof("+DataType().AddSubOperandTypeName()+")-1)*8"+memoryOrderScope+");\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    oldValues[tid] = host_atomic_fetch_add(&destMemory[0], (HostDataType)tid + 3, MemoryOrder());
    host_atomic_fetch_add(&destMemory[0], (HostDataType)tid + 3, MemoryOrder());
    host_atomic_fetch_add(&destMemory[0], (HostDataType)tid + 3, MemoryOrder());
    host_atomic_fetch_add(&destMemory[0], ((HostDataType)tid + 3) << (sizeof(HostDataType)-1)*8, MemoryOrder());
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    expected = StartValue();
    for(cl_uint i = 0; i < threadCount; i++)
      expected += ((HostDataType)i+3)*3+(((HostDataType)i + 3) << (sizeof(HostDataType)-1)*8);
    return true;
  }
};

int test_atomic_fetch_add_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestFetchAdd<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchAdd<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchAdd<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchAdd<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestFetchAdd<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchAdd<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchAdd<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchAdd<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestFetchAdd<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchAdd<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchAdd<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchAdd<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_fetch_add(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_add_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_fetch_add(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_add_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestFetchSub : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::DataType;
  CBasicTestFetchSub(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
  }
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScopeStr();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      "  oldValues[tid] = atomic_fetch_sub"+postfix+"(&destMemory[0], tid + 3 +((("+DataType().AddSubOperandTypeName()+")tid + 3) << (sizeof("+DataType().AddSubOperandTypeName()+")-1)*8)"+memoryOrderScope+");\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    oldValues[tid] = host_atomic_fetch_sub(&destMemory[0], (HostDataType)tid + 3+(((HostDataType)tid + 3) << (sizeof(HostDataType)-1)*8), MemoryOrder());
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    expected = StartValue();
    for(cl_uint i = 0; i < threadCount; i++)
      expected -= (HostDataType)i + 3 +(((HostDataType)i + 3) << (sizeof(HostDataType)-1)*8);
    return true;
  }
};

int test_atomic_fetch_sub_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestFetchSub<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchSub<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchSub<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchSub<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestFetchSub<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchSub<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchSub<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchSub<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestFetchSub<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchSub<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchSub<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchSub<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_fetch_sub(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_sub_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_fetch_sub(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_sub_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestFetchOr : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::DataType;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  CBasicTestFetchOr(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    StartValue(0);
  }
  virtual cl_uint NumResults(cl_uint threadCount, cl_device_id deviceID)
  {
    cl_uint numBits = DataType().Size(deviceID) * 8;

    return (threadCount + numBits - 1) / numBits;
  }
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScopeStr();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      std::string("    size_t numBits = sizeof(")+DataType().RegularTypeName()+") * 8;\n"
      "    int whichResult = tid / numBits;\n"
      "    int bitIndex = tid - (whichResult * numBits);\n"
      "\n"
      "    oldValues[tid] = atomic_fetch_or"+postfix+"(&destMemory[whichResult], (("+DataType().RegularTypeName()+")1 << bitIndex) "+memoryOrderScope+");\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    size_t numBits = sizeof(HostDataType) * 8;
    size_t whichResult = tid / numBits;
    size_t bitIndex = tid - (whichResult * numBits);

    oldValues[tid] = host_atomic_fetch_or(&destMemory[whichResult], ((HostDataType)1 << bitIndex), MemoryOrder());
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    cl_uint numValues = (threadCount + (sizeof(HostDataType)*8-1)) / (sizeof(HostDataType)*8);
    if(whichDestValue < numValues - 1)
    {
      expected = ~(HostDataType)0;
      return true;
    }
    // Last item doesn't get or'ed on every bit, so we have to mask away
    cl_uint numBits = threadCount - whichDestValue * (sizeof(HostDataType)*8);
    expected = StartValue();
    for(cl_uint i = 0; i < numBits; i++)
      expected |= ((HostDataType)1 << i);
    return true;
  }
};

int test_atomic_fetch_or_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestFetchOr<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchOr<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchOr<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchOr<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestFetchOr<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchOr<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchOr<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchOr<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestFetchOr<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchOr<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchOr<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchOr<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_fetch_or(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_or_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_fetch_or(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_or_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestFetchXor : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::DataType;
  CBasicTestFetchXor(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    StartValue((HostDataType)0x2f08ab418ba0541LL);
  }
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScopeStr();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      std::string("  int numBits = sizeof(")+DataType().RegularTypeName()+") * 8;\n"
      "  int bitIndex = (numBits-1)*(tid+1)/threadCount;\n"
      "\n"
      "  oldValues[tid] = atomic_fetch_xor"+postfix+"(&destMemory[0], (("+DataType().RegularTypeName()+")1 << bitIndex) "+memoryOrderScope+");\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    int numBits = sizeof(HostDataType) * 8;
    int bitIndex = (numBits-1)*(tid+1)/threadCount;

    oldValues[tid] = host_atomic_fetch_xor(&destMemory[0], ((HostDataType)1 << bitIndex), MemoryOrder());
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    int numBits = sizeof(HostDataType)*8;
    expected = StartValue();
    for(cl_uint i = 0; i < threadCount; i++)
    {
      int bitIndex = (numBits-1)*(i+1)/threadCount;
      expected ^= ((HostDataType)1 << bitIndex);
    }
    return true;
  }
};

int test_atomic_fetch_xor_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestFetchXor<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchXor<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchXor<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchXor<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestFetchXor<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchXor<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchXor<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchXor<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestFetchXor<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchXor<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchXor<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchXor<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_fetch_xor(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_xor_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_fetch_xor(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_xor_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestFetchAnd : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::DataType;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  CBasicTestFetchAnd(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    StartValue(~(HostDataType)0);
  }
  virtual cl_uint NumResults(cl_uint threadCount, cl_device_id deviceID)
  {
    cl_uint numBits = DataType().Size(deviceID) * 8;

    return (threadCount + numBits - 1) / numBits;
  }
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScopeStr();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      std::string("  size_t numBits = sizeof(")+DataType().RegularTypeName()+") * 8;\n"
      "  int whichResult = tid / numBits;\n"
      "  int bitIndex = tid - (whichResult * numBits);\n"
      "\n"
      "  oldValues[tid] = atomic_fetch_and"+postfix+"(&destMemory[whichResult], ~(("+DataType().RegularTypeName()+")1 << bitIndex) "+memoryOrderScope+");\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    size_t numBits = sizeof(HostDataType) * 8;
    size_t whichResult = tid / numBits;
    size_t bitIndex = tid - (whichResult * numBits);

    oldValues[tid] = host_atomic_fetch_and(&destMemory[whichResult], ~((HostDataType)1 << bitIndex), MemoryOrder());
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    cl_uint numValues = (threadCount + (sizeof(HostDataType)*8-1)) / (sizeof(HostDataType)*8);
    if(whichDestValue < numValues - 1)
    {
      expected = 0;
      return true;
    }
    // Last item doesn't get and'ed on every bit, so we have to mask away
    size_t numBits = threadCount - whichDestValue * (sizeof(HostDataType)*8);
    expected = StartValue();
    for(size_t i = 0; i < numBits; i++)
      expected &= ~((HostDataType)1 << i);
    return true;
  }
};

int test_atomic_fetch_and_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestFetchAnd<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchAnd<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchAnd<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchAnd<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestFetchAnd<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchAnd<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchAnd<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchAnd<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestFetchAnd<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchAnd<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchAnd<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchAnd<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_fetch_and(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_and_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_fetch_and(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_and_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestFetchOrAnd : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::DataType;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::Iterations;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::IterationsStr;
  CBasicTestFetchOrAnd(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    StartValue(0);
  }
  virtual cl_uint NumResults(cl_uint threadCount, cl_device_id deviceID)
  {
    return 1+(threadCount-1)/(DataType().Size(deviceID)*8);
  }
  // each thread modifies (with OR and AND operations) and verifies
  // only one bit in atomic variable
  // other bits are modified by other threads but it must not affect current thread operation
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScopeStr();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      std::string("  int bits = sizeof(")+DataType().RegularTypeName()+")*8;\n"+
      "  size_t valueInd = tid/bits;\n"
      "  "+DataType().RegularTypeName()+" value, bitMask = ("+DataType().RegularTypeName()+")1 << tid%bits;\n"
      "  oldValues[tid] = 0;\n"
      "  for(int i = 0; i < "+IterationsStr()+"; i++)\n"
      "  {\n"
      "    value = atomic_fetch_or"+postfix+"(destMemory+valueInd, bitMask"+memoryOrderScope+");\n"
      "    if(value & bitMask) // bit should be set to 0\n"
      "      oldValues[tid]++;\n"
      "    value = atomic_fetch_and"+postfix+"(destMemory+valueInd, ~bitMask"+memoryOrderScope+");\n"
      "    if(!(value & bitMask)) // bit should be set to 1\n"
      "      oldValues[tid]++;\n"
      "  }\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    int bits = sizeof(HostDataType)*8;
    size_t valueInd = tid/bits;
    HostDataType value, bitMask = (HostDataType)1 << tid%bits;
    oldValues[tid] = 0;
    for(int i = 0; i < Iterations(); i++)
    {
      value = host_atomic_fetch_or(destMemory+valueInd, bitMask, MemoryOrder());
      if(value & bitMask) // bit should be set to 0
        oldValues[tid]++;
      value = host_atomic_fetch_and(destMemory+valueInd, ~bitMask, MemoryOrder());
      if(!(value & bitMask)) // bit should be set to 1
        oldValues[tid]++;
    }
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    expected = 0;
    return true;
  }
  virtual bool VerifyRefs(bool &correct, cl_uint threadCount, HostDataType *refValues, HostAtomicType *finalValues)
  {
    correct = true;
    for(cl_uint i = 0; i < threadCount; i++)
    {
      if(refValues[i] > 0)
      {
        log_error("Thread %d found %d mismatch(es)\n", i, (cl_uint)refValues[i]);
        correct = false;
      }
    }
    return true;
  }
};

int test_atomic_fetch_orand_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestFetchOrAnd<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchOrAnd<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchOrAnd<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchOrAnd<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestFetchOrAnd<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchOrAnd<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchOrAnd<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchOrAnd<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestFetchOrAnd<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchOrAnd<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchOrAnd<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchOrAnd<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_fetch_orand(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_orand_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_fetch_orand(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_orand_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestFetchXor2 : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::DataType;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::Iterations;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::IterationsStr;
  CBasicTestFetchXor2(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    StartValue(0);
  }
  virtual cl_uint NumResults(cl_uint threadCount, cl_device_id deviceID)
  {
    return 1+(threadCount-1)/(DataType().Size(deviceID)*8);
  }
  // each thread modifies (with XOR operation) and verifies
  // only one bit in atomic variable
  // other bits are modified by other threads but it must not affect current thread operation
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScopeStr();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      std::string("  int bits = sizeof(")+DataType().RegularTypeName()+")*8;\n"+
      "  size_t valueInd = tid/bits;\n"
      "  "+DataType().RegularTypeName()+" value, bitMask = ("+DataType().RegularTypeName()+")1 << tid%bits;\n"
      "  oldValues[tid] = 0;\n"
      "  for(int i = 0; i < "+IterationsStr()+"; i++)\n"
      "  {\n"
      "    value = atomic_fetch_xor"+postfix+"(destMemory+valueInd, bitMask"+memoryOrderScope+");\n"
      "    if(value & bitMask) // bit should be set to 0\n"
      "      oldValues[tid]++;\n"
      "    value = atomic_fetch_xor"+postfix+"(destMemory+valueInd, bitMask"+memoryOrderScope+");\n"
      "    if(!(value & bitMask)) // bit should be set to 1\n"
      "      oldValues[tid]++;\n"
      "  }\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    int bits = sizeof(HostDataType)*8;
    size_t valueInd = tid/bits;
    HostDataType value, bitMask = (HostDataType)1 << tid%bits;
    oldValues[tid] = 0;
    for(int i = 0; i < Iterations(); i++)
    {
      value = host_atomic_fetch_xor(destMemory+valueInd, bitMask, MemoryOrder());
      if(value & bitMask) // bit should be set to 0
        oldValues[tid]++;
      value = host_atomic_fetch_xor(destMemory+valueInd, bitMask, MemoryOrder());
      if(!(value & bitMask)) // bit should be set to 1
        oldValues[tid]++;
    }
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    expected = 0;
    return true;
  }
  virtual bool VerifyRefs(bool &correct, cl_uint threadCount, HostDataType *refValues, HostAtomicType *finalValues)
  {
    correct = true;
    for(cl_uint i = 0; i < threadCount; i++)
    {
      if(refValues[i] > 0)
      {
        log_error("Thread %d found %d mismatches\n", i, (cl_uint)refValues[i]);
        correct = false;
      }
    }
    return true;
  }
};

int test_atomic_fetch_xor2_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestFetchXor2<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchXor2<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchXor2<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchXor2<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestFetchXor2<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchXor2<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchXor2<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchXor2<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestFetchXor2<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchXor2<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchXor2<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchXor2<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_fetch_xor2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_xor2_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_fetch_xor2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_xor2_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestFetchMin : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::DataType;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  CBasicTestFetchMin(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    StartValue(DataType().MaxValue());
  }
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScopeStr();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      "  oldValues[tid] = atomic_fetch_min"+postfix+"(&destMemory[0], oldValues[tid] "+memoryOrderScope+");\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    oldValues[tid] = host_atomic_fetch_min(&destMemory[0], oldValues[tid], MemoryOrder());
  }
  virtual bool GenerateRefs(cl_uint threadCount, HostDataType *startRefValues, MTdata d)
  {
    for(cl_uint i = 0; i < threadCount; i++)
    {
      startRefValues[i] = genrand_int32(d);
      if(sizeof(HostDataType) >= 8)
        startRefValues[i] |= (HostDataType)genrand_int32(d) << 16;
    }
    return true;
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    expected = StartValue();
    for(cl_uint i = 0; i < threadCount; i++)
    {
      if(startRefValues[ i ] < expected)
        expected = startRefValues[ i ];
    }
    return true;
  }
};

int test_atomic_fetch_min_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestFetchMin<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchMin<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchMin<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchMin<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestFetchMin<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchMin<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchMin<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchMin<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestFetchMin<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchMin<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchMin<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchMin<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_fetch_min(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_min_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_fetch_min(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_min_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestFetchMax : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::DataType;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  CBasicTestFetchMax(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    StartValue(DataType().MinValue());
  }
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScopeStr();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    return
      "  oldValues[tid] = atomic_fetch_max"+postfix+"(&destMemory[0], oldValues[tid] "+memoryOrderScope+");\n";
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    oldValues[tid] = host_atomic_fetch_max(&destMemory[0], oldValues[tid], MemoryOrder());
  }
  virtual bool GenerateRefs(cl_uint threadCount, HostDataType *startRefValues, MTdata d)
  {
    for(cl_uint i = 0; i < threadCount; i++)
    {
      startRefValues[i] = genrand_int32(d);
      if(sizeof(HostDataType) >= 8)
        startRefValues[i] |= (HostDataType)genrand_int32(d) << 16;
    }
    return true;
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    expected = StartValue();
    for(cl_uint i = 0; i < threadCount; i++)
    {
      if(startRefValues[ i ] > expected)
        expected = startRefValues[ i ];
    }
    return true;
  }
};

int test_atomic_fetch_max_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestFetchMax<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchMax<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchMax<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestFetchMax<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  if(AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestFetchMax<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchMax<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchMax<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchMax<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestFetchMax<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchMax<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchMax<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFetchMax<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_fetch_max(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_max_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_fetch_max(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fetch_max_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestFlag : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
  static const HostDataType CRITICAL_SECTION_NOT_VISITED = 1000000000;
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::OldValueCheck;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryScopeStr;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::UseSVM;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::LocalMemory;
  CBasicTestFlag(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    StartValue(0);
    OldValueCheck(false);
  }
  virtual cl_uint NumResults(cl_uint threadCount, cl_device_id deviceID)
  {
    return threadCount;
  }
  TExplicitMemoryOrderType MemoryOrderForClear()
  {
    // Memory ordering for atomic_flag_clear function
    // ("shall not be memory_order_acquire nor memory_order_acq_rel")
    if(MemoryOrder() == MEMORY_ORDER_ACQUIRE)
      return MEMORY_ORDER_RELAXED;
    if (MemoryOrder() == MEMORY_ORDER_ACQ_REL)
      return MEMORY_ORDER_RELEASE;
    return MemoryOrder();
  }
  std::string MemoryOrderScopeStrForClear()
  {
    std::string orderStr;
    if (MemoryOrder() != MEMORY_ORDER_EMPTY)
      orderStr = std::string(", ") + get_memory_order_type_name(MemoryOrderForClear());
    return orderStr + MemoryScopeStr();
  }

  virtual int ExecuteSingleTest(cl_device_id deviceID, cl_context context,
                                cl_command_queue queue)
  {
      // This test assumes support for the memory_scope_device scope in the case
      // that LocalMemory() == false. Therefore we should skip this test in that
      // configuration on a 3.0 driver since supporting the memory_scope_device
      // scope is optionaly.
      if (get_device_cl_version(deviceID) >= Version{ 3, 0 })
      {
          if (!LocalMemory()
              && !(gAtomicFenceCap & CL_DEVICE_ATOMIC_SCOPE_DEVICE))
          {
              log_info(
                  "Skipping atomic_flag test due to use of atomic_scope_device "
                  "which is optionally not supported on this device\n");
              return 0; // skip test - not applicable
          }
      }
      return CBasicTestMemOrderScope<HostAtomicType,
                                     HostDataType>::ExecuteSingleTest(deviceID,
                                                                      context,
                                                                      queue);
  }
  virtual std::string ProgramCore()
  {
    std::string memoryOrderScope = MemoryOrderScopeStr();
    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
    std::string program =
      "  uint cnt, stop = 0;\n"
      "  for(cnt = 0; !stop && cnt < threadCount; cnt++) // each thread must find critical section where it is the first visitor\n"
      "  {\n"
      "    bool set = atomic_flag_test_and_set" + postfix + "(&destMemory[cnt]" + memoryOrderScope + ");\n";
    if (MemoryOrder() == MEMORY_ORDER_RELAXED || MemoryOrder() == MEMORY_ORDER_RELEASE)
      program += "    atomic_work_item_fence(" +
                 std::string(LocalMemory() ? "CLK_LOCAL_MEM_FENCE, " : "CLK_GLOBAL_MEM_FENCE, ") +
                 "memory_order_acquire," +
                 std::string(LocalMemory() ? "memory_scope_work_group" : (UseSVM() ? "memory_scope_all_svm_devices" : "memory_scope_device") ) +
                 ");\n";

    program +=
      "    if (!set)\n"
      "    {\n";

    if (LocalMemory())
      program += "      uint csIndex = get_enqueued_local_size(0)*get_group_id(0)+cnt;\n";
    else
      program += "      uint csIndex = cnt;\n";

    std::ostringstream csNotVisited;
    csNotVisited << CRITICAL_SECTION_NOT_VISITED;
    program +=
      "      // verify that thread is the first visitor\n"
      "      if(oldValues[csIndex] == "+csNotVisited.str()+")\n"
      "      {\n"
      "        oldValues[csIndex] = tid; // set the winner id for this critical section\n"
      "        stop = 1;\n"
      "      }\n";

    if (MemoryOrder() == MEMORY_ORDER_ACQUIRE || MemoryOrder() == MEMORY_ORDER_RELAXED)
      program += "      atomic_work_item_fence(" +
                 std::string(LocalMemory() ? "CLK_LOCAL_MEM_FENCE, " : "CLK_GLOBAL_MEM_FENCE, ") +
                 "memory_order_release," +
                 std::string(LocalMemory() ? "memory_scope_work_group" : (UseSVM() ? "memory_scope_all_svm_devices" : "memory_scope_device") ) +
                 ");\n";

    program +=
      "      atomic_flag_clear" + postfix + "(&destMemory[cnt]" + MemoryOrderScopeStrForClear() + ");\n"
      "    }\n"
      "  }\n";
    return program;
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    cl_uint cnt, stop = 0;
    for (cnt = 0; !stop && cnt < threadCount; cnt++) // each thread must find critical section where it is the first visitor\n"
    {
      if (!host_atomic_flag_test_and_set(&destMemory[cnt], MemoryOrder()))
      {
        cl_uint csIndex = cnt;
        // verify that thread is the first visitor\n"
        if (oldValues[csIndex] == CRITICAL_SECTION_NOT_VISITED)
        {
          oldValues[csIndex] = tid; // set the winner id for this critical section\n"
          stop = 1;
        }
        host_atomic_flag_clear(&destMemory[cnt], MemoryOrderForClear());
      }
    }
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    expected = StartValue();
    return true;
  }
  virtual bool GenerateRefs(cl_uint threadCount, HostDataType *startRefValues, MTdata d)
  {
    for(cl_uint i = 0 ; i < threadCount; i++)
      startRefValues[i] = CRITICAL_SECTION_NOT_VISITED;
    return true;
  }
  virtual bool VerifyRefs(bool &correct, cl_uint threadCount, HostDataType *refValues, HostAtomicType *finalValues)
  {
    correct = true;
    /* We are expecting unique values from 0 to threadCount-1 (each critical section must be visited) */
    /* These values must be distributed across refValues array */
    std::vector<bool> tidFound(threadCount);
    cl_uint i;

    for (i = 0; i < threadCount; i++)
    {
      cl_uint value = (cl_uint)refValues[i];
      if (value == CRITICAL_SECTION_NOT_VISITED)
      {
        // Special initial value
        log_error("ERROR: Critical section %u not visited\n", i);
        correct = false;
        return true;
      }
      if (value >= threadCount)
      {
        log_error("ERROR: Reference value %u outside of valid range! (%u)\n", i, value);
        correct = false;
        return true;
      }
      if (tidFound[value])
      {
        log_error("ERROR: Value (%u) occurred more thane once\n", value);
        correct = false;
        return true;
      }
      tidFound[value] = true;
    }
    return true;
  }
};

int test_atomic_flag_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestFlag<HOST_ATOMIC_FLAG, HOST_FLAG> test_flag(TYPE_ATOMIC_FLAG, useSVM);
  EXECUTE_TEST(error, test_flag.Execute(deviceID, context, queue, num_elements));
  return error;
}

int test_atomic_flag(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_flag_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_flag(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_flag_generic(deviceID, context, queue, num_elements, true);
}

template<typename HostAtomicType, typename HostDataType>
class CBasicTestFence : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
  struct TestDefinition {
    bool op1IsFence;
    TExplicitMemoryOrderType op1MemOrder;
    bool op2IsFence;
    TExplicitMemoryOrderType op2MemOrder;
  };
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::StartValue;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::OldValueCheck;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryScope;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryScopeStr;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::DeclaredInProgram;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::UsedInFunction;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::DataType;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::CurrentGroupSize;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::UseSVM;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::LocalMemory;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::LocalRefValues;
  CBasicTestFence(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
    StartValue(0);
    OldValueCheck(false);
  }
  virtual cl_uint NumResults(cl_uint threadCount, cl_device_id deviceID)
  {
    return threadCount;
  }
  virtual cl_uint NumNonAtomicVariablesPerThread()
  {
    if (MemoryOrder() == MEMORY_ORDER_SEQ_CST)
      return 1;
    if (LocalMemory())
    {
      if (gIsEmbedded)
      {
        if (CurrentGroupSize() > 1024)
          CurrentGroupSize(1024);
        return 1; //1KB of local memory required by spec. Clamp group size to 1k and allow 1 variable per thread
      }
      else
        return 32 * 1024 / 8 / CurrentGroupSize() - 1; //32KB of local memory required by spec
    }
    return 256;
  }
  virtual std::string SingleTestName()
  {
    std::string testName;
    if (MemoryOrder() == MEMORY_ORDER_SEQ_CST)
      testName += "seq_cst fence, ";
    else
      testName += std::string(get_memory_order_type_name(_subCase.op1MemOrder)).substr(sizeof("memory_order"))
        + (_subCase.op1IsFence ? " fence" : " atomic") + " synchronizes-with "
        + std::string(get_memory_order_type_name(_subCase.op2MemOrder)).substr(sizeof("memory_order"))
        + (_subCase.op2IsFence ? " fence" : " atomic") + ", ";
    testName += CBasicTest<HostAtomicType, HostDataType>::SingleTestName();
    testName += std::string(", ") + std::string(get_memory_scope_type_name(MemoryScope())).substr(sizeof("memory"));
    return testName;
  }
  virtual bool SVMDataBufferAllSVMConsistent()
  {
      // Although memory_scope_all_devices doesn't mention SVM it is just an
      // alias for memory_scope_all_svm_devices.  So both scopes interact with
      // SVM allocations, on devices that support those, just the same.
      return MemoryScope() == MEMORY_SCOPE_ALL_DEVICES
          || MemoryScope() == MEMORY_SCOPE_ALL_SVM_DEVICES;
  }
  virtual int ExecuteForEachParameterSet(cl_device_id deviceID, cl_context context, cl_command_queue queue)
  {
    int error = 0;
    // execute 3 (maximum) sub cases for each memory order
    for (_subCaseId = 0; _subCaseId < 3; _subCaseId++)
    {
      EXECUTE_TEST(error, (CBasicTestMemOrderScope<HostAtomicType, HostDataType>::ExecuteForEachParameterSet(deviceID, context, queue)));
    }
    return error;
  }
  virtual int ExecuteSingleTest(cl_device_id deviceID, cl_context context, cl_command_queue queue)
  {
    if(DeclaredInProgram() || UsedInFunction())
      return 0; //skip test - not applicable - no overloaded fence functions for different address spaces
    if(MemoryOrder() == MEMORY_ORDER_EMPTY ||
      MemoryScope() == MEMORY_SCOPE_EMPTY) // empty 'scope' not required since opencl20-openclc-rev15
      return 0; //skip test - not applicable
    if((UseSVM() || gHost)
      && LocalMemory())
      return 0; // skip test - not applicable for SVM and local memory
    struct TestDefinition acqTests[] = {
      // {op1IsFence, op1MemOrder, op2IsFence, op2MemOrder}
      { false, MEMORY_ORDER_RELEASE, true, MEMORY_ORDER_ACQUIRE },
      { true, MEMORY_ORDER_RELEASE, true, MEMORY_ORDER_ACQUIRE },
      { true, MEMORY_ORDER_ACQ_REL, true, MEMORY_ORDER_ACQUIRE }
    };
    struct TestDefinition relTests[] = {
      { true, MEMORY_ORDER_RELEASE, false, MEMORY_ORDER_ACQUIRE },
      { true, MEMORY_ORDER_RELEASE, true, MEMORY_ORDER_ACQ_REL }
    };
    struct TestDefinition arTests[] = {
      { false, MEMORY_ORDER_RELEASE, true, MEMORY_ORDER_ACQ_REL },
      { true, MEMORY_ORDER_ACQ_REL, false, MEMORY_ORDER_ACQUIRE },
      { true, MEMORY_ORDER_ACQ_REL, true, MEMORY_ORDER_ACQ_REL }
    };
    switch (MemoryOrder())
    {
    case MEMORY_ORDER_ACQUIRE:
      if (_subCaseId >= sizeof(acqTests) / sizeof(struct TestDefinition))
        return 0;
      _subCase = acqTests[_subCaseId];
      break;
    case MEMORY_ORDER_RELEASE:
      if (_subCaseId >= sizeof(relTests) / sizeof(struct TestDefinition))
        return 0;
      _subCase = relTests[_subCaseId];
      break;
    case MEMORY_ORDER_ACQ_REL:
      if (_subCaseId >= sizeof(arTests) / sizeof(struct TestDefinition))
        return 0;
      _subCase = arTests[_subCaseId];
      break;
    case MEMORY_ORDER_SEQ_CST:
      if (_subCaseId != 0) // one special case only
        return 0;
      break;
    default:
      return 0;
    }
    LocalRefValues(LocalMemory());
    return CBasicTestMemOrderScope<HostAtomicType, HostDataType>::ExecuteSingleTest(deviceID, context, queue);
  }
  virtual std::string ProgramHeader(cl_uint maxNumDestItems)
  {
    std::string header;
    if(gOldAPI)
    {
      if(MemoryScope() == MEMORY_SCOPE_EMPTY)
      {
        header += "#define atomic_work_item_fence(x,y)                        mem_fence(x)\n";
      }
      else
      {
        header += "#define atomic_work_item_fence(x,y,z)                      mem_fence(x)\n";
      }
    }
    return header+CBasicTestMemOrderScope<HostAtomicType, HostDataType>::ProgramHeader(maxNumDestItems);
  }
  virtual std::string ProgramCore()
  {
    std::ostringstream naValues;
    naValues << NumNonAtomicVariablesPerThread();
    std::string program, fenceType, nonAtomic;
    if (LocalMemory())
    {
      program = "  size_t myId = get_local_id(0), hisId = get_local_size(0)-1-myId;\n";
      fenceType = "CLK_LOCAL_MEM_FENCE";
      nonAtomic = "localValues";
    }
    else
    {
      program = "  size_t myId = tid, hisId = threadCount-1-tid;\n";
      fenceType = "CLK_GLOBAL_MEM_FENCE";
      nonAtomic = "oldValues";
    }
    if (MemoryOrder() == MEMORY_ORDER_SEQ_CST)
    {
      // All threads are divided into pairs.
      // Each thread has its own atomic variable and performs the following actions:
      // - increments its own variable
      // - performs fence operation to propagate its value and to see value from other thread
      // - reads value from other thread's variable
      // - repeats the above steps when both values are the same (and less than 1000000)
      // - stores the last value read from other thread (in additional variable)
      // At the end of execution at least one thread should know the last value from other thread
      program += std::string("") +
        "  " + DataType().RegularTypeName() + " myValue = 0, hisValue; \n"
        "  do {\n"
        "    myValue++;\n"
        "    atomic_store_explicit(&destMemory[myId], myValue, memory_order_relaxed" + MemoryScopeStr() + ");\n"
        "    atomic_work_item_fence(" + fenceType + ", memory_order_seq_cst" + MemoryScopeStr() + "); \n"
        "    hisValue = atomic_load_explicit(&destMemory[hisId], memory_order_relaxed" + MemoryScopeStr() + ");\n"
        "  } while(myValue == hisValue && myValue < 1000000);\n"
        "  " + nonAtomic + "[myId] = hisValue; \n";
    }
    else
    {
      // Each thread modifies one of its non-atomic variables, increments value of its atomic variable
      // and reads values from another thread in typical synchronizes-with scenario with:
      // - non-atomic variable (at index A) modification (value change from 0 to A)
      // - release operation (additional fence or within atomic) + atomic variable modification (value A)
      // - atomic variable read (value B) + acquire operation (additional fence or within atomic)
      // - non-atomic variable (at index B) read (value C)
      // Each thread verifies dependency between atomic and non-atomic value read from another thread
      // The following condition must be true: B == C
      program += std::string("") +
        "  " + DataType().RegularTypeName() + " myValue = 0, hisAtomicValue, hisValue; \n"
        "  do {\n"
        "    myValue++;\n"
        "    " + nonAtomic + "[myId*" + naValues.str() +"+myValue] = myValue;\n";
      if (_subCase.op1IsFence)
        program += std::string("") +
        "    atomic_work_item_fence(" + fenceType + ", " + get_memory_order_type_name(_subCase.op1MemOrder) + MemoryScopeStr() + "); \n"
        "    atomic_store_explicit(&destMemory[myId], myValue, memory_order_relaxed" + MemoryScopeStr() + ");\n";
      else
        program += std::string("") +
        "    atomic_store_explicit(&destMemory[myId], myValue, " + get_memory_order_type_name(_subCase.op1MemOrder) + MemoryScopeStr() + ");\n";
      if (_subCase.op2IsFence)
        program += std::string("") +
        "    hisAtomicValue = atomic_load_explicit(&destMemory[hisId], memory_order_relaxed" + MemoryScopeStr() + ");\n"
        "    atomic_work_item_fence(" + fenceType + ", " + get_memory_order_type_name(_subCase.op2MemOrder) + MemoryScopeStr() + "); \n";
      else
        program += std::string("") +
        "    hisAtomicValue = atomic_load_explicit(&destMemory[hisId], " + get_memory_order_type_name(_subCase.op2MemOrder) + MemoryScopeStr() + ");\n";
      program +=
        "    hisValue = " + nonAtomic + "[hisId*" + naValues.str() + "+hisAtomicValue]; \n";
      if (LocalMemory())
        program += "    hisId = (hisId+1)%get_local_size(0);\n";
      else
        program += "    hisId = (hisId+1)%threadCount;\n";
      program +=
        "  } while(hisAtomicValue == hisValue && myValue < "+naValues.str()+"-1);\n"
        "  if(hisAtomicValue != hisValue)\n"
        "  { // fail\n"
        "    atomic_store(&destMemory[myId], myValue-1);\n";
      if (LocalMemory())
        program += "    hisId = (hisId+get_local_size(0)-1)%get_local_size(0);\n";
      else
        program += "    hisId = (hisId+threadCount-1)%threadCount;\n";
      program +=
        "    if(myValue+1 < " + naValues.str() + ")\n"
        "      " + nonAtomic + "[myId*" + naValues.str() + "+myValue+1] = hisId;\n"
        "    if(myValue+2 < " + naValues.str() + ")\n"
        "      " + nonAtomic + "[myId*" + naValues.str() + "+myValue+2] = hisAtomicValue;\n"
        "    if(myValue+3 < " + naValues.str() + ")\n"
        "      " + nonAtomic + "[myId*" + naValues.str() + "+myValue+3] = hisValue;\n";
      if (gDebug)
      {
        program +=
          "    printf(\"WI %d: atomic value (%d) at index %d is different than non-atomic value (%d)\\n\", tid, hisAtomicValue, hisId, hisValue);\n";
      }
      program +=
        "  }\n";
    }
    return program;
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    size_t myId = tid, hisId = threadCount - 1 - tid;
    if (MemoryOrder() == MEMORY_ORDER_SEQ_CST)
    {
      HostDataType myValue = 0, hisValue;
      // CPU thread typically starts faster - wait for GPU thread
      myValue++;
      host_atomic_store<HostAtomicType, HostDataType>(&destMemory[myId], myValue, MEMORY_ORDER_SEQ_CST);
      while (host_atomic_load<HostAtomicType, HostDataType>(&destMemory[hisId], MEMORY_ORDER_SEQ_CST) == 0);
      do {
        myValue++;
        host_atomic_store<HostAtomicType, HostDataType>(&destMemory[myId], myValue, MEMORY_ORDER_RELAXED);
        host_atomic_thread_fence(MemoryOrder());
        hisValue = host_atomic_load<HostAtomicType, HostDataType>(&destMemory[hisId], MEMORY_ORDER_RELAXED);
      } while (myValue == hisValue && hisValue < 1000000);
      oldValues[tid] = hisValue;
    }
    else
    {
      HostDataType myValue = 0, hisAtomicValue, hisValue;
      do {
        myValue++;
        oldValues[myId*NumNonAtomicVariablesPerThread()+myValue] = myValue;
        if (_subCase.op1IsFence)
        {
          host_atomic_thread_fence(_subCase.op1MemOrder);
          host_atomic_store<HostAtomicType, HostDataType>(&destMemory[myId], myValue, MEMORY_ORDER_RELAXED);
        }
        else
          host_atomic_store<HostAtomicType, HostDataType>(&destMemory[myId], myValue, _subCase.op1MemOrder);
        if (_subCase.op2IsFence)
        {
          hisAtomicValue = host_atomic_load<HostAtomicType, HostDataType>(&destMemory[hisId], MEMORY_ORDER_RELAXED);
          host_atomic_thread_fence(_subCase.op2MemOrder);
        }
        else
          hisAtomicValue = host_atomic_load<HostAtomicType, HostDataType>(&destMemory[hisId], _subCase.op2MemOrder);
        hisValue = oldValues[hisId*NumNonAtomicVariablesPerThread() + hisAtomicValue];
        hisId = (hisId + 1) % threadCount;
      } while(hisAtomicValue == hisValue && myValue < (HostDataType)NumNonAtomicVariablesPerThread()-1);
      if(hisAtomicValue != hisValue)
      { // fail
        host_atomic_store<HostAtomicType, HostDataType>(&destMemory[myId], myValue-1, MEMORY_ORDER_SEQ_CST);
        if (gDebug)
        {
          hisId = (hisId + threadCount - 1) % threadCount;
          printf("WI %d: atomic value (%d) at index %d is different than non-atomic value (%d)\n", tid, hisAtomicValue, hisId, hisValue);
        }
      }
    }
  }
  virtual bool GenerateRefs(cl_uint threadCount, HostDataType *startRefValues, MTdata d)
  {
    for(cl_uint i = 0 ; i < threadCount*NumNonAtomicVariablesPerThread(); i++)
      startRefValues[i] = 0;
    return true;
  }
  virtual bool VerifyRefs(bool &correct, cl_uint threadCount, HostDataType *refValues, HostAtomicType *finalValues)
  {
    correct = true;
    cl_uint workSize = LocalMemory() ? CurrentGroupSize() : threadCount;
    for(cl_uint workOffset = 0; workOffset < threadCount; workOffset+= workSize)
    {
      if(workOffset+workSize > threadCount)
        // last workgroup (host threads)
        workSize = threadCount-workOffset;
      for(cl_uint i = 0 ; i < workSize && workOffset+i < threadCount; i++)
      {
        HostAtomicType myValue = finalValues[workOffset + i];
        if (MemoryOrder() == MEMORY_ORDER_SEQ_CST)
        {
          HostDataType hisValue = refValues[workOffset + i];
          if (myValue == hisValue)
          {
            // a draw - both threads should reach final value 1000000
            if (myValue != 1000000)
            {
              log_error("ERROR: Invalid reference value #%u (%d instead of 1000000)\n", workOffset + i, myValue);
              correct = false;
              return true;
            }
          }
          else
          {
            //slower thread (in total order of seq_cst operations) must know last value written by faster thread
            HostAtomicType hisRealValue = finalValues[workOffset + workSize - 1 - i];
            HostDataType myValueReadByHim = refValues[workOffset + workSize - 1 - i];

            // who is the winner? - thread with lower private counter value
            if (myValue == hisRealValue) // forbidden result - fence doesn't work
            {
              log_error("ERROR: Atomic counter values #%u and #%u are the same (%u)\n", workOffset + i, workOffset + workSize - 1 - i, myValue);
              log_error("ERROR: Both threads have outdated values read from another thread (%u and %u)\n", hisValue, myValueReadByHim);
              correct = false;
              return true;
            }
            if (myValue > hisRealValue) // I'm slower
            {
              if (hisRealValue != hisValue)
              {
                log_error("ERROR: Invalid reference value #%u (%d instead of %d)\n", workOffset + i, hisValue, hisRealValue);
                log_error("ERROR: Slower thread #%u should know value written by faster thread #%u\n", workOffset + i, workOffset + workSize - 1 - i);
                correct = false;
                return true;
              }
            }
            else // I'm faster
            {
              if (myValueReadByHim != myValue)
              {
                log_error("ERROR: Invalid reference value #%u (%d instead of %d)\n", workOffset + workSize - 1 - i, myValueReadByHim, myValue);
                log_error("ERROR: Slower thread #%u should know value written by faster thread #%u\n", workOffset + workSize - 1 - i, workOffset + i);
                correct = false;
                return true;
              }
            }
          }
        }
        else
        {
          if (myValue != NumNonAtomicVariablesPerThread()-1)
          {
            log_error("ERROR: Invalid atomic value #%u (%d instead of %d)\n", workOffset + i, myValue, NumNonAtomicVariablesPerThread()-1);
            log_error("ERROR: Thread #%u observed invalid values in other thread's variables\n", workOffset + i, myValue);
            correct = false;
            return true;
          }
        }
      }
    }
    return true;
  }
private:
  int _subCaseId;
  struct TestDefinition _subCase;
};

int test_atomic_fence_generic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, bool useSVM)
{
  int error = 0;
  CBasicTestFence<HOST_ATOMIC_INT, HOST_INT> test_int(TYPE_ATOMIC_INT, useSVM);
  EXECUTE_TEST(error, test_int.Execute(deviceID, context, queue, num_elements));
  CBasicTestFence<HOST_ATOMIC_UINT, HOST_UINT> test_uint(TYPE_ATOMIC_UINT, useSVM);
  EXECUTE_TEST(error, test_uint.Execute(deviceID, context, queue, num_elements));
  CBasicTestFence<HOST_ATOMIC_LONG, HOST_LONG> test_long(TYPE_ATOMIC_LONG, useSVM);
  EXECUTE_TEST(error, test_long.Execute(deviceID, context, queue, num_elements));
  CBasicTestFence<HOST_ATOMIC_ULONG, HOST_ULONG> test_ulong(TYPE_ATOMIC_ULONG, useSVM);
  EXECUTE_TEST(error, test_ulong.Execute(deviceID, context, queue, num_elements));
  if (AtomicTypeInfo(TYPE_ATOMIC_SIZE_T).Size(deviceID) == 4)
  {
    CBasicTestFence<HOST_ATOMIC_INTPTR_T32, HOST_INTPTR_T32> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFence<HOST_ATOMIC_UINTPTR_T32, HOST_UINTPTR_T32> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFence<HOST_ATOMIC_SIZE_T32, HOST_SIZE_T32> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFence<HOST_ATOMIC_PTRDIFF_T32, HOST_PTRDIFF_T32> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  else
  {
    CBasicTestFence<HOST_ATOMIC_INTPTR_T64, HOST_INTPTR_T64> test_intptr_t(TYPE_ATOMIC_INTPTR_T, useSVM);
    EXECUTE_TEST(error, test_intptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFence<HOST_ATOMIC_UINTPTR_T64, HOST_UINTPTR_T64> test_uintptr_t(TYPE_ATOMIC_UINTPTR_T, useSVM);
    EXECUTE_TEST(error, test_uintptr_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFence<HOST_ATOMIC_SIZE_T64, HOST_SIZE_T64> test_size_t(TYPE_ATOMIC_SIZE_T, useSVM);
    EXECUTE_TEST(error, test_size_t.Execute(deviceID, context, queue, num_elements));
    CBasicTestFence<HOST_ATOMIC_PTRDIFF_T64, HOST_PTRDIFF_T64> test_ptrdiff_t(TYPE_ATOMIC_PTRDIFF_T, useSVM);
    EXECUTE_TEST(error, test_ptrdiff_t.Execute(deviceID, context, queue, num_elements));
  }
  return error;
}

int test_atomic_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fence_generic(deviceID, context, queue, num_elements, false);
}

int test_svm_atomic_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
  return test_atomic_fence_generic(deviceID, context, queue, num_elements, true);
}