	.version 1.4
	.target sm_11, map_f64_to_f32
	// compiled with ../../../External/3rdParty/NVIDIA/CUDA/win/64/bin64/../open64/lib//be.exe
	// nvopencc 2.3 built on 2009-07-14

	//-----------------------------------------------------------
	// Compiling C:\Users\dvaeng\AppData\Local\Temp/tmpxft_0001d904_00000000-9_SingleChannelBlur.cpp3.i (C:/Users/dvaeng/AppData/Local/Temp/ccBI#.a20356)
	//-----------------------------------------------------------

	//-----------------------------------------------------------
	// Options:
	//-----------------------------------------------------------
	//  Target:ptx, ISA:sm_11, Endian:little, Pointer Size:64
	//  -O3	(Optimization level)
	//  -g0	(Debug level)
	//  -m2	(Report advisories)
	//-----------------------------------------------------------

	.file	1	"C:\Users\dvaeng\AppData\Local\Temp/tmpxft_0001d904_00000000-8_SingleChannelBlur.cudafe2.gpu"
	.file	2	"C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\include\crtdefs.h"
	.file	3	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\crt/device_runtime.h"
	.file	4	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\host_defines.h"
	.file	5	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\builtin_types.h"
	.file	6	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\device_types.h"
	.file	7	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\driver_types.h"
	.file	8	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\texture_types.h"
	.file	9	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\vector_types.h"
	.file	10	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\host_defines.h"
	.file	11	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\device_launch_parameters.h"
	.file	12	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\crt\storage_class.h"
	.file	13	"C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\include\time.h"
	.file	14	"c:\scully64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/Utils.h"
	.file	15	"c:\scully64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/PixelRGB.h"
	.file	16	"c:/scully64/shared/adobe/MediaCore/Display/Src/CUDA/Effects/SingleChannelBlur.cu"
	.file	17	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\common_functions.h"
	.file	18	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\crt/func_macro.h"
	.file	19	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\math_functions.h"
	.file	20	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\device_functions.h"
	.file	21	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\math_constants.h"
	.file	22	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\sm_11_atomic_functions.h"
	.file	23	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\sm_12_atomic_functions.h"
	.file	24	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\sm_13_double_functions.h"
	.file	25	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\common_types.h"
	.file	26	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\texture_fetch_functions.h"
	.file	27	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\math_functions_dbl_ptx1.h"
	.file	28	"c:\scully64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/ColorSpaceConvert.h"

	.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 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 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};

	.entry VerticalRecursiveGaussianGRAYF32_kernel (
		.param .u64 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_pOut,
		.param .u64 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_pIn,
		.param .s32 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_width,
		.param .s32 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_height,
		.param .s32 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_pitch_f,
		.param .f32 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_n_plus0,
		.param .f32 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_n_plus1,
		.param .f32 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_neg_d_plus1,
		.param .f32 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_neg_d_plus2,
		.param .f32 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_n_minus1,
		.param .f32 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_n_minus2,
		.param .f32 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_neg_d_minus1,
		.param .f32 __cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_neg_d_minus2)
	{
	.reg .u32 %r<21>;
	.reg .u64 %rd<17>;
	.reg .f32 %f<25>;
	.reg .pred %p<7>;
	.loc	16	30	0
$LBB1_VerticalRecursiveGaussianGRAYF32_kernel:
	.loc	16	33	0
	cvt.s32.u16 	%r1, %ntid.x;
	cvt.s32.u16 	%r2, %ctaid.x;
	mul24.lo.s32 	%r3, %r1, %r2;
	cvt.u32.u16 	%r4, %tid.x;
	add.u32 	%r5, %r3, %r4;
	ld.param.s32 	%r6, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_width];
	setp.ge.s32 	%p1, %r6, %r5;
	@%p1 bra 	$Lt_0_3074;
	bra.uni 	$LBB14_VerticalRecursiveGaussianGRAYF32_kernel;
$Lt_0_3074:
	.loc	16	48	0
	mov.s32 	%r7, %r5;
	ld.param.s32 	%r8, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_height];
	mov.u32 	%r9, 0;
	setp.le.s32 	%p2, %r8, %r9;
	mov.s32 	%r10, 0;
	@%p2 bra 	$Lt_0_5634;
	mov.s32 	%r11, %r8;
	ld.param.s32 	%r12, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_pitch_f];
	cvt.s64.s32 	%rd1, %r12;
	cvt.u64.s32 	%rd2, %r5;
	mul.lo.u64 	%rd3, %rd2, 4;
	mul.lo.u64 	%rd4, %rd1, 4;
	ld.param.u64 	%rd5, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_pIn];
	add.u64 	%rd6, %rd5, %rd3;
	ld.param.u64 	%rd7, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_pOut];
	add.u64 	%rd8, %rd7, %rd3;
	ld.param.f32 	%f1, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_neg_d_plus2];
	ld.param.f32 	%f2, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_neg_d_plus1];
	ld.param.f32 	%f3, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_n_plus0];
	ld.param.f32 	%f4, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_n_plus1];
	mov.f32 	%f5, 0f00000000;     	// 0
	mov.f32 	%f6, 0f00000000;     	// 0
	mov.f32 	%f7, 0f00000000;     	// 0
	mov.s32 	%r13, %r11;
$Lt_0_4098:
 //<loop> Loop body line 48, nesting depth: 1, estimated iterations: unknown
	.loc	16	52	0
	ld.global.f32 	%f8, [%rd6+0];
	.loc	16	53	0
	mul.f32 	%f9, %f4, %f7;
	mad.f32 	%f10, %f3, %f8, %f9;
	mad.f32 	%f11, %f2, %f6, %f10;
	mad.f32 	%f12, %f1, %f5, %f11;
	.loc	16	55	0
	mov.f32 	%f7, %f8;
	.loc	16	56	0
	mov.f32 	%f5, %f6;
	.loc	16	57	0
	mov.f32 	%f6, %f12;
	.loc	16	59	0
	st.global.f32 	[%rd8+0], %f12;
	add.s32 	%r10, %r10, 1;
	add.u64 	%rd8, %rd8, %rd4;
	add.u64 	%rd6, %rd6, %rd4;
	setp.ne.s32 	%p3, %r8, %r10;
	@%p3 bra 	$Lt_0_4098;
	mov.s32 	%r10, %r8;
	mul.lo.s32 	%r14, %r12, %r8;
	add.s32 	%r7, %r5, %r14;
	bra.uni 	$Lt_0_3586;
$Lt_0_5634:
$Lt_0_3586:
	mov.u32 	%r15, 0;
	setp.le.s32 	%p4, %r10, %r15;
	@%p4 bra 	$LBB14_VerticalRecursiveGaussianGRAYF32_kernel;
	mov.s32 	%r16, %r10;
	ld.param.s32 	%r17, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_pitch_f];
	cvt.s64.s32 	%rd9, %r17;
	cvt.u64.s32 	%rd10, %r7;
	mul.lo.u64 	%rd11, %rd10, 4;
	mul.lo.u64 	%rd4, %rd9, 4;
	ld.param.u64 	%rd12, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_pIn];
	add.u64 	%rd13, %rd12, %rd11;
	ld.param.u64 	%rd14, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_pOut];
	add.u64 	%rd15, %rd14, %rd11;
	ld.param.f32 	%f13, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_neg_d_minus2];
	ld.param.f32 	%f14, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_neg_d_minus1];
	ld.param.f32 	%f15, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_n_minus1];
	ld.param.f32 	%f16, [__cudaparm_VerticalRecursiveGaussianGRAYF32_kernel_n_minus2];
	mov.f32 	%f17, 0f00000000;    	// 0
	mov.f32 	%f5, 0f00000000;     	// 0
	mov.f32 	%f6, 0f00000000;     	// 0
	mov.f32 	%f7, 0f00000000;     	// 0
	mov.s32 	%r18, %r16;
$Lt_0_5122:
 //<loop> Loop body line 59, nesting depth: 1, estimated iterations: unknown
	.loc	16	73	0
	sub.s32 	%r10, %r10, 1;
	.loc	16	74	0
	sub.u64 	%rd15, %rd15, %rd4;
	sub.u64 	%rd13, %rd13, %rd4;
	.loc	16	76	0
	mul.f32 	%f18, %f16, %f17;
	mad.f32 	%f19, %f15, %f7, %f18;
	mad.f32 	%f20, %f14, %f6, %f19;
	mad.f32 	%f21, %f13, %f5, %f20;
	.loc	16	78	0
	mov.f32 	%f17, %f7;
	.loc	16	79	0
	ld.global.f32 	%f7, [%rd13+0];
	.loc	16	80	0
	mov.f32 	%f5, %f6;
	.loc	16	81	0
	mov.f32 	%f6, %f21;
	.loc	16	83	0
	ld.global.f32 	%f22, [%rd15+0];
	add.f32 	%f23, %f22, %f21;
	st.global.f32 	[%rd15+0], %f23;
	mov.u32 	%r19, 0;
	setp.ne.s32 	%p5, %r10, %r19;
	@%p5 bra 	$Lt_0_5122;
$LBB14_VerticalRecursiveGaussianGRAYF32_kernel:
	.loc	16	85	0
	exit;
$LDWend_VerticalRecursiveGaussianGRAYF32_kernel:
	} // VerticalRecursiveGaussianGRAYF32_kernel

	.entry HorizontalRecursiveGaussianGrayF32_kernel (
		.param .u64 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pOut,
		.param .u64 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pIn,
		.param .s32 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_width,
		.param .s32 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_height,
		.param .s32 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f,
		.param .f32 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_n_plus0,
		.param .f32 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_n_plus1,
		.param .f32 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_neg_d_plus1,
		.param .f32 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_neg_d_plus2,
		.param .f32 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_n_minus1,
		.param .f32 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_n_minus2,
		.param .f32 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_neg_d_minus1,
		.param .f32 __cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_neg_d_minus2)
	{
	.reg .u32 %r<200>;
	.reg .u64 %rd<145>;
	.reg .f32 %f<55>;
	.reg .pred %p<29>;
	.shared .align 4 .b8 __cuda_smem124[4224];
	.loc	16	94	0
$LBB1_HorizontalRecursiveGaussianGrayF32_kernel:
	.loc	16	113	0
	ld.param.s32 	%r1, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_width];
	mov.u32 	%r2, 0;
	setp.le.s32 	%p1, %r1, %r2;
	mov.s32 	%r3, 0;
	mov.u64 	%rd1, __cuda_smem124;
	@%p1 bra 	$Lt_1_37634;
	add.s32 	%r4, %r1, 31;
	shr.s32 	%r5, %r4, 31;
	mov.s32 	%r6, 31;
	and.b32 	%r7, %r5, %r6;
	add.s32 	%r8, %r7, %r4;
	shr.s32 	%r9, %r8, 5;
	cvt.s32.u16 	%r10, %ctaid.y;
	mov.s32 	%r11, 32;
	mul24.lo.s32 	%r12, %r10, %r11;
	cvt.u32.u16 	%r13, %tid.y;
	mov.u32 	%r14, 0;
	setp.eq.u32 	%p2, %r13, %r14;
	add.u32 	%r15, %r12, %r13;
	ld.param.s32 	%r16, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_height];
	setp.gt.s32 	%p3, %r16, %r15;
	add.s32 	%r17, %r15, 8;
	add.s32 	%r18, %r15, 16;
	add.s32 	%r19, %r15, 24;
	setp.lt.s32 	%p4, %r17, %r16;
	setp.lt.s32 	%p5, %r18, %r16;
	setp.lt.s32 	%p6, %r19, %r16;
	mov.f32 	%f1, 0f00000000;     	// 0
	mov.f32 	%f2, 0f00000000;     	// 0
	mov.f32 	%f3, 0f00000000;     	// 0
	mov.s32 	%r20, %r9;
$Lt_1_16642:
 //<loop> Loop body line 113, nesting depth: 1, estimated iterations: unknown
	.loc	16	123	0
	@!%p3 bra 	$Lt_1_17410;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r22, %r3, %r21;
	setp.le.u32 	%p7, %r1, %r22;
	@%p7 bra 	$Lt_1_17666;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	131	0
	mov.s32 	%r23, 33;
	mul24.lo.s32 	%r24, %r13, %r23;
	ld.param.u64 	%rd2, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pIn];
	ld.param.s32 	%r25, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	mul.lo.s32 	%r26, %r25, %r15;
	add.s32 	%r27, %r3, %r26;
	add.u32 	%r28, %r21, %r27;
	cvt.u64.s32 	%rd3, %r28;
	mul.lo.u64 	%rd4, %rd3, 4;
	add.u64 	%rd5, %rd2, %rd4;
	ld.global.f32 	%f4, [%rd5+0];
	add.u32 	%r29, %r24, %r21;
	cvt.u64.s32 	%rd6, %r29;
	mul.lo.u64 	%rd7, %rd6, 4;
	add.u64 	%rd8, %rd1, %rd7;
	st.shared.f32 	[%rd8+0], %f4;
	bra.uni 	$Lt_1_17410;
$Lt_1_17666:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	133	0
	mov.s32 	%r30, 33;
	mul24.lo.s32 	%r24, %r13, %r30;
	mov.f32 	%f5, 0f00000000;     	// 0
	add.u32 	%r31, %r24, %r21;
	cvt.u64.s32 	%rd9, %r31;
	mul.lo.u64 	%rd10, %rd9, 4;
	add.u64 	%rd11, %rd1, %rd10;
	st.shared.f32 	[%rd11+0], %f5;
$Lt_1_17410:
$Lt_1_16898:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	@!%p4 bra 	$Lt_1_18434;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r32, %r3, %r21;
	setp.le.u32 	%p8, %r1, %r32;
	@%p8 bra 	$Lt_1_18690;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	131	0
	mov.s32 	%r33, 33;
	mul24.lo.s32 	%r24, %r13, %r33;
	ld.param.s32 	%r34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	ld.param.u64 	%rd12, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pIn];
	shl.b32 	%r35, %r34, 3;
	mul.lo.s32 	%r36, %r34, %r15;
	add.s32 	%r37, %r3, %r36;
	add.u32 	%r38, %r21, %r37;
	add.s32 	%r39, %r35, %r38;
	cvt.u64.s32 	%rd13, %r39;
	mul.lo.u64 	%rd14, %rd13, 4;
	add.u64 	%rd15, %rd12, %rd14;
	ld.global.f32 	%f6, [%rd15+0];
	add.u32 	%r40, %r24, %r21;
	cvt.u64.s32 	%rd16, %r40;
	mul.lo.u64 	%rd17, %rd16, 4;
	add.u64 	%rd18, %rd1, %rd17;
	st.shared.f32 	[%rd18+1056], %f6;
	bra.uni 	$Lt_1_18434;
$Lt_1_18690:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	133	0
	mov.s32 	%r41, 33;
	mul24.lo.s32 	%r24, %r13, %r41;
	mov.f32 	%f7, 0f00000000;     	// 0
	add.u32 	%r42, %r24, %r21;
	cvt.u64.s32 	%rd19, %r42;
	mul.lo.u64 	%rd20, %rd19, 4;
	add.u64 	%rd21, %rd1, %rd20;
	st.shared.f32 	[%rd21+1056], %f7;
$Lt_1_18434:
$Lt_1_17922:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	@!%p5 bra 	$Lt_1_19458;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r43, %r3, %r21;
	setp.le.u32 	%p9, %r1, %r43;
	@%p9 bra 	$Lt_1_19714;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	131	0
	ld.param.s32 	%r34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	shl.b32 	%r44, %r34, 3;
	mov.s32 	%r45, 33;
	mul24.lo.s32 	%r24, %r13, %r45;
	ld.param.u64 	%rd22, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pIn];
	mul.lo.s32 	%r46, %r34, %r15;
	add.s32 	%r47, %r3, %r46;
	add.u32 	%r48, %r21, %r47;
	add.s32 	%r49, %r44, %r48;
	add.s32 	%r50, %r44, %r49;
	cvt.u64.s32 	%rd23, %r50;
	mul.lo.u64 	%rd24, %rd23, 4;
	add.u64 	%rd25, %rd22, %rd24;
	ld.global.f32 	%f8, [%rd25+0];
	add.u32 	%r51, %r24, %r21;
	cvt.u64.s32 	%rd26, %r51;
	mul.lo.u64 	%rd27, %rd26, 4;
	add.u64 	%rd28, %rd1, %rd27;
	st.shared.f32 	[%rd28+2112], %f8;
	bra.uni 	$Lt_1_19458;
$Lt_1_19714:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	133	0
	mov.s32 	%r52, 33;
	mul24.lo.s32 	%r24, %r13, %r52;
	mov.f32 	%f9, 0f00000000;     	// 0
	add.u32 	%r53, %r24, %r21;
	cvt.u64.s32 	%rd29, %r53;
	mul.lo.u64 	%rd30, %rd29, 4;
	add.u64 	%rd31, %rd1, %rd30;
	st.shared.f32 	[%rd31+2112], %f9;
$Lt_1_19458:
$Lt_1_18946:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	@!%p6 bra 	$Lt_1_20482;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r54, %r3, %r21;
	setp.le.u32 	%p10, %r1, %r54;
	@%p10 bra 	$Lt_1_20738;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	131	0
	ld.param.s32 	%r34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	shl.b32 	%r44, %r34, 3;
	mov.s32 	%r55, 33;
	mul24.lo.s32 	%r24, %r13, %r55;
	ld.param.u64 	%rd32, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pIn];
	add.s32 	%r56, %r44, %r44;
	mul.lo.s32 	%r57, %r34, %r15;
	add.s32 	%r58, %r3, %r57;
	add.u32 	%r59, %r21, %r58;
	add.s32 	%r60, %r44, %r59;
	add.s32 	%r61, %r56, %r60;
	cvt.u64.s32 	%rd33, %r61;
	mul.lo.u64 	%rd34, %rd33, 4;
	add.u64 	%rd35, %rd32, %rd34;
	ld.global.f32 	%f10, [%rd35+0];
	add.u32 	%r62, %r24, %r21;
	cvt.u64.s32 	%rd36, %r62;
	mul.lo.u64 	%rd37, %rd36, 4;
	add.u64 	%rd38, %rd1, %rd37;
	st.shared.f32 	[%rd38+3168], %f10;
	bra.uni 	$Lt_1_20482;
$Lt_1_20738:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	133	0
	mov.s32 	%r63, 33;
	mul24.lo.s32 	%r24, %r13, %r63;
	mov.f32 	%f11, 0f00000000;    	// 0
	add.u32 	%r64, %r24, %r21;
	cvt.u64.s32 	%rd39, %r64;
	mul.lo.u64 	%rd40, %rd39, 4;
	add.u64 	%rd41, %rd1, %rd40;
	st.shared.f32 	[%rd41+3168], %f11;
$Lt_1_20482:
$Lt_1_19970:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	138	0
	bar.sync 	0;
	@!%p2 bra 	$Lt_1_20994;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	145	0
	cvt.u32.u16 	%r65, %tid.x;
	mov.s32 	%r66, 33;
	mul24.lo.s32 	%r67, %r65, %r66;
	mov.s32 	%r68, %r67;
	add.s32 	%r69, %r67, 32;
	cvt.u64.s32 	%rd42, %r67;
	mul.lo.u64 	%rd43, %rd42, 4;
	add.u64 	%rd44, %rd1, %rd43;
	ld.param.f32 	%f12, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_neg_d_plus2];
	ld.param.f32 	%f13, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_neg_d_plus1];
	ld.param.f32 	%f14, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_n_plus0];
	ld.param.f32 	%f15, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_n_plus1];
$Lt_1_22018:
 //<loop> Loop body line 145, nesting depth: 2, iterations: 32
	.loc	16	148	0
	ld.shared.f32 	%f16, [%rd44+0];
	.loc	16	149	0
	mul.f32 	%f17, %f15, %f3;
	mad.f32 	%f18, %f14, %f16, %f17;
	mad.f32 	%f19, %f13, %f2, %f18;
	mad.f32 	%f20, %f12, %f1, %f19;
	.loc	16	151	0
	mov.f32 	%f3, %f16;
	.loc	16	152	0
	mov.f32 	%f1, %f2;
	.loc	16	153	0
	mov.f32 	%f2, %f20;
	.loc	16	155	0
	st.shared.f32 	[%rd44+0], %f20;
	add.s32 	%r68, %r68, 1;
	add.u64 	%rd44, %rd44, 4;
	setp.ne.s32 	%p11, %r68, %r69;
	@%p11 bra 	$Lt_1_22018;
$Lt_1_20994:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	160	0
	bar.sync 	0;
	.loc	16	164	0
	@!%p3 bra 	$Lt_1_23042;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r70, %r3, %r21;
	setp.le.u32 	%p12, %r1, %r70;
	@%p12 bra 	$Lt_1_23042;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	172	0
	mov.s32 	%r71, 33;
	mul24.lo.s32 	%r72, %r13, %r71;
	add.u32 	%r73, %r21, %r72;
	cvt.u64.s32 	%rd45, %r73;
	mul.lo.u64 	%rd46, %rd45, 4;
	add.u64 	%rd47, %rd1, %rd46;
	ld.shared.f32 	%f21, [%rd47+0];
	ld.param.u64 	%rd48, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pOut];
	ld.param.s32 	%r74, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	mul.lo.s32 	%r75, %r74, %r15;
	add.s32 	%r76, %r3, %r75;
	add.u32 	%r77, %r21, %r76;
	cvt.u64.s32 	%rd49, %r77;
	mul.lo.u64 	%rd50, %rd49, 4;
	add.u64 	%rd51, %rd48, %rd50;
	st.global.f32 	[%rd51+0], %f21;
$Lt_1_23042:
$Lt_1_22530:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	@!%p4 bra 	$Lt_1_24066;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r78, %r3, %r21;
	setp.le.u32 	%p13, %r1, %r78;
	@%p13 bra 	$Lt_1_24066;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	ld.param.s32 	%r34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	mov.s32 	%r79, 33;
	mul24.lo.s32 	%r80, %r13, %r79;
	add.u32 	%r81, %r21, %r80;
	cvt.u64.s32 	%rd52, %r81;
	mul.lo.u64 	%rd53, %rd52, 4;
	add.u64 	%rd54, %rd1, %rd53;
	ld.shared.f32 	%f22, [%rd54+1056];
	ld.param.u64 	%rd55, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pOut];
	shl.b32 	%r82, %r34, 3;
	mul.lo.s32 	%r83, %r34, %r15;
	add.s32 	%r84, %r3, %r83;
	add.u32 	%r85, %r21, %r84;
	add.s32 	%r86, %r82, %r85;
	cvt.u64.s32 	%rd56, %r86;
	mul.lo.u64 	%rd57, %rd56, 4;
	add.u64 	%rd58, %rd55, %rd57;
	st.global.f32 	[%rd58+0], %f22;
$Lt_1_24066:
$Lt_1_23554:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	@!%p5 bra 	$Lt_1_25090;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r87, %r3, %r21;
	setp.le.u32 	%p14, %r1, %r87;
	@%p14 bra 	$Lt_1_25090;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	ld.param.s32 	%r34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	shl.b32 	%r44, %r34, 3;
	mov.s32 	%r88, 33;
	mul24.lo.s32 	%r89, %r13, %r88;
	add.u32 	%r90, %r21, %r89;
	cvt.u64.s32 	%rd59, %r90;
	mul.lo.u64 	%rd60, %rd59, 4;
	add.u64 	%rd61, %rd1, %rd60;
	ld.shared.f32 	%f23, [%rd61+2112];
	ld.param.u64 	%rd62, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pOut];
	mul.lo.s32 	%r91, %r34, %r15;
	add.s32 	%r92, %r3, %r91;
	add.u32 	%r93, %r21, %r92;
	add.s32 	%r94, %r44, %r93;
	add.s32 	%r95, %r44, %r94;
	cvt.u64.s32 	%rd63, %r95;
	mul.lo.u64 	%rd64, %rd63, 4;
	add.u64 	%rd65, %rd62, %rd64;
	st.global.f32 	[%rd65+0], %f23;
$Lt_1_25090:
$Lt_1_24578:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	@!%p6 bra 	$Lt_1_26114;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r96, %r3, %r21;
	setp.le.u32 	%p15, %r1, %r96;
	@%p15 bra 	$Lt_1_26114;
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	ld.param.s32 	%r34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	shl.b32 	%r44, %r34, 3;
	mov.s32 	%r97, 33;
	mul24.lo.s32 	%r98, %r13, %r97;
	add.u32 	%r99, %r21, %r98;
	cvt.u64.s32 	%rd66, %r99;
	mul.lo.u64 	%rd67, %rd66, 4;
	add.u64 	%rd68, %rd1, %rd67;
	ld.shared.f32 	%f24, [%rd68+3168];
	ld.param.u64 	%rd69, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pOut];
	add.s32 	%r100, %r44, %r44;
	mul.lo.s32 	%r101, %r34, %r15;
	add.s32 	%r102, %r3, %r101;
	add.u32 	%r103, %r21, %r102;
	add.s32 	%r104, %r44, %r103;
	add.s32 	%r105, %r100, %r104;
	cvt.u64.s32 	%rd70, %r105;
	mul.lo.u64 	%rd71, %rd70, 4;
	add.u64 	%rd72, %rd69, %rd71;
	st.global.f32 	[%rd72+0], %f24;
$Lt_1_26114:
$Lt_1_25602:
 //<loop> Part of loop body line 113, head labeled $Lt_1_16642
	.loc	16	178	0
	bar.sync 	0;
	add.s32 	%r3, %r3, 32;
	setp.gt.s32 	%p16, %r1, %r3;
	@%p16 bra 	$Lt_1_16642;
	bra.uni 	$Lt_1_16130;
$Lt_1_37634:
$Lt_1_16130:
	mov.u32 	%r106, 0;
	setp.le.s32 	%p17, %r3, %r106;
	@%p17 bra 	$Lt_1_26882;
	add.s32 	%r107, %r3, 31;
	shr.s32 	%r108, %r107, 31;
	mov.s32 	%r109, 31;
	and.b32 	%r110, %r108, %r109;
	add.s32 	%r111, %r110, %r107;
	shr.s32 	%r112, %r111, 5;
	cvt.s32.u16 	%r113, %ctaid.y;
	mov.s32 	%r114, 32;
	mul24.lo.s32 	%r115, %r113, %r114;
	cvt.u32.u16 	%r13, %tid.y;
	mov.u32 	%r116, 0;
	setp.eq.u32 	%p2, %r13, %r116;
	add.u32 	%r15, %r115, %r13;
	ld.param.s32 	%r117, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_height];
	setp.gt.s32 	%p3, %r117, %r15;
	add.s32 	%r118, %r15, 8;
	add.s32 	%r119, %r15, 16;
	add.s32 	%r120, %r15, 24;
	setp.lt.s32 	%p4, %r118, %r117;
	setp.lt.s32 	%p5, %r119, %r117;
	setp.lt.s32 	%p6, %r120, %r117;
	mov.f32 	%f25, 0f00000000;    	// 0
	mov.f32 	%f1, 0f00000000;     	// 0
	mov.f32 	%f2, 0f00000000;     	// 0
	mov.f32 	%f3, 0f00000000;     	// 0
	mov.s32 	%r121, %r112;
$Lt_1_27394:
 //<loop> Loop body line 178, nesting depth: 1, estimated iterations: unknown
	.loc	16	191	0
	sub.s32 	%r3, %r3, 32;
	.loc	16	194	0
	@!%p3 bra 	$Lt_1_28162;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r122, %r3, %r21;
	setp.le.u32 	%p18, %r1, %r122;
	@%p18 bra 	$Lt_1_28418;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	202	0
	mov.s32 	%r123, 33;
	mul24.lo.s32 	%r24, %r13, %r123;
	ld.param.u64 	%rd73, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pIn];
	ld.param.s32 	%r124, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	mul.lo.s32 	%r125, %r124, %r15;
	add.s32 	%r126, %r3, %r125;
	add.u32 	%r127, %r21, %r126;
	cvt.u64.s32 	%rd74, %r127;
	mul.lo.u64 	%rd75, %rd74, 4;
	add.u64 	%rd76, %rd73, %rd75;
	ld.global.f32 	%f26, [%rd76+0];
	add.u32 	%r128, %r24, %r21;
	cvt.u64.s32 	%rd77, %r128;
	mul.lo.u64 	%rd78, %rd77, 4;
	add.u64 	%rd79, %rd1, %rd78;
	st.shared.f32 	[%rd79+0], %f26;
	bra.uni 	$Lt_1_28162;
$Lt_1_28418:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	204	0
	mov.s32 	%r129, 33;
	mul24.lo.s32 	%r24, %r13, %r129;
	mov.f32 	%f27, 0f00000000;    	// 0
	add.u32 	%r130, %r24, %r21;
	cvt.u64.s32 	%rd80, %r130;
	mul.lo.u64 	%rd81, %rd80, 4;
	add.u64 	%rd82, %rd1, %rd81;
	st.shared.f32 	[%rd82+0], %f27;
$Lt_1_28162:
$Lt_1_27650:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	@!%p4 bra 	$Lt_1_29186;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r131, %r3, %r21;
	setp.le.u32 	%p19, %r1, %r131;
	@%p19 bra 	$Lt_1_29442;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	202	0
	mov.s32 	%r132, 33;
	mul24.lo.s32 	%r24, %r13, %r132;
	ld.param.s32 	%r34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	ld.param.u64 	%rd83, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pIn];
	shl.b32 	%r133, %r34, 3;
	mul.lo.s32 	%r134, %r34, %r15;
	add.s32 	%r135, %r3, %r134;
	add.u32 	%r136, %r21, %r135;
	add.s32 	%r137, %r133, %r136;
	cvt.u64.s32 	%rd84, %r137;
	mul.lo.u64 	%rd85, %rd84, 4;
	add.u64 	%rd86, %rd83, %rd85;
	ld.global.f32 	%f28, [%rd86+0];
	add.u32 	%r138, %r24, %r21;
	cvt.u64.s32 	%rd87, %r138;
	mul.lo.u64 	%rd88, %rd87, 4;
	add.u64 	%rd89, %rd1, %rd88;
	st.shared.f32 	[%rd89+1056], %f28;
	bra.uni 	$Lt_1_29186;
$Lt_1_29442:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	204	0
	mov.s32 	%r139, 33;
	mul24.lo.s32 	%r24, %r13, %r139;
	mov.f32 	%f29, 0f00000000;    	// 0
	add.u32 	%r140, %r24, %r21;
	cvt.u64.s32 	%rd90, %r140;
	mul.lo.u64 	%rd91, %rd90, 4;
	add.u64 	%rd92, %rd1, %rd91;
	st.shared.f32 	[%rd92+1056], %f29;
$Lt_1_29186:
$Lt_1_28674:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	@!%p5 bra 	$Lt_1_30210;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r141, %r3, %r21;
	setp.le.u32 	%p20, %r1, %r141;
	@%p20 bra 	$Lt_1_30466;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	202	0
	ld.param.s32 	%r34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	shl.b32 	%r44, %r34, 3;
	mov.s32 	%r142, 33;
	mul24.lo.s32 	%r24, %r13, %r142;
	ld.param.u64 	%rd93, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pIn];
	mul.lo.s32 	%r143, %r34, %r15;
	add.s32 	%r144, %r3, %r143;
	add.u32 	%r145, %r21, %r144;
	add.s32 	%r146, %r44, %r145;
	add.s32 	%r147, %r44, %r146;
	cvt.u64.s32 	%rd94, %r147;
	mul.lo.u64 	%rd95, %rd94, 4;
	add.u64 	%rd96, %rd93, %rd95;
	ld.global.f32 	%f30, [%rd96+0];
	add.u32 	%r148, %r24, %r21;
	cvt.u64.s32 	%rd97, %r148;
	mul.lo.u64 	%rd98, %rd97, 4;
	add.u64 	%rd99, %rd1, %rd98;
	st.shared.f32 	[%rd99+2112], %f30;
	bra.uni 	$Lt_1_30210;
$Lt_1_30466:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	204	0
	mov.s32 	%r149, 33;
	mul24.lo.s32 	%r24, %r13, %r149;
	mov.f32 	%f31, 0f00000000;    	// 0
	add.u32 	%r150, %r24, %r21;
	cvt.u64.s32 	%rd100, %r150;
	mul.lo.u64 	%rd101, %rd100, 4;
	add.u64 	%rd102, %rd1, %rd101;
	st.shared.f32 	[%rd102+2112], %f31;
$Lt_1_30210:
$Lt_1_29698:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	@!%p6 bra 	$Lt_1_31234;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r151, %r3, %r21;
	setp.le.u32 	%p21, %r1, %r151;
	@%p21 bra 	$Lt_1_31490;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	202	0
	ld.param.s32 	%r34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	shl.b32 	%r44, %r34, 3;
	mov.s32 	%r152, 33;
	mul24.lo.s32 	%r24, %r13, %r152;
	ld.param.u64 	%rd103, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pIn];
	add.s32 	%r153, %r44, %r44;
	mul.lo.s32 	%r154, %r34, %r15;
	add.s32 	%r155, %r3, %r154;
	add.u32 	%r156, %r21, %r155;
	add.s32 	%r157, %r44, %r156;
	add.s32 	%r158, %r153, %r157;
	cvt.u64.s32 	%rd104, %r158;
	mul.lo.u64 	%rd105, %rd104, 4;
	add.u64 	%rd106, %rd103, %rd105;
	ld.global.f32 	%f32, [%rd106+0];
	add.u32 	%r159, %r24, %r21;
	cvt.u64.s32 	%rd107, %r159;
	mul.lo.u64 	%rd108, %rd107, 4;
	add.u64 	%rd109, %rd1, %rd108;
	st.shared.f32 	[%rd109+3168], %f32;
	bra.uni 	$Lt_1_31234;
$Lt_1_31490:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	204	0
	mov.s32 	%r160, 33;
	mul24.lo.s32 	%r24, %r13, %r160;
	mov.f32 	%f33, 0f00000000;    	// 0
	add.u32 	%r161, %r24, %r21;
	cvt.u64.s32 	%rd110, %r161;
	mul.lo.u64 	%rd111, %rd110, 4;
	add.u64 	%rd112, %rd1, %rd111;
	st.shared.f32 	[%rd112+3168], %f33;
$Lt_1_31234:
$Lt_1_30722:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	209	0
	bar.sync 	0;
	@!%p2 bra 	$Lt_1_31746;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	214	0
	cvt.u32.u16 	%r162, %tid.x;
	mov.s32 	%r163, 33;
	mul24.lo.s32 	%r164, %r162, %r163;
	add.s32 	%r165, %r164, 31;
	cvt.u64.s32 	%rd113, %r165;
	mul.lo.u64 	%rd114, %rd113, 4;
	add.u64 	%rd115, %rd1, %rd114;
	ld.param.f32 	%f34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_neg_d_minus2];
	ld.param.f32 	%f35, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_neg_d_minus1];
	ld.param.f32 	%f36, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_n_minus1];
	ld.param.f32 	%f37, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_n_minus2];
	mov.s32 	%r166, 31;
$Lt_1_32770:
 //<loop> Loop body line 214, nesting depth: 1, iterations: 32
	.loc	16	217	0
	mul.f32 	%f38, %f37, %f25;
	mad.f32 	%f39, %f36, %f3, %f38;
	mad.f32 	%f40, %f35, %f2, %f39;
	mad.f32 	%f41, %f34, %f1, %f40;
	.loc	16	219	0
	mov.f32 	%f25, %f3;
	.loc	16	220	0
	ld.shared.f32 	%f3, [%rd115+0];
	.loc	16	221	0
	mov.f32 	%f1, %f2;
	.loc	16	222	0
	mov.f32 	%f2, %f41;
	.loc	16	224	0
	st.shared.f32 	[%rd115+0], %f41;
	.loc	16	225	0
	sub.u64 	%rd115, %rd115, 4;
	sub.s32 	%r166, %r166, 1;
	mov.u32 	%r167, -1;
	setp.ne.s32 	%p22, %r166, %r167;
	@%p22 bra 	$Lt_1_32770;
$Lt_1_31746:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	228	0
	bar.sync 	0;
	.loc	16	231	0
	@!%p3 bra 	$Lt_1_33794;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r168, %r3, %r21;
	setp.le.u32 	%p23, %r1, %r168;
	@%p23 bra 	$Lt_1_33794;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	239	0
	ld.param.s32 	%r169, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	mul.lo.s32 	%r170, %r169, %r15;
	ld.param.u64 	%rd116, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pOut];
	add.s32 	%r171, %r170, %r3;
	add.u32 	%r172, %r21, %r171;
	cvt.u64.s32 	%rd117, %r172;
	mul.lo.u64 	%rd118, %rd117, 4;
	add.u64 	%rd119, %rd116, %rd118;
	ld.global.f32 	%f42, [%rd119+0];
	mov.s32 	%r173, 33;
	mul24.lo.s32 	%r174, %r13, %r173;
	add.u32 	%r175, %r21, %r174;
	cvt.u64.s32 	%rd120, %r175;
	mul.lo.u64 	%rd121, %rd120, 4;
	add.u64 	%rd122, %rd1, %rd121;
	ld.shared.f32 	%f43, [%rd122+0];
	add.f32 	%f44, %f42, %f43;
	st.global.f32 	[%rd119+0], %f44;
$Lt_1_33794:
$Lt_1_33282:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	@!%p4 bra 	$Lt_1_34818;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r176, %r3, %r21;
	setp.le.u32 	%p24, %r1, %r176;
	@%p24 bra 	$Lt_1_34818;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	ld.param.s32 	%r34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	shl.b32 	%r44, %r34, 3;
	mul.lo.s32 	%r170, %r34, %r15;
	add.s32 	%r177, %r170, %r3;
	ld.param.u64 	%rd123, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pOut];
	add.u32 	%r178, %r177, %r21;
	add.s32 	%r179, %r44, %r178;
	cvt.u64.s32 	%rd124, %r179;
	mul.lo.u64 	%rd125, %rd124, 4;
	add.u64 	%rd126, %rd123, %rd125;
	ld.global.f32 	%f45, [%rd126+0];
	mov.s32 	%r180, 33;
	mul24.lo.s32 	%r181, %r13, %r180;
	add.u32 	%r182, %r21, %r181;
	cvt.u64.s32 	%rd127, %r182;
	mul.lo.u64 	%rd128, %rd127, 4;
	add.u64 	%rd129, %rd1, %rd128;
	ld.shared.f32 	%f46, [%rd129+1056];
	add.f32 	%f47, %f45, %f46;
	st.global.f32 	[%rd126+0], %f47;
$Lt_1_34818:
$Lt_1_34306:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	@!%p5 bra 	$Lt_1_35842;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r183, %r3, %r21;
	setp.le.u32 	%p25, %r1, %r183;
	@%p25 bra 	$Lt_1_35842;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	ld.param.s32 	%r34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	shl.b32 	%r44, %r34, 3;
	mul.lo.s32 	%r170, %r34, %r15;
	add.s32 	%r177, %r170, %r3;
	add.u32 	%r184, %r177, %r21;
	ld.param.u64 	%rd130, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pOut];
	add.s32 	%r185, %r44, %r184;
	add.s32 	%r186, %r44, %r185;
	cvt.u64.s32 	%rd131, %r186;
	mul.lo.u64 	%rd132, %rd131, 4;
	add.u64 	%rd133, %rd130, %rd132;
	ld.global.f32 	%f48, [%rd133+0];
	mov.s32 	%r187, 33;
	mul24.lo.s32 	%r188, %r13, %r187;
	add.u32 	%r189, %r21, %r188;
	cvt.u64.s32 	%rd134, %r189;
	mul.lo.u64 	%rd135, %rd134, 4;
	add.u64 	%rd136, %rd1, %rd135;
	ld.shared.f32 	%f49, [%rd136+2112];
	add.f32 	%f50, %f48, %f49;
	st.global.f32 	[%rd133+0], %f50;
$Lt_1_35842:
$Lt_1_35330:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	@!%p6 bra 	$Lt_1_36866;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	cvt.u32.u16 	%r21, %tid.x;
	add.u32 	%r190, %r3, %r21;
	setp.le.u32 	%p26, %r1, %r190;
	@%p26 bra 	$Lt_1_36866;
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	ld.param.s32 	%r34, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pitch_f];
	shl.b32 	%r44, %r34, 3;
	mul.lo.s32 	%r170, %r34, %r15;
	add.s32 	%r177, %r170, %r3;
	add.u32 	%r191, %r177, %r21;
	add.s32 	%r192, %r44, %r191;
	ld.param.u64 	%rd137, [__cudaparm_HorizontalRecursiveGaussianGrayF32_kernel_pOut];
	add.s32 	%r193, %r192, %r44;
	add.s32 	%r194, %r44, %r193;
	cvt.u64.s32 	%rd138, %r194;
	mul.lo.u64 	%rd139, %rd138, 4;
	add.u64 	%rd140, %rd137, %rd139;
	ld.global.f32 	%f51, [%rd140+0];
	mov.s32 	%r195, 33;
	mul24.lo.s32 	%r196, %r13, %r195;
	add.u32 	%r197, %r21, %r196;
	cvt.u64.s32 	%rd141, %r197;
	mul.lo.u64 	%rd142, %rd141, 4;
	add.u64 	%rd143, %rd1, %rd142;
	ld.shared.f32 	%f52, [%rd143+3168];
	add.f32 	%f53, %f51, %f52;
	st.global.f32 	[%rd140+0], %f53;
$Lt_1_36866:
$Lt_1_36354:
 //<loop> Part of loop body line 178, head labeled $Lt_1_27394
	.loc	16	245	0
	bar.sync 	0;
	mov.u32 	%r198, 0;
	setp.gt.s32 	%p27, %r3, %r198;
	@%p27 bra 	$Lt_1_27394;
$Lt_1_26882:
	.loc	16	247	0
	exit;
$LDWend_HorizontalRecursiveGaussianGrayF32_kernel:
	} // HorizontalRecursiveGaussianGrayF32_kernel

