Artem-B · October 11, 2021 20:24 · Artem-B · Oct 11, 2021
diff --git a/SROA-input-head.ll b/SROA-input-head.ll
 ;*** IR Dump After Straight line strength reduction (slsr) *** (function: _ZN8cuforces12forcesDeviceI13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSD_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyI18xsph_forces_paramsESF_I20volume_forces_paramsESF_I21grenier_forces_paramsESF_I25sa_boundary_forces_paramsESF_I28dummy_boundary_forces_paramsESF_I25water_depth_forces_paramsESF_I18keps_forces_paramsESF_I14tau_tex_paramsESF_I22eulerVel_forces_paramsESF_I29internal_energy_forces_paramsESF_I28effective_visc_forces_paramsEELS2_3ELS3_1ELS4_3ELS5_4ESC_Lm517ELSD_1ELSD_0EEEvT_)
 ; ModuleID = 'reduced.ll.ll'
 source_filename = "<stdin>"
 target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-nvidia-cuda"

 %struct.char3 = type { i8, i8, i8 }
 %struct.float4 = type { float, float, float, float }
 %struct.float3 = type { float, float, float }
 %struct.int3 = type { i32, i32, i32 }
 %struct.forces_params.415 = type { %struct.common_forces_params }
 %struct.common_forces_params = type { %struct.stage_common_forces_params.base, i16*, float, float, i32, float }
 %struct.stage_common_forces_params.base = type <{ %struct.pos_info_wrapper, %struct.vel_wrapper, %struct.float4*, i32*, i32*, i32, i32, float }>
 %struct.pos_info_wrapper = type { %struct.pos_wrapper, %struct.info_wrapper }
 %struct.pos_wrapper = type { %struct.float4* }
 %struct.info_wrapper = type { %struct.ushort4* }
 %struct.ushort4 = type { i16, i16, i16, i16 }
 %struct.vel_wrapper = type { %struct.float4* }
 %"class.cuneibs::neiblist_iterator.1" = type { %"class.cuneibs::neiblist_iterator_simple.base.3", [8 x i8], %"class.cuneibs::neiblist_iterator_core.base", [12 x i8] }
 %"class.cuneibs::neiblist_iterator_simple.base.3" = type { i32 (...)** }
 %"class.cuneibs::neiblist_iterator_core.base" = type <{ i32*, i16*, %struct.float4, %struct.int3, i32, %struct.float3, [4 x i8], i64, i32, i8, [3 x i8], i32 }>
 %struct.cell_params = type { %struct.cellStart_wrapper, %struct.cellEnd_wrapper }
 %struct.cellStart_wrapper = type { i32* }
 %struct.cellEnd_wrapper = type { i32* }
 %struct.uint4 = type { i32, i32, i32, i32 }
 %struct.jacobi_update_params = type <{ %struct.info_wrapper, %struct.float4*, float*, float*, i32, [4 x i8] }>
 %struct.sa_outgoing_bc_params = type { %struct.neibs_list_params.base, %struct.vel_wrapper, %struct.boundelements_wrapper, %struct.vertPos_params, %struct.uint4*, %struct.float4* }
 %struct.neibs_list_params.base = type <{ %struct.pos_info_wrapper, i32*, i32*, i16*, i32, float, float }>
 %struct.boundelements_wrapper = type { %struct.float4* }
 %struct.vertPos_params = type { %struct.float2.0*, %struct.float2.0*, %struct.float2.0* }
 %struct.float2.0 = type { float, float }
 %struct.reorder_params = type { %struct.reorder_data, %struct.reorder_data.9, %struct.reorder_data.30 }
 %struct.reorder_data = type { %struct.float4*, %struct.float4* }
 %struct.reorder_data.9 = type { %struct.float4*, %struct.float4* }
 %struct.reorder_data.30 = type { %struct.float4*, %struct.float4* }
 %"class.thrust::zip_iterator" = type { %"class.thrust::tuple.32" }
 %"class.thrust::tuple.32" = type { %"struct.thrust::detail::cons.33" }
 %"struct.thrust::detail::cons.33" = type { %"class.thrust::device_ptr", %"struct.thrust::detail::cons.35" }
 %"class.thrust::device_ptr" = type { %"class.thrust::pointer" }
 %"class.thrust::pointer" = type { %"class.thrust::iterator_adaptor" }
 %"class.thrust::iterator_adaptor" = type { i32* }
 %"struct.thrust::detail::cons.35" = type { %"class.thrust::device_ptr.36" }
 %"class.thrust::device_ptr.36" = type { %"class.thrust::pointer.37" }
 %"class.thrust::pointer.37" = type { %"class.thrust::iterator_adaptor.38" }
 %"class.thrust::iterator_adaptor.38" = type { %struct.ushort4* }
 %"class.thrust::tuple" = type { %"struct.thrust::detail::cons" }
 %"struct.thrust::detail::cons" = type { i32, [4 x i8], %"struct.thrust::detail::cons.31" }
 %"struct.thrust::detail::cons.31" = type { %struct.ushort4 }
 %struct.ptype_hash_compare = type { i8 }
 %"struct.thrust::cuda_cub::__transform::unary_transform_f" = type <{ %"class.thrust::device_ptr", i32*, %"struct.thrust::identity", %"struct.thrust::cuda_cub::__transform::always_true_predicate", [6 x i8] }>
 %"struct.thrust::identity" = type { i8 }
 %"struct.thrust::cuda_cub::__transform::always_true_predicate" = type { i8 }
 %"struct.thrust::cuda_cub::__transform::unary_transform_f.69" = type <{ %"class.thrust::device_ptr.36", %struct.ushort4*, %"struct.thrust::identity.70", %"struct.thrust::cuda_cub::__transform::always_true_predicate", [6 x i8] }>
 %"struct.thrust::identity.70" = type { i8 }
 %"struct.thrust::cuda_cub::__transform::unary_transform_f.72" = type <{ i32*, %"class.thrust::device_ptr", %"struct.thrust::identity", %"struct.thrust::cuda_cub::__transform::always_true_predicate", [6 x i8] }>
 %"struct.thrust::cuda_cub::__transform::unary_transform_f.74" = type <{ %struct.ushort4*, %"class.thrust::device_ptr.36", %"struct.thrust::identity.70", %"struct.thrust::cuda_cub::__transform::always_true_predicate", [6 x i8] }>
 %"struct.thrust::cuda_cub::__transform::unary_transform_f.79" = type <{ %"class.thrust::device_ptr.36", %"class.thrust::device_ptr.36", %"struct.thrust::identity.70", %"struct.thrust::cuda_cub::__transform::always_true_predicate", [6 x i8] }>
 %"struct.thrust::cuda_cub::__transform::unary_transform_f.81" = type <{ %"class.thrust::device_ptr", %"class.thrust::device_ptr", %"struct.thrust::identity", %"struct.thrust::cuda_cub::__transform::always_true_predicate", [6 x i8] }>
 %struct.buildneibs_params = type { %struct.common_buildneibs_params, %struct.planes_buildneibs_params }
 %struct.common_buildneibs_params = type { %struct.pos_info_wrapper, %struct.cell_params, i32*, i16*, i32, float }
 %struct.planes_buildneibs_params = type { %struct.int4* }
 %struct.int4 = type { i32, i32, i32, i32 }
 %struct.euler_params = type { %struct.common_euler_params.base, %struct.dummy_euler_params }
 %struct.common_euler_params.base = type <{ %struct.Pos_params, %struct.Vel_params, i32*, %struct.ushort4*, %struct.float4*, i32, float, float }>
 %struct.Pos_params = type { %struct.float4*, %struct.float4* }
 %struct.Vel_params = type { %struct.float4*, %struct.float4* }
 %struct.dummy_euler_params = type { %struct.float4* }
 %struct.euler_params.111 = type { %struct.common_euler_params.base, %struct.dummy_euler_params }
 %struct.euler_params.114 = type { %struct.common_euler_params.base, %struct.dummy_euler_params }
 %struct.euler_params.117 = type { %struct.common_euler_params.base, %struct.dummy_euler_params }
 %"struct.cub::ReduceByKeyScanTileState" = type { %"struct.cub::ScanTileState" }
 %"struct.cub::ScanTileState" = type { i8*, %"struct.cub::KeyValuePair"*, %"struct.cub::KeyValuePair"* }
 %"struct.cub::KeyValuePair" = type { i32, [12 x i8], %struct.float4 }
 %"class.thrust::device_ptr.120" = type { %"class.thrust::pointer.121" }
 %"class.thrust::pointer.121" = type { %"class.thrust::iterator_adaptor.122" }
 %"class.thrust::iterator_adaptor.122" = type { i32* }
 %"class.thrust::device_ptr.124" = type { %"class.thrust::pointer.125" }
 %"class.thrust::pointer.125" = type { %"class.thrust::iterator_adaptor.126" }
 %"class.thrust::iterator_adaptor.126" = type { %struct.float4* }
 %"struct.thrust::equal_to" = type { i8 }
 %"struct.thrust::plus" = type { i8 }
 %"struct.thrust::cuda_cub::__scan_by_key::DoNothing" = type { i8 }
 %"struct.thrust::cuda_cub::__transform::unary_transform_f.166" = type <{ %struct.float4*, %"class.thrust::device_ptr.124", %"struct.thrust::identity.167", %"struct.thrust::cuda_cub::__transform::always_true_predicate", [6 x i8] }>
 %"struct.thrust::identity.167" = type { i8 }
 %"struct.thrust::cuda_cub::__transform::unary_transform_f.169" = type <{ %"class.thrust::device_ptr.120", i32*, %"struct.thrust::identity", %"struct.thrust::cuda_cub::__transform::always_true_predicate", [6 x i8] }>
 %"struct.thrust::cuda_cub::__transform::unary_transform_f.171" = type <{ %"class.thrust::device_ptr.124", %struct.float4*, %"struct.thrust::identity.167", %"struct.thrust::cuda_cub::__transform::always_true_predicate", [6 x i8] }>
 %struct.density_diffusion_params = type { %struct.common_density_diffusion_params.base, [4 x i8] }
 %struct.common_density_diffusion_params.base = type <{ %struct.float4*, %struct.float4*, %struct.float4*, %struct.ushort4*, i32*, i32*, i16*, i32, float, float, float, float }>
 %struct.forces_params = type { %struct.common_forces_params }
 %struct.forces_params.223 = type { %struct.common_forces_params, %struct.dummy_boundary_forces_params }
 %struct.dummy_boundary_forces_params = type { %struct.float4* }
 %struct.finalize_forces_params = type { %struct.common_finalize_forces_params, %struct.planes_forces_params, %struct.dyndt_finalize_forces_params }
 %struct.common_finalize_forces_params = type { %struct.stage_common_forces_params.base, %struct.float4*, float, %struct.float4*, %struct.float4* }
 %struct.planes_forces_params = type { %struct.int4* }
 %struct.dyndt_finalize_forces_params = type { float*, float*, float*, i32, i32 }
 %struct.forces_params.236 = type { %struct.common_forces_params }
 %struct.forces_params.250 = type { %struct.common_forces_params, %struct.dummy_boundary_forces_params }
 %struct.forces_params.256 = type { %struct.common_forces_params }
 %struct.finalize_forces_params.262 = type { %struct.common_finalize_forces_params.263, %struct.planes_forces_params, %struct.dyndt_finalize_forces_params }
 %struct.common_finalize_forces_params.263 = type { %struct.stage_common_forces_params.base, %struct.float4*, %struct.float4*, %struct.float4* }
 %struct.neibs_interaction_params = type { %struct.neibs_list_params.base, %struct.vel_wrapper }
 %"struct.cupostprocess::testpoints_params" = type { %struct.neibs_interaction_params, %struct.float4* }
 %struct.reorder_params.278 = type { %struct.reorder_data, %struct.reorder_data.9, %struct.reorder_data.30 }
 %struct.buildneibs_params.279 = type { %struct.common_buildneibs_params }
 %struct.euler_params.282 = type { %struct.common_euler_params.base, %struct.dummy_euler_params }
 %struct.euler_params.285 = type { %struct.common_euler_params.base, %struct.dummy_euler_params }
 %struct.euler_params.288 = type { %struct.common_euler_params.base, %struct.dummy_euler_params }
 %struct.euler_params.291 = type { %struct.common_euler_params.base, %struct.dummy_euler_params }
 %struct.forces_params.294 = type { %struct.common_forces_params }
 %struct.forces_params.300 = type { %struct.common_forces_params, %struct.dummy_boundary_forces_params }
 %struct.finalize_forces_params.305 = type { %struct.common_finalize_forces_params, %struct.dyndt_finalize_forces_params }
 %struct.forces_params.312 = type { %struct.common_forces_params }
 %struct.forces_params.319 = type { %struct.common_forces_params, %struct.dummy_boundary_forces_params }
 %struct.forces_params.325 = type { %struct.common_forces_params }
 %struct.finalize_forces_params.331 = type { %struct.common_finalize_forces_params.263, %struct.dyndt_finalize_forces_params }
 %"struct.cupostprocess::testpoints_params.336" = type { %struct.neibs_interaction_params, %struct.float4* }
 %struct.density_diffusion_params.339 = type { %struct.common_density_diffusion_params.base, [4 x i8] }
 %struct.forces_params.341 = type { %struct.common_forces_params }
 %struct.forces_params.346 = type { %struct.common_forces_params, %struct.dummy_boundary_forces_params }
 %struct.forces_params.351 = type { %struct.common_forces_params }
 %struct.forces_params.356 = type { %struct.common_forces_params }
 %struct.forces_params.361 = type { %struct.common_forces_params, %struct.dummy_boundary_forces_params }
 %struct.forces_params.366 = type { %struct.common_forces_params }
 %struct.density_diffusion_params.371 = type { %struct.common_density_diffusion_params.base, [4 x i8] }
 %struct.forces_params.373 = type { %struct.common_forces_params }
 %struct.forces_params.378 = type { %struct.common_forces_params, %struct.dummy_boundary_forces_params }
 %struct.forces_params.383 = type { %struct.common_forces_params }
 %struct.forces_params.388 = type { %struct.common_forces_params }
 %struct.forces_params.393 = type { %struct.common_forces_params, %struct.dummy_boundary_forces_params }
 %struct.forces_params.398 = type { %struct.common_forces_params }
 %struct.density_diffusion_params.403 = type { %struct.common_density_diffusion_params.base, [4 x i8] }
 %struct.forces_params.405 = type { %struct.common_forces_params }
 %struct.forces_params.410 = type { %struct.common_forces_params, %struct.dummy_boundary_forces_params }
 %struct.forces_params.420 = type { %struct.common_forces_params }
 %struct.forces_params.425 = type { %struct.common_forces_params, %struct.dummy_boundary_forces_params }
 %struct.forces_params.430 = type { %struct.common_forces_params }

 @_ZN7cuneibs16d_cell_to_offsetE = external local_unnamed_addr addrspace(4) externally_initialized global [27 x %struct.char3], align 1
 @_ZN7cuneibs17d_neiblist_strideE = external local_unnamed_addr addrspace(4) externally_initialized global i64, align 8
 @_ZTVN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEE = linkonce_odr unnamed_addr addrspace(1) constant { [3 x i8*] } { [3 x i8*] [i8* inttoptr (i64 16 to i8*), i8* null, i8* null] }, align 8

 ; Function Attrs: argmemonly mustprogress nofree nounwind willreturn
 declare void @llvm.memcpy.p0i8.p0i8.i64(i8* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg) #0

 declare i1 @_ZL8isfinitef(float) local_unnamed_addr

 declare %struct.float4 @_ZL11make_float4ffff(float, float) local_unnamed_addr

 declare %struct.float3 @_ZL11make_float3fff(float, float) local_unnamed_addr

 declare %struct.float3 @_ZmlRK5char3RK6float3() local_unnamed_addr

 declare %struct.int3 @_ZL9make_int3iii(i32, i32, i32) local_unnamed_addr

 define void @_ZN8cuforces12forcesDeviceI13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSD_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyI18xsph_forces_paramsESF_I20volume_forces_paramsESF_I21grenier_forces_paramsESF_I25sa_boundary_forces_paramsESF_I28dummy_boundary_forces_paramsESF_I25water_depth_forces_paramsESF_I18keps_forces_paramsESF_I14tau_tex_paramsESF_I22eulerVel_forces_paramsESF_I29internal_energy_forces_paramsESF_I28effective_visc_forces_paramsEELS2_3ELS3_1ELS4_3ELS5_4ESC_Lm517ELSD_1ELSD_0EEEvT_(%struct.forces_params.415* nocapture readonly %0) local_unnamed_addr {
  %2 = addrspacecast %struct.forces_params.415* %0 to %struct.forces_params.415 addrspace(1)*
  %3 = alloca %"class.cuneibs::neiblist_iterator.1", align 16
  %4 = addrspacecast %"class.cuneibs::neiblist_iterator.1"* %3 to %"class.cuneibs::neiblist_iterator.1" addrspace(5)*
  %5 = getelementptr inbounds %struct.forces_params.415, %struct.forces_params.415 addrspace(1)* %2, i64 0, i32 0, i32 0, i32 0, i32 0, i32 0
  %6 = bitcast %struct.forces_params.415 addrspace(1)* %2 to i8* addrspace(1)*
  %7 = load i8*, i8* addrspace(1)* %6, align 8
  tail call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 1 undef, i8* align 1 %7, i64 undef, i1 false)
  %8 = getelementptr inbounds %struct.forces_params.415, %struct.forces_params.415 addrspace(1)* %2, i64 0, i32 0, i32 1
  %9 = load i16*, i16* addrspace(1)* %8, align 8
  %10 = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 1
  store i16* %9, i16* addrspace(5)* %10, align 8
  %11 = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 0, i32 0
  store i32 (...)** bitcast (i8** getelementptr ({ [3 x i8*] }, { [3 x i8*] }* addrspacecast ({ [3 x i8*] } addrspace(1)* @_ZTVN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEE to { [3 x i8*] }*), i64 0, i32 0, i64 3) to i32 (...)**), i32 (...)** addrspace(5)* %11, align 16, !tbaa !125
  %12 = bitcast %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4 to i8 addrspace(5)*
  %13 = getelementptr inbounds i16, i16* %9, i64 undef
  %14 = load i16, i16* %13, align 2
  %15 = icmp eq i16 %14, -1
  br i1 %15, label %.critedge, label %_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit.lr.ph

 _ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit.lr.ph: ; preds = %1
  %16 = bitcast %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4 to i8* addrspace(5)*
  %17 = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 0
  %18 = load i32*, i32* addrspace(5)* %17, align 16
  %19 = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 9
  %20 = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 8
  %21 = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 2, i32 0
  %.idx.val.i.i.i = load float, float addrspace(5)* %21, align 16
  %22 = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 2, i32 2
  %.idx5.val.i.i.i = load float, float addrspace(5)* %22, align 8
  %23 = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 5, i32 0
  %.sroa.215.0..sroa_idx16.i.i.i = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 5, i32 1
  %.sroa.3.0..sroa_idx17.i.i.i = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 5, i32 2
  %24 = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 3, i32 0
  %.idx8.val.i.i.i = load i32, i32 addrspace(5)* %24, align 16
  %.idx9.i.i.i = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 3, i32 1
  %.idx9.val.i.i.i = load i32, i32 addrspace(5)* %.idx9.i.i.i, align 4
  %.idx10.i.i.i = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 3, i32 2
  %.idx10.val.i.i.i = load i32, i32 addrspace(5)* %.idx10.i.i.i, align 8
  %25 = getelementptr inbounds i32, i32* %18, i64 undef
  %26 = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 7
  %27 = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 4
  %28 = load i32, i32 addrspace(5)* %27, align 4
  %29 = zext i32 %28 to i64
  %.pre = load i8*, i8* addrspace(5)* %16, align 16
  %30 = lshr i16 %14, 11
  %31 = trunc i16 %30 to i8
  %32 = add nsw i8 %31, -1
  store i8 %32, i8 addrspace(5)* %19, align 4
  %33 = tail call %struct.float3 @_ZmlRK5char3RK6float3()
  %34 = extractvalue %struct.float3 %33, 0
  %35 = extractvalue %struct.float3 %33, 2
  %36 = fsub contract float %.idx.val.i.i.i, %34
  %37 = fsub contract float %.idx5.val.i.i.i, %35
  %38 = tail call %struct.float3 @_ZL11make_float3fff(float %36, float %37)
  %oldret.i.i.i.i.peel = extractvalue %struct.float3 %38, 0
  %oldret1.i.i.i.i.peel = extractvalue %struct.float3 %38, 1
  %oldret3.i.i.i.i.peel = extractvalue %struct.float3 %38, 2
  store float %oldret.i.i.i.i.peel, float addrspace(5)* %23, align 16
  store float %oldret1.i.i.i.i.peel, float addrspace(5)* %.sroa.215.0..sroa_idx16.i.i.i, align 4
  store float %oldret3.i.i.i.i.peel, float addrspace(5)* %.sroa.3.0..sroa_idx17.i.i.i, align 8
  %39 = zext i8 %32 to i64
  %40 = getelementptr inbounds [27 x %struct.char3], [27 x %struct.char3] addrspace(4)* @_ZN7cuneibs16d_cell_to_offsetE, i64 0, i64 %39
  %.idx11.i.i.i.peel = getelementptr %struct.char3, %struct.char3 addrspace(4)* %40, i64 0, i32 0
  %.idx11.val.i.i.i.peel = load i8, i8 addrspace(4)* %.idx11.i.i.i.peel, align 1
  %.idx12.i.i.i.peel = getelementptr %struct.char3, %struct.char3 addrspace(4)* %40, i64 0, i32 1
  %.idx12.val.i.i.i.peel = load i8, i8 addrspace(4)* %.idx12.i.i.i.peel, align 1
  %.idx13.i.i.i.peel = getelementptr %struct.char3, %struct.char3 addrspace(4)* %40, i64 0, i32 2
  %.idx13.val.i.i.i.peel = load i8, i8 addrspace(4)* %.idx13.i.i.i.peel, align 1
  %41 = sext i8 %.idx11.val.i.i.i.peel to i32
  %42 = add nsw i32 %.idx8.val.i.i.i, %41
  %43 = sext i8 %.idx12.val.i.i.i.peel to i32
  %44 = add nsw i32 %.idx9.val.i.i.i, %43
  %45 = sext i8 %.idx13.val.i.i.i.peel to i32
  %46 = add nsw i32 %.idx10.val.i.i.i, %45
  %47 = tail call %struct.int3 @_ZL9make_int3iii(i32 %42, i32 %44, i32 %46)
  %48 = load i32, i32* %25, align 4
  store i32 %48, i32 addrspace(5)* %20, align 8
  %49 = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 11
  %50 = load i32, i32 addrspace(5)* %49, align 16
  %51 = load %struct.float4*, %struct.float4* addrspace(1)* %5, align 8
  %52 = zext i32 %50 to i64
  %53 = getelementptr inbounds %struct.float4, %struct.float4* %51, i64 %52
  %54 = bitcast %struct.float4* %53 to i8*
  tail call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 1 undef, i8* align 1 %54, i64 undef, i1 false)
  %.idx.i.peel = getelementptr inbounds %"class.cuneibs::neiblist_iterator.1", %"class.cuneibs::neiblist_iterator.1" addrspace(5)* %4, i64 0, i32 2, i32 5, i32 2
  %.idx.val.i.peel = load float, float addrspace(5)* %.idx.i.peel, align 8
  %.idx1.val.i.peel = load float, float* inttoptr (i64 8 to float*), align 8
  %55 = fsub contract float %.idx.val.i.peel, %.idx1.val.i.peel
  %56 = tail call %struct.float4 @_ZL11make_float4ffff(float %55, float undef)
  %57 = tail call i1 @_ZL8isfinitef(float undef)
  %58 = load i64, i64 addrspace(4)* @_ZN7cuneibs17d_neiblist_strideE, align 8, !tbaa !128
  %59 = load i64, i64 addrspace(5)* %26, align 16
  %60 = add i64 %59, %58
  store i64 %60, i64 addrspace(5)* %26, align 16, !tbaa !131
  %61 = add i64 %60, %29
  %62 = getelementptr inbounds i16, i16* %9, i64 %61
  %63 = load i16, i16* %62, align 2
  %64 = icmp eq i16 %63, -1
  br i1 %64, label %.critedge, label %_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit.lr.ph.peel.newph

 _ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit.lr.ph.peel.newph: ; preds = %_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit.lr.ph
  %65 = getelementptr i8, i8* %.pre, i64 -24
  %66 = bitcast i8* %65 to i64*
  br label %_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit

 _ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit: ; preds = %_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit, %_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit.lr.ph.peel.newph
  %67 = phi i16 [ %63, %_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit.lr.ph.peel.newph ], [ %105, %_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit ]
  %68 = lshr i16 %67, 11
  %69 = trunc i16 %68 to i8
  %70 = add nsw i8 %69, -1
  store i8 %70, i8 addrspace(5)* %19, align 4
  %71 = tail call %struct.float3 @_ZmlRK5char3RK6float3()
  %72 = extractvalue %struct.float3 %71, 0
  %73 = extractvalue %struct.float3 %71, 2
  %74 = fsub contract float %.idx.val.i.i.i, %72
  %75 = fsub contract float %.idx5.val.i.i.i, %73
  %76 = tail call %struct.float3 @_ZL11make_float3fff(float %74, float %75)
  %oldret.i.i.i.i = extractvalue %struct.float3 %76, 0
  %oldret1.i.i.i.i = extractvalue %struct.float3 %76, 1
  %oldret3.i.i.i.i = extractvalue %struct.float3 %76, 2
  store float %oldret.i.i.i.i, float addrspace(5)* %23, align 16
  store float %oldret1.i.i.i.i, float addrspace(5)* %.sroa.215.0..sroa_idx16.i.i.i, align 4
  store float %oldret3.i.i.i.i, float addrspace(5)* %.sroa.3.0..sroa_idx17.i.i.i, align 8
  %77 = zext i8 %70 to i64
  %78 = getelementptr inbounds [27 x %struct.char3], [27 x %struct.char3] addrspace(4)* @_ZN7cuneibs16d_cell_to_offsetE, i64 0, i64 %77
  %.idx11.i.i.i = getelementptr %struct.char3, %struct.char3 addrspace(4)* %78, i64 0, i32 0
  %.idx11.val.i.i.i = load i8, i8 addrspace(4)* %.idx11.i.i.i, align 1
  %.idx12.i.i.i = getelementptr %struct.char3, %struct.char3 addrspace(4)* %78, i64 0, i32 1
  %.idx12.val.i.i.i = load i8, i8 addrspace(4)* %.idx12.i.i.i, align 1
  %.idx13.i.i.i = getelementptr %struct.char3, %struct.char3 addrspace(4)* %78, i64 0, i32 2
  %.idx13.val.i.i.i = load i8, i8 addrspace(4)* %.idx13.i.i.i, align 1
  %79 = sext i8 %.idx11.val.i.i.i to i32
  %80 = add nsw i32 %.idx8.val.i.i.i, %79
  %81 = sext i8 %.idx12.val.i.i.i to i32
  %82 = add nsw i32 %.idx9.val.i.i.i, %81
  %83 = sext i8 %.idx13.val.i.i.i to i32
  %84 = add nsw i32 %.idx10.val.i.i.i, %83
  %85 = tail call %struct.int3 @_ZL9make_int3iii(i32 %80, i32 %82, i32 %84)
  %86 = load i32, i32* %25, align 4
  store i32 %86, i32 addrspace(5)* %20, align 8
  %87 = load i64, i64* %66, align 8
  %88 = getelementptr inbounds i8, i8 addrspace(5)* %12, i64 %87
  %89 = getelementptr inbounds i8, i8 addrspace(5)* %88, i64 80
  %90 = bitcast i8 addrspace(5)* %89 to i32 addrspace(5)*
  %91 = load i32, i32 addrspace(5)* %90, align 4
  %92 = load %struct.float4*, %struct.float4* addrspace(1)* %5, align 8
  %93 = zext i32 %91 to i64
  %94 = getelementptr inbounds %struct.float4, %struct.float4* %92, i64 %93
  %95 = bitcast %struct.float4* %94 to i8*
  tail call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 1 undef, i8* align 1 %95, i64 undef, i1 false)
  %.idx.i = getelementptr i8, i8 addrspace(5)* %88, i64 56
  %96 = bitcast i8 addrspace(5)* %.idx.i to float addrspace(5)*
  %.idx.val.i = load float, float addrspace(5)* %96, align 4
  %.idx1.val.i = load float, float* inttoptr (i64 8 to float*), align 8
  %97 = fsub contract float %.idx.val.i, %.idx1.val.i
  %98 = tail call %struct.float4 @_ZL11make_float4ffff(float %97, float undef)
  %99 = tail call i1 @_ZL8isfinitef(float undef)
  %100 = load i64, i64 addrspace(4)* @_ZN7cuneibs17d_neiblist_strideE, align 8, !tbaa !128
  %101 = load i64, i64 addrspace(5)* %26, align 16
  %102 = add i64 %101, %100
  store i64 %102, i64 addrspace(5)* %26, align 16, !tbaa !131
  %103 = add i64 %102, %29
  %104 = getelementptr inbounds i16, i16* %9, i64 %103
  %105 = load i16, i16* %104, align 2
  %106 = icmp eq i16 %105, -1
  br i1 %106, label %.critedge, label %_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit, !llvm.loop !139

 .critedge:                                        ; preds = %_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit.lr.ph, %_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv.exit, %1
  ret void
 }

 attributes #0 = { argmemonly mustprogress nofree nounwind willreturn }

 !nvvm.annotations = !{!0, !1, !2, !3, !4, !5, !6, !7, !8, !9, !10, !11, !12, !13, !14, !15, !16, !17, !18, !19, !20, !21, !22, !23, !24, !25, !26, !27, !28, !29, !30, !31, !32, !33, !34, !35, !36, !37, !38, !39, !40, !41, !42, !43, !44, !45, !46, !47, !48, !49, !50, !51, !52, !53, !54, !55, !56, !57, !57, !58, !58, !59, !59, !60, !60, !61, !62, !63, !64, !65, !66, !67, !68, !69, !70, !71, !72, !73, !74, !75, !76, !77, !78, !79, !80, !81, !79, !80, !81, !82, !83, !84, !84, !85, !85, !86, !87, !88, !89, !90, !91, !92, !92, !93, !93, !94, !94, !95, !95, !71, !96, !97, !98, !99, !100, !101, !102, !103, !84, !84, !85, !85, !104, !105, !106, !107, !104, !108, !109, !110, !111, !112, !113, !114, !111, !115, !116, !117, !118, !119, !120, !121, !118, !122, !123, !124}

 !0 = !{void (i32*, i32*, %struct.ushort4*, i32*, i32)* undef, !"kernel", i32 1}
 !1 = !{void (i32*, i32*, %struct.ushort4*, i32*, i32)* undef, !"maxntidx", i32 256}
 !2 = !{void (i32*, i32*, %struct.ushort4*, i32*, i32)* undef, !"minctasm", i32 6}
 !3 = !{void (%struct.cell_params*, i32)* undef, !"kernel", i32 1}
 !4 = !{void (%struct.float4*, %struct.ushort4*, i32*, %struct.uint4*, i32*, i16*, i32, float, float)* undef, !"kernel", i32 1}
 !5 = !{void (%struct.ushort4*, %struct.float4*, %struct.uint4*, i32)* undef, !"kernel", i32 1}
 !6 = !{void (%struct.ushort4*, %struct.float4*, %struct.float4*, i32, i32, float)* undef, !"kernel", i32 1}
 !7 = !{void (%struct.float4*, %struct.ushort4*, i32)* undef, !"kernel", i32 1}
 !8 = !{void (%struct.pos_info_wrapper*, float*, i32, float)* undef, !"kernel", i32 1}
 !9 = !{void (%struct.pos_info_wrapper*, float*, i32, float)* undef, !"maxntidx", i32 128}
 !10 = !{void (%struct.pos_info_wrapper*, float*, i32, float)* undef, !"minctasm", i32 6}
 !11 = !{void (%struct.jacobi_update_params*)* undef, !"kernel", i32 1}
 !12 = !{void (%struct.jacobi_update_params*)* undef, !"maxntidx", i32 128}
 !13 = !{void (%struct.jacobi_update_params*)* undef, !"minctasm", i32 6}
 !14 = !{void (%struct.ushort4*, %struct.float4*, %struct.float4*, float*, i32)* undef, !"kernel", i32 1}
 !15 = !{void (%struct.float4*, %struct.float4*, %struct.ushort4*, i32*, i32, i32, %struct.float4*)* undef, !"kernel", i32 1}
 !16 = !{void (%struct.float4*, i32, i32)* undef, !"kernel", i32 1}
 !17 = !{void ()* undef, !"kernel", i32 1}
 !18 = !{void (float*, %struct.float4*, i32)* undef, !"kernel", i32 1}
 !19 = !{void (%struct.float4*, %struct.float4*, %struct.float4*, %struct.float4*, %struct.ushort4*, i32*, i16*, i32*, i32, float, float)* undef, !"kernel", i32 1}
 !20 = !{void (%struct.float4*, %struct.float4*, %struct.float4*, %struct.float4*, %struct.ushort4*, i32*, i16*, i32*, i32, float, float)* undef, !"maxntidx", i32 128}
 !21 = !{void (%struct.float4*, %struct.float4*, %struct.float4*, %struct.float4*, %struct.ushort4*, i32*, i16*, i32*, i32, float, float)* undef, !"minctasm", i32 6}
 !22 = !{void (%struct.sa_outgoing_bc_params*)* undef, !"kernel", i32 1}
 !23 = !{void (%struct.float4*, %struct.uint4*, %struct.ushort4*, i32*, i32*, i16*, i32)* undef, !"kernel", i32 1}
 !24 = !{void (%struct.uint4*, i32*, %struct.ushort4*, i32*, i16*, %struct.float4*, i32)* undef, !"kernel", i32 1}
 !25 = !{void (%struct.uint4*, i32*, %struct.ushort4*, i32*, i16*, %struct.float4*, i32)* undef, !"maxntidx", i32 128}
 !26 = !{void (%struct.uint4*, i32*, %struct.ushort4*, i32*, i16*, %struct.float4*, i32)* undef, !"minctasm", i32 6}
 !27 = !{void (%struct.float4*, %struct.float4*, %struct.uint4*, i32*, %struct.ushort4*, i32*, i16*, %struct.float4*, i32, float)* undef, !"kernel", i32 1}
 !28 = !{void (%struct.float4*, %struct.float4*, %struct.uint4*, i32*, %struct.ushort4*, i32*, i16*, %struct.float4*, i32, float)* undef, !"maxntidx", i32 128}
 !29 = !{void (%struct.float4*, %struct.float4*, %struct.uint4*, i32*, %struct.ushort4*, i32*, i16*, %struct.float4*, i32, float)* undef, !"minctasm", i32 6}
 !30 = !{void (%struct.float4*, i32*, i32*, %struct.ushort4*, i32*, i32)* undef, !"kernel", i32 1}
 !31 = !{void (%struct.float4*, i32*, i32*, %struct.ushort4*, i32*, i32)* undef, !"maxntidx", i32 256}
 !32 = !{void (%struct.float4*, i32*, i32*, %struct.ushort4*, i32*, i32)* undef, !"minctasm", i32 6}
 !33 = !{void (%struct.reorder_params*, i32*, i32*, i32*, %struct.ushort4*, i32*, i32*, i32, i32*)* undef, !"kernel", i32 1}
 !34 = !{void (%struct.reorder_params*, i32*, i32*, i32*, %struct.ushort4*, i32*, i32*, i32, i32*)* undef, !"maxntidx", i32 256}
 !35 = !{void (%struct.reorder_params*, i32*, i32*, i32*, %struct.ushort4*, i32*, i32*, i32, i32*)* undef, !"minctasm", i32 6}
 !36 = !{void (i1, %"class.thrust::zip_iterator"*, %"class.thrust::device_ptr"*, i64, %"class.thrust::tuple"*, i32*, %struct.ptype_hash_compare*)* undef, !"kernel", i32 1}
 !37 = !{void (i1, %"class.thrust::zip_iterator"*, %"class.thrust::device_ptr"*, i64, %"class.thrust::tuple"*, i32*, %struct.ptype_hash_compare*)* undef, !"maxntidx", i32 256}
 !38 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f"*, i64)* undef, !"kernel", i32 1}
 !39 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f"*, i64)* undef, !"maxntidx", i32 256}
 !40 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.69"*, i64)* undef, !"kernel", i32 1}
 !41 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.69"*, i64)* undef, !"maxntidx", i32 256}
 !42 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.72"*, i64)* undef, !"kernel", i32 1}
 !43 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.72"*, i64)* undef, !"maxntidx", i32 256}
 !44 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.74"*, i64)* undef, !"kernel", i32 1}
 !45 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.74"*, i64)* undef, !"maxntidx", i32 256}
 !46 = !{void (i1, %"class.thrust::zip_iterator"*, %"class.thrust::tuple"*, i64, i64, i64*, %struct.ptype_hash_compare*, i64, i32)* undef, !"kernel", i32 1}
 !47 = !{void (i1, %"class.thrust::zip_iterator"*, %"class.thrust::tuple"*, i64, i64, i64*, %struct.ptype_hash_compare*, i64, i32)* undef, !"maxntidx", i32 256}
 !48 = !{void (i1, %"class.thrust::zip_iterator"*, %"class.thrust::device_ptr"*, i64, %"class.thrust::tuple"*, i32*, %struct.ptype_hash_compare*, i64*, i64)* undef, !"kernel", i32 1}
 !49 = !{void (i1, %"class.thrust::zip_iterator"*, %"class.thrust::device_ptr"*, i64, %"class.thrust::tuple"*, i32*, %struct.ptype_hash_compare*, i64*, i64)* undef, !"maxntidx", i32 256}
 !50 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.79"*, i64)* undef, !"kernel", i32 1}
 !51 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.79"*, i64)* undef, !"maxntidx", i32 256}
 !52 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.81"*, i64)* undef, !"kernel", i32 1}
 !53 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.81"*, i64)* undef, !"maxntidx", i32 256}
 !54 = !{void (%struct.buildneibs_params*)* undef, !"kernel", i32 1}
 !55 = !{void (%struct.buildneibs_params*)* undef, !"maxntidx", i32 256}
 !56 = !{void (%struct.buildneibs_params*)* undef, !"minctasm", i32 5}
 !57 = !{void (%struct.euler_params*)* undef, !"kernel", i32 1}
 !58 = !{void (%struct.euler_params.111*)* undef, !"kernel", i32 1}
 !59 = !{void (%struct.euler_params.114*)* undef, !"kernel", i32 1}
 !60 = !{void (%struct.euler_params.117*)* undef, !"kernel", i32 1}
 !61 = !{void (%"struct.cub::ReduceByKeyScanTileState"*, i64)* undef, !"kernel", i32 1}
 !62 = !{void (%"struct.cub::ReduceByKeyScanTileState"*, i64)* undef, !"maxntidx", i32 128}
 !63 = !{void (%"class.thrust::device_ptr.120"*, %"class.thrust::device_ptr.124"*, %"class.thrust::device_ptr.124"*, %"struct.thrust::equal_to"*, %"struct.thrust::plus"*, %"struct.cub::ReduceByKeyScanTileState"*, i32, %"struct.thrust::cuda_cub::__scan_by_key::DoNothing"*)* undef, !"kernel", i32 1}
 !64 = !{void (%"class.thrust::device_ptr.120"*, %"class.thrust::device_ptr.124"*, %"class.thrust::device_ptr.124"*, %"struct.thrust::equal_to"*, %"struct.thrust::plus"*, %"struct.cub::ReduceByKeyScanTileState"*, i32, %"struct.thrust::cuda_cub::__scan_by_key::DoNothing"*)* undef, !"maxntidx", i32 256}
 !65 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.166"*, i64)* undef, !"kernel", i32 1}
 !66 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.166"*, i64)* undef, !"maxntidx", i32 256}
 !67 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.169"*, i64)* undef, !"kernel", i32 1}
 !68 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.169"*, i64)* undef, !"maxntidx", i32 256}
 !69 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.171"*, i64)* undef, !"kernel", i32 1}
 !70 = !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.171"*, i64)* undef, !"maxntidx", i32 256}
 !71 = !{void (%struct.density_diffusion_params*)* undef, !"kernel", i32 1}
 !72 = !{void (%struct.forces_params*)* undef, !"kernel", i32 1}
 !73 = !{void (%struct.forces_params.223*)* undef, !"kernel", i32 1}
 !74 = !{void (%struct.finalize_forces_params*)* undef, !"kernel", i32 1}
 !75 = !{void (%struct.forces_params.236*)* undef, !"kernel", i32 1}
 !76 = !{void (%struct.forces_params.250*)* undef, !"kernel", i32 1}
 !77 = !{void (%struct.forces_params.256*)* undef, !"kernel", i32 1}
 !78 = !{void (%struct.finalize_forces_params.262*)* undef, !"kernel", i32 1}
 !79 = !{void (%struct.neibs_interaction_params*, %struct.float4*)* undef, !"kernel", i32 1}
 !80 = !{void (%struct.neibs_interaction_params*, %struct.float4*)* undef, !"maxntidx", i32 128}
 !81 = !{void (%struct.neibs_interaction_params*, %struct.float4*)* undef, !"minctasm", i32 6}
 !82 = !{void (%struct.neibs_interaction_params*, %struct.float3*)* undef, !"kernel", i32 1}
 !83 = !{void (%"struct.cupostprocess::testpoints_params"*)* undef, !"kernel", i32 1}
 !84 = !{void (%struct.neibs_interaction_params*, %struct.float4*, %struct.ushort4*)* undef, !"kernel", i32 1}
 !85 = !{void (%struct.neibs_interaction_params*, %struct.float4*, %struct.ushort4*, float)* undef, !"kernel", i32 1}
 !86 = !{void (%struct.reorder_params.278*, i32*, i32*, i32*, %struct.ushort4*, i32*, i32*, i32, i32*)* undef, !"kernel", i32 1}
 !87 = !{void (%struct.reorder_params.278*, i32*, i32*, i32*, %struct.ushort4*, i32*, i32*, i32, i32*)* undef, !"maxntidx", i32 256}
 !88 = !{void (%struct.reorder_params.278*, i32*, i32*, i32*, %struct.ushort4*, i32*, i32*, i32, i32*)* undef, !"minctasm", i32 6}
 !89 = !{void (%struct.buildneibs_params.279*)* undef, !"kernel", i32 1}
 !90 = !{void (%struct.buildneibs_params.279*)* undef, !"maxntidx", i32 256}
 !91 = !{void (%struct.buildneibs_params.279*)* undef, !"minctasm", i32 5}
 !92 = !{void (%struct.euler_params.282*)* undef, !"kernel", i32 1}
 !93 = !{void (%struct.euler_params.285*)* undef, !"kernel", i32 1}
 !94 = !{void (%struct.euler_params.288*)* undef, !"kernel", i32 1}
 !95 = !{void (%struct.euler_params.291*)* undef, !"kernel", i32 1}
 !96 = !{void (%struct.forces_params.294*)* undef, !"kernel", i32 1}
 !97 = !{void (%struct.forces_params.300*)* undef, !"kernel", i32 1}
 !98 = !{void (%struct.finalize_forces_params.305*)* undef, !"kernel", i32 1}
 !99 = !{void (%struct.forces_params.312*)* undef, !"kernel", i32 1}
 !100 = !{void (%struct.forces_params.319*)* undef, !"kernel", i32 1}
 !101 = !{void (%struct.forces_params.325*)* undef, !"kernel", i32 1}
 !102 = !{void (%struct.finalize_forces_params.331*)* undef, !"kernel", i32 1}
 !103 = !{void (%"struct.cupostprocess::testpoints_params.336"*)* undef, !"kernel", i32 1}
 !104 = !{void (%struct.density_diffusion_params.339*)* undef, !"kernel", i32 1}
 !105 = !{void (%struct.forces_params.341*)* undef, !"kernel", i32 1}
 !106 = !{void (%struct.forces_params.346*)* undef, !"kernel", i32 1}
 !107 = !{void (%struct.forces_params.351*)* undef, !"kernel", i32 1}
 !108 = !{void (%struct.forces_params.356*)* undef, !"kernel", i32 1}
 !109 = !{void (%struct.forces_params.361*)* undef, !"kernel", i32 1}
 !110 = !{void (%struct.forces_params.366*)* undef, !"kernel", i32 1}
 !111 = !{void (%struct.density_diffusion_params.371*)* undef, !"kernel", i32 1}
 !112 = !{void (%struct.forces_params.373*)* undef, !"kernel", i32 1}
 !113 = !{void (%struct.forces_params.378*)* undef, !"kernel", i32 1}
 !114 = !{void (%struct.forces_params.383*)* undef, !"kernel", i32 1}
 !115 = !{void (%struct.forces_params.388*)* undef, !"kernel", i32 1}
 !116 = !{void (%struct.forces_params.393*)* undef, !"kernel", i32 1}
 !117 = !{void (%struct.forces_params.398*)* undef, !"kernel", i32 1}
 !118 = !{void (%struct.density_diffusion_params.403*)* undef, !"kernel", i32 1}
 !119 = !{void (%struct.forces_params.405*)* undef, !"kernel", i32 1}
 !120 = !{void (%struct.forces_params.410*)* undef, !"kernel", i32 1}
 !121 = !{void (%struct.forces_params.415*)* @_ZN8cuforces12forcesDeviceI13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSD_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyI18xsph_forces_paramsESF_I20volume_forces_paramsESF_I21grenier_forces_paramsESF_I25sa_boundary_forces_paramsESF_I28dummy_boundary_forces_paramsESF_I25water_depth_forces_paramsESF_I18keps_forces_paramsESF_I14tau_tex_paramsESF_I22eulerVel_forces_paramsESF_I29internal_energy_forces_paramsESF_I28effective_visc_forces_paramsEELS2_3ELS3_1ELS4_3ELS5_4ESC_Lm517ELSD_1ELSD_0EEEvT_, !"kernel", i32 1}
 !122 = !{void (%struct.forces_params.420*)* undef, !"kernel", i32 1}
 !123 = !{void (%struct.forces_params.425*)* undef, !"kernel", i32 1}
 !124 = !{void (%struct.forces_params.430*)* undef, !"kernel", i32 1}
 !125 = !{!126, !126, i64 0}
 !126 = !{!"vtable pointer", !127, i64 0}
 !127 = !{!"Simple C++ TBAA"}
 !128 = !{!129, !129, i64 0}
 !129 = !{!"long", !130, i64 0}
 !130 = !{!"omnipotent char", !127, i64 0}
 !131 = !{!132, !129, i64 64}
 !132 = !{!"_ZTSN7cuneibs22neiblist_iterator_coreE", !133, i64 0, !133, i64 8, !134, i64 16, !136, i64 32, !137, i64 44, !138, i64 48, !129, i64 64, !137, i64 72, !130, i64 76, !137, i64 80}
 !133 = !{!"any pointer", !130, i64 0}
 !134 = !{!"_ZTS6float4", !135, i64 0, !135, i64 4, !135, i64 8, !135, i64 12}
 !135 = !{!"float", !130, i64 0}
 !136 = !{!"_ZTS4int3", !137, i64 0, !137, i64 4, !137, i64 8}
 !137 = !{!"int", !130, i64 0}
 !138 = !{!"_ZTS6float3", !135, i64 0, !135, i64 4, !135, i64 8}
 !139 = distinct !{!139, !140}
 !140 = !{!"llvm.loop.peeled.count", i32 1}