Skip to content

Instantly share code, notes, and snippets.

@karolherbst
Created June 16, 2023 20:41
Show Gist options
  • Save karolherbst/2965158c53342bb7b3f87b85d97f07e2 to your computer and use it in GitHub Desktop.
Save karolherbst/2965158c53342bb7b3f87b85d97f07e2 to your computer and use it in GitHub Desktop.
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Linker; 0
; Bound: 1098
; Schema: 0
OpCapability Addresses
OpCapability Kernel
OpCapability Int64
OpCapability GenericPointer
OpCapability Int8
OpCapability ExpectAssumeKHR
OpExtension "SPV_KHR_expect_assume"
%1 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInLocalInvocationId
OpEntryPoint Kernel %3 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInLocalInvocationId
OpEntryPoint Kernel %4 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInLocalInvocationId
OpEntryPoint Kernel %5 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInLocalInvocationId
OpEntryPoint Kernel %6 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInGlobalInvocationId
OpEntryPoint Kernel %7 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInLocalInvocationId
OpEntryPoint Kernel %8 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInLocalInvocationId
OpEntryPoint Kernel %9 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInLocalInvocationId
OpEntryPoint Kernel %10 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInGlobalInvocationId
OpEntryPoint Kernel %11 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInGlobalInvocationId
OpEntryPoint Kernel %12 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInLocalInvocationId
OpEntryPoint Kernel %13 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInGlobalInvocationId
OpEntryPoint Kernel %14 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInGlobalInvocationId
OpEntryPoint Kernel %15 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInGlobalInvocationId
OpEntryPoint Kernel %16 "_ZTSZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_EUlNS3_7nd_itemILi1EEEE_" %__spirv_BuiltInWorkgroupSize %__spirv_BuiltInGlobalLinearId %__spirv_BuiltInWorkgroupId %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalOffset %__spirv_BuiltInLocalInvocationId
OpSource OpenCL_CPP 100000
OpName %__spirv_BuiltInGlobalOffset "__spirv_BuiltInGlobalOffset"
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId"
OpName %__spirv_BuiltInWorkgroupId "__spirv_BuiltInWorkgroupId"
OpName %__spirv_BuiltInGlobalLinearId "__spirv_BuiltInGlobalLinearId"
OpName %__spirv_BuiltInWorkgroupSize "__spirv_BuiltInWorkgroupSize"
OpName %class_sycl___V1__id "class.sycl::_V1::id"
OpName %class_sycl___V1__detail__array "class.sycl::_V1::detail::array"
OpName %_arg_sync_flag_acc "_arg_sync_flag_acc"
OpName %_arg_data_acc "_arg_data_acc"
OpName %_arg_value "_arg_value"
OpName %_arg_order_write "_arg_order_write"
OpName %_arg_order_read "_arg_order_read"
OpName %_arg_res_acc "_arg_res_acc"
OpName %_arg_res_acc9 "_arg_res_acc9"
OpName %entry "entry"
OpName %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i "_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE.exit.i"
OpName %if_then_i "if.then.i"
OpName %sw_bb1_i_i_i "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i "for.cond.i"
OpName %for_body_i "for.body.i"
OpName %for_inc_i "for.inc.i"
OpName %cleanup_i "cleanup.i"
OpName %sw_bb1_i_i22_i "sw.bb1.i.i22.i"
OpName %sw_bb2_i_i23_i "sw.bb2.i.i23.i"
OpName %sw_bb3_i_i24_i "sw.bb3.i.i24.i"
OpName %sw_bb4_i_i25_i "sw.bb4.i.i25.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit27.i"
OpName %if_then8_i "if.then8.i"
OpName %if_then11_i "if.then11.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %__itt_offload_wi_start_wrapper "__itt_offload_wi_start_wrapper"
OpName %add_ptr_i "add.ptr.i"
OpName %arrayidx_ascast_i_i "arrayidx.ascast.i.i"
OpName %cmp_i5_i_i "cmp.i5.i.i"
OpName %sub_i_i_i_i_i "sub.i.i.i.i.i"
OpName %cmp_i_i_i "cmp.i.i.i"
OpName %SpvOrder_0_i_i_i "SpvOrder.0.i.i.i"
OpName %__itt_offload_atomic_op_start "__itt_offload_atomic_op_start"
OpName %object "object"
OpName %op_type "op_type"
OpName %mem_order "mem_order"
OpName %__itt_offload_atomic_op_finish "__itt_offload_atomic_op_finish"
OpName %object_0 "object"
OpName %op_type_0 "op_type"
OpName %mem_order_0 "mem_order"
OpName %inc_i "inc.i"
OpName %i_0_i "i.0.i"
OpName %cmp_i "cmp.i"
OpName %call3_i_i_i_i "call3.i.i.i.i"
OpName %cmp6_i "cmp6.i"
OpName %SpvOrder_0_i_i26_i "SpvOrder.0.i.i26.i"
OpName %cmp10_not_i "cmp10.not.i"
OpName %__itt_offload_wi_finish_wrapper "__itt_offload_wi_finish_wrapper"
OpName %_arg_sync_flag_acc_0 "_arg_sync_flag_acc"
OpName %_arg_data_acc_0 "_arg_data_acc"
OpName %_arg_value_0 "_arg_value"
OpName %_arg_order_write_0 "_arg_order_write"
OpName %_arg_order_read_0 "_arg_order_read"
OpName %_arg_res_acc_0 "_arg_res_acc"
OpName %_arg_res_acc9_0 "_arg_res_acc9"
OpName %entry_0 "entry"
OpName %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i "_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE.exit.i"
OpName %if_then_i_0 "if.then.i"
OpName %sw_bb1_i_i_i_0 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_0 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_0 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_0 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_0 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_0 "for.cond.i"
OpName %for_body_i_0 "for.body.i"
OpName %for_inc_i_0 "for.inc.i"
OpName %cleanup_i_0 "cleanup.i"
OpName %sw_bb1_i_i22_i_0 "sw.bb1.i.i22.i"
OpName %sw_bb2_i_i23_i_0 "sw.bb2.i.i23.i"
OpName %sw_bb3_i_i24_i_0 "sw.bb3.i.i24.i"
OpName %sw_bb4_i_i25_i_0 "sw.bb4.i.i25.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_0 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit27.i"
OpName %if_then8_i_0 "if.then8.i"
OpName %if_then11_i_0 "if.then11.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_0 "add.ptr.i"
OpName %arrayidx_ascast_i_i_0 "arrayidx.ascast.i.i"
OpName %cmp_i5_i_i_0 "cmp.i5.i.i"
OpName %sub_i_i_i_i_i_0 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_0 "cmp.i.i.i"
OpName %SpvOrder_0_i_i_i_0 "SpvOrder.0.i.i.i"
OpName %inc_i_0 "inc.i"
OpName %i_0_i_0 "i.0.i"
OpName %cmp_i_0 "cmp.i"
OpName %call3_i_i_i_i_0 "call3.i.i.i.i"
OpName %cmp6_i_0 "cmp6.i"
OpName %SpvOrder_0_i_i26_i_0 "SpvOrder.0.i.i26.i"
OpName %cmp10_not_i_0 "cmp10.not.i"
OpName %_arg_sync_flag_acc_1 "_arg_sync_flag_acc"
OpName %_arg_data_acc_1 "_arg_data_acc"
OpName %_arg_value_1 "_arg_value"
OpName %_arg_order_write_1 "_arg_order_write"
OpName %_arg_order_read_1 "_arg_order_read"
OpName %_arg_res_acc_1 "_arg_res_acc"
OpName %_arg_res_acc9_1 "_arg_res_acc9"
OpName %entry_1 "entry"
OpName %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i "_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE.exit.i"
OpName %if_then_i_1 "if.then.i"
OpName %sw_bb1_i_i_i_1 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_1 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_1 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_1 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_1 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_1 "for.cond.i"
OpName %for_body_i_1 "for.body.i"
OpName %for_inc_i_1 "for.inc.i"
OpName %cleanup_i_1 "cleanup.i"
OpName %sw_bb1_i_i22_i_1 "sw.bb1.i.i22.i"
OpName %sw_bb2_i_i23_i_1 "sw.bb2.i.i23.i"
OpName %sw_bb3_i_i24_i_1 "sw.bb3.i.i24.i"
OpName %sw_bb4_i_i25_i_1 "sw.bb4.i.i25.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_1 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit27.i"
OpName %if_then8_i_1 "if.then8.i"
OpName %if_then11_i_1 "if.then11.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_1 "add.ptr.i"
OpName %arrayidx_ascast_i_i_1 "arrayidx.ascast.i.i"
OpName %cmp_i5_i_i_1 "cmp.i5.i.i"
OpName %sub_i_i_i_i_i_1 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_1 "cmp.i.i.i"
OpName %SpvOrder_0_i_i_i_1 "SpvOrder.0.i.i.i"
OpName %inc_i_1 "inc.i"
OpName %i_0_i_1 "i.0.i"
OpName %cmp_i_1 "cmp.i"
OpName %call3_i_i_i_i_1 "call3.i.i.i.i"
OpName %cmp6_i_1 "cmp6.i"
OpName %SpvOrder_0_i_i26_i_1 "SpvOrder.0.i.i26.i"
OpName %cmp10_not_i_1 "cmp10.not.i"
OpName %_arg_sync_flag_acc_2 "_arg_sync_flag_acc"
OpName %_arg_data_acc_2 "_arg_data_acc"
OpName %_arg_value_2 "_arg_value"
OpName %_arg_order_write_2 "_arg_order_write"
OpName %_arg_order_read_2 "_arg_order_read"
OpName %_arg_res_acc_2 "_arg_res_acc"
OpName %_arg_res_acc9_2 "_arg_res_acc9"
OpName %entry_2 "entry"
OpName %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i "_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE.exit.i"
OpName %if_then_i_2 "if.then.i"
OpName %sw_bb1_i_i_i_2 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_2 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_2 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_2 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_2 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_2 "for.cond.i"
OpName %for_body_i_2 "for.body.i"
OpName %for_inc_i_2 "for.inc.i"
OpName %cleanup_i_2 "cleanup.i"
OpName %sw_bb1_i_i22_i_2 "sw.bb1.i.i22.i"
OpName %sw_bb2_i_i23_i_2 "sw.bb2.i.i23.i"
OpName %sw_bb3_i_i24_i_2 "sw.bb3.i.i24.i"
OpName %sw_bb4_i_i25_i_2 "sw.bb4.i.i25.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_2 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit27.i"
OpName %if_then8_i_2 "if.then8.i"
OpName %if_then11_i_2 "if.then11.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_2 "add.ptr.i"
OpName %arrayidx_ascast_i_i_2 "arrayidx.ascast.i.i"
OpName %cmp_i5_i_i_2 "cmp.i5.i.i"
OpName %sub_i_i_i_i_i_2 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_2 "cmp.i.i.i"
OpName %SpvOrder_0_i_i_i_2 "SpvOrder.0.i.i.i"
OpName %inc_i_2 "inc.i"
OpName %i_0_i_2 "i.0.i"
OpName %cmp_i_2 "cmp.i"
OpName %call3_i_i_i_i_2 "call3.i.i.i.i"
OpName %cmp6_i_2 "cmp6.i"
OpName %SpvOrder_0_i_i26_i_2 "SpvOrder.0.i.i26.i"
OpName %cmp10_not_i_2 "cmp10.not.i"
OpName %_arg_sync_flag_acc_3 "_arg_sync_flag_acc"
OpName %_arg_sync_flag_acc3 "_arg_sync_flag_acc3"
OpName %_arg_data_acc_3 "_arg_data_acc"
OpName %_arg_data_acc6 "_arg_data_acc6"
OpName %_arg_value_3 "_arg_value"
OpName %_arg_order_write_3 "_arg_order_write"
OpName %_arg_order_read_3 "_arg_order_read"
OpName %_arg_res_acc_3 "_arg_res_acc"
OpName %_arg_res_acc9_3 "_arg_res_acc9"
OpName %entry_3 "entry"
OpName %for_cond_i_preheader "for.cond.i.preheader"
OpName %if_then_i_3 "if.then.i"
OpName %sw_bb1_i_i_i_3 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_3 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_3 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_3 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_3 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_3 "for.cond.i"
OpName %for_body_i_3 "for.body.i"
OpName %for_inc_i_3 "for.inc.i"
OpName %cleanup_i_3 "cleanup.i"
OpName %sw_bb1_i_i28_i "sw.bb1.i.i28.i"
OpName %sw_bb2_i_i29_i "sw.bb2.i.i29.i"
OpName %sw_bb3_i_i30_i "sw.bb3.i.i30.i"
OpName %sw_bb4_i_i31_i "sw.bb4.i.i31.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit33.i"
OpName %if_then9_i "if.then9.i"
OpName %if_then12_i "if.then12.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_3 "add.ptr.i"
OpName %add_ptr_i39 "add.ptr.i39"
OpName %add_ptr_i53 "add.ptr.i53"
OpName %arrayidx_ascast_i_i_3 "arrayidx.ascast.i.i"
OpName %sub_i_i_i_i_i_3 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_3 "cmp.i.i.i"
OpName %cmp_i_i "cmp.i.i"
OpName %SpvOrder_0_i_i_i_3 "SpvOrder.0.i.i.i"
OpName %inc_i_3 "inc.i"
OpName %i_0_i_3 "i.0.i"
OpName %cmp_i_3 "cmp.i"
OpName %call3_i_i_i_i_3 "call3.i.i.i.i"
OpName %cmp7_i "cmp7.i"
OpName %SpvOrder_0_i_i32_i "SpvOrder.0.i.i32.i"
OpName %cmp11_not_i "cmp11.not.i"
OpName %_arg_sync_flag_acc_4 "_arg_sync_flag_acc"
OpName %_arg_data_acc_4 "_arg_data_acc"
OpName %_arg_value_4 "_arg_value"
OpName %_arg_order_write_4 "_arg_order_write"
OpName %_arg_order_read_4 "_arg_order_read"
OpName %_arg_res_acc_4 "_arg_res_acc"
OpName %_arg_res_acc9_4 "_arg_res_acc9"
OpName %entry_4 "entry"
OpName %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i "_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE.exit.i"
OpName %if_then_i_4 "if.then.i"
OpName %sw_bb1_i_i_i_4 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_4 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_4 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_4 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_4 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_4 "for.cond.i"
OpName %for_body_i_4 "for.body.i"
OpName %for_inc_i_4 "for.inc.i"
OpName %cleanup_i_4 "cleanup.i"
OpName %sw_bb1_i_i22_i_3 "sw.bb1.i.i22.i"
OpName %sw_bb2_i_i23_i_3 "sw.bb2.i.i23.i"
OpName %sw_bb3_i_i24_i_3 "sw.bb3.i.i24.i"
OpName %sw_bb4_i_i25_i_3 "sw.bb4.i.i25.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_3 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit27.i"
OpName %if_then8_i_3 "if.then8.i"
OpName %if_then11_i_3 "if.then11.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_4 "add.ptr.i"
OpName %arrayidx_ascast_i_i_4 "arrayidx.ascast.i.i"
OpName %cmp_i5_i_i_3 "cmp.i5.i.i"
OpName %sub_i_i_i_i_i_4 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_4 "cmp.i.i.i"
OpName %SpvOrder_0_i_i_i_4 "SpvOrder.0.i.i.i"
OpName %inc_i_4 "inc.i"
OpName %i_0_i_4 "i.0.i"
OpName %cmp_i_4 "cmp.i"
OpName %call3_i_i_i_i_4 "call3.i.i.i.i"
OpName %cmp6_i_3 "cmp6.i"
OpName %SpvOrder_0_i_i26_i_3 "SpvOrder.0.i.i26.i"
OpName %cmp10_not_i_3 "cmp10.not.i"
OpName %_arg_sync_flag_acc_5 "_arg_sync_flag_acc"
OpName %_arg_data_acc_5 "_arg_data_acc"
OpName %_arg_value_5 "_arg_value"
OpName %_arg_order_write_5 "_arg_order_write"
OpName %_arg_order_read_5 "_arg_order_read"
OpName %_arg_res_acc_5 "_arg_res_acc"
OpName %_arg_res_acc9_5 "_arg_res_acc9"
OpName %entry_5 "entry"
OpName %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i "_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE.exit.i"
OpName %if_then_i_5 "if.then.i"
OpName %sw_bb1_i_i_i_5 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_5 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_5 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_5 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_5 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_5 "for.cond.i"
OpName %for_body_i_5 "for.body.i"
OpName %for_inc_i_5 "for.inc.i"
OpName %cleanup_i_5 "cleanup.i"
OpName %sw_bb1_i_i22_i_4 "sw.bb1.i.i22.i"
OpName %sw_bb2_i_i23_i_4 "sw.bb2.i.i23.i"
OpName %sw_bb3_i_i24_i_4 "sw.bb3.i.i24.i"
OpName %sw_bb4_i_i25_i_4 "sw.bb4.i.i25.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_4 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit27.i"
OpName %if_then8_i_4 "if.then8.i"
OpName %if_then11_i_4 "if.then11.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_5 "add.ptr.i"
OpName %arrayidx_ascast_i_i_5 "arrayidx.ascast.i.i"
OpName %cmp_i5_i_i_4 "cmp.i5.i.i"
OpName %sub_i_i_i_i_i_5 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_5 "cmp.i.i.i"
OpName %SpvOrder_0_i_i_i_5 "SpvOrder.0.i.i.i"
OpName %inc_i_5 "inc.i"
OpName %i_0_i_5 "i.0.i"
OpName %cmp_i_5 "cmp.i"
OpName %call3_i_i_i_i_5 "call3.i.i.i.i"
OpName %cmp6_i_4 "cmp6.i"
OpName %SpvOrder_0_i_i26_i_4 "SpvOrder.0.i.i26.i"
OpName %cmp10_not_i_4 "cmp10.not.i"
OpName %_arg_sync_flag_acc_6 "_arg_sync_flag_acc"
OpName %_arg_data_acc_6 "_arg_data_acc"
OpName %_arg_value_6 "_arg_value"
OpName %_arg_order_write_6 "_arg_order_write"
OpName %_arg_order_read_6 "_arg_order_read"
OpName %_arg_res_acc_6 "_arg_res_acc"
OpName %_arg_res_acc9_6 "_arg_res_acc9"
OpName %entry_6 "entry"
OpName %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i "_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE.exit.i"
OpName %if_then_i_6 "if.then.i"
OpName %sw_bb1_i_i_i_6 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_6 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_6 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_6 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_6 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_6 "for.cond.i"
OpName %for_body_i_6 "for.body.i"
OpName %for_inc_i_6 "for.inc.i"
OpName %cleanup_i_6 "cleanup.i"
OpName %sw_bb1_i_i22_i_5 "sw.bb1.i.i22.i"
OpName %sw_bb2_i_i23_i_5 "sw.bb2.i.i23.i"
OpName %sw_bb3_i_i24_i_5 "sw.bb3.i.i24.i"
OpName %sw_bb4_i_i25_i_5 "sw.bb4.i.i25.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_5 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit27.i"
OpName %if_then8_i_5 "if.then8.i"
OpName %if_then11_i_5 "if.then11.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_6 "add.ptr.i"
OpName %arrayidx_ascast_i_i_6 "arrayidx.ascast.i.i"
OpName %cmp_i5_i_i_5 "cmp.i5.i.i"
OpName %sub_i_i_i_i_i_6 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_6 "cmp.i.i.i"
OpName %SpvOrder_0_i_i_i_6 "SpvOrder.0.i.i.i"
OpName %inc_i_6 "inc.i"
OpName %i_0_i_6 "i.0.i"
OpName %cmp_i_6 "cmp.i"
OpName %call3_i_i_i_i_6 "call3.i.i.i.i"
OpName %cmp6_i_5 "cmp6.i"
OpName %SpvOrder_0_i_i26_i_5 "SpvOrder.0.i.i26.i"
OpName %cmp10_not_i_5 "cmp10.not.i"
OpName %_arg_sync_flag_acc_7 "_arg_sync_flag_acc"
OpName %_arg_sync_flag_acc3_0 "_arg_sync_flag_acc3"
OpName %_arg_data_acc_7 "_arg_data_acc"
OpName %_arg_data_acc6_0 "_arg_data_acc6"
OpName %_arg_value_7 "_arg_value"
OpName %_arg_order_write_7 "_arg_order_write"
OpName %_arg_order_read_7 "_arg_order_read"
OpName %_arg_res_acc_7 "_arg_res_acc"
OpName %_arg_res_acc9_7 "_arg_res_acc9"
OpName %entry_7 "entry"
OpName %for_cond_i_preheader_0 "for.cond.i.preheader"
OpName %if_then_i_7 "if.then.i"
OpName %sw_bb1_i_i_i_7 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_7 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_7 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_7 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_7 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_7 "for.cond.i"
OpName %for_body_i_7 "for.body.i"
OpName %for_inc_i_7 "for.inc.i"
OpName %cleanup_i_7 "cleanup.i"
OpName %sw_bb1_i_i28_i_0 "sw.bb1.i.i28.i"
OpName %sw_bb2_i_i29_i_0 "sw.bb2.i.i29.i"
OpName %sw_bb3_i_i30_i_0 "sw.bb3.i.i30.i"
OpName %sw_bb4_i_i31_i_0 "sw.bb4.i.i31.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_0 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit33.i"
OpName %if_then9_i_0 "if.then9.i"
OpName %if_then12_i_0 "if.then12.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_7 "add.ptr.i"
OpName %add_ptr_i39_0 "add.ptr.i39"
OpName %add_ptr_i53_0 "add.ptr.i53"
OpName %arrayidx_ascast_i_i_7 "arrayidx.ascast.i.i"
OpName %sub_i_i_i_i_i_7 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_7 "cmp.i.i.i"
OpName %cmp_i_i_0 "cmp.i.i"
OpName %SpvOrder_0_i_i_i_7 "SpvOrder.0.i.i.i"
OpName %inc_i_7 "inc.i"
OpName %i_0_i_7 "i.0.i"
OpName %cmp_i_7 "cmp.i"
OpName %call3_i_i_i_i_7 "call3.i.i.i.i"
OpName %cmp7_i_0 "cmp7.i"
OpName %SpvOrder_0_i_i32_i_0 "SpvOrder.0.i.i32.i"
OpName %cmp11_not_i_0 "cmp11.not.i"
OpName %_arg_sync_flag_acc_8 "_arg_sync_flag_acc"
OpName %_arg_sync_flag_acc3_1 "_arg_sync_flag_acc3"
OpName %_arg_data_acc_8 "_arg_data_acc"
OpName %_arg_data_acc6_1 "_arg_data_acc6"
OpName %_arg_value_8 "_arg_value"
OpName %_arg_order_write_8 "_arg_order_write"
OpName %_arg_order_read_8 "_arg_order_read"
OpName %_arg_res_acc_8 "_arg_res_acc"
OpName %_arg_res_acc9_8 "_arg_res_acc9"
OpName %entry_8 "entry"
OpName %for_cond_i_preheader_1 "for.cond.i.preheader"
OpName %if_then_i_8 "if.then.i"
OpName %sw_bb1_i_i_i_8 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_8 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_8 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_8 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_8 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_8 "for.cond.i"
OpName %for_body_i_8 "for.body.i"
OpName %for_inc_i_8 "for.inc.i"
OpName %cleanup_i_8 "cleanup.i"
OpName %sw_bb1_i_i28_i_1 "sw.bb1.i.i28.i"
OpName %sw_bb2_i_i29_i_1 "sw.bb2.i.i29.i"
OpName %sw_bb3_i_i30_i_1 "sw.bb3.i.i30.i"
OpName %sw_bb4_i_i31_i_1 "sw.bb4.i.i31.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_1 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit33.i"
OpName %if_then9_i_1 "if.then9.i"
OpName %if_then12_i_1 "if.then12.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_8 "add.ptr.i"
OpName %add_ptr_i39_1 "add.ptr.i39"
OpName %add_ptr_i53_1 "add.ptr.i53"
OpName %arrayidx_ascast_i_i_8 "arrayidx.ascast.i.i"
OpName %sub_i_i_i_i_i_8 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_8 "cmp.i.i.i"
OpName %cmp_i_i_1 "cmp.i.i"
OpName %SpvOrder_0_i_i_i_8 "SpvOrder.0.i.i.i"
OpName %inc_i_8 "inc.i"
OpName %i_0_i_8 "i.0.i"
OpName %cmp_i_8 "cmp.i"
OpName %call3_i_i_i_i_8 "call3.i.i.i.i"
OpName %cmp7_i_1 "cmp7.i"
OpName %SpvOrder_0_i_i32_i_1 "SpvOrder.0.i.i32.i"
OpName %cmp11_not_i_1 "cmp11.not.i"
OpName %_arg_sync_flag_acc_9 "_arg_sync_flag_acc"
OpName %_arg_data_acc_9 "_arg_data_acc"
OpName %_arg_value_9 "_arg_value"
OpName %_arg_order_write_9 "_arg_order_write"
OpName %_arg_order_read_9 "_arg_order_read"
OpName %_arg_res_acc_9 "_arg_res_acc"
OpName %_arg_res_acc9_9 "_arg_res_acc9"
OpName %entry_9 "entry"
OpName %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i "_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE.exit.i"
OpName %if_then_i_9 "if.then.i"
OpName %sw_bb1_i_i_i_9 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_9 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_9 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_9 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_9 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_9 "for.cond.i"
OpName %for_body_i_9 "for.body.i"
OpName %for_inc_i_9 "for.inc.i"
OpName %cleanup_i_9 "cleanup.i"
OpName %sw_bb1_i_i22_i_6 "sw.bb1.i.i22.i"
OpName %sw_bb2_i_i23_i_6 "sw.bb2.i.i23.i"
OpName %sw_bb3_i_i24_i_6 "sw.bb3.i.i24.i"
OpName %sw_bb4_i_i25_i_6 "sw.bb4.i.i25.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_6 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit27.i"
OpName %if_then8_i_6 "if.then8.i"
OpName %if_then11_i_6 "if.then11.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_9 "add.ptr.i"
OpName %arrayidx_ascast_i_i_9 "arrayidx.ascast.i.i"
OpName %cmp_i5_i_i_6 "cmp.i5.i.i"
OpName %sub_i_i_i_i_i_9 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_9 "cmp.i.i.i"
OpName %SpvOrder_0_i_i_i_9 "SpvOrder.0.i.i.i"
OpName %inc_i_9 "inc.i"
OpName %i_0_i_9 "i.0.i"
OpName %cmp_i_9 "cmp.i"
OpName %call3_i_i_i_i_9 "call3.i.i.i.i"
OpName %cmp6_i_6 "cmp6.i"
OpName %SpvOrder_0_i_i26_i_6 "SpvOrder.0.i.i26.i"
OpName %cmp10_not_i_6 "cmp10.not.i"
OpName %_arg_sync_flag_acc_10 "_arg_sync_flag_acc"
OpName %_arg_sync_flag_acc3_2 "_arg_sync_flag_acc3"
OpName %_arg_data_acc_10 "_arg_data_acc"
OpName %_arg_data_acc6_2 "_arg_data_acc6"
OpName %_arg_value_10 "_arg_value"
OpName %_arg_order_write_10 "_arg_order_write"
OpName %_arg_order_read_10 "_arg_order_read"
OpName %_arg_res_acc_10 "_arg_res_acc"
OpName %_arg_res_acc9_10 "_arg_res_acc9"
OpName %entry_10 "entry"
OpName %for_cond_i_preheader_2 "for.cond.i.preheader"
OpName %if_then_i_10 "if.then.i"
OpName %sw_bb1_i_i_i_10 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_10 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_10 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_10 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_10 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_10 "for.cond.i"
OpName %for_body_i_10 "for.body.i"
OpName %for_inc_i_10 "for.inc.i"
OpName %cleanup_i_10 "cleanup.i"
OpName %sw_bb1_i_i28_i_2 "sw.bb1.i.i28.i"
OpName %sw_bb2_i_i29_i_2 "sw.bb2.i.i29.i"
OpName %sw_bb3_i_i30_i_2 "sw.bb3.i.i30.i"
OpName %sw_bb4_i_i31_i_2 "sw.bb4.i.i31.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_2 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit33.i"
OpName %if_then9_i_2 "if.then9.i"
OpName %if_then12_i_2 "if.then12.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_10 "add.ptr.i"
OpName %add_ptr_i39_2 "add.ptr.i39"
OpName %add_ptr_i53_2 "add.ptr.i53"
OpName %arrayidx_ascast_i_i_10 "arrayidx.ascast.i.i"
OpName %sub_i_i_i_i_i_10 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_10 "cmp.i.i.i"
OpName %cmp_i_i_2 "cmp.i.i"
OpName %SpvOrder_0_i_i_i_10 "SpvOrder.0.i.i.i"
OpName %inc_i_10 "inc.i"
OpName %i_0_i_10 "i.0.i"
OpName %cmp_i_10 "cmp.i"
OpName %call3_i_i_i_i_10 "call3.i.i.i.i"
OpName %cmp7_i_2 "cmp7.i"
OpName %SpvOrder_0_i_i32_i_2 "SpvOrder.0.i.i32.i"
OpName %cmp11_not_i_2 "cmp11.not.i"
OpName %_arg_sync_flag_acc_11 "_arg_sync_flag_acc"
OpName %_arg_sync_flag_acc3_3 "_arg_sync_flag_acc3"
OpName %_arg_data_acc_11 "_arg_data_acc"
OpName %_arg_data_acc6_3 "_arg_data_acc6"
OpName %_arg_value_11 "_arg_value"
OpName %_arg_order_write_11 "_arg_order_write"
OpName %_arg_order_read_11 "_arg_order_read"
OpName %_arg_res_acc_11 "_arg_res_acc"
OpName %_arg_res_acc9_11 "_arg_res_acc9"
OpName %entry_11 "entry"
OpName %for_cond_i_preheader_3 "for.cond.i.preheader"
OpName %if_then_i_11 "if.then.i"
OpName %sw_bb1_i_i_i_11 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_11 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_11 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_11 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_11 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_11 "for.cond.i"
OpName %for_body_i_11 "for.body.i"
OpName %for_inc_i_11 "for.inc.i"
OpName %cleanup_i_11 "cleanup.i"
OpName %sw_bb1_i_i28_i_3 "sw.bb1.i.i28.i"
OpName %sw_bb2_i_i29_i_3 "sw.bb2.i.i29.i"
OpName %sw_bb3_i_i30_i_3 "sw.bb3.i.i30.i"
OpName %sw_bb4_i_i31_i_3 "sw.bb4.i.i31.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_3 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit33.i"
OpName %if_then9_i_3 "if.then9.i"
OpName %if_then12_i_3 "if.then12.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_11 "add.ptr.i"
OpName %add_ptr_i39_3 "add.ptr.i39"
OpName %add_ptr_i53_3 "add.ptr.i53"
OpName %arrayidx_ascast_i_i_11 "arrayidx.ascast.i.i"
OpName %sub_i_i_i_i_i_11 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_11 "cmp.i.i.i"
OpName %cmp_i_i_3 "cmp.i.i"
OpName %SpvOrder_0_i_i_i_11 "SpvOrder.0.i.i.i"
OpName %inc_i_11 "inc.i"
OpName %i_0_i_11 "i.0.i"
OpName %cmp_i_11 "cmp.i"
OpName %call3_i_i_i_i_11 "call3.i.i.i.i"
OpName %cmp7_i_3 "cmp7.i"
OpName %SpvOrder_0_i_i32_i_3 "SpvOrder.0.i.i32.i"
OpName %cmp11_not_i_3 "cmp11.not.i"
OpName %_arg_sync_flag_acc_12 "_arg_sync_flag_acc"
OpName %_arg_sync_flag_acc3_4 "_arg_sync_flag_acc3"
OpName %_arg_data_acc_12 "_arg_data_acc"
OpName %_arg_data_acc6_4 "_arg_data_acc6"
OpName %_arg_value_12 "_arg_value"
OpName %_arg_order_write_12 "_arg_order_write"
OpName %_arg_order_read_12 "_arg_order_read"
OpName %_arg_res_acc_12 "_arg_res_acc"
OpName %_arg_res_acc9_12 "_arg_res_acc9"
OpName %entry_12 "entry"
OpName %for_cond_i_preheader_4 "for.cond.i.preheader"
OpName %if_then_i_12 "if.then.i"
OpName %sw_bb1_i_i_i_12 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_12 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_12 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_12 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_12 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_12 "for.cond.i"
OpName %for_body_i_12 "for.body.i"
OpName %for_inc_i_12 "for.inc.i"
OpName %cleanup_i_12 "cleanup.i"
OpName %sw_bb1_i_i28_i_4 "sw.bb1.i.i28.i"
OpName %sw_bb2_i_i29_i_4 "sw.bb2.i.i29.i"
OpName %sw_bb3_i_i30_i_4 "sw.bb3.i.i30.i"
OpName %sw_bb4_i_i31_i_4 "sw.bb4.i.i31.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_4 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit33.i"
OpName %if_then9_i_4 "if.then9.i"
OpName %if_then12_i_4 "if.then12.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_12 "add.ptr.i"
OpName %add_ptr_i39_4 "add.ptr.i39"
OpName %add_ptr_i53_4 "add.ptr.i53"
OpName %arrayidx_ascast_i_i_12 "arrayidx.ascast.i.i"
OpName %sub_i_i_i_i_i_12 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_12 "cmp.i.i.i"
OpName %cmp_i_i_4 "cmp.i.i"
OpName %SpvOrder_0_i_i_i_12 "SpvOrder.0.i.i.i"
OpName %inc_i_12 "inc.i"
OpName %i_0_i_12 "i.0.i"
OpName %cmp_i_12 "cmp.i"
OpName %call3_i_i_i_i_12 "call3.i.i.i.i"
OpName %cmp7_i_4 "cmp7.i"
OpName %SpvOrder_0_i_i32_i_4 "SpvOrder.0.i.i32.i"
OpName %cmp11_not_i_4 "cmp11.not.i"
OpName %_arg_sync_flag_acc_13 "_arg_sync_flag_acc"
OpName %_arg_data_acc_13 "_arg_data_acc"
OpName %_arg_value_13 "_arg_value"
OpName %_arg_order_write_13 "_arg_order_write"
OpName %_arg_order_read_13 "_arg_order_read"
OpName %_arg_res_acc_13 "_arg_res_acc"
OpName %_arg_res_acc9_13 "_arg_res_acc9"
OpName %entry_13 "entry"
OpName %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i "_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE.exit.i"
OpName %if_then_i_13 "if.then.i"
OpName %sw_bb1_i_i_i_13 "sw.bb1.i.i.i"
OpName %sw_bb2_i_i_i_13 "sw.bb2.i.i.i"
OpName %sw_bb3_i_i_i_13 "sw.bb3.i.i.i"
OpName %sw_bb4_i_i_i_13 "sw.bb4.i.i.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_13 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit.i"
OpName %for_cond_i_13 "for.cond.i"
OpName %for_body_i_13 "for.body.i"
OpName %for_inc_i_13 "for.inc.i"
OpName %cleanup_i_13 "cleanup.i"
OpName %sw_bb1_i_i22_i_7 "sw.bb1.i.i22.i"
OpName %sw_bb2_i_i23_i_7 "sw.bb2.i.i23.i"
OpName %sw_bb3_i_i24_i_7 "sw.bb3.i.i24.i"
OpName %sw_bb4_i_i25_i_7 "sw.bb4.i.i25.i"
OpName %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_7 "_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE.exit27.i"
OpName %if_then8_i_7 "if.then8.i"
OpName %if_then11_i_7 "if.then11.i"
OpName %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit "_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN_.exit"
OpName %add_ptr_i_13 "add.ptr.i"
OpName %arrayidx_ascast_i_i_13 "arrayidx.ascast.i.i"
OpName %cmp_i5_i_i_7 "cmp.i5.i.i"
OpName %sub_i_i_i_i_i_13 "sub.i.i.i.i.i"
OpName %cmp_i_i_i_13 "cmp.i.i.i"
OpName %SpvOrder_0_i_i_i_13 "SpvOrder.0.i.i.i"
OpName %inc_i_13 "inc.i"
OpName %i_0_i_13 "i.0.i"
OpName %cmp_i_13 "cmp.i"
OpName %call3_i_i_i_i_13 "call3.i.i.i.i"
OpName %cmp6_i_7 "cmp6.i"
OpName %SpvOrder_0_i_i26_i_7 "SpvOrder.0.i.i26.i"
OpName %cmp10_not_i_7 "cmp10.not.i"
OpName %entry_14 "entry"
OpName %if_then "if.then"
OpName %if_end "if.end"
OpName %call_i "call.i"
OpName %cmp_i_not "cmp.i.not"
OpName %__itt_offload_atomic_op_start_stub "__itt_offload_atomic_op_start_stub"
OpName %object_1 "object"
OpName %op_type_1 "op_type"
OpName %mem_order_1 "mem_order"
OpName %entry_15 "entry"
OpName %if_then_0 "if.then"
OpName %if_end_0 "if.end"
OpName %call_i_0 "call.i"
OpName %cmp_i_not_0 "cmp.i.not"
OpName %__itt_offload_atomic_op_finish_stub "__itt_offload_atomic_op_finish_stub"
OpName %object_2 "object"
OpName %op_type_2 "op_type"
OpName %mem_order_2 "mem_order"
OpName %entry_16 "entry"
OpName %if_end_1 "if.end"
OpName %return "return"
OpName %GroupID "GroupID"
OpName %call_i_1 "call.i"
OpName %cmp_i_not_1 "cmp.i.not"
OpName %arrayinit_begin5 "arrayinit.begin5"
OpName %arrayinit_begin "arrayinit.begin"
OpName %arrayinit_element6 "arrayinit.element6"
OpName %arrayinit_element17 "arrayinit.element17"
OpName %mul "mul"
OpName %mul2 "mul2"
OpName %conv "conv"
OpName %__itt_offload_wi_start_stub "__itt_offload_wi_start_stub"
OpName %group_id "group_id"
OpName %wi_id "wi_id"
OpName %wg_size "wg_size"
OpName %entry_17 "entry"
OpName %if_end_2 "if.end"
OpName %return_0 "return"
OpName %GroupID_0 "GroupID"
OpName %call_i_2 "call.i"
OpName %cmp_i_not_2 "cmp.i.not"
OpName %arrayinit_begin3 "arrayinit.begin3"
OpName %arrayinit_begin_0 "arrayinit.begin"
OpName %arrayinit_element4 "arrayinit.element4"
OpName %arrayinit_element15 "arrayinit.element15"
OpName %__itt_offload_wi_finish_stub "__itt_offload_wi_finish_stub"
OpName %group_id_0 "group_id"
OpName %wi_id_0 "wi_id"
OpName %entry_18 "entry"
OpName %group_id_addr "group_id.addr"
OpName %wi_id_addr "wi_id.addr"
OpName %wg_size_addr "wg_size.addr"
OpName %group_id_addr_ascast "group_id.addr.ascast"
OpName %wi_id_addr_ascast "wi_id.addr.ascast"
OpName %wg_size_addr_ascast "wg_size.addr.ascast"
OpName %entry_19 "entry"
OpName %group_id_addr_0 "group_id.addr"
OpName %wi_id_addr_0 "wi_id.addr"
OpName %group_id_addr_ascast_0 "group_id.addr.ascast"
OpName %wi_id_addr_ascast_0 "wi_id.addr.ascast"
OpName %entry_20 "entry"
OpName %object_addr "object.addr"
OpName %op_type_addr "op_type.addr"
OpName %mem_order_addr "mem_order.addr"
OpName %object_addr_ascast "object.addr.ascast"
OpName %op_type_addr_ascast "op_type.addr.ascast"
OpName %mem_order_addr_ascast "mem_order.addr.ascast"
OpName %entry_21 "entry"
OpName %object_addr_0 "object.addr"
OpName %op_type_addr_0 "op_type.addr"
OpName %mem_order_addr_0 "mem_order.addr"
OpName %object_addr_ascast_0 "object.addr.ascast"
OpName %op_type_addr_ascast_0 "op_type.addr.ascast"
OpName %mem_order_addr_ascast_0 "mem_order.addr.ascast"
OpModuleProcessed "Linked by SPIR-V Tools Linker"
OpDecorate %__spirv_BuiltInGlobalOffset Constant
OpDecorate %__spirv_BuiltInGlobalOffset BuiltIn GlobalOffset
OpDecorate %__spirv_BuiltInGlobalOffset Alignment 32
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Alignment 32
OpDecorate %__spirv_BuiltInLocalInvocationId Constant
OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId
OpDecorate %__spirv_BuiltInLocalInvocationId Alignment 32
OpDecorate %__spirv_BuiltInWorkgroupId Constant
OpDecorate %__spirv_BuiltInWorkgroupId BuiltIn WorkgroupId
OpDecorate %__spirv_BuiltInWorkgroupId Alignment 32
OpDecorate %__spirv_BuiltInGlobalLinearId Constant
OpDecorate %__spirv_BuiltInGlobalLinearId BuiltIn GlobalLinearId
OpDecorate %__spirv_BuiltInGlobalLinearId Alignment 8
OpDecorate %__spirv_BuiltInWorkgroupSize Constant
OpDecorate %__spirv_BuiltInWorkgroupSize BuiltIn WorkgroupSize
OpDecorate %__spirv_BuiltInWorkgroupSize Alignment 32
OpDecorate %_arg_sync_flag_acc Alignment 4
OpDecorate %_arg_data_acc Alignment 4
OpDecorate %_arg_res_acc Alignment 1
OpDecorate %_arg_res_acc9 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9 Alignment 8
OpDecorate %inc_i NoSignedWrap
OpDecorate %inc_i NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_0 Alignment 4
OpDecorate %_arg_data_acc_0 Alignment 4
OpDecorate %_arg_res_acc_0 Alignment 1
OpDecorate %_arg_res_acc9_0 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_0 Alignment 8
OpDecorate %inc_i_0 NoSignedWrap
OpDecorate %inc_i_0 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_1 Alignment 4
OpDecorate %_arg_data_acc_1 Alignment 4
OpDecorate %_arg_res_acc_1 Alignment 1
OpDecorate %_arg_res_acc9_1 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_1 Alignment 8
OpDecorate %inc_i_1 NoSignedWrap
OpDecorate %inc_i_1 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_2 Alignment 4
OpDecorate %_arg_data_acc_2 Alignment 4
OpDecorate %_arg_res_acc_2 Alignment 1
OpDecorate %_arg_res_acc9_2 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_2 Alignment 8
OpDecorate %inc_i_2 NoSignedWrap
OpDecorate %inc_i_2 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_3 Alignment 4
OpDecorate %_arg_sync_flag_acc3 FuncParamAttr ByVal
OpDecorate %_arg_sync_flag_acc3 Alignment 8
OpDecorate %_arg_data_acc_3 Alignment 4
OpDecorate %_arg_data_acc6 FuncParamAttr ByVal
OpDecorate %_arg_data_acc6 Alignment 8
OpDecorate %_arg_res_acc_3 Alignment 1
OpDecorate %_arg_res_acc9_3 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_3 Alignment 8
OpDecorate %inc_i_3 NoSignedWrap
OpDecorate %inc_i_3 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_4 Alignment 4
OpDecorate %_arg_data_acc_4 Alignment 4
OpDecorate %_arg_res_acc_4 Alignment 1
OpDecorate %_arg_res_acc9_4 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_4 Alignment 8
OpDecorate %inc_i_4 NoSignedWrap
OpDecorate %inc_i_4 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_5 Alignment 4
OpDecorate %_arg_data_acc_5 Alignment 4
OpDecorate %_arg_res_acc_5 Alignment 1
OpDecorate %_arg_res_acc9_5 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_5 Alignment 8
OpDecorate %inc_i_5 NoSignedWrap
OpDecorate %inc_i_5 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_6 Alignment 4
OpDecorate %_arg_data_acc_6 Alignment 4
OpDecorate %_arg_res_acc_6 Alignment 1
OpDecorate %_arg_res_acc9_6 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_6 Alignment 8
OpDecorate %inc_i_6 NoSignedWrap
OpDecorate %inc_i_6 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_7 Alignment 4
OpDecorate %_arg_sync_flag_acc3_0 FuncParamAttr ByVal
OpDecorate %_arg_sync_flag_acc3_0 Alignment 8
OpDecorate %_arg_data_acc_7 Alignment 4
OpDecorate %_arg_data_acc6_0 FuncParamAttr ByVal
OpDecorate %_arg_data_acc6_0 Alignment 8
OpDecorate %_arg_res_acc_7 Alignment 1
OpDecorate %_arg_res_acc9_7 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_7 Alignment 8
OpDecorate %inc_i_7 NoSignedWrap
OpDecorate %inc_i_7 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_8 Alignment 4
OpDecorate %_arg_sync_flag_acc3_1 FuncParamAttr ByVal
OpDecorate %_arg_sync_flag_acc3_1 Alignment 8
OpDecorate %_arg_data_acc_8 Alignment 4
OpDecorate %_arg_data_acc6_1 FuncParamAttr ByVal
OpDecorate %_arg_data_acc6_1 Alignment 8
OpDecorate %_arg_res_acc_8 Alignment 1
OpDecorate %_arg_res_acc9_8 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_8 Alignment 8
OpDecorate %inc_i_8 NoSignedWrap
OpDecorate %inc_i_8 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_9 Alignment 4
OpDecorate %_arg_data_acc_9 Alignment 4
OpDecorate %_arg_res_acc_9 Alignment 1
OpDecorate %_arg_res_acc9_9 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_9 Alignment 8
OpDecorate %inc_i_9 NoSignedWrap
OpDecorate %inc_i_9 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_10 Alignment 4
OpDecorate %_arg_sync_flag_acc3_2 FuncParamAttr ByVal
OpDecorate %_arg_sync_flag_acc3_2 Alignment 8
OpDecorate %_arg_data_acc_10 Alignment 4
OpDecorate %_arg_data_acc6_2 FuncParamAttr ByVal
OpDecorate %_arg_data_acc6_2 Alignment 8
OpDecorate %_arg_res_acc_10 Alignment 1
OpDecorate %_arg_res_acc9_10 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_10 Alignment 8
OpDecorate %inc_i_10 NoSignedWrap
OpDecorate %inc_i_10 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_11 Alignment 4
OpDecorate %_arg_sync_flag_acc3_3 FuncParamAttr ByVal
OpDecorate %_arg_sync_flag_acc3_3 Alignment 8
OpDecorate %_arg_data_acc_11 Alignment 4
OpDecorate %_arg_data_acc6_3 FuncParamAttr ByVal
OpDecorate %_arg_data_acc6_3 Alignment 8
OpDecorate %_arg_res_acc_11 Alignment 1
OpDecorate %_arg_res_acc9_11 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_11 Alignment 8
OpDecorate %inc_i_11 NoSignedWrap
OpDecorate %inc_i_11 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_12 Alignment 4
OpDecorate %_arg_sync_flag_acc3_4 FuncParamAttr ByVal
OpDecorate %_arg_sync_flag_acc3_4 Alignment 8
OpDecorate %_arg_data_acc_12 Alignment 4
OpDecorate %_arg_data_acc6_4 FuncParamAttr ByVal
OpDecorate %_arg_data_acc6_4 Alignment 8
OpDecorate %_arg_res_acc_12 Alignment 1
OpDecorate %_arg_res_acc9_12 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_12 Alignment 8
OpDecorate %inc_i_12 NoSignedWrap
OpDecorate %inc_i_12 NoUnsignedWrap
OpDecorate %_arg_sync_flag_acc_13 Alignment 4
OpDecorate %_arg_data_acc_13 Alignment 4
OpDecorate %_arg_res_acc_13 Alignment 1
OpDecorate %_arg_res_acc9_13 FuncParamAttr ByVal
OpDecorate %_arg_res_acc9_13 Alignment 8
OpDecorate %inc_i_13 NoSignedWrap
OpDecorate %inc_i_13 NoUnsignedWrap
OpDecorate %call_i SpecId 4285822057
OpDecorate %call_i_0 SpecId 4285822057
OpDecorate %GroupID Alignment 8
OpDecorate %call_i_1 SpecId 4285822057
OpDecorate %GroupID_0 Alignment 8
OpDecorate %call_i_2 SpecId 4285822057
OpDecorate %group_id_addr Alignment 8
OpDecorate %wi_id_addr Alignment 8
OpDecorate %wg_size_addr Alignment 4
OpDecorate %group_id_addr_0 Alignment 8
OpDecorate %wi_id_addr_0 Alignment 8
OpDecorate %object_addr Alignment 8
OpDecorate %op_type_addr Alignment 4
OpDecorate %mem_order_addr Alignment 4
OpDecorate %object_addr_0 Alignment 8
OpDecorate %op_type_addr_0 Alignment 4
OpDecorate %mem_order_addr_0 Alignment 4
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%uchar = OpTypeInt 8 0
%ulong_1 = OpConstant %ulong 1
%ulong_0 = OpConstant %ulong 0
%ulong_2147483648 = OpConstant %ulong 2147483648
%uint_912 = OpConstant %uint 912
%uint_904 = OpConstant %uint 904
%uint_900 = OpConstant %uint 900
%uint_898 = OpConstant %uint 898
%uint_896 = OpConstant %uint 896
%uint_2 = OpConstant %uint 2
%uint_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0
%ulong_256 = OpConstant %ulong 256
%uchar_0 = OpConstant %uchar 0
%call_i = OpSpecConstant %uchar 0
%call_i_0 = OpSpecConstant %uchar 0
%ulong_3 = OpConstant %ulong 3
%call_i_1 = OpSpecConstant %uchar 0
%ulong_2 = OpConstant %ulong 2
%call_i_2 = OpSpecConstant %uchar 0
%v3ulong = OpTypeVector %ulong 3
%_ptr_CrossWorkgroup_v3ulong = OpTypePointer CrossWorkgroup %v3ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%void = OpTypeVoid
%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
%_arr_ulong_ulong_1 = OpTypeArray %ulong %ulong_1
%class_sycl___V1__detail__array = OpTypeStruct %_arr_ulong_ulong_1
%class_sycl___V1__id = OpTypeStruct %class_sycl___V1__detail__array
%_ptr_Function_class_sycl___V1__id = OpTypePointer Function %class_sycl___V1__id
%759 = OpTypeFunction %void %_ptr_Workgroup_uint %_ptr_Workgroup_uint %uint %uint %uint %_ptr_CrossWorkgroup_uchar %_ptr_Function_class_sycl___V1__id
%760 = OpTypeFunction %void
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_uint = OpTypePointer Generic %uint
%bool = OpTypeBool
%_ptr_Generic_uchar = OpTypePointer Generic %uchar
%765 = OpTypeFunction %void %_ptr_Generic_uchar %uint %uint
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%767 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_Function_class_sycl___V1__id %_ptr_CrossWorkgroup_uint %_ptr_Function_class_sycl___V1__id %uint %uint %uint %_ptr_CrossWorkgroup_uchar %_ptr_Function_class_sycl___V1__id
%_arr_ulong_ulong_3 = OpTypeArray %ulong %ulong_3
%_ptr_Function__arr_ulong_ulong_3 = OpTypePointer Function %_arr_ulong_ulong_3
%_ptr_Function_uchar = OpTypePointer Function %uchar
%_ptr_Generic_ulong = OpTypePointer Generic %ulong
%772 = OpTypeFunction %void %_ptr_Generic_ulong %ulong %uint
%773 = OpTypeFunction %void %_ptr_Generic_ulong %ulong
%_ptr_Function__ptr_Generic_ulong = OpTypePointer Function %_ptr_Generic_ulong
%_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Generic__ptr_Generic_ulong = OpTypePointer Generic %_ptr_Generic_ulong
%_ptr_Function__ptr_Generic_uchar = OpTypePointer Function %_ptr_Generic_uchar
%_ptr_Generic__ptr_Generic_uchar = OpTypePointer Generic %_ptr_Generic_uchar
%__spirv_BuiltInGlobalOffset = OpVariable %_ptr_CrossWorkgroup_v3ulong CrossWorkgroup
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_CrossWorkgroup_v3ulong CrossWorkgroup
%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_CrossWorkgroup_v3ulong CrossWorkgroup
%__spirv_BuiltInWorkgroupId = OpVariable %_ptr_CrossWorkgroup_v3ulong CrossWorkgroup
%__spirv_BuiltInGlobalLinearId = OpVariable %_ptr_CrossWorkgroup_ulong CrossWorkgroup
%__spirv_BuiltInWorkgroupSize = OpVariable %_ptr_CrossWorkgroup_v3ulong CrossWorkgroup
%2 = OpFunction %void None %759
%_arg_sync_flag_acc = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_data_acc = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_value = OpFunctionParameter %uint
%_arg_order_write = OpFunctionParameter %uint
%_arg_order_read = OpFunctionParameter %uint
%_arg_res_acc = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry = OpLabel
%779 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%780 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9
%781 = OpLoad %ulong %780 Aligned 8
%add_ptr_i = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc %781
%arrayidx_ascast_i_i = OpPtrCastToGeneric %_ptr_Generic_uint %_arg_sync_flag_acc
%782 = OpLoad %v3ulong %__spirv_BuiltInLocalInvocationId Aligned 32
%783 = OpCompositeExtract %ulong %782 0
%cmp_i5_i_i = OpIEqual %bool %783 %ulong_0
OpBranchConditional %cmp_i5_i_i %if_then_i %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i = OpLabel
%784 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%785 = OpCompositeExtract %ulong %784 0
%786 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%787 = OpCompositeExtract %ulong %786 0
%sub_i_i_i_i_i = OpISub %ulong %787 %785
%cmp_i_i_i = OpULessThan %bool %sub_i_i_i_i_i %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i
%788 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i
OpBranch %for_cond_i
%if_then_i = OpLabel
OpStore %_arg_data_acc %_arg_value Aligned 4
OpSwitch %_arg_order_write %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i 5 %sw_bb4_i_i_i 2 %sw_bb1_i_i_i 1 %sw_bb1_i_i_i 3 %sw_bb2_i_i_i 4 %sw_bb3_i_i_i
%sw_bb1_i_i_i = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i
%sw_bb2_i_i_i = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i
%sw_bb3_i_i_i = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i
%sw_bb4_i_i_i = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i = OpLabel
%SpvOrder_0_i_i_i = OpPhi %uint %uint_912 %sw_bb4_i_i_i %uint_904 %sw_bb3_i_i_i %uint_900 %sw_bb2_i_i_i %uint_898 %sw_bb1_i_i_i %uint_896 %if_then_i
OpMemoryBarrier %uint_2 %SpvOrder_0_i_i_i
%789 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i
%790 = OpFunctionCall %void %__itt_offload_atomic_op_start %789 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i %uint_2 %uint_896 %uint_1
%791 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i
%792 = OpFunctionCall %void %__itt_offload_atomic_op_finish %791 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i = OpLabel
%i_0_i = OpPhi %ulong %inc_i %for_inc_i %ulong_0 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%cmp_i = OpULessThan %bool %i_0_i %ulong_256
OpBranchConditional %cmp_i %for_body_i %cleanup_i
%for_body_i = OpLabel
%793 = OpFunctionCall %void %__itt_offload_atomic_op_start %788 %uint_0 %uint_0
%call3_i_i_i_i = OpAtomicLoad %uint %arrayidx_ascast_i_i %uint_2 %uint_896
%794 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i
%795 = OpFunctionCall %void %__itt_offload_atomic_op_finish %794 %uint_0 %uint_0
%cmp6_i = OpIEqual %bool %call3_i_i_i_i %uint_1
OpBranchConditional %cmp6_i %cleanup_i %for_inc_i
%for_inc_i = OpLabel
%inc_i = OpIAdd %ulong %i_0_i %ulong_1
OpBranch %for_cond_i
%cleanup_i = OpLabel
OpSwitch %_arg_order_read %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i 5 %sw_bb4_i_i25_i 2 %sw_bb1_i_i22_i 1 %sw_bb1_i_i22_i 3 %sw_bb2_i_i23_i 4 %sw_bb3_i_i24_i
%sw_bb1_i_i22_i = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i
%sw_bb2_i_i23_i = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i
%sw_bb3_i_i24_i = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i
%sw_bb4_i_i25_i = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i = OpLabel
%SpvOrder_0_i_i26_i = OpPhi %uint %uint_912 %sw_bb4_i_i25_i %uint_904 %sw_bb3_i_i24_i %uint_900 %sw_bb2_i_i23_i %uint_898 %sw_bb1_i_i22_i %uint_896 %cleanup_i
OpMemoryBarrier %uint_2 %SpvOrder_0_i_i26_i
OpBranchConditional %cmp_i %if_then8_i %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then8_i = OpLabel
%796 = OpLoad %uint %_arg_data_acc Aligned 4
%cmp10_not_i = OpIEqual %bool %796 %_arg_value
OpBranchConditional %cmp10_not_i %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then11_i
%if_then11_i = OpLabel
OpStore %add_ptr_i %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%797 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%__itt_offload_wi_start_wrapper = OpFunction %void Inline %760
%entry_16 = OpLabel
%GroupID = OpVariable %_ptr_Function__arr_ulong_ulong_3 Function
%cmp_i_not_1 = OpIEqual %bool %call_i_1 %uchar_0
OpBranchConditional %cmp_i_not_1 %return %if_end_1
%if_end_1 = OpLabel
%798 = OpBitcast %_ptr_Function_uchar %GroupID
OpLifetimeStart %798 24
%arrayinit_begin5 = OpInBoundsPtrAccessChain %_ptr_Function_ulong %GroupID %ulong_0 %ulong_0
%arrayinit_begin = OpPtrCastToGeneric %_ptr_Generic_ulong %arrayinit_begin5
%799 = OpLoad %v3ulong %__spirv_BuiltInWorkgroupId Aligned 32
%800 = OpCompositeExtract %ulong %799 0
OpStore %arrayinit_begin5 %800 Aligned 8
%arrayinit_element6 = OpInBoundsPtrAccessChain %_ptr_Function_ulong %GroupID %ulong_0 %ulong_1
%801 = OpCompositeExtract %ulong %799 1
OpStore %arrayinit_element6 %801 Aligned 8
%arrayinit_element17 = OpInBoundsPtrAccessChain %_ptr_Function_ulong %GroupID %ulong_0 %ulong_2
%802 = OpCompositeExtract %ulong %799 2
OpStore %arrayinit_element17 %802 Aligned 8
%803 = OpLoad %ulong %__spirv_BuiltInGlobalLinearId Aligned 8
%804 = OpLoad %v3ulong %__spirv_BuiltInWorkgroupSize Aligned 32
%805 = OpCompositeExtract %ulong %804 0
%806 = OpLoad %v3ulong %__spirv_BuiltInWorkgroupSize Aligned 32
%807 = OpCompositeExtract %ulong %806 1
%mul = OpIMul %ulong %805 %807
%808 = OpLoad %v3ulong %__spirv_BuiltInWorkgroupSize Aligned 32
%809 = OpCompositeExtract %ulong %808 2
%mul2 = OpIMul %ulong %mul %809
%conv = OpUConvert %uint %mul2
%810 = OpFunctionCall %void %__itt_offload_wi_start_stub %arrayinit_begin %803 %conv
OpLifetimeStop %798 24
OpBranch %return
%return = OpLabel
OpReturn
OpFunctionEnd
%__itt_offload_atomic_op_start = OpFunction %void None %765
%object = OpFunctionParameter %_ptr_Generic_uchar
%op_type = OpFunctionParameter %uint
%mem_order = OpFunctionParameter %uint
%entry_14 = OpLabel
%cmp_i_not = OpIEqual %bool %call_i %uchar_0
OpBranchConditional %cmp_i_not %if_end %if_then
%if_then = OpLabel
%811 = OpFunctionCall %void %__itt_offload_atomic_op_start_stub %object %op_type %mem_order
OpBranch %if_end
%if_end = OpLabel
OpReturn
OpFunctionEnd
%__itt_offload_atomic_op_finish = OpFunction %void None %765
%object_0 = OpFunctionParameter %_ptr_Generic_uchar
%op_type_0 = OpFunctionParameter %uint
%mem_order_0 = OpFunctionParameter %uint
%entry_15 = OpLabel
%cmp_i_not_0 = OpIEqual %bool %call_i_0 %uchar_0
OpBranchConditional %cmp_i_not_0 %if_end_0 %if_then_0
%if_then_0 = OpLabel
%812 = OpFunctionCall %void %__itt_offload_atomic_op_finish_stub %object_0 %op_type_0 %mem_order_0
OpBranch %if_end_0
%if_end_0 = OpLabel
OpReturn
OpFunctionEnd
%__itt_offload_wi_finish_wrapper = OpFunction %void Inline %760
%entry_17 = OpLabel
%GroupID_0 = OpVariable %_ptr_Function__arr_ulong_ulong_3 Function
%cmp_i_not_2 = OpIEqual %bool %call_i_2 %uchar_0
OpBranchConditional %cmp_i_not_2 %return_0 %if_end_2
%if_end_2 = OpLabel
%813 = OpBitcast %_ptr_Function_uchar %GroupID_0
OpLifetimeStart %813 24
%arrayinit_begin3 = OpInBoundsPtrAccessChain %_ptr_Function_ulong %GroupID_0 %ulong_0 %ulong_0
%arrayinit_begin_0 = OpPtrCastToGeneric %_ptr_Generic_ulong %arrayinit_begin3
%814 = OpLoad %v3ulong %__spirv_BuiltInWorkgroupId Aligned 32
%815 = OpCompositeExtract %ulong %814 0
OpStore %arrayinit_begin3 %815 Aligned 8
%arrayinit_element4 = OpInBoundsPtrAccessChain %_ptr_Function_ulong %GroupID_0 %ulong_0 %ulong_1
%816 = OpCompositeExtract %ulong %814 1
OpStore %arrayinit_element4 %816 Aligned 8
%arrayinit_element15 = OpInBoundsPtrAccessChain %_ptr_Function_ulong %GroupID_0 %ulong_0 %ulong_2
%817 = OpCompositeExtract %ulong %814 2
OpStore %arrayinit_element15 %817 Aligned 8
%818 = OpLoad %ulong %__spirv_BuiltInGlobalLinearId Aligned 8
%819 = OpFunctionCall %void %__itt_offload_wi_finish_stub %arrayinit_begin_0 %818
OpLifetimeStop %813 24
OpBranch %return_0
%return_0 = OpLabel
OpReturn
OpFunctionEnd
%3 = OpFunction %void None %759
%_arg_sync_flag_acc_0 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_data_acc_0 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_value_0 = OpFunctionParameter %uint
%_arg_order_write_0 = OpFunctionParameter %uint
%_arg_order_read_0 = OpFunctionParameter %uint
%_arg_res_acc_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_0 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_0 = OpLabel
%820 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%821 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_0
%822 = OpLoad %ulong %821 Aligned 8
%add_ptr_i_0 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_0 %822
%arrayidx_ascast_i_i_0 = OpPtrCastToGeneric %_ptr_Generic_uint %_arg_sync_flag_acc_0
%823 = OpLoad %v3ulong %__spirv_BuiltInLocalInvocationId Aligned 32
%824 = OpCompositeExtract %ulong %823 0
%cmp_i5_i_i_0 = OpIEqual %bool %824 %ulong_0
OpBranchConditional %cmp_i5_i_i_0 %if_then_i_0 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i = OpLabel
%825 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%826 = OpCompositeExtract %ulong %825 0
%827 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%828 = OpCompositeExtract %ulong %827 0
%sub_i_i_i_i_i_0 = OpISub %ulong %828 %826
%cmp_i_i_i_0 = OpULessThan %bool %sub_i_i_i_i_i_0 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_0
%829 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_0
OpBranch %for_cond_i_0
%if_then_i_0 = OpLabel
OpStore %_arg_data_acc_0 %_arg_value_0 Aligned 4
OpSwitch %_arg_order_write_0 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_0 5 %sw_bb4_i_i_i_0 2 %sw_bb1_i_i_i_0 1 %sw_bb1_i_i_i_0 3 %sw_bb2_i_i_i_0 4 %sw_bb3_i_i_i_0
%sw_bb1_i_i_i_0 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_0
%sw_bb2_i_i_i_0 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_0
%sw_bb3_i_i_i_0 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_0
%sw_bb4_i_i_i_0 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_0
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_0 = OpLabel
%SpvOrder_0_i_i_i_0 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_0 %uint_904 %sw_bb3_i_i_i_0 %uint_900 %sw_bb2_i_i_i_0 %uint_898 %sw_bb1_i_i_i_0 %uint_896 %if_then_i_0
OpMemoryBarrier %uint_2 %SpvOrder_0_i_i_i_0
%830 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_0
%831 = OpFunctionCall %void %__itt_offload_atomic_op_start %830 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_0 %uint_2 %uint_896 %uint_1
%832 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_0
%833 = OpFunctionCall %void %__itt_offload_atomic_op_finish %832 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_0 = OpLabel
%i_0_i_0 = OpPhi %ulong %inc_i_0 %for_inc_i_0 %ulong_0 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%cmp_i_0 = OpULessThan %bool %i_0_i_0 %ulong_256
OpBranchConditional %cmp_i_0 %for_body_i_0 %cleanup_i_0
%for_body_i_0 = OpLabel
%834 = OpFunctionCall %void %__itt_offload_atomic_op_start %829 %uint_0 %uint_0
%call3_i_i_i_i_0 = OpAtomicLoad %uint %arrayidx_ascast_i_i_0 %uint_2 %uint_896
%835 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_0
%836 = OpFunctionCall %void %__itt_offload_atomic_op_finish %835 %uint_0 %uint_0
%cmp6_i_0 = OpIEqual %bool %call3_i_i_i_i_0 %uint_1
OpBranchConditional %cmp6_i_0 %cleanup_i_0 %for_inc_i_0
%for_inc_i_0 = OpLabel
%inc_i_0 = OpIAdd %ulong %i_0_i_0 %ulong_1
OpBranch %for_cond_i_0
%cleanup_i_0 = OpLabel
OpSwitch %_arg_order_read_0 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_0 5 %sw_bb4_i_i25_i_0 2 %sw_bb1_i_i22_i_0 1 %sw_bb1_i_i22_i_0 3 %sw_bb2_i_i23_i_0 4 %sw_bb3_i_i24_i_0
%sw_bb1_i_i22_i_0 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_0
%sw_bb2_i_i23_i_0 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_0
%sw_bb3_i_i24_i_0 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_0
%sw_bb4_i_i25_i_0 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_0
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_0 = OpLabel
%SpvOrder_0_i_i26_i_0 = OpPhi %uint %uint_912 %sw_bb4_i_i25_i_0 %uint_904 %sw_bb3_i_i24_i_0 %uint_900 %sw_bb2_i_i23_i_0 %uint_898 %sw_bb1_i_i22_i_0 %uint_896 %cleanup_i_0
OpMemoryBarrier %uint_2 %SpvOrder_0_i_i26_i_0
OpBranchConditional %cmp_i_0 %if_then8_i_0 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then8_i_0 = OpLabel
%837 = OpLoad %uint %_arg_data_acc_0 Aligned 4
%cmp10_not_i_0 = OpIEqual %bool %837 %_arg_value_0
OpBranchConditional %cmp10_not_i_0 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then11_i_0
%if_then11_i_0 = OpLabel
OpStore %add_ptr_i_0 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%838 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%4 = OpFunction %void None %759
%_arg_sync_flag_acc_1 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_data_acc_1 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_value_1 = OpFunctionParameter %uint
%_arg_order_write_1 = OpFunctionParameter %uint
%_arg_order_read_1 = OpFunctionParameter %uint
%_arg_res_acc_1 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_1 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_1 = OpLabel
%839 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%840 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_1
%841 = OpLoad %ulong %840 Aligned 8
%add_ptr_i_1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_1 %841
%arrayidx_ascast_i_i_1 = OpPtrCastToGeneric %_ptr_Generic_uint %_arg_sync_flag_acc_1
%842 = OpLoad %v3ulong %__spirv_BuiltInLocalInvocationId Aligned 32
%843 = OpCompositeExtract %ulong %842 0
%cmp_i5_i_i_1 = OpIEqual %bool %843 %ulong_0
OpBranchConditional %cmp_i5_i_i_1 %if_then_i_1 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i = OpLabel
%844 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%845 = OpCompositeExtract %ulong %844 0
%846 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%847 = OpCompositeExtract %ulong %846 0
%sub_i_i_i_i_i_1 = OpISub %ulong %847 %845
%cmp_i_i_i_1 = OpULessThan %bool %sub_i_i_i_i_i_1 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_1
%848 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_1
OpBranch %for_cond_i_1
%if_then_i_1 = OpLabel
OpStore %_arg_data_acc_1 %_arg_value_1 Aligned 4
OpSwitch %_arg_order_write_1 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_1 5 %sw_bb4_i_i_i_1 2 %sw_bb1_i_i_i_1 1 %sw_bb1_i_i_i_1 3 %sw_bb2_i_i_i_1 4 %sw_bb3_i_i_i_1
%sw_bb1_i_i_i_1 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_1
%sw_bb2_i_i_i_1 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_1
%sw_bb3_i_i_i_1 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_1
%sw_bb4_i_i_i_1 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_1
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_1 = OpLabel
%SpvOrder_0_i_i_i_1 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_1 %uint_904 %sw_bb3_i_i_i_1 %uint_900 %sw_bb2_i_i_i_1 %uint_898 %sw_bb1_i_i_i_1 %uint_896 %if_then_i_1
OpMemoryBarrier %uint_0 %SpvOrder_0_i_i_i_1
%849 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_1
%850 = OpFunctionCall %void %__itt_offload_atomic_op_start %849 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_1 %uint_2 %uint_896 %uint_1
%851 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_1
%852 = OpFunctionCall %void %__itt_offload_atomic_op_finish %851 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_1 = OpLabel
%i_0_i_1 = OpPhi %ulong %inc_i_1 %for_inc_i_1 %ulong_0 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%cmp_i_1 = OpULessThan %bool %i_0_i_1 %ulong_256
OpBranchConditional %cmp_i_1 %for_body_i_1 %cleanup_i_1
%for_body_i_1 = OpLabel
%853 = OpFunctionCall %void %__itt_offload_atomic_op_start %848 %uint_0 %uint_0
%call3_i_i_i_i_1 = OpAtomicLoad %uint %arrayidx_ascast_i_i_1 %uint_2 %uint_896
%854 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_1
%855 = OpFunctionCall %void %__itt_offload_atomic_op_finish %854 %uint_0 %uint_0
%cmp6_i_1 = OpIEqual %bool %call3_i_i_i_i_1 %uint_1
OpBranchConditional %cmp6_i_1 %cleanup_i_1 %for_inc_i_1
%for_inc_i_1 = OpLabel
%inc_i_1 = OpIAdd %ulong %i_0_i_1 %ulong_1
OpBranch %for_cond_i_1
%cleanup_i_1 = OpLabel
OpSwitch %_arg_order_read_1 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_1 5 %sw_bb4_i_i25_i_1 2 %sw_bb1_i_i22_i_1 1 %sw_bb1_i_i22_i_1 3 %sw_bb2_i_i23_i_1 4 %sw_bb3_i_i24_i_1
%sw_bb1_i_i22_i_1 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_1
%sw_bb2_i_i23_i_1 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_1
%sw_bb3_i_i24_i_1 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_1
%sw_bb4_i_i25_i_1 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_1
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_1 = OpLabel
%SpvOrder_0_i_i26_i_1 = OpPhi %uint %uint_912 %sw_bb4_i_i25_i_1 %uint_904 %sw_bb3_i_i24_i_1 %uint_900 %sw_bb2_i_i23_i_1 %uint_898 %sw_bb1_i_i22_i_1 %uint_896 %cleanup_i_1
OpMemoryBarrier %uint_0 %SpvOrder_0_i_i26_i_1
OpBranchConditional %cmp_i_1 %if_then8_i_1 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then8_i_1 = OpLabel
%856 = OpLoad %uint %_arg_data_acc_1 Aligned 4
%cmp10_not_i_1 = OpIEqual %bool %856 %_arg_value_1
OpBranchConditional %cmp10_not_i_1 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then11_i_1
%if_then11_i_1 = OpLabel
OpStore %add_ptr_i_1 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%857 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%5 = OpFunction %void None %759
%_arg_sync_flag_acc_2 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_data_acc_2 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_value_2 = OpFunctionParameter %uint
%_arg_order_write_2 = OpFunctionParameter %uint
%_arg_order_read_2 = OpFunctionParameter %uint
%_arg_res_acc_2 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_2 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_2 = OpLabel
%858 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%859 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_2
%860 = OpLoad %ulong %859 Aligned 8
%add_ptr_i_2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_2 %860
%arrayidx_ascast_i_i_2 = OpPtrCastToGeneric %_ptr_Generic_uint %_arg_sync_flag_acc_2
%861 = OpLoad %v3ulong %__spirv_BuiltInLocalInvocationId Aligned 32
%862 = OpCompositeExtract %ulong %861 0
%cmp_i5_i_i_2 = OpIEqual %bool %862 %ulong_0
OpBranchConditional %cmp_i5_i_i_2 %if_then_i_2 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i = OpLabel
%863 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%864 = OpCompositeExtract %ulong %863 0
%865 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%866 = OpCompositeExtract %ulong %865 0
%sub_i_i_i_i_i_2 = OpISub %ulong %866 %864
%cmp_i_i_i_2 = OpULessThan %bool %sub_i_i_i_i_i_2 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_2
%867 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_2
OpBranch %for_cond_i_2
%if_then_i_2 = OpLabel
OpStore %_arg_data_acc_2 %_arg_value_2 Aligned 4
OpSwitch %_arg_order_write_2 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_2 5 %sw_bb4_i_i_i_2 2 %sw_bb1_i_i_i_2 1 %sw_bb1_i_i_i_2 3 %sw_bb2_i_i_i_2 4 %sw_bb3_i_i_i_2
%sw_bb1_i_i_i_2 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_2
%sw_bb2_i_i_i_2 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_2
%sw_bb3_i_i_i_2 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_2
%sw_bb4_i_i_i_2 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_2
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_2 = OpLabel
%SpvOrder_0_i_i_i_2 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_2 %uint_904 %sw_bb3_i_i_i_2 %uint_900 %sw_bb2_i_i_i_2 %uint_898 %sw_bb1_i_i_i_2 %uint_896 %if_then_i_2
OpMemoryBarrier %uint_1 %SpvOrder_0_i_i_i_2
%868 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_2
%869 = OpFunctionCall %void %__itt_offload_atomic_op_start %868 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_2 %uint_2 %uint_896 %uint_1
%870 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_2
%871 = OpFunctionCall %void %__itt_offload_atomic_op_finish %870 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_2 = OpLabel
%i_0_i_2 = OpPhi %ulong %inc_i_2 %for_inc_i_2 %ulong_0 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%cmp_i_2 = OpULessThan %bool %i_0_i_2 %ulong_256
OpBranchConditional %cmp_i_2 %for_body_i_2 %cleanup_i_2
%for_body_i_2 = OpLabel
%872 = OpFunctionCall %void %__itt_offload_atomic_op_start %867 %uint_0 %uint_0
%call3_i_i_i_i_2 = OpAtomicLoad %uint %arrayidx_ascast_i_i_2 %uint_2 %uint_896
%873 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_2
%874 = OpFunctionCall %void %__itt_offload_atomic_op_finish %873 %uint_0 %uint_0
%cmp6_i_2 = OpIEqual %bool %call3_i_i_i_i_2 %uint_1
OpBranchConditional %cmp6_i_2 %cleanup_i_2 %for_inc_i_2
%for_inc_i_2 = OpLabel
%inc_i_2 = OpIAdd %ulong %i_0_i_2 %ulong_1
OpBranch %for_cond_i_2
%cleanup_i_2 = OpLabel
OpSwitch %_arg_order_read_2 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_2 5 %sw_bb4_i_i25_i_2 2 %sw_bb1_i_i22_i_2 1 %sw_bb1_i_i22_i_2 3 %sw_bb2_i_i23_i_2 4 %sw_bb3_i_i24_i_2
%sw_bb1_i_i22_i_2 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_2
%sw_bb2_i_i23_i_2 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_2
%sw_bb3_i_i24_i_2 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_2
%sw_bb4_i_i25_i_2 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_2
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_2 = OpLabel
%SpvOrder_0_i_i26_i_2 = OpPhi %uint %uint_912 %sw_bb4_i_i25_i_2 %uint_904 %sw_bb3_i_i24_i_2 %uint_900 %sw_bb2_i_i23_i_2 %uint_898 %sw_bb1_i_i22_i_2 %uint_896 %cleanup_i_2
OpMemoryBarrier %uint_1 %SpvOrder_0_i_i26_i_2
OpBranchConditional %cmp_i_2 %if_then8_i_2 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then8_i_2 = OpLabel
%875 = OpLoad %uint %_arg_data_acc_2 Aligned 4
%cmp10_not_i_2 = OpIEqual %bool %875 %_arg_value_2
OpBranchConditional %cmp10_not_i_2 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then11_i_2
%if_then11_i_2 = OpLabel
OpStore %add_ptr_i_2 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%876 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%6 = OpFunction %void None %767
%_arg_sync_flag_acc_3 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%_arg_sync_flag_acc3 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%_arg_data_acc_3 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%_arg_data_acc6 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%_arg_value_3 = OpFunctionParameter %uint
%_arg_order_write_3 = OpFunctionParameter %uint
%_arg_order_read_3 = OpFunctionParameter %uint
%_arg_res_acc_3 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_3 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_3 = OpLabel
%877 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%878 = OpBitcast %_ptr_Function_ulong %_arg_sync_flag_acc3
%879 = OpLoad %ulong %878 Aligned 8
%add_ptr_i_3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %_arg_sync_flag_acc_3 %879
%880 = OpBitcast %_ptr_Function_ulong %_arg_data_acc6
%881 = OpLoad %ulong %880 Aligned 8
%add_ptr_i39 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %_arg_data_acc_3 %881
%882 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_3
%883 = OpLoad %ulong %882 Aligned 8
%add_ptr_i53 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_3 %883
%884 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%885 = OpCompositeExtract %ulong %884 0
%886 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%887 = OpCompositeExtract %ulong %886 0
%arrayidx_ascast_i_i_3 = OpPtrCastToGeneric %_ptr_Generic_uint %add_ptr_i_3
%sub_i_i_i_i_i_3 = OpISub %ulong %885 %887
%cmp_i_i_i_3 = OpULessThan %bool %sub_i_i_i_i_i_3 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_3
%cmp_i_i = OpIEqual %bool %885 %887
OpBranchConditional %cmp_i_i %if_then_i_3 %for_cond_i_preheader
%for_cond_i_preheader = OpLabel
%888 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_3
OpBranch %for_cond_i_3
%if_then_i_3 = OpLabel
OpStore %add_ptr_i39 %_arg_value_3 Aligned 4
OpSwitch %_arg_order_write_3 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_3 5 %sw_bb4_i_i_i_3 2 %sw_bb1_i_i_i_3 1 %sw_bb1_i_i_i_3 3 %sw_bb2_i_i_i_3 4 %sw_bb3_i_i_i_3
%sw_bb1_i_i_i_3 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_3
%sw_bb2_i_i_i_3 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_3
%sw_bb3_i_i_i_3 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_3
%sw_bb4_i_i_i_3 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_3
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_3 = OpLabel
%SpvOrder_0_i_i_i_3 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_3 %uint_904 %sw_bb3_i_i_i_3 %uint_900 %sw_bb2_i_i_i_3 %uint_898 %sw_bb1_i_i_i_3 %uint_896 %if_then_i_3
OpMemoryBarrier %uint_0 %SpvOrder_0_i_i_i_3
%889 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_3
%890 = OpFunctionCall %void %__itt_offload_atomic_op_start %889 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_3 %uint_2 %uint_896 %uint_1
%891 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_3
%892 = OpFunctionCall %void %__itt_offload_atomic_op_finish %891 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_3 = OpLabel
%i_0_i_3 = OpPhi %ulong %inc_i_3 %for_inc_i_3 %ulong_0 %for_cond_i_preheader
%cmp_i_3 = OpULessThan %bool %i_0_i_3 %ulong_256
OpBranchConditional %cmp_i_3 %for_body_i_3 %cleanup_i_3
%for_body_i_3 = OpLabel
%893 = OpFunctionCall %void %__itt_offload_atomic_op_start %888 %uint_0 %uint_0
%call3_i_i_i_i_3 = OpAtomicLoad %uint %arrayidx_ascast_i_i_3 %uint_2 %uint_896
%894 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_3
%895 = OpFunctionCall %void %__itt_offload_atomic_op_finish %894 %uint_0 %uint_0
%cmp7_i = OpIEqual %bool %call3_i_i_i_i_3 %uint_1
OpBranchConditional %cmp7_i %cleanup_i_3 %for_inc_i_3
%for_inc_i_3 = OpLabel
%inc_i_3 = OpIAdd %ulong %i_0_i_3 %ulong_1
OpBranch %for_cond_i_3
%cleanup_i_3 = OpLabel
OpSwitch %_arg_order_read_3 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i 5 %sw_bb4_i_i31_i 2 %sw_bb1_i_i28_i 1 %sw_bb1_i_i28_i 3 %sw_bb2_i_i29_i 4 %sw_bb3_i_i30_i
%sw_bb1_i_i28_i = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i
%sw_bb2_i_i29_i = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i
%sw_bb3_i_i30_i = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i
%sw_bb4_i_i31_i = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i = OpLabel
%SpvOrder_0_i_i32_i = OpPhi %uint %uint_912 %sw_bb4_i_i31_i %uint_904 %sw_bb3_i_i30_i %uint_900 %sw_bb2_i_i29_i %uint_898 %sw_bb1_i_i28_i %uint_896 %cleanup_i_3
OpMemoryBarrier %uint_0 %SpvOrder_0_i_i32_i
OpBranchConditional %cmp_i_3 %if_then9_i %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then9_i = OpLabel
%896 = OpLoad %uint %add_ptr_i39 Aligned 4
%cmp11_not_i = OpIEqual %bool %896 %_arg_value_3
OpBranchConditional %cmp11_not_i %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then12_i
%if_then12_i = OpLabel
OpStore %add_ptr_i53 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%897 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%7 = OpFunction %void None %759
%_arg_sync_flag_acc_4 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_data_acc_4 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_value_4 = OpFunctionParameter %uint
%_arg_order_write_4 = OpFunctionParameter %uint
%_arg_order_read_4 = OpFunctionParameter %uint
%_arg_res_acc_4 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_4 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_4 = OpLabel
%898 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%899 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_4
%900 = OpLoad %ulong %899 Aligned 8
%add_ptr_i_4 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_4 %900
%arrayidx_ascast_i_i_4 = OpPtrCastToGeneric %_ptr_Generic_uint %_arg_sync_flag_acc_4
%901 = OpLoad %v3ulong %__spirv_BuiltInLocalInvocationId Aligned 32
%902 = OpCompositeExtract %ulong %901 0
%cmp_i5_i_i_3 = OpIEqual %bool %902 %ulong_0
OpBranchConditional %cmp_i5_i_i_3 %if_then_i_4 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i = OpLabel
%903 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%904 = OpCompositeExtract %ulong %903 0
%905 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%906 = OpCompositeExtract %ulong %905 0
%sub_i_i_i_i_i_4 = OpISub %ulong %906 %904
%cmp_i_i_i_4 = OpULessThan %bool %sub_i_i_i_i_i_4 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_4
%907 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_4
OpBranch %for_cond_i_4
%if_then_i_4 = OpLabel
OpStore %_arg_data_acc_4 %_arg_value_4 Aligned 4
OpSwitch %_arg_order_write_4 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_4 5 %sw_bb4_i_i_i_4 2 %sw_bb1_i_i_i_4 1 %sw_bb1_i_i_i_4 3 %sw_bb2_i_i_i_4 4 %sw_bb3_i_i_i_4
%sw_bb1_i_i_i_4 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_4
%sw_bb2_i_i_i_4 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_4
%sw_bb3_i_i_i_4 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_4
%sw_bb4_i_i_i_4 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_4
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_4 = OpLabel
%SpvOrder_0_i_i_i_4 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_4 %uint_904 %sw_bb3_i_i_i_4 %uint_900 %sw_bb2_i_i_i_4 %uint_898 %sw_bb1_i_i_i_4 %uint_896 %if_then_i_4
OpMemoryBarrier %uint_0 %SpvOrder_0_i_i_i_4
%908 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_4
%909 = OpFunctionCall %void %__itt_offload_atomic_op_start %908 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_4 %uint_2 %uint_896 %uint_1
%910 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_4
%911 = OpFunctionCall %void %__itt_offload_atomic_op_finish %910 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_4 = OpLabel
%i_0_i_4 = OpPhi %ulong %inc_i_4 %for_inc_i_4 %ulong_0 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%cmp_i_4 = OpULessThan %bool %i_0_i_4 %ulong_256
OpBranchConditional %cmp_i_4 %for_body_i_4 %cleanup_i_4
%for_body_i_4 = OpLabel
%912 = OpFunctionCall %void %__itt_offload_atomic_op_start %907 %uint_0 %uint_0
%call3_i_i_i_i_4 = OpAtomicLoad %uint %arrayidx_ascast_i_i_4 %uint_2 %uint_896
%913 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_4
%914 = OpFunctionCall %void %__itt_offload_atomic_op_finish %913 %uint_0 %uint_0
%cmp6_i_3 = OpIEqual %bool %call3_i_i_i_i_4 %uint_1
OpBranchConditional %cmp6_i_3 %cleanup_i_4 %for_inc_i_4
%for_inc_i_4 = OpLabel
%inc_i_4 = OpIAdd %ulong %i_0_i_4 %ulong_1
OpBranch %for_cond_i_4
%cleanup_i_4 = OpLabel
OpSwitch %_arg_order_read_4 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_3 5 %sw_bb4_i_i25_i_3 2 %sw_bb1_i_i22_i_3 1 %sw_bb1_i_i22_i_3 3 %sw_bb2_i_i23_i_3 4 %sw_bb3_i_i24_i_3
%sw_bb1_i_i22_i_3 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_3
%sw_bb2_i_i23_i_3 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_3
%sw_bb3_i_i24_i_3 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_3
%sw_bb4_i_i25_i_3 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_3
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_3 = OpLabel
%SpvOrder_0_i_i26_i_3 = OpPhi %uint %uint_912 %sw_bb4_i_i25_i_3 %uint_904 %sw_bb3_i_i24_i_3 %uint_900 %sw_bb2_i_i23_i_3 %uint_898 %sw_bb1_i_i22_i_3 %uint_896 %cleanup_i_4
OpMemoryBarrier %uint_0 %SpvOrder_0_i_i26_i_3
OpBranchConditional %cmp_i_4 %if_then8_i_3 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then8_i_3 = OpLabel
%915 = OpLoad %uint %_arg_data_acc_4 Aligned 4
%cmp10_not_i_3 = OpIEqual %bool %915 %_arg_value_4
OpBranchConditional %cmp10_not_i_3 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then11_i_3
%if_then11_i_3 = OpLabel
OpStore %add_ptr_i_4 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%916 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%8 = OpFunction %void None %759
%_arg_sync_flag_acc_5 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_data_acc_5 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_value_5 = OpFunctionParameter %uint
%_arg_order_write_5 = OpFunctionParameter %uint
%_arg_order_read_5 = OpFunctionParameter %uint
%_arg_res_acc_5 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_5 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_5 = OpLabel
%917 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%918 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_5
%919 = OpLoad %ulong %918 Aligned 8
%add_ptr_i_5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_5 %919
%arrayidx_ascast_i_i_5 = OpPtrCastToGeneric %_ptr_Generic_uint %_arg_sync_flag_acc_5
%920 = OpLoad %v3ulong %__spirv_BuiltInLocalInvocationId Aligned 32
%921 = OpCompositeExtract %ulong %920 0
%cmp_i5_i_i_4 = OpIEqual %bool %921 %ulong_0
OpBranchConditional %cmp_i5_i_i_4 %if_then_i_5 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i = OpLabel
%922 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%923 = OpCompositeExtract %ulong %922 0
%924 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%925 = OpCompositeExtract %ulong %924 0
%sub_i_i_i_i_i_5 = OpISub %ulong %925 %923
%cmp_i_i_i_5 = OpULessThan %bool %sub_i_i_i_i_i_5 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_5
%926 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_5
OpBranch %for_cond_i_5
%if_then_i_5 = OpLabel
OpStore %_arg_data_acc_5 %_arg_value_5 Aligned 4
OpSwitch %_arg_order_write_5 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_5 5 %sw_bb4_i_i_i_5 2 %sw_bb1_i_i_i_5 1 %sw_bb1_i_i_i_5 3 %sw_bb2_i_i_i_5 4 %sw_bb3_i_i_i_5
%sw_bb1_i_i_i_5 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_5
%sw_bb2_i_i_i_5 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_5
%sw_bb3_i_i_i_5 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_5
%sw_bb4_i_i_i_5 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_5
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_5 = OpLabel
%SpvOrder_0_i_i_i_5 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_5 %uint_904 %sw_bb3_i_i_i_5 %uint_900 %sw_bb2_i_i_i_5 %uint_898 %sw_bb1_i_i_i_5 %uint_896 %if_then_i_5
OpMemoryBarrier %uint_0 %SpvOrder_0_i_i_i_5
%927 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_5
%928 = OpFunctionCall %void %__itt_offload_atomic_op_start %927 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_5 %uint_2 %uint_896 %uint_1
%929 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_5
%930 = OpFunctionCall %void %__itt_offload_atomic_op_finish %929 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_5 = OpLabel
%i_0_i_5 = OpPhi %ulong %inc_i_5 %for_inc_i_5 %ulong_0 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%cmp_i_5 = OpULessThan %bool %i_0_i_5 %ulong_256
OpBranchConditional %cmp_i_5 %for_body_i_5 %cleanup_i_5
%for_body_i_5 = OpLabel
%931 = OpFunctionCall %void %__itt_offload_atomic_op_start %926 %uint_0 %uint_0
%call3_i_i_i_i_5 = OpAtomicLoad %uint %arrayidx_ascast_i_i_5 %uint_2 %uint_896
%932 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_5
%933 = OpFunctionCall %void %__itt_offload_atomic_op_finish %932 %uint_0 %uint_0
%cmp6_i_4 = OpIEqual %bool %call3_i_i_i_i_5 %uint_1
OpBranchConditional %cmp6_i_4 %cleanup_i_5 %for_inc_i_5
%for_inc_i_5 = OpLabel
%inc_i_5 = OpIAdd %ulong %i_0_i_5 %ulong_1
OpBranch %for_cond_i_5
%cleanup_i_5 = OpLabel
OpSwitch %_arg_order_read_5 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_4 5 %sw_bb4_i_i25_i_4 2 %sw_bb1_i_i22_i_4 1 %sw_bb1_i_i22_i_4 3 %sw_bb2_i_i23_i_4 4 %sw_bb3_i_i24_i_4
%sw_bb1_i_i22_i_4 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_4
%sw_bb2_i_i23_i_4 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_4
%sw_bb3_i_i24_i_4 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_4
%sw_bb4_i_i25_i_4 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_4
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_4 = OpLabel
%SpvOrder_0_i_i26_i_4 = OpPhi %uint %uint_912 %sw_bb4_i_i25_i_4 %uint_904 %sw_bb3_i_i24_i_4 %uint_900 %sw_bb2_i_i23_i_4 %uint_898 %sw_bb1_i_i22_i_4 %uint_896 %cleanup_i_5
OpMemoryBarrier %uint_0 %SpvOrder_0_i_i26_i_4
OpBranchConditional %cmp_i_5 %if_then8_i_4 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then8_i_4 = OpLabel
%934 = OpLoad %uint %_arg_data_acc_5 Aligned 4
%cmp10_not_i_4 = OpIEqual %bool %934 %_arg_value_5
OpBranchConditional %cmp10_not_i_4 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then11_i_4
%if_then11_i_4 = OpLabel
OpStore %add_ptr_i_5 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%935 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%9 = OpFunction %void None %759
%_arg_sync_flag_acc_6 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_data_acc_6 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_value_6 = OpFunctionParameter %uint
%_arg_order_write_6 = OpFunctionParameter %uint
%_arg_order_read_6 = OpFunctionParameter %uint
%_arg_res_acc_6 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_6 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_6 = OpLabel
%936 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%937 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_6
%938 = OpLoad %ulong %937 Aligned 8
%add_ptr_i_6 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_6 %938
%arrayidx_ascast_i_i_6 = OpPtrCastToGeneric %_ptr_Generic_uint %_arg_sync_flag_acc_6
%939 = OpLoad %v3ulong %__spirv_BuiltInLocalInvocationId Aligned 32
%940 = OpCompositeExtract %ulong %939 0
%cmp_i5_i_i_5 = OpIEqual %bool %940 %ulong_0
OpBranchConditional %cmp_i5_i_i_5 %if_then_i_6 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i = OpLabel
%941 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%942 = OpCompositeExtract %ulong %941 0
%943 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%944 = OpCompositeExtract %ulong %943 0
%sub_i_i_i_i_i_6 = OpISub %ulong %944 %942
%cmp_i_i_i_6 = OpULessThan %bool %sub_i_i_i_i_i_6 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_6
%945 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_6
OpBranch %for_cond_i_6
%if_then_i_6 = OpLabel
OpStore %_arg_data_acc_6 %_arg_value_6 Aligned 4
OpSwitch %_arg_order_write_6 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_6 5 %sw_bb4_i_i_i_6 2 %sw_bb1_i_i_i_6 1 %sw_bb1_i_i_i_6 3 %sw_bb2_i_i_i_6 4 %sw_bb3_i_i_i_6
%sw_bb1_i_i_i_6 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_6
%sw_bb2_i_i_i_6 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_6
%sw_bb3_i_i_i_6 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_6
%sw_bb4_i_i_i_6 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_6
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_6 = OpLabel
%SpvOrder_0_i_i_i_6 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_6 %uint_904 %sw_bb3_i_i_i_6 %uint_900 %sw_bb2_i_i_i_6 %uint_898 %sw_bb1_i_i_i_6 %uint_896 %if_then_i_6
OpMemoryBarrier %uint_1 %SpvOrder_0_i_i_i_6
%946 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_6
%947 = OpFunctionCall %void %__itt_offload_atomic_op_start %946 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_6 %uint_2 %uint_896 %uint_1
%948 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_6
%949 = OpFunctionCall %void %__itt_offload_atomic_op_finish %948 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_6 = OpLabel
%i_0_i_6 = OpPhi %ulong %inc_i_6 %for_inc_i_6 %ulong_0 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%cmp_i_6 = OpULessThan %bool %i_0_i_6 %ulong_256
OpBranchConditional %cmp_i_6 %for_body_i_6 %cleanup_i_6
%for_body_i_6 = OpLabel
%950 = OpFunctionCall %void %__itt_offload_atomic_op_start %945 %uint_0 %uint_0
%call3_i_i_i_i_6 = OpAtomicLoad %uint %arrayidx_ascast_i_i_6 %uint_2 %uint_896
%951 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_6
%952 = OpFunctionCall %void %__itt_offload_atomic_op_finish %951 %uint_0 %uint_0
%cmp6_i_5 = OpIEqual %bool %call3_i_i_i_i_6 %uint_1
OpBranchConditional %cmp6_i_5 %cleanup_i_6 %for_inc_i_6
%for_inc_i_6 = OpLabel
%inc_i_6 = OpIAdd %ulong %i_0_i_6 %ulong_1
OpBranch %for_cond_i_6
%cleanup_i_6 = OpLabel
OpSwitch %_arg_order_read_6 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_5 5 %sw_bb4_i_i25_i_5 2 %sw_bb1_i_i22_i_5 1 %sw_bb1_i_i22_i_5 3 %sw_bb2_i_i23_i_5 4 %sw_bb3_i_i24_i_5
%sw_bb1_i_i22_i_5 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_5
%sw_bb2_i_i23_i_5 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_5
%sw_bb3_i_i24_i_5 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_5
%sw_bb4_i_i25_i_5 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_5
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_5 = OpLabel
%SpvOrder_0_i_i26_i_5 = OpPhi %uint %uint_912 %sw_bb4_i_i25_i_5 %uint_904 %sw_bb3_i_i24_i_5 %uint_900 %sw_bb2_i_i23_i_5 %uint_898 %sw_bb1_i_i22_i_5 %uint_896 %cleanup_i_6
OpMemoryBarrier %uint_1 %SpvOrder_0_i_i26_i_5
OpBranchConditional %cmp_i_6 %if_then8_i_5 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then8_i_5 = OpLabel
%953 = OpLoad %uint %_arg_data_acc_6 Aligned 4
%cmp10_not_i_5 = OpIEqual %bool %953 %_arg_value_6
OpBranchConditional %cmp10_not_i_5 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then11_i_5
%if_then11_i_5 = OpLabel
OpStore %add_ptr_i_6 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%954 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%10 = OpFunction %void None %767
%_arg_sync_flag_acc_7 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%_arg_sync_flag_acc3_0 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%_arg_data_acc_7 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%_arg_data_acc6_0 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%_arg_value_7 = OpFunctionParameter %uint
%_arg_order_write_7 = OpFunctionParameter %uint
%_arg_order_read_7 = OpFunctionParameter %uint
%_arg_res_acc_7 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_7 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_7 = OpLabel
%955 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%956 = OpBitcast %_ptr_Function_ulong %_arg_sync_flag_acc3_0
%957 = OpLoad %ulong %956 Aligned 8
%add_ptr_i_7 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %_arg_sync_flag_acc_7 %957
%958 = OpBitcast %_ptr_Function_ulong %_arg_data_acc6_0
%959 = OpLoad %ulong %958 Aligned 8
%add_ptr_i39_0 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %_arg_data_acc_7 %959
%960 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_7
%961 = OpLoad %ulong %960 Aligned 8
%add_ptr_i53_0 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_7 %961
%962 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%963 = OpCompositeExtract %ulong %962 0
%964 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%965 = OpCompositeExtract %ulong %964 0
%arrayidx_ascast_i_i_7 = OpPtrCastToGeneric %_ptr_Generic_uint %add_ptr_i_7
%sub_i_i_i_i_i_7 = OpISub %ulong %963 %965
%cmp_i_i_i_7 = OpULessThan %bool %sub_i_i_i_i_i_7 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_7
%cmp_i_i_0 = OpIEqual %bool %963 %965
OpBranchConditional %cmp_i_i_0 %if_then_i_7 %for_cond_i_preheader_0
%for_cond_i_preheader_0 = OpLabel
%966 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_7
OpBranch %for_cond_i_7
%if_then_i_7 = OpLabel
OpStore %add_ptr_i39_0 %_arg_value_7 Aligned 4
OpSwitch %_arg_order_write_7 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_7 5 %sw_bb4_i_i_i_7 2 %sw_bb1_i_i_i_7 1 %sw_bb1_i_i_i_7 3 %sw_bb2_i_i_i_7 4 %sw_bb3_i_i_i_7
%sw_bb1_i_i_i_7 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_7
%sw_bb2_i_i_i_7 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_7
%sw_bb3_i_i_i_7 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_7
%sw_bb4_i_i_i_7 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_7
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_7 = OpLabel
%SpvOrder_0_i_i_i_7 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_7 %uint_904 %sw_bb3_i_i_i_7 %uint_900 %sw_bb2_i_i_i_7 %uint_898 %sw_bb1_i_i_i_7 %uint_896 %if_then_i_7
OpMemoryBarrier %uint_0 %SpvOrder_0_i_i_i_7
%967 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_7
%968 = OpFunctionCall %void %__itt_offload_atomic_op_start %967 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_7 %uint_2 %uint_896 %uint_1
%969 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_7
%970 = OpFunctionCall %void %__itt_offload_atomic_op_finish %969 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_7 = OpLabel
%i_0_i_7 = OpPhi %ulong %inc_i_7 %for_inc_i_7 %ulong_0 %for_cond_i_preheader_0
%cmp_i_7 = OpULessThan %bool %i_0_i_7 %ulong_256
OpBranchConditional %cmp_i_7 %for_body_i_7 %cleanup_i_7
%for_body_i_7 = OpLabel
%971 = OpFunctionCall %void %__itt_offload_atomic_op_start %966 %uint_0 %uint_0
%call3_i_i_i_i_7 = OpAtomicLoad %uint %arrayidx_ascast_i_i_7 %uint_2 %uint_896
%972 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_7
%973 = OpFunctionCall %void %__itt_offload_atomic_op_finish %972 %uint_0 %uint_0
%cmp7_i_0 = OpIEqual %bool %call3_i_i_i_i_7 %uint_1
OpBranchConditional %cmp7_i_0 %cleanup_i_7 %for_inc_i_7
%for_inc_i_7 = OpLabel
%inc_i_7 = OpIAdd %ulong %i_0_i_7 %ulong_1
OpBranch %for_cond_i_7
%cleanup_i_7 = OpLabel
OpSwitch %_arg_order_read_7 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_0 5 %sw_bb4_i_i31_i_0 2 %sw_bb1_i_i28_i_0 1 %sw_bb1_i_i28_i_0 3 %sw_bb2_i_i29_i_0 4 %sw_bb3_i_i30_i_0
%sw_bb1_i_i28_i_0 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_0
%sw_bb2_i_i29_i_0 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_0
%sw_bb3_i_i30_i_0 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_0
%sw_bb4_i_i31_i_0 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_0
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_0 = OpLabel
%SpvOrder_0_i_i32_i_0 = OpPhi %uint %uint_912 %sw_bb4_i_i31_i_0 %uint_904 %sw_bb3_i_i30_i_0 %uint_900 %sw_bb2_i_i29_i_0 %uint_898 %sw_bb1_i_i28_i_0 %uint_896 %cleanup_i_7
OpMemoryBarrier %uint_0 %SpvOrder_0_i_i32_i_0
OpBranchConditional %cmp_i_7 %if_then9_i_0 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then9_i_0 = OpLabel
%974 = OpLoad %uint %add_ptr_i39_0 Aligned 4
%cmp11_not_i_0 = OpIEqual %bool %974 %_arg_value_7
OpBranchConditional %cmp11_not_i_0 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then12_i_0
%if_then12_i_0 = OpLabel
OpStore %add_ptr_i53_0 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%975 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%11 = OpFunction %void None %767
%_arg_sync_flag_acc_8 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%_arg_sync_flag_acc3_1 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%_arg_data_acc_8 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%_arg_data_acc6_1 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%_arg_value_8 = OpFunctionParameter %uint
%_arg_order_write_8 = OpFunctionParameter %uint
%_arg_order_read_8 = OpFunctionParameter %uint
%_arg_res_acc_8 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_8 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_8 = OpLabel
%976 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%977 = OpBitcast %_ptr_Function_ulong %_arg_sync_flag_acc3_1
%978 = OpLoad %ulong %977 Aligned 8
%add_ptr_i_8 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %_arg_sync_flag_acc_8 %978
%979 = OpBitcast %_ptr_Function_ulong %_arg_data_acc6_1
%980 = OpLoad %ulong %979 Aligned 8
%add_ptr_i39_1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %_arg_data_acc_8 %980
%981 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_8
%982 = OpLoad %ulong %981 Aligned 8
%add_ptr_i53_1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_8 %982
%983 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%984 = OpCompositeExtract %ulong %983 0
%985 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%986 = OpCompositeExtract %ulong %985 0
%arrayidx_ascast_i_i_8 = OpPtrCastToGeneric %_ptr_Generic_uint %add_ptr_i_8
%sub_i_i_i_i_i_8 = OpISub %ulong %984 %986
%cmp_i_i_i_8 = OpULessThan %bool %sub_i_i_i_i_i_8 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_8
%cmp_i_i_1 = OpIEqual %bool %984 %986
OpBranchConditional %cmp_i_i_1 %if_then_i_8 %for_cond_i_preheader_1
%for_cond_i_preheader_1 = OpLabel
%987 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_8
OpBranch %for_cond_i_8
%if_then_i_8 = OpLabel
OpStore %add_ptr_i39_1 %_arg_value_8 Aligned 4
OpSwitch %_arg_order_write_8 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_8 5 %sw_bb4_i_i_i_8 2 %sw_bb1_i_i_i_8 1 %sw_bb1_i_i_i_8 3 %sw_bb2_i_i_i_8 4 %sw_bb3_i_i_i_8
%sw_bb1_i_i_i_8 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_8
%sw_bb2_i_i_i_8 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_8
%sw_bb3_i_i_i_8 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_8
%sw_bb4_i_i_i_8 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_8
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_8 = OpLabel
%SpvOrder_0_i_i_i_8 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_8 %uint_904 %sw_bb3_i_i_i_8 %uint_900 %sw_bb2_i_i_i_8 %uint_898 %sw_bb1_i_i_i_8 %uint_896 %if_then_i_8
OpMemoryBarrier %uint_1 %SpvOrder_0_i_i_i_8
%988 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_8
%989 = OpFunctionCall %void %__itt_offload_atomic_op_start %988 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_8 %uint_2 %uint_896 %uint_1
%990 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_8
%991 = OpFunctionCall %void %__itt_offload_atomic_op_finish %990 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_8 = OpLabel
%i_0_i_8 = OpPhi %ulong %inc_i_8 %for_inc_i_8 %ulong_0 %for_cond_i_preheader_1
%cmp_i_8 = OpULessThan %bool %i_0_i_8 %ulong_256
OpBranchConditional %cmp_i_8 %for_body_i_8 %cleanup_i_8
%for_body_i_8 = OpLabel
%992 = OpFunctionCall %void %__itt_offload_atomic_op_start %987 %uint_0 %uint_0
%call3_i_i_i_i_8 = OpAtomicLoad %uint %arrayidx_ascast_i_i_8 %uint_2 %uint_896
%993 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_8
%994 = OpFunctionCall %void %__itt_offload_atomic_op_finish %993 %uint_0 %uint_0
%cmp7_i_1 = OpIEqual %bool %call3_i_i_i_i_8 %uint_1
OpBranchConditional %cmp7_i_1 %cleanup_i_8 %for_inc_i_8
%for_inc_i_8 = OpLabel
%inc_i_8 = OpIAdd %ulong %i_0_i_8 %ulong_1
OpBranch %for_cond_i_8
%cleanup_i_8 = OpLabel
OpSwitch %_arg_order_read_8 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_1 5 %sw_bb4_i_i31_i_1 2 %sw_bb1_i_i28_i_1 1 %sw_bb1_i_i28_i_1 3 %sw_bb2_i_i29_i_1 4 %sw_bb3_i_i30_i_1
%sw_bb1_i_i28_i_1 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_1
%sw_bb2_i_i29_i_1 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_1
%sw_bb3_i_i30_i_1 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_1
%sw_bb4_i_i31_i_1 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_1
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_1 = OpLabel
%SpvOrder_0_i_i32_i_1 = OpPhi %uint %uint_912 %sw_bb4_i_i31_i_1 %uint_904 %sw_bb3_i_i30_i_1 %uint_900 %sw_bb2_i_i29_i_1 %uint_898 %sw_bb1_i_i28_i_1 %uint_896 %cleanup_i_8
OpMemoryBarrier %uint_1 %SpvOrder_0_i_i32_i_1
OpBranchConditional %cmp_i_8 %if_then9_i_1 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then9_i_1 = OpLabel
%995 = OpLoad %uint %add_ptr_i39_1 Aligned 4
%cmp11_not_i_1 = OpIEqual %bool %995 %_arg_value_8
OpBranchConditional %cmp11_not_i_1 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then12_i_1
%if_then12_i_1 = OpLabel
OpStore %add_ptr_i53_1 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%996 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%12 = OpFunction %void None %759
%_arg_sync_flag_acc_9 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_data_acc_9 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_value_9 = OpFunctionParameter %uint
%_arg_order_write_9 = OpFunctionParameter %uint
%_arg_order_read_9 = OpFunctionParameter %uint
%_arg_res_acc_9 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_9 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_9 = OpLabel
%997 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%998 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_9
%999 = OpLoad %ulong %998 Aligned 8
%add_ptr_i_9 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_9 %999
%arrayidx_ascast_i_i_9 = OpPtrCastToGeneric %_ptr_Generic_uint %_arg_sync_flag_acc_9
%1000 = OpLoad %v3ulong %__spirv_BuiltInLocalInvocationId Aligned 32
%1001 = OpCompositeExtract %ulong %1000 0
%cmp_i5_i_i_6 = OpIEqual %bool %1001 %ulong_0
OpBranchConditional %cmp_i5_i_i_6 %if_then_i_9 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i = OpLabel
%1002 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%1003 = OpCompositeExtract %ulong %1002 0
%1004 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%1005 = OpCompositeExtract %ulong %1004 0
%sub_i_i_i_i_i_9 = OpISub %ulong %1005 %1003
%cmp_i_i_i_9 = OpULessThan %bool %sub_i_i_i_i_i_9 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_9
%1006 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_9
OpBranch %for_cond_i_9
%if_then_i_9 = OpLabel
OpStore %_arg_data_acc_9 %_arg_value_9 Aligned 4
OpSwitch %_arg_order_write_9 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_9 5 %sw_bb4_i_i_i_9 2 %sw_bb1_i_i_i_9 1 %sw_bb1_i_i_i_9 3 %sw_bb2_i_i_i_9 4 %sw_bb3_i_i_i_9
%sw_bb1_i_i_i_9 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_9
%sw_bb2_i_i_i_9 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_9
%sw_bb3_i_i_i_9 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_9
%sw_bb4_i_i_i_9 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_9
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_9 = OpLabel
%SpvOrder_0_i_i_i_9 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_9 %uint_904 %sw_bb3_i_i_i_9 %uint_900 %sw_bb2_i_i_i_9 %uint_898 %sw_bb1_i_i_i_9 %uint_896 %if_then_i_9
OpMemoryBarrier %uint_1 %SpvOrder_0_i_i_i_9
%1007 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_9
%1008 = OpFunctionCall %void %__itt_offload_atomic_op_start %1007 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_9 %uint_2 %uint_896 %uint_1
%1009 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_9
%1010 = OpFunctionCall %void %__itt_offload_atomic_op_finish %1009 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_9 = OpLabel
%i_0_i_9 = OpPhi %ulong %inc_i_9 %for_inc_i_9 %ulong_0 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%cmp_i_9 = OpULessThan %bool %i_0_i_9 %ulong_256
OpBranchConditional %cmp_i_9 %for_body_i_9 %cleanup_i_9
%for_body_i_9 = OpLabel
%1011 = OpFunctionCall %void %__itt_offload_atomic_op_start %1006 %uint_0 %uint_0
%call3_i_i_i_i_9 = OpAtomicLoad %uint %arrayidx_ascast_i_i_9 %uint_2 %uint_896
%1012 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_9
%1013 = OpFunctionCall %void %__itt_offload_atomic_op_finish %1012 %uint_0 %uint_0
%cmp6_i_6 = OpIEqual %bool %call3_i_i_i_i_9 %uint_1
OpBranchConditional %cmp6_i_6 %cleanup_i_9 %for_inc_i_9
%for_inc_i_9 = OpLabel
%inc_i_9 = OpIAdd %ulong %i_0_i_9 %ulong_1
OpBranch %for_cond_i_9
%cleanup_i_9 = OpLabel
OpSwitch %_arg_order_read_9 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_6 5 %sw_bb4_i_i25_i_6 2 %sw_bb1_i_i22_i_6 1 %sw_bb1_i_i22_i_6 3 %sw_bb2_i_i23_i_6 4 %sw_bb3_i_i24_i_6
%sw_bb1_i_i22_i_6 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_6
%sw_bb2_i_i23_i_6 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_6
%sw_bb3_i_i24_i_6 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_6
%sw_bb4_i_i25_i_6 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_6
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_6 = OpLabel
%SpvOrder_0_i_i26_i_6 = OpPhi %uint %uint_912 %sw_bb4_i_i25_i_6 %uint_904 %sw_bb3_i_i24_i_6 %uint_900 %sw_bb2_i_i23_i_6 %uint_898 %sw_bb1_i_i22_i_6 %uint_896 %cleanup_i_9
OpMemoryBarrier %uint_1 %SpvOrder_0_i_i26_i_6
OpBranchConditional %cmp_i_9 %if_then8_i_6 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then8_i_6 = OpLabel
%1014 = OpLoad %uint %_arg_data_acc_9 Aligned 4
%cmp10_not_i_6 = OpIEqual %bool %1014 %_arg_value_9
OpBranchConditional %cmp10_not_i_6 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then11_i_6
%if_then11_i_6 = OpLabel
OpStore %add_ptr_i_9 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%1015 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%13 = OpFunction %void None %767
%_arg_sync_flag_acc_10 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%_arg_sync_flag_acc3_2 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%_arg_data_acc_10 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%_arg_data_acc6_2 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%_arg_value_10 = OpFunctionParameter %uint
%_arg_order_write_10 = OpFunctionParameter %uint
%_arg_order_read_10 = OpFunctionParameter %uint
%_arg_res_acc_10 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_10 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_10 = OpLabel
%1016 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%1017 = OpBitcast %_ptr_Function_ulong %_arg_sync_flag_acc3_2
%1018 = OpLoad %ulong %1017 Aligned 8
%add_ptr_i_10 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %_arg_sync_flag_acc_10 %1018
%1019 = OpBitcast %_ptr_Function_ulong %_arg_data_acc6_2
%1020 = OpLoad %ulong %1019 Aligned 8
%add_ptr_i39_2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %_arg_data_acc_10 %1020
%1021 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_10
%1022 = OpLoad %ulong %1021 Aligned 8
%add_ptr_i53_2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_10 %1022
%1023 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%1024 = OpCompositeExtract %ulong %1023 0
%1025 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%1026 = OpCompositeExtract %ulong %1025 0
%arrayidx_ascast_i_i_10 = OpPtrCastToGeneric %_ptr_Generic_uint %add_ptr_i_10
%sub_i_i_i_i_i_10 = OpISub %ulong %1024 %1026
%cmp_i_i_i_10 = OpULessThan %bool %sub_i_i_i_i_i_10 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_10
%cmp_i_i_2 = OpIEqual %bool %1024 %1026
OpBranchConditional %cmp_i_i_2 %if_then_i_10 %for_cond_i_preheader_2
%for_cond_i_preheader_2 = OpLabel
%1027 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_10
OpBranch %for_cond_i_10
%if_then_i_10 = OpLabel
OpStore %add_ptr_i39_2 %_arg_value_10 Aligned 4
OpSwitch %_arg_order_write_10 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_10 5 %sw_bb4_i_i_i_10 2 %sw_bb1_i_i_i_10 1 %sw_bb1_i_i_i_10 3 %sw_bb2_i_i_i_10 4 %sw_bb3_i_i_i_10
%sw_bb1_i_i_i_10 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_10
%sw_bb2_i_i_i_10 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_10
%sw_bb3_i_i_i_10 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_10
%sw_bb4_i_i_i_10 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_10
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_10 = OpLabel
%SpvOrder_0_i_i_i_10 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_10 %uint_904 %sw_bb3_i_i_i_10 %uint_900 %sw_bb2_i_i_i_10 %uint_898 %sw_bb1_i_i_i_10 %uint_896 %if_then_i_10
OpMemoryBarrier %uint_1 %SpvOrder_0_i_i_i_10
%1028 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_10
%1029 = OpFunctionCall %void %__itt_offload_atomic_op_start %1028 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_10 %uint_2 %uint_896 %uint_1
%1030 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_10
%1031 = OpFunctionCall %void %__itt_offload_atomic_op_finish %1030 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_10 = OpLabel
%i_0_i_10 = OpPhi %ulong %inc_i_10 %for_inc_i_10 %ulong_0 %for_cond_i_preheader_2
%cmp_i_10 = OpULessThan %bool %i_0_i_10 %ulong_256
OpBranchConditional %cmp_i_10 %for_body_i_10 %cleanup_i_10
%for_body_i_10 = OpLabel
%1032 = OpFunctionCall %void %__itt_offload_atomic_op_start %1027 %uint_0 %uint_0
%call3_i_i_i_i_10 = OpAtomicLoad %uint %arrayidx_ascast_i_i_10 %uint_2 %uint_896
%1033 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_10
%1034 = OpFunctionCall %void %__itt_offload_atomic_op_finish %1033 %uint_0 %uint_0
%cmp7_i_2 = OpIEqual %bool %call3_i_i_i_i_10 %uint_1
OpBranchConditional %cmp7_i_2 %cleanup_i_10 %for_inc_i_10
%for_inc_i_10 = OpLabel
%inc_i_10 = OpIAdd %ulong %i_0_i_10 %ulong_1
OpBranch %for_cond_i_10
%cleanup_i_10 = OpLabel
OpSwitch %_arg_order_read_10 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_2 5 %sw_bb4_i_i31_i_2 2 %sw_bb1_i_i28_i_2 1 %sw_bb1_i_i28_i_2 3 %sw_bb2_i_i29_i_2 4 %sw_bb3_i_i30_i_2
%sw_bb1_i_i28_i_2 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_2
%sw_bb2_i_i29_i_2 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_2
%sw_bb3_i_i30_i_2 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_2
%sw_bb4_i_i31_i_2 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_2
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_2 = OpLabel
%SpvOrder_0_i_i32_i_2 = OpPhi %uint %uint_912 %sw_bb4_i_i31_i_2 %uint_904 %sw_bb3_i_i30_i_2 %uint_900 %sw_bb2_i_i29_i_2 %uint_898 %sw_bb1_i_i28_i_2 %uint_896 %cleanup_i_10
OpMemoryBarrier %uint_1 %SpvOrder_0_i_i32_i_2
OpBranchConditional %cmp_i_10 %if_then9_i_2 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then9_i_2 = OpLabel
%1035 = OpLoad %uint %add_ptr_i39_2 Aligned 4
%cmp11_not_i_2 = OpIEqual %bool %1035 %_arg_value_10
OpBranchConditional %cmp11_not_i_2 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then12_i_2
%if_then12_i_2 = OpLabel
OpStore %add_ptr_i53_2 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_5EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%1036 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%14 = OpFunction %void None %767
%_arg_sync_flag_acc_11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%_arg_sync_flag_acc3_3 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%_arg_data_acc_11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%_arg_data_acc6_3 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%_arg_value_11 = OpFunctionParameter %uint
%_arg_order_write_11 = OpFunctionParameter %uint
%_arg_order_read_11 = OpFunctionParameter %uint
%_arg_res_acc_11 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_11 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_11 = OpLabel
%1037 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%1038 = OpBitcast %_ptr_Function_ulong %_arg_sync_flag_acc3_3
%1039 = OpLoad %ulong %1038 Aligned 8
%add_ptr_i_11 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %_arg_sync_flag_acc_11 %1039
%1040 = OpBitcast %_ptr_Function_ulong %_arg_data_acc6_3
%1041 = OpLoad %ulong %1040 Aligned 8
%add_ptr_i39_3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %_arg_data_acc_11 %1041
%1042 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_11
%1043 = OpLoad %ulong %1042 Aligned 8
%add_ptr_i53_3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_11 %1043
%1044 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%1045 = OpCompositeExtract %ulong %1044 0
%1046 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%1047 = OpCompositeExtract %ulong %1046 0
%arrayidx_ascast_i_i_11 = OpPtrCastToGeneric %_ptr_Generic_uint %add_ptr_i_11
%sub_i_i_i_i_i_11 = OpISub %ulong %1045 %1047
%cmp_i_i_i_11 = OpULessThan %bool %sub_i_i_i_i_i_11 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_11
%cmp_i_i_3 = OpIEqual %bool %1045 %1047
OpBranchConditional %cmp_i_i_3 %if_then_i_11 %for_cond_i_preheader_3
%for_cond_i_preheader_3 = OpLabel
%1048 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_11
OpBranch %for_cond_i_11
%if_then_i_11 = OpLabel
OpStore %add_ptr_i39_3 %_arg_value_11 Aligned 4
OpSwitch %_arg_order_write_11 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_11 5 %sw_bb4_i_i_i_11 2 %sw_bb1_i_i_i_11 1 %sw_bb1_i_i_i_11 3 %sw_bb2_i_i_i_11 4 %sw_bb3_i_i_i_11
%sw_bb1_i_i_i_11 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_11
%sw_bb2_i_i_i_11 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_11
%sw_bb3_i_i_i_11 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_11
%sw_bb4_i_i_i_11 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_11
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_11 = OpLabel
%SpvOrder_0_i_i_i_11 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_11 %uint_904 %sw_bb3_i_i_i_11 %uint_900 %sw_bb2_i_i_i_11 %uint_898 %sw_bb1_i_i_i_11 %uint_896 %if_then_i_11
OpMemoryBarrier %uint_1 %SpvOrder_0_i_i_i_11
%1049 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_11
%1050 = OpFunctionCall %void %__itt_offload_atomic_op_start %1049 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_11 %uint_2 %uint_896 %uint_1
%1051 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_11
%1052 = OpFunctionCall %void %__itt_offload_atomic_op_finish %1051 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_11 = OpLabel
%i_0_i_11 = OpPhi %ulong %inc_i_11 %for_inc_i_11 %ulong_0 %for_cond_i_preheader_3
%cmp_i_11 = OpULessThan %bool %i_0_i_11 %ulong_256
OpBranchConditional %cmp_i_11 %for_body_i_11 %cleanup_i_11
%for_body_i_11 = OpLabel
%1053 = OpFunctionCall %void %__itt_offload_atomic_op_start %1048 %uint_0 %uint_0
%call3_i_i_i_i_11 = OpAtomicLoad %uint %arrayidx_ascast_i_i_11 %uint_2 %uint_896
%1054 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_11
%1055 = OpFunctionCall %void %__itt_offload_atomic_op_finish %1054 %uint_0 %uint_0
%cmp7_i_3 = OpIEqual %bool %call3_i_i_i_i_11 %uint_1
OpBranchConditional %cmp7_i_3 %cleanup_i_11 %for_inc_i_11
%for_inc_i_11 = OpLabel
%inc_i_11 = OpIAdd %ulong %i_0_i_11 %ulong_1
OpBranch %for_cond_i_11
%cleanup_i_11 = OpLabel
OpSwitch %_arg_order_read_11 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_3 5 %sw_bb4_i_i31_i_3 2 %sw_bb1_i_i28_i_3 1 %sw_bb1_i_i28_i_3 3 %sw_bb2_i_i29_i_3 4 %sw_bb3_i_i30_i_3
%sw_bb1_i_i28_i_3 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_3
%sw_bb2_i_i29_i_3 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_3
%sw_bb3_i_i30_i_3 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_3
%sw_bb4_i_i31_i_3 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_3
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_3 = OpLabel
%SpvOrder_0_i_i32_i_3 = OpPhi %uint %uint_912 %sw_bb4_i_i31_i_3 %uint_904 %sw_bb3_i_i30_i_3 %uint_900 %sw_bb2_i_i29_i_3 %uint_898 %sw_bb1_i_i28_i_3 %uint_896 %cleanup_i_11
OpMemoryBarrier %uint_1 %SpvOrder_0_i_i32_i_3
OpBranchConditional %cmp_i_11 %if_then9_i_3 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then9_i_3 = OpLabel
%1056 = OpLoad %uint %add_ptr_i39_3 Aligned 4
%cmp11_not_i_3 = OpIEqual %bool %1056 %_arg_value_11
OpBranchConditional %cmp11_not_i_3 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then12_i_3
%if_then12_i_3 = OpLabel
OpStore %add_ptr_i53_3 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_3EES1_INS3_12memory_scopeELS6_3EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%1057 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%15 = OpFunction %void None %767
%_arg_sync_flag_acc_12 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%_arg_sync_flag_acc3_4 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%_arg_data_acc_12 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%_arg_data_acc6_4 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%_arg_value_12 = OpFunctionParameter %uint
%_arg_order_write_12 = OpFunctionParameter %uint
%_arg_order_read_12 = OpFunctionParameter %uint
%_arg_res_acc_12 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_12 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_12 = OpLabel
%1058 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%1059 = OpBitcast %_ptr_Function_ulong %_arg_sync_flag_acc3_4
%1060 = OpLoad %ulong %1059 Aligned 8
%add_ptr_i_12 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %_arg_sync_flag_acc_12 %1060
%1061 = OpBitcast %_ptr_Function_ulong %_arg_data_acc6_4
%1062 = OpLoad %ulong %1061 Aligned 8
%add_ptr_i39_4 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %_arg_data_acc_12 %1062
%1063 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_12
%1064 = OpLoad %ulong %1063 Aligned 8
%add_ptr_i53_4 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_12 %1064
%1065 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%1066 = OpCompositeExtract %ulong %1065 0
%1067 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%1068 = OpCompositeExtract %ulong %1067 0
%arrayidx_ascast_i_i_12 = OpPtrCastToGeneric %_ptr_Generic_uint %add_ptr_i_12
%sub_i_i_i_i_i_12 = OpISub %ulong %1066 %1068
%cmp_i_i_i_12 = OpULessThan %bool %sub_i_i_i_i_i_12 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_12
%cmp_i_i_4 = OpIEqual %bool %1066 %1068
OpBranchConditional %cmp_i_i_4 %if_then_i_12 %for_cond_i_preheader_4
%for_cond_i_preheader_4 = OpLabel
%1069 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_12
OpBranch %for_cond_i_12
%if_then_i_12 = OpLabel
OpStore %add_ptr_i39_4 %_arg_value_12 Aligned 4
OpSwitch %_arg_order_write_12 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_12 5 %sw_bb4_i_i_i_12 2 %sw_bb1_i_i_i_12 1 %sw_bb1_i_i_i_12 3 %sw_bb2_i_i_i_12 4 %sw_bb3_i_i_i_12
%sw_bb1_i_i_i_12 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_12
%sw_bb2_i_i_i_12 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_12
%sw_bb3_i_i_i_12 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_12
%sw_bb4_i_i_i_12 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_12
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_12 = OpLabel
%SpvOrder_0_i_i_i_12 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_12 %uint_904 %sw_bb3_i_i_i_12 %uint_900 %sw_bb2_i_i_i_12 %uint_898 %sw_bb1_i_i_i_12 %uint_896 %if_then_i_12
OpMemoryBarrier %uint_0 %SpvOrder_0_i_i_i_12
%1070 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_12
%1071 = OpFunctionCall %void %__itt_offload_atomic_op_start %1070 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_12 %uint_2 %uint_896 %uint_1
%1072 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_12
%1073 = OpFunctionCall %void %__itt_offload_atomic_op_finish %1072 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_12 = OpLabel
%i_0_i_12 = OpPhi %ulong %inc_i_12 %for_inc_i_12 %ulong_0 %for_cond_i_preheader_4
%cmp_i_12 = OpULessThan %bool %i_0_i_12 %ulong_256
OpBranchConditional %cmp_i_12 %for_body_i_12 %cleanup_i_12
%for_body_i_12 = OpLabel
%1074 = OpFunctionCall %void %__itt_offload_atomic_op_start %1069 %uint_0 %uint_0
%call3_i_i_i_i_12 = OpAtomicLoad %uint %arrayidx_ascast_i_i_12 %uint_2 %uint_896
%1075 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_12
%1076 = OpFunctionCall %void %__itt_offload_atomic_op_finish %1075 %uint_0 %uint_0
%cmp7_i_4 = OpIEqual %bool %call3_i_i_i_i_12 %uint_1
OpBranchConditional %cmp7_i_4 %cleanup_i_12 %for_inc_i_12
%for_inc_i_12 = OpLabel
%inc_i_12 = OpIAdd %ulong %i_0_i_12 %ulong_1
OpBranch %for_cond_i_12
%cleanup_i_12 = OpLabel
OpSwitch %_arg_order_read_12 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_4 5 %sw_bb4_i_i31_i_4 2 %sw_bb1_i_i28_i_4 1 %sw_bb1_i_i28_i_4 3 %sw_bb2_i_i29_i_4 4 %sw_bb3_i_i30_i_4
%sw_bb1_i_i28_i_4 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_4
%sw_bb2_i_i29_i_4 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_4
%sw_bb3_i_i30_i_4 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_4
%sw_bb4_i_i31_i_4 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_4
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit33_i_4 = OpLabel
%SpvOrder_0_i_i32_i_4 = OpPhi %uint %uint_912 %sw_bb4_i_i31_i_4 %uint_904 %sw_bb3_i_i30_i_4 %uint_900 %sw_bb2_i_i29_i_4 %uint_898 %sw_bb1_i_i28_i_4 %uint_896 %cleanup_i_12
OpMemoryBarrier %uint_0 %SpvOrder_0_i_i32_i_4
OpBranchConditional %cmp_i_12 %if_then9_i_4 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then9_i_4 = OpLabel
%1077 = OpLoad %uint %add_ptr_i39_4 Aligned 4
%cmp11_not_i_4 = OpIEqual %bool %1077 %_arg_value_12
OpBranchConditional %cmp11_not_i_4 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then12_i_4
%if_then12_i_4 = OpLabel
OpStore %add_ptr_i53_4 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_4EES1_INS_9test_typeELS8_2EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%1078 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%16 = OpFunction %void None %759
%_arg_sync_flag_acc_13 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_data_acc_13 = OpFunctionParameter %_ptr_Workgroup_uint
%_arg_value_13 = OpFunctionParameter %uint
%_arg_order_write_13 = OpFunctionParameter %uint
%_arg_order_read_13 = OpFunctionParameter %uint
%_arg_res_acc_13 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%_arg_res_acc9_13 = OpFunctionParameter %_ptr_Function_class_sycl___V1__id
%entry_13 = OpLabel
%1079 = OpFunctionCall %void %__itt_offload_wi_start_wrapper
%1080 = OpBitcast %_ptr_Function_ulong %_arg_res_acc9_13
%1081 = OpLoad %ulong %1080 Aligned 8
%add_ptr_i_13 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %_arg_res_acc_13 %1081
%arrayidx_ascast_i_i_13 = OpPtrCastToGeneric %_ptr_Generic_uint %_arg_sync_flag_acc_13
%1082 = OpLoad %v3ulong %__spirv_BuiltInLocalInvocationId Aligned 32
%1083 = OpCompositeExtract %ulong %1082 0
%cmp_i5_i_i_7 = OpIEqual %bool %1083 %ulong_0
OpBranchConditional %cmp_i5_i_i_7 %if_then_i_13 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i = OpLabel
%1084 = OpLoad %v3ulong %__spirv_BuiltInGlobalOffset Aligned 32
%1085 = OpCompositeExtract %ulong %1084 0
%1086 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%1087 = OpCompositeExtract %ulong %1086 0
%sub_i_i_i_i_i_13 = OpISub %ulong %1087 %1085
%cmp_i_i_i_13 = OpULessThan %bool %sub_i_i_i_i_i_13 %ulong_2147483648
OpAssumeTrueKHR %cmp_i_i_i_13
%1088 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_13
OpBranch %for_cond_i_13
%if_then_i_13 = OpLabel
OpStore %_arg_data_acc_13 %_arg_value_13 Aligned 4
OpSwitch %_arg_order_write_13 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_13 5 %sw_bb4_i_i_i_13 2 %sw_bb1_i_i_i_13 1 %sw_bb1_i_i_i_13 3 %sw_bb2_i_i_i_13 4 %sw_bb3_i_i_i_13
%sw_bb1_i_i_i_13 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_13
%sw_bb2_i_i_i_13 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_13
%sw_bb3_i_i_i_13 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_13
%sw_bb4_i_i_i_13 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_13
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit_i_13 = OpLabel
%SpvOrder_0_i_i_i_13 = OpPhi %uint %uint_912 %sw_bb4_i_i_i_13 %uint_904 %sw_bb3_i_i_i_13 %uint_900 %sw_bb2_i_i_i_13 %uint_898 %sw_bb1_i_i_i_13 %uint_896 %if_then_i_13
OpMemoryBarrier %uint_2 %SpvOrder_0_i_i_i_13
%1089 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_13
%1090 = OpFunctionCall %void %__itt_offload_atomic_op_start %1089 %uint_1 %uint_0
OpAtomicStore %arrayidx_ascast_i_i_13 %uint_2 %uint_896 %uint_1
%1091 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_13
%1092 = OpFunctionCall %void %__itt_offload_atomic_op_finish %1091 %uint_1 %uint_0
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%for_cond_i_13 = OpLabel
%i_0_i_13 = OpPhi %ulong %inc_i_13 %for_inc_i_13 %ulong_0 %_ZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEE27is_specified_item_in_kernelERKNS3_7nd_itemILi1EEE_exit_i
%cmp_i_13 = OpULessThan %bool %i_0_i_13 %ulong_256
OpBranchConditional %cmp_i_13 %for_body_i_13 %cleanup_i_13
%for_body_i_13 = OpLabel
%1093 = OpFunctionCall %void %__itt_offload_atomic_op_start %1088 %uint_0 %uint_0
%call3_i_i_i_i_13 = OpAtomicLoad %uint %arrayidx_ascast_i_i_13 %uint_2 %uint_896
%1094 = OpBitcast %_ptr_Generic_uchar %arrayidx_ascast_i_i_13
%1095 = OpFunctionCall %void %__itt_offload_atomic_op_finish %1094 %uint_0 %uint_0
%cmp6_i_7 = OpIEqual %bool %call3_i_i_i_i_13 %uint_1
OpBranchConditional %cmp6_i_7 %cleanup_i_13 %for_inc_i_13
%for_inc_i_13 = OpLabel
%inc_i_13 = OpIAdd %ulong %i_0_i_13 %ulong_1
OpBranch %for_cond_i_13
%cleanup_i_13 = OpLabel
OpSwitch %_arg_order_read_13 %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_7 5 %sw_bb4_i_i25_i_7 2 %sw_bb1_i_i22_i_7 1 %sw_bb1_i_i22_i_7 3 %sw_bb2_i_i23_i_7 4 %sw_bb3_i_i24_i_7
%sw_bb1_i_i22_i_7 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_7
%sw_bb2_i_i23_i_7 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_7
%sw_bb3_i_i24_i_7 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_7
%sw_bb4_i_i25_i_7 = OpLabel
OpBranch %_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_7
%_ZN4sycl3_V1L12atomic_fenceENS0_12memory_orderENS0_12memory_scopeE_exit27_i_7 = OpLabel
%SpvOrder_0_i_i26_i_7 = OpPhi %uint %uint_912 %sw_bb4_i_i25_i_7 %uint_904 %sw_bb3_i_i24_i_7 %uint_900 %sw_bb2_i_i23_i_7 %uint_898 %sw_bb1_i_i22_i_7 %uint_896 %cleanup_i_13
OpMemoryBarrier %uint_2 %SpvOrder_0_i_i26_i_7
OpBranchConditional %cmp_i_13 %if_then8_i_7 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%if_then8_i_7 = OpLabel
%1096 = OpLoad %uint %_arg_data_acc_13 Aligned 4
%cmp10_not_i_7 = OpIEqual %bool %1096 %_arg_value_13
OpBranchConditional %cmp10_not_i_7 %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit %if_then11_i_7
%if_then11_i_7 = OpLabel
OpStore %add_ptr_i_13 %uchar_0 Aligned 1
OpBranch %_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit
%_ZZZN17atomic_fence_test16run_atomic_fenceISt17integral_constantIN4sycl3_V112memory_orderELS4_4EES1_INS3_12memory_scopeELS6_2EES1_INS_9test_typeELS8_1EEEclERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESI_SI_ENKUlRNS3_7handlerEE_clESK_ENKUlNS3_7nd_itemILi1EEEE_clESN__exit = OpLabel
%1097 = OpFunctionCall %void %__itt_offload_wi_finish_wrapper
OpReturn
OpFunctionEnd
%__itt_offload_atomic_op_start_stub = OpFunction %void DontInline %765
%object_1 = OpFunctionParameter %_ptr_Generic_uchar
%op_type_1 = OpFunctionParameter %uint
%mem_order_1 = OpFunctionParameter %uint
%entry_20 = OpLabel
%object_addr = OpVariable %_ptr_Function__ptr_Generic_uchar Function
%op_type_addr = OpVariable %_ptr_Function_uint Function
%mem_order_addr = OpVariable %_ptr_Function_uint Function
%object_addr_ascast = OpPtrCastToGeneric %_ptr_Generic__ptr_Generic_uchar %object_addr
%op_type_addr_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %op_type_addr
%mem_order_addr_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %mem_order_addr
OpStore %object_addr_ascast %object_1 Aligned 8
OpStore %op_type_addr_ascast %op_type_1 Aligned 4
OpStore %mem_order_addr_ascast %mem_order_1 Aligned 4
OpReturn
OpFunctionEnd
%__itt_offload_atomic_op_finish_stub = OpFunction %void DontInline %765
%object_2 = OpFunctionParameter %_ptr_Generic_uchar
%op_type_2 = OpFunctionParameter %uint
%mem_order_2 = OpFunctionParameter %uint
%entry_21 = OpLabel
%object_addr_0 = OpVariable %_ptr_Function__ptr_Generic_uchar Function
%op_type_addr_0 = OpVariable %_ptr_Function_uint Function
%mem_order_addr_0 = OpVariable %_ptr_Function_uint Function
%object_addr_ascast_0 = OpPtrCastToGeneric %_ptr_Generic__ptr_Generic_uchar %object_addr_0
%op_type_addr_ascast_0 = OpPtrCastToGeneric %_ptr_Generic_uint %op_type_addr_0
%mem_order_addr_ascast_0 = OpPtrCastToGeneric %_ptr_Generic_uint %mem_order_addr_0
OpStore %object_addr_ascast_0 %object_2 Aligned 8
OpStore %op_type_addr_ascast_0 %op_type_2 Aligned 4
OpStore %mem_order_addr_ascast_0 %mem_order_2 Aligned 4
OpReturn
OpFunctionEnd
%__itt_offload_wi_start_stub = OpFunction %void DontInline %772
%group_id = OpFunctionParameter %_ptr_Generic_ulong
%wi_id = OpFunctionParameter %ulong
%wg_size = OpFunctionParameter %uint
%entry_18 = OpLabel
%group_id_addr = OpVariable %_ptr_Function__ptr_Generic_ulong Function
%wi_id_addr = OpVariable %_ptr_Function_ulong Function
%wg_size_addr = OpVariable %_ptr_Function_uint Function
%group_id_addr_ascast = OpPtrCastToGeneric %_ptr_Generic__ptr_Generic_ulong %group_id_addr
%wi_id_addr_ascast = OpPtrCastToGeneric %_ptr_Generic_ulong %wi_id_addr
%wg_size_addr_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %wg_size_addr
OpStore %group_id_addr_ascast %group_id Aligned 8
OpStore %wi_id_addr_ascast %wi_id Aligned 8
OpStore %wg_size_addr_ascast %wg_size Aligned 4
OpReturn
OpFunctionEnd
%__itt_offload_wi_finish_stub = OpFunction %void DontInline %773
%group_id_0 = OpFunctionParameter %_ptr_Generic_ulong
%wi_id_0 = OpFunctionParameter %ulong
%entry_19 = OpLabel
%group_id_addr_0 = OpVariable %_ptr_Function__ptr_Generic_ulong Function
%wi_id_addr_0 = OpVariable %_ptr_Function_ulong Function
%group_id_addr_ascast_0 = OpPtrCastToGeneric %_ptr_Generic__ptr_Generic_ulong %group_id_addr_0
%wi_id_addr_ascast_0 = OpPtrCastToGeneric %_ptr_Generic_ulong %wi_id_addr_0
OpStore %group_id_addr_ascast_0 %group_id_0 Aligned 8
OpStore %wi_id_addr_ascast_0 %wi_id_0 Aligned 8
OpReturn
OpFunctionEnd
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment