diff --git a/cuda/slonczewski.go b/cuda/slonczewski.go index cd35d9e98..bab29cfe8 100644 --- a/cuda/slonczewski.go +++ b/cuda/slonczewski.go @@ -1,17 +1,15 @@ package cuda import ( - "unsafe" - "github.com/mumax/3/data" ) // Add Slonczewski ST torque to torque (Tesla). // see slonczewski.cu -func AddSlonczewskiTorque2(torque, m *data.Slice, Msat, J, fixedP, alpha, pol, λ, ε_prime MSlice, flp float64, mesh *data.Mesh) { +func AddSlonczewskiTorque2(torque, m *data.Slice, Msat, J, fixedP, alpha, pol, λ, ε_prime MSlice, thickness MSlice, flp float64, mesh *data.Mesh) { N := torque.Len() cfg := make1DConf(N) - flt := float32(flp * mesh.WorldSize()[Z]) + meshThickness := mesh.WorldSize()[Z] k_addslonczewskitorque2_async( torque.DevPtr(X), torque.DevPtr(Y), torque.DevPtr(Z), @@ -25,6 +23,8 @@ func AddSlonczewskiTorque2(torque, m *data.Slice, Msat, J, fixedP, alpha, pol, pol.DevPtr(0), pol.Mul(0), λ.DevPtr(0), λ.Mul(0), ε_prime.DevPtr(0), ε_prime.Mul(0), - unsafe.Pointer(uintptr(0)), flt, + thickness.DevPtr(0), thickness.Mul(0), + float32(meshThickness), + float32(flp), N, cfg) } diff --git a/cuda/slonczewski2.cu b/cuda/slonczewski2.cu index a7d969baa..f6548ee5f 100644 --- a/cuda/slonczewski2.cu +++ b/cuda/slonczewski2.cu @@ -9,16 +9,18 @@ extern "C" __global__ void addslonczewskitorque2(float* __restrict__ tx, float* __restrict__ ty, float* __restrict__ tz, float* __restrict__ mx, float* __restrict__ my, float* __restrict__ mz, - float* __restrict__ Ms_, float Ms_mul, - float* __restrict__ jz_, float jz_mul, - float* __restrict__ px_, float px_mul, - float* __restrict__ py_, float py_mul, - float* __restrict__ pz_, float pz_mul, - float* __restrict__ alpha_, float alpha_mul, - float* __restrict__ pol_, float pol_mul, - float* __restrict__ lambda_, float lambda_mul, - float* __restrict__ epsPrime_,float epsPrime_mul, - float* __restrict__ flt_, float flt_mul, + float* __restrict__ Ms_, float Ms_mul, + float* __restrict__ jz_, float jz_mul, + float* __restrict__ px_, float px_mul, + float* __restrict__ py_, float py_mul, + float* __restrict__ pz_, float pz_mul, + float* __restrict__ alpha_, float alpha_mul, + float* __restrict__ pol_, float pol_mul, + float* __restrict__ lambda_, float lambda_mul, + float* __restrict__ epsPrime_, float epsPrime_mul, + float* __restrict__ thickness_, float thickness_mul, + float meshThickness, + float freeLayerPosition, int N) { int i = ( blockIdx.y*gridDim.x + blockIdx.x ) * blockDim.x + threadIdx.x; @@ -29,16 +31,21 @@ addslonczewskitorque2(float* __restrict__ tx, float* __restrict__ ty, float* __r float3 p = normalized(vmul(px_, py_, pz_, px_mul, py_mul, pz_mul, i)); float Ms = amul(Ms_, Ms_mul, i); float alpha = amul(alpha_, alpha_mul, i); - float flt = amul(flt_, flt_mul, i); float pol = amul(pol_, pol_mul, i); float lambda = amul(lambda_, lambda_mul, i); float epsilonPrime = amul(epsPrime_, epsPrime_mul, i); + float thickness = amul(thickness_, thickness_mul, i); + if (thickness == 0.0) { // if thickness is not set, use the thickness of the mesh instead + thickness = meshThickness; + } + thickness *= freeLayerPosition; // switch sign if fixedlayer is at the bottom + if (J == 0.0f || Ms == 0.0f) { return; } - float beta = (HBAR / QE) * (J / (flt*Ms) ); + float beta = (HBAR / QE) * (J / (thickness*Ms) ); float lambda2 = lambda * lambda; float epsilon = pol * lambda2 / ((lambda2 + 1.0f) + (lambda2 - 1.0f) * dot(p, m)); diff --git a/cuda/slonczewski2_wrapper.go b/cuda/slonczewski2_wrapper.go index 389c6cec4..ade952068 100644 --- a/cuda/slonczewski2_wrapper.go +++ b/cuda/slonczewski2_wrapper.go @@ -17,34 +17,36 @@ var addslonczewskitorque2_code cu.Function // Stores the arguments for addslonczewskitorque2 kernel invocation type addslonczewskitorque2_args_t struct { - arg_tx unsafe.Pointer - arg_ty unsafe.Pointer - arg_tz unsafe.Pointer - arg_mx unsafe.Pointer - arg_my unsafe.Pointer - arg_mz unsafe.Pointer - arg_Ms_ unsafe.Pointer - arg_Ms_mul float32 - arg_jz_ unsafe.Pointer - arg_jz_mul float32 - arg_px_ unsafe.Pointer - arg_px_mul float32 - arg_py_ unsafe.Pointer - arg_py_mul float32 - arg_pz_ unsafe.Pointer - arg_pz_mul float32 - arg_alpha_ unsafe.Pointer - arg_alpha_mul float32 - arg_pol_ unsafe.Pointer - arg_pol_mul float32 - arg_lambda_ unsafe.Pointer - arg_lambda_mul float32 - arg_epsPrime_ unsafe.Pointer - arg_epsPrime_mul float32 - arg_flt_ unsafe.Pointer - arg_flt_mul float32 - arg_N int - argptr [27]unsafe.Pointer + arg_tx unsafe.Pointer + arg_ty unsafe.Pointer + arg_tz unsafe.Pointer + arg_mx unsafe.Pointer + arg_my unsafe.Pointer + arg_mz unsafe.Pointer + arg_Ms_ unsafe.Pointer + arg_Ms_mul float32 + arg_jz_ unsafe.Pointer + arg_jz_mul float32 + arg_px_ unsafe.Pointer + arg_px_mul float32 + arg_py_ unsafe.Pointer + arg_py_mul float32 + arg_pz_ unsafe.Pointer + arg_pz_mul float32 + arg_alpha_ unsafe.Pointer + arg_alpha_mul float32 + arg_pol_ unsafe.Pointer + arg_pol_mul float32 + arg_lambda_ unsafe.Pointer + arg_lambda_mul float32 + arg_epsPrime_ unsafe.Pointer + arg_epsPrime_mul float32 + arg_thickness_ unsafe.Pointer + arg_thickness_mul float32 + arg_meshThickness float32 + arg_freeLayerPosition float32 + arg_N int + argptr [29]unsafe.Pointer sync.Mutex } @@ -77,13 +79,15 @@ func init() { addslonczewskitorque2_args.argptr[21] = unsafe.Pointer(&addslonczewskitorque2_args.arg_lambda_mul) addslonczewskitorque2_args.argptr[22] = unsafe.Pointer(&addslonczewskitorque2_args.arg_epsPrime_) addslonczewskitorque2_args.argptr[23] = unsafe.Pointer(&addslonczewskitorque2_args.arg_epsPrime_mul) - addslonczewskitorque2_args.argptr[24] = unsafe.Pointer(&addslonczewskitorque2_args.arg_flt_) - addslonczewskitorque2_args.argptr[25] = unsafe.Pointer(&addslonczewskitorque2_args.arg_flt_mul) - addslonczewskitorque2_args.argptr[26] = unsafe.Pointer(&addslonczewskitorque2_args.arg_N) + addslonczewskitorque2_args.argptr[24] = unsafe.Pointer(&addslonczewskitorque2_args.arg_thickness_) + addslonczewskitorque2_args.argptr[25] = unsafe.Pointer(&addslonczewskitorque2_args.arg_thickness_mul) + addslonczewskitorque2_args.argptr[26] = unsafe.Pointer(&addslonczewskitorque2_args.arg_meshThickness) + addslonczewskitorque2_args.argptr[27] = unsafe.Pointer(&addslonczewskitorque2_args.arg_freeLayerPosition) + addslonczewskitorque2_args.argptr[28] = unsafe.Pointer(&addslonczewskitorque2_args.arg_N) } // Wrapper for addslonczewskitorque2 CUDA kernel, asynchronous. -func k_addslonczewskitorque2_async(tx unsafe.Pointer, ty unsafe.Pointer, tz unsafe.Pointer, mx unsafe.Pointer, my unsafe.Pointer, mz unsafe.Pointer, Ms_ unsafe.Pointer, Ms_mul float32, jz_ unsafe.Pointer, jz_mul float32, px_ unsafe.Pointer, px_mul float32, py_ unsafe.Pointer, py_mul float32, pz_ unsafe.Pointer, pz_mul float32, alpha_ unsafe.Pointer, alpha_mul float32, pol_ unsafe.Pointer, pol_mul float32, lambda_ unsafe.Pointer, lambda_mul float32, epsPrime_ unsafe.Pointer, epsPrime_mul float32, flt_ unsafe.Pointer, flt_mul float32, N int, cfg *config) { +func k_addslonczewskitorque2_async(tx unsafe.Pointer, ty unsafe.Pointer, tz unsafe.Pointer, mx unsafe.Pointer, my unsafe.Pointer, mz unsafe.Pointer, Ms_ unsafe.Pointer, Ms_mul float32, jz_ unsafe.Pointer, jz_mul float32, px_ unsafe.Pointer, px_mul float32, py_ unsafe.Pointer, py_mul float32, pz_ unsafe.Pointer, pz_mul float32, alpha_ unsafe.Pointer, alpha_mul float32, pol_ unsafe.Pointer, pol_mul float32, lambda_ unsafe.Pointer, lambda_mul float32, epsPrime_ unsafe.Pointer, epsPrime_mul float32, thickness_ unsafe.Pointer, thickness_mul float32, meshThickness float32, freeLayerPosition float32, N int, cfg *config) { if Synchronous { // debug Sync() timer.Start("addslonczewskitorque2") @@ -120,8 +124,10 @@ func k_addslonczewskitorque2_async(tx unsafe.Pointer, ty unsafe.Pointer, tz unsa addslonczewskitorque2_args.arg_lambda_mul = lambda_mul addslonczewskitorque2_args.arg_epsPrime_ = epsPrime_ addslonczewskitorque2_args.arg_epsPrime_mul = epsPrime_mul - addslonczewskitorque2_args.arg_flt_ = flt_ - addslonczewskitorque2_args.arg_flt_mul = flt_mul + addslonczewskitorque2_args.arg_thickness_ = thickness_ + addslonczewskitorque2_args.arg_thickness_mul = thickness_mul + addslonczewskitorque2_args.arg_meshThickness = meshThickness + addslonczewskitorque2_args.arg_freeLayerPosition = freeLayerPosition addslonczewskitorque2_args.arg_N = N args := addslonczewskitorque2_args.argptr[:] @@ -185,11 +191,13 @@ const ( .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -202,26 +210,28 @@ const ( ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -247,8 +257,8 @@ const ( cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -256,8 +266,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -265,8 +275,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -274,31 +284,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -306,113 +316,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; @@ -454,11 +467,13 @@ BB0_25: .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -471,26 +486,28 @@ BB0_25: ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -516,8 +533,8 @@ BB0_25: cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.nc.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.nc.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -525,8 +542,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.nc.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.nc.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -534,8 +551,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.nc.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.nc.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -543,31 +560,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.nc.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.nc.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.nc.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.nc.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -575,113 +592,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.nc.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.nc.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.nc.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.nc.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.nc.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.nc.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.nc.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.nc.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.nc.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.nc.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; @@ -723,11 +743,13 @@ BB0_25: .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -740,26 +762,28 @@ BB0_25: ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -785,8 +809,8 @@ BB0_25: cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.nc.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.nc.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -794,8 +818,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.nc.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.nc.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -803,8 +827,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.nc.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.nc.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -812,31 +836,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.nc.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.nc.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.nc.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.nc.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -844,113 +868,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.nc.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.nc.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.nc.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.nc.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.nc.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.nc.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.nc.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.nc.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.nc.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.nc.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; @@ -992,11 +1019,13 @@ BB0_25: .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -1009,26 +1038,28 @@ BB0_25: ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -1054,8 +1085,8 @@ BB0_25: cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.nc.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.nc.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -1063,8 +1094,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.nc.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.nc.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -1072,8 +1103,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.nc.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.nc.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -1081,31 +1112,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.nc.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.nc.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.nc.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.nc.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -1113,113 +1144,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.nc.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.nc.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.nc.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.nc.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.nc.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.nc.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.nc.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.nc.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.nc.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.nc.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; @@ -1261,11 +1295,13 @@ BB0_25: .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -1278,26 +1314,28 @@ BB0_25: ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -1323,8 +1361,8 @@ BB0_25: cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.nc.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.nc.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -1332,8 +1370,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.nc.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.nc.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -1341,8 +1379,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.nc.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.nc.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -1350,31 +1388,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.nc.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.nc.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.nc.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.nc.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -1382,113 +1420,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.nc.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.nc.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.nc.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.nc.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.nc.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.nc.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.nc.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.nc.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.nc.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.nc.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; @@ -1530,11 +1571,13 @@ BB0_25: .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -1547,26 +1590,28 @@ BB0_25: ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -1592,8 +1637,8 @@ BB0_25: cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.nc.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.nc.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -1601,8 +1646,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.nc.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.nc.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -1610,8 +1655,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.nc.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.nc.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -1619,31 +1664,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.nc.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.nc.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.nc.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.nc.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -1651,113 +1696,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.nc.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.nc.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.nc.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.nc.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.nc.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.nc.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.nc.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.nc.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.nc.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.nc.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; @@ -1799,11 +1847,13 @@ BB0_25: .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -1816,26 +1866,28 @@ BB0_25: ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -1861,8 +1913,8 @@ BB0_25: cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.nc.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.nc.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -1870,8 +1922,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.nc.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.nc.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -1879,8 +1931,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.nc.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.nc.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -1888,31 +1940,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.nc.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.nc.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.nc.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.nc.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -1920,113 +1972,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.nc.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.nc.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.nc.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.nc.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.nc.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.nc.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.nc.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.nc.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.nc.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.nc.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; @@ -2068,11 +2123,13 @@ BB0_25: .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -2085,26 +2142,28 @@ BB0_25: ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -2130,8 +2189,8 @@ BB0_25: cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.nc.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.nc.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -2139,8 +2198,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.nc.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.nc.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -2148,8 +2207,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.nc.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.nc.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -2157,31 +2216,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.nc.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.nc.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.nc.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.nc.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -2189,113 +2248,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.nc.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.nc.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.nc.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.nc.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.nc.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.nc.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.nc.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.nc.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.nc.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.nc.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; @@ -2337,11 +2399,13 @@ BB0_25: .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -2354,26 +2418,28 @@ BB0_25: ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -2399,8 +2465,8 @@ BB0_25: cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.nc.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.nc.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -2408,8 +2474,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.nc.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.nc.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -2417,8 +2483,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.nc.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.nc.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -2426,31 +2492,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.nc.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.nc.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.nc.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.nc.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -2458,113 +2524,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.nc.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.nc.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.nc.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.nc.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.nc.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.nc.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.nc.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.nc.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.nc.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.nc.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; @@ -2606,11 +2675,13 @@ BB0_25: .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -2623,26 +2694,28 @@ BB0_25: ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -2668,8 +2741,8 @@ BB0_25: cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.nc.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.nc.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -2677,8 +2750,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.nc.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.nc.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -2686,8 +2759,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.nc.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.nc.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -2695,31 +2768,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.nc.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.nc.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.nc.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.nc.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -2727,113 +2800,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.nc.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.nc.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.nc.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.nc.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.nc.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.nc.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.nc.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.nc.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.nc.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.nc.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; @@ -2875,11 +2951,13 @@ BB0_25: .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -2892,26 +2970,28 @@ BB0_25: ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -2937,8 +3017,8 @@ BB0_25: cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.nc.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.nc.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -2946,8 +3026,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.nc.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.nc.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -2955,8 +3035,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.nc.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.nc.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -2964,31 +3044,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.nc.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.nc.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.nc.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.nc.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -2996,113 +3076,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.nc.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.nc.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.nc.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.nc.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.nc.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.nc.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.nc.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.nc.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.nc.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.nc.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; @@ -3144,11 +3227,13 @@ BB0_25: .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -3161,26 +3246,28 @@ BB0_25: ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -3206,8 +3293,8 @@ BB0_25: cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.nc.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.nc.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -3215,8 +3302,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.nc.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.nc.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -3224,8 +3311,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.nc.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.nc.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -3233,31 +3320,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.nc.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.nc.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.nc.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.nc.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -3265,113 +3352,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.nc.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.nc.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.nc.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.nc.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.nc.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.nc.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.nc.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.nc.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.nc.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.nc.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; @@ -3413,11 +3503,13 @@ BB0_25: .param .f32 addslonczewskitorque2_param_23, .param .u64 addslonczewskitorque2_param_24, .param .f32 addslonczewskitorque2_param_25, - .param .u32 addslonczewskitorque2_param_26 + .param .f32 addslonczewskitorque2_param_26, + .param .f32 addslonczewskitorque2_param_27, + .param .u32 addslonczewskitorque2_param_28 ) { - .reg .pred %p<16>; - .reg .f32 %f<116>; + .reg .pred %p<17>; + .reg .f32 %f<120>; .reg .b32 %r<86>; .reg .f64 %fd<3>; .reg .b64 %rd<61>; @@ -3430,26 +3522,28 @@ BB0_25: ld.param.u64 %rd5, [addslonczewskitorque2_param_4]; ld.param.u64 %rd6, [addslonczewskitorque2_param_5]; ld.param.u64 %rd7, [addslonczewskitorque2_param_6]; - ld.param.f32 %f110, [addslonczewskitorque2_param_7]; + ld.param.f32 %f114, [addslonczewskitorque2_param_7]; ld.param.u64 %rd8, [addslonczewskitorque2_param_8]; - ld.param.f32 %f105, [addslonczewskitorque2_param_9]; + ld.param.f32 %f109, [addslonczewskitorque2_param_9]; ld.param.u64 %rd9, [addslonczewskitorque2_param_10]; - ld.param.f32 %f106, [addslonczewskitorque2_param_11]; + ld.param.f32 %f110, [addslonczewskitorque2_param_11]; ld.param.u64 %rd10, [addslonczewskitorque2_param_12]; - ld.param.f32 %f107, [addslonczewskitorque2_param_13]; + ld.param.f32 %f111, [addslonczewskitorque2_param_13]; ld.param.u64 %rd11, [addslonczewskitorque2_param_14]; - ld.param.f32 %f108, [addslonczewskitorque2_param_15]; + ld.param.f32 %f112, [addslonczewskitorque2_param_15]; ld.param.u64 %rd12, [addslonczewskitorque2_param_16]; - ld.param.f32 %f111, [addslonczewskitorque2_param_17]; + ld.param.f32 %f115, [addslonczewskitorque2_param_17]; ld.param.u64 %rd13, [addslonczewskitorque2_param_18]; - ld.param.f32 %f113, [addslonczewskitorque2_param_19]; + ld.param.f32 %f116, [addslonczewskitorque2_param_19]; ld.param.u64 %rd14, [addslonczewskitorque2_param_20]; - ld.param.f32 %f114, [addslonczewskitorque2_param_21]; + ld.param.f32 %f117, [addslonczewskitorque2_param_21]; ld.param.u64 %rd15, [addslonczewskitorque2_param_22]; - ld.param.f32 %f115, [addslonczewskitorque2_param_23]; + ld.param.f32 %f118, [addslonczewskitorque2_param_23]; ld.param.u64 %rd16, [addslonczewskitorque2_param_24]; - ld.param.f32 %f112, [addslonczewskitorque2_param_25]; - ld.param.u32 %r2, [addslonczewskitorque2_param_26]; + ld.param.f32 %f119, [addslonczewskitorque2_param_25]; + ld.param.f32 %f40, [addslonczewskitorque2_param_26]; + ld.param.f32 %f41, [addslonczewskitorque2_param_27]; + ld.param.u32 %r2, [addslonczewskitorque2_param_28]; mov.u32 %r3, %nctaid.x; mov.u32 %r4, %ctaid.y; mov.u32 %r5, %ctaid.x; @@ -3475,8 +3569,8 @@ BB0_25: cvta.to.global.u64 %rd24, %rd8; add.s64 %rd26, %rd24, %rd18; - ld.global.nc.f32 %f40, [%rd26]; - mul.f32 %f105, %f40, %f105; + ld.global.nc.f32 %f42, [%rd26]; + mul.f32 %f109, %f42, %f109; BB0_3: setp.eq.s64 %p3, %rd9, 0; @@ -3484,8 +3578,8 @@ BB0_3: cvta.to.global.u64 %rd27, %rd9; add.s64 %rd29, %rd27, %rd18; - ld.global.nc.f32 %f41, [%rd29]; - mul.f32 %f106, %f41, %f106; + ld.global.nc.f32 %f43, [%rd29]; + mul.f32 %f110, %f43, %f110; BB0_5: setp.eq.s64 %p4, %rd10, 0; @@ -3493,8 +3587,8 @@ BB0_5: cvta.to.global.u64 %rd30, %rd10; add.s64 %rd32, %rd30, %rd18; - ld.global.nc.f32 %f42, [%rd32]; - mul.f32 %f107, %f42, %f107; + ld.global.nc.f32 %f44, [%rd32]; + mul.f32 %f111, %f44, %f111; BB0_7: setp.eq.s64 %p5, %rd11, 0; @@ -3502,31 +3596,31 @@ BB0_7: cvta.to.global.u64 %rd33, %rd11; add.s64 %rd35, %rd33, %rd18; - ld.global.nc.f32 %f43, [%rd35]; - mul.f32 %f108, %f43, %f108; + ld.global.nc.f32 %f45, [%rd35]; + mul.f32 %f112, %f45, %f112; BB0_9: - mul.f32 %f45, %f107, %f107; - fma.rn.f32 %f46, %f106, %f106, %f45; - fma.rn.f32 %f47, %f108, %f108, %f46; - sqrt.rn.f32 %f12, %f47; - mov.f32 %f109, 0f00000000; + mul.f32 %f47, %f111, %f111; + fma.rn.f32 %f48, %f110, %f110, %f47; + fma.rn.f32 %f49, %f112, %f112, %f48; + sqrt.rn.f32 %f12, %f49; + mov.f32 %f113, 0f00000000; setp.eq.f32 %p6, %f12, 0f00000000; @%p6 bra BB0_11; - rcp.rn.f32 %f109, %f12; + rcp.rn.f32 %f113, %f12; BB0_11: - mul.f32 %f15, %f106, %f109; - mul.f32 %f16, %f107, %f109; - mul.f32 %f17, %f108, %f109; + mul.f32 %f15, %f110, %f113; + mul.f32 %f16, %f111, %f113; + mul.f32 %f17, %f112, %f113; setp.eq.s64 %p7, %rd7, 0; @%p7 bra BB0_13; cvta.to.global.u64 %rd36, %rd7; add.s64 %rd38, %rd36, %rd18; - ld.global.nc.f32 %f48, [%rd38]; - mul.f32 %f110, %f48, %f110; + ld.global.nc.f32 %f50, [%rd38]; + mul.f32 %f114, %f50, %f114; BB0_13: setp.eq.s64 %p8, %rd12, 0; @@ -3534,113 +3628,116 @@ BB0_13: cvta.to.global.u64 %rd39, %rd12; add.s64 %rd41, %rd39, %rd18; - ld.global.nc.f32 %f49, [%rd41]; - mul.f32 %f111, %f49, %f111; + ld.global.nc.f32 %f51, [%rd41]; + mul.f32 %f115, %f51, %f115; BB0_15: - setp.eq.s64 %p9, %rd16, 0; + setp.eq.s64 %p9, %rd13, 0; @%p9 bra BB0_17; - cvta.to.global.u64 %rd42, %rd16; + cvta.to.global.u64 %rd42, %rd13; add.s64 %rd44, %rd42, %rd18; - ld.global.nc.f32 %f50, [%rd44]; - mul.f32 %f112, %f50, %f112; + ld.global.nc.f32 %f52, [%rd44]; + mul.f32 %f116, %f52, %f116; BB0_17: - setp.eq.s64 %p10, %rd13, 0; + setp.eq.s64 %p10, %rd14, 0; @%p10 bra BB0_19; - cvta.to.global.u64 %rd45, %rd13; + cvta.to.global.u64 %rd45, %rd14; add.s64 %rd47, %rd45, %rd18; - ld.global.nc.f32 %f51, [%rd47]; - mul.f32 %f113, %f51, %f113; + ld.global.nc.f32 %f53, [%rd47]; + mul.f32 %f117, %f53, %f117; BB0_19: - setp.eq.s64 %p11, %rd14, 0; + setp.eq.s64 %p11, %rd15, 0; @%p11 bra BB0_21; - cvta.to.global.u64 %rd48, %rd14; + cvta.to.global.u64 %rd48, %rd15; add.s64 %rd50, %rd48, %rd18; - ld.global.nc.f32 %f52, [%rd50]; - mul.f32 %f114, %f52, %f114; + ld.global.nc.f32 %f54, [%rd50]; + mul.f32 %f118, %f54, %f118; BB0_21: - setp.eq.s64 %p12, %rd15, 0; + setp.eq.s64 %p12, %rd16, 0; @%p12 bra BB0_23; - cvta.to.global.u64 %rd51, %rd15; + cvta.to.global.u64 %rd51, %rd16; add.s64 %rd53, %rd51, %rd18; - ld.global.nc.f32 %f53, [%rd53]; - mul.f32 %f115, %f53, %f115; + ld.global.nc.f32 %f55, [%rd53]; + mul.f32 %f119, %f55, %f119; BB0_23: - setp.eq.f32 %p13, %f110, 0f00000000; - setp.eq.f32 %p14, %f105, 0f00000000; + setp.eq.f32 %p13, %f114, 0f00000000; + setp.eq.f32 %p14, %f109, 0f00000000; or.pred %p15, %p14, %p13; @%p15 bra BB0_25; - mul.f32 %f54, %f110, %f112; - div.rn.f32 %f55, %f105, %f54; - cvt.f64.f32 %fd1, %f55; + setp.eq.f32 %p16, %f119, 0f00000000; + selp.f32 %f56, %f40, %f119, %p16; + mul.f32 %f57, %f56, %f41; + mul.f32 %f58, %f114, %f57; + div.rn.f32 %f59, %f109, %f58; + cvt.f64.f32 %fd1, %f59; mul.f64 %fd2, %fd1, 0d3CC7B6EF14E9250C; - cvt.rn.f32.f64 %f56, %fd2; - mul.f32 %f57, %f114, %f114; - mul.f32 %f58, %f113, %f57; - add.f32 %f59, %f57, 0f3F800000; - add.f32 %f60, %f57, 0fBF800000; - mul.f32 %f61, %f2, %f16; - fma.rn.f32 %f62, %f1, %f15, %f61; - fma.rn.f32 %f63, %f3, %f17, %f62; - fma.rn.f32 %f64, %f63, %f60, %f59; - div.rn.f32 %f65, %f58, %f64; - mul.f32 %f66, %f65, %f56; - mul.f32 %f67, %f115, %f56; - fma.rn.f32 %f68, %f111, %f111, 0f3F800000; - rcp.rn.f32 %f69, %f68; - fma.rn.f32 %f70, %f111, %f67, %f66; - mul.f32 %f71, %f69, %f70; - mul.f32 %f72, %f111, %f66; - sub.f32 %f73, %f67, %f72; - mul.f32 %f74, %f69, %f73; - mul.f32 %f75, %f2, %f17; - mul.f32 %f76, %f3, %f16; - sub.f32 %f77, %f76, %f75; - mul.f32 %f78, %f3, %f15; - mul.f32 %f79, %f1, %f17; - sub.f32 %f80, %f79, %f78; - mul.f32 %f81, %f1, %f16; - mul.f32 %f82, %f2, %f15; - sub.f32 %f83, %f82, %f81; - mul.f32 %f84, %f2, %f83; - mul.f32 %f85, %f3, %f80; - sub.f32 %f86, %f84, %f85; - mul.f32 %f87, %f3, %f77; - mul.f32 %f88, %f1, %f83; - sub.f32 %f89, %f87, %f88; - mul.f32 %f90, %f1, %f80; - mul.f32 %f91, %f2, %f77; - sub.f32 %f92, %f90, %f91; - mul.f32 %f93, %f77, %f74; - fma.rn.f32 %f94, %f86, %f71, %f93; + cvt.rn.f32.f64 %f60, %fd2; + mul.f32 %f61, %f117, %f117; + mul.f32 %f62, %f116, %f61; + add.f32 %f63, %f61, 0f3F800000; + add.f32 %f64, %f61, 0fBF800000; + mul.f32 %f65, %f2, %f16; + fma.rn.f32 %f66, %f1, %f15, %f65; + fma.rn.f32 %f67, %f3, %f17, %f66; + fma.rn.f32 %f68, %f67, %f64, %f63; + div.rn.f32 %f69, %f62, %f68; + mul.f32 %f70, %f69, %f60; + mul.f32 %f71, %f118, %f60; + fma.rn.f32 %f72, %f115, %f115, 0f3F800000; + rcp.rn.f32 %f73, %f72; + fma.rn.f32 %f74, %f115, %f71, %f70; + mul.f32 %f75, %f73, %f74; + mul.f32 %f76, %f115, %f70; + sub.f32 %f77, %f71, %f76; + mul.f32 %f78, %f73, %f77; + mul.f32 %f79, %f2, %f17; + mul.f32 %f80, %f3, %f16; + sub.f32 %f81, %f80, %f79; + mul.f32 %f82, %f3, %f15; + mul.f32 %f83, %f1, %f17; + sub.f32 %f84, %f83, %f82; + mul.f32 %f85, %f1, %f16; + mul.f32 %f86, %f2, %f15; + sub.f32 %f87, %f86, %f85; + mul.f32 %f88, %f2, %f87; + mul.f32 %f89, %f3, %f84; + sub.f32 %f90, %f88, %f89; + mul.f32 %f91, %f3, %f81; + mul.f32 %f92, %f1, %f87; + sub.f32 %f93, %f91, %f92; + mul.f32 %f94, %f1, %f84; + mul.f32 %f95, %f2, %f81; + sub.f32 %f96, %f94, %f95; + mul.f32 %f97, %f81, %f78; + fma.rn.f32 %f98, %f90, %f75, %f97; cvta.to.global.u64 %rd54, %rd1; add.s64 %rd56, %rd54, %rd18; - ld.global.f32 %f95, [%rd56]; - add.f32 %f96, %f95, %f94; - st.global.f32 [%rd56], %f96; - mul.f32 %f97, %f80, %f74; - fma.rn.f32 %f98, %f89, %f71, %f97; + ld.global.f32 %f99, [%rd56]; + add.f32 %f100, %f99, %f98; + st.global.f32 [%rd56], %f100; + mul.f32 %f101, %f84, %f78; + fma.rn.f32 %f102, %f93, %f75, %f101; cvta.to.global.u64 %rd57, %rd2; add.s64 %rd58, %rd57, %rd18; - ld.global.f32 %f99, [%rd58]; - add.f32 %f100, %f99, %f98; - st.global.f32 [%rd58], %f100; - mul.f32 %f101, %f83, %f74; - fma.rn.f32 %f102, %f92, %f71, %f101; + ld.global.f32 %f103, [%rd58]; + add.f32 %f104, %f103, %f102; + st.global.f32 [%rd58], %f104; + mul.f32 %f105, %f87, %f78; + fma.rn.f32 %f106, %f96, %f75, %f105; cvta.to.global.u64 %rd59, %rd3; add.s64 %rd60, %rd59, %rd18; - ld.global.f32 %f103, [%rd60]; - add.f32 %f104, %f103, %f102; - st.global.f32 [%rd60], %f104; + ld.global.f32 %f107, [%rd60]; + add.f32 %f108, %f107, %f106; + st.global.f32 [%rd60], %f108; BB0_25: ret; diff --git a/doc/templates/api-template.html b/doc/templates/api-template.html index 6ac6bb526..36766235a 100644 --- a/doc/templates/api-template.html +++ b/doc/templates/api-template.html @@ -314,7 +314,7 @@

Slonczewski model


- {{range .FilterName "epsilonprime" "Lambda" "Pol" "xi" "J" "fixedlayer" "fixedlayerposition" "fixedlayer_top" "fixedlayer_bottom" "DisableSlonczewskiTorque" "DisableZhangLiTorque" }} {{template "entry" .}} {{end}} + {{range .FilterName "epsilonprime" "Lambda" "Pol" "xi" "J" "FreeLayerThickness" "fixedlayer" "fixedlayerposition" "fixedlayer_top" "fixedlayer_bottom" "DisableSlonczewskiTorque" "DisableZhangLiTorque" }} {{template "entry" .}} {{end}} diff --git a/engine/torque.go b/engine/torque.go index 12a26d7ed..9b1a18245 100644 --- a/engine/torque.go +++ b/engine/torque.go @@ -1,10 +1,11 @@ package engine import ( + "reflect" + "github.com/mumax/3/cuda" "github.com/mumax/3/data" "github.com/mumax/3/util" - "reflect" ) var ( @@ -14,6 +15,7 @@ var ( Lambda = NewScalarParam("Lambda", "", "Slonczewski Λ parameter") EpsilonPrime = NewScalarParam("EpsilonPrime", "", "Slonczewski secondairy STT term ε'") FrozenSpins = NewScalarParam("frozenspins", "", "Defines spins that should be fixed") // 1 - frozen, 0 - free. TODO: check if it only contains 0/1 values + FreeLayerThickness = NewScalarParam("FreeLayerThickness", "m", "Slonczewski free layer thickness (if set to zero (default), then the thickness will be deduced from the mesh size)") FixedLayer = NewExcitation("FixedLayer", "", "Slonczewski fixed layer polarization") Torque = NewVectorField("torque", "T", "Total torque/γ0", SetTorque) LLTorque = NewVectorField("LLtorque", "T", "Landau-Lifshitz torque/γ0", SetLLTorque) @@ -100,8 +102,11 @@ func AddSTTorque(dst *data.Slice) { defer lambda.Recycle() epsPrime := EpsilonPrime.MSlice() defer epsPrime.Recycle() + thickness := FreeLayerThickness.MSlice() + defer thickness.Recycle() cuda.AddSlonczewskiTorque2(dst, M.Buffer(), msat, j, fixedP, alpha, pol, lambda, epsPrime, + thickness, CurrentSignFromFixedLayerPosition[fixedLayerPosition], Mesh()) } diff --git a/test/freelayerthickness.mx3 b/test/freelayerthickness.mx3 new file mode 100644 index 000000000..3e13350af --- /dev/null +++ b/test/freelayerthickness.mx3 @@ -0,0 +1,38 @@ +/* + Test the FreeLayerThickness for the Slonczewski torque +*/ + +SetGridSize(1,1,2) +SetCellSize(1e-9,1e-9,1e-9) + +// Arbitrary material parameters which result in a non zero Slonczewski Torque +Msat = 4.3 +J = vector(0,0,5.4) +Pol = 0.6 +FixedLayer = vector(1,0,1) + +// If the freeLayerThickness is zero (the default), then the thickness of the free +// layer will be deduced from the mesh thickness (1e-9m in this case) +FreeLayerThickness = 0 + +// Computation of the Slonczewski torque in the bottom cell +torqueWanted := Vector( + STTorque.hostCopy().Get(0,0,0,0), + STTorque.hostCopy().Get(1,0,0,0), + STTorque.hostCopy().Get(2,0,0,0)) + + + +// Adding empty layers to the system should not change the result +// if we set the freeLayerThickness by hand. +Setgridsize(1,1,4) +SetGeom(Layers(0,2)) +FreeLayerThickness = 1e-9 + +// Computation of the Slonczewski torque in the bottom cell +torqueResult:= Vector( + STTorque.hostCopy().Get(0,0,0,0), + STTorque.hostCopy().Get(1,0,0,0), + STTorque.hostCopy().Get(2,0,0,0)) + +expectv("sttorque",torqueResult,torqueWanted,1e-5 ) \ No newline at end of file