//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Wed Jul 10 12:41:20 2013 (1373485280)
// Cuda compilation tools, release 5.5, V5.5.0
//

.version 3.2
.target sm_30
.address_size 64

	.file	1 "C:/ame2014Iupdate/releases/2014.03/external/adobe/MediaCore/GPUFoundation/Src/ImageProcessing/SingleChannelBlur.cu", 1405710183, 12690
	.file	2 "c:\\ame2014iupdate\\releases\\2014.03\\external\\adobe\\mediacore\\external\\3rdparty\\nvidia\\cuda\\win\\include\\device_functions.h", 1405710127, 191626
.const .align 4 .b8 kRGB32f_To_601YPbPr[36] = {135, 22, 153, 62, 162, 69, 22, 63, 213, 120, 233, 61, 33, 201, 44, 190, 111, 155, 169, 190, 0, 0, 0, 63, 0, 0, 0, 63, 70, 94, 214, 190, 232, 134, 166, 189};
.const .align 4 .b8 k601YPbPr_To_RGB32f[36] = {0, 0, 128, 63, 0, 0, 0, 0, 188, 116, 179, 63, 0, 0, 128, 63, 152, 50, 176, 190, 158, 209, 54, 191, 0, 0, 128, 63, 229, 208, 226, 63, 0, 0, 0, 0};
.const .align 4 .b8 kRGB32f_To_601YCbCr[36] = {70, 246, 130, 66, 145, 141, 0, 67, 94, 186, 199, 65, 33, 48, 23, 194, 240, 103, 148, 194, 0, 0, 224, 66, 0, 0, 224, 66, 111, 146, 187, 194, 70, 182, 145, 193};
.const .align 4 .b8 k601YCbCr_To_RGB32f[36] = {37, 160, 149, 59, 0, 0, 0, 0, 182, 23, 205, 59, 37, 160, 149, 59, 40, 15, 201, 186, 156, 239, 80, 187, 37, 160, 149, 59, 236, 155, 1, 60, 0, 0, 0, 0};
.const .align 4 .b8 kRGB8u_To_601YCbCr[36] = {219, 121, 131, 62, 152, 14, 1, 63, 18, 131, 200, 61, 174, 199, 23, 190, 238, 252, 148, 190, 197, 224, 224, 62, 197, 224, 224, 62, 217, 78, 188, 190, 174, 71, 146, 189};
.const .align 4 .b8 k601YCbCr_To_RGB8u[36] = {127, 10, 149, 63, 0, 0, 0, 0, 160, 74, 204, 63, 127, 10, 149, 63, 254, 148, 200, 190, 184, 30, 80, 191, 127, 10, 149, 63, 78, 26, 1, 64, 0, 0, 0, 0};
.const .align 4 .b8 kRGB8u_To_601YCbCrFullRange[36] = {135, 22, 153, 62, 162, 69, 22, 63, 213, 120, 233, 61, 166, 27, 44, 190, 39, 241, 168, 190, 250, 254, 254, 62, 250, 254, 254, 62, 43, 135, 213, 190, 59, 223, 165, 189};
.const .align 4 .b8 k601YCbCrFullRange_To_RGB8u[36] = {0, 0, 128, 63, 0, 0, 0, 0, 72, 193, 178, 63, 0, 0, 128, 63, 143, 130, 175, 190, 225, 26, 54, 191, 0, 0, 128, 63, 20, 238, 225, 63, 0, 0, 0, 0};
.const .align 4 .b8 kRGB32f_To_601YCbCrFullRange[36] = {113, 125, 152, 66, 92, 175, 21, 67, 92, 143, 232, 65, 158, 111, 43, 194, 49, 72, 168, 194, 0, 0, 254, 66, 0, 0, 254, 66, 170, 177, 212, 194, 88, 57, 165, 193};
.const .align 4 .b8 k601YCbCrFullRange_To_RGB32f[36] = {129, 128, 128, 59, 0, 0, 0, 0, 188, 116, 179, 59, 129, 128, 128, 59, 194, 50, 176, 186, 179, 209, 54, 187, 129, 128, 128, 59, 229, 208, 226, 59, 0, 0, 0, 0};
.const .align 4 .b8 kRGB32f_To_709YPbPr[36] = {208, 179, 89, 62, 89, 23, 55, 63, 152, 221, 147, 61, 186, 164, 234, 189, 210, 86, 197, 190, 0, 0, 0, 63, 0, 0, 0, 63, 190, 134, 232, 190, 16, 202, 59, 189};
.const .align 4 .b8 k709YPbPr_To_RGB32f[36] = {0, 0, 128, 63, 0, 0, 0, 0, 12, 147, 201, 63, 0, 0, 128, 63, 221, 209, 63, 190, 243, 173, 239, 190, 0, 0, 128, 63, 77, 132, 237, 63, 0, 0, 0, 0};
.const .align 4 .b8 kRGB32f_To_709YCbCr[36] = {106, 60, 58, 66, 6, 161, 28, 67, 244, 253, 124, 65, 223, 79, 205, 193, 8, 172, 172, 194, 0, 0, 224, 66, 0, 0, 224, 66, 195, 117, 203, 194, 236, 81, 36, 193};
.const .align 4 .b8 k709YCbCr_To_RGB32f[36] = {37, 160, 149, 59, 0, 0, 0, 0, 239, 94, 230, 59, 37, 160, 149, 59, 33, 57, 91, 186, 178, 245, 8, 187, 37, 160, 149, 59, 82, 185, 7, 60, 0, 0, 0, 0};
.const .align 4 .b8 k709YCbCrFullRange_To_RGB32f[36] = {131, 128, 128, 59, 0, 0, 0, 0, 28, 147, 201, 59, 131, 128, 128, 59, 61, 210, 63, 186, 248, 173, 239, 186, 131, 128, 128, 59, 82, 132, 237, 59, 0, 0, 0, 0};
.const .align 4 .b8 kRGB8u_To_709YCbCr[36] = {207, 247, 58, 62, 53, 62, 29, 63, 231, 251, 125, 61, 147, 24, 206, 61, 23, 89, 173, 190, 197, 224, 224, 62, 197, 224, 224, 62, 12, 66, 204, 190, 195, 245, 36, 189};
.const .align 4 .b8 k709YCbCr_To_RGB8u[36] = {127, 10, 149, 63, 0, 0, 0, 0, 147, 120, 229, 63, 127, 10, 149, 63, 53, 94, 90, 190, 205, 108, 8, 191, 127, 10, 149, 63, 154, 49, 7, 64, 0, 0, 0, 0};
.const .align 4 .b8 k709YCbCr_To_601YCbCr[36] = {0, 0, 128, 63, 23, 100, 203, 61, 1, 77, 68, 62, 0, 0, 0, 0, 18, 103, 125, 63, 10, 158, 226, 189, 0, 0, 0, 0, 61, 98, 148, 189, 249, 191, 123, 63};
.const .align 4 .b8 k601YCbCr_To_709YCbCr[36] = {0, 0, 128, 63, 122, 165, 236, 189, 179, 237, 84, 190, 0, 0, 0, 0, 204, 98, 130, 63, 216, 188, 234, 61, 0, 0, 0, 0, 74, 179, 153, 61, 234, 61, 131, 63};
.const .align 4 .b8 kYCbCrOffset[12] = {0, 0, 128, 65, 0, 0, 0, 67, 0, 0, 0, 67};
.const .align 4 .b8 kYCbCrFullRangeOffset[12] = {0, 0, 0, 0, 0, 0, 0, 67, 0, 0, 0, 67};
// HorizontalRecursiveGaussianGray_kernel$__cuda_local_var_169763_363_non_const_smem has been demoted
.global .align 1 .b8 $str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90, 0};

.visible .func  (.param .b32 func_retval0) _Z9ReadFloatPKfi17DevicePixelFormat(
	.param .b64 _Z9ReadFloatPKfi17DevicePixelFormat_param_0,
	.param .b32 _Z9ReadFloatPKfi17DevicePixelFormat_param_1,
	.param .b32 _Z9ReadFloatPKfi17DevicePixelFormat_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .s16 	%rs<2>;
	.reg .s32 	%r<3>;
	.reg .f32 	%f<7>;
	.reg .s64 	%rd<6>;


	ld.param.u64 	%rd1, [_Z9ReadFloatPKfi17DevicePixelFormat_param_0];
	ld.param.u32 	%r1, [_Z9ReadFloatPKfi17DevicePixelFormat_param_1];
	ld.param.u32 	%r2, [_Z9ReadFloatPKfi17DevicePixelFormat_param_2];
	.loc 1 25 1
	setp.eq.s32	%p1, %r2, 0;
	@%p1 bra 	BB0_2;

	mul.wide.s32 	%rd2, %r1, 4;
	add.s64 	%rd3, %rd1, %rd2;
	.loc 1 25 1
	ld.f32 	%f4, [%rd3];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f6, %f4;
	bra.uni 	BB0_3;

BB0_2:
	mul.wide.s32 	%rd4, %r1, 2;
	add.s64 	%rd5, %rd1, %rd4;
	.loc 1 25 1
	ld.u16 	%rs1, [%rd5];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs1;
	cvt.f32.f16 	%f5, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f6, %f5;

BB0_3:
	st.param.f32	[func_retval0+0], %f6;
	.loc 1 25 40
	ret;
}

.visible .entry VerticalRecursiveGaussianGray_kernel(
	.param .u64 VerticalRecursiveGaussianGray_kernel_param_0,
	.param .u64 VerticalRecursiveGaussianGray_kernel_param_1,
	.param .u32 VerticalRecursiveGaussianGray_kernel_param_2,
	.param .u32 VerticalRecursiveGaussianGray_kernel_param_3,
	.param .u32 VerticalRecursiveGaussianGray_kernel_param_4,
	.param .u32 VerticalRecursiveGaussianGray_kernel_param_5,
	.param .f32 VerticalRecursiveGaussianGray_kernel_param_6,
	.param .f32 VerticalRecursiveGaussianGray_kernel_param_7,
	.param .f32 VerticalRecursiveGaussianGray_kernel_param_8,
	.param .f32 VerticalRecursiveGaussianGray_kernel_param_9,
	.param .f32 VerticalRecursiveGaussianGray_kernel_param_10,
	.param .f32 VerticalRecursiveGaussianGray_kernel_param_11,
	.param .f32 VerticalRecursiveGaussianGray_kernel_param_12,
	.param .f32 VerticalRecursiveGaussianGray_kernel_param_13,
	.param .f32 VerticalRecursiveGaussianGray_kernel_param_14
)
{
	.reg .pred 	%p<13>;
	.reg .s16 	%rs<8>;
	.reg .s32 	%r<25>;
	.reg .f32 	%f<77>;
	.reg .s64 	%rd<37>;


	ld.param.u64 	%rd12, [VerticalRecursiveGaussianGray_kernel_param_0];
	ld.param.u64 	%rd13, [VerticalRecursiveGaussianGray_kernel_param_1];
	ld.param.u32 	%r19, [VerticalRecursiveGaussianGray_kernel_param_2];
	ld.param.u32 	%r16, [VerticalRecursiveGaussianGray_kernel_param_3];
	ld.param.u32 	%r17, [VerticalRecursiveGaussianGray_kernel_param_4];
	ld.param.u32 	%r18, [VerticalRecursiveGaussianGray_kernel_param_5];
	ld.param.f32 	%f31, [VerticalRecursiveGaussianGray_kernel_param_6];
	ld.param.f32 	%f32, [VerticalRecursiveGaussianGray_kernel_param_7];
	ld.param.f32 	%f33, [VerticalRecursiveGaussianGray_kernel_param_8];
	ld.param.f32 	%f34, [VerticalRecursiveGaussianGray_kernel_param_9];
	ld.param.f32 	%f35, [VerticalRecursiveGaussianGray_kernel_param_10];
	ld.param.f32 	%f36, [VerticalRecursiveGaussianGray_kernel_param_11];
	ld.param.f32 	%f37, [VerticalRecursiveGaussianGray_kernel_param_12];
	ld.param.f32 	%f38, [VerticalRecursiveGaussianGray_kernel_param_13];
	ld.param.f32 	%f39, [VerticalRecursiveGaussianGray_kernel_param_14];
	cvta.to.global.u64 	%rd1, %rd13;
	.loc 1 93 1
	mov.u32 	%r1, %ntid.x;
	mov.u32 	%r2, %ctaid.x;
	mov.u32 	%r3, %tid.x;
	mad.lo.s32 	%r24, %r1, %r2, %r3;
	.loc 1 93 1
	setp.gt.s32	%p1, %r24, %r19;
	@%p1 bra 	BB1_29;

	.loc 1 25 1
	setp.eq.s32	%p2, %r18, 0;
	@%p2 bra 	BB1_3;

	mul.wide.s32 	%rd14, %r24, 4;
	add.s64 	%rd15, %rd1, %rd14;
	.loc 1 25 1
	ld.global.f32 	%f40, [%rd15];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f57, %f40;
	bra.uni 	BB1_4;

BB1_3:
	mul.wide.s32 	%rd16, %r24, 2;
	add.s64 	%rd17, %rd1, %rd16;
	.loc 1 25 1
	ld.global.u16 	%rs1, [%rd17];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs1;
	cvt.f32.f16 	%f41, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f57, %f41;

BB1_4:
	.loc 1 93 50
	mul.ftz.f32 	%f63, %f57, %f39;
	.loc 1 93 1
	mul.ftz.f32 	%f61, %f63, 0f3F000000;
	.loc 1 93 1
	setp.gt.s32	%p3, %r16, 0;
	@%p3 bra 	BB1_6;

	mov.u32 	%r23, 0;
	bra.uni 	BB1_14;

BB1_6:
	.loc 1 93 1
	mad.lo.s32 	%r22, %r1, %r2, %r3;
	mul.wide.s32 	%rd36, %r22, 4;
	mul.wide.s32 	%rd35, %r22, 2;
	mul.wide.s32 	%rd4, %r17, 2;
	mov.u32 	%r23, 0;
	mov.f32 	%f60, %f61;

BB1_7:
	.loc 1 25 1
	mov.f32 	%f8, %f63;
	mov.f32 	%f6, %f60;
	mov.f32 	%f60, %f61;
	@%p2 bra 	BB1_9;

	.loc 1 93 1
	add.s64 	%rd18, %rd1, %rd36;
	.loc 1 25 1
	ld.global.f32 	%f42, [%rd18];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f64, %f42;
	bra.uni 	BB1_10;

BB1_9:
	.loc 1 93 1
	add.s64 	%rd19, %rd1, %rd35;
	.loc 1 25 1
	ld.global.u16 	%rs2, [%rd19];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs2;
	cvt.f32.f16 	%f43, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f64, %f43;

BB1_10:
	.loc 1 25 40
	mov.f32 	%f63, %f64;
	.loc 1 93 1
	mul.ftz.f32 	%f44, %f8, %f32;
	fma.rn.ftz.f32 	%f45, %f63, %f31, %f44;
	fma.rn.ftz.f32 	%f46, %f60, %f33, %f45;
	fma.rn.ftz.f32 	%f61, %f6, %f34, %f46;
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f13, %f61;
	.loc 1 93 1
	@%p2 bra 	BB1_12;

	cvta.to.global.u64 	%rd20, %rd12;
	.loc 1 93 1
	add.s64 	%rd21, %rd20, %rd36;
	.loc 1 93 59
	st.global.f32 	[%rd21], %f13;
	bra.uni 	BB1_13;

BB1_12:
	cvta.to.global.u64 	%rd22, %rd12;
	.loc 1 93 1
	add.s64 	%rd23, %rd22, %rd35;
	.loc 2 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f13;
	mov.b16 	%rs3, %temp;
}
	.loc 1 93 79
	st.global.u16 	[%rd23], %rs3;

BB1_13:
	.loc 1 93 1
	add.s32 	%r24, %r24, %r17;
	.loc 1 93 1
	mul.wide.s32 	%rd24, %r17, 4;
	add.s64 	%rd36, %rd36, %rd24;
	add.s64 	%rd35, %rd35, %rd4;
	.loc 1 93 60
	add.s32 	%r23, %r23, 1;
	.loc 1 93 1
	setp.lt.s32	%p6, %r23, %r16;
	@%p6 bra 	BB1_7;

BB1_14:
	.loc 1 93 1
	sub.s32 	%r11, %r24, %r17;
	.loc 1 25 1
	@%p2 bra 	BB1_16;

	mul.wide.s32 	%rd25, %r11, 4;
	add.s64 	%rd26, %rd1, %rd25;
	.loc 1 25 1
	ld.global.f32 	%f47, [%rd26];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f65, %f47;
	bra.uni 	BB1_17;

BB1_16:
	mul.wide.s32 	%rd27, %r11, 2;
	add.s64 	%rd28, %rd1, %rd27;
	.loc 1 25 1
	ld.global.u16 	%rs4, [%rd28];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs4;
	cvt.f32.f16 	%f48, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f65, %f48;

BB1_17:
	.loc 1 93 50
	mul.ftz.f32 	%f74, %f65, %f39;
	.loc 1 93 1
	mul.ftz.f32 	%f69, %f74, 0f3F000000;
	.loc 1 93 1
	setp.lt.s32	%p8, %r23, 1;
	@%p8 bra 	BB1_29;

	cvta.to.global.u64 	%rd9, %rd12;
	mov.f32 	%f68, %f69;
	mov.f32 	%f73, %f74;

BB1_19:
	.loc 1 93 1
	mov.f32 	%f21, %f73;
	mov.f32 	%f73, %f74;
	mov.f32 	%f19, %f68;
	mov.f32 	%f68, %f69;
	mul.ftz.f32 	%f49, %f21, %f36;
	fma.rn.ftz.f32 	%f50, %f73, %f35, %f49;
	fma.rn.ftz.f32 	%f51, %f68, %f37, %f50;
	fma.rn.ftz.f32 	%f69, %f19, %f38, %f51;
	sub.s32 	%r24, %r24, %r17;
	.loc 1 25 1
	@%p2 bra 	BB1_21;

	mul.wide.s32 	%rd29, %r24, 4;
	add.s64 	%rd30, %rd1, %rd29;
	.loc 1 25 1
	ld.global.f32 	%f52, [%rd30];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f75, %f52;
	bra.uni 	BB1_22;

BB1_21:
	mul.wide.s32 	%rd31, %r24, 2;
	add.s64 	%rd32, %rd1, %rd31;
	.loc 1 25 1
	ld.global.u16 	%rs5, [%rd32];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs5;
	cvt.f32.f16 	%f53, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f75, %f53;

BB1_22:
	.loc 1 25 40
	mov.f32 	%f74, %f75;
	mul.wide.s32 	%rd33, %r24, 4;
	add.s64 	%rd10, %rd9, %rd33;
	mul.wide.s32 	%rd34, %r24, 2;
	add.s64 	%rd11, %rd9, %rd34;
	.loc 1 25 1
	@%p2 bra 	BB1_24;

	.loc 1 25 1
	ld.global.f32 	%f54, [%rd10];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f76, %f54;
	bra.uni 	BB1_25;

BB1_24:
	.loc 1 25 1
	ld.global.u16 	%rs6, [%rd11];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs6;
	cvt.f32.f16 	%f55, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f76, %f55;

BB1_25:
	.loc 1 93 58
	add.ftz.f32 	%f56, %f69, %f76;
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f30, %f56;
	.loc 1 93 1
	@%p2 bra 	BB1_27;

	.loc 1 93 18
	st.global.f32 	[%rd10], %f30;
	bra.uni 	BB1_28;

BB1_27:
	.loc 2 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f30;
	mov.b16 	%rs7, %temp;
}
	.loc 1 93 38
	st.global.u16 	[%rd11], %rs7;

BB1_28:
	.loc 1 93 55
	add.s32 	%r23, %r23, -1;
	.loc 1 93 1
	setp.gt.s32	%p12, %r23, 0;
	@%p12 bra 	BB1_19;

BB1_29:
	.loc 1 93 2
	ret;
}

.visible .entry HorizontalRecursiveGaussianGray_kernel(
	.param .u64 HorizontalRecursiveGaussianGray_kernel_param_0,
	.param .u64 HorizontalRecursiveGaussianGray_kernel_param_1,
	.param .u32 HorizontalRecursiveGaussianGray_kernel_param_2,
	.param .u32 HorizontalRecursiveGaussianGray_kernel_param_3,
	.param .u32 HorizontalRecursiveGaussianGray_kernel_param_4,
	.param .u32 HorizontalRecursiveGaussianGray_kernel_param_5,
	.param .f32 HorizontalRecursiveGaussianGray_kernel_param_6,
	.param .f32 HorizontalRecursiveGaussianGray_kernel_param_7,
	.param .f32 HorizontalRecursiveGaussianGray_kernel_param_8,
	.param .f32 HorizontalRecursiveGaussianGray_kernel_param_9,
	.param .f32 HorizontalRecursiveGaussianGray_kernel_param_10,
	.param .f32 HorizontalRecursiveGaussianGray_kernel_param_11,
	.param .f32 HorizontalRecursiveGaussianGray_kernel_param_12,
	.param .f32 HorizontalRecursiveGaussianGray_kernel_param_13,
	.param .f32 HorizontalRecursiveGaussianGray_kernel_param_14
)
{
	.reg .pred 	%p<85>;
	.reg .s16 	%rs<23>;
	.reg .s32 	%r<210>;
	.reg .f32 	%f<272>;
	.reg .s64 	%rd<112>;
	// demoted variable
	.shared .align 4 .b8 HorizontalRecursiveGaussianGray_kernel$__cuda_local_var_169763_363_non_const_smem[4224];

	ld.param.u64 	%rd49, [HorizontalRecursiveGaussianGray_kernel_param_0];
	ld.param.u64 	%rd50, [HorizontalRecursiveGaussianGray_kernel_param_1];
	ld.param.u32 	%r36, [HorizontalRecursiveGaussianGray_kernel_param_2];
	ld.param.u32 	%r37, [HorizontalRecursiveGaussianGray_kernel_param_3];
	ld.param.u32 	%r38, [HorizontalRecursiveGaussianGray_kernel_param_4];
	ld.param.u32 	%r39, [HorizontalRecursiveGaussianGray_kernel_param_5];
	ld.param.f32 	%f85, [HorizontalRecursiveGaussianGray_kernel_param_6];
	ld.param.f32 	%f86, [HorizontalRecursiveGaussianGray_kernel_param_7];
	ld.param.f32 	%f87, [HorizontalRecursiveGaussianGray_kernel_param_8];
	ld.param.f32 	%f88, [HorizontalRecursiveGaussianGray_kernel_param_9];
	ld.param.f32 	%f89, [HorizontalRecursiveGaussianGray_kernel_param_10];
	ld.param.f32 	%f90, [HorizontalRecursiveGaussianGray_kernel_param_11];
	ld.param.f32 	%f91, [HorizontalRecursiveGaussianGray_kernel_param_12];
	ld.param.f32 	%f92, [HorizontalRecursiveGaussianGray_kernel_param_13];
	ld.param.f32 	%f93, [HorizontalRecursiveGaussianGray_kernel_param_14];
	cvta.to.global.u64 	%rd1, %rd50;
	.loc 1 190 1
	mov.u32 	%r40, %ctaid.y;
	shl.b32 	%r41, %r40, 5;
	mov.u32 	%r1, %tid.y;
	add.s32 	%r42, %r41, %r1;
	setp.eq.s32	%p7, %r1, 0;
	mov.u32 	%r2, %tid.x;
	add.s32 	%r3, %r2, %r42;
	setp.lt.s32	%p8, %r3, %r37;
	and.pred  	%p9, %p7, %p8;
	setp.gt.ftz.f32	%p10, %f93, 0f00000000;
	and.pred  	%p11, %p9, %p10;
	.loc 1 190 1
	@%p11 bra 	BB2_2;

	mov.f32 	%f230, 0f00000000;
	bra.uni 	BB2_6;

BB2_2:
	.loc 1 190 1
	mul.lo.s32 	%r4, %r3, %r38;
	.loc 1 25 1
	setp.eq.s32	%p12, %r39, 0;
	@%p12 bra 	BB2_4;

	mul.wide.s32 	%rd51, %r4, 4;
	add.s64 	%rd52, %rd1, %rd51;
	.loc 1 25 1
	ld.global.f32 	%f95, [%rd52];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f210, %f95;
	bra.uni 	BB2_5;

BB2_4:
	mul.wide.s32 	%rd53, %r4, 2;
	add.s64 	%rd54, %rd1, %rd53;
	.loc 1 25 1
	ld.global.u16 	%rs1, [%rd54];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs1;
	cvt.f32.f16 	%f96, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f210, %f96;

BB2_5:
	mul.ftz.f32 	%f230, %f210, %f93;

BB2_6:
	.loc 1 190 1
	mul.ftz.f32 	%f228, %f230, 0f3F000000;
	.loc 1 190 1
	mad.lo.s32 	%r43, %r1, 33, %r2;
	mul.wide.s32 	%rd55, %r43, 4;
	mov.u64 	%rd56, HorizontalRecursiveGaussianGray_kernel$__cuda_local_var_169763_363_non_const_smem;
	add.s64 	%rd2, %rd56, %rd55;
	mov.u32 	%r208, 0;
	.loc 1 190 1
	setp.gt.s32	%p13, %r36, 0;
	@%p13 bra 	BB2_7;
	bra.uni 	BB2_56;

BB2_7:
	mov.u32 	%r205, %r208;
	mov.f32 	%f229, %f228;

BB2_8:
	.loc 1 190 1
	mov.f32 	%f227, %f229;
	mov.f32 	%f225, %f228;
	add.s32 	%r50, %r1, %r41;
	mad.lo.s32 	%r52, %r50, %r38, %r2;
	.loc 1 190 1
	shl.b32 	%r53, %r205, 5;
	add.s32 	%r7, %r52, %r53;
	mul.wide.s32 	%rd58, %r7, 4;
	add.s64 	%rd3, %rd1, %rd58;
	shl.b32 	%r8, %r7, 1;
	cvt.s64.s32	%rd59, %r8;
	add.s64 	%rd4, %rd1, %rd59;
	.loc 1 190 1
	setp.ge.s32	%p14, %r50, %r37;
	@%p14 bra 	BB2_15;

	.loc 1 190 1
	add.s32 	%r55, %r2, %r208;
	setp.lt.s32	%p15, %r55, %r36;
	.loc 1 190 1
	@%p15 bra 	BB2_11;

	mov.u32 	%r56, 0;
	.loc 1 190 1
	st.shared.u32 	[%rd2], %r56;
	bra.uni 	BB2_15;

BB2_11:
	setp.eq.s32	%p16, %r39, 0;
	.loc 1 25 1
	@%p16 bra 	BB2_13;

	.loc 1 25 1
	ld.global.f32 	%f97, [%rd3];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f211, %f97;
	bra.uni 	BB2_14;

BB2_13:
	.loc 1 25 1
	ld.global.u16 	%rs2, [%rd4];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs2;
	cvt.f32.f16 	%f98, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f211, %f98;

BB2_14:
	.loc 1 190 102
	st.shared.f32 	[%rd2], %f211;

BB2_15:
	.loc 1 190 1
	add.s32 	%r61, %r50, 8;
	shl.b32 	%r62, %r38, 3;
	add.s32 	%r63, %r7, %r62;
	mul.wide.s32 	%rd61, %r63, 4;
	add.s64 	%rd5, %rd1, %rd61;
	shl.b32 	%r9, %r63, 1;
	cvt.s64.s32	%rd62, %r9;
	add.s64 	%rd6, %rd1, %rd62;
	.loc 1 190 1
	setp.ge.s32	%p17, %r61, %r37;
	@%p17 bra 	BB2_22;

	.loc 1 190 1
	add.s32 	%r65, %r2, %r208;
	setp.lt.s32	%p18, %r65, %r36;
	.loc 1 190 1
	@%p18 bra 	BB2_18;

	mov.u32 	%r66, 0;
	.loc 1 190 1
	st.shared.u32 	[%rd2+1056], %r66;
	bra.uni 	BB2_22;

BB2_18:
	setp.eq.s32	%p19, %r39, 0;
	.loc 1 25 1
	@%p19 bra 	BB2_20;

	.loc 1 25 1
	ld.global.f32 	%f99, [%rd5];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f212, %f99;
	bra.uni 	BB2_21;

BB2_20:
	.loc 1 25 1
	ld.global.u16 	%rs3, [%rd6];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs3;
	cvt.f32.f16 	%f100, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f212, %f100;

BB2_21:
	.loc 1 190 102
	st.shared.f32 	[%rd2+1056], %f212;

BB2_22:
	.loc 1 190 1
	add.s32 	%r71, %r50, 16;
	shl.b32 	%r72, %r38, 4;
	add.s32 	%r73, %r7, %r72;
	mul.wide.s32 	%rd64, %r73, 4;
	add.s64 	%rd7, %rd1, %rd64;
	shl.b32 	%r74, %r73, 1;
	cvt.s64.s32	%rd65, %r74;
	add.s64 	%rd8, %rd1, %rd65;
	.loc 1 190 1
	setp.ge.s32	%p20, %r71, %r37;
	@%p20 bra 	BB2_29;

	.loc 1 190 1
	add.s32 	%r76, %r2, %r208;
	setp.lt.s32	%p21, %r76, %r36;
	.loc 1 190 1
	@%p21 bra 	BB2_25;

	mov.u32 	%r77, 0;
	.loc 1 190 1
	st.shared.u32 	[%rd2+2112], %r77;
	bra.uni 	BB2_29;

BB2_25:
	setp.eq.s32	%p22, %r39, 0;
	.loc 1 25 1
	@%p22 bra 	BB2_27;

	.loc 1 25 1
	ld.global.f32 	%f101, [%rd7];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f213, %f101;
	bra.uni 	BB2_28;

BB2_27:
	.loc 1 25 1
	ld.global.u16 	%rs4, [%rd8];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs4;
	cvt.f32.f16 	%f102, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f213, %f102;

BB2_28:
	.loc 1 190 102
	st.shared.f32 	[%rd2+2112], %f213;

BB2_29:
	.loc 1 190 1
	add.s32 	%r82, %r50, 24;
	mad.lo.s32 	%r83, %r38, 24, %r7;
	mul.wide.s32 	%rd67, %r83, 4;
	add.s64 	%rd9, %rd1, %rd67;
	shl.b32 	%r10, %r83, 1;
	cvt.s64.s32	%rd68, %r10;
	add.s64 	%rd10, %rd1, %rd68;
	.loc 1 190 1
	setp.ge.s32	%p23, %r82, %r37;
	@%p23 bra 	BB2_36;

	.loc 1 190 1
	add.s32 	%r85, %r2, %r208;
	setp.lt.s32	%p24, %r85, %r36;
	.loc 1 190 1
	@%p24 bra 	BB2_32;

	mov.u32 	%r86, 0;
	.loc 1 190 1
	st.shared.u32 	[%rd2+3168], %r86;
	bra.uni 	BB2_36;

BB2_32:
	setp.eq.s32	%p25, %r39, 0;
	.loc 1 25 1
	@%p25 bra 	BB2_34;

	.loc 1 25 1
	ld.global.f32 	%f103, [%rd9];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f214, %f103;
	bra.uni 	BB2_35;

BB2_34:
	.loc 1 25 1
	ld.global.u16 	%rs5, [%rd10];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs5;
	cvt.f32.f16 	%f104, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f214, %f104;

BB2_35:
	.loc 1 190 102
	st.shared.f32 	[%rd2+3168], %f214;

BB2_36:
	mad.lo.s32 	%r87, %r2, 33, 4;
	mul.wide.s32 	%rd69, %r87, 4;
	add.s64 	%rd110, %rd56, %rd69;
	.loc 1 190 1
	bar.sync 	0;
	.loc 1 190 1
	@!%p7 bra 	BB2_39;
	bra.uni 	BB2_37;

BB2_37:
	mov.u32 	%r206, 0;
	mov.f32 	%f224, %f225;
	mov.f32 	%f226, %f227;

BB2_38:
	.loc 1 190 1
	ld.shared.f32 	%f105, [%rd110+-16];
	mul.ftz.f32 	%f106, %f230, %f86;
	fma.rn.ftz.f32 	%f107, %f105, %f85, %f106;
	fma.rn.ftz.f32 	%f108, %f224, %f87, %f107;
	fma.rn.ftz.f32 	%f109, %f226, %f88, %f108;
	ld.shared.f32 	%f110, [%rd110+-12];
	ld.shared.f32 	%f111, [%rd110+-8];
	ld.shared.f32 	%f112, [%rd110+-4];
	st.shared.f32 	[%rd110+-16], %f109;
	mul.ftz.f32 	%f113, %f105, %f86;
	fma.rn.ftz.f32 	%f114, %f110, %f85, %f113;
	fma.rn.ftz.f32 	%f115, %f109, %f87, %f114;
	fma.rn.ftz.f32 	%f116, %f224, %f88, %f115;
	st.shared.f32 	[%rd110+-12], %f116;
	mul.ftz.f32 	%f117, %f110, %f86;
	fma.rn.ftz.f32 	%f118, %f111, %f85, %f117;
	fma.rn.ftz.f32 	%f119, %f116, %f87, %f118;
	fma.rn.ftz.f32 	%f120, %f109, %f88, %f119;
	st.shared.f32 	[%rd110+-8], %f120;
	mul.ftz.f32 	%f121, %f111, %f86;
	fma.rn.ftz.f32 	%f122, %f112, %f85, %f121;
	fma.rn.ftz.f32 	%f123, %f120, %f87, %f122;
	fma.rn.ftz.f32 	%f124, %f116, %f88, %f123;
	st.shared.f32 	[%rd110+-4], %f124;
	ld.shared.f32 	%f125, [%rd110];
	mul.ftz.f32 	%f126, %f112, %f86;
	fma.rn.ftz.f32 	%f127, %f125, %f85, %f126;
	fma.rn.ftz.f32 	%f128, %f124, %f87, %f127;
	fma.rn.ftz.f32 	%f129, %f120, %f88, %f128;
	ld.shared.f32 	%f130, [%rd110+4];
	ld.shared.f32 	%f131, [%rd110+8];
	ld.shared.f32 	%f230, [%rd110+12];
	st.shared.f32 	[%rd110], %f129;
	mul.ftz.f32 	%f132, %f125, %f86;
	fma.rn.ftz.f32 	%f133, %f130, %f85, %f132;
	fma.rn.ftz.f32 	%f134, %f129, %f87, %f133;
	fma.rn.ftz.f32 	%f135, %f124, %f88, %f134;
	st.shared.f32 	[%rd110+4], %f135;
	mul.ftz.f32 	%f136, %f130, %f86;
	fma.rn.ftz.f32 	%f137, %f131, %f85, %f136;
	fma.rn.ftz.f32 	%f138, %f135, %f87, %f137;
	fma.rn.ftz.f32 	%f227, %f129, %f88, %f138;
	st.shared.f32 	[%rd110+8], %f227;
	mul.ftz.f32 	%f139, %f131, %f86;
	fma.rn.ftz.f32 	%f140, %f230, %f85, %f139;
	fma.rn.ftz.f32 	%f141, %f227, %f87, %f140;
	fma.rn.ftz.f32 	%f225, %f135, %f88, %f141;
	st.shared.f32 	[%rd110+12], %f225;
	add.s64 	%rd110, %rd110, 32;
	.loc 1 190 18
	add.s32 	%r206, %r206, 8;
	.loc 1 190 1
	setp.ne.s32	%p26, %r206, 32;
	mov.f32 	%f226, %f227;
	mov.f32 	%f224, %f225;
	@%p26 bra 	BB2_38;

BB2_39:
	.loc 1 190 1
	mov.f32 	%f229, %f227;
	mov.f32 	%f228, %f225;
	.loc 1 190 1
	add.s32 	%r89, %r2, %r208;
	setp.lt.s32	%p2, %r89, %r36;
	setp.lt.s32	%p3, %r50, %r37;
	cvta.to.global.u64 	%rd15, %rd49;
	.loc 1 190 1
	bar.sync 	0;
	mul.wide.s32 	%rd71, %r7, 4;
	add.s64 	%rd16, %rd15, %rd71;
	add.s64 	%rd17, %rd15, %rd59;
	.loc 1 190 1
	not.pred 	%p27, %p3;
	not.pred 	%p28, %p2;
	or.pred  	%p29, %p27, %p28;
	@%p29 bra 	BB2_43;

	setp.eq.s32	%p30, %r39, 0;
	.loc 1 190 1
	ld.shared.f32 	%f142, [%rd2];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f31, %f142;
	.loc 1 190 1
	@%p30 bra 	BB2_42;

	.loc 1 190 59
	st.global.f32 	[%rd16], %f31;
	bra.uni 	BB2_43;

BB2_42:
	.loc 2 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f31;
	mov.b16 	%rs6, %temp;
}
	.loc 1 190 79
	st.global.u16 	[%rd17], %rs6;

BB2_43:
	.loc 1 190 1
	shl.b32 	%r195, %r38, 3;
	add.s32 	%r194, %r7, %r195;
	mul.wide.s32 	%rd73, %r194, 4;
	add.s64 	%rd19, %rd15, %rd73;
	add.s64 	%rd20, %rd15, %rd62;
	.loc 1 190 1
	setp.ge.s32	%p32, %r89, %r36;
	or.pred  	%p33, %p17, %p32;
	@%p33 bra 	BB2_47;

	setp.eq.s32	%p34, %r39, 0;
	.loc 1 190 1
	ld.shared.f32 	%f143, [%rd2+1056];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f32, %f143;
	.loc 1 190 1
	@%p34 bra 	BB2_46;

	.loc 1 190 59
	st.global.f32 	[%rd19], %f32;
	bra.uni 	BB2_47;

BB2_46:
	.loc 2 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f32;
	mov.b16 	%rs7, %temp;
}
	.loc 1 190 79
	st.global.u16 	[%rd20], %rs7;

BB2_47:
	.loc 1 190 1
	shl.b32 	%r197, %r38, 4;
	add.s32 	%r196, %r7, %r197;
	mul.wide.s32 	%rd75, %r196, 4;
	add.s64 	%rd21, %rd15, %rd75;
	add.s64 	%rd22, %rd15, %rd65;
	.loc 1 190 1
	or.pred  	%p37, %p20, %p32;
	@%p37 bra 	BB2_51;

	setp.eq.s32	%p38, %r39, 0;
	.loc 1 190 1
	ld.shared.f32 	%f144, [%rd2+2112];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f33, %f144;
	.loc 1 190 1
	@%p38 bra 	BB2_50;

	.loc 1 190 59
	st.global.f32 	[%rd21], %f33;
	bra.uni 	BB2_51;

BB2_50:
	.loc 2 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f33;
	mov.b16 	%rs8, %temp;
}
	.loc 1 190 79
	st.global.u16 	[%rd22], %rs8;

BB2_51:
	.loc 1 190 1
	mad.lo.s32 	%r198, %r38, 24, %r7;
	mul.wide.s32 	%rd77, %r198, 4;
	add.s64 	%rd23, %rd15, %rd77;
	add.s64 	%rd24, %rd15, %rd68;
	.loc 1 190 1
	or.pred  	%p41, %p23, %p32;
	@%p41 bra 	BB2_55;

	setp.eq.s32	%p42, %r39, 0;
	.loc 1 190 1
	ld.shared.f32 	%f145, [%rd2+3168];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f34, %f145;
	.loc 1 190 1
	@%p42 bra 	BB2_54;

	.loc 1 190 59
	st.global.f32 	[%rd23], %f34;
	bra.uni 	BB2_55;

BB2_54:
	.loc 2 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f34;
	mov.b16 	%rs9, %temp;
}
	.loc 1 190 79
	st.global.u16 	[%rd24], %rs9;

BB2_55:
	.loc 1 190 1
	bar.sync 	0;
	.loc 1 190 1
	add.s32 	%r208, %r208, 32;
	.loc 1 190 1
	setp.lt.s32	%p43, %r208, %r36;
	add.s32 	%r205, %r205, 1;
	@%p43 bra 	BB2_8;

BB2_56:
	.loc 1 190 1
	@%p11 bra 	BB2_58;

	mov.f32 	%f267, 0f00000000;
	bra.uni 	BB2_62;

BB2_58:
	.loc 1 190 1
	add.s32 	%r199, %r41, %r1;
	mad.lo.s32 	%r136, %r199, %r38, %r36;
	add.s32 	%r137, %r136, -1;
	mad.lo.s32 	%r20, %r2, %r38, %r137;
	.loc 1 25 1
	setp.eq.s32	%p49, %r39, 0;
	@%p49 bra 	BB2_60;

	mul.wide.s32 	%rd80, %r20, 4;
	add.s64 	%rd81, %rd1, %rd80;
	.loc 1 25 1
	ld.global.f32 	%f147, [%rd81];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f231, %f147;
	bra.uni 	BB2_61;

BB2_60:
	mul.wide.s32 	%rd83, %r20, 2;
	add.s64 	%rd84, %rd1, %rd83;
	.loc 1 25 1
	ld.global.u16 	%rs10, [%rd84];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs10;
	cvt.f32.f16 	%f148, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f231, %f148;

BB2_61:
	ld.param.f32 	%f209, [HorizontalRecursiveGaussianGray_kernel_param_14];
	mul.ftz.f32 	%f267, %f231, %f209;

BB2_62:
	.loc 1 190 1
	mov.f32 	%f265, %f267;
	mul.ftz.f32 	%f249, %f265, 0f3F000000;
	.loc 1 190 1
	setp.lt.s32	%p50, %r208, 1;
	@%p50 bra 	BB2_124;

	add.s32 	%r142, %r1, %r41;
	add.s32 	%r144, %r2, %r208;
	mad.lo.s32 	%r145, %r142, %r38, %r144;
	add.s32 	%r21, %r145, -32;
	mov.u32 	%r207, 0;
	mov.f32 	%f250, %f249;
	mov.f32 	%f266, %f265;

BB2_64:
	.loc 1 190 1
	mov.f32 	%f264, %f266;
	mov.f32 	%f262, %f265;
	mov.f32 	%f248, %f250;
	mov.f32 	%f246, %f249;
	mad.lo.s32 	%r24, %r207, -32, %r21;
	add.s32 	%r208, %r208, -32;
	mul.wide.s32 	%rd86, %r24, 4;
	add.s64 	%rd25, %rd1, %rd86;
	shl.b32 	%r26, %r24, 1;
	cvt.s64.s32	%rd87, %r26;
	add.s64 	%rd26, %rd1, %rd87;
	.loc 1 190 1
	setp.ge.s32	%p51, %r142, %r37;
	@%p51 bra 	BB2_71;

	.loc 1 190 1
	add.s32 	%r151, %r2, %r208;
	setp.lt.s32	%p52, %r151, %r36;
	.loc 1 190 1
	@%p52 bra 	BB2_67;

	mov.u32 	%r152, 0;
	.loc 1 190 1
	st.shared.u32 	[%rd2], %r152;
	bra.uni 	BB2_71;

BB2_67:
	setp.eq.s32	%p53, %r39, 0;
	.loc 1 25 1
	@%p53 bra 	BB2_69;

	.loc 1 25 1
	ld.global.f32 	%f149, [%rd25];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f232, %f149;
	bra.uni 	BB2_70;

BB2_69:
	.loc 1 25 1
	ld.global.u16 	%rs11, [%rd26];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs11;
	cvt.f32.f16 	%f150, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f232, %f150;

BB2_70:
	.loc 1 190 103
	st.shared.f32 	[%rd2], %f232;

BB2_71:
	.loc 1 190 1
	add.s32 	%r157, %r142, 8;
	shl.b32 	%r158, %r38, 3;
	add.s32 	%r27, %r24, %r158;
	mul.wide.s32 	%rd89, %r27, 4;
	add.s64 	%rd27, %rd1, %rd89;
	shl.b32 	%r28, %r27, 1;
	cvt.s64.s32	%rd90, %r28;
	add.s64 	%rd28, %rd1, %rd90;
	.loc 1 190 1
	setp.ge.s32	%p54, %r157, %r37;
	@%p54 bra 	BB2_78;

	.loc 1 190 1
	add.s32 	%r160, %r2, %r208;
	setp.lt.s32	%p55, %r160, %r36;
	.loc 1 190 1
	@%p55 bra 	BB2_74;

	mov.u32 	%r161, 0;
	.loc 1 190 1
	st.shared.u32 	[%rd2+1056], %r161;
	bra.uni 	BB2_78;

BB2_74:
	setp.eq.s32	%p56, %r39, 0;
	.loc 1 25 1
	@%p56 bra 	BB2_76;

	.loc 1 25 1
	ld.global.f32 	%f151, [%rd27];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f233, %f151;
	bra.uni 	BB2_77;

BB2_76:
	.loc 1 25 1
	ld.global.u16 	%rs12, [%rd28];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs12;
	cvt.f32.f16 	%f152, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f233, %f152;

BB2_77:
	.loc 1 190 103
	st.shared.f32 	[%rd2+1056], %f233;

BB2_78:
	.loc 1 190 1
	add.s32 	%r166, %r142, 16;
	shl.b32 	%r167, %r38, 4;
	add.s32 	%r29, %r24, %r167;
	mul.wide.s32 	%rd92, %r29, 4;
	add.s64 	%rd29, %rd1, %rd92;
	shl.b32 	%r168, %r29, 1;
	cvt.s64.s32	%rd93, %r168;
	add.s64 	%rd30, %rd1, %rd93;
	.loc 1 190 1
	setp.ge.s32	%p57, %r166, %r37;
	@%p57 bra 	BB2_85;

	.loc 1 190 1
	add.s32 	%r170, %r2, %r208;
	setp.lt.s32	%p58, %r170, %r36;
	.loc 1 190 1
	@%p58 bra 	BB2_81;

	mov.u32 	%r171, 0;
	.loc 1 190 1
	st.shared.u32 	[%rd2+2112], %r171;
	bra.uni 	BB2_85;

BB2_81:
	setp.eq.s32	%p59, %r39, 0;
	.loc 1 25 1
	@%p59 bra 	BB2_83;

	.loc 1 25 1
	ld.global.f32 	%f153, [%rd29];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f234, %f153;
	bra.uni 	BB2_84;

BB2_83:
	.loc 1 25 1
	ld.global.u16 	%rs13, [%rd30];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs13;
	cvt.f32.f16 	%f154, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f234, %f154;

BB2_84:
	.loc 1 190 103
	st.shared.f32 	[%rd2+2112], %f234;

BB2_85:
	.loc 1 190 1
	add.s32 	%r176, %r142, 24;
	mad.lo.s32 	%r30, %r38, 24, %r24;
	setp.ge.s32	%p60, %r176, %r37;
	@%p60 bra 	BB2_92;

	.loc 1 190 1
	add.s32 	%r178, %r2, %r208;
	setp.lt.s32	%p61, %r178, %r36;
	.loc 1 190 1
	@%p61 bra 	BB2_88;

	mov.u32 	%r179, 0;
	.loc 1 190 1
	st.shared.u32 	[%rd2+3168], %r179;
	bra.uni 	BB2_92;

BB2_88:
	setp.eq.s32	%p62, %r39, 0;
	.loc 1 25 1
	@%p62 bra 	BB2_90;

	mul.wide.s32 	%rd95, %r30, 4;
	add.s64 	%rd96, %rd1, %rd95;
	.loc 1 25 1
	ld.global.f32 	%f155, [%rd96];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f235, %f155;
	bra.uni 	BB2_91;

BB2_90:
	shl.b32 	%r180, %r30, 1;
	cvt.s64.s32	%rd98, %r180;
	add.s64 	%rd99, %rd1, %rd98;
	.loc 1 25 1
	ld.global.u16 	%rs14, [%rd99];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs14;
	cvt.f32.f16 	%f156, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f235, %f156;

BB2_91:
	.loc 1 190 103
	st.shared.f32 	[%rd2+3168], %f235;

BB2_92:
	mad.lo.s32 	%r182, %r2, 33, 31;
	mul.wide.s32 	%rd100, %r182, 4;
	add.s64 	%rd111, %rd56, %rd100;
	.loc 1 190 1
	bar.sync 	0;
	.loc 1 190 1
	@!%p7 bra 	BB2_95;
	bra.uni 	BB2_93;

BB2_93:
	mov.u32 	%r209, 0;
	mov.f32 	%f245, %f246;
	mov.f32 	%f247, %f248;
	mov.f32 	%f261, %f262;
	mov.f32 	%f263, %f264;

BB2_94:
	.loc 1 190 1
	mul.ftz.f32 	%f157, %f263, %f90;
	fma.rn.ftz.f32 	%f158, %f261, %f89, %f157;
	fma.rn.ftz.f32 	%f159, %f245, %f91, %f158;
	fma.rn.ftz.f32 	%f160, %f247, %f92, %f159;
	ld.shared.f32 	%f161, [%rd111];
	ld.shared.f32 	%f162, [%rd111+-4];
	ld.shared.f32 	%f163, [%rd111+-8];
	ld.shared.f32 	%f164, [%rd111+-12];
	st.shared.f32 	[%rd111], %f160;
	mul.ftz.f32 	%f165, %f261, %f90;
	fma.rn.ftz.f32 	%f166, %f161, %f89, %f165;
	fma.rn.ftz.f32 	%f167, %f160, %f91, %f166;
	fma.rn.ftz.f32 	%f168, %f245, %f92, %f167;
	st.shared.f32 	[%rd111+-4], %f168;
	mul.ftz.f32 	%f169, %f161, %f90;
	fma.rn.ftz.f32 	%f170, %f162, %f89, %f169;
	fma.rn.ftz.f32 	%f171, %f168, %f91, %f170;
	fma.rn.ftz.f32 	%f172, %f160, %f92, %f171;
	st.shared.f32 	[%rd111+-8], %f172;
	mul.ftz.f32 	%f173, %f162, %f90;
	fma.rn.ftz.f32 	%f174, %f163, %f89, %f173;
	fma.rn.ftz.f32 	%f175, %f172, %f91, %f174;
	fma.rn.ftz.f32 	%f176, %f168, %f92, %f175;
	st.shared.f32 	[%rd111+-12], %f176;
	mul.ftz.f32 	%f177, %f163, %f90;
	fma.rn.ftz.f32 	%f178, %f164, %f89, %f177;
	fma.rn.ftz.f32 	%f179, %f176, %f91, %f178;
	fma.rn.ftz.f32 	%f180, %f172, %f92, %f179;
	ld.shared.f32 	%f181, [%rd111+-16];
	ld.shared.f32 	%f182, [%rd111+-20];
	ld.shared.f32 	%f264, [%rd111+-24];
	ld.shared.f32 	%f262, [%rd111+-28];
	st.shared.f32 	[%rd111+-16], %f180;
	mul.ftz.f32 	%f183, %f164, %f90;
	fma.rn.ftz.f32 	%f184, %f181, %f89, %f183;
	fma.rn.ftz.f32 	%f185, %f180, %f91, %f184;
	fma.rn.ftz.f32 	%f186, %f176, %f92, %f185;
	st.shared.f32 	[%rd111+-20], %f186;
	mul.ftz.f32 	%f187, %f181, %f90;
	fma.rn.ftz.f32 	%f188, %f182, %f89, %f187;
	fma.rn.ftz.f32 	%f189, %f186, %f91, %f188;
	fma.rn.ftz.f32 	%f248, %f180, %f92, %f189;
	st.shared.f32 	[%rd111+-24], %f248;
	mul.ftz.f32 	%f190, %f182, %f90;
	fma.rn.ftz.f32 	%f191, %f264, %f89, %f190;
	fma.rn.ftz.f32 	%f192, %f248, %f91, %f191;
	fma.rn.ftz.f32 	%f246, %f186, %f92, %f192;
	st.shared.f32 	[%rd111+-28], %f246;
	add.s64 	%rd111, %rd111, -32;
	.loc 1 190 1
	add.s32 	%r209, %r209, 8;
	setp.ne.s32	%p63, %r209, 32;
	mov.f32 	%f247, %f248;
	mov.f32 	%f245, %f246;
	mov.f32 	%f263, %f264;
	mov.f32 	%f261, %f262;
	@%p63 bra 	BB2_94;

BB2_95:
	.loc 1 190 1
	mov.f32 	%f266, %f264;
	mov.f32 	%f265, %f262;
	mov.f32 	%f250, %f248;
	mov.f32 	%f249, %f246;
	.loc 1 190 1
	add.s32 	%r33, %r2, %r208;
	setp.lt.s32	%p5, %r33, %r36;
	setp.lt.s32	%p6, %r142, %r37;
	cvta.to.global.u64 	%rd35, %rd49;
	.loc 1 190 1
	bar.sync 	0;
	mul.wide.s32 	%rd102, %r24, 4;
	add.s64 	%rd36, %rd35, %rd102;
	add.s64 	%rd37, %rd35, %rd87;
	.loc 1 190 1
	not.pred 	%p64, %p6;
	not.pred 	%p65, %p5;
	or.pred  	%p66, %p64, %p65;
	@%p66 bra 	BB2_102;

	setp.eq.s32	%p67, %r39, 0;
	.loc 1 25 1
	@%p67 bra 	BB2_98;

	.loc 1 25 1
	ld.global.f32 	%f193, [%rd36];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f268, %f193;
	bra.uni 	BB2_99;

BB2_98:
	.loc 1 25 1
	ld.global.u16 	%rs15, [%rd37];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs15;
	cvt.f32.f16 	%f194, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f268, %f194;

BB2_99:
	.loc 1 190 13
	ld.shared.f32 	%f195, [%rd2];
	add.ftz.f32 	%f196, %f268, %f195;
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f72, %f196;
	.loc 1 190 1
	@%p67 bra 	BB2_101;

	.loc 1 190 20
	st.global.f32 	[%rd36], %f72;
	bra.uni 	BB2_102;

BB2_101:
	.loc 2 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f72;
	mov.b16 	%rs16, %temp;
}
	.loc 1 190 40
	st.global.u16 	[%rd37], %rs16;

BB2_102:
	.loc 1 190 1
	shl.b32 	%r201, %r38, 3;
	add.s32 	%r200, %r24, %r201;
	mul.wide.s32 	%rd104, %r200, 4;
	add.s64 	%rd40, %rd35, %rd104;
	add.s64 	%rd41, %rd35, %rd90;
	.loc 1 190 1
	setp.ge.s32	%p70, %r33, %r36;
	or.pred  	%p71, %p54, %p70;
	@%p71 bra 	BB2_109;

	setp.eq.s32	%p72, %r39, 0;
	.loc 1 25 1
	@%p72 bra 	BB2_105;

	.loc 1 25 1
	ld.global.f32 	%f197, [%rd40];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f269, %f197;
	bra.uni 	BB2_106;

BB2_105:
	.loc 1 25 1
	ld.global.u16 	%rs17, [%rd41];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs17;
	cvt.f32.f16 	%f198, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f269, %f198;

BB2_106:
	.loc 1 190 13
	ld.shared.f32 	%f199, [%rd2+1056];
	add.ftz.f32 	%f200, %f269, %f199;
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f76, %f200;
	.loc 1 190 1
	@%p72 bra 	BB2_108;

	.loc 1 190 20
	st.global.f32 	[%rd40], %f76;
	bra.uni 	BB2_109;

BB2_108:
	.loc 2 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f76;
	mov.b16 	%rs18, %temp;
}
	.loc 1 190 40
	st.global.u16 	[%rd41], %rs18;

BB2_109:
	.loc 1 190 1
	shl.b32 	%r203, %r38, 4;
	add.s32 	%r202, %r24, %r203;
	mul.wide.s32 	%rd106, %r202, 4;
	add.s64 	%rd43, %rd35, %rd106;
	add.s64 	%rd44, %rd35, %rd93;
	.loc 1 190 1
	or.pred  	%p76, %p57, %p70;
	@%p76 bra 	BB2_116;

	setp.eq.s32	%p77, %r39, 0;
	.loc 1 25 1
	@%p77 bra 	BB2_112;

	.loc 1 25 1
	ld.global.f32 	%f201, [%rd43];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f270, %f201;
	bra.uni 	BB2_113;

BB2_112:
	.loc 1 25 1
	ld.global.u16 	%rs19, [%rd44];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs19;
	cvt.f32.f16 	%f202, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f270, %f202;

BB2_113:
	.loc 1 190 13
	ld.shared.f32 	%f203, [%rd2+2112];
	add.ftz.f32 	%f204, %f270, %f203;
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f80, %f204;
	.loc 1 190 1
	@%p77 bra 	BB2_115;

	.loc 1 190 20
	st.global.f32 	[%rd43], %f80;
	bra.uni 	BB2_116;

BB2_115:
	.loc 2 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f80;
	mov.b16 	%rs20, %temp;
}
	.loc 1 190 40
	st.global.u16 	[%rd44], %rs20;

BB2_116:
	.loc 1 190 1
	mad.lo.s32 	%r204, %r38, 24, %r24;
	mul.wide.s32 	%rd108, %r204, 4;
	add.s64 	%rd46, %rd35, %rd108;
	shl.b32 	%r193, %r204, 1;
	cvt.s64.s32	%rd109, %r193;
	add.s64 	%rd47, %rd35, %rd109;
	.loc 1 190 1
	or.pred  	%p81, %p60, %p70;
	@%p81 bra 	BB2_123;

	setp.eq.s32	%p82, %r39, 0;
	.loc 1 25 1
	@%p82 bra 	BB2_119;

	.loc 1 25 1
	ld.global.f32 	%f205, [%rd46];
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f271, %f205;
	bra.uni 	BB2_120;

BB2_119:
	.loc 1 25 1
	ld.global.u16 	%rs21, [%rd47];
	.loc 2 3518 10
	{
	.reg .b16 %temp;
	mov.b16 	%temp, %rs21;
	cvt.f32.f16 	%f206, %temp;
	}
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f271, %f206;

BB2_120:
	.loc 1 190 13
	ld.shared.f32 	%f207, [%rd2+3168];
	add.ftz.f32 	%f208, %f271, %f207;
	.loc 2 2820 10
	cvt.ftz.sat.f32.f32	%f84, %f208;
	.loc 1 190 1
	@%p82 bra 	BB2_122;

	.loc 1 190 20
	st.global.f32 	[%rd46], %f84;
	bra.uni 	BB2_123;

BB2_122:
	.loc 2 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f84;
	mov.b16 	%rs22, %temp;
}
	.loc 1 190 40
	st.global.u16 	[%rd47], %rs22;

BB2_123:
	.loc 1 190 1
	bar.sync 	0;
	.loc 1 190 1
	add.s32 	%r207, %r207, 1;
	setp.gt.s32	%p84, %r208, 0;
	@%p84 bra 	BB2_64;

BB2_124:
	.loc 1 190 2
	ret;
}


