	.version 2.3
	.target sm_20
	.address_size 64
	// compiled with C:\mightykilt64/shared/adobe/MediaCore//External/3rdParty/NVIDIA/CUDA/win/bin/../open64/lib//be.exe
	// nvopencc 4.0 built on 2011-05-13

	.visible .func (.param .f32 __cudaretf__Z9ReadAlphaPK6float4i17DevicePixelFormat) _Z9ReadAlphaPK6float4i17DevicePixelFormat (.param .u64 __cudaparmf1__Z9ReadAlphaPK6float4i17DevicePixelFormat, .param .s32 __cudaparmf2__Z9ReadAlphaPK6float4i17DevicePixelFormat, .param .u32 __cudaparmf3__Z9ReadAlphaPK6float4i17DevicePixelFormat)

	.visible .func _Z6MinRowiP6float2Pf (.param .s32 __cudaparmf1__Z6MinRowiP6float2Pf, .param .u64 __cudaparmf2__Z6MinRowiP6float2Pf, .param .u64 __cudaparmf3__Z6MinRowiP6float2Pf)

	.visible .func (.param .align 8 .b8 __cudaretf__Z9MinColumniP6float2[8]) _Z9MinColumniP6float2 (.param .s32 __cudaparmf1__Z9MinColumniP6float2, .param .u64 __cudaparmf2__Z9MinColumniP6float2)

	.visible .func _Z6SumRowiP6float2Pf (.param .s32 __cudaparmf1__Z6SumRowiP6float2Pf, .param .u64 __cudaparmf2__Z6SumRowiP6float2Pf, .param .u64 __cudaparmf3__Z6SumRowiP6float2Pf)

	.visible .func (.param .f32 __cudaretf__Z9SumColumniP6float2f) _Z9SumColumniP6float2f (.param .s32 __cudaparmf1__Z9SumColumniP6float2f, .param .u64 __cudaparmf2__Z9SumColumniP6float2f, .param .f32 __cudaparmf3__Z9SumColumniP6float2f)

	//-----------------------------------------------------------
	// Compiling C:/Users/dvaeng/AppData/Local/Temp/tmpxft_0000307c_00000000-11_UltraKey.cpp3.i (C:/Users/dvaeng/AppData/Local/Temp/ccBI#.a11792)
	//-----------------------------------------------------------

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

	.file	1	"C:/Users/dvaeng/AppData/Local/Temp/tmpxft_0000307c_00000000-10_UltraKey.cudafe2.gpu"
	.file	2	"C:\mightykilt64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/DevicePixelFormat.h"
	.file	3	"C:\mightykilt64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/PixelRGB.h"
	.file	4	"C:\mightykilt64\shared\adobe\MediaCore\Renderers\RendererGPU\Inc\Effects/UltraKey.h"
	.file	5	"c:\program files (x86)\microsoft visual studio 10.0\vc\include\codeanalysis\sourceannotations.h"
	.file	6	"C:\mightykilt64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\include\crt/device_runtime.h"
	.file	7	"C:\mightykilt64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\include\host_defines.h"
	.file	8	"C:\mightykilt64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\include\builtin_types.h"
	.file	9	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\device_types.h"
	.file	10	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\driver_types.h"
	.file	11	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\surface_types.h"
	.file	12	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\texture_types.h"
	.file	13	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\vector_types.h"
	.file	14	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\builtin_types.h"
	.file	15	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\host_defines.h"
	.file	16	"C:\mightykilt64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\include\device_launch_parameters.h"
	.file	17	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\crt\storage_class.h"
	.file	18	"C:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\INCLUDE\time.h"
	.file	19	"C:/mightykilt64/shared/adobe/MediaCore/Renderers/RendererGPU/Src/Effects/UltraKey.cu"
	.file	20	"C:\mightykilt64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\include\common_functions.h"
	.file	21	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\math_functions.h"
	.file	22	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\math_constants.h"
	.file	23	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\device_functions.h"
	.file	24	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\sm_11_atomic_functions.h"
	.file	25	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\sm_12_atomic_functions.h"
	.file	26	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\sm_13_double_functions.h"
	.file	27	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\sm_20_atomic_functions.h"
	.file	28	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\sm_20_intrinsics.h"
	.file	29	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\surface_functions.h"
	.file	30	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\texture_fetch_functions.h"
	.file	31	"c:\mightykilt64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\math_functions_dbl_ptx3.h"


	.visible .func (.param .f32 __cudaretf__Z9ReadAlphaPK6float4i17DevicePixelFormat) _Z9ReadAlphaPK6float4i17DevicePixelFormat (.param .u64 __cudaparmf1__Z9ReadAlphaPK6float4i17DevicePixelFormat, .param .s32 __cudaparmf2__Z9ReadAlphaPK6float4i17DevicePixelFormat, .param .u32 __cudaparmf3__Z9ReadAlphaPK6float4i17DevicePixelFormat)
	{
	.reg .u32 %r<8>;
	.reg .u64 %rd<9>;
	.reg .f32 %f<4>;
	.reg .pred %p<3>;
	.loc	19	19	0
$LDWbegin__Z9ReadAlphaPK6float4i17DevicePixelFormat:
	ld.param.u64 	%rd1, [__cudaparmf1__Z9ReadAlphaPK6float4i17DevicePixelFormat];
	mov.s64 	%rd2, %rd1;
	ld.param.u32 	%r1, [__cudaparmf2__Z9ReadAlphaPK6float4i17DevicePixelFormat];
	mov.s32 	%r2, %r1;
	ld.param.u32 	%r3, [__cudaparmf3__Z9ReadAlphaPK6float4i17DevicePixelFormat];
	mov.s32 	%r4, %r3;
	cvt.s64.s32 	%rd3, %r2;
	mov.u32 	%r5, 0;
	setp.ne.s32 	%p1, %r4, %r5;
	@%p1 bra 	$Lt_0_1026;
	mul.lo.u64 	%rd4, %rd3, 8;
	add.u64 	%rd5, %rd2, %rd4;
	ld.u16 	%r6, [%rd5+6];
	{ .reg .b32 %b1;
	mov.b32		%b1, %r6;
	cvt.ftz.f32.f16	%f1, %b1; }
	bra.uni 	$LBB4__Z9ReadAlphaPK6float4i17DevicePixelFormat;
$Lt_0_1026:
	mul.lo.u64 	%rd6, %rd3, 16;
	add.u64 	%rd7, %rd2, %rd6;
	ld.f32 	%f1, [%rd7+12];
$LBB4__Z9ReadAlphaPK6float4i17DevicePixelFormat:
	mov.f32 	%f2, %f1;
	st.param.f32 	[__cudaretf__Z9ReadAlphaPK6float4i17DevicePixelFormat], %f2;
	ret;
$LDWend__Z9ReadAlphaPK6float4i17DevicePixelFormat:
	} // _Z9ReadAlphaPK6float4i17DevicePixelFormat

	.visible .func _Z6MinRowiP6float2Pf (.param .s32 __cudaparmf1__Z6MinRowiP6float2Pf, .param .u64 __cudaparmf2__Z6MinRowiP6float2Pf, .param .u64 __cudaparmf3__Z6MinRowiP6float2Pf)
	{
	.reg .u32 %r<16>;
	.reg .u64 %rd<13>;
	.reg .f32 %f<10>;
	.reg .pred %p<4>;
	.loc	19	88	0
$LDWbegin__Z6MinRowiP6float2Pf:
	ld.param.u32 	%r1, [__cudaparmf1__Z6MinRowiP6float2Pf];
	mov.s32 	%r2, %r1;
	ld.param.u64 	%rd1, [__cudaparmf2__Z6MinRowiP6float2Pf];
	mov.s64 	%rd2, %rd1;
	ld.param.u64 	%rd3, [__cudaparmf3__Z6MinRowiP6float2Pf];
	mov.s64 	%rd4, %rd3;
	neg.s32 	%r3, %r2;
	cvt.s64.s32 	%rd5, %r3;
	mul.wide.s32 	%rd6, %r3, 4;
	add.u64 	%rd7, %rd4, %rd6;
	ld.f32 	%f1, [%rd7+0];
	mul.lo.s32 	%r4, %r2, 2;
	mov.u32 	%r5, 0;
	setp.le.s32 	%p1, %r4, %r5;
	@%p1 bra 	$Lt_1_1282;
	add.s32 	%r6, %r4, 1;
	shr.s32 	%r7, %r6, 31;
	mov.s32 	%r8, 1;
	and.b32 	%r9, %r7, %r8;
	add.s32 	%r10, %r9, %r6;
	shr.s32 	%r11, %r10, 1;
	mov.s32 	%r12, %r3;
	sub.s32 	%r13, %r4, %r2;
	mov.s64 	%rd8, %rd7;
	mov.s32 	%r14, %r11;
$Lt_1_1794:
 //<loop> Loop body line 88, nesting depth: 1, estimated iterations: unknown
	ld.f32 	%f2, [%rd8+4];
	min.ftz.f32 	%f3, %f2, %f1;
	ld.f32 	%f4, [%rd8+8];
	min.ftz.f32 	%f1, %f4, %f3;
	add.s32 	%r12, %r12, 2;
	add.u64 	%rd8, %rd8, 8;
	setp.lt.s32 	%p2, %r12, %r13;
	@%p2 bra 	$Lt_1_1794;
$Lt_1_1282:
	ld.f32 	%f5, [%rd7+-4];
	cvt.s64.s32 	%rd9, %r2;
	mul.wide.s32 	%rd10, %r2, 4;
	add.u64 	%rd11, %rd4, %rd10;
	ld.f32 	%f6, [%rd11+4];
	min.ftz.f32 	%f7, %f5, %f6;
	min.ftz.f32 	%f8, %f7, %f1;
	st.v2.f32 	[%rd2+0], {%f8,%f1};
	ret;
$LDWend__Z6MinRowiP6float2Pf:
	} // _Z6MinRowiP6float2Pf

	.visible .func (.param .align 8 .b8 __cudaretf__Z9MinColumniP6float2[8]) _Z9MinColumniP6float2 (.param .s32 __cudaparmf1__Z9MinColumniP6float2, .param .u64 __cudaparmf2__Z9MinColumniP6float2)
	{
	.reg .u32 %r<16>;
	.reg .u64 %rd<11>;
	.reg .f32 %f<15>;
	.reg .pred %p<4>;
$LDWbegin__Z9MinColumniP6float2:
	ld.param.u32 	%r1, [__cudaparmf1__Z9MinColumniP6float2];
	mov.s32 	%r2, %r1;
	ld.param.u64 	%rd1, [__cudaparmf2__Z9MinColumniP6float2];
	mov.s64 	%rd2, %rd1;
	neg.s32 	%r3, %r2;
	cvt.s64.s32 	%rd3, %r3;
	mul.wide.s32 	%rd4, %r3, 8;
	add.u64 	%rd5, %rd2, %rd4;
	ld.f32 	%f1, [%rd5+-8];
	ld.v2.f32 	{%f2,%f3}, [%rd5+0];
	min.ftz.f32 	%f4, %f1, %f2;
	mul.lo.s32 	%r4, %r2, 2;
	mov.u32 	%r5, 0;
	setp.le.s32 	%p1, %r4, %r5;
	@%p1 bra 	$Lt_2_1282;
	add.s32 	%r6, %r4, 1;
	shr.s32 	%r7, %r6, 31;
	mov.s32 	%r8, 1;
	and.b32 	%r9, %r7, %r8;
	add.s32 	%r10, %r9, %r6;
	shr.s32 	%r11, %r10, 1;
	mov.s32 	%r12, %r3;
	sub.s32 	%r13, %r4, %r2;
	mov.s64 	%rd6, %rd5;
	mov.s32 	%r14, %r11;
$Lt_2_1794:
 //<loop> Loop body line 88, nesting depth: 1, estimated iterations: unknown
	ld.v2.f32 	{%f5,%f6}, [%rd6+8];
	min.ftz.f32 	%f7, %f5, %f4;
	min.ftz.f32 	%f8, %f6, %f3;
	ld.v2.f32 	{%f9,%f10}, [%rd6+16];
	min.ftz.f32 	%f4, %f9, %f7;
	min.ftz.f32 	%f3, %f10, %f8;
	add.s32 	%r12, %r12, 2;
	add.u64 	%rd6, %rd6, 16;
	setp.lt.s32 	%p2, %r12, %r13;
	@%p2 bra 	$Lt_2_1794;
$Lt_2_1282:
	cvt.s64.s32 	%rd7, %r2;
	mul.wide.s32 	%rd8, %r2, 8;
	add.u64 	%rd9, %rd2, %rd8;
	ld.f32 	%f11, [%rd9+8];
	min.ftz.f32 	%f4, %f11, %f4;
	mov.f32 	%f12, %f4;
	st.param.f32 	[__cudaretf__Z9MinColumniP6float2+0], %f12;
	mov.f32 	%f13, %f3;
	st.param.f32 	[__cudaretf__Z9MinColumniP6float2+4], %f13;
	ret;
$LDWend__Z9MinColumniP6float2:
	} // _Z9MinColumniP6float2

	.visible .func _Z6SumRowiP6float2Pf (.param .s32 __cudaparmf1__Z6SumRowiP6float2Pf, .param .u64 __cudaparmf2__Z6SumRowiP6float2Pf, .param .u64 __cudaparmf3__Z6SumRowiP6float2Pf)
	{
	.reg .u32 %r<16>;
	.reg .u64 %rd<13>;
	.reg .f32 %f<10>;
	.reg .pred %p<4>;
$LDWbegin__Z6SumRowiP6float2Pf:
	ld.param.u32 	%r1, [__cudaparmf1__Z6SumRowiP6float2Pf];
	mov.s32 	%r2, %r1;
	ld.param.u64 	%rd1, [__cudaparmf2__Z6SumRowiP6float2Pf];
	mov.s64 	%rd2, %rd1;
	ld.param.u64 	%rd3, [__cudaparmf3__Z6SumRowiP6float2Pf];
	mov.s64 	%rd4, %rd3;
	neg.s32 	%r3, %r2;
	cvt.s64.s32 	%rd5, %r3;
	mul.wide.s32 	%rd6, %r3, 4;
	add.u64 	%rd7, %rd4, %rd6;
	ld.f32 	%f1, [%rd7+0];
	mul.lo.s32 	%r4, %r2, 2;
	mov.u32 	%r5, 0;
	setp.le.s32 	%p1, %r4, %r5;
	@%p1 bra 	$Lt_3_1282;
	add.s32 	%r6, %r4, 1;
	shr.s32 	%r7, %r6, 31;
	mov.s32 	%r8, 1;
	and.b32 	%r9, %r7, %r8;
	add.s32 	%r10, %r9, %r6;
	shr.s32 	%r11, %r10, 1;
	mov.s32 	%r12, %r3;
	sub.s32 	%r13, %r4, %r2;
	mov.s64 	%rd8, %rd7;
	mov.s32 	%r14, %r11;
$Lt_3_1794:
 //<loop> Loop body line 88, nesting depth: 1, estimated iterations: unknown
	ld.f32 	%f2, [%rd8+4];
	add.ftz.f32 	%f3, %f2, %f1;
	ld.f32 	%f4, [%rd8+8];
	add.ftz.f32 	%f1, %f4, %f3;
	add.s32 	%r12, %r12, 2;
	add.u64 	%rd8, %rd8, 8;
	setp.lt.s32 	%p2, %r12, %r13;
	@%p2 bra 	$Lt_3_1794;
$Lt_3_1282:
	ld.f32 	%f5, [%rd7+-4];
	cvt.s64.s32 	%rd9, %r2;
	mul.wide.s32 	%rd10, %r2, 4;
	add.u64 	%rd11, %rd4, %rd10;
	ld.f32 	%f6, [%rd11+4];
	add.ftz.f32 	%f7, %f5, %f6;
	add.ftz.f32 	%f8, %f7, %f1;
	st.v2.f32 	[%rd2+0], {%f8,%f1};
	ret;
$LDWend__Z6SumRowiP6float2Pf:
	} // _Z6SumRowiP6float2Pf

	.visible .func (.param .f32 __cudaretf__Z9SumColumniP6float2f) _Z9SumColumniP6float2f (.param .s32 __cudaparmf1__Z9SumColumniP6float2f, .param .u64 __cudaparmf2__Z9SumColumniP6float2f, .param .f32 __cudaparmf3__Z9SumColumniP6float2f)
	{
	.reg .u32 %r<19>;
	.reg .u64 %rd<11>;
	.reg .f32 %f<23>;
	.reg .pred %p<4>;
$LDWbegin__Z9SumColumniP6float2f:
	ld.param.u32 	%r1, [__cudaparmf1__Z9SumColumniP6float2f];
	mov.s32 	%r2, %r1;
	ld.param.u64 	%rd1, [__cudaparmf2__Z9SumColumniP6float2f];
	mov.s64 	%rd2, %rd1;
	ld.param.f32 	%f1, [__cudaparmf3__Z9SumColumniP6float2f];
	mov.f32 	%f2, %f1;
	neg.s32 	%r3, %r2;
	cvt.s64.s32 	%rd3, %r3;
	mul.wide.s32 	%rd4, %r3, 8;
	add.u64 	%rd5, %rd2, %rd4;
	ld.f32 	%f3, [%rd5+-8];
	ld.v2.f32 	{%f4,%f5}, [%rd5+0];
	add.ftz.f32 	%f6, %f3, %f4;
	mul.lo.s32 	%r4, %r2, 2;
	add.s32 	%r5, %r4, 1;
	mov.u32 	%r6, 0;
	setp.le.s32 	%p1, %r4, %r6;
	@%p1 bra 	$Lt_4_1282;
	shr.s32 	%r7, %r5, 31;
	mov.s32 	%r8, 1;
	and.b32 	%r9, %r7, %r8;
	add.s32 	%r10, %r9, %r5;
	shr.s32 	%r11, %r10, 1;
	mov.s32 	%r12, %r3;
	sub.s32 	%r13, %r4, %r2;
	mov.s64 	%rd6, %rd5;
	mov.s32 	%r14, %r11;
$Lt_4_1794:
 //<loop> Loop body line 88, nesting depth: 1, estimated iterations: unknown
	ld.v2.f32 	{%f7,%f8}, [%rd6+8];
	add.ftz.f32 	%f9, %f7, %f6;
	add.ftz.f32 	%f10, %f8, %f5;
	ld.v2.f32 	{%f11,%f12}, [%rd6+16];
	add.ftz.f32 	%f6, %f11, %f9;
	add.ftz.f32 	%f5, %f12, %f10;
	add.s32 	%r12, %r12, 2;
	add.u64 	%rd6, %rd6, 16;
	setp.lt.s32 	%p2, %r12, %r13;
	@%p2 bra 	$Lt_4_1794;
$Lt_4_1282:
	cvt.s64.s32 	%rd7, %r2;
	mul.wide.s32 	%rd8, %r2, 8;
	add.u64 	%rd9, %rd2, %rd8;
	ld.f32 	%f13, [%rd9+8];
	add.ftz.f32 	%f6, %f13, %f6;
	add.s32 	%r15, %r4, 3;
	mul.lo.s32 	%r16, %r15, %r15;
	cvt.rn.f32.s32 	%f14, %r16;
	div.approx.ftz.f32 	%f15, %f2, %f14;
	mul.ftz.f32 	%f16, %f6, %f15;
	mov.f32 	%f17, 0f3f800000;    	// 1
	sub.ftz.f32 	%f18, %f17, %f2;
	mul.lo.s32 	%r17, %r5, %r5;
	cvt.rn.f32.s32 	%f19, %r17;
	div.approx.ftz.f32 	%f20, %f18, %f19;
	fma.rn.ftz.f32 	%f5, %f5, %f20, %f16;
	mov.f32 	%f21, %f5;
	st.param.f32 	[__cudaretf__Z9SumColumniP6float2f], %f21;
	ret;
$LDWend__Z9SumColumniP6float2f:
	} // _Z9SumColumniP6float2f

	.entry cuda_kernel_chokefilter_x (
		.param .u64 __cudaparm_cuda_kernel_chokefilter_x_id,
		.param .u64 __cudaparm_cuda_kernel_chokefilter_x_od,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_x_idPitch,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_x_odPitch,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_x_w,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_x_h,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_x_r,
		.param .u32 __cudaparm_cuda_kernel_chokefilter_x_inDeviceFormat)
	{
	.reg .u32 %r<55>;
	.reg .u64 %rd<45>;
	.reg .f32 %f<19>;
	.reg .pred %p<13>;
	.shared .align 4 .b8 __cuda___cuda_local_var_104714_3342_non_const_data208[1152];
$LDWbegin_cuda_kernel_chokefilter_x:
	mov.u32 	%r1, %ctaid.x;
	mov.u32 	%r2, %ntid.x;
	mul.lo.u32 	%r3, %r1, %r2;
	mov.u32 	%r4, %tid.x;
	add.u32 	%r5, %r4, %r3;
	ld.param.s32 	%r6, [__cudaparm_cuda_kernel_chokefilter_x_w];
	setp.le.s32 	%p1, %r6, %r5;
	@%p1 bra 	$Lt_5_14594;
	ld.param.u32 	%r7, [__cudaparm_cuda_kernel_chokefilter_x_inDeviceFormat];
	mov.s32 	%r8, 0;
	setp.eq.s32 	%p2, %r7, %r8;
	ld.param.s32 	%r9, [__cudaparm_cuda_kernel_chokefilter_x_idPitch];
	cvt.s32.u32 	%r10, %ctaid.y;
	mul.lo.s32 	%r11, %r9, %r10;
	add.s32 	%r12, %r5, %r11;
	cvt.s64.s32 	%rd1, %r12;
	ld.param.u64 	%rd2, [__cudaparm_cuda_kernel_chokefilter_x_id];
	@!%p2 bra 	$Lt_5_12034;
	.loc	19	19	0
	mul.lo.u64 	%rd3, %rd1, 8;
	add.u64 	%rd4, %rd2, %rd3;
	ld.global.u16 	%r13, [%rd4+6];
	{ .reg .b32 %b1;
	mov.b32		%b1, %r13;
	cvt.ftz.f32.f16	%f1, %b1; }
	bra.uni 	$LDWendi__Z9ReadAlphaPK6float4i17DevicePixelFormat_186_5;
$Lt_5_12034:
	mul.lo.u64 	%rd5, %rd1, 16;
	add.u64 	%rd6, %rd2, %rd5;
	ld.global.f32 	%f1, [%rd6+12];
$LDWendi__Z9ReadAlphaPK6float4i17DevicePixelFormat_186_5:
	.loc	19	88	0
	mov.u64 	%rd7, __cuda___cuda_local_var_104714_3342_non_const_data208;
	cvt.s32.u32 	%r14, %tid.x;
	cvt.s64.s32 	%rd8, %r14;
	mul.wide.s32 	%rd9, %r14, 4;
	add.u64 	%rd10, %rd7, %rd9;
	st.shared.f32 	[%rd10+64], %f1;
	ld.param.s32 	%r15, [__cudaparm_cuda_kernel_chokefilter_x_r];
	add.s32 	%r16, %r15, 1;
	setp.le.s32 	%p3, %r16, %r14;
	@%p3 bra 	$Lt_5_12802;
	neg.s32 	%r17, %r15;
	cvt.s64.s32 	%rd11, %r17;
	mul.wide.s32 	%rd12, %r17, 4;
	add.u64 	%rd13, %rd10, %rd12;
	setp.le.s32 	%p4, %r16, %r5;
	@%p4 bra 	$Lt_5_13314;
	st.shared.f32 	[%rd13+60], %f1;
	bra.uni 	$Lt_5_14594;
$Lt_5_13314:
	sub.s32 	%r18, %r12, %r15;
	cvt.s64.s32 	%rd14, %r18;
	@!%p2 bra 	$Lt_5_13570;
	.loc	19	19	0
	mul.lo.u64 	%rd15, %rd14, 8;
	add.u64 	%rd16, %rd2, %rd15;
	ld.global.u16 	%r19, [%rd16+-2];
	{ .reg .b32 %b1;
	mov.b32		%b1, %r19;
	cvt.ftz.f32.f16	%f2, %b1; }
	bra.uni 	$LDWendi__Z9ReadAlphaPK6float4i17DevicePixelFormat_186_3;
$Lt_5_13570:
	mul.lo.u64 	%rd17, %rd14, 16;
	add.u64 	%rd18, %rd2, %rd17;
	ld.global.f32 	%f2, [%rd18+-4];
$LDWendi__Z9ReadAlphaPK6float4i17DevicePixelFormat_186_3:
	.loc	19	88	0
	st.shared.f32 	[%rd13+60], %f2;
	bra.uni 	$Lt_5_14594;
$Lt_5_12802:
	cvt.s32.u32 	%r20, %ntid.x;
	sub.s32 	%r21, %r20, %r15;
	sub.s32 	%r22, %r21, 1;
	setp.lt.s32 	%p5, %r14, %r22;
	@%p5 bra 	$Lt_5_14594;
	cvt.s64.s32 	%rd19, %r15;
	mul.wide.s32 	%rd20, %r15, 4;
	add.u64 	%rd21, %rd10, %rd20;
	sub.s32 	%r23, %r6, %r15;
	sub.s32 	%r24, %r23, 1;
	setp.lt.s32 	%p6, %r5, %r24;
	@%p6 bra 	$Lt_5_14850;
	st.shared.f32 	[%rd21+68], %f1;
	bra.uni 	$Lt_5_14594;
$Lt_5_14850:
	add.s32 	%r25, %r12, %r15;
	cvt.s64.s32 	%rd22, %r25;
	@!%p2 bra 	$Lt_5_15106;
	.loc	19	19	0
	mul.lo.u64 	%rd23, %rd22, 8;
	add.u64 	%rd24, %rd2, %rd23;
	ld.global.u16 	%r26, [%rd24+14];
	{ .reg .b32 %b1;
	mov.b32		%b1, %r26;
	cvt.ftz.f32.f16	%f3, %b1; }
	bra.uni 	$LDWendi__Z9ReadAlphaPK6float4i17DevicePixelFormat_186_1;
$Lt_5_15106:
	mul.lo.u64 	%rd25, %rd22, 16;
	add.u64 	%rd26, %rd2, %rd25;
	ld.global.f32 	%f3, [%rd26+28];
$LDWendi__Z9ReadAlphaPK6float4i17DevicePixelFormat_186_1:
	.loc	19	88	0
	st.shared.f32 	[%rd21+68], %f3;
$Lt_5_14594:
$Lt_5_14082:
$Lt_5_12546:
$Lt_5_11522:
	mov.u64 	%rd7, __cuda___cuda_local_var_104714_3342_non_const_data208;
	bar.sync 	0;
	setp.gt.s32 	%p7, %r6, %r5;
	@%p7 bra 	$Lt_5_15618;
	bra.uni 	$LBB34_cuda_kernel_chokefilter_x;
$Lt_5_15618:
	cvt.s32.u32 	%r27, %tid.x;
	cvt.s64.s32 	%rd27, %r27;
	ld.param.s32 	%r15, [__cudaparm_cuda_kernel_chokefilter_x_r];
	neg.s32 	%r28, %r15;
	mul.lo.s32 	%r29, %r15, 2;
	cvt.s64.s32 	%rd28, %r15;
	mul.wide.s32 	%rd29, %r27, 4;
	cvt.s64.s32 	%rd30, %r28;
	mov.s32 	%r30, 0;
	setp.gt.s32 	%p8, %r29, %r30;
	mul.wide.s32 	%rd31, %r15, 4;
	add.u64 	%rd32, %rd29, %rd7;
	mul.wide.s32 	%rd33, %r28, 4;
	add.u64 	%rd34, %rd31, %rd32;
	add.u64 	%rd35, %rd33, %rd32;
	ld.shared.f32 	%f4, [%rd34+68];
	ld.shared.f32 	%f5, [%rd35+64];
	ld.shared.f32 	%f6, [%rd35+60];
	ld.param.u64 	%rd36, [__cudaparm_cuda_kernel_chokefilter_x_od];
	ld.param.s32 	%r31, [__cudaparm_cuda_kernel_chokefilter_x_odPitch];
	cvt.s32.u32 	%r32, %ctaid.y;
	mul.lo.s32 	%r33, %r31, %r32;
	add.s32 	%r34, %r5, %r33;
	cvt.s64.s32 	%rd37, %r34;
	mul.wide.s32 	%rd38, %r34, 8;
	add.u64 	%rd39, %rd36, %rd38;
	min.ftz.f32 	%f7, %f4, %f6;
	mov.u32 	%r35, 3;
	setp.gt.u32 	%p9, %r15, %r35;
	@%p9 bra 	$Lt_5_16386;
	mov.f32 	%f8, %f5;
	@!%p8 bra 	$Lt_5_16642;
	add.s32 	%r36, %r29, 1;
	shr.s32 	%r37, %r36, 31;
	mov.s32 	%r38, 1;
	and.b32 	%r39, %r37, %r38;
	add.s32 	%r40, %r39, %r36;
	shr.s32 	%r41, %r40, 1;
	mov.s32 	%r42, %r28;
	sub.s32 	%r43, %r29, %r15;
	add.s64 	%rd40, %rd33, %rd29;
	add.u64 	%rd41, %rd7, %rd40;
	mov.s32 	%r44, %r41;
$Lt_5_17154:
 //<loop> Loop body line 88, nesting depth: 1, estimated iterations: unknown
	ld.shared.f32 	%f9, [%rd41+68];
	min.ftz.f32 	%f10, %f9, %f8;
	ld.shared.f32 	%f11, [%rd41+72];
	min.ftz.f32 	%f8, %f11, %f10;
	add.s32 	%r42, %r42, 2;
	add.u64 	%rd41, %rd41, 8;
	setp.lt.s32 	%p10, %r42, %r43;
	@%p10 bra 	$Lt_5_17154;
$Lt_5_16642:
	min.ftz.f32 	%f12, %f7, %f8;
	st.global.v2.f32 	[%rd39+0], {%f12,%f8};
	bra.uni 	$LBB34_cuda_kernel_chokefilter_x;
$Lt_5_16386:
	mov.f32 	%f13, %f5;
	@!%p8 bra 	$Lt_5_17666;
	add.s32 	%r45, %r29, 1;
	shr.s32 	%r46, %r45, 31;
	mov.s32 	%r47, 1;
	and.b32 	%r48, %r46, %r47;
	add.s32 	%r49, %r48, %r45;
	shr.s32 	%r50, %r49, 1;
	mov.s32 	%r51, %r28;
	sub.s32 	%r52, %r29, %r15;
	add.s64 	%rd42, %rd33, %rd29;
	add.u64 	%rd43, %rd7, %rd42;
	mov.s32 	%r53, %r50;
$Lt_5_18178:
 //<loop> Loop body line 88, nesting depth: 1, estimated iterations: unknown
	ld.shared.f32 	%f14, [%rd43+68];
	min.ftz.f32 	%f15, %f14, %f13;
	ld.shared.f32 	%f16, [%rd43+72];
	min.ftz.f32 	%f13, %f16, %f15;
	add.s32 	%r51, %r51, 2;
	add.u64 	%rd43, %rd43, 8;
	setp.lt.s32 	%p11, %r51, %r52;
	@%p11 bra 	$Lt_5_18178;
$Lt_5_17666:
	min.ftz.f32 	%f17, %f7, %f13;
	st.global.v2.f32 	[%rd39+0], {%f17,%f13};
$LBB34_cuda_kernel_chokefilter_x:
	exit;
$LDWend_cuda_kernel_chokefilter_x:
	} // cuda_kernel_chokefilter_x
	.global .align 8 .b8 inParams[164];

	.entry cuda_kernel_chokefilter_y (
		.param .u64 __cudaparm_cuda_kernel_chokefilter_y_id,
		.param .u64 __cudaparm_cuda_kernel_chokefilter_y_od,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_y_idPitch,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_y_odPitch,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_y_w,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_y_h,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_y_r,
		.param .f32 __cudaparm_cuda_kernel_chokefilter_y_f,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_y_remap,
		.param .u32 __cudaparm_cuda_kernel_chokefilter_y_inDeviceFormat,
		.param .u64 __cudaparm_cuda_kernel_chokefilter_y_ConstantinParams)
	{
	.reg .u32 %r<78>;
	.reg .u64 %rd<40>;
	.reg .f32 %f<39>;
	.reg .pred %p<14>;
	.shared .align 8 .b8 __cuda___cuda_local_var_104714_4756_non_const_data1416[3200];
$LDWbegin_cuda_kernel_chokefilter_y:
	mov.u32 	%r1, %ctaid.x;
	mov.u32 	%r2, %ntid.x;
	mul.lo.u32 	%r3, %r1, %r2;
	mov.u32 	%r4, %ctaid.y;
	mov.u32 	%r5, %ntid.y;
	mul.lo.u32 	%r6, %r4, %r5;
	mov.u32 	%r7, %tid.x;
	add.u32 	%r8, %r7, %r3;
	mov.u32 	%r9, %tid.y;
	add.u32 	%r10, %r9, %r6;
	ld.param.s32 	%r11, [__cudaparm_cuda_kernel_chokefilter_y_h];
	ld.param.s32 	%r12, [__cudaparm_cuda_kernel_chokefilter_y_w];
	set.gt.u32.s32 	%r13, %r11, %r10;
	neg.s32 	%r14, %r13;
	set.gt.u32.s32 	%r15, %r12, %r8;
	neg.s32 	%r16, %r15;
	and.b32 	%r17, %r14, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_6_11778;
	mov.u64 	%rd1, __cuda___cuda_local_var_104714_4756_non_const_data1416;
	cvt.s32.u32 	%r19, %tid.x;
	mul.lo.s32 	%r20, %r19, 25;
	cvt.s32.u32 	%r21, %tid.y;
	add.s32 	%r22, %r20, %r21;
	ld.param.s32 	%r23, [__cudaparm_cuda_kernel_chokefilter_y_idPitch];
	cvt.s64.s32 	%rd2, %r22;
	mul.wide.s32 	%rd3, %r22, 8;
	add.u64 	%rd4, %rd3, %rd1;
	ld.param.u64 	%rd5, [__cudaparm_cuda_kernel_chokefilter_y_id];
	mul.lo.s32 	%r24, %r23, %r10;
	add.s32 	%r25, %r8, %r24;
	cvt.s64.s32 	%rd6, %r25;
	mul.wide.s32 	%rd7, %r25, 8;
	add.u64 	%rd8, %rd5, %rd7;
	ld.global.v2.f32 	{%f1,%f2}, [%rd8+0];
	st.shared.v2.f32 	[%rd4+64], {%f1,%f2};
	ld.param.s32 	%r26, [__cudaparm_cuda_kernel_chokefilter_y_r];
	add.s32 	%r27, %r26, 1;
	setp.le.s32 	%p2, %r27, %r21;
	@%p2 bra 	$Lt_6_10754;
	neg.s32 	%r28, %r26;
	cvt.s64.s32 	%rd9, %r28;
	mul.wide.s32 	%rd10, %r28, 8;
	add.u64 	%rd11, %rd4, %rd10;
	setp.le.s32 	%p3, %r27, %r10;
	@%p3 bra 	$Lt_6_11010;
	ld.shared.f32 	%f3, [%rd4+68];
	st.shared.v2.f32 	[%rd11+56], {%f1,%f3};
	bra.uni 	$Lt_6_10754;
$Lt_6_11010:
	mul.lo.s32 	%r29, %r27, %r23;
	neg.s32 	%r30, %r29;
	cvt.s64.s32 	%rd12, %r30;
	mul.wide.s32 	%rd13, %r30, 8;
	add.u64 	%rd14, %rd8, %rd13;
	ld.global.v2.f32 	{%f4,%f5}, [%rd14+0];
	st.shared.v2.f32 	[%rd11+56], {%f4,%f5};
$Lt_6_10754:
$Lt_6_10242:
	cvt.s32.u32 	%r31, %ntid.y;
	sub.s32 	%r32, %r31, %r26;
	sub.s32 	%r33, %r32, 1;
	setp.lt.s32 	%p4, %r21, %r33;
	@%p4 bra 	$Lt_6_11778;
	cvt.s64.s32 	%rd15, %r26;
	mul.wide.s32 	%rd16, %r26, 8;
	add.u64 	%rd17, %rd4, %rd16;
	sub.s32 	%r34, %r11, %r26;
	sub.s32 	%r35, %r34, 1;
	setp.lt.s32 	%p5, %r10, %r35;
	@%p5 bra 	$Lt_6_12034;
	ld.shared.v2.f32 	{%f6,%f7}, [%rd4+64];
	st.shared.v2.f32 	[%rd17+72], {%f6,%f7};
	bra.uni 	$Lt_6_11778;
$Lt_6_12034:
	mul.lo.s32 	%r36, %r27, %r23;
	cvt.s64.s32 	%rd18, %r36;
	mul.wide.s32 	%rd19, %r36, 8;
	add.u64 	%rd20, %rd8, %rd19;
	ld.global.v2.f32 	{%f8,%f9}, [%rd20+0];
	st.shared.v2.f32 	[%rd17+72], {%f8,%f9};
$Lt_6_11778:
$Lt_6_11266:
$Lt_6_9730:
	mov.u64 	%rd1, __cuda___cuda_local_var_104714_4756_non_const_data1416;
	bar.sync 	0;
	set.le.u32.s32 	%r37, %r11, %r10;
	neg.s32 	%r38, %r37;
	set.le.u32.s32 	%r39, %r12, %r8;
	neg.s32 	%r40, %r39;
	or.b32 	%r41, %r38, %r40;
	mov.u32 	%r42, 0;
	setp.eq.s32 	%p6, %r41, %r42;
	@%p6 bra 	$Lt_6_12290;
	bra.uni 	$LBB29_cuda_kernel_chokefilter_y;
$Lt_6_12290:
	cvt.s32.u32 	%r43, %tid.x;
	mul.lo.s32 	%r44, %r43, 25;
	ld.param.s32 	%r26, [__cudaparm_cuda_kernel_chokefilter_y_r];
	neg.s32 	%r45, %r26;
	mul.lo.s32 	%r46, %r26, 2;
	cvt.s64.s32 	%rd21, %r26;
	cvt.s32.u32 	%r47, %tid.y;
	add.s32 	%r48, %r47, %r44;
	cvt.s64.s32 	%rd22, %r45;
	mov.s32 	%r49, 0;
	setp.gt.s32 	%p7, %r46, %r49;
	mul.wide.s32 	%rd23, %r26, 8;
	cvt.s64.s32 	%rd24, %r48;
	mul.wide.s32 	%rd25, %r45, 8;
	mul.wide.s32 	%rd3, %r48, 8;
	add.u64 	%rd26, %rd3, %rd1;
	add.u64 	%rd27, %rd25, %rd26;
	add.u64 	%rd28, %rd23, %rd26;
	ld.shared.f32 	%f10, [%rd27+56];
	ld.shared.v2.f32 	{%f11,%f12}, [%rd27+64];
	ld.shared.f32 	%f13, [%rd28+72];
	min.ftz.f32 	%f14, %f10, %f11;
	mov.u32 	%r50, 3;
	setp.gt.u32 	%p8, %r26, %r50;
	@%p8 bra 	$Lt_6_13058;
	mov.f32 	%f15, %f14;
	mov.f32 	%f16, %f12;
	@!%p7 bra 	$Lt_6_13314;
	add.s32 	%r51, %r46, 1;
	shr.s32 	%r52, %r51, 31;
	mov.s32 	%r53, 1;
	and.b32 	%r54, %r52, %r53;
	add.s32 	%r55, %r54, %r51;
	shr.s32 	%r56, %r55, 1;
	mov.s32 	%r57, %r45;
	sub.s32 	%r58, %r46, %r26;
	add.s64 	%rd29, %rd25, %rd3;
	add.u64 	%rd30, %rd1, %rd29;
	mov.s32 	%r59, %r56;
$Lt_6_13826:
 //<loop> Loop body line 88, nesting depth: 1, estimated iterations: unknown
	ld.shared.v2.f32 	{%f17,%f18}, [%rd30+72];
	min.ftz.f32 	%f19, %f17, %f15;
	min.ftz.f32 	%f20, %f18, %f16;
	ld.shared.v2.f32 	{%f21,%f22}, [%rd30+80];
	min.ftz.f32 	%f15, %f21, %f19;
	min.ftz.f32 	%f16, %f22, %f20;
	add.s32 	%r57, %r57, 2;
	add.u64 	%rd30, %rd30, 16;
	setp.lt.s32 	%p9, %r57, %r58;
	@%p9 bra 	$Lt_6_13826;
$Lt_6_13314:
	min.ftz.f32 	%f15, %f13, %f15;
	mov.f32 	%f23, %f15;
	mov.f32 	%f24, %f16;
	bra.uni 	$Lt_6_12802;
$Lt_6_13058:
	mov.f32 	%f23, %f14;
	mov.f32 	%f24, %f12;
	@!%p7 bra 	$Lt_6_14338;
	add.s32 	%r60, %r46, 1;
	shr.s32 	%r61, %r60, 31;
	mov.s32 	%r62, 1;
	and.b32 	%r63, %r61, %r62;
	add.s32 	%r64, %r63, %r60;
	shr.s32 	%r65, %r64, 1;
	mov.s32 	%r66, %r45;
	sub.s32 	%r67, %r46, %r26;
	add.s64 	%rd31, %rd25, %rd3;
	add.u64 	%rd32, %rd1, %rd31;
	mov.s32 	%r68, %r65;
$Lt_6_14850:
 //<loop> Loop body line 88, nesting depth: 1, estimated iterations: unknown
	ld.shared.v2.f32 	{%f25,%f26}, [%rd32+72];
	min.ftz.f32 	%f27, %f25, %f23;
	min.ftz.f32 	%f28, %f26, %f24;
	ld.shared.v2.f32 	{%f29,%f30}, [%rd32+80];
	min.ftz.f32 	%f23, %f29, %f27;
	min.ftz.f32 	%f24, %f30, %f28;
	add.s32 	%r66, %r66, 2;
	add.u64 	%rd32, %rd32, 16;
	setp.lt.s32 	%p10, %r66, %r67;
	@%p10 bra 	$Lt_6_14850;
$Lt_6_14338:
	min.ftz.f32 	%f23, %f13, %f23;
$Lt_6_12802:
	ld.param.f32 	%f31, [__cudaparm_cuda_kernel_chokefilter_y_f];
	mov.f32 	%f32, 0f3f800000;    	// 1
	sub.ftz.f32 	%f33, %f32, %f31;
	mul.ftz.f32 	%f34, %f24, %f33;
	fma.rn.ftz.f32 	%f23, %f31, %f23, %f34;
	ld.param.s32 	%r69, [__cudaparm_cuda_kernel_chokefilter_y_remap];
	mov.u32 	%r70, 0;
	setp.eq.s32 	%p11, %r69, %r70;
	@%p11 bra 	$Lt_6_15362;
	ld.global.f32 	%f35, [inParams+144];
	ld.global.f32 	%f36, [inParams+140];
	fma.rn.ftz.f32 	%f37, %f23, %f36, %f35;
	cvt.ftz.sat.f32.f32 	%f23, %f37;
$Lt_6_15362:
	ld.param.s32 	%r71, [__cudaparm_cuda_kernel_chokefilter_y_odPitch];
	mul.lo.s32 	%r72, %r71, %r10;
	add.s32 	%r73, %r8, %r72;
	cvt.s64.s32 	%rd33, %r73;
	ld.param.u64 	%rd34, [__cudaparm_cuda_kernel_chokefilter_y_od];
	ld.param.u32 	%r74, [__cudaparm_cuda_kernel_chokefilter_y_inDeviceFormat];
	mov.u32 	%r75, 0;
	setp.ne.s32 	%p12, %r74, %r75;
	@%p12 bra 	$Lt_6_16130;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f23;
	mov.b32		%r76, %b1; }
	mul.lo.u64 	%rd35, %rd33, 8;
	add.u64 	%rd36, %rd34, %rd35;
	st.global.u16 	[%rd36+6], %r76;
	bra.uni 	$LBB29_cuda_kernel_chokefilter_y;
$Lt_6_16130:
	mul.lo.u64 	%rd37, %rd33, 16;
	add.u64 	%rd38, %rd34, %rd37;
	st.global.f32 	[%rd38+12], %f23;
$LBB29_cuda_kernel_chokefilter_y:
	exit;
$LDWend_cuda_kernel_chokefilter_y:
	} // cuda_kernel_chokefilter_y

	.entry cuda_kernel_boxfilter_x (
		.param .u64 __cudaparm_cuda_kernel_boxfilter_x_id,
		.param .u64 __cudaparm_cuda_kernel_boxfilter_x_od,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_x_idPitch,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_x_odPitch,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_x_w,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_x_h,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_x_r,
		.param .u32 __cudaparm_cuda_kernel_boxfilter_x_inDeviceFormat)
	{
	.reg .u32 %r<47>;
	.reg .u64 %rd<41>;
	.reg .f32 %f<17>;
	.reg .pred %p<12>;
	.shared .align 4 .b8 __cuda___cuda_local_var_104714_6169_non_const_data4656[1152];
$LDWbegin_cuda_kernel_boxfilter_x:
	mov.u32 	%r1, %ctaid.x;
	mov.u32 	%r2, %ntid.x;
	mul.lo.u32 	%r3, %r1, %r2;
	mov.u32 	%r4, %tid.x;
	add.u32 	%r5, %r4, %r3;
	ld.param.s32 	%r6, [__cudaparm_cuda_kernel_boxfilter_x_w];
	setp.le.s32 	%p1, %r6, %r5;
	@%p1 bra 	$Lt_7_14594;
	ld.param.u32 	%r7, [__cudaparm_cuda_kernel_boxfilter_x_inDeviceFormat];
	mov.s32 	%r8, 0;
	setp.eq.s32 	%p2, %r7, %r8;
	ld.param.s32 	%r9, [__cudaparm_cuda_kernel_boxfilter_x_idPitch];
	cvt.s32.u32 	%r10, %ctaid.y;
	mul.lo.s32 	%r11, %r9, %r10;
	add.s32 	%r12, %r5, %r11;
	cvt.s64.s32 	%rd1, %r12;
	ld.param.u64 	%rd2, [__cudaparm_cuda_kernel_boxfilter_x_id];
	@!%p2 bra 	$Lt_7_12034;
	.loc	19	19	0
	mul.lo.u64 	%rd3, %rd1, 8;
	add.u64 	%rd4, %rd2, %rd3;
	ld.global.u16 	%r13, [%rd4+6];
	{ .reg .b32 %b1;
	mov.b32		%b1, %r13;
	cvt.ftz.f32.f16	%f1, %b1; }
	bra.uni 	$LDWendi__Z9ReadAlphaPK6float4i17DevicePixelFormat_188_5;
$Lt_7_12034:
	mul.lo.u64 	%rd5, %rd1, 16;
	add.u64 	%rd6, %rd2, %rd5;
	ld.global.f32 	%f1, [%rd6+12];
$LDWendi__Z9ReadAlphaPK6float4i17DevicePixelFormat_188_5:
	.loc	19	88	0
	mov.u64 	%rd7, __cuda___cuda_local_var_104714_6169_non_const_data4656;
	cvt.s32.u32 	%r14, %tid.x;
	cvt.s64.s32 	%rd8, %r14;
	mul.wide.s32 	%rd9, %r14, 4;
	add.u64 	%rd10, %rd7, %rd9;
	st.shared.f32 	[%rd10+64], %f1;
	ld.param.s32 	%r15, [__cudaparm_cuda_kernel_boxfilter_x_r];
	add.s32 	%r16, %r15, 1;
	setp.le.s32 	%p3, %r16, %r14;
	@%p3 bra 	$Lt_7_12802;
	neg.s32 	%r17, %r15;
	cvt.s64.s32 	%rd11, %r17;
	mul.wide.s32 	%rd12, %r17, 4;
	add.u64 	%rd13, %rd10, %rd12;
	setp.le.s32 	%p4, %r16, %r5;
	@%p4 bra 	$Lt_7_13314;
	st.shared.f32 	[%rd13+60], %f1;
	bra.uni 	$Lt_7_14594;
$Lt_7_13314:
	sub.s32 	%r18, %r12, %r15;
	cvt.s64.s32 	%rd14, %r18;
	@!%p2 bra 	$Lt_7_13570;
	.loc	19	19	0
	mul.lo.u64 	%rd15, %rd14, 8;
	add.u64 	%rd16, %rd2, %rd15;
	ld.global.u16 	%r19, [%rd16+-2];
	{ .reg .b32 %b1;
	mov.b32		%b1, %r19;
	cvt.ftz.f32.f16	%f2, %b1; }
	bra.uni 	$LDWendi__Z9ReadAlphaPK6float4i17DevicePixelFormat_188_3;
$Lt_7_13570:
	mul.lo.u64 	%rd17, %rd14, 16;
	add.u64 	%rd18, %rd2, %rd17;
	ld.global.f32 	%f2, [%rd18+-4];
$LDWendi__Z9ReadAlphaPK6float4i17DevicePixelFormat_188_3:
	.loc	19	88	0
	st.shared.f32 	[%rd13+60], %f2;
	bra.uni 	$Lt_7_14594;
$Lt_7_12802:
	cvt.s32.u32 	%r20, %ntid.x;
	sub.s32 	%r21, %r20, %r15;
	sub.s32 	%r22, %r21, 1;
	setp.lt.s32 	%p5, %r14, %r22;
	@%p5 bra 	$Lt_7_14594;
	cvt.s64.s32 	%rd19, %r15;
	mul.wide.s32 	%rd20, %r15, 4;
	add.u64 	%rd21, %rd10, %rd20;
	sub.s32 	%r23, %r6, %r15;
	sub.s32 	%r24, %r23, 1;
	setp.lt.s32 	%p6, %r5, %r24;
	@%p6 bra 	$Lt_7_14850;
	st.shared.f32 	[%rd21+68], %f1;
	bra.uni 	$Lt_7_14594;
$Lt_7_14850:
	add.s32 	%r25, %r12, %r15;
	cvt.s64.s32 	%rd22, %r25;
	@!%p2 bra 	$Lt_7_15106;
	.loc	19	19	0
	mul.lo.u64 	%rd23, %rd22, 8;
	add.u64 	%rd24, %rd2, %rd23;
	ld.global.u16 	%r26, [%rd24+14];
	{ .reg .b32 %b1;
	mov.b32		%b1, %r26;
	cvt.ftz.f32.f16	%f3, %b1; }
	bra.uni 	$LDWendi__Z9ReadAlphaPK6float4i17DevicePixelFormat_188_1;
$Lt_7_15106:
	mul.lo.u64 	%rd25, %rd22, 16;
	add.u64 	%rd26, %rd2, %rd25;
	ld.global.f32 	%f3, [%rd26+28];
$LDWendi__Z9ReadAlphaPK6float4i17DevicePixelFormat_188_1:
	.loc	19	88	0
	st.shared.f32 	[%rd21+68], %f3;
$Lt_7_14594:
$Lt_7_14082:
$Lt_7_12546:
$Lt_7_11522:
	mov.u64 	%rd7, __cuda___cuda_local_var_104714_6169_non_const_data4656;
	bar.sync 	0;
	setp.gt.s32 	%p7, %r6, %r5;
	@%p7 bra 	$Lt_7_15618;
	bra.uni 	$LBB30_cuda_kernel_boxfilter_x;
$Lt_7_15618:
	cvt.s32.u32 	%r27, %tid.x;
	cvt.s64.s32 	%rd27, %r27;
	mul.wide.s32 	%rd28, %r27, 4;
	add.u64 	%rd10, %rd28, %rd7;
	ld.param.u64 	%rd29, [__cudaparm_cuda_kernel_boxfilter_x_od];
	ld.param.s32 	%r28, [__cudaparm_cuda_kernel_boxfilter_x_odPitch];
	cvt.s32.u32 	%r29, %ctaid.y;
	mul.lo.s32 	%r30, %r28, %r29;
	add.s32 	%r31, %r5, %r30;
	cvt.s64.s32 	%rd30, %r31;
	mul.wide.s32 	%rd31, %r31, 8;
	add.u64 	%rd32, %rd29, %rd31;
	ld.param.s32 	%r15, [__cudaparm_cuda_kernel_boxfilter_x_r];
	sub.u32 	%r32, %r15, 1;
	mov.u32 	%r33, 2;
	setp.gt.u32 	%p8, %r32, %r33;
	@%p8 bra 	$Lt_7_16386;
	neg.s32 	%r34, %r15;
	cvt.s64.s32 	%rd33, %r34;
	mul.wide.s32 	%rd34, %r34, 4;
	add.u64 	%rd13, %rd34, %rd10;
	ld.shared.f32 	%f4, [%rd13+64];
	mul.lo.s32 	%r35, %r15, 2;
	mov.u32 	%r36, 0;
	setp.le.s32 	%p9, %r35, %r36;
	@%p9 bra 	$Lt_7_16642;
	add.s32 	%r37, %r35, 1;
	shr.s32 	%r38, %r37, 31;
	mov.s32 	%r39, 1;
	and.b32 	%r40, %r38, %r39;
	add.s32 	%r41, %r40, %r37;
	shr.s32 	%r42, %r41, 1;
	mov.s32 	%r43, %r34;
	sub.s32 	%r44, %r35, %r15;
	add.s64 	%rd35, %rd34, %rd28;
	add.u64 	%rd36, %rd7, %rd35;
	mov.s32 	%r45, %r42;
$Lt_7_17154:
 //<loop> Loop body line 88, nesting depth: 1, estimated iterations: unknown
	ld.shared.f32 	%f5, [%rd36+68];
	add.ftz.f32 	%f6, %f5, %f4;
	ld.shared.f32 	%f7, [%rd36+72];
	add.ftz.f32 	%f4, %f7, %f6;
	add.s32 	%r43, %r43, 2;
	add.u64 	%rd36, %rd36, 8;
	setp.lt.s32 	%p10, %r43, %r44;
	@%p10 bra 	$Lt_7_17154;
$Lt_7_16642:
	ld.shared.f32 	%f8, [%rd13+60];
	cvt.s64.s32 	%rd37, %r15;
	mul.wide.s32 	%rd38, %r15, 4;
	add.u64 	%rd39, %rd10, %rd38;
	ld.shared.f32 	%f9, [%rd39+68];
	add.ftz.f32 	%f10, %f8, %f9;
	add.ftz.f32 	%f11, %f4, %f10;
	st.global.v2.f32 	[%rd32+0], {%f11,%f4};
	bra.uni 	$LBB30_cuda_kernel_boxfilter_x;
$Lt_7_16386:
	ld.shared.f32 	%f4, [%rd10+64];
	ld.shared.f32 	%f12, [%rd10+60];
	ld.shared.f32 	%f13, [%rd10+68];
	add.ftz.f32 	%f14, %f12, %f13;
	add.ftz.f32 	%f15, %f4, %f14;
	st.global.v2.f32 	[%rd32+0], {%f15,%f4};
$LBB30_cuda_kernel_boxfilter_x:
	exit;
$LDWend_cuda_kernel_boxfilter_x:
	} // cuda_kernel_boxfilter_x

	.entry cuda_kernel_boxfilter_y (
		.param .u64 __cudaparm_cuda_kernel_boxfilter_y_id,
		.param .u64 __cudaparm_cuda_kernel_boxfilter_y_od,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_y_idPitch,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_y_odPitch,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_y_w,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_y_h,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_y_r,
		.param .f32 __cudaparm_cuda_kernel_boxfilter_y_f,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_y_remap,
		.param .u32 __cudaparm_cuda_kernel_boxfilter_y_inDeviceFormat,
		.param .u64 __cudaparm_cuda_kernel_boxfilter_y_ConstantinParams)
	{
	.reg .u32 %r<73>;
	.reg .u64 %rd<37>;
	.reg .f32 %f<44>;
	.reg .pred %p<13>;
	.shared .align 8 .b8 __cuda___cuda_local_var_104714_7466_non_const_data5864[3200];
$LDWbegin_cuda_kernel_boxfilter_y:
	mov.u32 	%r1, %ctaid.x;
	mov.u32 	%r2, %ntid.x;
	mul.lo.u32 	%r3, %r1, %r2;
	mov.u32 	%r4, %ctaid.y;
	mov.u32 	%r5, %ntid.y;
	mul.lo.u32 	%r6, %r4, %r5;
	mov.u32 	%r7, %tid.x;
	add.u32 	%r8, %r7, %r3;
	mov.u32 	%r9, %tid.y;
	add.u32 	%r10, %r9, %r6;
	ld.param.s32 	%r11, [__cudaparm_cuda_kernel_boxfilter_y_h];
	ld.param.s32 	%r12, [__cudaparm_cuda_kernel_boxfilter_y_w];
	set.gt.u32.s32 	%r13, %r11, %r10;
	neg.s32 	%r14, %r13;
	set.gt.u32.s32 	%r15, %r12, %r8;
	neg.s32 	%r16, %r15;
	and.b32 	%r17, %r14, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_8_11778;
	mov.u64 	%rd1, __cuda___cuda_local_var_104714_7466_non_const_data5864;
	cvt.s32.u32 	%r19, %tid.x;
	mul.lo.s32 	%r20, %r19, 25;
	cvt.s32.u32 	%r21, %tid.y;
	add.s32 	%r22, %r20, %r21;
	ld.param.s32 	%r23, [__cudaparm_cuda_kernel_boxfilter_y_idPitch];
	cvt.s64.s32 	%rd2, %r22;
	mul.wide.s32 	%rd3, %r22, 8;
	add.u64 	%rd4, %rd3, %rd1;
	ld.param.u64 	%rd5, [__cudaparm_cuda_kernel_boxfilter_y_id];
	mul.lo.s32 	%r24, %r23, %r10;
	add.s32 	%r25, %r8, %r24;
	cvt.s64.s32 	%rd6, %r25;
	mul.wide.s32 	%rd7, %r25, 8;
	add.u64 	%rd8, %rd5, %rd7;
	ld.global.v2.f32 	{%f1,%f2}, [%rd8+0];
	st.shared.v2.f32 	[%rd4+64], {%f1,%f2};
	ld.param.s32 	%r26, [__cudaparm_cuda_kernel_boxfilter_y_r];
	add.s32 	%r27, %r26, 1;
	setp.le.s32 	%p2, %r27, %r21;
	@%p2 bra 	$Lt_8_10498;
	neg.s32 	%r28, %r26;
	cvt.s64.s32 	%rd9, %r28;
	mul.wide.s32 	%rd10, %r28, 8;
	add.u64 	%rd11, %rd4, %rd10;
	setp.le.s32 	%p3, %r27, %r10;
	@%p3 bra 	$Lt_8_11010;
	ld.shared.f32 	%f3, [%rd4+68];
	st.shared.v2.f32 	[%rd11+56], {%f1,%f3};
	bra.uni 	$Lt_8_11778;
$Lt_8_11010:
	mul.lo.s32 	%r29, %r27, %r23;
	neg.s32 	%r30, %r29;
	cvt.s64.s32 	%rd12, %r30;
	mul.wide.s32 	%rd13, %r30, 8;
	add.u64 	%rd14, %rd8, %rd13;
	ld.global.v2.f32 	{%f4,%f5}, [%rd14+0];
	st.shared.v2.f32 	[%rd11+56], {%f4,%f5};
	bra.uni 	$Lt_8_11778;
$Lt_8_10498:
	cvt.s32.u32 	%r31, %ntid.y;
	sub.s32 	%r32, %r31, %r26;
	sub.s32 	%r33, %r32, 1;
	setp.lt.s32 	%p4, %r21, %r33;
	@%p4 bra 	$Lt_8_11778;
	cvt.s64.s32 	%rd15, %r26;
	mul.wide.s32 	%rd16, %r26, 8;
	add.u64 	%rd17, %rd4, %rd16;
	sub.s32 	%r34, %r11, %r26;
	sub.s32 	%r35, %r34, 1;
	setp.lt.s32 	%p5, %r10, %r35;
	@%p5 bra 	$Lt_8_12034;
	ld.shared.f32 	%f6, [%rd4+68];
	st.shared.v2.f32 	[%rd17+72], {%f1,%f6};
	bra.uni 	$Lt_8_11778;
$Lt_8_12034:
	mul.lo.s32 	%r36, %r27, %r23;
	cvt.s64.s32 	%rd18, %r36;
	mul.wide.s32 	%rd19, %r36, 8;
	add.u64 	%rd20, %rd8, %rd19;
	ld.global.v2.f32 	{%f7,%f8}, [%rd20+0];
	st.shared.v2.f32 	[%rd17+72], {%f7,%f8};
$Lt_8_11778:
$Lt_8_11266:
$Lt_8_10242:
$Lt_8_9730:
	mov.u64 	%rd1, __cuda___cuda_local_var_104714_7466_non_const_data5864;
	bar.sync 	0;
	set.le.u32.s32 	%r37, %r11, %r10;
	neg.s32 	%r38, %r37;
	set.le.u32.s32 	%r39, %r12, %r8;
	neg.s32 	%r40, %r39;
	or.b32 	%r41, %r38, %r40;
	mov.u32 	%r42, 0;
	setp.eq.s32 	%p6, %r41, %r42;
	@%p6 bra 	$Lt_8_12290;
	bra.uni 	$LBB26_cuda_kernel_boxfilter_y;
$Lt_8_12290:
	cvt.s32.u32 	%r43, %tid.x;
	mul.lo.s32 	%r44, %r43, 25;
	ld.param.f32 	%f9, [__cudaparm_cuda_kernel_boxfilter_y_f];
	mov.f32 	%f10, 0f3f800000;    	// 1
	sub.ftz.f32 	%f11, %f10, %f9;
	cvt.s32.u32 	%r45, %tid.y;
	add.s32 	%r46, %r45, %r44;
	cvt.s64.s32 	%rd21, %r46;
	mul.wide.s32 	%rd3, %r46, 8;
	add.u64 	%rd4, %rd3, %rd1;
	ld.param.s32 	%r26, [__cudaparm_cuda_kernel_boxfilter_y_r];
	sub.u32 	%r47, %r26, 1;
	mov.u32 	%r48, 2;
	setp.gt.u32 	%p7, %r47, %r48;
	@%p7 bra 	$Lt_8_13058;
	neg.s32 	%r49, %r26;
	cvt.s64.s32 	%rd22, %r49;
	mul.wide.s32 	%rd23, %r49, 8;
	add.u64 	%rd24, %rd23, %rd4;
	ld.shared.f32 	%f12, [%rd24+56];
	ld.shared.v2.f32 	{%f13,%f14}, [%rd24+64];
	add.ftz.f32 	%f15, %f12, %f13;
	mul.lo.s32 	%r50, %r26, 2;
	add.s32 	%r51, %r50, 1;
	mov.u32 	%r52, 0;
	setp.le.s32 	%p8, %r50, %r52;
	@%p8 bra 	$Lt_8_13314;
	shr.s32 	%r53, %r51, 31;
	mov.s32 	%r54, 1;
	and.b32 	%r55, %r53, %r54;
	add.s32 	%r56, %r55, %r51;
	shr.s32 	%r57, %r56, 1;
	mov.s32 	%r58, %r49;
	sub.s32 	%r59, %r50, %r26;
	add.s64 	%rd25, %rd23, %rd3;
	add.u64 	%rd26, %rd1, %rd25;
	mov.s32 	%r60, %r57;
$Lt_8_13826:
 //<loop> Loop body line 88, nesting depth: 1, estimated iterations: unknown
	ld.shared.v2.f32 	{%f16,%f17}, [%rd26+72];
	add.ftz.f32 	%f18, %f16, %f15;
	add.ftz.f32 	%f19, %f17, %f14;
	ld.shared.v2.f32 	{%f20,%f21}, [%rd26+80];
	add.ftz.f32 	%f15, %f20, %f18;
	add.ftz.f32 	%f14, %f21, %f19;
	add.s32 	%r58, %r58, 2;
	add.u64 	%rd26, %rd26, 16;
	setp.lt.s32 	%p9, %r58, %r59;
	@%p9 bra 	$Lt_8_13826;
$Lt_8_13314:
	cvt.s64.s32 	%rd27, %r26;
	mul.wide.s32 	%rd28, %r26, 8;
	add.u64 	%rd29, %rd4, %rd28;
	ld.shared.f32 	%f22, [%rd29+72];
	add.ftz.f32 	%f15, %f22, %f15;
	add.s32 	%r61, %r50, 3;
	mul.lo.s32 	%r62, %r61, %r61;
	cvt.rn.f32.s32 	%f23, %r62;
	div.approx.ftz.f32 	%f24, %f9, %f23;
	mul.ftz.f32 	%f25, %f15, %f24;
	mul.lo.s32 	%r63, %r51, %r51;
	cvt.rn.f32.s32 	%f26, %r63;
	div.approx.ftz.f32 	%f27, %f11, %f26;
	fma.rn.ftz.f32 	%f14, %f14, %f27, %f25;
	mov.f32 	%f28, %f14;
	bra.uni 	$Lt_8_12802;
$Lt_8_13058:
	ld.shared.f32 	%f29, [%rd4+72];
	ld.shared.f32 	%f30, [%rd4+56];
	ld.shared.v2.f32 	{%f31,%f32}, [%rd4+64];
	add.ftz.f32 	%f33, %f30, %f31;
	add.ftz.f32 	%f34, %f29, %f33;
	mov.f32 	%f35, 0f41100000;    	// 9
	div.approx.ftz.f32 	%f36, %f9, %f35;
	mul.ftz.f32 	%f37, %f34, %f36;
	mov.f32 	%f38, 0f3f800000;    	// 1
	div.approx.ftz.f32 	%f39, %f11, %f38;
	fma.rn.ftz.f32 	%f28, %f32, %f39, %f37;
$Lt_8_12802:
	ld.param.s32 	%r64, [__cudaparm_cuda_kernel_boxfilter_y_remap];
	mov.u32 	%r65, 0;
	setp.eq.s32 	%p10, %r64, %r65;
	@%p10 bra 	$Lt_8_15362;
	ld.global.f32 	%f40, [inParams+144];
	ld.global.f32 	%f41, [inParams+140];
	fma.rn.ftz.f32 	%f42, %f28, %f41, %f40;
	cvt.ftz.sat.f32.f32 	%f28, %f42;
$Lt_8_15362:
	ld.param.s32 	%r66, [__cudaparm_cuda_kernel_boxfilter_y_odPitch];
	mul.lo.s32 	%r67, %r66, %r10;
	add.s32 	%r68, %r8, %r67;
	cvt.s64.s32 	%rd30, %r68;
	ld.param.u64 	%rd31, [__cudaparm_cuda_kernel_boxfilter_y_od];
	ld.param.u32 	%r69, [__cudaparm_cuda_kernel_boxfilter_y_inDeviceFormat];
	mov.u32 	%r70, 0;
	setp.ne.s32 	%p11, %r69, %r70;
	@%p11 bra 	$Lt_8_16130;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f28;
	mov.b32		%r71, %b1; }
	mul.lo.u64 	%rd32, %rd30, 8;
	add.u64 	%rd33, %rd31, %rd32;
	st.global.u16 	[%rd33+6], %r71;
	bra.uni 	$LBB26_cuda_kernel_boxfilter_y;
$Lt_8_16130:
	mul.lo.u64 	%rd34, %rd30, 16;
	add.u64 	%rd35, %rd31, %rd34;
	st.global.f32 	[%rd35+12], %f28;
$LBB26_cuda_kernel_boxfilter_y:
	exit;
$LDWend_cuda_kernel_boxfilter_y:
	} // cuda_kernel_boxfilter_y

	.entry cuda_kernel_composite (
		.param .u64 __cudaparm_cuda_kernel_composite_srcFrame,
		.param .u64 __cudaparm_cuda_kernel_composite_dstFrame,
		.param .s32 __cudaparm_cuda_kernel_composite_width,
		.param .s32 __cudaparm_cuda_kernel_composite_height,
		.param .s32 __cudaparm_cuda_kernel_composite_srcPitch,
		.param .s32 __cudaparm_cuda_kernel_composite_dstPitch,
		.param .s32 __cudaparm_cuda_kernel_composite_remap,
		.param .u32 __cudaparm_cuda_kernel_composite_inDeviceFormat,
		.param .u64 __cudaparm_cuda_kernel_composite_ConstantinParams)
	{
	.reg .u32 %r<38>;
	.reg .u64 %rd<14>;
	.reg .f32 %f<244>;
	.reg .pred %p<19>;
$LDWbegin_cuda_kernel_composite:
	mov.u32 	%r1, %ctaid.x;
	mov.u32 	%r2, %ntid.x;
	mul.lo.u32 	%r3, %r1, %r2;
	mov.u32 	%r4, %ctaid.y;
	mov.u32 	%r5, %ntid.y;
	mul.lo.u32 	%r6, %r4, %r5;
	mov.u32 	%r7, %tid.x;
	add.u32 	%r8, %r7, %r3;
	mov.u32 	%r9, %tid.y;
	add.u32 	%r10, %r9, %r6;
	ld.param.s32 	%r11, [__cudaparm_cuda_kernel_composite_width];
	set.le.u32.s32 	%r12, %r11, %r8;
	neg.s32 	%r13, %r12;
	ld.param.s32 	%r14, [__cudaparm_cuda_kernel_composite_height];
	set.le.u32.s32 	%r15, %r14, %r10;
	neg.s32 	%r16, %r15;
	or.b32 	%r17, %r13, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_9_14338;
	bra.uni 	$LBB51_cuda_kernel_composite;
$Lt_9_14338:
	ld.param.u32 	%r19, [__cudaparm_cuda_kernel_composite_inDeviceFormat];
	mov.s32 	%r20, 0;
	setp.eq.s32 	%p2, %r19, %r20;
	ld.param.s32 	%r21, [__cudaparm_cuda_kernel_composite_srcPitch];
	mul.lo.s32 	%r22, %r21, %r10;
	add.s32 	%r23, %r8, %r22;
	cvt.s64.s32 	%rd1, %r23;
	ld.param.u64 	%rd2, [__cudaparm_cuda_kernel_composite_srcFrame];
	@!%p2 bra 	$Lt_9_15106;
	mul.lo.u64 	%rd3, %rd1, 8;
	add.u64 	%rd4, %rd2, %rd3;
	ld.global.v4.u16 	{%r24,%r25,%r26,%r27}, [%rd4+0];
	{ .reg .b32 %b1;
	mov.b32		%b1, %r24;
	cvt.ftz.f32.f16	%f1, %b1; }
	{ .reg .b32 %b1;
	mov.b32		%b1, %r25;
	cvt.ftz.f32.f16	%f2, %b1; }
	{ .reg .b32 %b1;
	mov.b32		%b1, %r26;
	cvt.ftz.f32.f16	%f3, %b1; }
	{ .reg .b32 %b1;
	mov.b32		%b1, %r27;
	cvt.ftz.f32.f16	%f4, %b1; }
	bra.uni 	$Lt_9_14850;
$Lt_9_15106:
	mul.lo.u64 	%rd5, %rd1, 16;
	add.u64 	%rd6, %rd2, %rd5;
	ld.global.v4.f32 	{%f1,%f2,%f3,%f4}, [%rd6+0];
$Lt_9_14850:
	cvt.ftz.sat.f32.f32 	%f5, %f1;
	mov.f32 	%f6, 0f3d25aee6;     	// 0.04045
	setp.gt.ftz.f32 	%p3, %f5, %f6;
	@!%p3 bra 	$Lt_9_15618;
	mov.f32 	%f7, 0f3d6147ae;     	// 0.055
	add.ftz.f32 	%f8, %f5, %f7;
	mov.f32 	%f9, 0f3f870a3d;     	// 1.055
	div.approx.ftz.f32 	%f10, %f8, %f9;
	mov.f32 	%f11, 0f00000000;    	// 0
	setp.ge.ftz.f32 	%p4, %f10, %f11;
	@!%p4 bra 	$Lt_9_16130;
	lg2.approx.ftz.f32 	%f12, %f10;
	mov.f32 	%f13, 0f4019999a;    	// 2.4
	mul.ftz.f32 	%f14, %f12, %f13;
	ex2.approx.ftz.f32 	%f15, %f14;
	bra.uni 	$Lt_9_15874;
$Lt_9_16130:
	neg.ftz.f32 	%f16, %f10;
	lg2.approx.ftz.f32 	%f17, %f16;
	mov.f32 	%f18, 0f4019999a;    	// 2.4
	mul.ftz.f32 	%f19, %f17, %f18;
	ex2.approx.ftz.f32 	%f20, %f19;
	neg.ftz.f32 	%f15, %f20;
$Lt_9_15874:
	mov.f32 	%f21, %f15;
	bra.uni 	$Lt_9_15362;
$Lt_9_15618:
	mov.f32 	%f22, 0f414eb852;    	// 12.92
	div.approx.ftz.f32 	%f21, %f5, %f22;
$Lt_9_15362:
	cvt.ftz.sat.f32.f32 	%f23, %f2;
	mov.f32 	%f24, 0f3d25aee6;    	// 0.04045
	setp.gt.ftz.f32 	%p5, %f23, %f24;
	@!%p5 bra 	$Lt_9_16642;
	mov.f32 	%f25, 0f3d6147ae;    	// 0.055
	add.ftz.f32 	%f26, %f23, %f25;
	mov.f32 	%f27, 0f3f870a3d;    	// 1.055
	div.approx.ftz.f32 	%f28, %f26, %f27;
	mov.f32 	%f29, 0f00000000;    	// 0
	setp.ge.ftz.f32 	%p6, %f28, %f29;
	@!%p6 bra 	$Lt_9_17154;
	lg2.approx.ftz.f32 	%f30, %f28;
	mov.f32 	%f31, 0f4019999a;    	// 2.4
	mul.ftz.f32 	%f32, %f30, %f31;
	ex2.approx.ftz.f32 	%f33, %f32;
	bra.uni 	$Lt_9_16898;
$Lt_9_17154:
	neg.ftz.f32 	%f34, %f28;
	lg2.approx.ftz.f32 	%f35, %f34;
	mov.f32 	%f36, 0f4019999a;    	// 2.4
	mul.ftz.f32 	%f37, %f35, %f36;
	ex2.approx.ftz.f32 	%f38, %f37;
	neg.ftz.f32 	%f33, %f38;
$Lt_9_16898:
	mov.f32 	%f39, %f33;
	bra.uni 	$Lt_9_16386;
$Lt_9_16642:
	mov.f32 	%f40, 0f414eb852;    	// 12.92
	div.approx.ftz.f32 	%f39, %f23, %f40;
$Lt_9_16386:
	cvt.ftz.sat.f32.f32 	%f41, %f3;
	mov.f32 	%f42, 0f3d25aee6;    	// 0.04045
	setp.gt.ftz.f32 	%p7, %f41, %f42;
	@!%p7 bra 	$Lt_9_17666;
	mov.f32 	%f43, 0f3d6147ae;    	// 0.055
	add.ftz.f32 	%f44, %f41, %f43;
	mov.f32 	%f45, 0f3f870a3d;    	// 1.055
	div.approx.ftz.f32 	%f46, %f44, %f45;
	mov.f32 	%f47, 0f00000000;    	// 0
	setp.ge.ftz.f32 	%p8, %f46, %f47;
	@!%p8 bra 	$Lt_9_18178;
	lg2.approx.ftz.f32 	%f48, %f46;
	mov.f32 	%f49, 0f4019999a;    	// 2.4
	mul.ftz.f32 	%f50, %f48, %f49;
	ex2.approx.ftz.f32 	%f51, %f50;
	bra.uni 	$Lt_9_17922;
$Lt_9_18178:
	neg.ftz.f32 	%f52, %f46;
	lg2.approx.ftz.f32 	%f53, %f52;
	mov.f32 	%f54, 0f4019999a;    	// 2.4
	mul.ftz.f32 	%f55, %f53, %f54;
	ex2.approx.ftz.f32 	%f56, %f55;
	neg.ftz.f32 	%f51, %f56;
$Lt_9_17922:
	mov.f32 	%f57, %f51;
	bra.uni 	$Lt_9_17410;
$Lt_9_17666:
	mov.f32 	%f58, 0f414eb852;    	// 12.92
	div.approx.ftz.f32 	%f57, %f41, %f58;
$Lt_9_17410:
	ld.global.f32 	%f59, [inParams+52];
	mul.ftz.f32 	%f60, %f59, %f57;
	ld.global.f32 	%f61, [inParams+64];
	mul.ftz.f32 	%f62, %f61, %f57;
	ld.global.f32 	%f63, [inParams+40];
	mul.ftz.f32 	%f64, %f63, %f57;
	ld.global.f32 	%f65, [inParams+56];
	fma.rn.ftz.f32 	%f66, %f39, %f65, %f60;
	ld.global.f32 	%f67, [inParams+68];
	fma.rn.ftz.f32 	%f68, %f39, %f67, %f62;
	ld.global.f32 	%f69, [inParams+44];
	fma.rn.ftz.f32 	%f70, %f39, %f69, %f64;
	ld.global.f32 	%f71, [inParams+60];
	fma.rn.ftz.f32 	%f72, %f21, %f71, %f66;
	ld.global.f32 	%f73, [inParams+72];
	fma.rn.ftz.f32 	%f74, %f21, %f73, %f68;
	ld.global.f32 	%f75, [inParams+48];
	fma.rn.ftz.f32 	%f76, %f21, %f75, %f70;
	mov.f32 	%f77, 0f02081cea;    	// 1e-037
	max.ftz.f32 	%f78, %f74, %f77;
	div.approx.ftz.f32 	%f79, %f72, %f78;
	div.approx.ftz.f32 	%f80, %f76, %f78;
	ld.global.f32 	%f81, [inParams+152];
	sub.ftz.f32 	%f82, %f79, %f81;
	ld.global.f32 	%f83, [inParams+148];
	sub.ftz.f32 	%f84, %f80, %f83;
	mul.ftz.f32 	%f85, %f82, %f82;
	fma.rn.ftz.f32 	%f86, %f84, %f84, %f85;
	mov.f32 	%f87, 0f02081cea;    	// 1e-037
	max.ftz.f32 	%f88, %f86, %f87;
	rsqrt.approx.ftz.f32 	%f89, %f88;
	ld.global.f32 	%f90, [inParams+160];
	mul.ftz.f32 	%f91, %f82, %f81;
	fma.rn.ftz.f32 	%f92, %f83, %f84, %f91;
	mul.ftz.f32 	%f93, %f89, %f90;
	mul.ftz.f32 	%f94, %f92, %f93;
	ld.global.f32 	%f95, [inParams+8];
	mov.f32 	%f96, 0f00000000;    	// 0
	setp.gt.ftz.f32 	%p9, %f94, %f96;
	@!%p9 bra 	$Lt_9_18690;
	mul.ftz.f32 	%f97, %f94, %f95;
	mov.f32 	%f98, %f95;
	bra.uni 	$Lt_9_18434;
$Lt_9_18690:
	mov.f32 	%f99, 0f3f800000;    	// 1
	add.ftz.f32 	%f100, %f94, %f99;
	mul.ftz.f32 	%f98, %f95, %f100;
	mov.f32 	%f97, 0f00000000;    	// 0
$Lt_9_18434:
	.loc	23	544	0
	mov.f32 	%f101, 0f3f800000;   	// 1
	div.approx.ftz.f32 	%f102, %f101, %f89;
	.loc	19	88	0
	mul.ftz.f32 	%f103, %f102, %f90;
	mov.f32 	%f104, 0f3f800000;   	// 1
	sub.ftz.f32 	%f105, %f104, %f97;
	mul.ftz.f32 	%f97, %f103, %f105;
	mov.f32 	%f106, 0f3f800000;   	// 1
	sub.ftz.f32 	%f107, %f106, %f98;
	mul.ftz.f32 	%f98, %f103, %f107;
	ld.global.f32 	%f108, [inParams+16];
	mul.ftz.f32 	%f109, %f108, %f97;
	cvt.ftz.sat.f32.f32 	%f110, %f109;
	ld.global.f32 	%f111, [inParams+156];
	sub.ftz.f32 	%f112, %f78, %f111;
	div.approx.ftz.f32 	%f113, %f112, %f111;
	mov.f32 	%f114, 0f00000000;   	// 0
	setp.lt.ftz.f32 	%p10, %f113, %f114;
	ld.global.f32 	%f115, [inParams+20];
	mul.ftz.f32 	%f116, %f115, %f113;
	ld.global.f32 	%f117, [inParams+24];
	mul.ftz.f32 	%f118, %f117, %f113;
	selp.f32 	%f119, %f116, %f118, %p10;
	cvt.ftz.sat.f32.f32 	%f120, %f119;
	add.ftz.f32 	%f121, %f110, %f120;
	mul.ftz.f32 	%f122, %f120, %f110;
	sub.ftz.f32 	%f123, %f121, %f122;
	ld.global.f32 	%f124, [inParams+28];
	mov.f32 	%f125, 0f3f800000;   	// 1
	mov.f32 	%f126, 0f02081cea;   	// 1e-037
	max.ftz.f32 	%f127, %f123, %f126;
	div.approx.ftz.f32 	%f128, %f125, %f127;
	mov.f32 	%f129, 0fbf800000;   	// -1
	add.ftz.f32 	%f130, %f128, %f129;
	mul.ftz.f32 	%f131, %f124, %f130;
	fma.rn.ftz.f32 	%f132, %f131, %f112, %f78;
	ld.global.f32 	%f133, [inParams+32];
	mul.ftz.f32 	%f134, %f133, %f123;
	ld.global.f32 	%f135, [inParams+36];
	sub.ftz.f32 	%f136, %f134, %f135;
	cvt.ftz.sat.f32.f32 	%f137, %f136;
	ld.global.f32 	%f138, [inParams+12];
	mul.ftz.f32 	%f139, %f138, %f98;
	ld.global.f32 	%f140, [inParams+0];
	mul.ftz.f32 	%f141, %f140, %f98;
	cvt.ftz.sat.f32.f32 	%f142, %f139;
	cvt.ftz.sat.f32.f32 	%f143, %f141;
	mov.f32 	%f144, 0f02081cea;   	// 1e-037
	max.ftz.f32 	%f145, %f143, %f144;
	mov.f32 	%f146, 0f3f800000;   	// 1
	div.approx.ftz.f32 	%f147, %f146, %f145;
	mov.f32 	%f148, 0fbf800000;   	// -1
	add.ftz.f32 	%f149, %f147, %f148;
	ld.global.f32 	%f150, [inParams+4];
	fma.rn.ftz.f32 	%f151, %f149, %f150, %f149;
	fma.rn.ftz.f32 	%f152, %f151, %f82, %f79;
	fma.rn.ftz.f32 	%f153, %f84, %f151, %f80;
	mul.ftz.f32 	%f154, %f142, %f152;
	mul.ftz.f32 	%f155, %f142, %f153;
	mul.ftz.f32 	%f156, %f154, %f132;
	mul.ftz.f32 	%f157, %f155, %f132;
	ld.global.f32 	%f158, [inParams+112];
	ld.global.v2.f32 	{%f159,%f160}, [inParams+104];
	mul.ftz.f32 	%f161, %f159, %f156;
	ld.global.f32 	%f162, [inParams+100];
	fma.rn.ftz.f32 	%f163, %f157, %f162, %f161;
	fma.rn.ftz.f32 	%f164, %f132, %f160, %f163;
	add.ftz.f32 	%f165, %f158, %f164;
	cvt.ftz.sat.f32.f32 	%f166, %f165;
	mov.f32 	%f167, 0f3b4d2e1c;   	// 0.0031308
	setp.lt.ftz.f32 	%p11, %f166, %f167;
	@!%p11 bra 	$Lt_9_19202;
	mov.f32 	%f168, 0f414eb852;   	// 12.92
	mul.ftz.f32 	%f166, %f166, %f168;
	bra.uni 	$Lt_9_18946;
$Lt_9_19202:
	mov.f32 	%f169, 0f00000000;   	// 0
	setp.ge.ftz.f32 	%p12, %f166, %f169;
	@!%p12 bra 	$Lt_9_19714;
	mov.f32 	%f170, 0fbd6147ae;   	// -0.055
	lg2.approx.ftz.f32 	%f171, %f166;
	mov.f32 	%f172, 0f3ed55476;   	// 0.41666
	mul.ftz.f32 	%f173, %f171, %f172;
	ex2.approx.ftz.f32 	%f174, %f173;
	mov.f32 	%f175, 0f3f870a3d;   	// 1.055
	fma.rn.ftz.f32 	%f176, %f174, %f175, %f170;
	bra.uni 	$Lt_9_19458;
$Lt_9_19714:
	mov.f32 	%f177, 0fbd6147ae;   	// -0.055
	neg.ftz.f32 	%f178, %f166;
	lg2.approx.ftz.f32 	%f179, %f178;
	mov.f32 	%f180, 0f3ed55476;   	// 0.41666
	mul.ftz.f32 	%f181, %f179, %f180;
	ex2.approx.ftz.f32 	%f182, %f181;
	neg.ftz.f32 	%f183, %f182;
	mov.f32 	%f184, 0f3f870a3d;   	// 1.055
	fma.rn.ftz.f32 	%f176, %f183, %f184, %f177;
$Lt_9_19458:
	mov.f32 	%f166, %f176;
$Lt_9_18946:
	ld.global.f32 	%f185, [inParams+116];
	ld.global.v2.f32 	{%f186,%f187}, [inParams+88];
	mul.ftz.f32 	%f188, %f187, %f156;
	fma.rn.ftz.f32 	%f189, %f157, %f186, %f188;
	ld.global.f32 	%f190, [inParams+96];
	fma.rn.ftz.f32 	%f191, %f132, %f190, %f189;
	add.ftz.f32 	%f192, %f185, %f191;
	cvt.ftz.sat.f32.f32 	%f193, %f192;
	mov.f32 	%f194, 0f3b4d2e1c;   	// 0.0031308
	setp.lt.ftz.f32 	%p13, %f193, %f194;
	@!%p13 bra 	$Lt_9_20226;
	mov.f32 	%f195, 0f414eb852;   	// 12.92
	mul.ftz.f32 	%f193, %f193, %f195;
	bra.uni 	$Lt_9_19970;
$Lt_9_20226:
	mov.f32 	%f196, 0f00000000;   	// 0
	setp.ge.ftz.f32 	%p14, %f193, %f196;
	@!%p14 bra 	$Lt_9_20738;
	mov.f32 	%f197, 0fbd6147ae;   	// -0.055
	lg2.approx.ftz.f32 	%f198, %f193;
	mov.f32 	%f199, 0f3ed55476;   	// 0.41666
	mul.ftz.f32 	%f200, %f198, %f199;
	ex2.approx.ftz.f32 	%f201, %f200;
	mov.f32 	%f202, 0f3f870a3d;   	// 1.055
	fma.rn.ftz.f32 	%f203, %f201, %f202, %f197;
	bra.uni 	$Lt_9_20482;
$Lt_9_20738:
	mov.f32 	%f204, 0fbd6147ae;   	// -0.055
	neg.ftz.f32 	%f205, %f193;
	lg2.approx.ftz.f32 	%f206, %f205;
	mov.f32 	%f207, 0f3ed55476;   	// 0.41666
	mul.ftz.f32 	%f208, %f206, %f207;
	ex2.approx.ftz.f32 	%f209, %f208;
	neg.ftz.f32 	%f210, %f209;
	mov.f32 	%f211, 0f3f870a3d;   	// 1.055
	fma.rn.ftz.f32 	%f203, %f210, %f211, %f204;
$Lt_9_20482:
	mov.f32 	%f193, %f203;
$Lt_9_19970:
	ld.global.f32 	%f212, [inParams+120];
	ld.global.v2.f32 	{%f213,%f214}, [inParams+80];
	mul.ftz.f32 	%f215, %f213, %f156;
	ld.global.f32 	%f216, [inParams+76];
	fma.rn.ftz.f32 	%f217, %f157, %f216, %f215;
	fma.rn.ftz.f32 	%f218, %f132, %f214, %f217;
	add.ftz.f32 	%f219, %f212, %f218;
	cvt.ftz.sat.f32.f32 	%f220, %f219;
	mov.f32 	%f221, 0f3b4d2e1c;   	// 0.0031308
	setp.lt.ftz.f32 	%p15, %f220, %f221;
	@!%p15 bra 	$Lt_9_21250;
	mov.f32 	%f222, 0f414eb852;   	// 12.92
	mul.ftz.f32 	%f220, %f220, %f222;
	bra.uni 	$Lt_9_20994;
$Lt_9_21250:
	mov.f32 	%f223, 0f00000000;   	// 0
	setp.ge.ftz.f32 	%p16, %f220, %f223;
	@!%p16 bra 	$Lt_9_21762;
	mov.f32 	%f224, 0fbd6147ae;   	// -0.055
	lg2.approx.ftz.f32 	%f225, %f220;
	mov.f32 	%f226, 0f3ed55476;   	// 0.41666
	mul.ftz.f32 	%f227, %f225, %f226;
	ex2.approx.ftz.f32 	%f228, %f227;
	mov.f32 	%f229, 0f3f870a3d;   	// 1.055
	fma.rn.ftz.f32 	%f230, %f228, %f229, %f224;
	bra.uni 	$Lt_9_21506;
$Lt_9_21762:
	mov.f32 	%f231, 0fbd6147ae;   	// -0.055
	neg.ftz.f32 	%f232, %f220;
	lg2.approx.ftz.f32 	%f233, %f232;
	mov.f32 	%f234, 0f3ed55476;   	// 0.41666
	mul.ftz.f32 	%f235, %f233, %f234;
	ex2.approx.ftz.f32 	%f236, %f235;
	neg.ftz.f32 	%f237, %f236;
	mov.f32 	%f238, 0f3f870a3d;   	// 1.055
	fma.rn.ftz.f32 	%f230, %f237, %f238, %f231;
$Lt_9_21506:
	mov.f32 	%f220, %f230;
$Lt_9_20994:
	mul.ftz.f32 	%f239, %f137, %f4;
	ld.param.s32 	%r28, [__cudaparm_cuda_kernel_composite_remap];
	mov.u32 	%r29, 0;
	setp.eq.s32 	%p17, %r28, %r29;
	@%p17 bra 	$Lt_9_22018;
	ld.global.f32 	%f240, [inParams+144];
	ld.global.f32 	%f241, [inParams+140];
	fma.rn.ftz.f32 	%f242, %f239, %f241, %f240;
	cvt.ftz.sat.f32.f32 	%f239, %f242;
$Lt_9_22018:
	ld.param.s32 	%r30, [__cudaparm_cuda_kernel_composite_dstPitch];
	mul.lo.s32 	%r31, %r30, %r10;
	add.s32 	%r32, %r8, %r31;
	cvt.s64.s32 	%rd7, %r32;
	ld.param.u64 	%rd8, [__cudaparm_cuda_kernel_composite_dstFrame];
	@!%p2 bra 	$Lt_9_22786;
	mul.lo.u64 	%rd9, %rd7, 8;
	add.u64 	%rd10, %rd8, %rd9;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f166;
	mov.b32		%r33, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f193;
	mov.b32		%r34, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f220;
	mov.b32		%r35, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f239;
	mov.b32		%r36, %b1; }
	st.global.v4.u16 	[%rd10+0], {%r33,%r34,%r35,%r36};
	bra.uni 	$LBB51_cuda_kernel_composite;
$Lt_9_22786:
	mul.lo.u64 	%rd11, %rd7, 16;
	add.u64 	%rd12, %rd8, %rd11;
	st.global.v4.f32 	[%rd12+0], {%f166,%f193,%f220,%f239};
$LBB51_cuda_kernel_composite:
	exit;
$LDWend_cuda_kernel_composite:
	} // cuda_kernel_composite

	.entry cuda_kernel_showAlpha (
		.param .u64 __cudaparm_cuda_kernel_showAlpha_srcFrame,
		.param .u64 __cudaparm_cuda_kernel_showAlpha_dstFrame,
		.param .s32 __cudaparm_cuda_kernel_showAlpha_width,
		.param .s32 __cudaparm_cuda_kernel_showAlpha_height,
		.param .s32 __cudaparm_cuda_kernel_showAlpha_srcPitch,
		.param .s32 __cudaparm_cuda_kernel_showAlpha_dstPitch,
		.param .u32 __cudaparm_cuda_kernel_showAlpha_inDeviceFormat)
	{
	.reg .u32 %r<32>;
	.reg .u64 %rd<14>;
	.reg .f32 %f<7>;
	.reg .pred %p<4>;
$LDWbegin_cuda_kernel_showAlpha:
	mov.u32 	%r1, %ctaid.x;
	mov.u32 	%r2, %ntid.x;
	mul.lo.u32 	%r3, %r1, %r2;
	mov.u32 	%r4, %ctaid.y;
	mov.u32 	%r5, %ntid.y;
	mul.lo.u32 	%r6, %r4, %r5;
	mov.u32 	%r7, %tid.x;
	add.u32 	%r8, %r7, %r3;
	mov.u32 	%r9, %tid.y;
	add.u32 	%r10, %r9, %r6;
	ld.param.s32 	%r11, [__cudaparm_cuda_kernel_showAlpha_width];
	set.le.u32.s32 	%r12, %r11, %r8;
	neg.s32 	%r13, %r12;
	ld.param.s32 	%r14, [__cudaparm_cuda_kernel_showAlpha_height];
	set.le.u32.s32 	%r15, %r14, %r10;
	neg.s32 	%r16, %r15;
	or.b32 	%r17, %r13, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_10_2818;
	bra.uni 	$LBB10_cuda_kernel_showAlpha;
$Lt_10_2818:
	ld.param.u32 	%r19, [__cudaparm_cuda_kernel_showAlpha_inDeviceFormat];
	mov.s32 	%r20, 0;
	setp.eq.s32 	%p2, %r19, %r20;
	ld.param.s32 	%r21, [__cudaparm_cuda_kernel_showAlpha_srcPitch];
	mul.lo.s32 	%r22, %r21, %r10;
	add.s32 	%r23, %r8, %r22;
	cvt.s64.s32 	%rd1, %r23;
	ld.param.u64 	%rd2, [__cudaparm_cuda_kernel_showAlpha_srcFrame];
	@!%p2 bra 	$Lt_10_3586;
	mul.lo.u64 	%rd3, %rd1, 8;
	add.u64 	%rd4, %rd2, %rd3;
	ld.global.u16 	%r24, [%rd4+6];
	{ .reg .b32 %b1;
	mov.b32		%b1, %r24;
	cvt.ftz.f32.f16	%f1, %b1; }
	bra.uni 	$Lt_10_3330;
$Lt_10_3586:
	mul.lo.u64 	%rd5, %rd1, 16;
	add.u64 	%rd6, %rd2, %rd5;
	ld.global.f32 	%f1, [%rd6+12];
$Lt_10_3330:
	ld.param.s32 	%r25, [__cudaparm_cuda_kernel_showAlpha_dstPitch];
	mul.lo.s32 	%r26, %r25, %r10;
	add.s32 	%r27, %r8, %r26;
	cvt.s64.s32 	%rd7, %r27;
	ld.param.u64 	%rd8, [__cudaparm_cuda_kernel_showAlpha_dstFrame];
	@!%p2 bra 	$Lt_10_4098;
	cvt.ftz.sat.f32.f32 	%f2, %f1;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f2;
	mov.b32		%r28, %b1; }
	mov.s32 	%r29, %r28;
	mul.lo.u64 	%rd9, %rd7, 8;
	add.u64 	%rd10, %rd8, %rd9;
	mov.f32 	%f3, 0f3f800000;     	// 1
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f3;
	mov.b32		%r30, %b1; }
	st.global.v4.u16 	[%rd10+0], {%r29,%r29,%r29,%r30};
	bra.uni 	$LBB10_cuda_kernel_showAlpha;
$Lt_10_4098:
	cvt.ftz.sat.f32.f32 	%f4, %f1;
	mul.lo.u64 	%rd11, %rd7, 16;
	add.u64 	%rd12, %rd8, %rd11;
	mov.f32 	%f5, 0f3f800000;     	// 1
	st.global.v4.f32 	[%rd12+0], {%f4,%f4,%f4,%f5};
$LBB10_cuda_kernel_showAlpha:
	exit;
$LDWend_cuda_kernel_showAlpha:
	} // cuda_kernel_showAlpha

	.entry cuda_kernel_showColor (
		.param .u64 __cudaparm_cuda_kernel_showColor_srcFrame,
		.param .u64 __cudaparm_cuda_kernel_showColor_dstFrame,
		.param .s32 __cudaparm_cuda_kernel_showColor_width,
		.param .s32 __cudaparm_cuda_kernel_showColor_height,
		.param .s32 __cudaparm_cuda_kernel_showColor_srcPitch,
		.param .s32 __cudaparm_cuda_kernel_showColor_dstPitch,
		.param .u32 __cudaparm_cuda_kernel_showColor_inDeviceFormat)
	{
	.reg .u32 %r<36>;
	.reg .u64 %rd<14>;
	.reg .f32 %f<16>;
	.reg .pred %p<4>;
$LDWbegin_cuda_kernel_showColor:
	mov.u32 	%r1, %ctaid.x;
	mov.u32 	%r2, %ntid.x;
	mul.lo.u32 	%r3, %r1, %r2;
	mov.u32 	%r4, %ctaid.y;
	mov.u32 	%r5, %ntid.y;
	mul.lo.u32 	%r6, %r4, %r5;
	mov.u32 	%r7, %tid.x;
	add.u32 	%r8, %r7, %r3;
	mov.u32 	%r9, %tid.y;
	add.u32 	%r10, %r9, %r6;
	ld.param.s32 	%r11, [__cudaparm_cuda_kernel_showColor_width];
	set.le.u32.s32 	%r12, %r11, %r8;
	neg.s32 	%r13, %r12;
	ld.param.s32 	%r14, [__cudaparm_cuda_kernel_showColor_height];
	set.le.u32.s32 	%r15, %r14, %r10;
	neg.s32 	%r16, %r15;
	or.b32 	%r17, %r13, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_11_2818;
	bra.uni 	$LBB10_cuda_kernel_showColor;
$Lt_11_2818:
	ld.param.u32 	%r19, [__cudaparm_cuda_kernel_showColor_inDeviceFormat];
	mov.s32 	%r20, 0;
	setp.eq.s32 	%p2, %r19, %r20;
	ld.param.s32 	%r21, [__cudaparm_cuda_kernel_showColor_srcPitch];
	mul.lo.s32 	%r22, %r21, %r10;
	add.s32 	%r23, %r8, %r22;
	cvt.s64.s32 	%rd1, %r23;
	ld.param.u64 	%rd2, [__cudaparm_cuda_kernel_showColor_srcFrame];
	@!%p2 bra 	$Lt_11_3586;
	mul.lo.u64 	%rd3, %rd1, 8;
	add.u64 	%rd4, %rd2, %rd3;
	ld.global.v4.u16 	{%r24,%r25,%r26,%r27}, [%rd4+0];
	{ .reg .b32 %b1;
	mov.b32		%b1, %r24;
	cvt.ftz.f32.f16	%f1, %b1; }
	{ .reg .b32 %b1;
	mov.b32		%b1, %r25;
	cvt.ftz.f32.f16	%f2, %b1; }
	{ .reg .b32 %b1;
	mov.b32		%b1, %r26;
	cvt.ftz.f32.f16	%f3, %b1; }
	{ .reg .b32 %b1;
	mov.b32		%b1, %r27;
	cvt.ftz.f32.f16	%f4, %b1; }
	bra.uni 	$Lt_11_3330;
$Lt_11_3586:
	mul.lo.u64 	%rd5, %rd1, 16;
	add.u64 	%rd6, %rd2, %rd5;
	ld.global.v4.f32 	{%f1,%f2,%f3,%f4}, [%rd6+0];
$Lt_11_3330:
	ld.param.s32 	%r28, [__cudaparm_cuda_kernel_showColor_dstPitch];
	mul.lo.s32 	%r29, %r28, %r10;
	add.s32 	%r30, %r8, %r29;
	cvt.s64.s32 	%rd7, %r30;
	ld.param.u64 	%rd8, [__cudaparm_cuda_kernel_showColor_dstFrame];
	@!%p2 bra 	$Lt_11_4098;
	cvt.ftz.sat.f32.f32 	%f5, %f4;
	mul.lo.u64 	%rd9, %rd7, 8;
	add.u64 	%rd10, %rd8, %rd9;
	mul.ftz.f32 	%f6, %f5, %f1;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f6;
	mov.b32		%r31, %b1; }
	mul.ftz.f32 	%f7, %f5, %f2;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f7;
	mov.b32		%r32, %b1; }
	mul.ftz.f32 	%f8, %f5, %f3;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f8;
	mov.b32		%r33, %b1; }
	mov.f32 	%f9, 0f3f800000;     	// 1
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f9;
	mov.b32		%r34, %b1; }
	st.global.v4.u16 	[%rd10+0], {%r31,%r32,%r33,%r34};
	bra.uni 	$LBB10_cuda_kernel_showColor;
$Lt_11_4098:
	cvt.ftz.sat.f32.f32 	%f10, %f4;
	mul.lo.u64 	%rd11, %rd7, 16;
	add.u64 	%rd12, %rd8, %rd11;
	mul.ftz.f32 	%f11, %f10, %f1;
	mul.ftz.f32 	%f12, %f10, %f2;
	mul.ftz.f32 	%f13, %f10, %f3;
	mov.f32 	%f14, 0f3f800000;    	// 1
	st.global.v4.f32 	[%rd12+0], {%f11,%f12,%f13,%f14};
$LBB10_cuda_kernel_showColor:
	exit;
$LDWend_cuda_kernel_showColor:
	} // cuda_kernel_showColor

