-
-
Save karolherbst/2965158c53342bb7b3f87b85d97f07e2 to your computer and use it in GitHub Desktop.
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
; 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