Skip to content

Instantly share code, notes, and snippets.

@Artem-B

Artem-B/SROA-reproducer.ll

Created October 11, 2021 20:08
Show Gist options
  • Star 0 You must be signed in to star a gist
  • Fork 0 You must be signed in to fork a gist
  • Save Artem-B/c8d048ce7666f5693a8c899458829f5a to your computer and use it in GitHub Desktop.
Save Artem-B/c8d048ce7666f5693a8c899458829f5a to your computer and use it in GitHub Desktop.
Download ZIP
Raw
SROA-reproducer.ll
; Compile with:
;
; clang "-cc1" "-triple" "nvptx64-nvidia-cuda" "-aux-triple" "x86_64-pc-linux-gnu"
; "-S" "-disable-llvm-verifier" "-discard-value-names" "-main-file-name"
; "DamBreak3D.cu" "-mrelocation-model" "static" "-mframe-pointer=all"
; "-fno-rounding-math" "-fno-verbose-asm" "-no-integrated-as" "-aux-target-cpu"
; "x86-64" "-fcuda-is-device"
; "-target-feature" "+ptx70" "-target-sdk-version=11.0" "-target-cpu" "sm_70"
; "-O3" "-x" "ir"
%struct.uint3 = type { i32, i32, i32 }
%struct.char3 = type { i8, i8, i8 }
%struct.int3 = type { i32, i32, i32 }
%"class.cuneibs::neiblist_iterator_core" = type <{ i32*, i16*, %struct.float4, %struct.int3, i32, %struct.float3, [4 x i8], i64, i32, i8, [3 x i8], i32, [12 x i8] }>
%struct.float4 = type { float, float, float, float }
%struct.float3 = type { float, float, float }
%struct.pos_wrapper = type { %struct.float4* }
%struct.ushort4 = type { i16, i16, i16, i16 }
%struct.info_wrapper = type { %struct.ushort4* }
%"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 }>
%"class.cuneibs::neiblist_iterator_simple.2" = type { i32 (...)**, [8 x i8], %"class.cuneibs::neiblist_iterator_core.base", [12 x i8] }
%struct.vel_wrapper = type { %struct.float4* }
%"struct.cuforces::common_particle_data" = type <{ i32, [4 x i8], %struct.ushort4*, %struct.float4*, %struct.int3, [4 x i8] }>
%"struct.cuforces::rb_particle_data" = type { 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.cuforces::forces_particle_data.416" = type { %"struct.cuforces::common_particle_data.base", %"struct.cuforces::rb_particle_data", [8 x i8], %"struct.cuforces::vel_particle_data.base", %"struct.cuforces::p_precalc_particle_data", [8 x i8] }
%"struct.cuforces::common_particle_data.base" = type <{ i32, [4 x i8], %struct.ushort4*, %struct.float4*, %struct.int3 }>
%"struct.cuforces::vel_particle_data.base" = type { %struct.float4, float }
%"struct.cuforces::p_precalc_particle_data" = type { float }
%"struct.cuforces::vel_particle_data" = type { %struct.float4, float, [12 x i8] }
%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 }
$_ZNK7cuneibs22neiblist_iterator_core10neib_indexEv = comdat any
$_ZNK11pos_wrapper8fetchPosEj = comdat any
$_ZN7cuneibs22neiblist_iterator_core17update_neib_indexEt = comdat any
$_ZN7cuneibs12getNeibIndexERK6float4R6float3PKjtRK4int3RhRj = comdat any
$_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv = comdat any
$_ZNK7cuneibs22neiblist_iterator_core6relPosERK6float4 = comdat any
$_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE5resetEv = comdat any
$_ZN7cuneibs15neib_list_startIL12ParticleType0EEEmv = comdat any
$_ZN7cuneibs14neib_list_stepIL12ParticleType0EEEmv = comdat any
$_ZN7cuneibs14first_neib_locIL12ParticleType0EEEjv = comdat any
$_ZTVN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEE = comdat any
$_ZTTN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEE = comdat any
$_ZTCN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEE0_NS_24neiblist_iterator_simpleILS1_0EEE = comdat any
@_ZN7cuneibs10d_gridSizeE = external addrspace(4) externally_initialized global %struct.uint3, align 4
@_ZN7cuneibs16d_cell_to_offsetE = external addrspace(4) externally_initialized global [27 x %struct.char3], align 1
@_ZN7cuneibs17d_neiblist_strideE = external addrspace(4) externally_initialized global i64, align 8
@_ZTVN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEE = linkonce_odr unnamed_addr constant { [3 x i8*] } { [3 x i8*] [i8* inttoptr (i64 16 to i8*), i8* null, i8* null] }, comdat, align 8
@_ZTTN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEE = linkonce_odr unnamed_addr constant [2 x i8*] [i8* bitcast (i8** getelementptr inbounds ({ [3 x i8*] }, { [3 x i8*] }* @_ZTVN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEE, i32 0, inrange i32 0, i32 3) to i8*), i8* bitcast (i8** getelementptr inbounds ({ [3 x i8*] }, { [3 x i8*] }* @_ZTCN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEE0_NS_24neiblist_iterator_simpleILS1_0EEE, i32 0, inrange i32 0, i32 3) to i8*)], comdat, align 8
@_ZTCN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEE0_NS_24neiblist_iterator_simpleILS1_0EEE = linkonce_odr unnamed_addr constant { [3 x i8*] } { [3 x i8*] [i8* inttoptr (i64 16 to i8*), i8* null, i8* null] }, comdat, align 8
; Function Attrs: argmemonly nofree nosync nounwind willreturn
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #0
; Function Attrs: argmemonly nofree nosync nounwind willreturn
declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #0
declare i32 @_ZN25__cuda_builtin_blockIdx_t17__fetch_builtin_xEv() align 2
declare i32 @_ZN25__cuda_builtin_blockDim_t17__fetch_builtin_xEv() align 2
declare i32 @_ZN26__cuda_builtin_threadIdx_t17__fetch_builtin_xEv() align 2
; Function Attrs: argmemonly nofree nounwind willreturn
declare void @llvm.memcpy.p0i8.p0i8.i64(i8* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg) #1
define internal i16 @_ZL6objectRK7ushort4() {
ret i16 undef
}
define internal i16* @_ZL4typeRK7ushort4() {
ret i16* undef
}
declare %struct.int3 @_ZN7cuneibs27calcGridPosFromParticleHashEj()
define linkonce_odr i32* @_ZNK7cuneibs22neiblist_iterator_core10neib_indexEv(%"class.cuneibs::neiblist_iterator_core"* %0) comdat align 2 {
%2 = alloca %"class.cuneibs::neiblist_iterator_core"*, i32 undef, align 8
store %"class.cuneibs::neiblist_iterator_core"* %0, %"class.cuneibs::neiblist_iterator_core"** %2, align 8
%3 = load %"class.cuneibs::neiblist_iterator_core"*, %"class.cuneibs::neiblist_iterator_core"** %2, align 8
%4 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %3, i32 undef, i32 11
ret i32* %4
}
define internal i32 @_ZL2idRK7ushort4() {
ret i32 undef
}
declare i1 @_ZL8isfinitef(float)
define linkonce_odr %struct.float4 @_ZNK11pos_wrapper8fetchPosEj(%struct.pos_wrapper* %0, i32 %1) comdat align 2 {
%3 = alloca %struct.float4, i32 undef, align 16
%4 = alloca %struct.pos_wrapper*, i32 undef, align 8
%5 = alloca i32, i32 undef, align 4
store %struct.pos_wrapper* %0, %struct.pos_wrapper** %4, align 8
store i32 %1, i32* %5, align 4
%6 = load %struct.pos_wrapper*, %struct.pos_wrapper** %4, align 8
%7 = getelementptr inbounds %struct.pos_wrapper, %struct.pos_wrapper* %6, i32 undef, i32 0
%8 = load %struct.float4*, %struct.float4** %7, align 8
%9 = load i32, i32* %5, align 4
%10 = zext i32 %9 to i64
%11 = getelementptr inbounds %struct.float4, %struct.float4* %8, i64 %10
%12 = bitcast %struct.float4* undef to i8*
%13 = bitcast %struct.float4* %11 to i8*
call void @llvm.memcpy.p0i8.p0i8.i64(i8* undef, i8* %13, i64 undef, i1 false)
%14 = load %struct.float4, %struct.float4* undef, align 16
ret %struct.float4 undef
}
define linkonce_odr %struct.ushort4 @_ZNK12info_wrapper9fetchInfoEj(%struct.info_wrapper* %0) align 2 {
ret %struct.ushort4 undef
}
define internal i16 @_ZL9fluid_numRK7ushort4() {
ret i16 undef
}
declare float @_ZN6cuphys10soundSpeedEft()
declare %struct.float4 @_ZL11make_float4ffff(float, float)
declare %struct.float3 @_ZL11make_float3fff(float, float)
define linkonce_odr void @_ZN7cuneibs22neiblist_iterator_coreC2EjRK6float4RK4int3PKjPKt(%"class.cuneibs::neiblist_iterator_core"* %0, i16* %1) unnamed_addr align 2 {
%3 = alloca %"class.cuneibs::neiblist_iterator_core"*, i32 undef, align 8
%4 = alloca i16*, i32 undef, align 8
store %"class.cuneibs::neiblist_iterator_core"* %0, %"class.cuneibs::neiblist_iterator_core"** %3, align 8
store i16* %1, i16** %4, align 8
%5 = load %"class.cuneibs::neiblist_iterator_core"*, %"class.cuneibs::neiblist_iterator_core"** %3, align 8
%6 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %5, i32 undef, i32 1
%7 = load i16*, i16** %4, align 8
store i16* %7, i16** %6, align 8
ret void
}
define internal %struct.float3 @_ZL11make_float3RKf() {
ret %struct.float3 undef
}
define linkonce_odr void @_ZN7cuneibs22neiblist_iterator_core17update_neib_indexEt(%"class.cuneibs::neiblist_iterator_core"* %0, i16 %1) comdat align 2 {
%3 = alloca %"class.cuneibs::neiblist_iterator_core"*, i32 undef, align 8
%4 = alloca i16, i32 undef, align 2
store %"class.cuneibs::neiblist_iterator_core"* %0, %"class.cuneibs::neiblist_iterator_core"** %3, align 8
store i16 %1, i16* %4, align 2
%5 = load %"class.cuneibs::neiblist_iterator_core"*, %"class.cuneibs::neiblist_iterator_core"** %3, align 8
%6 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %5, i32 undef, i32 2
%7 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %5, i32 undef, i32 5
%8 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %5, i32 undef, i32 0
%9 = load i32*, i32** %8, align 16
%10 = load i16, i16* %4, align 2
%11 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %5, i32 undef, i32 3
%12 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %5, i32 undef, i32 9
%13 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %5, i32 undef, i32 8
%14 = call i32 @_ZN7cuneibs12getNeibIndexERK6float4R6float3PKjtRK4int3RhRj(%struct.float4* %6, %struct.float3* %7, i32* %9, i16 %10, %struct.int3* %11, i8* %12, i32* %13)
ret void
}
define linkonce_odr i32 @_ZN7cuneibs12getNeibIndexERK6float4R6float3PKjtRK4int3RhRj(%struct.float4* %0, %struct.float3* %1, i32* %2, i16 %3, %struct.int3* %4, i8* %5, i32* %6) comdat {
%8 = alloca %struct.float4*, i32 undef, align 8
%9 = alloca %struct.float3*, i32 undef, align 8
%10 = alloca i32*, i32 undef, align 8
%11 = alloca i16, i32 undef, align 2
%12 = alloca %struct.int3*, i32 undef, align 8
%13 = alloca i8*, i32 undef, align 8
%14 = alloca i32*, i32 undef, align 8
%15 = alloca %struct.float3, align 4
%16 = alloca %struct.float3, align 4
store %struct.float4* %0, %struct.float4** %8, align 8
store %struct.float3* %1, %struct.float3** %9, align 8
store i32* %2, i32** %10, align 8
store i16 %3, i16* %11, align 2
store %struct.int3* %4, %struct.int3** %12, align 8
store i8* %5, i8** %13, align 8
store i32* %6, i32** %14, align 8
br label %17
17: ; preds = %7
%18 = load i16, i16* %11, align 2
%19 = zext i16 %18 to i32
%20 = ashr i32 %19, 11
%21 = sub nsw i32 %20, 1
%22 = trunc i32 %21 to i8
%23 = load i8*, i8** %13, align 8
store i8 %22, i8* %23, align 1
%24 = load %struct.float4*, %struct.float4** %8, align 8
%25 = call %struct.float3* @_ZL9as_float3RK6float4(%struct.float4* %24)
%26 = call %struct.float3 @_ZmlRK5char3RK6float3()
%27 = getelementptr inbounds %struct.float3, %struct.float3* %16, i32 undef, i32 0
%28 = extractvalue %struct.float3 %26, 0
store float %28, float* %27, align 4
%29 = getelementptr inbounds %struct.float3, %struct.float3* %16, i32 undef, i32 2
%30 = extractvalue %struct.float3 %26, 2
store float %30, float* %29, align 4
%31 = call %struct.float3 @_ZmiRK6float3S1_(%struct.float3* %25, %struct.float3* %16)
%32 = getelementptr inbounds %struct.float3, %struct.float3* %15, i32 undef, i32 0
%33 = extractvalue %struct.float3 %31, 0
store float %33, float* %32, align 4
%34 = getelementptr inbounds %struct.float3, %struct.float3* %15, i32 undef, i32 1
%35 = extractvalue %struct.float3 %31, 1
store float %35, float* %34, align 4
%36 = getelementptr inbounds %struct.float3, %struct.float3* %15, i32 undef, i32 2
%37 = extractvalue %struct.float3 %31, 2
store float %37, float* %36, align 4
%38 = load %struct.float3*, %struct.float3** %9, align 8
%39 = bitcast %struct.float3* %38 to i8*
%40 = bitcast %struct.float3* %15 to i8*
call void @llvm.memcpy.p0i8.p0i8.i64(i8* %39, i8* %40, i64 12, i1 false)
%41 = load i32*, i32** %10, align 8
%42 = load %struct.int3*, %struct.int3** %12, align 8
%43 = load i8*, i8** %13, align 8
%44 = load i8, i8* %43, align 1
%45 = zext i8 %44 to i64
%46 = getelementptr inbounds [27 x %struct.char3], [27 x %struct.char3]* addrspacecast ([27 x %struct.char3] addrspace(4)* @_ZN7cuneibs16d_cell_to_offsetE to [27 x %struct.char3]*), i64 0, i64 %45
%47 = call %struct.int3 @_ZplRK4int3RK5char3(%struct.int3* %42, %struct.char3* %46)
%48 = getelementptr inbounds i32, i32* %41, i64 undef
%49 = load i32, i32* %48, align 4
%50 = load i32*, i32** %14, align 8
store i32 %49, i32* %50, align 4
ret i32 undef
}
define internal %struct.float3 @_ZmiRK6float3S1_(%struct.float3* %0, %struct.float3* %1) {
%3 = alloca %struct.float3, align 4
%4 = alloca %struct.float3*, i32 undef, align 8
%5 = alloca %struct.float3*, i32 undef, align 8
store %struct.float3* %0, %struct.float3** %4, align 8
store %struct.float3* %1, %struct.float3** %5, align 8
%6 = load %struct.float3*, %struct.float3** %4, align 8
%7 = getelementptr inbounds %struct.float3, %struct.float3* %6, i32 undef, i32 0
%8 = load float, float* %7, align 4
%9 = load %struct.float3*, %struct.float3** %5, align 8
%10 = getelementptr inbounds %struct.float3, %struct.float3* %9, i32 undef, i32 0
%11 = load float, float* %10, align 4
%12 = fsub contract float %8, %11
%13 = load %struct.float3*, %struct.float3** %4, align 8
%14 = getelementptr inbounds %struct.float3, %struct.float3* %13, i32 undef, i32 2
%15 = load float, float* %14, align 4
%16 = load %struct.float3*, %struct.float3** %5, align 8
%17 = getelementptr inbounds %struct.float3, %struct.float3* %16, i32 undef, i32 2
%18 = load float, float* %17, align 4
%19 = fsub contract float %15, %18
%20 = call %struct.float3 @_ZL11make_float3fff(float %12, float %19)
%21 = getelementptr inbounds %struct.float3, %struct.float3* %3, i32 0, i32 2
%22 = extractvalue %struct.float3 %20, 2
store float %22, float* %21, align 4
%23 = load %struct.float3, %struct.float3* %3, align 4
ret %struct.float3 %23
}
define internal %struct.float3* @_ZL9as_float3RK6float4(%struct.float4* %0) {
%2 = alloca %struct.float4*, i32 undef, align 8
store %struct.float4* %0, %struct.float4** %2, align 8
%3 = load %struct.float4*, %struct.float4** %2, align 8
%4 = bitcast %struct.float4* %3 to %struct.float3*
ret %struct.float3* %4
}
declare %struct.float3 @_ZmlRK5char3RK6float3()
define linkonce_odr i32 @_ZN7cuneibs20calcGridHashPeriodicE4int3() {
ret i32 undef
}
define internal %struct.int3 @_ZplRK4int3RK5char3(%struct.int3* %0, %struct.char3* %1) {
%3 = alloca %struct.int3*, i32 undef, align 8
%4 = alloca %struct.char3*, i32 undef, align 8
store %struct.int3* %0, %struct.int3** %3, align 8
store %struct.char3* %1, %struct.char3** %4, align 8
%5 = load %struct.int3*, %struct.int3** %3, align 8
%6 = getelementptr inbounds %struct.int3, %struct.int3* %5, i32 undef, i32 0
%7 = load i32, i32* %6, align 4
%8 = load %struct.char3*, %struct.char3** %4, align 8
%9 = getelementptr inbounds %struct.char3, %struct.char3* %8, i32 undef, i32 0
%10 = load i8, i8* %9, align 1
%11 = sext i8 %10 to i32
%12 = add nsw i32 %7, %11
%13 = load %struct.int3*, %struct.int3** %3, align 8
%14 = getelementptr inbounds %struct.int3, %struct.int3* %13, i32 undef, i32 1
%15 = load i32, i32* %14, align 4
%16 = load %struct.char3*, %struct.char3** %4, align 8
%17 = getelementptr inbounds %struct.char3, %struct.char3* %16, i32 undef, i32 1
%18 = load i8, i8* %17, align 1
%19 = sext i8 %18 to i32
%20 = add nsw i32 %15, %19
%21 = load %struct.int3*, %struct.int3** %3, align 8
%22 = getelementptr inbounds %struct.int3, %struct.int3* %21, i32 undef, i32 2
%23 = load i32, i32* %22, align 4
%24 = load %struct.char3*, %struct.char3** %4, align 8
%25 = getelementptr inbounds %struct.char3, %struct.char3* %24, i32 undef, i32 2
%26 = load i8, i8* %25, align 1
%27 = sext i8 %26 to i32
%28 = add nsw i32 %23, %27
%29 = call %struct.int3 @_ZL9make_int3iii(i32 %12, i32 %20, i32 %28)
ret %struct.int3 undef
}
define linkonce_odr i32 @_ZN7cuneibs12calcGridHashERK4int3() {
ret i32 undef
}
declare %struct.int3 @_ZL9make_int3iii(i32, i32, i32)
define linkonce_odr void @_ZN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEC1EjRK6float4RK4int3PKjPKt(%"class.cuneibs::neiblist_iterator.1"* %0, i16* %1) unnamed_addr align 2 {
%3 = alloca %"class.cuneibs::neiblist_iterator.1"*, i32 undef, align 8
%4 = alloca i16*, i32 undef, align 8
store %"class.cuneibs::neiblist_iterator.1"* %0, %"class.cuneibs::neiblist_iterator.1"** %3, align 8
store i16* %1, i16** %4, align 8
%5 = load %"class.cuneibs::neiblist_iterator.1"*, %"class.cuneibs::neiblist_iterator.1"** %3, align 8
%6 = bitcast %"class.cuneibs::neiblist_iterator.1"* %5 to i8*
%7 = getelementptr inbounds i8, i8* %6, i64 16
%8 = bitcast i8* %7 to %"class.cuneibs::neiblist_iterator_core"*
%9 = load i16*, i16** %4, align 8
call void @_ZN7cuneibs22neiblist_iterator_coreC2EjRK6float4RK4int3PKjPKt(%"class.cuneibs::neiblist_iterator_core"* %8, i16* %9)
%10 = bitcast %"class.cuneibs::neiblist_iterator.1"* %5 to i32 (...)***
store i32 (...)** bitcast (i8** getelementptr inbounds ({ [3 x i8*] }, { [3 x i8*] }* @_ZTVN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEE, i32 0, inrange i32 0, i32 3) to i32 (...)**), i32 (...)*** %10, align 8, !tbaa !146
ret void
}
define linkonce_odr i1 @_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv(%"class.cuneibs::neiblist_iterator_simple.2"* %0) comdat align 2 {
%2 = alloca i1, i32 undef, align 1
%3 = alloca %"class.cuneibs::neiblist_iterator_simple.2"*, i32 undef, align 8
%4 = alloca i16, align 2
store %"class.cuneibs::neiblist_iterator_simple.2"* %0, %"class.cuneibs::neiblist_iterator_simple.2"** %3, align 8
%5 = load %"class.cuneibs::neiblist_iterator_simple.2"*, %"class.cuneibs::neiblist_iterator_simple.2"** %3, align 8
%6 = call i64 @_ZN7cuneibs14neib_list_stepIL12ParticleType0EEEmv()
%7 = bitcast %"class.cuneibs::neiblist_iterator_simple.2"* %5 to i8**
%8 = load i8*, i8** %7, align 8, !tbaa !146
%9 = getelementptr i8, i8* %8, i64 -24
%10 = bitcast i8* %9 to i64*
%11 = load i64, i64* %10, align 8
%12 = bitcast %"class.cuneibs::neiblist_iterator_simple.2"* %5 to i8*
%13 = getelementptr inbounds i8, i8* %12, i64 %11
%14 = bitcast i8* %13 to %"class.cuneibs::neiblist_iterator_core"*
%15 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %14, i32 undef, i32 7
%16 = load i64, i64* %15, align 16
%17 = add i64 %16, %6
store i64 %17, i64* %15, align 16, !tbaa !149
%18 = bitcast %"class.cuneibs::neiblist_iterator_simple.2"* %5 to i8**
%19 = load i8*, i8** %18, align 8, !tbaa !146
%20 = getelementptr i8, i8* %19, i64 -24
%21 = bitcast i8* %20 to i64*
%22 = load i64, i64* %21, align 8
%23 = bitcast %"class.cuneibs::neiblist_iterator_simple.2"* %5 to i8*
%24 = getelementptr inbounds i8, i8* %23, i64 %22
%25 = bitcast i8* %24 to %"class.cuneibs::neiblist_iterator_core"*
%26 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %25, i32 undef, i32 1
%27 = load i16*, i16** %26, align 8
%28 = bitcast %"class.cuneibs::neiblist_iterator_simple.2"* %5 to i8**
%29 = load i8*, i8** %28, align 8
%30 = getelementptr i8, i8* %29, i64 -24
%31 = bitcast i8* %30 to i64*
%32 = load i64, i64* %31, align 8
%33 = bitcast %"class.cuneibs::neiblist_iterator_simple.2"* %5 to i8*
%34 = getelementptr inbounds i8, i8* %33, i64 %32
%35 = bitcast i8* %34 to %"class.cuneibs::neiblist_iterator_core"*
%36 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %35, i32 undef, i32 7
%37 = load i64, i64* %36, align 16
%38 = bitcast %"class.cuneibs::neiblist_iterator_simple.2"* %5 to i8**
%39 = load i8*, i8** %38, align 8
%40 = getelementptr i8, i8* %39, i64 -24
%41 = bitcast i8* %40 to i64*
%42 = load i64, i64* %41, align 8
%43 = bitcast %"class.cuneibs::neiblist_iterator_simple.2"* %5 to i8*
%44 = getelementptr inbounds i8, i8* %43, i64 %42
%45 = bitcast i8* %44 to %"class.cuneibs::neiblist_iterator_core"*
%46 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %45, i32 undef, i32 4
%47 = load i32, i32* %46, align 4
%48 = zext i32 %47 to i64
%49 = add i64 %37, %48
%50 = getelementptr inbounds i16, i16* %27, i64 %49
%51 = load i16, i16* %50, align 2
store i16 %51, i16* %4, align 2
%52 = load i16, i16* %4, align 2
%53 = zext i16 %52 to i32
%54 = icmp eq i32 %53, 65535
br i1 %54, label %55, label %56
55: ; preds = %1
store i1 false, i1* %2, align 1
br label %66
56: ; preds = %1
%57 = bitcast %"class.cuneibs::neiblist_iterator_simple.2"* %5 to i8**
%58 = load i8*, i8** %57, align 8
%59 = getelementptr i8, i8* %58, i64 -24
%60 = bitcast i8* %59 to i64*
%61 = load i64, i64* %60, align 8
%62 = bitcast %"class.cuneibs::neiblist_iterator_simple.2"* %5 to i8*
%63 = getelementptr inbounds i8, i8* %62, i64 %61
%64 = bitcast i8* %63 to %"class.cuneibs::neiblist_iterator_core"*
%65 = load i16, i16* %4, align 2
call void @_ZN7cuneibs22neiblist_iterator_core17update_neib_indexEt(%"class.cuneibs::neiblist_iterator_core"* %64, i16 %65)
store i1 true, i1* %2, align 1
br label %66
66: ; preds = %56, %55
%67 = load i1, i1* %2, align 1
ret i1 %67
}
define linkonce_odr %struct.float4 @_ZNK7cuneibs22neiblist_iterator_core6relPosERK6float4(%"class.cuneibs::neiblist_iterator_core"* %0, %struct.float4* %1) comdat align 2 {
%3 = alloca %"class.cuneibs::neiblist_iterator_core"*, i32 undef, align 8
%4 = alloca %struct.float4*, i32 undef, align 8
store %"class.cuneibs::neiblist_iterator_core"* %0, %"class.cuneibs::neiblist_iterator_core"** %3, align 8
store %struct.float4* %1, %struct.float4** %4, align 8
%5 = load %"class.cuneibs::neiblist_iterator_core"*, %"class.cuneibs::neiblist_iterator_core"** %3, align 8
%6 = getelementptr inbounds %"class.cuneibs::neiblist_iterator_core", %"class.cuneibs::neiblist_iterator_core"* %5, i32 undef, i32 5
%7 = load %struct.float4*, %struct.float4** %4, align 8
%8 = call %struct.float4 @_ZmiRK6float3RK6float4(%struct.float3* %6, %struct.float4* %7)
ret %struct.float4 undef
}
declare float @_ZN6cuphys1PEft()
declare float @_ZN6cuphys16physical_densityEft()
define linkonce_odr void @_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EEC2EjRK6float4RK4int3PKjPKt(%"class.cuneibs::neiblist_iterator_simple.2"* %0, i8** %1) unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE5resetEv(%"class.cuneibs::neiblist_iterator_simple.2"* %0) comdat align 2 {
ret void
}
define linkonce_odr i64 @_ZN7cuneibs15neib_list_startIL12ParticleType0EEEmv() comdat {
ret i64 undef
}
define linkonce_odr i64 @_ZN7cuneibs14neib_list_stepIL12ParticleType0EEEmv() comdat {
%1 = load i64, i64* addrspacecast (i64 addrspace(4)* @_ZN7cuneibs17d_neiblist_strideE to i64*), align 8, !tbaa !159
ret i64 %1
}
define linkonce_odr i32 @_ZN7cuneibs14first_neib_locIL12ParticleType0EEEjv() comdat {
ret i32 undef
}
define internal %struct.float4 @_ZmiRK6float3RK6float4(%struct.float3* %0, %struct.float4* %1) {
%3 = alloca %struct.float3*, i32 undef, align 8
%4 = alloca %struct.float4*, i32 undef, align 8
store %struct.float3* %0, %struct.float3** %3, align 8
store %struct.float4* %1, %struct.float4** %4, align 8
%5 = load %struct.float3*, %struct.float3** %3, align 8
%6 = getelementptr inbounds %struct.float3, %struct.float3* %5, i32 undef, i32 2
%7 = load float, float* %6, align 4
%8 = load %struct.float4*, %struct.float4** %4, align 8
%9 = getelementptr inbounds %struct.float4, %struct.float4* %8, i32 undef, i32 2
%10 = load float, float* %9, align 8
%11 = fsub contract float %7, %10
%12 = load %struct.float4*, %struct.float4** undef, align 8
%13 = getelementptr inbounds %struct.float4, %struct.float4* undef, i32 undef, i32 3
%14 = load float, float* undef, align 4
%15 = call %struct.float4 @_ZL11make_float4ffff(float %11, float undef)
ret %struct.float4 undef
}
define linkonce_odr %struct.float4 @_ZNK11vel_wrapper8fetchVelEj(%struct.vel_wrapper* %0) align 2 {
ret %struct.float4 undef
}
define linkonce_odr void @_ZN8cuforces20common_particle_dataC2EjRK6float4RK7ushort4PKj(%"struct.cuforces::common_particle_data"* %0) unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN8cuforces16rb_particle_dataC2ERK7ushort4(%"struct.cuforces::rb_particle_data"* %0) unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN5emptyIN8cuforces25gamma_cfl_particle_outputEEC2Ev() unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN5emptyIN8cuforces20keps_particle_outputEEC2Ev() unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN5emptyIN8cuforces20xsph_particle_outputEEC2Ev() unnamed_addr align 2 {
ret void
}
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* %0) {
%2 = alloca %"class.cuneibs::neiblist_iterator.1", align 16
%3 = alloca i32, i32 undef, align 4
%4 = alloca %struct.float4, i32 undef, align 16
br label %5
5: ; preds = %1
br label %6
6: ; preds = %5
br label %7
7: ; preds = %6
%8 = bitcast %struct.forces_params.415* %0 to %struct.pos_wrapper*
%9 = call %struct.float4 @_ZNK11pos_wrapper8fetchPosEj(%struct.pos_wrapper* %8, i32 undef)
br label %10
10: ; preds = %7
br label %11
11: ; preds = %10
%12 = bitcast %struct.forces_params.415* %0 to %struct.common_forces_params*
%13 = getelementptr inbounds %struct.common_forces_params, %struct.common_forces_params* %12, i32 undef, i32 1
%14 = load i16*, i16** %13, align 8
call void @_ZN7cuneibs17neiblist_iteratorIJL12ParticleType0EEEC1EjRK6float4RK4int3PKjPKt(%"class.cuneibs::neiblist_iterator.1"* %2, i16* %14)
br label %15
15: ; preds = %48, %11
%16 = bitcast %"class.cuneibs::neiblist_iterator.1"* %2 to %"class.cuneibs::neiblist_iterator_simple.2"*
%17 = call i1 @_ZN7cuneibs24neiblist_iterator_simpleIL12ParticleType0EE4nextEv(%"class.cuneibs::neiblist_iterator_simple.2"* %16)
br i1 %17, label %19, label %18
18: ; preds = %15
ret void
19: ; preds = %15
%20 = bitcast %"class.cuneibs::neiblist_iterator.1"* %2 to i8**
%21 = load i8*, i8** %20, align 16, !tbaa !146
%22 = getelementptr i8, i8* %21, i64 -24
%23 = bitcast i8* %22 to i64*
%24 = load i64, i64* %23, align 8
%25 = bitcast %"class.cuneibs::neiblist_iterator.1"* %2 to i8*
%26 = getelementptr inbounds i8, i8* %25, i64 %24
%27 = bitcast i8* %26 to %"class.cuneibs::neiblist_iterator_core"*
%28 = call i32* @_ZNK7cuneibs22neiblist_iterator_core10neib_indexEv(%"class.cuneibs::neiblist_iterator_core"* %27)
%29 = load i32, i32* %28, align 4
store i32 %29, i32* %3, align 4
%30 = bitcast %"class.cuneibs::neiblist_iterator.1"* %2 to i8**
%31 = load i8*, i8** %30, align 16
%32 = getelementptr i8, i8* %31, i64 -24
%33 = bitcast i8* %32 to i64*
%34 = load i64, i64* %33, align 8
%35 = bitcast %"class.cuneibs::neiblist_iterator.1"* %2 to i8*
%36 = getelementptr inbounds i8, i8* %35, i64 %34
%37 = bitcast i8* %36 to %"class.cuneibs::neiblist_iterator_core"*
%38 = bitcast %struct.forces_params.415* %0 to %struct.pos_wrapper*
%39 = load i32, i32* %3, align 4
%40 = call %struct.float4 @_ZNK11pos_wrapper8fetchPosEj(%struct.pos_wrapper* %38, i32 %39)
%41 = getelementptr inbounds %struct.float4, %struct.float4* %4, i32 undef, i32 2
%42 = extractvalue %struct.float4 undef, 2
store float undef, float* %41, align 8
%43 = getelementptr inbounds %struct.float4, %struct.float4* %4, i32 undef, i32 3
%44 = extractvalue %struct.float4 undef, 3
store float undef, float* %43, align 4
%45 = call %struct.float4 @_ZNK7cuneibs22neiblist_iterator_core6relPosERK6float4(%"class.cuneibs::neiblist_iterator_core"* %37, %struct.float4* %4)
%46 = call i1 @_ZL8isfinitef(float undef)
br label %47
47: ; preds = %19
br label %48
48: ; preds = %47
switch i32 undef, label %49 [
i32 5, label %15
]
49: ; preds = %48
unreachable
}
define linkonce_odr void @_ZN8cuforces20forces_particle_dataIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSC_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyINS_20volume_particle_dataEESE_INS_21grenier_particle_dataEESE_INS_22eulerVel_particle_dataEESE_INS_25sa_boundary_particle_dataEESE_INS_18keps_particle_dataEESE_INS_17sps_particle_dataEESE_INS_28effective_visc_particle_dataEESE_INS_26keps_precalc_particle_dataEEEC1I13forces_paramsILS1_3ELS2_1ELS3_3ELS4_4ESB_Lm517ELSC_1ELSC_0ELSD_1ELb0ELb0ELb0ELb0ESE_I18xsph_forces_paramsESE_I20volume_forces_paramsESE_I21grenier_forces_paramsESE_I25sa_boundary_forces_paramsESE_I28dummy_boundary_forces_paramsESE_I25water_depth_forces_paramsESE_I18keps_forces_paramsESE_I14tau_tex_paramsESE_I22eulerVel_forces_paramsESE_I29internal_energy_forces_paramsESE_I28effective_visc_forces_paramsEEEEjRK6float4RK7ushort4RKT_(%"struct.cuforces::forces_particle_data.416"* %0, %struct.forces_params.415* %1) unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN8cuforces22forces_particle_outputIL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELS9_0ELb0EEC1I13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3ELS1_4ES8_Lm517ELS9_1ELS9_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyI18xsph_forces_paramsESH_I20volume_forces_paramsESH_I21grenier_forces_paramsESH_I25sa_boundary_forces_paramsESH_I28dummy_boundary_forces_paramsESH_I25water_depth_forces_paramsESH_I18keps_forces_paramsESH_I14tau_tex_paramsESH_I22eulerVel_forces_paramsESH_I29internal_energy_forces_paramsESH_I28effective_visc_forces_paramsEEEERKT_j() unnamed_addr align 2 {
ret void
}
define linkonce_odr i1 @_ZN8cuforces13skip_neiblistI13forces_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_paramsEENS_20forces_particle_dataILS2_3ELS3_1ELS4_3ELS5_4ESC_Lm517ELSD_1ELSD_0ELSE_1ELb0ELb0ELb0ELb0ESF_INS_20volume_particle_dataEESF_INS_21grenier_particle_dataEESF_INS_22eulerVel_particle_dataEESF_INS_25sa_boundary_particle_dataEESF_INS_18keps_particle_dataEESF_INS_17sps_particle_dataEESF_INS_28effective_visc_particle_dataEESF_INS_26keps_precalc_particle_dataEEEEEENSt9enable_ifIXoooonesrT_12boundarytypeLS5_2EsrS1M_8has_kepsnesrS1M_6cptypeLSD_2EEbE4typeERKS1M_RKT0_() {
ret i1 undef
}
define linkonce_odr void @_ZN8cuforces20forces_particle_dataIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSC_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyINS_20volume_particle_dataEESE_INS_21grenier_particle_dataEESE_INS_22eulerVel_particle_dataEESE_INS_25sa_boundary_particle_dataEESE_INS_18keps_particle_dataEESE_INS_17sps_particle_dataEESE_INS_28effective_visc_particle_dataEESE_INS_26keps_precalc_particle_dataEEEC2I13forces_paramsILS1_3ELS2_1ELS3_3ELS4_4ESB_Lm517ELSC_1ELSC_0ELSD_1ELb0ELb0ELb0ELb0ESE_I18xsph_forces_paramsESE_I20volume_forces_paramsESE_I21grenier_forces_paramsESE_I25sa_boundary_forces_paramsESE_I28dummy_boundary_forces_paramsESE_I25water_depth_forces_paramsESE_I18keps_forces_paramsESE_I14tau_tex_paramsESE_I22eulerVel_forces_paramsESE_I29internal_energy_forces_paramsESE_I28effective_visc_forces_paramsEEEEjRK6float4RK7ushort4RKT_(%"struct.cuforces::forces_particle_data.416"* %0, %struct.forces_params.415* %1) unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN8cuforces17vel_particle_dataC2I13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSE_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyI18xsph_forces_paramsESG_I20volume_forces_paramsESG_I21grenier_forces_paramsESG_I25sa_boundary_forces_paramsESG_I28dummy_boundary_forces_paramsESG_I25water_depth_forces_paramsESG_I18keps_forces_paramsESG_I14tau_tex_paramsESG_I22eulerVel_forces_paramsESG_I29internal_energy_forces_paramsESG_I28effective_visc_forces_paramsEEEERKT_jRK7ushort4(%"struct.cuforces::vel_particle_data"* %0, %struct.forces_params.415* %1) unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN5emptyIN8cuforces20volume_particle_dataEEC2IJj13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSG_0EL7RunMode1ELb0ELb0ELb0ELb0ES_I18xsph_forces_paramsES_I20volume_forces_paramsES_I21grenier_forces_paramsES_I25sa_boundary_forces_paramsES_I28dummy_boundary_forces_paramsES_I25water_depth_forces_paramsES_I18keps_forces_paramsES_I14tau_tex_paramsES_I22eulerVel_forces_paramsES_I29internal_energy_forces_paramsES_I28effective_visc_forces_paramsEEEEEDpRKT_() unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN5emptyIN8cuforces21grenier_particle_dataEEC2IJj13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSG_0EL7RunMode1ELb0ELb0ELb0ELb0ES_I18xsph_forces_paramsES_I20volume_forces_paramsES_I21grenier_forces_paramsES_I25sa_boundary_forces_paramsES_I28dummy_boundary_forces_paramsES_I25water_depth_forces_paramsES_I18keps_forces_paramsES_I14tau_tex_paramsES_I22eulerVel_forces_paramsES_I29internal_energy_forces_paramsES_I28effective_visc_forces_paramsEEEEEDpRKT_() unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN5emptyIN8cuforces22eulerVel_particle_dataEEC2IJj13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSG_0EL7RunMode1ELb0ELb0ELb0ELb0ES_I18xsph_forces_paramsES_I20volume_forces_paramsES_I21grenier_forces_paramsES_I25sa_boundary_forces_paramsES_I28dummy_boundary_forces_paramsES_I25water_depth_forces_paramsES_I18keps_forces_paramsES_I14tau_tex_paramsES_I22eulerVel_forces_paramsES_I29internal_energy_forces_paramsES_I28effective_visc_forces_paramsEEEEEDpRKT_() unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN5emptyIN8cuforces25sa_boundary_particle_dataEEC2IJj7ushort413forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSH_0EL7RunMode1ELb0ELb0ELb0ELb0ES_I18xsph_forces_paramsES_I20volume_forces_paramsES_I21grenier_forces_paramsES_I25sa_boundary_forces_paramsES_I28dummy_boundary_forces_paramsES_I25water_depth_forces_paramsES_I18keps_forces_paramsES_I14tau_tex_paramsES_I22eulerVel_forces_paramsES_I29internal_energy_forces_paramsES_I28effective_visc_forces_paramsEEEEEDpRKT_() unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN5emptyIN8cuforces18keps_particle_dataEEC2IJj7ushort413forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSH_0EL7RunMode1ELb0ELb0ELb0ELb0ES_I18xsph_forces_paramsES_I20volume_forces_paramsES_I21grenier_forces_paramsES_I25sa_boundary_forces_paramsES_I28dummy_boundary_forces_paramsES_I25water_depth_forces_paramsES_I18keps_forces_paramsES_I14tau_tex_paramsES_I22eulerVel_forces_paramsES_I29internal_energy_forces_paramsES_I28effective_visc_forces_paramsEEEEEDpRKT_() unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN5emptyIN8cuforces17sps_particle_dataEEC2IJj13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSG_0EL7RunMode1ELb0ELb0ELb0ELb0ES_I18xsph_forces_paramsES_I20volume_forces_paramsES_I21grenier_forces_paramsES_I25sa_boundary_forces_paramsES_I28dummy_boundary_forces_paramsES_I25water_depth_forces_paramsES_I18keps_forces_paramsES_I14tau_tex_paramsES_I22eulerVel_forces_paramsES_I29internal_energy_forces_paramsES_I28effective_visc_forces_paramsEEEEEDpRKT_() unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN8cuforces23p_precalc_particle_dataC2INS_20forces_particle_dataIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSE_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyINS_20volume_particle_dataEESG_INS_21grenier_particle_dataEESG_INS_22eulerVel_particle_dataEESG_INS_25sa_boundary_particle_dataEESG_INS_18keps_particle_dataEESG_INS_17sps_particle_dataEESG_INS_28effective_visc_particle_dataEESG_INS_26keps_precalc_particle_dataEEEEEERKT_(%"struct.cuforces::p_precalc_particle_data"* %0) unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN5emptyIN8cuforces28effective_visc_particle_dataEEC2IJj13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSG_0EL7RunMode1ELb0ELb0ELb0ELb0ES_I18xsph_forces_paramsES_I20volume_forces_paramsES_I21grenier_forces_paramsES_I25sa_boundary_forces_paramsES_I28dummy_boundary_forces_paramsES_I25water_depth_forces_paramsES_I18keps_forces_paramsES_I14tau_tex_paramsES_I22eulerVel_forces_paramsES_I29internal_energy_forces_paramsES_I28effective_visc_forces_paramsEEEEEDpRKT_() unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN5emptyIN8cuforces26keps_precalc_particle_dataEEC2IJfNS0_20forces_particle_dataIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSG_0EL7RunMode1ELb0ELb0ELb0ELb0ES_INS0_20volume_particle_dataEES_INS0_21grenier_particle_dataEES_INS0_22eulerVel_particle_dataEES_INS0_25sa_boundary_particle_dataEES_INS0_18keps_particle_dataEES_INS0_17sps_particle_dataEES_INS0_28effective_visc_particle_dataEES2_EESW_EEEDpRKT_() unnamed_addr align 2 {
ret void
}
define linkonce_odr float @_ZN8cuforces16precalc_pressureINS_20forces_particle_dataIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSD_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyINS_20volume_particle_dataEESF_INS_21grenier_particle_dataEESF_INS_22eulerVel_particle_dataEESF_INS_25sa_boundary_particle_dataEESF_INS_18keps_particle_dataEESF_INS_17sps_particle_dataEESF_INS_28effective_visc_particle_dataEESF_INS_26keps_precalc_particle_dataEEEEEENSt9enable_ifIXeqsrT_15sph_formulationLS3_1EEfE4typeERKSY_() {
ret float undef
}
define linkonce_odr float @_ZN6cuphys16physical_densityIN8cuforces20forces_particle_dataIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSE_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyINS1_20volume_particle_dataEESG_INS1_21grenier_particle_dataEESG_INS1_22eulerVel_particle_dataEESG_INS1_25sa_boundary_particle_dataEESG_INS1_18keps_particle_dataEESG_INS1_17sps_particle_dataEESG_INS1_28effective_visc_particle_dataEESG_INS1_26keps_precalc_particle_dataEEEEEEfRKT_() {
ret float undef
}
define linkonce_odr float @_ZN8cuforces20pressure_for_precalcINS_20forces_particle_dataIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSD_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyINS_20volume_particle_dataEESF_INS_21grenier_particle_dataEESF_INS_22eulerVel_particle_dataEESF_INS_25sa_boundary_particle_dataEESF_INS_18keps_particle_dataEESF_INS_17sps_particle_dataEESF_INS_28effective_visc_particle_dataEESF_INS_26keps_precalc_particle_dataEEEEEENSt9enable_ifIXntsrT_8has_kepsEfE4typeERKSY_() {
ret float undef
}
define linkonce_odr float @_ZN6cuphys16particle_densityIN8cuforces20forces_particle_dataIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSE_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyINS1_20volume_particle_dataEESG_INS1_21grenier_particle_dataEESG_INS1_22eulerVel_particle_dataEESG_INS1_25sa_boundary_particle_dataEESG_INS1_18keps_particle_dataEESG_INS1_17sps_particle_dataEESG_INS1_28effective_visc_particle_dataEESG_INS1_26keps_precalc_particle_dataEEEEEENSt9enable_ifIXntcl10has_relVelIT_EEEfE4typeERKSZ_() {
ret float undef
}
define linkonce_odr float @_ZN8cuforces17particle_pressureINS_20forces_particle_dataIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSD_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyINS_20volume_particle_dataEESF_INS_21grenier_particle_dataEESF_INS_22eulerVel_particle_dataEESF_INS_25sa_boundary_particle_dataEESF_INS_18keps_particle_dataEESF_INS_17sps_particle_dataEESF_INS_28effective_visc_particle_dataEESF_INS_26keps_precalc_particle_dataEEEEEEfRKT_() {
ret float undef
}
define linkonce_odr void @_ZN8cuforces22forces_particle_outputIL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELS9_0ELb0EEC2I13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3ELS1_4ES8_Lm517ELS9_1ELS9_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyI18xsph_forces_paramsESH_I20volume_forces_paramsESH_I21grenier_forces_paramsESH_I25sa_boundary_forces_paramsESH_I28dummy_boundary_forces_paramsESH_I25water_depth_forces_paramsESH_I18keps_forces_paramsESH_I14tau_tex_paramsESH_I22eulerVel_forces_paramsESH_I29internal_energy_forces_paramsESH_I28effective_visc_forces_paramsEEEERKT_j() unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN8cuforces22common_particle_outputC2I13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSE_0EL7RunMode1ELb0ELb0ELb0ELb0E5emptyI18xsph_forces_paramsESG_I20volume_forces_paramsESG_I21grenier_forces_paramsESG_I25sa_boundary_forces_paramsESG_I28dummy_boundary_forces_paramsESG_I25water_depth_forces_paramsESG_I18keps_forces_paramsESG_I14tau_tex_paramsESG_I22eulerVel_forces_paramsESG_I29internal_energy_forces_paramsESG_I28effective_visc_forces_paramsEEEERKT_j() unnamed_addr align 2 {
ret void
}
define linkonce_odr void @_ZN5emptyIN8cuforces31internal_energy_particle_outputEEC2IJ13forces_paramsIL10KernelType3EL14SPHFormulation1EL20DensityDiffusionType3EL12BoundaryType4E12FullViscSpecIL12RheologyType0EL15TurbulenceModel1EL26ComputationalViscosityType0EL12ViscousModel0EL15AverageOperator0ELm517ELb0EELm517EL12ParticleType1ELSG_0EL7RunMode1ELb0ELb0ELb0ELb0ES_I18xsph_forces_paramsES_I20volume_forces_paramsES_I21grenier_forces_paramsES_I25sa_boundary_forces_paramsES_I28dummy_boundary_forces_paramsES_I25water_depth_forces_paramsES_I18keps_forces_paramsES_I14tau_tex_paramsES_I22eulerVel_forces_paramsES_I29internal_energy_forces_paramsES_I28effective_visc_forces_paramsEEjEEEDpRKT_() unnamed_addr align 2 {
ret void
}
attributes #0 = { argmemonly nofree nosync nounwind willreturn }
attributes #1 = { argmemonly 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, !58, !59, !60, !61, !62, !63, !64, !65, !66, !67, !68, !69, !70, !71, !72, !73, !74, !75, !76, !77, !78, !79, !80, !81, !82, !83, !84, !85, !86, !87, !88, !89, !90, !91, !92, !93, !94, !95, !96, !97, !98, !99, !100, !101, !102, !103, !104, !105, !106, !107, !108, !109, !110, !111, !112, !113, !114, !115, !116, !117, !118, !119, !120, !121, !122, !123, !124, !125, !126, !127, !128, !129, !130, !131, !132, !133, !134, !135, !136, !137, !138, !139, !140, !141, !142, !143, !144, !145}
!0 = distinct !{void (i32*, i32*, %struct.ushort4*, i32*, i32)* undef, !"kernel", i32 1}
!1 = distinct !{void (i32*, i32*, %struct.ushort4*, i32*, i32)* undef, !"maxntidx", i32 256}
!2 = distinct !{void (i32*, i32*, %struct.ushort4*, i32*, i32)* undef, !"minctasm", i32 6}
!3 = distinct !{void (%struct.cell_params*, i32)* undef, !"kernel", i32 1}
!4 = distinct !{void (%struct.float4*, %struct.ushort4*, i32*, %struct.uint4*, i32*, i16*, i32, float, float)* undef, !"kernel", i32 1}
!5 = distinct !{void (%struct.ushort4*, %struct.float4*, %struct.uint4*, i32)* undef, !"kernel", i32 1}
!6 = distinct !{void (%struct.ushort4*, %struct.float4*, %struct.float4*, i32, i32, float)* undef, !"kernel", i32 1}
!7 = distinct !{void (%struct.float4*, %struct.ushort4*, i32)* undef, !"kernel", i32 1}
!8 = distinct !{void (%struct.pos_info_wrapper*, float*, i32, float)* undef, !"kernel", i32 1}
!9 = distinct !{void (%struct.pos_info_wrapper*, float*, i32, float)* undef, !"maxntidx", i32 128}
!10 = distinct !{void (%struct.pos_info_wrapper*, float*, i32, float)* undef, !"minctasm", i32 6}
!11 = distinct !{void (%struct.jacobi_update_params*)* undef, !"kernel", i32 1}
!12 = distinct !{void (%struct.jacobi_update_params*)* undef, !"maxntidx", i32 128}
!13 = distinct !{void (%struct.jacobi_update_params*)* undef, !"minctasm", i32 6}
!14 = distinct !{void (%struct.ushort4*, %struct.float4*, %struct.float4*, float*, i32)* undef, !"kernel", i32 1}
!15 = distinct !{void (%struct.float4*, %struct.float4*, %struct.ushort4*, i32*, i32, i32, %struct.float4*)* undef, !"kernel", i32 1}
!16 = distinct !{void (%struct.float4*, i32, i32)* undef, !"kernel", i32 1}
!17 = distinct !{void ()* undef, !"kernel", i32 1}
!18 = distinct !{void (float*, %struct.float4*, i32)* undef, !"kernel", i32 1}
!19 = distinct !{void (%struct.float4*, %struct.float4*, %struct.float4*, %struct.float4*, %struct.ushort4*, i32*, i16*, i32*, i32, float, float)* undef, !"kernel", i32 1}
!20 = distinct !{void (%struct.float4*, %struct.float4*, %struct.float4*, %struct.float4*, %struct.ushort4*, i32*, i16*, i32*, i32, float, float)* undef, !"maxntidx", i32 128}
!21 = distinct !{void (%struct.float4*, %struct.float4*, %struct.float4*, %struct.float4*, %struct.ushort4*, i32*, i16*, i32*, i32, float, float)* undef, !"minctasm", i32 6}
!22 = distinct !{void (%struct.sa_outgoing_bc_params*)* undef, !"kernel", i32 1}
!23 = distinct !{void (%struct.float4*, %struct.uint4*, %struct.ushort4*, i32*, i32*, i16*, i32)* undef, !"kernel", i32 1}
!24 = distinct !{void (%struct.uint4*, i32*, %struct.ushort4*, i32*, i16*, %struct.float4*, i32)* undef, !"kernel", i32 1}
!25 = distinct !{void (%struct.uint4*, i32*, %struct.ushort4*, i32*, i16*, %struct.float4*, i32)* undef, !"maxntidx", i32 128}
!26 = distinct !{void (%struct.uint4*, i32*, %struct.ushort4*, i32*, i16*, %struct.float4*, i32)* undef, !"minctasm", i32 6}
!27 = distinct !{void (%struct.float4*, %struct.float4*, %struct.uint4*, i32*, %struct.ushort4*, i32*, i16*, %struct.float4*, i32, float)* undef, !"kernel", i32 1}
!28 = distinct !{void (%struct.float4*, %struct.float4*, %struct.uint4*, i32*, %struct.ushort4*, i32*, i16*, %struct.float4*, i32, float)* undef, !"maxntidx", i32 128}
!29 = distinct !{void (%struct.float4*, %struct.float4*, %struct.uint4*, i32*, %struct.ushort4*, i32*, i16*, %struct.float4*, i32, float)* undef, !"minctasm", i32 6}
!30 = distinct !{void (%struct.float4*, i32*, i32*, %struct.ushort4*, i32*, i32)* undef, !"kernel", i32 1}
!31 = distinct !{void (%struct.float4*, i32*, i32*, %struct.ushort4*, i32*, i32)* undef, !"maxntidx", i32 256}
!32 = distinct !{void (%struct.float4*, i32*, i32*, %struct.ushort4*, i32*, i32)* undef, !"minctasm", i32 6}
!33 = distinct !{void (%struct.reorder_params*, i32*, i32*, i32*, %struct.ushort4*, i32*, i32*, i32, i32*)* undef, !"kernel", i32 1}
!34 = distinct !{void (%struct.reorder_params*, i32*, i32*, i32*, %struct.ushort4*, i32*, i32*, i32, i32*)* undef, !"maxntidx", i32 256}
!35 = distinct !{void (%struct.reorder_params*, i32*, i32*, i32*, %struct.ushort4*, i32*, i32*, i32, i32*)* undef, !"minctasm", i32 6}
!36 = distinct !{void (i1, %"class.thrust::zip_iterator"*, %"class.thrust::device_ptr"*, i64, %"class.thrust::tuple"*, i32*, %struct.ptype_hash_compare*)* undef, !"kernel", i32 1}
!37 = distinct !{void (i1, %"class.thrust::zip_iterator"*, %"class.thrust::device_ptr"*, i64, %"class.thrust::tuple"*, i32*, %struct.ptype_hash_compare*)* undef, !"maxntidx", i32 256}
!38 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f"*, i64)* undef, !"kernel", i32 1}
!39 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f"*, i64)* undef, !"maxntidx", i32 256}
!40 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.69"*, i64)* undef, !"kernel", i32 1}
!41 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.69"*, i64)* undef, !"maxntidx", i32 256}
!42 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.72"*, i64)* undef, !"kernel", i32 1}
!43 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.72"*, i64)* undef, !"maxntidx", i32 256}
!44 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.74"*, i64)* undef, !"kernel", i32 1}
!45 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.74"*, i64)* undef, !"maxntidx", i32 256}
!46 = distinct !{void (i1, %"class.thrust::zip_iterator"*, %"class.thrust::tuple"*, i64, i64, i64*, %struct.ptype_hash_compare*, i64, i32)* undef, !"kernel", i32 1}
!47 = distinct !{void (i1, %"class.thrust::zip_iterator"*, %"class.thrust::tuple"*, i64, i64, i64*, %struct.ptype_hash_compare*, i64, i32)* undef, !"maxntidx", i32 256}
!48 = distinct !{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 = distinct !{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 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.79"*, i64)* undef, !"kernel", i32 1}
!51 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.79"*, i64)* undef, !"maxntidx", i32 256}
!52 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.81"*, i64)* undef, !"kernel", i32 1}
!53 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.81"*, i64)* undef, !"maxntidx", i32 256}
!54 = distinct !{void (%struct.buildneibs_params*)* undef, !"kernel", i32 1}
!55 = distinct !{void (%struct.buildneibs_params*)* undef, !"maxntidx", i32 256}
!56 = distinct !{void (%struct.buildneibs_params*)* undef, !"minctasm", i32 5}
!57 = distinct !{void (%struct.euler_params*)* undef, !"kernel", i32 1}
!58 = distinct !{void (%struct.euler_params*)* undef, !"kernel", i32 1}
!59 = distinct !{void (%struct.euler_params.111*)* undef, !"kernel", i32 1}
!60 = distinct !{void (%struct.euler_params.111*)* undef, !"kernel", i32 1}
!61 = distinct !{void (%struct.euler_params.114*)* undef, !"kernel", i32 1}
!62 = distinct !{void (%struct.euler_params.114*)* undef, !"kernel", i32 1}
!63 = distinct !{void (%struct.euler_params.117*)* undef, !"kernel", i32 1}
!64 = distinct !{void (%struct.euler_params.117*)* undef, !"kernel", i32 1}
!65 = distinct !{void (%"struct.cub::ReduceByKeyScanTileState"*, i64)* undef, !"kernel", i32 1}
!66 = distinct !{void (%"struct.cub::ReduceByKeyScanTileState"*, i64)* undef, !"maxntidx", i32 128}
!67 = distinct !{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}
!68 = distinct !{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}
!69 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.166"*, i64)* undef, !"kernel", i32 1}
!70 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.166"*, i64)* undef, !"maxntidx", i32 256}
!71 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.169"*, i64)* undef, !"kernel", i32 1}
!72 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.169"*, i64)* undef, !"maxntidx", i32 256}
!73 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.171"*, i64)* undef, !"kernel", i32 1}
!74 = distinct !{void (%"struct.thrust::cuda_cub::__transform::unary_transform_f.171"*, i64)* undef, !"maxntidx", i32 256}
!75 = distinct !{void (%struct.density_diffusion_params*)* undef, !"kernel", i32 1}
!76 = distinct !{void (%struct.forces_params*)* undef, !"kernel", i32 1}
!77 = distinct !{void (%struct.forces_params.223*)* undef, !"kernel", i32 1}
!78 = distinct !{void (%struct.finalize_forces_params*)* undef, !"kernel", i32 1}
!79 = distinct !{void (%struct.forces_params.236*)* undef, !"kernel", i32 1}
!80 = distinct !{void (%struct.forces_params.250*)* undef, !"kernel", i32 1}
!81 = distinct !{void (%struct.forces_params.256*)* undef, !"kernel", i32 1}
!82 = distinct !{void (%struct.finalize_forces_params.262*)* undef, !"kernel", i32 1}
!83 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*)* undef, !"kernel", i32 1}
!84 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*)* undef, !"maxntidx", i32 128}
!85 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*)* undef, !"minctasm", i32 6}
!86 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*)* undef, !"kernel", i32 1}
!87 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*)* undef, !"maxntidx", i32 128}
!88 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*)* undef, !"minctasm", i32 6}
!89 = distinct !{void (%struct.neibs_interaction_params*, %struct.float3*)* undef, !"kernel", i32 1}
!90 = distinct !{void (%"struct.cupostprocess::testpoints_params"*)* undef, !"kernel", i32 1}
!91 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*, %struct.ushort4*)* undef, !"kernel", i32 1}
!92 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*, %struct.ushort4*)* undef, !"kernel", i32 1}
!93 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*, %struct.ushort4*, float)* undef, !"kernel", i32 1}
!94 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*, %struct.ushort4*, float)* undef, !"kernel", i32 1}
!95 = distinct !{void (%struct.reorder_params.278*, i32*, i32*, i32*, %struct.ushort4*, i32*, i32*, i32, i32*)* undef, !"kernel", i32 1}
!96 = distinct !{void (%struct.reorder_params.278*, i32*, i32*, i32*, %struct.ushort4*, i32*, i32*, i32, i32*)* undef, !"maxntidx", i32 256}
!97 = distinct !{void (%struct.reorder_params.278*, i32*, i32*, i32*, %struct.ushort4*, i32*, i32*, i32, i32*)* undef, !"minctasm", i32 6}
!98 = distinct !{void (%struct.buildneibs_params.279*)* undef, !"kernel", i32 1}
!99 = distinct !{void (%struct.buildneibs_params.279*)* undef, !"maxntidx", i32 256}
!100 = distinct !{void (%struct.buildneibs_params.279*)* undef, !"minctasm", i32 5}
!101 = distinct !{void (%struct.euler_params.282*)* undef, !"kernel", i32 1}
!102 = distinct !{void (%struct.euler_params.282*)* undef, !"kernel", i32 1}
!103 = distinct !{void (%struct.euler_params.285*)* undef, !"kernel", i32 1}
!104 = distinct !{void (%struct.euler_params.285*)* undef, !"kernel", i32 1}
!105 = distinct !{void (%struct.euler_params.288*)* undef, !"kernel", i32 1}
!106 = distinct !{void (%struct.euler_params.288*)* undef, !"kernel", i32 1}
!107 = distinct !{void (%struct.euler_params.291*)* undef, !"kernel", i32 1}
!108 = distinct !{void (%struct.euler_params.291*)* undef, !"kernel", i32 1}
!109 = distinct !{void (%struct.density_diffusion_params*)* undef, !"kernel", i32 1}
!110 = distinct !{void (%struct.forces_params.294*)* undef, !"kernel", i32 1}
!111 = distinct !{void (%struct.forces_params.300*)* undef, !"kernel", i32 1}
!112 = distinct !{void (%struct.finalize_forces_params.305*)* undef, !"kernel", i32 1}
!113 = distinct !{void (%struct.forces_params.312*)* undef, !"kernel", i32 1}
!114 = distinct !{void (%struct.forces_params.319*)* undef, !"kernel", i32 1}
!115 = distinct !{void (%struct.forces_params.325*)* undef, !"kernel", i32 1}
!116 = distinct !{void (%struct.finalize_forces_params.331*)* undef, !"kernel", i32 1}
!117 = distinct !{void (%"struct.cupostprocess::testpoints_params.336"*)* undef, !"kernel", i32 1}
!118 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*, %struct.ushort4*)* undef, !"kernel", i32 1}
!119 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*, %struct.ushort4*)* undef, !"kernel", i32 1}
!120 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*, %struct.ushort4*, float)* undef, !"kernel", i32 1}
!121 = distinct !{void (%struct.neibs_interaction_params*, %struct.float4*, %struct.ushort4*, float)* undef, !"kernel", i32 1}
!122 = distinct !{void (%struct.density_diffusion_params.339*)* undef, !"kernel", i32 1}
!123 = distinct !{void (%struct.forces_params.341*)* undef, !"kernel", i32 1}
!124 = distinct !{void (%struct.forces_params.346*)* undef, !"kernel", i32 1}
!125 = distinct !{void (%struct.forces_params.351*)* undef, !"kernel", i32 1}
!126 = distinct !{void (%struct.density_diffusion_params.339*)* undef, !"kernel", i32 1}
!127 = distinct !{void (%struct.forces_params.356*)* undef, !"kernel", i32 1}
!128 = distinct !{void (%struct.forces_params.361*)* undef, !"kernel", i32 1}
!129 = distinct !{void (%struct.forces_params.366*)* undef, !"kernel", i32 1}
!130 = distinct !{void (%struct.density_diffusion_params.371*)* undef, !"kernel", i32 1}
!131 = distinct !{void (%struct.forces_params.373*)* undef, !"kernel", i32 1}
!132 = distinct !{void (%struct.forces_params.378*)* undef, !"kernel", i32 1}
!133 = distinct !{void (%struct.forces_params.383*)* undef, !"kernel", i32 1}
!134 = distinct !{void (%struct.density_diffusion_params.371*)* undef, !"kernel", i32 1}
!135 = distinct !{void (%struct.forces_params.388*)* undef, !"kernel", i32 1}
!136 = distinct !{void (%struct.forces_params.393*)* undef, !"kernel", i32 1}
!137 = distinct !{void (%struct.forces_params.398*)* undef, !"kernel", i32 1}
!138 = distinct !{void (%struct.density_diffusion_params.403*)* undef, !"kernel", i32 1}
!139 = distinct !{void (%struct.forces_params.405*)* undef, !"kernel", i32 1}
!140 = distinct !{void (%struct.forces_params.410*)* undef, !"kernel", i32 1}
!141 = !{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}
!142 = distinct !{void (%struct.density_diffusion_params.403*)* undef, !"kernel", i32 1}
!143 = distinct !{void (%struct.forces_params.420*)* undef, !"kernel", i32 1}
!144 = distinct !{void (%struct.forces_params.425*)* undef, !"kernel", i32 1}
!145 = distinct !{void (%struct.forces_params.430*)* undef, !"kernel", i32 1}
!146 = !{!147, !147, i64 0}
!147 = !{!"vtable pointer", !148, i64 0}
!148 = !{!"Simple C++ TBAA"}
!149 = !{!150, !158, i64 64}
!150 = !{!"_ZTSN7cuneibs22neiblist_iterator_coreE", !151, i64 0, !151, i64 8, !153, i64 16, !155, i64 32, !156, i64 44, !157, i64 48, !158, i64 64, !156, i64 72, !152, i64 76, !156, i64 80}
!151 = !{!"any pointer", !152, i64 0}
!152 = !{!"omnipotent char", !148, i64 0}
!153 = !{!"_ZTS6float4", !154, i64 0, !154, i64 4, !154, i64 8, !154, i64 12}
!154 = !{!"float", !152, i64 0}
!155 = !{!"_ZTS4int3", !156, i64 0, !156, i64 4, !156, i64 8}
!156 = !{!"int", !152, i64 0}
!157 = !{!"_ZTS6float3", !154, i64 0, !154, i64 4, !154, i64 8}
!158 = !{!"long", !152, i64 0}
!159 = !{!158, !158, i64 0}
@Artem-B
Copy link
Author

Artem-B commented Oct 11, 2021

llvm-reduce test script:

#! /bin/bash

T=path-to-the-test-dir
D=path-to-clang-head-dir/
FC=$D/bin/FileCheck
CLANG_11=clang-11
CLANG_HEAD=$D/bin/clang

ARGS=( "-cc1" "-triple" "nvptx64-nvidia-cuda" "-aux-triple"
  "x86_64-pc-linux-gnu" "-S" "-aux-target-cpu" "x86-64" "-fcuda-is-device"
  "-target-feature" "+ptx70" "-target-sdk-version=11.0" "-target-cpu" "sm_70"
  "-O3"  "-x" "ir" )

set -e

# We need to give old clang compatible bitcode, so dsassemble it w/ llvm-dis-11
LL_11=$1.ll
sed -e 's/seq_cst, align 4/seq_cst/' < $1 | opt-11 --strip-nonlinetable-debuginfo -S -o $LL_11

$CLANG_HEAD "${ARGS[@]}" $LL_11 -o clang-head.ptx
$FC --input-file=clang-head.ptx --check-prefixes=COMMON,HEAD $T/test.sh
$CLANG_11 "${ARGS[@]}" $LL_11 -o clang-11.ptx
$FC --input-file=clang-11.ptx --check-prefixes=COMMON,OLD $T/test.sh

exit 0

# COMMON-LABEL: .visible .entry _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_(
#
# HEAD:  __local_depot{{.*}}[112];
# OLD-NOT: __local_depot
# COMMON: ret;

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment