exchange_wrapper.go

package cuda

/*
 THIS FILE IS AUTO-GENERATED BY CUDA2GO.
 EDITING IS FUTILE.
*/

import (
	"github.com/kuchkin/mumax3-gneb/cuda/cu"
	"github.com/kuchkin/mumax3-gneb/timer"
	"sync"
	"unsafe"
)

// CUDA handle for addexchange kernel
var addexchange_code cu.Function

// Stores the arguments for addexchange kernel invocation
type addexchange_args_t struct {
	arg_Bx      unsafe.Pointer
	arg_By      unsafe.Pointer
	arg_Bz      unsafe.Pointer
	arg_mx      unsafe.Pointer
	arg_my      unsafe.Pointer
	arg_mz      unsafe.Pointer
	arg_Ms_     unsafe.Pointer
	arg_Ms_mul  float32
	arg_aLUT2d  unsafe.Pointer
	arg_regions unsafe.Pointer
	arg_wx      float32
	arg_wy      float32
	arg_wz      float32
	arg_Nx      int
	arg_Ny      int
	arg_Nz      int
	arg_PBC     byte
	arg_JZ      float32
	argptr      [18]unsafe.Pointer
	sync.Mutex
}

// Stores the arguments for addexchange kernel invocation
var addexchange_args addexchange_args_t

func init() {
	// CUDA driver kernel call wants pointers to arguments, set them up once.
	addexchange_args.argptr[0] = unsafe.Pointer(&addexchange_args.arg_Bx)
	addexchange_args.argptr[1] = unsafe.Pointer(&addexchange_args.arg_By)
	addexchange_args.argptr[2] = unsafe.Pointer(&addexchange_args.arg_Bz)
	addexchange_args.argptr[3] = unsafe.Pointer(&addexchange_args.arg_mx)
	addexchange_args.argptr[4] = unsafe.Pointer(&addexchange_args.arg_my)
	addexchange_args.argptr[5] = unsafe.Pointer(&addexchange_args.arg_mz)
	addexchange_args.argptr[6] = unsafe.Pointer(&addexchange_args.arg_Ms_)
	addexchange_args.argptr[7] = unsafe.Pointer(&addexchange_args.arg_Ms_mul)
	addexchange_args.argptr[8] = unsafe.Pointer(&addexchange_args.arg_aLUT2d)
	addexchange_args.argptr[9] = unsafe.Pointer(&addexchange_args.arg_regions)
	addexchange_args.argptr[10] = unsafe.Pointer(&addexchange_args.arg_wx)
	addexchange_args.argptr[11] = unsafe.Pointer(&addexchange_args.arg_wy)
	addexchange_args.argptr[12] = unsafe.Pointer(&addexchange_args.arg_wz)
	addexchange_args.argptr[13] = unsafe.Pointer(&addexchange_args.arg_Nx)
	addexchange_args.argptr[14] = unsafe.Pointer(&addexchange_args.arg_Ny)
	addexchange_args.argptr[15] = unsafe.Pointer(&addexchange_args.arg_Nz)
	addexchange_args.argptr[16] = unsafe.Pointer(&addexchange_args.arg_PBC)
	addexchange_args.argptr[17] = unsafe.Pointer(&addexchange_args.arg_JZ)
}

// Wrapper for addexchange CUDA kernel, asynchronous.
func k_addexchange_async(Bx unsafe.Pointer, By unsafe.Pointer, Bz unsafe.Pointer, mx unsafe.Pointer, my unsafe.Pointer, mz unsafe.Pointer, Ms_ unsafe.Pointer, Ms_mul float32, aLUT2d unsafe.Pointer, regions unsafe.Pointer, wx float32, wy float32, wz float32, Nx int, Ny int, Nz int, PBC byte, JZ float32, cfg *config) {
	if Synchronous { // debug
		Sync()
		timer.Start("addexchange")
	}

	addexchange_args.Lock()
	defer addexchange_args.Unlock()

	if addexchange_code == 0 {
		addexchange_code = fatbinLoad(addexchange_map, "addexchange")
	}

	addexchange_args.arg_Bx = Bx
	addexchange_args.arg_By = By
	addexchange_args.arg_Bz = Bz
	addexchange_args.arg_mx = mx
	addexchange_args.arg_my = my
	addexchange_args.arg_mz = mz
	addexchange_args.arg_Ms_ = Ms_
	addexchange_args.arg_Ms_mul = Ms_mul
	addexchange_args.arg_aLUT2d = aLUT2d
	addexchange_args.arg_regions = regions
	addexchange_args.arg_wx = wx
	addexchange_args.arg_wy = wy
	addexchange_args.arg_wz = wz
	addexchange_args.arg_Nx = Nx
	addexchange_args.arg_Ny = Ny
	addexchange_args.arg_Nz = Nz
	addexchange_args.arg_PBC = PBC
	addexchange_args.arg_JZ = JZ

	args := addexchange_args.argptr[:]
	cu.LaunchKernel(addexchange_code, cfg.Grid.X, cfg.Grid.Y, cfg.Grid.Z, cfg.Block.X, cfg.Block.Y, cfg.Block.Z, 0, stream0, args)

	if Synchronous { // debug
		Sync()
		timer.Stop("addexchange")
	}
}

// maps compute capability on PTX code for addexchange kernel.
var addexchange_map = map[int]string{0: "",
	50: addexchange_ptx_50}

// addexchange PTX code for various compute capabilities.
const (
	addexchange_ptx_50 = `
.version 7.5
.target sm_50
.address_size 64

	// .globl	addexchange

.visible .entry addexchange(
	.param .u64 addexchange_param_0,
	.param .u64 addexchange_param_1,
	.param .u64 addexchange_param_2,
	.param .u64 addexchange_param_3,
	.param .u64 addexchange_param_4,
	.param .u64 addexchange_param_5,
	.param .u64 addexchange_param_6,
	.param .f32 addexchange_param_7,
	.param .u64 addexchange_param_8,
	.param .u64 addexchange_param_9,
	.param .f32 addexchange_param_10,
	.param .f32 addexchange_param_11,
	.param .f32 addexchange_param_12,
	.param .u32 addexchange_param_13,
	.param .u32 addexchange_param_14,
	.param .u32 addexchange_param_15,
	.param .u8 addexchange_param_16,
	.param .f32 addexchange_param_17
)
{
	.reg .pred 	%p<22>;
	.reg .b16 	%rs<36>;
	.reg .f32 	%f<138>;
	.reg .b32 	%r<111>;
	.reg .b64 	%rd<79>;


	ld.param.u8 	%rs5, [addexchange_param_16];
	ld.param.u64 	%rd6, [addexchange_param_0];
	ld.param.u64 	%rd7, [addexchange_param_1];
	ld.param.u64 	%rd8, [addexchange_param_2];
	ld.param.u64 	%rd10, [addexchange_param_3];
	ld.param.u64 	%rd11, [addexchange_param_4];
	ld.param.u64 	%rd12, [addexchange_param_5];
	ld.param.u64 	%rd9, [addexchange_param_6];
	ld.param.f32 	%f136, [addexchange_param_7];
	ld.param.u64 	%rd13, [addexchange_param_8];
	ld.param.u64 	%rd14, [addexchange_param_9];
	ld.param.f32 	%f34, [addexchange_param_10];
	ld.param.f32 	%f35, [addexchange_param_11];
	ld.param.f32 	%f36, [addexchange_param_12];
	ld.param.u32 	%r31, [addexchange_param_13];
	ld.param.u32 	%r32, [addexchange_param_14];
	ld.param.u32 	%r33, [addexchange_param_15];
	ld.param.f32 	%f37, [addexchange_param_17];
	cvta.to.global.u64 	%rd1, %rd13;
	cvta.to.global.u64 	%rd2, %rd14;
	cvta.to.global.u64 	%rd3, %rd12;
	cvta.to.global.u64 	%rd4, %rd11;
	cvta.to.global.u64 	%rd5, %rd10;
	mov.u32 	%r34, %ntid.x;
	mov.u32 	%r35, %ctaid.x;
	mov.u32 	%r36, %tid.x;
	mad.lo.s32 	%r1, %r35, %r34, %r36;
	mov.u32 	%r37, %ntid.y;
	mov.u32 	%r38, %ctaid.y;
	mov.u32 	%r39, %tid.y;
	mad.lo.s32 	%r2, %r38, %r37, %r39;
	mov.u32 	%r40, %ntid.z;
	mov.u32 	%r41, %ctaid.z;
	mov.u32 	%r42, %tid.z;
	mad.lo.s32 	%r3, %r41, %r40, %r42;
	setp.ge.s32 	%p1, %r1, %r31;
	setp.ge.s32 	%p2, %r2, %r32;
	or.pred  	%p3, %p1, %p2;
	setp.ge.s32 	%p4, %r3, %r33;
	or.pred  	%p5, %p3, %p4;
	@%p5 bra 	$L__BB0_27;

	mul.lo.s32 	%r4, %r3, %r32;
	add.s32 	%r43, %r4, %r2;
	mul.lo.s32 	%r5, %r43, %r31;
	add.s32 	%r6, %r5, %r1;
	mul.wide.s32 	%rd15, %r6, 4;
	add.s64 	%rd16, %rd5, %rd15;
	add.s64 	%rd17, %rd4, %rd15;
	add.s64 	%rd18, %rd3, %rd15;
	ld.global.nc.f32 	%f1, [%rd16];
	ld.global.nc.f32 	%f2, [%rd17];
	ld.global.nc.f32 	%f3, [%rd18];
	mul.f32 	%f38, %f2, %f2;
	fma.rn.f32 	%f39, %f1, %f1, %f38;
	fma.rn.f32 	%f40, %f3, %f3, %f39;
	setp.eq.f32 	%p6, %f40, 0f00000000;
	@%p6 bra 	$L__BB0_27;

	cvt.s64.s32 	%rd19, %r6;
	add.s64 	%rd20, %rd2, %rd19;
	ld.global.nc.u8 	%rs1, [%rd20];
	and.b16  	%rs2, %rs5, 1;
	setp.eq.s16 	%p7, %rs2, 0;
	add.s32 	%r7, %r1, -1;
	@%p7 bra 	$L__BB0_4;
	bra.uni 	$L__BB0_3;

$L__BB0_4:
	max.s32 	%r105, %r7, 0;
	bra.uni 	$L__BB0_5;

$L__BB0_3:
	rem.s32 	%r44, %r7, %r31;
	add.s32 	%r45, %r44, %r31;
	rem.s32 	%r105, %r45, %r31;

$L__BB0_5:
	add.s32 	%r46, %r105, %r5;
	cvt.s64.s32 	%rd21, %r46;
	mul.wide.s32 	%rd22, %r46, 4;
	add.s64 	%rd23, %rd5, %rd22;
	add.s64 	%rd24, %rd4, %rd22;
	add.s64 	%rd25, %rd3, %rd22;
	ld.global.nc.f32 	%f41, [%rd25];
	ld.global.nc.f32 	%f42, [%rd23];
	ld.global.nc.f32 	%f43, [%rd24];
	mul.f32 	%f44, %f43, %f43;
	fma.rn.f32 	%f45, %f42, %f42, %f44;
	fma.rn.f32 	%f46, %f41, %f41, %f45;
	setp.eq.f32 	%p8, %f46, 0f00000000;
	selp.f32 	%f47, %f3, %f41, %p8;
	selp.f32 	%f48, %f2, %f43, %p8;
	selp.f32 	%f49, %f1, %f42, %p8;
	add.s64 	%rd26, %rd2, %rd21;
	ld.global.nc.u8 	%rs6, [%rd26];
	min.u16 	%rs9, %rs6, %rs1;
	cvt.u32.u16 	%r47, %rs9;
	max.u16 	%rs10, %rs6, %rs1;
	cvt.u32.u16 	%r48, %rs10;
	add.s32 	%r49, %r48, 1;
	mul.lo.s32 	%r50, %r49, %r48;
	shr.u32 	%r51, %r50, 1;
	add.s32 	%r52, %r51, %r47;
	mul.wide.s32 	%rd27, %r52, 4;
	add.s64 	%rd28, %rd1, %rd27;
	ld.global.nc.f32 	%f50, [%rd28];
	mul.f32 	%f51, %f50, %f34;
	sub.f32 	%f52, %f49, %f1;
	sub.f32 	%f53, %f48, %f2;
	sub.f32 	%f54, %f47, %f3;
	fma.rn.f32 	%f7, %f51, %f52, 0f00000000;
	fma.rn.f32 	%f8, %f51, %f53, 0f00000000;
	fma.rn.f32 	%f9, %f51, %f54, 0f00000000;
	add.s32 	%r11, %r1, 1;
	@%p7 bra 	$L__BB0_7;
	bra.uni 	$L__BB0_6;

$L__BB0_7:
	add.s32 	%r55, %r31, -1;
	min.s32 	%r106, %r11, %r55;
	bra.uni 	$L__BB0_8;

$L__BB0_6:
	rem.s32 	%r53, %r11, %r31;
	add.s32 	%r54, %r53, %r31;
	rem.s32 	%r106, %r54, %r31;

$L__BB0_8:
	add.s32 	%r56, %r106, %r5;
	cvt.s64.s32 	%rd29, %r56;
	mul.wide.s32 	%rd30, %r56, 4;
	add.s64 	%rd31, %rd5, %rd30;
	add.s64 	%rd32, %rd4, %rd30;
	add.s64 	%rd33, %rd3, %rd30;
	ld.global.nc.f32 	%f55, [%rd33];
	ld.global.nc.f32 	%f56, [%rd31];
	ld.global.nc.f32 	%f57, [%rd32];
	mul.f32 	%f58, %f57, %f57;
	fma.rn.f32 	%f59, %f56, %f56, %f58;
	fma.rn.f32 	%f60, %f55, %f55, %f59;
	setp.eq.f32 	%p10, %f60, 0f00000000;
	selp.f32 	%f61, %f3, %f55, %p10;
	selp.f32 	%f62, %f2, %f57, %p10;
	selp.f32 	%f63, %f1, %f56, %p10;
	add.s64 	%rd34, %rd2, %rd29;
	ld.global.nc.u8 	%rs11, [%rd34];
	min.u16 	%rs14, %rs11, %rs1;
	cvt.u32.u16 	%r57, %rs14;
	max.u16 	%rs15, %rs11, %rs1;
	cvt.u32.u16 	%r58, %rs15;
	add.s32 	%r59, %r58, 1;
	mul.lo.s32 	%r60, %r59, %r58;
	shr.u32 	%r61, %r60, 1;
	add.s32 	%r62, %r61, %r57;
	mul.wide.s32 	%rd35, %r62, 4;
	add.s64 	%rd36, %rd1, %rd35;
	ld.global.nc.f32 	%f64, [%rd36];
	mul.f32 	%f65, %f64, %f34;
	sub.f32 	%f66, %f63, %f1;
	sub.f32 	%f67, %f62, %f2;
	sub.f32 	%f68, %f61, %f3;
	fma.rn.f32 	%f10, %f65, %f66, %f7;
	fma.rn.f32 	%f11, %f65, %f67, %f8;
	fma.rn.f32 	%f12, %f65, %f68, %f9;
	and.b16  	%rs3, %rs5, 2;
	setp.eq.s16 	%p11, %rs3, 0;
	add.s32 	%r15, %r2, -1;
	@%p11 bra 	$L__BB0_10;
	bra.uni 	$L__BB0_9;

$L__BB0_10:
	max.s32 	%r107, %r15, 0;
	bra.uni 	$L__BB0_11;

$L__BB0_9:
	rem.s32 	%r63, %r15, %r32;
	add.s32 	%r64, %r63, %r32;
	rem.s32 	%r107, %r64, %r32;

$L__BB0_11:
	add.s32 	%r65, %r107, %r4;
	mad.lo.s32 	%r66, %r65, %r31, %r1;
	cvt.s64.s32 	%rd37, %r66;
	mul.wide.s32 	%rd38, %r66, 4;
	add.s64 	%rd39, %rd5, %rd38;
	add.s64 	%rd40, %rd4, %rd38;
	add.s64 	%rd41, %rd3, %rd38;
	ld.global.nc.f32 	%f69, [%rd41];
	ld.global.nc.f32 	%f70, [%rd39];
	ld.global.nc.f32 	%f71, [%rd40];
	mul.f32 	%f72, %f71, %f71;
	fma.rn.f32 	%f73, %f70, %f70, %f72;
	fma.rn.f32 	%f74, %f69, %f69, %f73;
	setp.eq.f32 	%p12, %f74, 0f00000000;
	selp.f32 	%f75, %f3, %f69, %p12;
	selp.f32 	%f76, %f2, %f71, %p12;
	selp.f32 	%f77, %f1, %f70, %p12;
	add.s64 	%rd42, %rd2, %rd37;
	ld.global.nc.u8 	%rs16, [%rd42];
	min.u16 	%rs19, %rs16, %rs1;
	cvt.u32.u16 	%r67, %rs19;
	max.u16 	%rs20, %rs16, %rs1;
	cvt.u32.u16 	%r68, %rs20;
	add.s32 	%r69, %r68, 1;
	mul.lo.s32 	%r70, %r69, %r68;
	shr.u32 	%r71, %r70, 1;
	add.s32 	%r72, %r71, %r67;
	mul.wide.s32 	%rd43, %r72, 4;
	add.s64 	%rd44, %rd1, %rd43;
	ld.global.nc.f32 	%f78, [%rd44];
	mul.f32 	%f79, %f78, %f35;
	sub.f32 	%f80, %f77, %f1;
	sub.f32 	%f81, %f76, %f2;
	sub.f32 	%f82, %f75, %f3;
	fma.rn.f32 	%f13, %f79, %f80, %f10;
	fma.rn.f32 	%f14, %f79, %f81, %f11;
	fma.rn.f32 	%f15, %f79, %f82, %f12;
	add.s32 	%r19, %r2, 1;
	@%p11 bra 	$L__BB0_13;
	bra.uni 	$L__BB0_12;

$L__BB0_13:
	add.s32 	%r75, %r32, -1;
	min.s32 	%r108, %r19, %r75;
	bra.uni 	$L__BB0_14;

$L__BB0_12:
	rem.s32 	%r73, %r19, %r32;
	add.s32 	%r74, %r73, %r32;
	rem.s32 	%r108, %r74, %r32;

$L__BB0_14:
	add.s32 	%r76, %r108, %r4;
	mad.lo.s32 	%r77, %r76, %r31, %r1;
	cvt.s64.s32 	%rd45, %r77;
	mul.wide.s32 	%rd46, %r77, 4;
	add.s64 	%rd47, %rd5, %rd46;
	add.s64 	%rd48, %rd4, %rd46;
	add.s64 	%rd49, %rd3, %rd46;
	ld.global.nc.f32 	%f83, [%rd49];
	ld.global.nc.f32 	%f84, [%rd47];
	ld.global.nc.f32 	%f85, [%rd48];
	mul.f32 	%f86, %f85, %f85;
	fma.rn.f32 	%f87, %f84, %f84, %f86;
	fma.rn.f32 	%f88, %f83, %f83, %f87;
	setp.eq.f32 	%p14, %f88, 0f00000000;
	selp.f32 	%f89, %f3, %f83, %p14;
	selp.f32 	%f90, %f2, %f85, %p14;
	selp.f32 	%f91, %f1, %f84, %p14;
	add.s64 	%rd50, %rd2, %rd45;
	ld.global.nc.u8 	%rs21, [%rd50];
	min.u16 	%rs24, %rs21, %rs1;
	cvt.u32.u16 	%r78, %rs24;
	max.u16 	%rs25, %rs21, %rs1;
	cvt.u32.u16 	%r79, %rs25;
	add.s32 	%r80, %r79, 1;
	mul.lo.s32 	%r81, %r80, %r79;
	shr.u32 	%r82, %r81, 1;
	add.s32 	%r83, %r82, %r78;
	mul.wide.s32 	%rd51, %r83, 4;
	add.s64 	%rd52, %rd1, %rd51;
	ld.global.nc.f32 	%f92, [%rd52];
	mul.f32 	%f93, %f92, %f35;
	sub.f32 	%f94, %f91, %f1;
	sub.f32 	%f95, %f90, %f2;
	sub.f32 	%f96, %f89, %f3;
	fma.rn.f32 	%f135, %f93, %f94, %f13;
	fma.rn.f32 	%f134, %f93, %f95, %f14;
	fma.rn.f32 	%f133, %f93, %f96, %f15;
	setp.eq.s32 	%p15, %r33, 1;
	@%p15 bra 	$L__BB0_22;

	and.b16  	%rs4, %rs5, 4;
	setp.eq.s16 	%p16, %rs4, 0;
	add.s32 	%r23, %r3, -1;
	@%p16 bra 	$L__BB0_17;
	bra.uni 	$L__BB0_16;

$L__BB0_17:
	max.s32 	%r109, %r23, 0;
	bra.uni 	$L__BB0_18;

$L__BB0_16:
	rem.s32 	%r84, %r23, %r33;
	add.s32 	%r85, %r84, %r33;
	rem.s32 	%r109, %r85, %r33;

$L__BB0_18:
	mad.lo.s32 	%r86, %r109, %r32, %r2;
	mad.lo.s32 	%r87, %r86, %r31, %r1;
	cvt.s64.s32 	%rd53, %r87;
	mul.wide.s32 	%rd54, %r87, 4;
	add.s64 	%rd55, %rd5, %rd54;
	add.s64 	%rd56, %rd4, %rd54;
	add.s64 	%rd57, %rd3, %rd54;
	ld.global.nc.f32 	%f97, [%rd57];
	ld.global.nc.f32 	%f98, [%rd55];
	ld.global.nc.f32 	%f99, [%rd56];
	mul.f32 	%f100, %f99, %f99;
	fma.rn.f32 	%f101, %f98, %f98, %f100;
	fma.rn.f32 	%f102, %f97, %f97, %f101;
	setp.eq.f32 	%p17, %f102, 0f00000000;
	selp.f32 	%f103, %f3, %f97, %p17;
	selp.f32 	%f104, %f2, %f99, %p17;
	selp.f32 	%f105, %f1, %f98, %p17;
	add.s64 	%rd58, %rd2, %rd53;
	ld.global.nc.u8 	%rs26, [%rd58];
	min.u16 	%rs29, %rs26, %rs1;
	cvt.u32.u16 	%r88, %rs29;
	max.u16 	%rs30, %rs26, %rs1;
	cvt.u32.u16 	%r89, %rs30;
	add.s32 	%r90, %r89, 1;
	mul.lo.s32 	%r91, %r90, %r89;
	shr.u32 	%r92, %r91, 1;
	add.s32 	%r93, %r92, %r88;
	mul.wide.s32 	%rd59, %r93, 4;
	add.s64 	%rd60, %rd1, %rd59;
	mul.f32 	%f19, %f36, %f37;
	ld.global.nc.f32 	%f106, [%rd60];
	mul.f32 	%f107, %f19, %f106;
	sub.f32 	%f108, %f105, %f1;
	sub.f32 	%f109, %f104, %f2;
	sub.f32 	%f110, %f103, %f3;
	fma.rn.f32 	%f20, %f107, %f108, %f135;
	fma.rn.f32 	%f21, %f107, %f109, %f134;
	fma.rn.f32 	%f22, %f107, %f110, %f133;
	add.s32 	%r27, %r3, 1;
	@%p16 bra 	$L__BB0_20;
	bra.uni 	$L__BB0_19;

$L__BB0_20:
	add.s32 	%r96, %r33, -1;
	min.s32 	%r110, %r27, %r96;
	bra.uni 	$L__BB0_21;

$L__BB0_19:
	rem.s32 	%r94, %r27, %r33;
	add.s32 	%r95, %r94, %r33;
	rem.s32 	%r110, %r95, %r33;

$L__BB0_21:
	mad.lo.s32 	%r97, %r110, %r32, %r2;
	mad.lo.s32 	%r98, %r97, %r31, %r1;
	cvt.s64.s32 	%rd61, %r98;
	mul.wide.s32 	%rd62, %r98, 4;
	add.s64 	%rd63, %rd5, %rd62;
	add.s64 	%rd64, %rd4, %rd62;
	add.s64 	%rd65, %rd3, %rd62;
	ld.global.nc.f32 	%f111, [%rd65];
	ld.global.nc.f32 	%f112, [%rd63];
	ld.global.nc.f32 	%f113, [%rd64];
	mul.f32 	%f114, %f113, %f113;
	fma.rn.f32 	%f115, %f112, %f112, %f114;
	fma.rn.f32 	%f116, %f111, %f111, %f115;
	setp.eq.f32 	%p19, %f116, 0f00000000;
	selp.f32 	%f117, %f3, %f111, %p19;
	selp.f32 	%f118, %f2, %f113, %p19;
	selp.f32 	%f119, %f1, %f112, %p19;
	add.s64 	%rd66, %rd2, %rd61;
	ld.global.nc.u8 	%rs31, [%rd66];
	min.u16 	%rs34, %rs31, %rs1;
	cvt.u32.u16 	%r99, %rs34;
	max.u16 	%rs35, %rs31, %rs1;
	cvt.u32.u16 	%r100, %rs35;
	add.s32 	%r101, %r100, 1;
	mul.lo.s32 	%r102, %r101, %r100;
	shr.u32 	%r103, %r102, 1;
	add.s32 	%r104, %r103, %r99;
	mul.wide.s32 	%rd67, %r104, 4;
	add.s64 	%rd68, %rd1, %rd67;
	ld.global.nc.f32 	%f120, [%rd68];
	mul.f32 	%f121, %f19, %f120;
	sub.f32 	%f122, %f119, %f1;
	sub.f32 	%f123, %f118, %f2;
	sub.f32 	%f124, %f117, %f3;
	fma.rn.f32 	%f135, %f121, %f122, %f20;
	fma.rn.f32 	%f134, %f121, %f123, %f21;
	fma.rn.f32 	%f133, %f121, %f124, %f22;

$L__BB0_22:
	setp.eq.s64 	%p20, %rd9, 0;
	@%p20 bra 	$L__BB0_24;

	cvta.to.global.u64 	%rd69, %rd9;
	add.s64 	%rd71, %rd69, %rd15;
	ld.global.nc.f32 	%f125, [%rd71];
	mul.f32 	%f136, %f125, %f136;

$L__BB0_24:
	setp.eq.f32 	%p21, %f136, 0f00000000;
	mov.f32 	%f137, 0f00000000;
	@%p21 bra 	$L__BB0_26;

	rcp.rn.f32 	%f137, %f136;

$L__BB0_26:
	cvta.to.global.u64 	%rd72, %rd6;
	add.s64 	%rd74, %rd72, %rd15;
	ld.global.f32 	%f127, [%rd74];
	fma.rn.f32 	%f128, %f135, %f137, %f127;
	st.global.f32 	[%rd74], %f128;
	cvta.to.global.u64 	%rd75, %rd7;
	add.s64 	%rd76, %rd75, %rd15;
	ld.global.f32 	%f129, [%rd76];
	fma.rn.f32 	%f130, %f134, %f137, %f129;
	st.global.f32 	[%rd76], %f130;
	cvta.to.global.u64 	%rd77, %rd8;
	add.s64 	%rd78, %rd77, %rd15;
	ld.global.f32 	%f131, [%rd78];
	fma.rn.f32 	%f132, %f133, %f137, %f131;
	st.global.f32 	[%rd78], %f132;

$L__BB0_27:
	ret;

}

`
)