; ModuleID = 'before_gvn.txt'
source_filename = "llvm-link"
target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda-sycldevice"

%"class._ZTSN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EEE.cl::sycl::accessor" = type { %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice", i32 addrspace(1)* }
%"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" = type { %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" }
%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] }
%"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon" = type { %"class._ZTSN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EEE.cl::sycl::accessor" }

$"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE11hpar_kernel" = comdat any

@"_ZZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_5groupILi1EEEE_clES5_E9wg_offset" = internal unnamed_addr addrspace(3) global i64 0, align 8
@WGCopy1.0 = internal unnamed_addr addrspace(3) global %"class._ZTSN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EEE.cl::sycl::accessor"* undef, align 8
@WGCopy = internal unnamed_addr addrspace(3) global %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon"* undef, align 8

; Function Attrs: noinline
define weak_odr dso_local void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE11hpar_kernel"(i32 addrspace(1)* %0, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %2, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %3) local_unnamed_addr #0 comdat !kernel_arg_addr_space !11 !kernel_arg_access_qual !12 !kernel_arg_type !13 !kernel_arg_base_type !13 !kernel_arg_type_qual !14 {
  %5 = alloca %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon", align 8
  %6 = bitcast %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon"* %5 to i8*
  call void @llvm.lifetime.start.p0i8(i64 32, i8* nonnull %6) #4
  %7 = getelementptr inbounds %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon", %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon"* %5, i64 0, i32 0, i32 0, i32 0, i32 0, i32 0, i64 0
  %8 = getelementptr inbounds %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon", %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon"* %5, i64 0, i32 0, i32 0, i32 1, i32 0, i32 0, i64 0
  %9 = getelementptr inbounds %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon", %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon"* %5, i64 0, i32 0, i32 0, i32 2, i32 0, i32 0, i64 0
  %10 = getelementptr inbounds %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %1, i64 0, i32 0, i32 0, i64 0
  %11 = load i64, i64* %10, align 8
  %12 = getelementptr inbounds %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %2, i64 0, i32 0, i32 0, i64 0
  %13 = load i64, i64* %12, align 8
  %14 = getelementptr inbounds %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %3, i64 0, i32 0, i32 0, i64 0
  %15 = load i64, i64* %14, align 8
  %16 = getelementptr inbounds %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon", %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon"* %5, i64 0, i32 0, i32 1
  store i64 %15, i64* %7, align 8, !tbaa !15
  store i64 %11, i64* %8, align 8, !tbaa !15
  store i64 %13, i64* %9, align 8, !tbaa !15
  %17 = getelementptr inbounds i32, i32 addrspace(1)* %0, i64 %15
  store i32 addrspace(1)* %17, i32 addrspace(1)** %16, align 8, !tbaa !19
  %18 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x() #4, !range !25
  %19 = zext i32 %18 to i64
  %20 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.y() #4, !range !25
  %21 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.z() #4, !range !26
  %22 = zext i32 %21 to i64
  %23 = tail call i32 @llvm.nvvm.read.ptx.sreg.ntid.y() #4, !range !27
  %24 = tail call i32 @llvm.nvvm.read.ptx.sreg.ntid.z() #4, !range !28
  %25 = mul nuw nsw i32 %23, %18
  %26 = add nuw nsw i32 %25, %20
  %27 = mul nuw nsw i32 %26, %24
  %28 = zext i32 %27 to i64
  %29 = sub nsw i64 0, %22
  %30 = icmp eq i64 %28, %29
  tail call void @llvm.nvvm.barrier0() #4
  br i1 %30, label %31, label %._crit_edge

._crit_edge:                                      ; preds = %4
  %.pre = load %"class._ZTSN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EEE.cl::sycl::accessor"*, %"class._ZTSN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EEE.cl::sycl::accessor"* addrspace(3)* @WGCopy1.0, align 8
  br label %37

31:                                               ; preds = %4
  %32 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #4, !range !29
  %33 = getelementptr inbounds %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon", %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon"* %5, i64 0, i32 0
  %34 = ptrtoint %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon"* %5 to i64
  %35 = mul nuw nsw i32 %32, 3
  %36 = zext i32 %35 to i64
  store i64 %36, i64* addrspacecast (i64 addrspace(3)* @"_ZZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_5groupILi1EEEE_clES5_E9wg_offset" to i64*), align 8, !tbaa !15
  store %"class._ZTSN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EEE.cl::sycl::accessor"* %33, %"class._ZTSN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EEE.cl::sycl::accessor"* addrspace(3)* @WGCopy1.0, align 8
  store i64 %34, i64 addrspace(3)* bitcast (%"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_5groupILi1EEEE_.anon"* addrspace(3)* @WGCopy to i64 addrspace(3)*), align 8
  br label %37

37:                                               ; preds = %._crit_edge, %31
  %38 = phi %"class._ZTSN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EEE.cl::sycl::accessor"* [ %.pre, %._crit_edge ], [ %33, %31 ]
  call void @llvm.nvvm.barrier0() #4
  call void @llvm.nvvm.barrier0() #4
  call void @llvm.nvvm.barrier0() #4
  %39 = load i64, i64* addrspacecast (i64 addrspace(3)* @"_ZZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_5groupILi1EEEE_clES5_E9wg_offset" to i64*), align 8, !tbaa !15
  %40 = add i64 %39, %19
  %41 = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EEE.cl::sycl::accessor"* %38, i64 0, i32 1
  %42 = load i32 addrspace(1)*, i32 addrspace(1)** %41, align 8, !tbaa !19
  %43 = getelementptr inbounds i32, i32 addrspace(1)* %42, i64 %40
  %44 = addrspacecast i32 addrspace(1)* %43 to i32*
  store i32 10, i32* %44, align 4, !tbaa !30
  call void @llvm.nvvm.barrier0() #4
  call void @llvm.lifetime.end.p0i8(i64 32, i8* nonnull %6) #4
  ret void
}

; Function Attrs: argmemonly nounwind willreturn
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1

; Function Attrs: nounwind readnone
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() #2

; Function Attrs: nounwind readnone
declare i32 @llvm.nvvm.read.ptx.sreg.tid.y() #2

; Function Attrs: nounwind readnone
declare i32 @llvm.nvvm.read.ptx.sreg.tid.z() #2

; Function Attrs: nounwind readnone
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.y() #2

; Function Attrs: nounwind readnone
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.z() #2

; Function Attrs: convergent nounwind
declare void @llvm.nvvm.barrier0() #3

; Function Attrs: nounwind readnone
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #2

; Function Attrs: argmemonly nounwind willreturn
declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1

attributes #0 = { noinline "correctly-rounded-divide-sqrt-fp-math"="false" "denormal-fp-math"="ieee,ieee" "denormal-fp-math-f32"="ieee,ieee" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "sycl-module-id"="test_ptx.cpp" "target-cpu"="sm_30" "target-features"="+ptx64,+sm_30" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #1 = { argmemonly nounwind willreturn }
attributes #2 = { nounwind readnone }
attributes #3 = { convergent nounwind }
attributes #4 = { nounwind }

!nvvm.annotations = !{!0, !1, !2, !1, !3, !3, !3, !3, !4, !4, !3}
!opencl.spir.version = !{!5}
!spirv.Source = !{!6}
!llvm.ident = !{!7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7}
!nvvmir.version = !{!8}
!llvm.module.flags = !{!9, !10}

!0 = !{void (i32 addrspace(1)*, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"*, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"*, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"*)* @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE11hpar_kernel", !"kernel", i32 1}
!1 = !{null, !"align", i32 8}
!2 = !{null, !"align", i32 8, !"align", i32 65544, !"align", i32 131080}
!3 = !{null, !"align", i32 16}
!4 = !{null, !"align", i32 16, !"align", i32 65552, !"align", i32 131088}
!5 = !{i32 1, i32 2}
!6 = !{i32 4, i32 100000}
!7 = !{!"clang version 11.0.0"}
!8 = !{i32 1, i32 4}
!9 = !{i32 2, !"SDK Version", [2 x i32] [i32 10, i32 1]}
!10 = !{i32 1, !"wchar_size", i32 4}
!11 = !{i32 0, i32 0, i32 0, i32 0}
!12 = !{!"none", !"none", !"none", !"none"}
!13 = !{!"int*", !"cl::sycl::range<1>", !"cl::sycl::range<1>", !"cl::sycl::id<1>"}
!14 = !{!"", !"", !"", !""}
!15 = !{!16, !16, i64 0}
!16 = !{!"long", !17, i64 0}
!17 = !{!"omnipotent char", !18, i64 0}
!18 = !{!"Simple C++ TBAA"}
!19 = !{!20, !24, i64 24}
!20 = !{!"_ZTSN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EEE", !21, i64 0, !24, i64 24}
!21 = !{!"_ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE", !22, i64 0, !23, i64 8, !23, i64 16}
!22 = !{!"_ZTSN2cl4sycl2idILi1EEE"}
!23 = !{!"_ZTSN2cl4sycl5rangeILi1EEE"}
!24 = !{!"any pointer", !17, i64 0}
!25 = !{i32 0, i32 1024}
!26 = !{i32 0, i32 64}
!27 = !{i32 1, i32 1025}
!28 = !{i32 1, i32 65}
!29 = !{i32 0, i32 65535}
!30 = !{!31, !31, i64 0}
!31 = !{!"int", !17, i64 0}