	.version 2.2
	.target sm_20
	// compiled with ../../../External/3rdParty/NVIDIA/CUDA/win/bin/../open64/lib//be.exe
	// nvopencc 3.2 built on 2010-11-04

	.visible .func (.param .f32 __cudaretf__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf) _Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf (.param .u64 __cudaparmf1__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf, .param .u64 __cudaparmf2__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf, .param .u64 __cudaparmf3__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf, .param .u64 __cudaparmf4__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf, .param .f32 __cudaparmf5__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf)

	.visible .func (.param .s32 __cudaretf__Z15IntegerMultiplyii) _Z15IntegerMultiplyii (.param .s32 __cudaparmf1__Z15IntegerMultiplyii, .param .s32 __cudaparmf2__Z15IntegerMultiplyii)

	.visible .func (.param .s32 __cudaretf__Z17Standard2DKernelXv) _Z17Standard2DKernelXv ()

	.visible .func (.param .s32 __cudaretf__Z17Standard2DKernelYv) _Z17Standard2DKernelYv ()

	.visible .func (.param .align 16 .b8 __cudaretf__Z13Half4ToFloat47ushort4[16]) _Z13Half4ToFloat47ushort4 (.param .align 8 .b8 __cudaparmf1__Z13Half4ToFloat47ushort4[8])

	.visible .func (.param .align 8 .b8 __cudaretf__Z13Float4ToHalf46float4[8]) _Z13Float4ToHalf46float4 (.param .align 16 .b8 __cudaparmf1__Z13Float4ToHalf46float4[16])

	.visible .func (.param .u32 __cudaretf__Z4Mix3RjS_S_) _Z4Mix3RjS_S_ (.param .u64 __cudaparmf1__Z4Mix3RjS_S_, .param .u64 __cudaparmf2__Z4Mix3RjS_S_, .param .u64 __cudaparmf3__Z4Mix3RjS_S_)

	.visible .func (.param .s32 __cudaretf__Z4Randj) _Z4Randj (.param .u32 __cudaparmf1__Z4Randj)

	.visible .func (.param .s32 __cudaretf__Z6Rand2Djjj) _Z6Rand2Djjj (.param .u32 __cudaparmf1__Z6Rand2Djjj, .param .u32 __cudaparmf2__Z6Rand2Djjj, .param .u32 __cudaparmf3__Z6Rand2Djjj)

	.visible .func (.param .s32 __cudaretf__Z6Rand2Dj) _Z6Rand2Dj (.param .u32 __cudaparmf1__Z6Rand2Dj)

	.visible .func (.param .align 8 .b8 __cudaretf__Z6Read2DI7ushort4ET_PKS1_iii[8]) _Z6Read2DI7ushort4ET_PKS1_iii (.param .u64 __cudaparmf1__Z6Read2DI7ushort4ET_PKS1_iii, .param .s32 __cudaparmf2__Z6Read2DI7ushort4ET_PKS1_iii, .param .s32 __cudaparmf3__Z6Read2DI7ushort4ET_PKS1_iii, .param .s32 __cudaparmf4__Z6Read2DI7ushort4ET_PKS1_iii)

	.visible .func (.param .align 16 .b8 __cudaretf__Z6Read2DI6float4ET_PKS1_iii[16]) _Z6Read2DI6float4ET_PKS1_iii (.param .u64 __cudaparmf1__Z6Read2DI6float4ET_PKS1_iii, .param .s32 __cudaparmf2__Z6Read2DI6float4ET_PKS1_iii, .param .s32 __cudaparmf3__Z6Read2DI6float4ET_PKS1_iii, .param .s32 __cudaparmf4__Z6Read2DI6float4ET_PKS1_iii)

	.visible .func _Z7Write2DI7ushort4EvT_PS1_iii (.param .align 8 .b8 __cudaparmf1__Z7Write2DI7ushort4EvT_PS1_iii[8], .param .u64 __cudaparmf2__Z7Write2DI7ushort4EvT_PS1_iii, .param .s32 __cudaparmf3__Z7Write2DI7ushort4EvT_PS1_iii, .param .s32 __cudaparmf4__Z7Write2DI7ushort4EvT_PS1_iii, .param .s32 __cudaparmf5__Z7Write2DI7ushort4EvT_PS1_iii)

	.visible .func _Z7Write2DI6float4EvT_PS1_iii (.param .align 16 .b8 __cudaparmf1__Z7Write2DI6float4EvT_PS1_iii[16], .param .u64 __cudaparmf2__Z7Write2DI6float4EvT_PS1_iii, .param .s32 __cudaparmf3__Z7Write2DI6float4EvT_PS1_iii, .param .s32 __cudaparmf4__Z7Write2DI6float4EvT_PS1_iii, .param .s32 __cudaparmf5__Z7Write2DI6float4EvT_PS1_iii)

	.visible .func (.param .align 16 .b8 __cudaretf__Z18UnpremultiplyPixel8PixelRGB[16]) _Z18UnpremultiplyPixel8PixelRGB (.param .align 16 .b8 __cudaparmf1__Z18UnpremultiplyPixel8PixelRGB[16])

	.visible .func (.param .f32 __cudaretf__Z13ToLinearColorf) _Z13ToLinearColorf (.param .f32 __cudaparmf1__Z13ToLinearColorf)

	.visible .func (.param .f32 __cudaretf__Z15FromLinearColorf) _Z15FromLinearColorf (.param .f32 __cudaparmf1__Z15FromLinearColorf)

	.visible .func (.param .align 16 .b8 __cudaretf__Z25PremultiplyLinearizePixel8PixelRGB[16]) _Z25PremultiplyLinearizePixel8PixelRGB (.param .align 16 .b8 __cudaparmf1__Z25PremultiplyLinearizePixel8PixelRGB[16])

	.visible .func (.param .align 16 .b8 __cudaretf__Z29UnpremultiplyUnlinearizePixel8PixelRGB[16]) _Z29UnpremultiplyUnlinearizePixel8PixelRGB (.param .align 16 .b8 __cudaparmf1__Z29UnpremultiplyUnlinearizePixel8PixelRGB[16])

	.visible .func (.param .align 16 .b8 __cudaretf__Z20PremultiplyLinearize6float4[16]) _Z20PremultiplyLinearize6float4 (.param .align 16 .b8 __cudaparmf1__Z20PremultiplyLinearize6float4[16])

	.visible .func (.param .align 16 .b8 __cudaretf__Z24UnpremultiplyUnlinearize6float4[16]) _Z24UnpremultiplyUnlinearize6float4 (.param .align 16 .b8 __cudaparmf1__Z24UnpremultiplyUnlinearize6float4[16])

	.visible .func (.param .align 16 .b8 __cudaretf__Z18SwapComponentOrderI6float4ET_RKS1_[16]) _Z18SwapComponentOrderI6float4ET_RKS1_ (.param .u64 __cudaparmf1__Z18SwapComponentOrderI6float4ET_RKS1_)

	.visible .func (.param .align 16 .b8 __cudaretf__Z13CalculateBGRA8PixelRGBf[16]) _Z13CalculateBGRA8PixelRGBf (.param .align 16 .b8 __cudaparmf1__Z13CalculateBGRA8PixelRGBf[16], .param .f32 __cudaparmf2__Z13CalculateBGRA8PixelRGBf)

	.visible .func (.param .f32 __cudaretf__Z4LERPIfET_S0_S0_S0_) _Z4LERPIfET_S0_S0_S0_ (.param .f32 __cudaparmf1__Z4LERPIfET_S0_S0_S0_, .param .f32 __cudaparmf2__Z4LERPIfET_S0_S0_S0_, .param .f32 __cudaparmf3__Z4LERPIfET_S0_S0_S0_)

	.visible .func (.param .f32 __cudaretf__Z6Read2DIfET_PKS0_iii) _Z6Read2DIfET_PKS0_iii (.param .u64 __cudaparmf1__Z6Read2DIfET_PKS0_iii, .param .s32 __cudaparmf2__Z6Read2DIfET_PKS0_iii, .param .s32 __cudaparmf3__Z6Read2DIfET_PKS0_iii, .param .s32 __cudaparmf4__Z6Read2DIfET_PKS0_iii)

	//-----------------------------------------------------------
	// Compiling C:/Users/dvaeng/AppData/Local/Temp/tmpxft_00003870_00000000-11_RGBCurves.cpp3.i (C:/Users/dvaeng/AppData/Local/Temp/ccBI#.a15064)
	//-----------------------------------------------------------

	//-----------------------------------------------------------
	// 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_00003870_00000000-10_RGBCurves.cudafe2.gpu"
	.file	2	"c:\Mulder64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/PixelFormat.h"
	.file	3	"c:\Mulder64\shared\adobe\MediaCore\Display\Inc\CUDA/Effects/ColorCorrector_Curves.h"
	.file	4	"c:\Mulder64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/PixelRGB.h"
	.file	5	"c:\Mulder64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/PixelYUV.h"
	.file	6	"C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\include\crtdefs.h"
	.file	7	"c:\Mulder64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\include\crt/device_runtime.h"
	.file	8	"c:\Mulder64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\include\host_defines.h"
	.file	9	"c:\Mulder64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\include\builtin_types.h"
	.file	10	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\device_types.h"
	.file	11	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\driver_types.h"
	.file	12	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\surface_types.h"
	.file	13	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\texture_types.h"
	.file	14	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\vector_types.h"
	.file	15	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\builtin_types.h"
	.file	16	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\host_defines.h"
	.file	17	"c:\Mulder64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\include\device_launch_parameters.h"
	.file	18	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\crt\storage_class.h"
	.file	19	"C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\include\time.h"
	.file	20	"c:\Mulder64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/Utils.h"
	.file	21	"c:\Mulder64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/VectorUtils.h"
	.file	22	"c:/Mulder64/shared/adobe/MediaCore/Display/Src/CUDA/Effects/RGBCurves.cu"
	.file	23	"c:\Mulder64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/Numeric.h"
	.file	24	"c:\Mulder64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\include\common_functions.h"
	.file	25	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\math_functions.h"
	.file	26	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\math_constants.h"
	.file	27	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\device_functions.h"
	.file	28	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\sm_11_atomic_functions.h"
	.file	29	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\sm_12_atomic_functions.h"
	.file	30	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\sm_13_double_functions.h"
	.file	31	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\sm_20_atomic_functions.h"
	.file	32	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\sm_20_intrinsics.h"
	.file	33	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\surface_functions.h"
	.file	34	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\texture_fetch_functions.h"
	.file	35	"c:\mulder64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\include\math_functions_dbl_ptx3.h"
	.file	36	"c:\Mulder64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/ColorSpaceConvert.h"


	.visible .func (.param .f32 __cudaretf__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf) _Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf (.param .u64 __cudaparmf1__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf, .param .u64 __cudaparmf2__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf, .param .u64 __cudaparmf3__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf, .param .u64 __cudaparmf4__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf, .param .f32 __cudaparmf5__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf)
	{
	.reg .u64 %rd<28>;
	.reg .f32 %f<33>;
	.reg .pred %p<6>;
	.loc	3	53	0
$LDWbegin__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf:
	ld.param.u64 	%rd1, [__cudaparmf1__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf];
	mov.s64 	%rd2, %rd1;
	ld.param.u64 	%rd3, [__cudaparmf2__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf];
	mov.s64 	%rd4, %rd3;
	ld.param.u64 	%rd5, [__cudaparmf3__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf];
	mov.s64 	%rd6, %rd5;
	ld.param.u64 	%rd7, [__cudaparmf4__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf];
	mov.s64 	%rd8, %rd7;
	ld.param.f32 	%f1, [__cudaparmf5__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf];
	mov.f32 	%f2, %f1;
	.loc	3	54	0
	mov.f32 	%f3, %f2;
	.loc	3	58	0
	sub.u64 	%rd9, %rd8, 1;
	mov.s64 	%rd10, %rd9;
	mov.u64 	%rd11, 1;
	setp.le.u64 	%p1, %rd9, %rd11;
	@%p1 bra 	$Lt_0_4866;
	mov.u64 	%rd12, 0;
$Lt_0_3330:
 //<loop> Loop body line 58
	add.u64 	%rd13, %rd10, %rd12;
	shr.u64 	%rd14, %rd13, 1;
	mul.lo.u64 	%rd15, %rd14, 4;
	add.u64 	%rd16, %rd2, %rd15;
	ld.f32 	%f4, [%rd16+0];
	setp.gt.ftz.f32 	%p2, %f4, %f2;
	@!%p2 bra 	$Lt_0_3842;
 //<loop> Part of loop body line 58, head labeled $Lt_0_3330
	.loc	3	67	0
	mov.s64 	%rd10, %rd14;
	bra.uni 	$Lt_0_3586;
$Lt_0_3842:
 //<loop> Part of loop body line 58, head labeled $Lt_0_3330
	.loc	3	71	0
	mov.s64 	%rd12, %rd14;
$Lt_0_3586:
 //<loop> Part of loop body line 58, head labeled $Lt_0_3330
	sub.u64 	%rd17, %rd10, %rd12;
	mov.u64 	%rd18, 1;
	setp.gt.u64 	%p3, %rd17, %rd18;
	@%p3 bra 	$Lt_0_3330;
	bra.uni 	$Lt_0_2818;
$Lt_0_4866:
	mov.u64 	%rd12, 0;
$Lt_0_2818:
	.loc	3	75	0
	mul.lo.u64 	%rd19, %rd10, 4;
	mul.lo.u64 	%rd20, %rd12, 4;
	add.u64 	%rd21, %rd19, %rd2;
	add.u64 	%rd22, %rd20, %rd2;
	ld.f32 	%f5, [%rd21+0];
	ld.f32 	%f6, [%rd22+0];
	sub.ftz.f32 	%f7, %f5, %f6;
	mov.f32 	%f8, 0f00000000;     	// 0
	setp.neu.ftz.f32 	%p4, %f7, %f8;
	@!%p4 bra 	$Lt_0_4354;
	.loc	3	78	0
	sub.ftz.f32 	%f9, %f5, %f2;
	.loc	3	79	0
	sub.ftz.f32 	%f10, %f2, %f6;
	.loc	3	80	0
	div.approx.ftz.f32 	%f11, %f10, %f7;
	div.approx.ftz.f32 	%f12, %f9, %f7;
	add.u64 	%rd23, %rd19, %rd6;
	ld.f32 	%f13, [%rd23+0];
	mul.ftz.f32 	%f14, %f11, %f11;
	mul.ftz.f32 	%f15, %f11, %f14;
	sub.ftz.f32 	%f16, %f15, %f11;
	mul.ftz.f32 	%f17, %f13, %f16;
	add.u64 	%rd24, %rd20, %rd6;
	ld.f32 	%f18, [%rd24+0];
	mul.ftz.f32 	%f19, %f12, %f12;
	mul.ftz.f32 	%f20, %f12, %f19;
	sub.ftz.f32 	%f21, %f20, %f12;
	fma.rn.ftz.f32 	%f22, %f18, %f21, %f17;
	mul.ftz.f32 	%f23, %f7, %f7;
	mul.ftz.f32 	%f24, %f22, %f23;
	mov.f32 	%f25, 0f40c00000;    	// 6
	div.approx.ftz.f32 	%f26, %f24, %f25;
	add.u64 	%rd25, %rd19, %rd4;
	ld.f32 	%f27, [%rd25+0];
	mul.ftz.f32 	%f28, %f27, %f11;
	add.u64 	%rd26, %rd20, %rd4;
	ld.f32 	%f29, [%rd26+0];
	fma.rn.ftz.f32 	%f30, %f29, %f12, %f28;
	add.ftz.f32 	%f3, %f26, %f30;
$Lt_0_4354:
	.loc	3	83	0
	mov.f32 	%f31, %f3;
	st.param.f32 	[__cudaretf__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf], %f31;
	ret;
$LDWend__Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf:
	} // _Z31InterpolateNaturalCubicSpline2DPKfS0_S0_yf

	.visible .func (.param .s32 __cudaretf__Z15IntegerMultiplyii) _Z15IntegerMultiplyii (.param .s32 __cudaparmf1__Z15IntegerMultiplyii, .param .s32 __cudaparmf2__Z15IntegerMultiplyii)
	{
	.reg .u32 %r<7>;
	.loc	20	60	0
$LDWbegin__Z15IntegerMultiplyii:
	ld.param.u32 	%r1, [__cudaparmf1__Z15IntegerMultiplyii];
	mov.s32 	%r2, %r1;
	ld.param.u32 	%r3, [__cudaparmf2__Z15IntegerMultiplyii];
	mov.s32 	%r4, %r3;
	.loc	20	64	0
	mul.lo.s32 	%r5, %r2, %r4;
	st.param.s32 	[__cudaretf__Z15IntegerMultiplyii], %r5;
	ret;
$LDWend__Z15IntegerMultiplyii:
	} // _Z15IntegerMultiplyii

	.visible .func (.param .s32 __cudaretf__Z17Standard2DKernelXv) _Z17Standard2DKernelXv ()
	{
	.reg .u32 %r<7>;
	.loc	20	73	0
$LDWbegin__Z17Standard2DKernelXv:
	.loc	20	74	0
	mov.u32 	%r1, %tid.x;
	cvt.s32.u32 	%r2, %ctaid.x;
	cvt.s32.u32 	%r3, %ntid.x;
	mul.lo.s32 	%r4, %r2, %r3;
	add.u32 	%r5, %r1, %r4;
	st.param.s32 	[__cudaretf__Z17Standard2DKernelXv], %r5;
	ret;
$LDWend__Z17Standard2DKernelXv:
	} // _Z17Standard2DKernelXv

	.visible .func (.param .s32 __cudaretf__Z17Standard2DKernelYv) _Z17Standard2DKernelYv ()
	{
	.reg .u32 %r<7>;
	.loc	20	77	0
$LDWbegin__Z17Standard2DKernelYv:
	.loc	20	78	0
	mov.u32 	%r1, %tid.y;
	cvt.s32.u32 	%r2, %ctaid.y;
	cvt.s32.u32 	%r3, %ntid.y;
	mul.lo.s32 	%r4, %r2, %r3;
	add.u32 	%r5, %r1, %r4;
	st.param.s32 	[__cudaretf__Z17Standard2DKernelYv], %r5;
	ret;
$LDWend__Z17Standard2DKernelYv:
	} // _Z17Standard2DKernelYv

	.visible .func (.param .align 16 .b8 __cudaretf__Z13Half4ToFloat47ushort4[16]) _Z13Half4ToFloat47ushort4 (.param .align 8 .b8 __cudaparmf1__Z13Half4ToFloat47ushort4[8])
	{
	.reg .u32 %r<14>;
	.reg .f32 %f<9>;
	.loc	20	86	0
$LDWbegin__Z13Half4ToFloat47ushort4:
	ld.param.u16 	%r1, [__cudaparmf1__Z13Half4ToFloat47ushort4+0];
	mov.s32 	%r2, %r1;
	ld.param.u16 	%r3, [__cudaparmf1__Z13Half4ToFloat47ushort4+2];
	mov.s32 	%r4, %r3;
	ld.param.u16 	%r5, [__cudaparmf1__Z13Half4ToFloat47ushort4+4];
	mov.s32 	%r6, %r5;
	ld.param.u16 	%r7, [__cudaparmf1__Z13Half4ToFloat47ushort4+6];
	mov.s32 	%r8, %r7;
	.loc	20	87	0
	cvt.u16.u32 	%r9, %r4;
	{ .reg .b32 %b1;
	mov.b32		%b1, %r9;
	cvt.ftz.f32.f16	%f1, %b1; }
	cvt.u16.u32 	%r10, %r6;
	{ .reg .b32 %b1;
	mov.b32		%b1, %r10;
	cvt.ftz.f32.f16	%f2, %b1; }
	cvt.u16.u32 	%r11, %r8;
	{ .reg .b32 %b1;
	mov.b32		%b1, %r11;
	cvt.ftz.f32.f16	%f3, %b1; }
	cvt.u16.u32 	%r12, %r2;
	{ .reg .b32 %b1;
	mov.b32		%b1, %r12;
	cvt.ftz.f32.f16	%f4, %b1; }
	st.param.f32 	[__cudaretf__Z13Half4ToFloat47ushort4+0], %f4;
	mov.f32 	%f5, %f1;
	st.param.f32 	[__cudaretf__Z13Half4ToFloat47ushort4+4], %f5;
	mov.f32 	%f6, %f2;
	st.param.f32 	[__cudaretf__Z13Half4ToFloat47ushort4+8], %f6;
	mov.f32 	%f7, %f3;
	st.param.f32 	[__cudaretf__Z13Half4ToFloat47ushort4+12], %f7;
	ret;
$LDWend__Z13Half4ToFloat47ushort4:
	} // _Z13Half4ToFloat47ushort4

	.visible .func (.param .align 8 .b8 __cudaretf__Z13Float4ToHalf46float4[8]) _Z13Float4ToHalf46float4 (.param .align 16 .b8 __cudaparmf1__Z13Float4ToHalf46float4[16])
	{
	.reg .u32 %r<13>;
	.reg .f32 %f<10>;
	.loc	20	95	0
$LDWbegin__Z13Float4ToHalf46float4:
	ld.param.f32 	%f1, [__cudaparmf1__Z13Float4ToHalf46float4+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z13Float4ToHalf46float4+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z13Float4ToHalf46float4+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z13Float4ToHalf46float4+12];
	mov.f32 	%f8, %f7;
	.loc	20	96	0
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f4;
	mov.b32		%r1, %b1; }
	cvt.u16.u32 	%r2, %r1;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f6;
	mov.b32		%r3, %b1; }
	cvt.u16.u32 	%r4, %r3;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f8;
	mov.b32		%r5, %b1; }
	cvt.u16.u32 	%r6, %r5;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f2;
	mov.b32		%r7, %b1; }
	cvt.u16.u32 	%r8, %r7;
	st.param.u16 	[__cudaretf__Z13Float4ToHalf46float4+0], %r8;
	mov.s32 	%r9, %r2;
	st.param.u16 	[__cudaretf__Z13Float4ToHalf46float4+2], %r9;
	mov.s32 	%r10, %r4;
	st.param.u16 	[__cudaretf__Z13Float4ToHalf46float4+4], %r10;
	mov.s32 	%r11, %r6;
	st.param.u16 	[__cudaretf__Z13Float4ToHalf46float4+6], %r11;
	ret;
$LDWend__Z13Float4ToHalf46float4:
	} // _Z13Float4ToHalf46float4

	.visible .func (.param .u32 __cudaretf__Z4Mix3RjS_S_) _Z4Mix3RjS_S_ (.param .u64 __cudaparmf1__Z4Mix3RjS_S_, .param .u64 __cudaparmf2__Z4Mix3RjS_S_, .param .u64 __cudaparmf3__Z4Mix3RjS_S_)
	{
	.reg .u32 %r<75>;
	.reg .u64 %rd<8>;
	.loc	20	138	0
$LDWbegin__Z4Mix3RjS_S_:
	ld.param.u64 	%rd1, [__cudaparmf1__Z4Mix3RjS_S_];
	mov.s64 	%rd2, %rd1;
	ld.param.u64 	%rd3, [__cudaparmf2__Z4Mix3RjS_S_];
	mov.s64 	%rd4, %rd3;
	ld.param.u64 	%rd5, [__cudaparmf3__Z4Mix3RjS_S_];
	mov.s64 	%rd6, %rd5;
	.loc	20	139	0
	ld.u32 	%r1, [%rd2+0];
	ld.u32 	%r2, [%rd4+0];
	sub.u32 	%r3, %r1, %r2;
	st.u32 	[%rd2+0], %r3;
	ld.u32 	%r4, [%rd6+0];
	sub.u32 	%r5, %r3, %r4;
	st.u32 	[%rd2+0], %r5;
	ld.u32 	%r6, [%rd6+0];
	shr.u32 	%r7, %r6, 13;
	xor.b32 	%r8, %r5, %r7;
	st.u32 	[%rd2+0], %r8;
	.loc	20	140	0
	ld.u32 	%r9, [%rd4+0];
	ld.u32 	%r10, [%rd6+0];
	sub.u32 	%r11, %r9, %r10;
	st.u32 	[%rd4+0], %r11;
	ld.u32 	%r12, [%rd2+0];
	sub.u32 	%r13, %r11, %r12;
	st.u32 	[%rd4+0], %r13;
	ld.u32 	%r14, [%rd2+0];
	shl.b32 	%r15, %r14, 8;
	xor.b32 	%r16, %r13, %r15;
	st.u32 	[%rd4+0], %r16;
	.loc	20	141	0
	ld.u32 	%r17, [%rd6+0];
	ld.u32 	%r18, [%rd2+0];
	sub.u32 	%r19, %r17, %r18;
	st.u32 	[%rd6+0], %r19;
	ld.u32 	%r20, [%rd4+0];
	sub.u32 	%r21, %r19, %r20;
	st.u32 	[%rd6+0], %r21;
	ld.u32 	%r22, [%rd4+0];
	shr.u32 	%r23, %r22, 13;
	xor.b32 	%r24, %r21, %r23;
	st.u32 	[%rd6+0], %r24;
	.loc	20	142	0
	ld.u32 	%r25, [%rd2+0];
	ld.u32 	%r26, [%rd4+0];
	sub.u32 	%r27, %r25, %r26;
	st.u32 	[%rd2+0], %r27;
	ld.u32 	%r28, [%rd6+0];
	sub.u32 	%r29, %r27, %r28;
	st.u32 	[%rd2+0], %r29;
	ld.u32 	%r30, [%rd6+0];
	shr.u32 	%r31, %r30, 12;
	xor.b32 	%r32, %r29, %r31;
	st.u32 	[%rd2+0], %r32;
	.loc	20	143	0
	ld.u32 	%r33, [%rd4+0];
	ld.u32 	%r34, [%rd6+0];
	sub.u32 	%r35, %r33, %r34;
	st.u32 	[%rd4+0], %r35;
	ld.u32 	%r36, [%rd2+0];
	sub.u32 	%r37, %r35, %r36;
	st.u32 	[%rd4+0], %r37;
	ld.u32 	%r38, [%rd2+0];
	shl.b32 	%r39, %r38, 16;
	xor.b32 	%r40, %r37, %r39;
	st.u32 	[%rd4+0], %r40;
	.loc	20	144	0
	ld.u32 	%r41, [%rd6+0];
	ld.u32 	%r42, [%rd2+0];
	sub.u32 	%r43, %r41, %r42;
	st.u32 	[%rd6+0], %r43;
	ld.u32 	%r44, [%rd4+0];
	sub.u32 	%r45, %r43, %r44;
	st.u32 	[%rd6+0], %r45;
	ld.u32 	%r46, [%rd4+0];
	shr.u32 	%r47, %r46, 5;
	xor.b32 	%r48, %r45, %r47;
	st.u32 	[%rd6+0], %r48;
	.loc	20	145	0
	ld.u32 	%r49, [%rd2+0];
	ld.u32 	%r50, [%rd4+0];
	sub.u32 	%r51, %r49, %r50;
	st.u32 	[%rd2+0], %r51;
	ld.u32 	%r52, [%rd6+0];
	sub.u32 	%r53, %r51, %r52;
	st.u32 	[%rd2+0], %r53;
	ld.u32 	%r54, [%rd6+0];
	shr.u32 	%r55, %r54, 3;
	xor.b32 	%r56, %r53, %r55;
	st.u32 	[%rd2+0], %r56;
	.loc	20	146	0
	ld.u32 	%r57, [%rd4+0];
	ld.u32 	%r58, [%rd6+0];
	sub.u32 	%r59, %r57, %r58;
	st.u32 	[%rd4+0], %r59;
	ld.u32 	%r60, [%rd2+0];
	sub.u32 	%r61, %r59, %r60;
	st.u32 	[%rd4+0], %r61;
	ld.u32 	%r62, [%rd2+0];
	shl.b32 	%r63, %r62, 10;
	xor.b32 	%r64, %r61, %r63;
	st.u32 	[%rd4+0], %r64;
	.loc	20	147	0
	ld.u32 	%r65, [%rd6+0];
	ld.u32 	%r66, [%rd2+0];
	sub.u32 	%r67, %r65, %r66;
	st.u32 	[%rd6+0], %r67;
	ld.u32 	%r68, [%rd4+0];
	sub.u32 	%r69, %r67, %r68;
	st.u32 	[%rd6+0], %r69;
	ld.u32 	%r70, [%rd4+0];
	shr.u32 	%r71, %r70, 15;
	xor.b32 	%r72, %r69, %r71;
	st.u32 	[%rd6+0], %r72;
	.loc	20	148	0
	mov.s32 	%r73, %r72;
	st.param.u32 	[__cudaretf__Z4Mix3RjS_S_], %r73;
	ret;
$LDWend__Z4Mix3RjS_S_:
	} // _Z4Mix3RjS_S_

	.visible .func (.param .s32 __cudaretf__Z4Randj) _Z4Randj (.param .u32 __cudaparmf1__Z4Randj)
	{
	.reg .u32 %r<14>;
	.loc	20	152	0
$LDWbegin__Z4Randj:
	ld.param.u32 	%r1, [__cudaparmf1__Z4Randj];
	mov.s32 	%r2, %r1;
	.loc	20	163	0
	mul.lo.u32 	%r3, %r2, 1103515245;
	add.u32 	%r4, %r3, 12345;
	shr.u32 	%r5, %r4, 16;
	and.b32 	%r6, %r5, 255;
	shl.b32 	%r7, %r6, 7;
	mul.lo.u32 	%r8, %r2, -1029531031;
	sub.u32 	%r9, %r8, 740551042;
	shr.u32 	%r10, %r9, 16;
	and.b32 	%r11, %r10, 255;
	xor.b32 	%r12, %r7, %r11;
	st.param.s32 	[__cudaretf__Z4Randj], %r12;
	ret;
$LDWend__Z4Randj:
	} // _Z4Randj

	.visible .func (.param .s32 __cudaretf__Z6Rand2Djjj) _Z6Rand2Djjj (.param .u32 __cudaparmf1__Z6Rand2Djjj, .param .u32 __cudaparmf2__Z6Rand2Djjj, .param .u32 __cudaparmf3__Z6Rand2Djjj)
	{
	.reg .u32 %r<54>;
	.loc	20	169	0
$LDWbegin__Z6Rand2Djjj:
	ld.param.u32 	%r1, [__cudaparmf1__Z6Rand2Djjj];
	mov.s32 	%r2, %r1;
	ld.param.u32 	%r3, [__cudaparmf2__Z6Rand2Djjj];
	mov.s32 	%r4, %r3;
	ld.param.u32 	%r5, [__cudaparmf3__Z6Rand2Djjj];
	mov.s32 	%r6, %r5;
	.loc	20	139	0
	sub.u32 	%r7, %r2, %r4;
	sub.u32 	%r8, %r7, %r6;
	shr.u32 	%r9, %r6, 13;
	xor.b32 	%r10, %r8, %r9;
	.loc	20	140	0
	sub.u32 	%r11, %r4, %r6;
	sub.u32 	%r12, %r11, %r10;
	shl.b32 	%r13, %r10, 8;
	xor.b32 	%r14, %r12, %r13;
	.loc	20	141	0
	sub.u32 	%r15, %r6, %r10;
	sub.u32 	%r16, %r15, %r14;
	shr.u32 	%r17, %r14, 13;
	xor.b32 	%r18, %r16, %r17;
	.loc	20	142	0
	sub.u32 	%r19, %r10, %r14;
	sub.u32 	%r20, %r19, %r18;
	shr.u32 	%r21, %r18, 12;
	xor.b32 	%r22, %r20, %r21;
	.loc	20	143	0
	sub.u32 	%r23, %r14, %r18;
	sub.u32 	%r24, %r23, %r22;
	shl.b32 	%r25, %r22, 16;
	xor.b32 	%r26, %r24, %r25;
	.loc	20	144	0
	sub.u32 	%r27, %r18, %r22;
	sub.u32 	%r28, %r27, %r26;
	shr.u32 	%r29, %r26, 5;
	xor.b32 	%r30, %r28, %r29;
	.loc	20	145	0
	sub.u32 	%r31, %r22, %r26;
	sub.u32 	%r32, %r31, %r30;
	shr.u32 	%r33, %r30, 3;
	xor.b32 	%r34, %r32, %r33;
	.loc	20	146	0
	sub.u32 	%r35, %r26, %r30;
	sub.u32 	%r36, %r35, %r34;
	shl.b32 	%r37, %r34, 10;
	xor.b32 	%r38, %r36, %r37;
	.loc	20	147	0
	sub.u32 	%r39, %r30, %r34;
	sub.u32 	%r40, %r39, %r38;
	shr.u32 	%r41, %r38, 15;
	xor.b32 	%r42, %r40, %r41;
	.loc	20	170	0
	mul.lo.u32 	%r43, %r42, 1103515245;
	add.u32 	%r44, %r43, 12345;
	shr.u32 	%r45, %r44, 16;
	and.b32 	%r46, %r45, 255;
	shl.b32 	%r47, %r46, 7;
	mul.lo.u32 	%r48, %r42, -1029531031;
	sub.u32 	%r49, %r48, 740551042;
	shr.u32 	%r50, %r49, 16;
	and.b32 	%r51, %r50, 255;
	xor.b32 	%r52, %r47, %r51;
	st.param.s32 	[__cudaretf__Z6Rand2Djjj], %r52;
	ret;
$LDWend__Z6Rand2Djjj:
	} // _Z6Rand2Djjj

	.visible .func (.param .s32 __cudaretf__Z6Rand2Dj) _Z6Rand2Dj (.param .u32 __cudaparmf1__Z6Rand2Dj)
	{
	.reg .u32 %r<60>;
	.loc	20	175	0
$LDWbegin__Z6Rand2Dj:
	ld.param.u32 	%r1, [__cudaparmf1__Z6Rand2Dj];
	mov.s32 	%r2, %r1;
	.loc	20	143	0
	cvt.s32.u32 	%r3, %ctaid.y;
	cvt.s32.u32 	%r4, %ntid.y;
	mul.lo.s32 	%r5, %r3, %r4;
	cvt.s32.u32 	%r6, %ctaid.x;
	cvt.s32.u32 	%r7, %ntid.x;
	mul.lo.s32 	%r8, %r6, %r7;
	mov.u32 	%r9, %tid.y;
	add.u32 	%r10, %r5, %r9;
	mov.u32 	%r11, %tid.x;
	add.u32 	%r12, %r8, %r11;
	shr.u32 	%r13, %r10, 13;
	sub.u32 	%r14, %r2, %r12;
	sub.u32 	%r15, %r12, %r10;
	sub.u32 	%r16, %r14, %r10;
	xor.b32 	%r17, %r13, %r16;
	shl.b32 	%r18, %r17, 8;
	sub.u32 	%r19, %r15, %r17;
	sub.u32 	%r20, %r10, %r17;
	xor.b32 	%r21, %r18, %r19;
	shr.u32 	%r22, %r21, 13;
	sub.u32 	%r23, %r20, %r21;
	sub.u32 	%r24, %r17, %r21;
	xor.b32 	%r25, %r22, %r23;
	shr.u32 	%r26, %r25, 12;
	sub.u32 	%r27, %r24, %r25;
	xor.b32 	%r28, %r26, %r27;
	sub.u32 	%r29, %r21, %r25;
	sub.u32 	%r30, %r29, %r28;
	shl.b32 	%r31, %r28, 16;
	xor.b32 	%r32, %r30, %r31;
	.loc	20	144	0
	sub.u32 	%r33, %r25, %r28;
	sub.u32 	%r34, %r33, %r32;
	shr.u32 	%r35, %r32, 5;
	xor.b32 	%r36, %r34, %r35;
	.loc	20	145	0
	sub.u32 	%r37, %r28, %r32;
	sub.u32 	%r38, %r37, %r36;
	shr.u32 	%r39, %r36, 3;
	xor.b32 	%r40, %r38, %r39;
	.loc	20	146	0
	sub.u32 	%r41, %r32, %r36;
	sub.u32 	%r42, %r41, %r40;
	shl.b32 	%r43, %r40, 10;
	xor.b32 	%r44, %r42, %r43;
	.loc	20	147	0
	sub.u32 	%r45, %r36, %r40;
	sub.u32 	%r46, %r45, %r44;
	shr.u32 	%r47, %r44, 15;
	xor.b32 	%r48, %r46, %r47;
	.loc	20	176	0
	mul.lo.u32 	%r49, %r48, 1103515245;
	add.u32 	%r50, %r49, 12345;
	shr.u32 	%r51, %r50, 16;
	and.b32 	%r52, %r51, 255;
	shl.b32 	%r53, %r52, 7;
	mul.lo.u32 	%r54, %r48, -1029531031;
	sub.u32 	%r55, %r54, 740551042;
	shr.u32 	%r56, %r55, 16;
	and.b32 	%r57, %r56, 255;
	xor.b32 	%r58, %r53, %r57;
	st.param.s32 	[__cudaretf__Z6Rand2Dj], %r58;
	ret;
$LDWend__Z6Rand2Dj:
	} // _Z6Rand2Dj

	.visible .func (.param .align 8 .b8 __cudaretf__Z6Read2DI7ushort4ET_PKS1_iii[8]) _Z6Read2DI7ushort4ET_PKS1_iii (.param .u64 __cudaparmf1__Z6Read2DI7ushort4ET_PKS1_iii, .param .s32 __cudaparmf2__Z6Read2DI7ushort4ET_PKS1_iii, .param .s32 __cudaparmf3__Z6Read2DI7ushort4ET_PKS1_iii, .param .s32 __cudaparmf4__Z6Read2DI7ushort4ET_PKS1_iii)
	{
	.reg .u32 %r<14>;
	.reg .u64 %rd<7>;
	.loc	20	114	0
$LDWbegin__Z6Read2DI7ushort4ET_PKS1_iii:
	ld.param.u64 	%rd1, [__cudaparmf1__Z6Read2DI7ushort4ET_PKS1_iii];
	mov.s64 	%rd2, %rd1;
	ld.param.u32 	%r1, [__cudaparmf2__Z6Read2DI7ushort4ET_PKS1_iii];
	mov.s32 	%r2, %r1;
	ld.param.u32 	%r3, [__cudaparmf3__Z6Read2DI7ushort4ET_PKS1_iii];
	mov.s32 	%r4, %r3;
	ld.param.u32 	%r5, [__cudaparmf4__Z6Read2DI7ushort4ET_PKS1_iii];
	mov.s32 	%r6, %r5;
	.loc	20	115	0
	mul.lo.s32 	%r7, %r2, %r6;
	add.s32 	%r8, %r4, %r7;
	cvt.s64.s32 	%rd3, %r8;
	mul.wide.s32 	%rd4, %r8, 8;
	add.u64 	%rd5, %rd2, %rd4;
	ld.v4.u16 	{%r9,%r10,%r11,%r12}, [%rd5+0];
	st.param.u16 	[__cudaretf__Z6Read2DI7ushort4ET_PKS1_iii+0], %r9;
	st.param.u16 	[__cudaretf__Z6Read2DI7ushort4ET_PKS1_iii+2], %r10;
	st.param.u16 	[__cudaretf__Z6Read2DI7ushort4ET_PKS1_iii+4], %r11;
	st.param.u16 	[__cudaretf__Z6Read2DI7ushort4ET_PKS1_iii+6], %r12;
	ret;
$LDWend__Z6Read2DI7ushort4ET_PKS1_iii:
	} // _Z6Read2DI7ushort4ET_PKS1_iii

	.visible .func (.param .align 16 .b8 __cudaretf__Z6Read2DI6float4ET_PKS1_iii[16]) _Z6Read2DI6float4ET_PKS1_iii (.param .u64 __cudaparmf1__Z6Read2DI6float4ET_PKS1_iii, .param .s32 __cudaparmf2__Z6Read2DI6float4ET_PKS1_iii, .param .s32 __cudaparmf3__Z6Read2DI6float4ET_PKS1_iii, .param .s32 __cudaparmf4__Z6Read2DI6float4ET_PKS1_iii)
	{
	.reg .u32 %r<10>;
	.reg .u64 %rd<7>;
	.reg .f32 %f<6>;
	.loc	20	114	0
$LDWbegin__Z6Read2DI6float4ET_PKS1_iii:
	ld.param.u64 	%rd1, [__cudaparmf1__Z6Read2DI6float4ET_PKS1_iii];
	mov.s64 	%rd2, %rd1;
	ld.param.u32 	%r1, [__cudaparmf2__Z6Read2DI6float4ET_PKS1_iii];
	mov.s32 	%r2, %r1;
	ld.param.u32 	%r3, [__cudaparmf3__Z6Read2DI6float4ET_PKS1_iii];
	mov.s32 	%r4, %r3;
	ld.param.u32 	%r5, [__cudaparmf4__Z6Read2DI6float4ET_PKS1_iii];
	mov.s32 	%r6, %r5;
	.loc	20	115	0
	mul.lo.s32 	%r7, %r2, %r6;
	add.s32 	%r8, %r4, %r7;
	cvt.s64.s32 	%rd3, %r8;
	mul.wide.s32 	%rd4, %r8, 16;
	add.u64 	%rd5, %rd2, %rd4;
	ld.v4.f32 	{%f1,%f2,%f3,%f4}, [%rd5+0];
	st.param.f32 	[__cudaretf__Z6Read2DI6float4ET_PKS1_iii+0], %f1;
	st.param.f32 	[__cudaretf__Z6Read2DI6float4ET_PKS1_iii+4], %f2;
	st.param.f32 	[__cudaretf__Z6Read2DI6float4ET_PKS1_iii+8], %f3;
	st.param.f32 	[__cudaretf__Z6Read2DI6float4ET_PKS1_iii+12], %f4;
	ret;
$LDWend__Z6Read2DI6float4ET_PKS1_iii:
	} // _Z6Read2DI6float4ET_PKS1_iii

	.visible .func _Z7Write2DI7ushort4EvT_PS1_iii (.param .align 8 .b8 __cudaparmf1__Z7Write2DI7ushort4EvT_PS1_iii[8], .param .u64 __cudaparmf2__Z7Write2DI7ushort4EvT_PS1_iii, .param .s32 __cudaparmf3__Z7Write2DI7ushort4EvT_PS1_iii, .param .s32 __cudaparmf4__Z7Write2DI7ushort4EvT_PS1_iii, .param .s32 __cudaparmf5__Z7Write2DI7ushort4EvT_PS1_iii)
	{
	.reg .u32 %r<18>;
	.reg .u64 %rd<7>;
	.loc	20	125	0
$LDWbegin__Z7Write2DI7ushort4EvT_PS1_iii:
	ld.param.u16 	%r1, [__cudaparmf1__Z7Write2DI7ushort4EvT_PS1_iii+0];
	mov.s32 	%r2, %r1;
	ld.param.u16 	%r3, [__cudaparmf1__Z7Write2DI7ushort4EvT_PS1_iii+2];
	mov.s32 	%r4, %r3;
	ld.param.u16 	%r5, [__cudaparmf1__Z7Write2DI7ushort4EvT_PS1_iii+4];
	mov.s32 	%r6, %r5;
	ld.param.u16 	%r7, [__cudaparmf1__Z7Write2DI7ushort4EvT_PS1_iii+6];
	mov.s32 	%r8, %r7;
	ld.param.u64 	%rd1, [__cudaparmf2__Z7Write2DI7ushort4EvT_PS1_iii];
	mov.s64 	%rd2, %rd1;
	ld.param.u32 	%r9, [__cudaparmf3__Z7Write2DI7ushort4EvT_PS1_iii];
	mov.s32 	%r10, %r9;
	ld.param.u32 	%r11, [__cudaparmf4__Z7Write2DI7ushort4EvT_PS1_iii];
	mov.s32 	%r12, %r11;
	ld.param.u32 	%r13, [__cudaparmf5__Z7Write2DI7ushort4EvT_PS1_iii];
	mov.s32 	%r14, %r13;
	.loc	20	126	0
	mul.lo.s32 	%r15, %r10, %r14;
	add.s32 	%r16, %r12, %r15;
	cvt.s64.s32 	%rd3, %r16;
	mul.wide.s32 	%rd4, %r16, 8;
	add.u64 	%rd5, %rd2, %rd4;
	st.v4.u16 	[%rd5+0], {%r2,%r4,%r6,%r8};
	.loc	20	127	0
	ret;
$LDWend__Z7Write2DI7ushort4EvT_PS1_iii:
	} // _Z7Write2DI7ushort4EvT_PS1_iii

	.visible .func _Z7Write2DI6float4EvT_PS1_iii (.param .align 16 .b8 __cudaparmf1__Z7Write2DI6float4EvT_PS1_iii[16], .param .u64 __cudaparmf2__Z7Write2DI6float4EvT_PS1_iii, .param .s32 __cudaparmf3__Z7Write2DI6float4EvT_PS1_iii, .param .s32 __cudaparmf4__Z7Write2DI6float4EvT_PS1_iii, .param .s32 __cudaparmf5__Z7Write2DI6float4EvT_PS1_iii)
	{
	.reg .u32 %r<10>;
	.reg .u64 %rd<7>;
	.reg .f32 %f<10>;
	.loc	20	125	0
$LDWbegin__Z7Write2DI6float4EvT_PS1_iii:
	ld.param.f32 	%f1, [__cudaparmf1__Z7Write2DI6float4EvT_PS1_iii+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z7Write2DI6float4EvT_PS1_iii+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z7Write2DI6float4EvT_PS1_iii+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z7Write2DI6float4EvT_PS1_iii+12];
	mov.f32 	%f8, %f7;
	ld.param.u64 	%rd1, [__cudaparmf2__Z7Write2DI6float4EvT_PS1_iii];
	mov.s64 	%rd2, %rd1;
	ld.param.u32 	%r1, [__cudaparmf3__Z7Write2DI6float4EvT_PS1_iii];
	mov.s32 	%r2, %r1;
	ld.param.u32 	%r3, [__cudaparmf4__Z7Write2DI6float4EvT_PS1_iii];
	mov.s32 	%r4, %r3;
	ld.param.u32 	%r5, [__cudaparmf5__Z7Write2DI6float4EvT_PS1_iii];
	mov.s32 	%r6, %r5;
	.loc	20	126	0
	mul.lo.s32 	%r7, %r2, %r6;
	add.s32 	%r8, %r4, %r7;
	cvt.s64.s32 	%rd3, %r8;
	mul.wide.s32 	%rd4, %r8, 16;
	add.u64 	%rd5, %rd2, %rd4;
	st.v4.f32 	[%rd5+0], {%f2,%f4,%f6,%f8};
	.loc	20	127	0
	ret;
$LDWend__Z7Write2DI6float4EvT_PS1_iii:
	} // _Z7Write2DI6float4EvT_PS1_iii

	.visible .func (.param .align 16 .b8 __cudaretf__Z18UnpremultiplyPixel8PixelRGB[16]) _Z18UnpremultiplyPixel8PixelRGB (.param .align 16 .b8 __cudaparmf1__Z18UnpremultiplyPixel8PixelRGB[16])
	{
	.reg .f32 %f<23>;
	.reg .pred %p<3>;
	.loc	4	206	0
$LDWbegin__Z18UnpremultiplyPixel8PixelRGB:
	ld.param.f32 	%f1, [__cudaparmf1__Z18UnpremultiplyPixel8PixelRGB+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z18UnpremultiplyPixel8PixelRGB+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z18UnpremultiplyPixel8PixelRGB+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z18UnpremultiplyPixel8PixelRGB+12];
	mov.f32 	%f8, %f7;
	.loc	4	208	0
	cvt.ftz.sat.f32.f32 	%f9, %f8;
	mov.f32 	%f10, %f9;
	mov.f32 	%f11, 0fb70637bd;    	// -8e-006
	add.ftz.f32 	%f12, %f9, %f11;
	mov.f32 	%f13, 0f00000000;    	// 0
	setp.le.ftz.f32 	%p1, %f12, %f13;
	@%p1 bra 	$Lt_14_1282;
	.loc	4	213	0
	rcp.approx.ftz.f32 	%f14, %f9;
	mul.ftz.f32 	%f15, %f14, %f6;
	.loc	4	214	0
	mul.ftz.f32 	%f16, %f14, %f4;
	.loc	4	215	0
	mul.ftz.f32 	%f17, %f14, %f2;
	bra.uni 	$Lt_14_1026;
$Lt_14_1282:
	.loc	4	219	0
	mov.f32 	%f15, 0f00000000;    	// 0
	mov.f32 	%f16, 0f00000000;    	// 0
	mov.f32 	%f17, 0f00000000;    	// 0
	mov.f32 	%f10, 0f00000000;    	// 0
$Lt_14_1026:
	.loc	4	224	0
	mov.f32 	%f18, %f17;
	st.param.f32 	[__cudaretf__Z18UnpremultiplyPixel8PixelRGB+0], %f18;
	mov.f32 	%f19, %f16;
	st.param.f32 	[__cudaretf__Z18UnpremultiplyPixel8PixelRGB+4], %f19;
	mov.f32 	%f20, %f15;
	st.param.f32 	[__cudaretf__Z18UnpremultiplyPixel8PixelRGB+8], %f20;
	mov.f32 	%f21, %f10;
	st.param.f32 	[__cudaretf__Z18UnpremultiplyPixel8PixelRGB+12], %f21;
	ret;
$LDWend__Z18UnpremultiplyPixel8PixelRGB:
	} // _Z18UnpremultiplyPixel8PixelRGB

	.visible .func (.param .f32 __cudaretf__Z13ToLinearColorf) _Z13ToLinearColorf (.param .f32 __cudaparmf1__Z13ToLinearColorf)
	{
	.reg .f32 %f<15>;
	.reg .pred %p<3>;
	.loc	4	231	0
$LDWbegin__Z13ToLinearColorf:
	ld.param.f32 	%f1, [__cudaparmf1__Z13ToLinearColorf];
	mov.f32 	%f2, %f1;
	mov.f32 	%f3, 0f00000000;     	// 0
	setp.lt.ftz.f32 	%p1, %f2, %f3;
	@!%p1 bra 	$Lt_15_1026;
	.loc	4	234	0
	neg.ftz.f32 	%f4, %f2;
	lg2.approx.ftz.f32 	%f5, %f4;
	mov.f32 	%f6, 0f400ccccd;     	// 2.2
	mul.ftz.f32 	%f7, %f5, %f6;
	ex2.approx.ftz.f32 	%f8, %f7;
	neg.ftz.f32 	%f9, %f8;
	bra.uni 	$LBB4__Z13ToLinearColorf;
$Lt_15_1026:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f10, %f2;
	mov.f32 	%f11, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f12, %f10, %f11;
	ex2.approx.ftz.f32 	%f9, %f12;
$LBB4__Z13ToLinearColorf:
	mov.f32 	%f13, %f9;
	st.param.f32 	[__cudaretf__Z13ToLinearColorf], %f13;
	ret;
$LDWend__Z13ToLinearColorf:
	} // _Z13ToLinearColorf

	.visible .func (.param .f32 __cudaretf__Z15FromLinearColorf) _Z15FromLinearColorf (.param .f32 __cudaparmf1__Z15FromLinearColorf)
	{
	.reg .f32 %f<15>;
	.reg .pred %p<3>;
	.loc	4	239	0
$LDWbegin__Z15FromLinearColorf:
	ld.param.f32 	%f1, [__cudaparmf1__Z15FromLinearColorf];
	mov.f32 	%f2, %f1;
	mov.f32 	%f3, 0f00000000;     	// 0
	setp.lt.ftz.f32 	%p1, %f2, %f3;
	@!%p1 bra 	$Lt_16_1026;
	.loc	4	242	0
	neg.ftz.f32 	%f4, %f2;
	lg2.approx.ftz.f32 	%f5, %f4;
	mov.f32 	%f6, 0f3ee8ba2e;     	// 0.454545
	mul.ftz.f32 	%f7, %f5, %f6;
	ex2.approx.ftz.f32 	%f8, %f7;
	neg.ftz.f32 	%f9, %f8;
	bra.uni 	$LBB4__Z15FromLinearColorf;
$Lt_16_1026:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f10, %f2;
	mov.f32 	%f11, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f12, %f10, %f11;
	ex2.approx.ftz.f32 	%f9, %f12;
$LBB4__Z15FromLinearColorf:
	mov.f32 	%f13, %f9;
	st.param.f32 	[__cudaretf__Z15FromLinearColorf], %f13;
	ret;
$LDWend__Z15FromLinearColorf:
	} // _Z15FromLinearColorf

	.visible .func (.param .align 16 .b8 __cudaretf__Z25PremultiplyLinearizePixel8PixelRGB[16]) _Z25PremultiplyLinearizePixel8PixelRGB (.param .align 16 .b8 __cudaparmf1__Z25PremultiplyLinearizePixel8PixelRGB[16])
	{
	.reg .f32 %f<47>;
	.reg .pred %p<5>;
	.loc	4	252	0
$LDWbegin__Z25PremultiplyLinearizePixel8PixelRGB:
	ld.param.f32 	%f1, [__cudaparmf1__Z25PremultiplyLinearizePixel8PixelRGB+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z25PremultiplyLinearizePixel8PixelRGB+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z25PremultiplyLinearizePixel8PixelRGB+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z25PremultiplyLinearizePixel8PixelRGB+12];
	mov.f32 	%f8, %f7;
	.loc	4	254	0
	cvt.ftz.sat.f32.f32 	%f9, %f8;
	.loc	4	255	0
	mov.f32 	%f10, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p1, %f2, %f10;
	@!%p1 bra 	$Lt_17_4098;
	.loc	4	234	0
	neg.ftz.f32 	%f11, %f2;
	lg2.approx.ftz.f32 	%f12, %f11;
	mov.f32 	%f13, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f14, %f12, %f13;
	ex2.approx.ftz.f32 	%f15, %f14;
	neg.ftz.f32 	%f16, %f15;
	bra.uni 	$LDWendi___log2f_194_5;
$Lt_17_4098:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f17, %f2;
	mov.f32 	%f18, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f19, %f17, %f18;
	ex2.approx.ftz.f32 	%f16, %f19;
$LDWendi___log2f_194_5:
	.loc	4	256	0
	mov.f32 	%f20, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p2, %f4, %f20;
	@!%p2 bra 	$Lt_17_4610;
	.loc	4	234	0
	neg.ftz.f32 	%f21, %f4;
	lg2.approx.ftz.f32 	%f22, %f21;
	mov.f32 	%f23, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f24, %f22, %f23;
	ex2.approx.ftz.f32 	%f25, %f24;
	neg.ftz.f32 	%f26, %f25;
	bra.uni 	$LDWendi___log2f_194_3;
$Lt_17_4610:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f27, %f4;
	mov.f32 	%f28, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f29, %f27, %f28;
	ex2.approx.ftz.f32 	%f26, %f29;
$LDWendi___log2f_194_3:
	.loc	4	257	0
	mov.f32 	%f30, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p3, %f6, %f30;
	@!%p3 bra 	$Lt_17_5122;
	.loc	4	234	0
	neg.ftz.f32 	%f31, %f6;
	lg2.approx.ftz.f32 	%f32, %f31;
	mov.f32 	%f33, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f34, %f32, %f33;
	ex2.approx.ftz.f32 	%f35, %f34;
	neg.ftz.f32 	%f36, %f35;
	bra.uni 	$LDWendi___log2f_194_1;
$Lt_17_5122:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f37, %f6;
	mov.f32 	%f38, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f39, %f37, %f38;
	ex2.approx.ftz.f32 	%f36, %f39;
$LDWendi___log2f_194_1:
	.loc	4	259	0
	mul.ftz.f32 	%f40, %f36, %f9;
	mul.ftz.f32 	%f41, %f26, %f9;
	mul.ftz.f32 	%f42, %f16, %f9;
	st.param.f32 	[__cudaretf__Z25PremultiplyLinearizePixel8PixelRGB+0], %f42;
	mov.f32 	%f43, %f41;
	st.param.f32 	[__cudaretf__Z25PremultiplyLinearizePixel8PixelRGB+4], %f43;
	mov.f32 	%f44, %f40;
	st.param.f32 	[__cudaretf__Z25PremultiplyLinearizePixel8PixelRGB+8], %f44;
	mov.f32 	%f45, %f9;
	st.param.f32 	[__cudaretf__Z25PremultiplyLinearizePixel8PixelRGB+12], %f45;
	ret;
$LDWend__Z25PremultiplyLinearizePixel8PixelRGB:
	} // _Z25PremultiplyLinearizePixel8PixelRGB

	.visible .func (.param .align 16 .b8 __cudaretf__Z29UnpremultiplyUnlinearizePixel8PixelRGB[16]) _Z29UnpremultiplyUnlinearizePixel8PixelRGB (.param .align 16 .b8 __cudaparmf1__Z29UnpremultiplyUnlinearizePixel8PixelRGB[16])
	{
	.reg .f32 %f<53>;
	.reg .pred %p<6>;
	.loc	4	263	0
$LDWbegin__Z29UnpremultiplyUnlinearizePixel8PixelRGB:
	ld.param.f32 	%f1, [__cudaparmf1__Z29UnpremultiplyUnlinearizePixel8PixelRGB+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z29UnpremultiplyUnlinearizePixel8PixelRGB+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z29UnpremultiplyUnlinearizePixel8PixelRGB+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z29UnpremultiplyUnlinearizePixel8PixelRGB+12];
	mov.f32 	%f8, %f7;
	.loc	4	208	0
	cvt.ftz.sat.f32.f32 	%f9, %f8;
	mov.f32 	%f10, %f9;
	mov.f32 	%f11, 0fb70637bd;    	// -8e-006
	add.ftz.f32 	%f12, %f9, %f11;
	mov.f32 	%f13, 0f00000000;    	// 0
	setp.le.ftz.f32 	%p1, %f12, %f13;
	@%p1 bra 	$Lt_18_5122;
	.loc	4	213	0
	rcp.approx.ftz.f32 	%f14, %f9;
	mul.ftz.f32 	%f15, %f14, %f6;
	.loc	4	214	0
	mul.ftz.f32 	%f16, %f14, %f4;
	.loc	4	215	0
	mul.ftz.f32 	%f17, %f14, %f2;
	bra.uni 	$Lt_18_4866;
$Lt_18_5122:
	.loc	4	219	0
	mov.f32 	%f15, 0f00000000;    	// 0
	mov.f32 	%f16, 0f00000000;    	// 0
	mov.f32 	%f17, 0f00000000;    	// 0
	mov.f32 	%f10, 0f00000000;    	// 0
$Lt_18_4866:
	.loc	4	266	0
	mov.f32 	%f18, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p2, %f17, %f18;
	@!%p2 bra 	$Lt_18_5378;
	.loc	4	242	0
	neg.ftz.f32 	%f19, %f17;
	lg2.approx.ftz.f32 	%f20, %f19;
	mov.f32 	%f21, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f22, %f20, %f21;
	ex2.approx.ftz.f32 	%f23, %f22;
	neg.ftz.f32 	%f24, %f23;
	bra.uni 	$LDWendi___log2f_195_5;
$Lt_18_5378:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f25, %f17;
	mov.f32 	%f26, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f27, %f25, %f26;
	ex2.approx.ftz.f32 	%f24, %f27;
$LDWendi___log2f_195_5:
	.loc	4	267	0
	mov.f32 	%f28, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p3, %f16, %f28;
	@!%p3 bra 	$Lt_18_5890;
	.loc	4	242	0
	neg.ftz.f32 	%f29, %f16;
	lg2.approx.ftz.f32 	%f30, %f29;
	mov.f32 	%f31, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f32, %f30, %f31;
	ex2.approx.ftz.f32 	%f33, %f32;
	neg.ftz.f32 	%f34, %f33;
	bra.uni 	$LDWendi___log2f_195_3;
$Lt_18_5890:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f35, %f16;
	mov.f32 	%f36, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f37, %f35, %f36;
	ex2.approx.ftz.f32 	%f34, %f37;
$LDWendi___log2f_195_3:
	.loc	4	268	0
	mov.f32 	%f38, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p4, %f15, %f38;
	@!%p4 bra 	$Lt_18_6402;
	.loc	4	242	0
	neg.ftz.f32 	%f39, %f15;
	lg2.approx.ftz.f32 	%f40, %f39;
	mov.f32 	%f41, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f42, %f40, %f41;
	ex2.approx.ftz.f32 	%f43, %f42;
	neg.ftz.f32 	%f44, %f43;
	bra.uni 	$LDWendi___log2f_195_1;
$Lt_18_6402:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f45, %f15;
	mov.f32 	%f46, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f47, %f45, %f46;
	ex2.approx.ftz.f32 	%f44, %f47;
$LDWendi___log2f_195_1:
	.loc	4	269	0
	mov.f32 	%f48, %f24;
	st.param.f32 	[__cudaretf__Z29UnpremultiplyUnlinearizePixel8PixelRGB+0], %f48;
	mov.f32 	%f49, %f34;
	st.param.f32 	[__cudaretf__Z29UnpremultiplyUnlinearizePixel8PixelRGB+4], %f49;
	mov.f32 	%f50, %f44;
	st.param.f32 	[__cudaretf__Z29UnpremultiplyUnlinearizePixel8PixelRGB+8], %f50;
	mov.f32 	%f51, %f10;
	st.param.f32 	[__cudaretf__Z29UnpremultiplyUnlinearizePixel8PixelRGB+12], %f51;
	ret;
$LDWend__Z29UnpremultiplyUnlinearizePixel8PixelRGB:
	} // _Z29UnpremultiplyUnlinearizePixel8PixelRGB

	.visible .func (.param .align 16 .b8 __cudaretf__Z20PremultiplyLinearize6float4[16]) _Z20PremultiplyLinearize6float4 (.param .align 16 .b8 __cudaparmf1__Z20PremultiplyLinearize6float4[16])
	{
	.reg .f32 %f<47>;
	.reg .pred %p<5>;
	.loc	4	277	0
$LDWbegin__Z20PremultiplyLinearize6float4:
	ld.param.f32 	%f1, [__cudaparmf1__Z20PremultiplyLinearize6float4+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z20PremultiplyLinearize6float4+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z20PremultiplyLinearize6float4+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z20PremultiplyLinearize6float4+12];
	mov.f32 	%f8, %f7;
	.loc	4	254	0
	cvt.ftz.sat.f32.f32 	%f9, %f8;
	.loc	4	255	0
	mov.f32 	%f10, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p1, %f2, %f10;
	@!%p1 bra 	$Lt_19_4098;
	.loc	4	234	0
	neg.ftz.f32 	%f11, %f2;
	lg2.approx.ftz.f32 	%f12, %f11;
	mov.f32 	%f13, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f14, %f12, %f13;
	ex2.approx.ftz.f32 	%f15, %f14;
	neg.ftz.f32 	%f16, %f15;
	bra.uni 	$LDWendi___log2f_196_5;
$Lt_19_4098:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f17, %f2;
	mov.f32 	%f18, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f19, %f17, %f18;
	ex2.approx.ftz.f32 	%f16, %f19;
$LDWendi___log2f_196_5:
	.loc	4	256	0
	mov.f32 	%f20, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p2, %f4, %f20;
	@!%p2 bra 	$Lt_19_4610;
	.loc	4	234	0
	neg.ftz.f32 	%f21, %f4;
	lg2.approx.ftz.f32 	%f22, %f21;
	mov.f32 	%f23, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f24, %f22, %f23;
	ex2.approx.ftz.f32 	%f25, %f24;
	neg.ftz.f32 	%f26, %f25;
	bra.uni 	$LDWendi___log2f_196_3;
$Lt_19_4610:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f27, %f4;
	mov.f32 	%f28, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f29, %f27, %f28;
	ex2.approx.ftz.f32 	%f26, %f29;
$LDWendi___log2f_196_3:
	.loc	4	257	0
	mov.f32 	%f30, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p3, %f6, %f30;
	@!%p3 bra 	$Lt_19_5122;
	.loc	4	234	0
	neg.ftz.f32 	%f31, %f6;
	lg2.approx.ftz.f32 	%f32, %f31;
	mov.f32 	%f33, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f34, %f32, %f33;
	ex2.approx.ftz.f32 	%f35, %f34;
	neg.ftz.f32 	%f36, %f35;
	bra.uni 	$LDWendi___log2f_196_1;
$Lt_19_5122:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f37, %f6;
	mov.f32 	%f38, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f39, %f37, %f38;
	ex2.approx.ftz.f32 	%f36, %f39;
$LDWendi___log2f_196_1:
	.loc	4	259	0
	mul.ftz.f32 	%f40, %f36, %f9;
	mul.ftz.f32 	%f41, %f26, %f9;
	.loc	4	278	0
	mul.ftz.f32 	%f42, %f16, %f9;
	st.param.f32 	[__cudaretf__Z20PremultiplyLinearize6float4+0], %f42;
	mov.f32 	%f43, %f41;
	st.param.f32 	[__cudaretf__Z20PremultiplyLinearize6float4+4], %f43;
	mov.f32 	%f44, %f40;
	st.param.f32 	[__cudaretf__Z20PremultiplyLinearize6float4+8], %f44;
	mov.f32 	%f45, %f9;
	st.param.f32 	[__cudaretf__Z20PremultiplyLinearize6float4+12], %f45;
	ret;
$LDWend__Z20PremultiplyLinearize6float4:
	} // _Z20PremultiplyLinearize6float4

	.visible .func (.param .align 16 .b8 __cudaretf__Z24UnpremultiplyUnlinearize6float4[16]) _Z24UnpremultiplyUnlinearize6float4 (.param .align 16 .b8 __cudaparmf1__Z24UnpremultiplyUnlinearize6float4[16])
	{
	.reg .f32 %f<53>;
	.reg .pred %p<6>;
	.loc	4	284	0
$LDWbegin__Z24UnpremultiplyUnlinearize6float4:
	ld.param.f32 	%f1, [__cudaparmf1__Z24UnpremultiplyUnlinearize6float4+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z24UnpremultiplyUnlinearize6float4+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z24UnpremultiplyUnlinearize6float4+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z24UnpremultiplyUnlinearize6float4+12];
	mov.f32 	%f8, %f7;
	.loc	4	208	0
	cvt.ftz.sat.f32.f32 	%f9, %f8;
	mov.f32 	%f10, %f9;
	mov.f32 	%f11, 0fb70637bd;    	// -8e-006
	add.ftz.f32 	%f12, %f9, %f11;
	mov.f32 	%f13, 0f00000000;    	// 0
	setp.le.ftz.f32 	%p1, %f12, %f13;
	@%p1 bra 	$Lt_20_5122;
	.loc	4	213	0
	rcp.approx.ftz.f32 	%f14, %f9;
	mul.ftz.f32 	%f15, %f14, %f6;
	.loc	4	214	0
	mul.ftz.f32 	%f16, %f14, %f4;
	.loc	4	215	0
	mul.ftz.f32 	%f17, %f14, %f2;
	bra.uni 	$Lt_20_4866;
$Lt_20_5122:
	.loc	4	219	0
	mov.f32 	%f15, 0f00000000;    	// 0
	mov.f32 	%f16, 0f00000000;    	// 0
	mov.f32 	%f17, 0f00000000;    	// 0
	mov.f32 	%f10, 0f00000000;    	// 0
$Lt_20_4866:
	.loc	4	266	0
	mov.f32 	%f18, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p2, %f17, %f18;
	@!%p2 bra 	$Lt_20_5378;
	.loc	4	242	0
	neg.ftz.f32 	%f19, %f17;
	lg2.approx.ftz.f32 	%f20, %f19;
	mov.f32 	%f21, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f22, %f20, %f21;
	ex2.approx.ftz.f32 	%f23, %f22;
	neg.ftz.f32 	%f24, %f23;
	bra.uni 	$LDWendi___log2f_197_5;
$Lt_20_5378:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f25, %f17;
	mov.f32 	%f26, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f27, %f25, %f26;
	ex2.approx.ftz.f32 	%f24, %f27;
$LDWendi___log2f_197_5:
	.loc	4	267	0
	mov.f32 	%f28, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p3, %f16, %f28;
	@!%p3 bra 	$Lt_20_5890;
	.loc	4	242	0
	neg.ftz.f32 	%f29, %f16;
	lg2.approx.ftz.f32 	%f30, %f29;
	mov.f32 	%f31, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f32, %f30, %f31;
	ex2.approx.ftz.f32 	%f33, %f32;
	neg.ftz.f32 	%f34, %f33;
	bra.uni 	$LDWendi___log2f_197_3;
$Lt_20_5890:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f35, %f16;
	mov.f32 	%f36, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f37, %f35, %f36;
	ex2.approx.ftz.f32 	%f34, %f37;
$LDWendi___log2f_197_3:
	.loc	4	268	0
	mov.f32 	%f38, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p4, %f15, %f38;
	@!%p4 bra 	$Lt_20_6402;
	.loc	4	242	0
	neg.ftz.f32 	%f39, %f15;
	lg2.approx.ftz.f32 	%f40, %f39;
	mov.f32 	%f41, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f42, %f40, %f41;
	ex2.approx.ftz.f32 	%f43, %f42;
	neg.ftz.f32 	%f44, %f43;
	bra.uni 	$LDWendi___log2f_197_1;
$Lt_20_6402:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f45, %f15;
	mov.f32 	%f46, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f47, %f45, %f46;
	ex2.approx.ftz.f32 	%f44, %f47;
$LDWendi___log2f_197_1:
	.loc	4	285	0
	mov.f32 	%f48, %f24;
	st.param.f32 	[__cudaretf__Z24UnpremultiplyUnlinearize6float4+0], %f48;
	mov.f32 	%f49, %f34;
	st.param.f32 	[__cudaretf__Z24UnpremultiplyUnlinearize6float4+4], %f49;
	mov.f32 	%f50, %f44;
	st.param.f32 	[__cudaretf__Z24UnpremultiplyUnlinearize6float4+8], %f50;
	mov.f32 	%f51, %f10;
	st.param.f32 	[__cudaretf__Z24UnpremultiplyUnlinearize6float4+12], %f51;
	ret;
$LDWend__Z24UnpremultiplyUnlinearize6float4:
	} // _Z24UnpremultiplyUnlinearize6float4

	.visible .func (.param .align 16 .b8 __cudaretf__Z18SwapComponentOrderI6float4ET_RKS1_[16]) _Z18SwapComponentOrderI6float4ET_RKS1_ (.param .u64 __cudaparmf1__Z18SwapComponentOrderI6float4ET_RKS1_)
	{
	.reg .u64 %rd<4>;
	.reg .f32 %f<6>;
	.loc	21	264	0
$LDWbegin__Z18SwapComponentOrderI6float4ET_RKS1_:
	ld.param.u64 	%rd1, [__cudaparmf1__Z18SwapComponentOrderI6float4ET_RKS1_];
	mov.s64 	%rd2, %rd1;
	.loc	21	270	0
	ld.f32 	%f1, [%rd2+12];
	st.param.f32 	[__cudaretf__Z18SwapComponentOrderI6float4ET_RKS1_+0], %f1;
	ld.f32 	%f2, [%rd2+8];
	st.param.f32 	[__cudaretf__Z18SwapComponentOrderI6float4ET_RKS1_+4], %f2;
	ld.f32 	%f3, [%rd2+4];
	st.param.f32 	[__cudaretf__Z18SwapComponentOrderI6float4ET_RKS1_+8], %f3;
	ld.f32 	%f4, [%rd2+0];
	st.param.f32 	[__cudaretf__Z18SwapComponentOrderI6float4ET_RKS1_+12], %f4;
	ret;
$LDWend__Z18SwapComponentOrderI6float4ET_RKS1_:
	} // _Z18SwapComponentOrderI6float4ET_RKS1_
	.global .align 8 .b8 sSpline[800];

	.visible .func (.param .align 16 .b8 __cudaretf__Z13CalculateBGRA8PixelRGBf[16]) _Z13CalculateBGRA8PixelRGBf (.param .align 16 .b8 __cudaparmf1__Z13CalculateBGRA8PixelRGBf[16], .param .f32 __cudaparmf2__Z13CalculateBGRA8PixelRGBf)
	{
	.reg .u64 %rd<76>;
	.reg .f32 %f<217>;
	.reg .pred %p<39>;
	.loc	22	39	0
$LDWbegin__Z13CalculateBGRA8PixelRGBf:
	ld.param.f32 	%f1, [__cudaparmf1__Z13CalculateBGRA8PixelRGBf+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z13CalculateBGRA8PixelRGBf+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z13CalculateBGRA8PixelRGBf+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z13CalculateBGRA8PixelRGBf+12];
	mov.f32 	%f8, %f7;
	ld.param.f32 	%f9, [__cudaparmf2__Z13CalculateBGRA8PixelRGBf];
	mov.f32 	%f10, %f9;
	.loc	22	40	0
	mov.f32 	%f11, %f2;
	mov.f32 	%f12, %f4;
	mov.f32 	%f13, %f6;
	.loc	22	41	0
	mov.f32 	%f14, 0f3a83126f;    	// 0.001
	setp.gt.ftz.f32 	%p1, %f10, %f14;
	@!%p1 bra 	$Lt_22_25602;
	mov.f32 	%f15, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p2, %f6, %f15;
	@!%p2 bra 	$Lt_22_26370;
	.loc	22	46	0
	ld.global.f32 	%f16, [sSpline+264];
	add.ftz.f32 	%f17, %f16, %f6;
	mov.u64 	%rd1, sSpline;
	bra.uni 	$Lt_22_26626;
$Lt_22_26370:
	ld.global.u64 	%rd2, [sSpline+392];
	mov.f32 	%f18, 0f3f800000;    	// 1
	setp.gt.ftz.f32 	%p3, %f6, %f18;
	@!%p3 bra 	$Lt_22_26882;
	mov.u64 	%rd1, sSpline;
	mov.f32 	%f19, 0f3f800000;    	// 1
	mul.lo.u64 	%rd3, %rd2, 4;
	add.u64 	%rd4, %rd1, %rd3;
	ld.global.f32 	%f20, [%rd4+260];
	sub.ftz.f32 	%f21, %f19, %f20;
	sub.ftz.f32 	%f17, %f6, %f21;
	bra.uni 	$Lt_22_26626;
$Lt_22_26882:
	.loc	3	54	0
	mov.f32 	%f22, %f6;
	.loc	3	58	0
	sub.u64 	%rd5, %rd2, 1;
	mov.s64 	%rd6, %rd5;
	mov.u64 	%rd7, 1;
	setp.le.u64 	%p4, %rd5, %rd7;
	@%p4 bra 	$Lt_22_44546;
	mov.u64 	%rd8, 0;
	mov.u64 	%rd1, sSpline;
$Lt_22_27650:
 //<loop> Loop body line 58
	add.u64 	%rd9, %rd6, %rd8;
	shr.u64 	%rd10, %rd9, 1;
	mul.lo.u64 	%rd11, %rd10, 4;
	add.u64 	%rd12, %rd1, %rd11;
	ld.global.f32 	%f23, [%rd12+200];
	setp.gt.ftz.f32 	%p5, %f23, %f6;
	@!%p5 bra 	$Lt_22_28162;
 //<loop> Part of loop body line 58, head labeled $Lt_22_27650
	.loc	3	67	0
	mov.s64 	%rd6, %rd10;
	bra.uni 	$Lt_22_27906;
$Lt_22_28162:
 //<loop> Part of loop body line 58, head labeled $Lt_22_27650
	.loc	3	71	0
	mov.s64 	%rd8, %rd10;
$Lt_22_27906:
 //<loop> Part of loop body line 58, head labeled $Lt_22_27650
	sub.u64 	%rd13, %rd6, %rd8;
	mov.u64 	%rd14, 1;
	setp.gt.u64 	%p6, %rd13, %rd14;
	@%p6 bra 	$Lt_22_27650;
	bra.uni 	$Lt_22_27138;
$Lt_22_44546:
	mov.u64 	%rd8, 0;
	mov.u64 	%rd1, sSpline;
$Lt_22_27138:
	mul.lo.u64 	%rd15, %rd6, 4;
	add.u64 	%rd16, %rd1, %rd15;
	mul.lo.u64 	%rd17, %rd8, 4;
	add.u64 	%rd18, %rd1, %rd17;
	ld.global.f32 	%f24, [%rd16+200];
	ld.global.f32 	%f25, [%rd18+200];
	sub.ftz.f32 	%f26, %f24, %f25;
	mov.f32 	%f27, 0f00000000;    	// 0
	setp.neu.ftz.f32 	%p7, %f26, %f27;
	@!%p7 bra 	$Lt_22_28674;
	.loc	27	529	0
	sub.ftz.f32 	%f28, %f24, %f6;
	sub.ftz.f32 	%f29, %f6, %f25;
	div.approx.ftz.f32 	%f30, %f28, %f26;
	div.approx.ftz.f32 	%f31, %f29, %f26;
	ld.global.f32 	%f32, [%rd16+328];
	mul.ftz.f32 	%f33, %f31, %f31;
	mul.ftz.f32 	%f34, %f31, %f33;
	sub.ftz.f32 	%f35, %f34, %f31;
	mul.ftz.f32 	%f36, %f32, %f35;
	ld.global.f32 	%f37, [%rd18+328];
	mul.ftz.f32 	%f38, %f30, %f30;
	mul.ftz.f32 	%f39, %f30, %f38;
	sub.ftz.f32 	%f40, %f39, %f30;
	fma.rn.ftz.f32 	%f41, %f37, %f40, %f36;
	mul.ftz.f32 	%f42, %f26, %f26;
	mul.ftz.f32 	%f43, %f41, %f42;
	mov.f32 	%f44, 0f40c00000;    	// 6
	div.approx.ftz.f32 	%f45, %f43, %f44;
	.loc	3	80	0
	ld.global.f32 	%f46, [%rd16+264];
	mul.ftz.f32 	%f47, %f46, %f31;
	ld.global.f32 	%f48, [%rd18+264];
	fma.rn.ftz.f32 	%f49, %f48, %f30, %f47;
	add.ftz.f32 	%f22, %f45, %f49;
$Lt_22_28674:
	.loc	22	46	0
	mov.f32 	%f17, %f22;
$Lt_22_26626:
$Lt_22_26114:
	mov.f32 	%f50, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p8, %f17, %f50;
	@!%p8 bra 	$Lt_22_29442;
	.loc	22	47	0
	ld.global.f32 	%f51, [sSpline+64];
	add.ftz.f32 	%f52, %f51, %f17;
	bra.uni 	$Lt_22_29698;
$Lt_22_29442:
	ld.global.u64 	%rd19, [sSpline+192];
	mov.f32 	%f53, 0f3f800000;    	// 1
	setp.gt.ftz.f32 	%p9, %f17, %f53;
	@!%p9 bra 	$Lt_22_29954;
	mov.f32 	%f54, 0f3f800000;    	// 1
	mul.lo.u64 	%rd20, %rd19, 4;
	add.u64 	%rd21, %rd1, %rd20;
	ld.global.f32 	%f55, [%rd21+60];
	sub.ftz.f32 	%f56, %f54, %f55;
	sub.ftz.f32 	%f52, %f17, %f56;
	bra.uni 	$Lt_22_29698;
$Lt_22_29954:
	.loc	3	54	0
	mov.f32 	%f22, %f17;
	.loc	3	58	0
	sub.u64 	%rd22, %rd19, 1;
	mov.s64 	%rd6, %rd22;
	mov.u64 	%rd23, 1;
	setp.le.u64 	%p10, %rd22, %rd23;
	@%p10 bra 	$Lt_22_44802;
	mov.u64 	%rd8, 0;
$Lt_22_30722:
 //<loop> Loop body line 58
	add.u64 	%rd24, %rd6, %rd8;
	shr.u64 	%rd10, %rd24, 1;
	mul.lo.u64 	%rd25, %rd10, 4;
	add.u64 	%rd26, %rd1, %rd25;
	ld.global.f32 	%f57, [%rd26+0];
	setp.gt.ftz.f32 	%p11, %f57, %f17;
	@!%p11 bra 	$Lt_22_31234;
 //<loop> Part of loop body line 58, head labeled $Lt_22_30722
	.loc	3	67	0
	mov.s64 	%rd6, %rd10;
	bra.uni 	$Lt_22_30978;
$Lt_22_31234:
 //<loop> Part of loop body line 58, head labeled $Lt_22_30722
	.loc	3	71	0
	mov.s64 	%rd8, %rd10;
$Lt_22_30978:
 //<loop> Part of loop body line 58, head labeled $Lt_22_30722
	sub.u64 	%rd27, %rd6, %rd8;
	mov.u64 	%rd28, 1;
	setp.gt.u64 	%p12, %rd27, %rd28;
	@%p12 bra 	$Lt_22_30722;
	bra.uni 	$Lt_22_30210;
$Lt_22_44802:
	mov.u64 	%rd8, 0;
$Lt_22_30210:
	mul.lo.u64 	%rd29, %rd6, 4;
	add.u64 	%rd16, %rd1, %rd29;
	mul.lo.u64 	%rd30, %rd8, 4;
	add.u64 	%rd18, %rd1, %rd30;
	ld.global.f32 	%f58, [%rd16+0];
	ld.global.f32 	%f59, [%rd18+0];
	sub.ftz.f32 	%f60, %f58, %f59;
	mov.f32 	%f61, 0f00000000;    	// 0
	setp.neu.ftz.f32 	%p13, %f60, %f61;
	@!%p13 bra 	$Lt_22_31746;
	.loc	27	529	0
	sub.ftz.f32 	%f62, %f58, %f17;
	sub.ftz.f32 	%f63, %f17, %f59;
	div.approx.ftz.f32 	%f64, %f62, %f60;
	div.approx.ftz.f32 	%f65, %f63, %f60;
	ld.global.f32 	%f66, [%rd16+128];
	mul.ftz.f32 	%f67, %f65, %f65;
	mul.ftz.f32 	%f68, %f65, %f67;
	sub.ftz.f32 	%f69, %f68, %f65;
	mul.ftz.f32 	%f70, %f66, %f69;
	ld.global.f32 	%f71, [%rd18+128];
	mul.ftz.f32 	%f72, %f64, %f64;
	mul.ftz.f32 	%f73, %f64, %f72;
	sub.ftz.f32 	%f74, %f73, %f64;
	fma.rn.ftz.f32 	%f75, %f71, %f74, %f70;
	mul.ftz.f32 	%f76, %f60, %f60;
	mul.ftz.f32 	%f77, %f75, %f76;
	mov.f32 	%f78, 0f40c00000;    	// 6
	div.approx.ftz.f32 	%f45, %f77, %f78;
	.loc	3	80	0
	ld.global.f32 	%f79, [%rd16+64];
	mul.ftz.f32 	%f80, %f79, %f65;
	ld.global.f32 	%f81, [%rd18+64];
	fma.rn.ftz.f32 	%f82, %f81, %f64, %f80;
	add.ftz.f32 	%f22, %f45, %f82;
$Lt_22_31746:
	.loc	22	47	0
	mov.f32 	%f52, %f22;
$Lt_22_29698:
$Lt_22_29186:
	.loc	22	48	0
	sub.ftz.f32 	%f83, %f52, %f6;
	fma.rn.ftz.f32 	%f13, %f10, %f83, %f6;
	mov.f32 	%f84, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p14, %f4, %f84;
	@!%p14 bra 	$Lt_22_32514;
	.loc	22	51	0
	ld.global.f32 	%f85, [sSpline+464];
	add.ftz.f32 	%f86, %f85, %f4;
	bra.uni 	$Lt_22_32770;
$Lt_22_32514:
	ld.global.u64 	%rd31, [sSpline+592];
	mov.f32 	%f87, 0f3f800000;    	// 1
	setp.gt.ftz.f32 	%p15, %f4, %f87;
	@!%p15 bra 	$Lt_22_33026;
	mov.f32 	%f88, 0f3f800000;    	// 1
	mul.lo.u64 	%rd32, %rd31, 4;
	add.u64 	%rd33, %rd1, %rd32;
	ld.global.f32 	%f89, [%rd33+460];
	sub.ftz.f32 	%f90, %f88, %f89;
	sub.ftz.f32 	%f86, %f4, %f90;
	bra.uni 	$Lt_22_32770;
$Lt_22_33026:
	.loc	3	54	0
	mov.f32 	%f22, %f4;
	.loc	3	58	0
	sub.u64 	%rd34, %rd31, 1;
	mov.s64 	%rd6, %rd34;
	mov.u64 	%rd35, 1;
	setp.le.u64 	%p16, %rd34, %rd35;
	@%p16 bra 	$Lt_22_45058;
	mov.u64 	%rd8, 0;
$Lt_22_33794:
 //<loop> Loop body line 58
	add.u64 	%rd36, %rd6, %rd8;
	shr.u64 	%rd10, %rd36, 1;
	mul.lo.u64 	%rd37, %rd10, 4;
	add.u64 	%rd38, %rd1, %rd37;
	ld.global.f32 	%f91, [%rd38+400];
	setp.gt.ftz.f32 	%p17, %f91, %f4;
	@!%p17 bra 	$Lt_22_34306;
 //<loop> Part of loop body line 58, head labeled $Lt_22_33794
	.loc	3	67	0
	mov.s64 	%rd6, %rd10;
	bra.uni 	$Lt_22_34050;
$Lt_22_34306:
 //<loop> Part of loop body line 58, head labeled $Lt_22_33794
	.loc	3	71	0
	mov.s64 	%rd8, %rd10;
$Lt_22_34050:
 //<loop> Part of loop body line 58, head labeled $Lt_22_33794
	sub.u64 	%rd39, %rd6, %rd8;
	mov.u64 	%rd40, 1;
	setp.gt.u64 	%p18, %rd39, %rd40;
	@%p18 bra 	$Lt_22_33794;
	bra.uni 	$Lt_22_33282;
$Lt_22_45058:
	mov.u64 	%rd8, 0;
$Lt_22_33282:
	mul.lo.u64 	%rd41, %rd6, 4;
	add.u64 	%rd16, %rd1, %rd41;
	mul.lo.u64 	%rd42, %rd8, 4;
	add.u64 	%rd18, %rd1, %rd42;
	ld.global.f32 	%f92, [%rd16+400];
	ld.global.f32 	%f93, [%rd18+400];
	sub.ftz.f32 	%f94, %f92, %f93;
	mov.f32 	%f95, 0f00000000;    	// 0
	setp.neu.ftz.f32 	%p19, %f94, %f95;
	@!%p19 bra 	$Lt_22_34818;
	.loc	27	529	0
	sub.ftz.f32 	%f96, %f92, %f4;
	sub.ftz.f32 	%f97, %f4, %f93;
	div.approx.ftz.f32 	%f98, %f96, %f94;
	div.approx.ftz.f32 	%f99, %f97, %f94;
	ld.global.f32 	%f100, [%rd16+528];
	mul.ftz.f32 	%f101, %f99, %f99;
	mul.ftz.f32 	%f102, %f99, %f101;
	sub.ftz.f32 	%f103, %f102, %f99;
	mul.ftz.f32 	%f104, %f100, %f103;
	ld.global.f32 	%f105, [%rd18+528];
	mul.ftz.f32 	%f106, %f98, %f98;
	mul.ftz.f32 	%f107, %f98, %f106;
	sub.ftz.f32 	%f108, %f107, %f98;
	fma.rn.ftz.f32 	%f109, %f105, %f108, %f104;
	mul.ftz.f32 	%f110, %f94, %f94;
	mul.ftz.f32 	%f111, %f109, %f110;
	mov.f32 	%f112, 0f40c00000;   	// 6
	div.approx.ftz.f32 	%f45, %f111, %f112;
	.loc	3	80	0
	ld.global.f32 	%f113, [%rd16+464];
	mul.ftz.f32 	%f114, %f113, %f99;
	ld.global.f32 	%f115, [%rd18+464];
	fma.rn.ftz.f32 	%f116, %f115, %f98, %f114;
	add.ftz.f32 	%f22, %f45, %f116;
$Lt_22_34818:
	.loc	22	51	0
	mov.f32 	%f86, %f22;
$Lt_22_32770:
$Lt_22_32258:
	mov.f32 	%f117, 0f00000000;   	// 0
	setp.lt.ftz.f32 	%p20, %f86, %f117;
	@!%p20 bra 	$Lt_22_35586;
	.loc	22	52	0
	ld.global.f32 	%f118, [sSpline+64];
	add.ftz.f32 	%f119, %f118, %f86;
	bra.uni 	$Lt_22_35842;
$Lt_22_35586:
	ld.global.u64 	%rd19, [sSpline+192];
	mov.f32 	%f120, 0f3f800000;   	// 1
	setp.gt.ftz.f32 	%p21, %f86, %f120;
	@!%p21 bra 	$Lt_22_36098;
	mov.f32 	%f121, 0f3f800000;   	// 1
	mul.lo.u64 	%rd43, %rd19, 4;
	add.u64 	%rd44, %rd1, %rd43;
	ld.global.f32 	%f122, [%rd44+60];
	sub.ftz.f32 	%f123, %f121, %f122;
	sub.ftz.f32 	%f119, %f86, %f123;
	bra.uni 	$Lt_22_35842;
$Lt_22_36098:
	.loc	3	54	0
	mov.f32 	%f22, %f86;
	.loc	3	58	0
	sub.u64 	%rd22, %rd19, 1;
	mov.s64 	%rd6, %rd22;
	mov.u64 	%rd45, 1;
	setp.le.u64 	%p22, %rd22, %rd45;
	@%p22 bra 	$Lt_22_45314;
	mov.u64 	%rd8, 0;
$Lt_22_36866:
 //<loop> Loop body line 58
	add.u64 	%rd46, %rd6, %rd8;
	shr.u64 	%rd10, %rd46, 1;
	mul.lo.u64 	%rd47, %rd10, 4;
	add.u64 	%rd48, %rd1, %rd47;
	ld.global.f32 	%f124, [%rd48+0];
	setp.gt.ftz.f32 	%p23, %f124, %f86;
	@!%p23 bra 	$Lt_22_37378;
 //<loop> Part of loop body line 58, head labeled $Lt_22_36866
	.loc	3	67	0
	mov.s64 	%rd6, %rd10;
	bra.uni 	$Lt_22_37122;
$Lt_22_37378:
 //<loop> Part of loop body line 58, head labeled $Lt_22_36866
	.loc	3	71	0
	mov.s64 	%rd8, %rd10;
$Lt_22_37122:
 //<loop> Part of loop body line 58, head labeled $Lt_22_36866
	sub.u64 	%rd49, %rd6, %rd8;
	mov.u64 	%rd50, 1;
	setp.gt.u64 	%p24, %rd49, %rd50;
	@%p24 bra 	$Lt_22_36866;
	bra.uni 	$Lt_22_36354;
$Lt_22_45314:
	mov.u64 	%rd8, 0;
$Lt_22_36354:
	mul.lo.u64 	%rd51, %rd6, 4;
	add.u64 	%rd16, %rd1, %rd51;
	mul.lo.u64 	%rd52, %rd8, 4;
	add.u64 	%rd18, %rd1, %rd52;
	ld.global.f32 	%f58, [%rd16+0];
	ld.global.f32 	%f59, [%rd18+0];
	sub.ftz.f32 	%f60, %f58, %f59;
	mov.f32 	%f125, 0f00000000;   	// 0
	setp.neu.ftz.f32 	%p25, %f60, %f125;
	@!%p25 bra 	$Lt_22_37890;
	.loc	27	529	0
	sub.ftz.f32 	%f126, %f58, %f86;
	sub.ftz.f32 	%f127, %f86, %f59;
	div.approx.ftz.f32 	%f128, %f126, %f60;
	div.approx.ftz.f32 	%f129, %f127, %f60;
	ld.global.f32 	%f130, [%rd16+128];
	mul.ftz.f32 	%f131, %f129, %f129;
	mul.ftz.f32 	%f132, %f129, %f131;
	sub.ftz.f32 	%f133, %f132, %f129;
	mul.ftz.f32 	%f134, %f130, %f133;
	ld.global.f32 	%f135, [%rd18+128];
	mul.ftz.f32 	%f136, %f128, %f128;
	mul.ftz.f32 	%f137, %f128, %f136;
	sub.ftz.f32 	%f138, %f137, %f128;
	fma.rn.ftz.f32 	%f139, %f135, %f138, %f134;
	mul.ftz.f32 	%f140, %f60, %f60;
	mul.ftz.f32 	%f141, %f139, %f140;
	mov.f32 	%f142, 0f40c00000;   	// 6
	div.approx.ftz.f32 	%f45, %f141, %f142;
	.loc	3	80	0
	ld.global.f32 	%f143, [%rd16+64];
	mul.ftz.f32 	%f144, %f143, %f129;
	ld.global.f32 	%f145, [%rd18+64];
	fma.rn.ftz.f32 	%f146, %f145, %f128, %f144;
	add.ftz.f32 	%f22, %f45, %f146;
$Lt_22_37890:
	.loc	22	52	0
	mov.f32 	%f119, %f22;
$Lt_22_35842:
$Lt_22_35330:
	.loc	22	53	0
	sub.ftz.f32 	%f147, %f119, %f4;
	fma.rn.ftz.f32 	%f12, %f10, %f147, %f4;
	mov.f32 	%f148, 0f00000000;   	// 0
	setp.lt.ftz.f32 	%p26, %f2, %f148;
	@!%p26 bra 	$Lt_22_38658;
	.loc	22	56	0
	ld.global.f32 	%f149, [sSpline+664];
	add.ftz.f32 	%f150, %f2, %f149;
	bra.uni 	$Lt_22_38914;
$Lt_22_38658:
	ld.global.u64 	%rd53, [sSpline+792];
	mov.f32 	%f151, 0f3f800000;   	// 1
	setp.gt.ftz.f32 	%p27, %f2, %f151;
	@!%p27 bra 	$Lt_22_39170;
	mov.f32 	%f152, 0f3f800000;   	// 1
	mul.lo.u64 	%rd54, %rd53, 4;
	add.u64 	%rd55, %rd1, %rd54;
	ld.global.f32 	%f153, [%rd55+660];
	sub.ftz.f32 	%f154, %f152, %f153;
	sub.ftz.f32 	%f150, %f2, %f154;
	bra.uni 	$Lt_22_38914;
$Lt_22_39170:
	.loc	3	54	0
	mov.f32 	%f22, %f2;
	.loc	3	58	0
	sub.u64 	%rd56, %rd53, 1;
	mov.s64 	%rd6, %rd56;
	mov.u64 	%rd57, 1;
	setp.le.u64 	%p28, %rd56, %rd57;
	@%p28 bra 	$Lt_22_45570;
	mov.u64 	%rd8, 0;
$Lt_22_39938:
 //<loop> Loop body line 58
	add.u64 	%rd58, %rd6, %rd8;
	shr.u64 	%rd10, %rd58, 1;
	mul.lo.u64 	%rd59, %rd10, 4;
	add.u64 	%rd60, %rd1, %rd59;
	ld.global.f32 	%f155, [%rd60+600];
	setp.gt.ftz.f32 	%p29, %f155, %f2;
	@!%p29 bra 	$Lt_22_40450;
 //<loop> Part of loop body line 58, head labeled $Lt_22_39938
	.loc	3	67	0
	mov.s64 	%rd6, %rd10;
	bra.uni 	$Lt_22_40194;
$Lt_22_40450:
 //<loop> Part of loop body line 58, head labeled $Lt_22_39938
	.loc	3	71	0
	mov.s64 	%rd8, %rd10;
$Lt_22_40194:
 //<loop> Part of loop body line 58, head labeled $Lt_22_39938
	sub.u64 	%rd61, %rd6, %rd8;
	mov.u64 	%rd62, 1;
	setp.gt.u64 	%p30, %rd61, %rd62;
	@%p30 bra 	$Lt_22_39938;
	bra.uni 	$Lt_22_39426;
$Lt_22_45570:
	mov.u64 	%rd8, 0;
$Lt_22_39426:
	mul.lo.u64 	%rd63, %rd6, 4;
	add.u64 	%rd16, %rd1, %rd63;
	mul.lo.u64 	%rd64, %rd8, 4;
	add.u64 	%rd18, %rd1, %rd64;
	ld.global.f32 	%f156, [%rd16+600];
	ld.global.f32 	%f157, [%rd18+600];
	sub.ftz.f32 	%f158, %f156, %f157;
	mov.f32 	%f159, 0f00000000;   	// 0
	setp.neu.ftz.f32 	%p31, %f158, %f159;
	@!%p31 bra 	$Lt_22_40962;
	.loc	27	529	0
	sub.ftz.f32 	%f160, %f156, %f2;
	sub.ftz.f32 	%f161, %f2, %f157;
	div.approx.ftz.f32 	%f162, %f160, %f158;
	div.approx.ftz.f32 	%f163, %f161, %f158;
	ld.global.f32 	%f164, [%rd16+728];
	mul.ftz.f32 	%f165, %f163, %f163;
	mul.ftz.f32 	%f166, %f163, %f165;
	sub.ftz.f32 	%f167, %f166, %f163;
	mul.ftz.f32 	%f168, %f164, %f167;
	ld.global.f32 	%f169, [%rd18+728];
	mul.ftz.f32 	%f170, %f162, %f162;
	mul.ftz.f32 	%f171, %f162, %f170;
	sub.ftz.f32 	%f172, %f171, %f162;
	fma.rn.ftz.f32 	%f173, %f169, %f172, %f168;
	mul.ftz.f32 	%f174, %f158, %f158;
	mul.ftz.f32 	%f175, %f173, %f174;
	mov.f32 	%f176, 0f40c00000;   	// 6
	div.approx.ftz.f32 	%f45, %f175, %f176;
	.loc	3	80	0
	ld.global.f32 	%f177, [%rd16+664];
	mul.ftz.f32 	%f178, %f177, %f163;
	ld.global.f32 	%f179, [%rd18+664];
	fma.rn.ftz.f32 	%f180, %f179, %f162, %f178;
	add.ftz.f32 	%f22, %f45, %f180;
$Lt_22_40962:
	.loc	22	56	0
	mov.f32 	%f150, %f22;
$Lt_22_38914:
$Lt_22_38402:
	mov.f32 	%f181, 0f00000000;   	// 0
	setp.lt.ftz.f32 	%p32, %f150, %f181;
	@!%p32 bra 	$Lt_22_41730;
	.loc	22	57	0
	ld.global.f32 	%f182, [sSpline+64];
	add.ftz.f32 	%f183, %f182, %f150;
	bra.uni 	$Lt_22_41986;
$Lt_22_41730:
	ld.global.u64 	%rd19, [sSpline+192];
	mov.f32 	%f184, 0f3f800000;   	// 1
	setp.gt.ftz.f32 	%p33, %f150, %f184;
	@!%p33 bra 	$Lt_22_42242;
	mov.f32 	%f185, 0f3f800000;   	// 1
	mul.lo.u64 	%rd65, %rd19, 4;
	add.u64 	%rd66, %rd1, %rd65;
	ld.global.f32 	%f186, [%rd66+60];
	sub.ftz.f32 	%f187, %f185, %f186;
	sub.ftz.f32 	%f183, %f150, %f187;
	bra.uni 	$Lt_22_41986;
$Lt_22_42242:
	.loc	3	54	0
	mov.f32 	%f22, %f150;
	.loc	3	58	0
	sub.u64 	%rd22, %rd19, 1;
	mov.s64 	%rd6, %rd22;
	mov.u64 	%rd67, 1;
	setp.le.u64 	%p34, %rd22, %rd67;
	@%p34 bra 	$Lt_22_45826;
	mov.u64 	%rd8, 0;
$Lt_22_43010:
 //<loop> Loop body line 58
	add.u64 	%rd68, %rd6, %rd8;
	shr.u64 	%rd10, %rd68, 1;
	mul.lo.u64 	%rd69, %rd10, 4;
	add.u64 	%rd70, %rd1, %rd69;
	ld.global.f32 	%f188, [%rd70+0];
	setp.gt.ftz.f32 	%p35, %f188, %f150;
	@!%p35 bra 	$Lt_22_43522;
 //<loop> Part of loop body line 58, head labeled $Lt_22_43010
	.loc	3	67	0
	mov.s64 	%rd6, %rd10;
	bra.uni 	$Lt_22_43266;
$Lt_22_43522:
 //<loop> Part of loop body line 58, head labeled $Lt_22_43010
	.loc	3	71	0
	mov.s64 	%rd8, %rd10;
$Lt_22_43266:
 //<loop> Part of loop body line 58, head labeled $Lt_22_43010
	sub.u64 	%rd71, %rd6, %rd8;
	mov.u64 	%rd72, 1;
	setp.gt.u64 	%p36, %rd71, %rd72;
	@%p36 bra 	$Lt_22_43010;
	bra.uni 	$Lt_22_42498;
$Lt_22_45826:
	mov.u64 	%rd8, 0;
$Lt_22_42498:
	mul.lo.u64 	%rd73, %rd6, 4;
	add.u64 	%rd16, %rd1, %rd73;
	mul.lo.u64 	%rd74, %rd8, 4;
	add.u64 	%rd18, %rd1, %rd74;
	ld.global.f32 	%f58, [%rd16+0];
	ld.global.f32 	%f59, [%rd18+0];
	sub.ftz.f32 	%f60, %f58, %f59;
	mov.f32 	%f189, 0f00000000;   	// 0
	setp.neu.ftz.f32 	%p37, %f60, %f189;
	@!%p37 bra 	$Lt_22_44034;
	.loc	27	529	0
	sub.ftz.f32 	%f190, %f58, %f150;
	sub.ftz.f32 	%f191, %f150, %f59;
	div.approx.ftz.f32 	%f192, %f190, %f60;
	div.approx.ftz.f32 	%f193, %f191, %f60;
	ld.global.f32 	%f194, [%rd16+128];
	mul.ftz.f32 	%f195, %f193, %f193;
	mul.ftz.f32 	%f196, %f193, %f195;
	sub.ftz.f32 	%f197, %f196, %f193;
	mul.ftz.f32 	%f198, %f194, %f197;
	ld.global.f32 	%f199, [%rd18+128];
	mul.ftz.f32 	%f200, %f192, %f192;
	mul.ftz.f32 	%f201, %f192, %f200;
	sub.ftz.f32 	%f202, %f201, %f192;
	fma.rn.ftz.f32 	%f203, %f199, %f202, %f198;
	mul.ftz.f32 	%f204, %f60, %f60;
	mul.ftz.f32 	%f205, %f203, %f204;
	mov.f32 	%f206, 0f40c00000;   	// 6
	div.approx.ftz.f32 	%f45, %f205, %f206;
	.loc	3	80	0
	ld.global.f32 	%f207, [%rd16+64];
	mul.ftz.f32 	%f208, %f207, %f193;
	ld.global.f32 	%f209, [%rd18+64];
	fma.rn.ftz.f32 	%f210, %f209, %f192, %f208;
	add.ftz.f32 	%f22, %f45, %f210;
$Lt_22_44034:
	.loc	22	57	0
	mov.f32 	%f183, %f22;
$Lt_22_41986:
$Lt_22_41474:
	.loc	22	58	0
	sub.ftz.f32 	%f211, %f183, %f2;
	fma.rn.ftz.f32 	%f11, %f10, %f211, %f2;
$Lt_22_25602:
	.loc	22	61	0
	mov.f32 	%f212, %f11;
	st.param.f32 	[__cudaretf__Z13CalculateBGRA8PixelRGBf+0], %f212;
	mov.f32 	%f213, %f12;
	st.param.f32 	[__cudaretf__Z13CalculateBGRA8PixelRGBf+4], %f213;
	mov.f32 	%f214, %f13;
	st.param.f32 	[__cudaretf__Z13CalculateBGRA8PixelRGBf+8], %f214;
	mov.f32 	%f215, %f8;
	st.param.f32 	[__cudaretf__Z13CalculateBGRA8PixelRGBf+12], %f215;
	ret;
$LDWend__Z13CalculateBGRA8PixelRGBf:
	} // _Z13CalculateBGRA8PixelRGBf

	.visible .func (.param .f32 __cudaretf__Z4LERPIfET_S0_S0_S0_) _Z4LERPIfET_S0_S0_S0_ (.param .f32 __cudaparmf1__Z4LERPIfET_S0_S0_S0_, .param .f32 __cudaparmf2__Z4LERPIfET_S0_S0_S0_, .param .f32 __cudaparmf3__Z4LERPIfET_S0_S0_S0_)
	{
	.reg .f32 %f<10>;
	.loc	23	78	0
$LDWbegin__Z4LERPIfET_S0_S0_S0_:
	ld.param.f32 	%f1, [__cudaparmf1__Z4LERPIfET_S0_S0_S0_];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf2__Z4LERPIfET_S0_S0_S0_];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf3__Z4LERPIfET_S0_S0_S0_];
	mov.f32 	%f6, %f5;
	.loc	23	79	0
	sub.ftz.f32 	%f7, %f4, %f2;
	fma.rn.ftz.f32 	%f8, %f6, %f7, %f2;
	st.param.f32 	[__cudaretf__Z4LERPIfET_S0_S0_S0_], %f8;
	ret;
$LDWend__Z4LERPIfET_S0_S0_S0_:
	} // _Z4LERPIfET_S0_S0_S0_

	.visible .func (.param .f32 __cudaretf__Z6Read2DIfET_PKS0_iii) _Z6Read2DIfET_PKS0_iii (.param .u64 __cudaparmf1__Z6Read2DIfET_PKS0_iii, .param .s32 __cudaparmf2__Z6Read2DIfET_PKS0_iii, .param .s32 __cudaparmf3__Z6Read2DIfET_PKS0_iii, .param .s32 __cudaparmf4__Z6Read2DIfET_PKS0_iii)
	{
	.reg .u32 %r<10>;
	.reg .u64 %rd<7>;
	.reg .f32 %f<3>;
	.loc	20	114	0
$LDWbegin__Z6Read2DIfET_PKS0_iii:
	ld.param.u64 	%rd1, [__cudaparmf1__Z6Read2DIfET_PKS0_iii];
	mov.s64 	%rd2, %rd1;
	ld.param.u32 	%r1, [__cudaparmf2__Z6Read2DIfET_PKS0_iii];
	mov.s32 	%r2, %r1;
	ld.param.u32 	%r3, [__cudaparmf3__Z6Read2DIfET_PKS0_iii];
	mov.s32 	%r4, %r3;
	ld.param.u32 	%r5, [__cudaparmf4__Z6Read2DIfET_PKS0_iii];
	mov.s32 	%r6, %r5;
	.loc	20	115	0
	mul.lo.s32 	%r7, %r2, %r6;
	add.s32 	%r8, %r4, %r7;
	cvt.s64.s32 	%rd3, %r8;
	mul.wide.s32 	%rd4, %r8, 4;
	add.u64 	%rd5, %rd2, %rd4;
	ld.f32 	%f1, [%rd5+0];
	st.param.f32 	[__cudaretf__Z6Read2DIfET_PKS0_iii], %f1;
	ret;
$LDWend__Z6Read2DIfET_PKS0_iii:
	} // _Z6Read2DIfET_PKS0_iii

	.entry RGBCurves_MaskKernel (
		.param .u64 __cudaparm_RGBCurves_MaskKernel_inImage,
		.param .s32 __cudaparm_RGBCurves_MaskKernel_inPitch,
		.param .u64 __cudaparm_RGBCurves_MaskKernel_inSecondaryMask,
		.param .s32 __cudaparm_RGBCurves_MaskKernel_inSecondaryPitch,
		.param .u32 __cudaparm_RGBCurves_MaskKernel_inDeviceFormat,
		.param .s32 __cudaparm_RGBCurves_MaskKernel_inWidth,
		.param .s32 __cudaparm_RGBCurves_MaskKernel_inHeight)
	{
	.reg .u32 %r<31>;
	.reg .u64 %rd<13>;
	.reg .f32 %f<5>;
	.reg .pred %p<5>;
	.loc	22	74	0
$LDWbegin_RGBCurves_MaskKernel:
	.loc	22	77	0
	cvt.s32.u32 	%r1, %ctaid.x;
	cvt.s32.u32 	%r2, %ntid.x;
	mul.lo.s32 	%r3, %r1, %r2;
	cvt.s32.u32 	%r4, %ctaid.y;
	cvt.s32.u32 	%r5, %ntid.y;
	mul.lo.s32 	%r6, %r4, %r5;
	mov.u32 	%r7, %tid.x;
	add.u32 	%r8, %r3, %r7;
	mov.u32 	%r9, %tid.y;
	add.u32 	%r10, %r6, %r9;
	ld.param.s32 	%r11, [__cudaparm_RGBCurves_MaskKernel_inWidth];
	set.gt.u32.s32 	%r12, %r11, %r8;
	neg.s32 	%r13, %r12;
	ld.param.s32 	%r14, [__cudaparm_RGBCurves_MaskKernel_inHeight];
	set.gt.u32.s32 	%r15, %r14, %r10;
	neg.s32 	%r16, %r15;
	and.b32 	%r17, %r13, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_25_3842;
	ld.param.u64 	%rd1, [__cudaparm_RGBCurves_MaskKernel_inSecondaryMask];
	mov.u64 	%rd2, 0;
	setp.eq.u64 	%p2, %rd1, %rd2;
	@%p2 bra 	$Lt_25_4610;
	.loc	22	82	0
	ld.param.s32 	%r19, [__cudaparm_RGBCurves_MaskKernel_inSecondaryPitch];
	mul.lo.s32 	%r20, %r19, %r10;
	add.s32 	%r21, %r8, %r20;
	cvt.s64.s32 	%rd3, %r21;
	mul.wide.s32 	%rd4, %r21, 4;
	add.u64 	%rd5, %rd1, %rd4;
	ld.global.f32 	%f1, [%rd5+0];
	bra.uni 	$Lt_25_4354;
$Lt_25_4610:
	mov.f32 	%f1, 0f3f800000;     	// 1
$Lt_25_4354:
	ld.param.s32 	%r22, [__cudaparm_RGBCurves_MaskKernel_inPitch];
	mul.lo.s32 	%r23, %r22, %r10;
	add.s32 	%r24, %r8, %r23;
	cvt.s64.s32 	%rd6, %r24;
	ld.param.u64 	%rd7, [__cudaparm_RGBCurves_MaskKernel_inImage];
	ld.param.s32 	%r25, [__cudaparm_RGBCurves_MaskKernel_inDeviceFormat];
	mov.u32 	%r26, 0;
	setp.ne.s32 	%p3, %r25, %r26;
	@%p3 bra 	$Lt_25_5122;
	.loc	20	126	0
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f1;
	mov.b32		%r27, %b1; }
	mov.s32 	%r28, %r27;
	mul.lo.u64 	%rd8, %rd6, 8;
	add.u64 	%rd9, %rd7, %rd8;
	mov.f32 	%f2, 0f3f800000;     	// 1
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f2;
	mov.b32		%r29, %b1; }
	st.global.v4.u16 	[%rd9+0], {%r28,%r28,%r28,%r29};
	.loc	22	86	0
	bra.uni 	$Lt_25_4866;
$Lt_25_5122:
	.loc	20	126	0
	mul.lo.u64 	%rd10, %rd6, 16;
	add.u64 	%rd11, %rd7, %rd10;
	mov.f32 	%f3, 0f3f800000;     	// 1
	st.global.v4.f32 	[%rd11+0], {%f1,%f1,%f1,%f3};
$Lt_25_4866:
$Lt_25_3842:
	.loc	22	88	0
	exit;
$LDWend_RGBCurves_MaskKernel:
	} // RGBCurves_MaskKernel
	.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};

	.entry RGBCurves_LumaKernel (
		.param .u64 __cudaparm_RGBCurves_LumaKernel_inImage,
		.param .s32 __cudaparm_RGBCurves_LumaKernel_inPitch,
		.param .u64 __cudaparm_RGBCurves_LumaKernel_inSecondaryMask,
		.param .s32 __cudaparm_RGBCurves_LumaKernel_inSecondaryPitch,
		.param .u32 __cudaparm_RGBCurves_LumaKernel_inDeviceFormat,
		.param .s32 __cudaparm_RGBCurves_LumaKernel_inWidth,
		.param .s32 __cudaparm_RGBCurves_LumaKernel_inHeight)
	{
	.reg .u32 %r<35>;
	.reg .u64 %rd<91>;
	.reg .f32 %f<217>;
	.reg .pred %p<42>;
	.loc	22	100	0
$LDWbegin_RGBCurves_LumaKernel:
	.loc	22	103	0
	cvt.s32.u32 	%r1, %ctaid.x;
	cvt.s32.u32 	%r2, %ntid.x;
	mul.lo.s32 	%r3, %r1, %r2;
	cvt.s32.u32 	%r4, %ctaid.y;
	cvt.s32.u32 	%r5, %ntid.y;
	mul.lo.s32 	%r6, %r4, %r5;
	mov.u32 	%r7, %tid.x;
	add.u32 	%r8, %r3, %r7;
	mov.u32 	%r9, %tid.y;
	add.u32 	%r10, %r6, %r9;
	ld.param.s32 	%r11, [__cudaparm_RGBCurves_LumaKernel_inWidth];
	set.gt.u32.s32 	%r12, %r11, %r8;
	neg.s32 	%r13, %r12;
	ld.param.s32 	%r14, [__cudaparm_RGBCurves_LumaKernel_inHeight];
	set.gt.u32.s32 	%r15, %r14, %r10;
	neg.s32 	%r16, %r15;
	and.b32 	%r17, %r13, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_26_49410;
	ld.param.s32 	%r19, [__cudaparm_RGBCurves_LumaKernel_inDeviceFormat];
	mov.s32 	%r20, 0;
	setp.eq.s32 	%p2, %r19, %r20;
	ld.param.s32 	%r21, [__cudaparm_RGBCurves_LumaKernel_inPitch];
	mul.lo.s32 	%r22, %r21, %r10;
	add.s32 	%r23, %r8, %r22;
	cvt.s64.s32 	%rd1, %r23;
	ld.param.u64 	%rd2, [__cudaparm_RGBCurves_LumaKernel_inImage];
	@!%p2 bra 	$Lt_26_29698;
	.loc	20	115	0
	mul.lo.u64 	%rd3, %rd1, 8;
	add.u64 	%rd4, %rd2, %rd3;
	ld.global.v4.u16 	{%r24,%r25,%r26,%r27}, [%rd4+0];
	.loc	22	106	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_26_29442;
$Lt_26_29698:
	mul.lo.u64 	%rd5, %rd1, 16;
	add.u64 	%rd6, %rd2, %rd5;
	ld.global.v4.f32 	{%f1,%f2,%f3,%f4}, [%rd6+0];
$Lt_26_29442:
	ld.param.u64 	%rd7, [__cudaparm_RGBCurves_LumaKernel_inSecondaryMask];
	mov.u64 	%rd8, 0;
	setp.eq.u64 	%p3, %rd7, %rd8;
	@%p3 bra 	$Lt_26_30210;
	.loc	22	107	0
	ld.param.s32 	%r28, [__cudaparm_RGBCurves_LumaKernel_inSecondaryPitch];
	mul.lo.s32 	%r29, %r28, %r10;
	add.s32 	%r30, %r8, %r29;
	cvt.s64.s32 	%rd9, %r30;
	mul.wide.s32 	%rd10, %r30, 4;
	add.u64 	%rd11, %rd7, %rd10;
	ld.global.f32 	%f5, [%rd11+0];
	bra.uni 	$Lt_26_29954;
$Lt_26_30210:
	mov.f32 	%f5, 0f3f800000;     	// 1
$Lt_26_29954:
	.loc	22	40	0
	mov.f32 	%f6, %f1;
	mov.f32 	%f7, %f6;
	mov.f32 	%f8, %f2;
	mov.f32 	%f9, %f8;
	mov.f32 	%f10, %f3;
	mov.f32 	%f11, %f10;
	.loc	22	41	0
	mov.f32 	%f12, 0f3a83126f;    	// 0.001
	setp.gt.ftz.f32 	%p4, %f5, %f12;
	@!%p4 bra 	$Lt_26_30466;
	mov.f32 	%f13, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p5, %f10, %f13;
	@!%p5 bra 	$Lt_26_31234;
	.loc	22	46	0
	ld.global.f32 	%f14, [sSpline+264];
	add.ftz.f32 	%f15, %f14, %f10;
	mov.u64 	%rd12, sSpline;
	bra.uni 	$Lt_26_31490;
$Lt_26_31234:
	ld.global.u64 	%rd13, [sSpline+392];
	mov.f32 	%f16, 0f3f800000;    	// 1
	setp.gt.ftz.f32 	%p6, %f10, %f16;
	@!%p6 bra 	$Lt_26_31746;
	mov.u64 	%rd12, sSpline;
	mov.f32 	%f17, 0f3f800000;    	// 1
	mul.lo.u64 	%rd14, %rd13, 4;
	add.u64 	%rd15, %rd12, %rd14;
	ld.global.f32 	%f18, [%rd15+260];
	sub.ftz.f32 	%f19, %f17, %f18;
	sub.ftz.f32 	%f15, %f10, %f19;
	bra.uni 	$Lt_26_31490;
$Lt_26_31746:
	.loc	3	54	0
	mov.f32 	%f20, %f10;
	.loc	3	58	0
	sub.u64 	%rd16, %rd13, 1;
	mov.s64 	%rd17, %rd16;
	mov.u64 	%rd18, 1;
	setp.le.u64 	%p7, %rd16, %rd18;
	@%p7 bra 	$Lt_26_49922;
	mov.u64 	%rd19, 0;
	mov.u64 	%rd12, sSpline;
$Lt_26_32514:
 //<loop> Loop body line 58
	add.u64 	%rd20, %rd17, %rd19;
	shr.u64 	%rd21, %rd20, 1;
	mul.lo.u64 	%rd22, %rd21, 4;
	add.u64 	%rd23, %rd12, %rd22;
	ld.global.f32 	%f21, [%rd23+200];
	setp.gt.ftz.f32 	%p8, %f21, %f10;
	@!%p8 bra 	$Lt_26_33026;
 //<loop> Part of loop body line 58, head labeled $Lt_26_32514
	.loc	3	67	0
	mov.s64 	%rd17, %rd21;
	bra.uni 	$Lt_26_32770;
$Lt_26_33026:
 //<loop> Part of loop body line 58, head labeled $Lt_26_32514
	.loc	3	71	0
	mov.s64 	%rd19, %rd21;
$Lt_26_32770:
 //<loop> Part of loop body line 58, head labeled $Lt_26_32514
	sub.u64 	%rd24, %rd17, %rd19;
	mov.u64 	%rd25, 1;
	setp.gt.u64 	%p9, %rd24, %rd25;
	@%p9 bra 	$Lt_26_32514;
	bra.uni 	$Lt_26_32002;
$Lt_26_49922:
	mov.u64 	%rd19, 0;
	mov.u64 	%rd12, sSpline;
$Lt_26_32002:
	mul.lo.u64 	%rd26, %rd17, 4;
	add.u64 	%rd27, %rd12, %rd26;
	mul.lo.u64 	%rd28, %rd19, 4;
	add.u64 	%rd29, %rd12, %rd28;
	ld.global.f32 	%f22, [%rd27+200];
	ld.global.f32 	%f23, [%rd29+200];
	sub.ftz.f32 	%f24, %f22, %f23;
	mov.f32 	%f25, 0f00000000;    	// 0
	setp.neu.ftz.f32 	%p10, %f24, %f25;
	@!%p10 bra 	$Lt_26_33538;
	.loc	27	529	0
	sub.ftz.f32 	%f26, %f22, %f10;
	sub.ftz.f32 	%f27, %f10, %f23;
	div.approx.ftz.f32 	%f28, %f26, %f24;
	div.approx.ftz.f32 	%f29, %f27, %f24;
	ld.global.f32 	%f30, [%rd27+328];
	mul.ftz.f32 	%f31, %f29, %f29;
	mul.ftz.f32 	%f32, %f29, %f31;
	sub.ftz.f32 	%f33, %f32, %f29;
	mul.ftz.f32 	%f34, %f30, %f33;
	ld.global.f32 	%f35, [%rd29+328];
	mul.ftz.f32 	%f36, %f28, %f28;
	mul.ftz.f32 	%f37, %f28, %f36;
	sub.ftz.f32 	%f38, %f37, %f28;
	fma.rn.ftz.f32 	%f39, %f35, %f38, %f34;
	mul.ftz.f32 	%f40, %f24, %f24;
	mul.ftz.f32 	%f41, %f39, %f40;
	mov.f32 	%f42, 0f40c00000;    	// 6
	div.approx.ftz.f32 	%f43, %f41, %f42;
	.loc	3	80	0
	ld.global.f32 	%f44, [%rd27+264];
	mul.ftz.f32 	%f45, %f44, %f29;
	ld.global.f32 	%f46, [%rd29+264];
	fma.rn.ftz.f32 	%f47, %f46, %f28, %f45;
	add.ftz.f32 	%f20, %f43, %f47;
$Lt_26_33538:
	.loc	22	46	0
	mov.f32 	%f15, %f20;
$Lt_26_31490:
$Lt_26_30978:
	mov.f32 	%f48, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p11, %f15, %f48;
	@!%p11 bra 	$Lt_26_34306;
	.loc	22	47	0
	ld.global.f32 	%f49, [sSpline+64];
	add.ftz.f32 	%f50, %f49, %f15;
	bra.uni 	$Lt_26_34562;
$Lt_26_34306:
	ld.global.u64 	%rd30, [sSpline+192];
	mov.f32 	%f51, 0f3f800000;    	// 1
	setp.gt.ftz.f32 	%p12, %f15, %f51;
	@!%p12 bra 	$Lt_26_34818;
	mov.f32 	%f52, 0f3f800000;    	// 1
	mul.lo.u64 	%rd31, %rd30, 4;
	add.u64 	%rd32, %rd12, %rd31;
	ld.global.f32 	%f53, [%rd32+60];
	sub.ftz.f32 	%f54, %f52, %f53;
	sub.ftz.f32 	%f50, %f15, %f54;
	bra.uni 	$Lt_26_34562;
$Lt_26_34818:
	.loc	3	54	0
	mov.f32 	%f20, %f15;
	.loc	3	58	0
	sub.u64 	%rd33, %rd30, 1;
	mov.s64 	%rd17, %rd33;
	mov.u64 	%rd34, 1;
	setp.le.u64 	%p13, %rd33, %rd34;
	@%p13 bra 	$Lt_26_50178;
	mov.u64 	%rd19, 0;
$Lt_26_35586:
 //<loop> Loop body line 58
	add.u64 	%rd35, %rd17, %rd19;
	shr.u64 	%rd21, %rd35, 1;
	mul.lo.u64 	%rd36, %rd21, 4;
	add.u64 	%rd37, %rd12, %rd36;
	ld.global.f32 	%f55, [%rd37+0];
	setp.gt.ftz.f32 	%p14, %f55, %f15;
	@!%p14 bra 	$Lt_26_36098;
 //<loop> Part of loop body line 58, head labeled $Lt_26_35586
	.loc	3	67	0
	mov.s64 	%rd17, %rd21;
	bra.uni 	$Lt_26_35842;
$Lt_26_36098:
 //<loop> Part of loop body line 58, head labeled $Lt_26_35586
	.loc	3	71	0
	mov.s64 	%rd19, %rd21;
$Lt_26_35842:
 //<loop> Part of loop body line 58, head labeled $Lt_26_35586
	sub.u64 	%rd38, %rd17, %rd19;
	mov.u64 	%rd39, 1;
	setp.gt.u64 	%p15, %rd38, %rd39;
	@%p15 bra 	$Lt_26_35586;
	bra.uni 	$Lt_26_35074;
$Lt_26_50178:
	mov.u64 	%rd19, 0;
$Lt_26_35074:
	mul.lo.u64 	%rd40, %rd17, 4;
	add.u64 	%rd27, %rd12, %rd40;
	mul.lo.u64 	%rd41, %rd19, 4;
	add.u64 	%rd29, %rd12, %rd41;
	ld.global.f32 	%f56, [%rd27+0];
	ld.global.f32 	%f57, [%rd29+0];
	sub.ftz.f32 	%f58, %f56, %f57;
	mov.f32 	%f59, 0f00000000;    	// 0
	setp.neu.ftz.f32 	%p16, %f58, %f59;
	@!%p16 bra 	$Lt_26_36610;
	.loc	27	529	0
	sub.ftz.f32 	%f60, %f56, %f15;
	sub.ftz.f32 	%f61, %f15, %f57;
	div.approx.ftz.f32 	%f62, %f60, %f58;
	div.approx.ftz.f32 	%f63, %f61, %f58;
	ld.global.f32 	%f64, [%rd27+128];
	mul.ftz.f32 	%f65, %f63, %f63;
	mul.ftz.f32 	%f66, %f63, %f65;
	sub.ftz.f32 	%f67, %f66, %f63;
	mul.ftz.f32 	%f68, %f64, %f67;
	ld.global.f32 	%f69, [%rd29+128];
	mul.ftz.f32 	%f70, %f62, %f62;
	mul.ftz.f32 	%f71, %f62, %f70;
	sub.ftz.f32 	%f72, %f71, %f62;
	fma.rn.ftz.f32 	%f73, %f69, %f72, %f68;
	mul.ftz.f32 	%f74, %f58, %f58;
	mul.ftz.f32 	%f75, %f73, %f74;
	mov.f32 	%f76, 0f40c00000;    	// 6
	div.approx.ftz.f32 	%f43, %f75, %f76;
	.loc	3	80	0
	ld.global.f32 	%f77, [%rd27+64];
	mul.ftz.f32 	%f78, %f77, %f63;
	ld.global.f32 	%f79, [%rd29+64];
	fma.rn.ftz.f32 	%f80, %f79, %f62, %f78;
	add.ftz.f32 	%f20, %f43, %f80;
$Lt_26_36610:
	.loc	22	47	0
	mov.f32 	%f50, %f20;
$Lt_26_34562:
$Lt_26_34050:
	.loc	22	48	0
	sub.ftz.f32 	%f81, %f50, %f10;
	fma.rn.ftz.f32 	%f11, %f5, %f81, %f10;
	mov.f32 	%f82, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p17, %f8, %f82;
	@!%p17 bra 	$Lt_26_37378;
	.loc	22	51	0
	ld.global.f32 	%f83, [sSpline+464];
	add.ftz.f32 	%f84, %f83, %f8;
	bra.uni 	$Lt_26_37634;
$Lt_26_37378:
	ld.global.u64 	%rd42, [sSpline+592];
	mov.f32 	%f85, 0f3f800000;    	// 1
	setp.gt.ftz.f32 	%p18, %f8, %f85;
	@!%p18 bra 	$Lt_26_37890;
	mov.f32 	%f86, 0f3f800000;    	// 1
	mul.lo.u64 	%rd43, %rd42, 4;
	add.u64 	%rd44, %rd12, %rd43;
	ld.global.f32 	%f87, [%rd44+460];
	sub.ftz.f32 	%f88, %f86, %f87;
	sub.ftz.f32 	%f84, %f8, %f88;
	bra.uni 	$Lt_26_37634;
$Lt_26_37890:
	.loc	3	54	0
	mov.f32 	%f20, %f8;
	.loc	3	58	0
	sub.u64 	%rd45, %rd42, 1;
	mov.s64 	%rd17, %rd45;
	mov.u64 	%rd46, 1;
	setp.le.u64 	%p19, %rd45, %rd46;
	@%p19 bra 	$Lt_26_50434;
	mov.u64 	%rd19, 0;
$Lt_26_38658:
 //<loop> Loop body line 58
	add.u64 	%rd47, %rd17, %rd19;
	shr.u64 	%rd21, %rd47, 1;
	mul.lo.u64 	%rd48, %rd21, 4;
	add.u64 	%rd49, %rd12, %rd48;
	ld.global.f32 	%f89, [%rd49+400];
	setp.gt.ftz.f32 	%p20, %f89, %f8;
	@!%p20 bra 	$Lt_26_39170;
 //<loop> Part of loop body line 58, head labeled $Lt_26_38658
	.loc	3	67	0
	mov.s64 	%rd17, %rd21;
	bra.uni 	$Lt_26_38914;
$Lt_26_39170:
 //<loop> Part of loop body line 58, head labeled $Lt_26_38658
	.loc	3	71	0
	mov.s64 	%rd19, %rd21;
$Lt_26_38914:
 //<loop> Part of loop body line 58, head labeled $Lt_26_38658
	sub.u64 	%rd50, %rd17, %rd19;
	mov.u64 	%rd51, 1;
	setp.gt.u64 	%p21, %rd50, %rd51;
	@%p21 bra 	$Lt_26_38658;
	bra.uni 	$Lt_26_38146;
$Lt_26_50434:
	mov.u64 	%rd19, 0;
$Lt_26_38146:
	mul.lo.u64 	%rd52, %rd17, 4;
	add.u64 	%rd27, %rd12, %rd52;
	mul.lo.u64 	%rd53, %rd19, 4;
	add.u64 	%rd29, %rd12, %rd53;
	ld.global.f32 	%f90, [%rd27+400];
	ld.global.f32 	%f91, [%rd29+400];
	sub.ftz.f32 	%f92, %f90, %f91;
	mov.f32 	%f93, 0f00000000;    	// 0
	setp.neu.ftz.f32 	%p22, %f92, %f93;
	@!%p22 bra 	$Lt_26_39682;
	.loc	27	529	0
	sub.ftz.f32 	%f94, %f90, %f8;
	sub.ftz.f32 	%f95, %f8, %f91;
	div.approx.ftz.f32 	%f96, %f94, %f92;
	div.approx.ftz.f32 	%f97, %f95, %f92;
	ld.global.f32 	%f98, [%rd27+528];
	mul.ftz.f32 	%f99, %f97, %f97;
	mul.ftz.f32 	%f100, %f97, %f99;
	sub.ftz.f32 	%f101, %f100, %f97;
	mul.ftz.f32 	%f102, %f98, %f101;
	ld.global.f32 	%f103, [%rd29+528];
	mul.ftz.f32 	%f104, %f96, %f96;
	mul.ftz.f32 	%f105, %f96, %f104;
	sub.ftz.f32 	%f106, %f105, %f96;
	fma.rn.ftz.f32 	%f107, %f103, %f106, %f102;
	mul.ftz.f32 	%f108, %f92, %f92;
	mul.ftz.f32 	%f109, %f107, %f108;
	mov.f32 	%f110, 0f40c00000;   	// 6
	div.approx.ftz.f32 	%f43, %f109, %f110;
	.loc	3	80	0
	ld.global.f32 	%f111, [%rd27+464];
	mul.ftz.f32 	%f112, %f111, %f97;
	ld.global.f32 	%f113, [%rd29+464];
	fma.rn.ftz.f32 	%f114, %f113, %f96, %f112;
	add.ftz.f32 	%f20, %f43, %f114;
$Lt_26_39682:
	.loc	22	51	0
	mov.f32 	%f84, %f20;
$Lt_26_37634:
$Lt_26_37122:
	mov.f32 	%f115, 0f00000000;   	// 0
	setp.lt.ftz.f32 	%p23, %f84, %f115;
	@!%p23 bra 	$Lt_26_40450;
	.loc	22	52	0
	ld.global.f32 	%f116, [sSpline+64];
	add.ftz.f32 	%f117, %f116, %f84;
	bra.uni 	$Lt_26_40706;
$Lt_26_40450:
	ld.global.u64 	%rd30, [sSpline+192];
	mov.f32 	%f118, 0f3f800000;   	// 1
	setp.gt.ftz.f32 	%p24, %f84, %f118;
	@!%p24 bra 	$Lt_26_40962;
	mov.f32 	%f119, 0f3f800000;   	// 1
	mul.lo.u64 	%rd54, %rd30, 4;
	add.u64 	%rd55, %rd12, %rd54;
	ld.global.f32 	%f120, [%rd55+60];
	sub.ftz.f32 	%f121, %f119, %f120;
	sub.ftz.f32 	%f117, %f84, %f121;
	bra.uni 	$Lt_26_40706;
$Lt_26_40962:
	.loc	3	54	0
	mov.f32 	%f20, %f84;
	.loc	3	58	0
	sub.u64 	%rd33, %rd30, 1;
	mov.s64 	%rd17, %rd33;
	mov.u64 	%rd56, 1;
	setp.le.u64 	%p25, %rd33, %rd56;
	@%p25 bra 	$Lt_26_50690;
	mov.u64 	%rd19, 0;
$Lt_26_41730:
 //<loop> Loop body line 58
	add.u64 	%rd57, %rd17, %rd19;
	shr.u64 	%rd21, %rd57, 1;
	mul.lo.u64 	%rd58, %rd21, 4;
	add.u64 	%rd59, %rd12, %rd58;
	ld.global.f32 	%f122, [%rd59+0];
	setp.gt.ftz.f32 	%p26, %f122, %f84;
	@!%p26 bra 	$Lt_26_42242;
 //<loop> Part of loop body line 58, head labeled $Lt_26_41730
	.loc	3	67	0
	mov.s64 	%rd17, %rd21;
	bra.uni 	$Lt_26_41986;
$Lt_26_42242:
 //<loop> Part of loop body line 58, head labeled $Lt_26_41730
	.loc	3	71	0
	mov.s64 	%rd19, %rd21;
$Lt_26_41986:
 //<loop> Part of loop body line 58, head labeled $Lt_26_41730
	sub.u64 	%rd60, %rd17, %rd19;
	mov.u64 	%rd61, 1;
	setp.gt.u64 	%p27, %rd60, %rd61;
	@%p27 bra 	$Lt_26_41730;
	bra.uni 	$Lt_26_41218;
$Lt_26_50690:
	mov.u64 	%rd19, 0;
$Lt_26_41218:
	mul.lo.u64 	%rd62, %rd17, 4;
	add.u64 	%rd27, %rd12, %rd62;
	mul.lo.u64 	%rd63, %rd19, 4;
	add.u64 	%rd29, %rd12, %rd63;
	ld.global.f32 	%f56, [%rd27+0];
	ld.global.f32 	%f57, [%rd29+0];
	sub.ftz.f32 	%f58, %f56, %f57;
	mov.f32 	%f123, 0f00000000;   	// 0
	setp.neu.ftz.f32 	%p28, %f58, %f123;
	@!%p28 bra 	$Lt_26_42754;
	.loc	27	529	0
	sub.ftz.f32 	%f124, %f56, %f84;
	sub.ftz.f32 	%f125, %f84, %f57;
	div.approx.ftz.f32 	%f126, %f124, %f58;
	div.approx.ftz.f32 	%f127, %f125, %f58;
	ld.global.f32 	%f128, [%rd27+128];
	mul.ftz.f32 	%f129, %f127, %f127;
	mul.ftz.f32 	%f130, %f127, %f129;
	sub.ftz.f32 	%f131, %f130, %f127;
	mul.ftz.f32 	%f132, %f128, %f131;
	ld.global.f32 	%f133, [%rd29+128];
	mul.ftz.f32 	%f134, %f126, %f126;
	mul.ftz.f32 	%f135, %f126, %f134;
	sub.ftz.f32 	%f136, %f135, %f126;
	fma.rn.ftz.f32 	%f137, %f133, %f136, %f132;
	mul.ftz.f32 	%f138, %f58, %f58;
	mul.ftz.f32 	%f139, %f137, %f138;
	mov.f32 	%f140, 0f40c00000;   	// 6
	div.approx.ftz.f32 	%f43, %f139, %f140;
	.loc	3	80	0
	ld.global.f32 	%f141, [%rd27+64];
	mul.ftz.f32 	%f142, %f141, %f127;
	ld.global.f32 	%f143, [%rd29+64];
	fma.rn.ftz.f32 	%f144, %f143, %f126, %f142;
	add.ftz.f32 	%f20, %f43, %f144;
$Lt_26_42754:
	.loc	22	52	0
	mov.f32 	%f117, %f20;
$Lt_26_40706:
$Lt_26_40194:
	.loc	22	53	0
	sub.ftz.f32 	%f145, %f117, %f8;
	fma.rn.ftz.f32 	%f9, %f5, %f145, %f8;
	mov.f32 	%f146, 0f00000000;   	// 0
	setp.lt.ftz.f32 	%p29, %f6, %f146;
	@!%p29 bra 	$Lt_26_43522;
	.loc	22	56	0
	ld.global.f32 	%f147, [sSpline+664];
	add.ftz.f32 	%f148, %f6, %f147;
	bra.uni 	$Lt_26_43778;
$Lt_26_43522:
	ld.global.u64 	%rd64, [sSpline+792];
	mov.f32 	%f149, 0f3f800000;   	// 1
	setp.gt.ftz.f32 	%p30, %f6, %f149;
	@!%p30 bra 	$Lt_26_44034;
	mov.f32 	%f150, 0f3f800000;   	// 1
	mul.lo.u64 	%rd65, %rd64, 4;
	add.u64 	%rd66, %rd12, %rd65;
	ld.global.f32 	%f151, [%rd66+660];
	sub.ftz.f32 	%f152, %f150, %f151;
	sub.ftz.f32 	%f148, %f6, %f152;
	bra.uni 	$Lt_26_43778;
$Lt_26_44034:
	.loc	3	54	0
	mov.f32 	%f20, %f6;
	.loc	3	58	0
	sub.u64 	%rd67, %rd64, 1;
	mov.s64 	%rd17, %rd67;
	mov.u64 	%rd68, 1;
	setp.le.u64 	%p31, %rd67, %rd68;
	@%p31 bra 	$Lt_26_50946;
	mov.u64 	%rd19, 0;
$Lt_26_44802:
 //<loop> Loop body line 58
	add.u64 	%rd69, %rd17, %rd19;
	shr.u64 	%rd21, %rd69, 1;
	mul.lo.u64 	%rd70, %rd21, 4;
	add.u64 	%rd71, %rd12, %rd70;
	ld.global.f32 	%f153, [%rd71+600];
	setp.gt.ftz.f32 	%p32, %f153, %f6;
	@!%p32 bra 	$Lt_26_45314;
 //<loop> Part of loop body line 58, head labeled $Lt_26_44802
	.loc	3	67	0
	mov.s64 	%rd17, %rd21;
	bra.uni 	$Lt_26_45058;
$Lt_26_45314:
 //<loop> Part of loop body line 58, head labeled $Lt_26_44802
	.loc	3	71	0
	mov.s64 	%rd19, %rd21;
$Lt_26_45058:
 //<loop> Part of loop body line 58, head labeled $Lt_26_44802
	sub.u64 	%rd72, %rd17, %rd19;
	mov.u64 	%rd73, 1;
	setp.gt.u64 	%p33, %rd72, %rd73;
	@%p33 bra 	$Lt_26_44802;
	bra.uni 	$Lt_26_44290;
$Lt_26_50946:
	mov.u64 	%rd19, 0;
$Lt_26_44290:
	mul.lo.u64 	%rd74, %rd17, 4;
	add.u64 	%rd27, %rd12, %rd74;
	mul.lo.u64 	%rd75, %rd19, 4;
	add.u64 	%rd29, %rd12, %rd75;
	ld.global.f32 	%f154, [%rd27+600];
	ld.global.f32 	%f155, [%rd29+600];
	sub.ftz.f32 	%f156, %f154, %f155;
	mov.f32 	%f157, 0f00000000;   	// 0
	setp.neu.ftz.f32 	%p34, %f156, %f157;
	@!%p34 bra 	$Lt_26_45826;
	.loc	27	529	0
	sub.ftz.f32 	%f158, %f154, %f6;
	sub.ftz.f32 	%f159, %f6, %f155;
	div.approx.ftz.f32 	%f160, %f158, %f156;
	div.approx.ftz.f32 	%f161, %f159, %f156;
	ld.global.f32 	%f162, [%rd27+728];
	mul.ftz.f32 	%f163, %f161, %f161;
	mul.ftz.f32 	%f164, %f161, %f163;
	sub.ftz.f32 	%f165, %f164, %f161;
	mul.ftz.f32 	%f166, %f162, %f165;
	ld.global.f32 	%f167, [%rd29+728];
	mul.ftz.f32 	%f168, %f160, %f160;
	mul.ftz.f32 	%f169, %f160, %f168;
	sub.ftz.f32 	%f170, %f169, %f160;
	fma.rn.ftz.f32 	%f171, %f167, %f170, %f166;
	mul.ftz.f32 	%f172, %f156, %f156;
	mul.ftz.f32 	%f173, %f171, %f172;
	mov.f32 	%f174, 0f40c00000;   	// 6
	div.approx.ftz.f32 	%f43, %f173, %f174;
	.loc	3	80	0
	ld.global.f32 	%f175, [%rd27+664];
	mul.ftz.f32 	%f176, %f175, %f161;
	ld.global.f32 	%f177, [%rd29+664];
	fma.rn.ftz.f32 	%f178, %f177, %f160, %f176;
	add.ftz.f32 	%f20, %f43, %f178;
$Lt_26_45826:
	.loc	22	56	0
	mov.f32 	%f148, %f20;
$Lt_26_43778:
$Lt_26_43266:
	mov.f32 	%f179, 0f00000000;   	// 0
	setp.lt.ftz.f32 	%p35, %f148, %f179;
	@!%p35 bra 	$Lt_26_46594;
	.loc	22	57	0
	ld.global.f32 	%f180, [sSpline+64];
	add.ftz.f32 	%f181, %f180, %f148;
	bra.uni 	$Lt_26_46850;
$Lt_26_46594:
	ld.global.u64 	%rd30, [sSpline+192];
	mov.f32 	%f182, 0f3f800000;   	// 1
	setp.gt.ftz.f32 	%p36, %f148, %f182;
	@!%p36 bra 	$Lt_26_47106;
	mov.f32 	%f183, 0f3f800000;   	// 1
	mul.lo.u64 	%rd76, %rd30, 4;
	add.u64 	%rd77, %rd12, %rd76;
	ld.global.f32 	%f184, [%rd77+60];
	sub.ftz.f32 	%f185, %f183, %f184;
	sub.ftz.f32 	%f181, %f148, %f185;
	bra.uni 	$Lt_26_46850;
$Lt_26_47106:
	.loc	3	54	0
	mov.f32 	%f20, %f148;
	.loc	3	58	0
	sub.u64 	%rd33, %rd30, 1;
	mov.s64 	%rd17, %rd33;
	mov.u64 	%rd78, 1;
	setp.le.u64 	%p37, %rd33, %rd78;
	@%p37 bra 	$Lt_26_51202;
	mov.u64 	%rd19, 0;
$Lt_26_47874:
 //<loop> Loop body line 58
	add.u64 	%rd79, %rd17, %rd19;
	shr.u64 	%rd21, %rd79, 1;
	mul.lo.u64 	%rd80, %rd21, 4;
	add.u64 	%rd81, %rd12, %rd80;
	ld.global.f32 	%f186, [%rd81+0];
	setp.gt.ftz.f32 	%p38, %f186, %f148;
	@!%p38 bra 	$Lt_26_48386;
 //<loop> Part of loop body line 58, head labeled $Lt_26_47874
	.loc	3	67	0
	mov.s64 	%rd17, %rd21;
	bra.uni 	$Lt_26_48130;
$Lt_26_48386:
 //<loop> Part of loop body line 58, head labeled $Lt_26_47874
	.loc	3	71	0
	mov.s64 	%rd19, %rd21;
$Lt_26_48130:
 //<loop> Part of loop body line 58, head labeled $Lt_26_47874
	sub.u64 	%rd82, %rd17, %rd19;
	mov.u64 	%rd83, 1;
	setp.gt.u64 	%p39, %rd82, %rd83;
	@%p39 bra 	$Lt_26_47874;
	bra.uni 	$Lt_26_47362;
$Lt_26_51202:
	mov.u64 	%rd19, 0;
$Lt_26_47362:
	mul.lo.u64 	%rd84, %rd17, 4;
	add.u64 	%rd27, %rd12, %rd84;
	mul.lo.u64 	%rd85, %rd19, 4;
	add.u64 	%rd29, %rd12, %rd85;
	ld.global.f32 	%f56, [%rd27+0];
	ld.global.f32 	%f57, [%rd29+0];
	sub.ftz.f32 	%f58, %f56, %f57;
	mov.f32 	%f187, 0f00000000;   	// 0
	setp.neu.ftz.f32 	%p40, %f58, %f187;
	@!%p40 bra 	$Lt_26_48898;
	.loc	27	529	0
	sub.ftz.f32 	%f188, %f56, %f148;
	sub.ftz.f32 	%f189, %f148, %f57;
	div.approx.ftz.f32 	%f190, %f188, %f58;
	div.approx.ftz.f32 	%f191, %f189, %f58;
	ld.global.f32 	%f192, [%rd27+128];
	mul.ftz.f32 	%f193, %f191, %f191;
	mul.ftz.f32 	%f194, %f191, %f193;
	sub.ftz.f32 	%f195, %f194, %f191;
	mul.ftz.f32 	%f196, %f192, %f195;
	ld.global.f32 	%f197, [%rd29+128];
	mul.ftz.f32 	%f198, %f190, %f190;
	mul.ftz.f32 	%f199, %f190, %f198;
	sub.ftz.f32 	%f200, %f199, %f190;
	fma.rn.ftz.f32 	%f201, %f197, %f200, %f196;
	mul.ftz.f32 	%f202, %f58, %f58;
	mul.ftz.f32 	%f203, %f201, %f202;
	mov.f32 	%f204, 0f40c00000;   	// 6
	div.approx.ftz.f32 	%f43, %f203, %f204;
	.loc	3	80	0
	ld.global.f32 	%f205, [%rd27+64];
	mul.ftz.f32 	%f206, %f205, %f191;
	ld.global.f32 	%f207, [%rd29+64];
	fma.rn.ftz.f32 	%f208, %f207, %f190, %f206;
	add.ftz.f32 	%f20, %f43, %f208;
$Lt_26_48898:
	.loc	22	57	0
	mov.f32 	%f181, %f20;
$Lt_26_46850:
$Lt_26_46338:
	.loc	22	58	0
	sub.ftz.f32 	%f209, %f181, %f6;
	fma.rn.ftz.f32 	%f7, %f5, %f209, %f6;
$Lt_26_30466:
	.loc	22	111	0
	ld.const.f32 	%f210, [kRGB32f_To_601YPbPr+4];
	mul.ftz.f32 	%f211, %f210, %f9;
	ld.const.f32 	%f212, [kRGB32f_To_601YPbPr+0];
	fma.rn.ftz.f32 	%f213, %f212, %f11, %f211;
	ld.const.f32 	%f214, [kRGB32f_To_601YPbPr+8];
	fma.rn.ftz.f32 	%f215, %f214, %f7, %f213;
	@!%p2 bra 	$Lt_26_49666;
	.loc	20	126	0
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f215;
	mov.b32		%r31, %b1; }
	mov.s32 	%r32, %r31;
	mul.lo.u64 	%rd86, %rd1, 8;
	add.u64 	%rd87, %rd2, %rd86;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f4;
	mov.b32		%r33, %b1; }
	st.global.v4.u16 	[%rd87+0], {%r32,%r32,%r32,%r33};
	.loc	22	115	0
	bra.uni 	$Lt_26_49410;
$Lt_26_49666:
	.loc	20	126	0
	mul.lo.u64 	%rd88, %rd1, 16;
	add.u64 	%rd89, %rd2, %rd88;
	st.global.v4.f32 	[%rd89+0], {%f215,%f215,%f215,%f4};
$Lt_26_49410:
$Lt_26_28930:
	.loc	22	117	0
	exit;
$LDWend_RGBCurves_LumaKernel:
	} // RGBCurves_LumaKernel

	.entry RGBCurves_CompositeKernel (
		.param .u64 __cudaparm_RGBCurves_CompositeKernel_inImage,
		.param .s32 __cudaparm_RGBCurves_CompositeKernel_inPitch,
		.param .u64 __cudaparm_RGBCurves_CompositeKernel_inSecondaryMask,
		.param .s32 __cudaparm_RGBCurves_CompositeKernel_inSecondaryPitch,
		.param .u32 __cudaparm_RGBCurves_CompositeKernel_inDeviceFormat,
		.param .s32 __cudaparm_RGBCurves_CompositeKernel_inWidth,
		.param .s32 __cudaparm_RGBCurves_CompositeKernel_inHeight)
	{
	.reg .u32 %r<36>;
	.reg .u64 %rd<91>;
	.reg .f32 %f<211>;
	.reg .pred %p<42>;
	.loc	22	129	0
$LDWbegin_RGBCurves_CompositeKernel:
	.loc	22	132	0
	cvt.s32.u32 	%r1, %ctaid.x;
	cvt.s32.u32 	%r2, %ntid.x;
	mul.lo.s32 	%r3, %r1, %r2;
	cvt.s32.u32 	%r4, %ctaid.y;
	cvt.s32.u32 	%r5, %ntid.y;
	mul.lo.s32 	%r6, %r4, %r5;
	mov.u32 	%r7, %tid.x;
	add.u32 	%r8, %r3, %r7;
	mov.u32 	%r9, %tid.y;
	add.u32 	%r10, %r6, %r9;
	ld.param.s32 	%r11, [__cudaparm_RGBCurves_CompositeKernel_inWidth];
	set.gt.u32.s32 	%r12, %r11, %r8;
	neg.s32 	%r13, %r12;
	ld.param.s32 	%r14, [__cudaparm_RGBCurves_CompositeKernel_inHeight];
	set.gt.u32.s32 	%r15, %r14, %r10;
	neg.s32 	%r16, %r15;
	and.b32 	%r17, %r13, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_27_49410;
	ld.param.s32 	%r19, [__cudaparm_RGBCurves_CompositeKernel_inDeviceFormat];
	mov.s32 	%r20, 0;
	setp.eq.s32 	%p2, %r19, %r20;
	ld.param.s32 	%r21, [__cudaparm_RGBCurves_CompositeKernel_inPitch];
	mul.lo.s32 	%r22, %r21, %r10;
	add.s32 	%r23, %r8, %r22;
	cvt.s64.s32 	%rd1, %r23;
	ld.param.u64 	%rd2, [__cudaparm_RGBCurves_CompositeKernel_inImage];
	@!%p2 bra 	$Lt_27_29698;
	.loc	20	115	0
	mul.lo.u64 	%rd3, %rd1, 8;
	add.u64 	%rd4, %rd2, %rd3;
	ld.global.v4.u16 	{%r24,%r25,%r26,%r27}, [%rd4+0];
	.loc	22	135	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_27_29442;
$Lt_27_29698:
	mul.lo.u64 	%rd5, %rd1, 16;
	add.u64 	%rd6, %rd2, %rd5;
	ld.global.v4.f32 	{%f1,%f2,%f3,%f4}, [%rd6+0];
$Lt_27_29442:
	.loc	22	136	0
	ld.param.u64 	%rd7, [__cudaparm_RGBCurves_CompositeKernel_inSecondaryMask];
	mov.u64 	%rd8, 0;
	setp.eq.u64 	%p3, %rd7, %rd8;
	@%p3 bra 	$Lt_27_30210;
	.loc	22	137	0
	ld.param.s32 	%r28, [__cudaparm_RGBCurves_CompositeKernel_inSecondaryPitch];
	mul.lo.s32 	%r29, %r28, %r10;
	add.s32 	%r30, %r8, %r29;
	cvt.s64.s32 	%rd9, %r30;
	mul.wide.s32 	%rd10, %r30, 4;
	add.u64 	%rd11, %rd7, %rd10;
	ld.global.f32 	%f5, [%rd11+0];
	bra.uni 	$Lt_27_29954;
$Lt_27_30210:
	mov.f32 	%f5, 0f3f800000;     	// 1
$Lt_27_29954:
	.loc	22	40	0
	mov.f32 	%f6, %f1;
	mov.f32 	%f7, %f6;
	mov.f32 	%f8, %f2;
	mov.f32 	%f9, %f8;
	mov.f32 	%f10, %f3;
	mov.f32 	%f11, %f10;
	.loc	22	41	0
	mov.f32 	%f12, 0f3a83126f;    	// 0.001
	setp.gt.ftz.f32 	%p4, %f5, %f12;
	@!%p4 bra 	$Lt_27_30466;
	mov.f32 	%f13, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p5, %f10, %f13;
	@!%p5 bra 	$Lt_27_31234;
	.loc	22	46	0
	ld.global.f32 	%f14, [sSpline+264];
	add.ftz.f32 	%f15, %f14, %f10;
	mov.u64 	%rd12, sSpline;
	bra.uni 	$Lt_27_31490;
$Lt_27_31234:
	ld.global.u64 	%rd13, [sSpline+392];
	mov.f32 	%f16, 0f3f800000;    	// 1
	setp.gt.ftz.f32 	%p6, %f10, %f16;
	@!%p6 bra 	$Lt_27_31746;
	mov.u64 	%rd12, sSpline;
	mov.f32 	%f17, 0f3f800000;    	// 1
	mul.lo.u64 	%rd14, %rd13, 4;
	add.u64 	%rd15, %rd12, %rd14;
	ld.global.f32 	%f18, [%rd15+260];
	sub.ftz.f32 	%f19, %f17, %f18;
	sub.ftz.f32 	%f15, %f10, %f19;
	bra.uni 	$Lt_27_31490;
$Lt_27_31746:
	.loc	3	54	0
	mov.f32 	%f20, %f10;
	.loc	3	58	0
	sub.u64 	%rd16, %rd13, 1;
	mov.s64 	%rd17, %rd16;
	mov.u64 	%rd18, 1;
	setp.le.u64 	%p7, %rd16, %rd18;
	@%p7 bra 	$Lt_27_49922;
	mov.u64 	%rd19, 0;
	mov.u64 	%rd12, sSpline;
$Lt_27_32514:
 //<loop> Loop body line 58
	add.u64 	%rd20, %rd17, %rd19;
	shr.u64 	%rd21, %rd20, 1;
	mul.lo.u64 	%rd22, %rd21, 4;
	add.u64 	%rd23, %rd12, %rd22;
	ld.global.f32 	%f21, [%rd23+200];
	setp.gt.ftz.f32 	%p8, %f21, %f10;
	@!%p8 bra 	$Lt_27_33026;
 //<loop> Part of loop body line 58, head labeled $Lt_27_32514
	.loc	3	67	0
	mov.s64 	%rd17, %rd21;
	bra.uni 	$Lt_27_32770;
$Lt_27_33026:
 //<loop> Part of loop body line 58, head labeled $Lt_27_32514
	.loc	3	71	0
	mov.s64 	%rd19, %rd21;
$Lt_27_32770:
 //<loop> Part of loop body line 58, head labeled $Lt_27_32514
	sub.u64 	%rd24, %rd17, %rd19;
	mov.u64 	%rd25, 1;
	setp.gt.u64 	%p9, %rd24, %rd25;
	@%p9 bra 	$Lt_27_32514;
	bra.uni 	$Lt_27_32002;
$Lt_27_49922:
	mov.u64 	%rd19, 0;
	mov.u64 	%rd12, sSpline;
$Lt_27_32002:
	mul.lo.u64 	%rd26, %rd17, 4;
	add.u64 	%rd27, %rd12, %rd26;
	mul.lo.u64 	%rd28, %rd19, 4;
	add.u64 	%rd29, %rd12, %rd28;
	ld.global.f32 	%f22, [%rd27+200];
	ld.global.f32 	%f23, [%rd29+200];
	sub.ftz.f32 	%f24, %f22, %f23;
	mov.f32 	%f25, 0f00000000;    	// 0
	setp.neu.ftz.f32 	%p10, %f24, %f25;
	@!%p10 bra 	$Lt_27_33538;
	.loc	27	529	0
	sub.ftz.f32 	%f26, %f22, %f10;
	sub.ftz.f32 	%f27, %f10, %f23;
	div.approx.ftz.f32 	%f28, %f26, %f24;
	div.approx.ftz.f32 	%f29, %f27, %f24;
	ld.global.f32 	%f30, [%rd27+328];
	mul.ftz.f32 	%f31, %f29, %f29;
	mul.ftz.f32 	%f32, %f29, %f31;
	sub.ftz.f32 	%f33, %f32, %f29;
	mul.ftz.f32 	%f34, %f30, %f33;
	ld.global.f32 	%f35, [%rd29+328];
	mul.ftz.f32 	%f36, %f28, %f28;
	mul.ftz.f32 	%f37, %f28, %f36;
	sub.ftz.f32 	%f38, %f37, %f28;
	fma.rn.ftz.f32 	%f39, %f35, %f38, %f34;
	mul.ftz.f32 	%f40, %f24, %f24;
	mul.ftz.f32 	%f41, %f39, %f40;
	mov.f32 	%f42, 0f40c00000;    	// 6
	div.approx.ftz.f32 	%f43, %f41, %f42;
	.loc	3	80	0
	ld.global.f32 	%f44, [%rd27+264];
	mul.ftz.f32 	%f45, %f44, %f29;
	ld.global.f32 	%f46, [%rd29+264];
	fma.rn.ftz.f32 	%f47, %f46, %f28, %f45;
	add.ftz.f32 	%f20, %f43, %f47;
$Lt_27_33538:
	.loc	22	46	0
	mov.f32 	%f15, %f20;
$Lt_27_31490:
$Lt_27_30978:
	mov.f32 	%f48, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p11, %f15, %f48;
	@!%p11 bra 	$Lt_27_34306;
	.loc	22	47	0
	ld.global.f32 	%f49, [sSpline+64];
	add.ftz.f32 	%f50, %f49, %f15;
	bra.uni 	$Lt_27_34562;
$Lt_27_34306:
	ld.global.u64 	%rd30, [sSpline+192];
	mov.f32 	%f51, 0f3f800000;    	// 1
	setp.gt.ftz.f32 	%p12, %f15, %f51;
	@!%p12 bra 	$Lt_27_34818;
	mov.f32 	%f52, 0f3f800000;    	// 1
	mul.lo.u64 	%rd31, %rd30, 4;
	add.u64 	%rd32, %rd12, %rd31;
	ld.global.f32 	%f53, [%rd32+60];
	sub.ftz.f32 	%f54, %f52, %f53;
	sub.ftz.f32 	%f50, %f15, %f54;
	bra.uni 	$Lt_27_34562;
$Lt_27_34818:
	.loc	3	54	0
	mov.f32 	%f20, %f15;
	.loc	3	58	0
	sub.u64 	%rd33, %rd30, 1;
	mov.s64 	%rd17, %rd33;
	mov.u64 	%rd34, 1;
	setp.le.u64 	%p13, %rd33, %rd34;
	@%p13 bra 	$Lt_27_50178;
	mov.u64 	%rd19, 0;
$Lt_27_35586:
 //<loop> Loop body line 58
	add.u64 	%rd35, %rd17, %rd19;
	shr.u64 	%rd21, %rd35, 1;
	mul.lo.u64 	%rd36, %rd21, 4;
	add.u64 	%rd37, %rd12, %rd36;
	ld.global.f32 	%f55, [%rd37+0];
	setp.gt.ftz.f32 	%p14, %f55, %f15;
	@!%p14 bra 	$Lt_27_36098;
 //<loop> Part of loop body line 58, head labeled $Lt_27_35586
	.loc	3	67	0
	mov.s64 	%rd17, %rd21;
	bra.uni 	$Lt_27_35842;
$Lt_27_36098:
 //<loop> Part of loop body line 58, head labeled $Lt_27_35586
	.loc	3	71	0
	mov.s64 	%rd19, %rd21;
$Lt_27_35842:
 //<loop> Part of loop body line 58, head labeled $Lt_27_35586
	sub.u64 	%rd38, %rd17, %rd19;
	mov.u64 	%rd39, 1;
	setp.gt.u64 	%p15, %rd38, %rd39;
	@%p15 bra 	$Lt_27_35586;
	bra.uni 	$Lt_27_35074;
$Lt_27_50178:
	mov.u64 	%rd19, 0;
$Lt_27_35074:
	mul.lo.u64 	%rd40, %rd17, 4;
	add.u64 	%rd27, %rd12, %rd40;
	mul.lo.u64 	%rd41, %rd19, 4;
	add.u64 	%rd29, %rd12, %rd41;
	ld.global.f32 	%f56, [%rd27+0];
	ld.global.f32 	%f57, [%rd29+0];
	sub.ftz.f32 	%f58, %f56, %f57;
	mov.f32 	%f59, 0f00000000;    	// 0
	setp.neu.ftz.f32 	%p16, %f58, %f59;
	@!%p16 bra 	$Lt_27_36610;
	.loc	27	529	0
	sub.ftz.f32 	%f60, %f56, %f15;
	sub.ftz.f32 	%f61, %f15, %f57;
	div.approx.ftz.f32 	%f62, %f60, %f58;
	div.approx.ftz.f32 	%f63, %f61, %f58;
	ld.global.f32 	%f64, [%rd27+128];
	mul.ftz.f32 	%f65, %f63, %f63;
	mul.ftz.f32 	%f66, %f63, %f65;
	sub.ftz.f32 	%f67, %f66, %f63;
	mul.ftz.f32 	%f68, %f64, %f67;
	ld.global.f32 	%f69, [%rd29+128];
	mul.ftz.f32 	%f70, %f62, %f62;
	mul.ftz.f32 	%f71, %f62, %f70;
	sub.ftz.f32 	%f72, %f71, %f62;
	fma.rn.ftz.f32 	%f73, %f69, %f72, %f68;
	mul.ftz.f32 	%f74, %f58, %f58;
	mul.ftz.f32 	%f75, %f73, %f74;
	mov.f32 	%f76, 0f40c00000;    	// 6
	div.approx.ftz.f32 	%f43, %f75, %f76;
	.loc	3	80	0
	ld.global.f32 	%f77, [%rd27+64];
	mul.ftz.f32 	%f78, %f77, %f63;
	ld.global.f32 	%f79, [%rd29+64];
	fma.rn.ftz.f32 	%f80, %f79, %f62, %f78;
	add.ftz.f32 	%f20, %f43, %f80;
$Lt_27_36610:
	.loc	22	47	0
	mov.f32 	%f50, %f20;
$Lt_27_34562:
$Lt_27_34050:
	.loc	22	48	0
	sub.ftz.f32 	%f81, %f50, %f10;
	fma.rn.ftz.f32 	%f11, %f5, %f81, %f10;
	mov.f32 	%f82, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p17, %f8, %f82;
	@!%p17 bra 	$Lt_27_37378;
	.loc	22	51	0
	ld.global.f32 	%f83, [sSpline+464];
	add.ftz.f32 	%f84, %f83, %f8;
	bra.uni 	$Lt_27_37634;
$Lt_27_37378:
	ld.global.u64 	%rd42, [sSpline+592];
	mov.f32 	%f85, 0f3f800000;    	// 1
	setp.gt.ftz.f32 	%p18, %f8, %f85;
	@!%p18 bra 	$Lt_27_37890;
	mov.f32 	%f86, 0f3f800000;    	// 1
	mul.lo.u64 	%rd43, %rd42, 4;
	add.u64 	%rd44, %rd12, %rd43;
	ld.global.f32 	%f87, [%rd44+460];
	sub.ftz.f32 	%f88, %f86, %f87;
	sub.ftz.f32 	%f84, %f8, %f88;
	bra.uni 	$Lt_27_37634;
$Lt_27_37890:
	.loc	3	54	0
	mov.f32 	%f20, %f8;
	.loc	3	58	0
	sub.u64 	%rd45, %rd42, 1;
	mov.s64 	%rd17, %rd45;
	mov.u64 	%rd46, 1;
	setp.le.u64 	%p19, %rd45, %rd46;
	@%p19 bra 	$Lt_27_50434;
	mov.u64 	%rd19, 0;
$Lt_27_38658:
 //<loop> Loop body line 58
	add.u64 	%rd47, %rd17, %rd19;
	shr.u64 	%rd21, %rd47, 1;
	mul.lo.u64 	%rd48, %rd21, 4;
	add.u64 	%rd49, %rd12, %rd48;
	ld.global.f32 	%f89, [%rd49+400];
	setp.gt.ftz.f32 	%p20, %f89, %f8;
	@!%p20 bra 	$Lt_27_39170;
 //<loop> Part of loop body line 58, head labeled $Lt_27_38658
	.loc	3	67	0
	mov.s64 	%rd17, %rd21;
	bra.uni 	$Lt_27_38914;
$Lt_27_39170:
 //<loop> Part of loop body line 58, head labeled $Lt_27_38658
	.loc	3	71	0
	mov.s64 	%rd19, %rd21;
$Lt_27_38914:
 //<loop> Part of loop body line 58, head labeled $Lt_27_38658
	sub.u64 	%rd50, %rd17, %rd19;
	mov.u64 	%rd51, 1;
	setp.gt.u64 	%p21, %rd50, %rd51;
	@%p21 bra 	$Lt_27_38658;
	bra.uni 	$Lt_27_38146;
$Lt_27_50434:
	mov.u64 	%rd19, 0;
$Lt_27_38146:
	mul.lo.u64 	%rd52, %rd17, 4;
	add.u64 	%rd27, %rd12, %rd52;
	mul.lo.u64 	%rd53, %rd19, 4;
	add.u64 	%rd29, %rd12, %rd53;
	ld.global.f32 	%f90, [%rd27+400];
	ld.global.f32 	%f91, [%rd29+400];
	sub.ftz.f32 	%f92, %f90, %f91;
	mov.f32 	%f93, 0f00000000;    	// 0
	setp.neu.ftz.f32 	%p22, %f92, %f93;
	@!%p22 bra 	$Lt_27_39682;
	.loc	27	529	0
	sub.ftz.f32 	%f94, %f90, %f8;
	sub.ftz.f32 	%f95, %f8, %f91;
	div.approx.ftz.f32 	%f96, %f94, %f92;
	div.approx.ftz.f32 	%f97, %f95, %f92;
	ld.global.f32 	%f98, [%rd27+528];
	mul.ftz.f32 	%f99, %f97, %f97;
	mul.ftz.f32 	%f100, %f97, %f99;
	sub.ftz.f32 	%f101, %f100, %f97;
	mul.ftz.f32 	%f102, %f98, %f101;
	ld.global.f32 	%f103, [%rd29+528];
	mul.ftz.f32 	%f104, %f96, %f96;
	mul.ftz.f32 	%f105, %f96, %f104;
	sub.ftz.f32 	%f106, %f105, %f96;
	fma.rn.ftz.f32 	%f107, %f103, %f106, %f102;
	mul.ftz.f32 	%f108, %f92, %f92;
	mul.ftz.f32 	%f109, %f107, %f108;
	mov.f32 	%f110, 0f40c00000;   	// 6
	div.approx.ftz.f32 	%f43, %f109, %f110;
	.loc	3	80	0
	ld.global.f32 	%f111, [%rd27+464];
	mul.ftz.f32 	%f112, %f111, %f97;
	ld.global.f32 	%f113, [%rd29+464];
	fma.rn.ftz.f32 	%f114, %f113, %f96, %f112;
	add.ftz.f32 	%f20, %f43, %f114;
$Lt_27_39682:
	.loc	22	51	0
	mov.f32 	%f84, %f20;
$Lt_27_37634:
$Lt_27_37122:
	mov.f32 	%f115, 0f00000000;   	// 0
	setp.lt.ftz.f32 	%p23, %f84, %f115;
	@!%p23 bra 	$Lt_27_40450;
	.loc	22	52	0
	ld.global.f32 	%f116, [sSpline+64];
	add.ftz.f32 	%f117, %f116, %f84;
	bra.uni 	$Lt_27_40706;
$Lt_27_40450:
	ld.global.u64 	%rd30, [sSpline+192];
	mov.f32 	%f118, 0f3f800000;   	// 1
	setp.gt.ftz.f32 	%p24, %f84, %f118;
	@!%p24 bra 	$Lt_27_40962;
	mov.f32 	%f119, 0f3f800000;   	// 1
	mul.lo.u64 	%rd54, %rd30, 4;
	add.u64 	%rd55, %rd12, %rd54;
	ld.global.f32 	%f120, [%rd55+60];
	sub.ftz.f32 	%f121, %f119, %f120;
	sub.ftz.f32 	%f117, %f84, %f121;
	bra.uni 	$Lt_27_40706;
$Lt_27_40962:
	.loc	3	54	0
	mov.f32 	%f20, %f84;
	.loc	3	58	0
	sub.u64 	%rd33, %rd30, 1;
	mov.s64 	%rd17, %rd33;
	mov.u64 	%rd56, 1;
	setp.le.u64 	%p25, %rd33, %rd56;
	@%p25 bra 	$Lt_27_50690;
	mov.u64 	%rd19, 0;
$Lt_27_41730:
 //<loop> Loop body line 58
	add.u64 	%rd57, %rd17, %rd19;
	shr.u64 	%rd21, %rd57, 1;
	mul.lo.u64 	%rd58, %rd21, 4;
	add.u64 	%rd59, %rd12, %rd58;
	ld.global.f32 	%f122, [%rd59+0];
	setp.gt.ftz.f32 	%p26, %f122, %f84;
	@!%p26 bra 	$Lt_27_42242;
 //<loop> Part of loop body line 58, head labeled $Lt_27_41730
	.loc	3	67	0
	mov.s64 	%rd17, %rd21;
	bra.uni 	$Lt_27_41986;
$Lt_27_42242:
 //<loop> Part of loop body line 58, head labeled $Lt_27_41730
	.loc	3	71	0
	mov.s64 	%rd19, %rd21;
$Lt_27_41986:
 //<loop> Part of loop body line 58, head labeled $Lt_27_41730
	sub.u64 	%rd60, %rd17, %rd19;
	mov.u64 	%rd61, 1;
	setp.gt.u64 	%p27, %rd60, %rd61;
	@%p27 bra 	$Lt_27_41730;
	bra.uni 	$Lt_27_41218;
$Lt_27_50690:
	mov.u64 	%rd19, 0;
$Lt_27_41218:
	mul.lo.u64 	%rd62, %rd17, 4;
	add.u64 	%rd27, %rd12, %rd62;
	mul.lo.u64 	%rd63, %rd19, 4;
	add.u64 	%rd29, %rd12, %rd63;
	ld.global.f32 	%f56, [%rd27+0];
	ld.global.f32 	%f57, [%rd29+0];
	sub.ftz.f32 	%f58, %f56, %f57;
	mov.f32 	%f123, 0f00000000;   	// 0
	setp.neu.ftz.f32 	%p28, %f58, %f123;
	@!%p28 bra 	$Lt_27_42754;
	.loc	27	529	0
	sub.ftz.f32 	%f124, %f56, %f84;
	sub.ftz.f32 	%f125, %f84, %f57;
	div.approx.ftz.f32 	%f126, %f124, %f58;
	div.approx.ftz.f32 	%f127, %f125, %f58;
	ld.global.f32 	%f128, [%rd27+128];
	mul.ftz.f32 	%f129, %f127, %f127;
	mul.ftz.f32 	%f130, %f127, %f129;
	sub.ftz.f32 	%f131, %f130, %f127;
	mul.ftz.f32 	%f132, %f128, %f131;
	ld.global.f32 	%f133, [%rd29+128];
	mul.ftz.f32 	%f134, %f126, %f126;
	mul.ftz.f32 	%f135, %f126, %f134;
	sub.ftz.f32 	%f136, %f135, %f126;
	fma.rn.ftz.f32 	%f137, %f133, %f136, %f132;
	mul.ftz.f32 	%f138, %f58, %f58;
	mul.ftz.f32 	%f139, %f137, %f138;
	mov.f32 	%f140, 0f40c00000;   	// 6
	div.approx.ftz.f32 	%f43, %f139, %f140;
	.loc	3	80	0
	ld.global.f32 	%f141, [%rd27+64];
	mul.ftz.f32 	%f142, %f141, %f127;
	ld.global.f32 	%f143, [%rd29+64];
	fma.rn.ftz.f32 	%f144, %f143, %f126, %f142;
	add.ftz.f32 	%f20, %f43, %f144;
$Lt_27_42754:
	.loc	22	52	0
	mov.f32 	%f117, %f20;
$Lt_27_40706:
$Lt_27_40194:
	.loc	22	53	0
	sub.ftz.f32 	%f145, %f117, %f8;
	fma.rn.ftz.f32 	%f9, %f5, %f145, %f8;
	mov.f32 	%f146, 0f00000000;   	// 0
	setp.lt.ftz.f32 	%p29, %f6, %f146;
	@!%p29 bra 	$Lt_27_43522;
	.loc	22	56	0
	ld.global.f32 	%f147, [sSpline+664];
	add.ftz.f32 	%f148, %f6, %f147;
	bra.uni 	$Lt_27_43778;
$Lt_27_43522:
	ld.global.u64 	%rd64, [sSpline+792];
	mov.f32 	%f149, 0f3f800000;   	// 1
	setp.gt.ftz.f32 	%p30, %f6, %f149;
	@!%p30 bra 	$Lt_27_44034;
	mov.f32 	%f150, 0f3f800000;   	// 1
	mul.lo.u64 	%rd65, %rd64, 4;
	add.u64 	%rd66, %rd12, %rd65;
	ld.global.f32 	%f151, [%rd66+660];
	sub.ftz.f32 	%f152, %f150, %f151;
	sub.ftz.f32 	%f148, %f6, %f152;
	bra.uni 	$Lt_27_43778;
$Lt_27_44034:
	.loc	3	54	0
	mov.f32 	%f20, %f6;
	.loc	3	58	0
	sub.u64 	%rd67, %rd64, 1;
	mov.s64 	%rd17, %rd67;
	mov.u64 	%rd68, 1;
	setp.le.u64 	%p31, %rd67, %rd68;
	@%p31 bra 	$Lt_27_50946;
	mov.u64 	%rd19, 0;
$Lt_27_44802:
 //<loop> Loop body line 58
	add.u64 	%rd69, %rd17, %rd19;
	shr.u64 	%rd21, %rd69, 1;
	mul.lo.u64 	%rd70, %rd21, 4;
	add.u64 	%rd71, %rd12, %rd70;
	ld.global.f32 	%f153, [%rd71+600];
	setp.gt.ftz.f32 	%p32, %f153, %f6;
	@!%p32 bra 	$Lt_27_45314;
 //<loop> Part of loop body line 58, head labeled $Lt_27_44802
	.loc	3	67	0
	mov.s64 	%rd17, %rd21;
	bra.uni 	$Lt_27_45058;
$Lt_27_45314:
 //<loop> Part of loop body line 58, head labeled $Lt_27_44802
	.loc	3	71	0
	mov.s64 	%rd19, %rd21;
$Lt_27_45058:
 //<loop> Part of loop body line 58, head labeled $Lt_27_44802
	sub.u64 	%rd72, %rd17, %rd19;
	mov.u64 	%rd73, 1;
	setp.gt.u64 	%p33, %rd72, %rd73;
	@%p33 bra 	$Lt_27_44802;
	bra.uni 	$Lt_27_44290;
$Lt_27_50946:
	mov.u64 	%rd19, 0;
$Lt_27_44290:
	mul.lo.u64 	%rd74, %rd17, 4;
	add.u64 	%rd27, %rd12, %rd74;
	mul.lo.u64 	%rd75, %rd19, 4;
	add.u64 	%rd29, %rd12, %rd75;
	ld.global.f32 	%f154, [%rd27+600];
	ld.global.f32 	%f155, [%rd29+600];
	sub.ftz.f32 	%f156, %f154, %f155;
	mov.f32 	%f157, 0f00000000;   	// 0
	setp.neu.ftz.f32 	%p34, %f156, %f157;
	@!%p34 bra 	$Lt_27_45826;
	.loc	27	529	0
	sub.ftz.f32 	%f158, %f154, %f6;
	sub.ftz.f32 	%f159, %f6, %f155;
	div.approx.ftz.f32 	%f160, %f158, %f156;
	div.approx.ftz.f32 	%f161, %f159, %f156;
	ld.global.f32 	%f162, [%rd27+728];
	mul.ftz.f32 	%f163, %f161, %f161;
	mul.ftz.f32 	%f164, %f161, %f163;
	sub.ftz.f32 	%f165, %f164, %f161;
	mul.ftz.f32 	%f166, %f162, %f165;
	ld.global.f32 	%f167, [%rd29+728];
	mul.ftz.f32 	%f168, %f160, %f160;
	mul.ftz.f32 	%f169, %f160, %f168;
	sub.ftz.f32 	%f170, %f169, %f160;
	fma.rn.ftz.f32 	%f171, %f167, %f170, %f166;
	mul.ftz.f32 	%f172, %f156, %f156;
	mul.ftz.f32 	%f173, %f171, %f172;
	mov.f32 	%f174, 0f40c00000;   	// 6
	div.approx.ftz.f32 	%f43, %f173, %f174;
	.loc	3	80	0
	ld.global.f32 	%f175, [%rd27+664];
	mul.ftz.f32 	%f176, %f175, %f161;
	ld.global.f32 	%f177, [%rd29+664];
	fma.rn.ftz.f32 	%f178, %f177, %f160, %f176;
	add.ftz.f32 	%f20, %f43, %f178;
$Lt_27_45826:
	.loc	22	56	0
	mov.f32 	%f148, %f20;
$Lt_27_43778:
$Lt_27_43266:
	mov.f32 	%f179, 0f00000000;   	// 0
	setp.lt.ftz.f32 	%p35, %f148, %f179;
	@!%p35 bra 	$Lt_27_46594;
	.loc	22	57	0
	ld.global.f32 	%f180, [sSpline+64];
	add.ftz.f32 	%f181, %f180, %f148;
	bra.uni 	$Lt_27_46850;
$Lt_27_46594:
	ld.global.u64 	%rd30, [sSpline+192];
	mov.f32 	%f182, 0f3f800000;   	// 1
	setp.gt.ftz.f32 	%p36, %f148, %f182;
	@!%p36 bra 	$Lt_27_47106;
	mov.f32 	%f183, 0f3f800000;   	// 1
	mul.lo.u64 	%rd76, %rd30, 4;
	add.u64 	%rd77, %rd12, %rd76;
	ld.global.f32 	%f184, [%rd77+60];
	sub.ftz.f32 	%f185, %f183, %f184;
	sub.ftz.f32 	%f181, %f148, %f185;
	bra.uni 	$Lt_27_46850;
$Lt_27_47106:
	.loc	3	54	0
	mov.f32 	%f20, %f148;
	.loc	3	58	0
	sub.u64 	%rd33, %rd30, 1;
	mov.s64 	%rd17, %rd33;
	mov.u64 	%rd78, 1;
	setp.le.u64 	%p37, %rd33, %rd78;
	@%p37 bra 	$Lt_27_51202;
	mov.u64 	%rd19, 0;
$Lt_27_47874:
 //<loop> Loop body line 58
	add.u64 	%rd79, %rd17, %rd19;
	shr.u64 	%rd21, %rd79, 1;
	mul.lo.u64 	%rd80, %rd21, 4;
	add.u64 	%rd81, %rd12, %rd80;
	ld.global.f32 	%f186, [%rd81+0];
	setp.gt.ftz.f32 	%p38, %f186, %f148;
	@!%p38 bra 	$Lt_27_48386;
 //<loop> Part of loop body line 58, head labeled $Lt_27_47874
	.loc	3	67	0
	mov.s64 	%rd17, %rd21;
	bra.uni 	$Lt_27_48130;
$Lt_27_48386:
 //<loop> Part of loop body line 58, head labeled $Lt_27_47874
	.loc	3	71	0
	mov.s64 	%rd19, %rd21;
$Lt_27_48130:
 //<loop> Part of loop body line 58, head labeled $Lt_27_47874
	sub.u64 	%rd82, %rd17, %rd19;
	mov.u64 	%rd83, 1;
	setp.gt.u64 	%p39, %rd82, %rd83;
	@%p39 bra 	$Lt_27_47874;
	bra.uni 	$Lt_27_47362;
$Lt_27_51202:
	mov.u64 	%rd19, 0;
$Lt_27_47362:
	mul.lo.u64 	%rd84, %rd17, 4;
	add.u64 	%rd27, %rd12, %rd84;
	mul.lo.u64 	%rd85, %rd19, 4;
	add.u64 	%rd29, %rd12, %rd85;
	ld.global.f32 	%f56, [%rd27+0];
	ld.global.f32 	%f57, [%rd29+0];
	sub.ftz.f32 	%f58, %f56, %f57;
	mov.f32 	%f187, 0f00000000;   	// 0
	setp.neu.ftz.f32 	%p40, %f58, %f187;
	@!%p40 bra 	$Lt_27_48898;
	.loc	27	529	0
	sub.ftz.f32 	%f188, %f56, %f148;
	sub.ftz.f32 	%f189, %f148, %f57;
	div.approx.ftz.f32 	%f190, %f188, %f58;
	div.approx.ftz.f32 	%f191, %f189, %f58;
	ld.global.f32 	%f192, [%rd27+128];
	mul.ftz.f32 	%f193, %f191, %f191;
	mul.ftz.f32 	%f194, %f191, %f193;
	sub.ftz.f32 	%f195, %f194, %f191;
	mul.ftz.f32 	%f196, %f192, %f195;
	ld.global.f32 	%f197, [%rd29+128];
	mul.ftz.f32 	%f198, %f190, %f190;
	mul.ftz.f32 	%f199, %f190, %f198;
	sub.ftz.f32 	%f200, %f199, %f190;
	fma.rn.ftz.f32 	%f201, %f197, %f200, %f196;
	mul.ftz.f32 	%f202, %f58, %f58;
	mul.ftz.f32 	%f203, %f201, %f202;
	mov.f32 	%f204, 0f40c00000;   	// 6
	div.approx.ftz.f32 	%f43, %f203, %f204;
	.loc	3	80	0
	ld.global.f32 	%f205, [%rd27+64];
	mul.ftz.f32 	%f206, %f205, %f191;
	ld.global.f32 	%f207, [%rd29+64];
	fma.rn.ftz.f32 	%f208, %f207, %f190, %f206;
	add.ftz.f32 	%f20, %f43, %f208;
$Lt_27_48898:
	.loc	22	57	0
	mov.f32 	%f181, %f20;
$Lt_27_46850:
$Lt_27_46338:
	.loc	22	58	0
	sub.ftz.f32 	%f209, %f181, %f6;
	fma.rn.ftz.f32 	%f7, %f5, %f209, %f6;
$Lt_27_30466:
	.loc	22	139	0
	@!%p2 bra 	$Lt_27_49666;
	.loc	20	126	0
	mul.lo.u64 	%rd86, %rd1, 8;
	add.u64 	%rd87, %rd2, %rd86;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f7;
	mov.b32		%r31, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f9;
	mov.b32		%r32, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f11;
	mov.b32		%r33, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f4;
	mov.b32		%r34, %b1; }
	st.global.v4.u16 	[%rd87+0], {%r31,%r32,%r33,%r34};
	.loc	22	139	0
	bra.uni 	$Lt_27_49410;
$Lt_27_49666:
	.loc	20	126	0
	mul.lo.u64 	%rd88, %rd1, 16;
	add.u64 	%rd89, %rd2, %rd88;
	st.global.v4.f32 	[%rd89+0], {%f7,%f9,%f11,%f4};
$Lt_27_49410:
$Lt_27_28930:
	.loc	22	147	0
	exit;
$LDWend_RGBCurves_CompositeKernel:
	} // RGBCurves_CompositeKernel
	.const .align 4 .b8 k601YPbPr_To_RGB32f[36] = {0,0,128,63,0,0,0,0,188,116,179,63,0,0,128,63,152,50,176,190,158,209,54,191,0,0,128,63,229,208,226,63,0,0,0,0};
	.const .align 4 .b8 kRGB32f_To_601YCbCr[36] = {70,246,130,66,145,141,0,67,94,186,199,65,33,48,23,194,240,103,148,194,0,0,224,66,0,0,224,66,111,146,187,194,70,182,145,193};
	.const .align 4 .b8 k601YCbCr_To_RGB32f[36] = {37,160,149,59,0,0,0,0,182,23,205,59,37,160,149,59,40,15,201,186,156,239,80,187,37,160,149,59,236,155,1,60,0,0,0,0};
	.const .align 4 .b8 kRGB8u_To_601YCbCr[36] = {219,121,131,62,152,14,1,63,18,131,200,61,174,199,23,190,238,252,148,190,197,224,224,62,197,224,224,62,217,78,188,190,174,71,146,189};
	.const .align 4 .b8 k601YCbCr_To_RGB8u[36] = {127,10,149,63,0,0,0,0,160,74,204,63,127,10,149,63,254,148,200,190,184,30,80,191,127,10,149,63,78,26,1,64,0,0,0,0};
	.const .align 4 .b8 kRGB8u_To_601YCbCrFullRange[36] = {135,22,153,62,162,69,22,63,213,120,233,61,166,27,44,190,39,241,168,190,250,254,254,62,250,254,254,62,43,135,213,190,59,223,165,189};
	.const .align 4 .b8 k601YCbCrFullRange_To_RGB8u[36] = {0,0,128,63,0,0,0,0,72,193,178,63,0,0,128,63,143,130,175,190,225,26,54,191,0,0,128,63,20,238,225,63,0,0,0,0};
	.const .align 4 .b8 kRGB32f_To_601YCbCrFullRange[36] = {113,125,152,66,92,175,21,67,92,143,232,65,158,111,43,194,49,72,168,194,0,0,254,66,0,0,254,66,170,177,212,194,88,57,165,193};
	.const .align 4 .b8 k601YCbCrFullRange_To_RGB32f[36] = {129,128,128,59,0,0,0,0,188,116,179,59,129,128,128,59,194,50,176,186,179,209,54,187,129,128,128,59,229,208,226,59,0,0,0,0};
	.const .align 4 .b8 kRGB32f_To_709YPbPr[36] = {208,179,89,62,89,23,55,63,152,221,147,61,186,164,234,189,210,86,197,190,0,0,0,63,0,0,0,63,190,134,232,190,16,202,59,189};
	.const .align 4 .b8 k709YPbPr_To_RGB32f[36] = {0,0,128,63,0,0,0,0,12,147,201,63,0,0,128,63,221,209,63,190,243,173,239,190,0,0,128,63,77,132,237,63,0,0,0,0};
	.const .align 4 .b8 kRGB32f_To_709YCbCr[36] = {106,60,58,66,6,161,28,67,244,253,124,65,223,79,205,193,8,172,172,194,0,0,224,66,0,0,224,66,195,117,203,194,236,81,36,193};
	.const .align 4 .b8 k709YCbCr_To_RGB32f[36] = {37,160,149,59,0,0,0,0,239,94,230,59,37,160,149,59,33,57,91,186,178,245,8,187,37,160,149,59,82,185,7,60,0,0,0,0};
	.const .align 4 .b8 kRGB8u_To_709YCbCr[36] = {207,247,58,62,53,62,29,63,231,251,125,61,147,24,206,61,23,89,173,190,197,224,224,62,197,224,224,62,12,66,204,190,195,245,36,189};
	.const .align 4 .b8 k709YCbCr_To_RGB8u[36] = {127,10,149,63,0,0,0,0,147,120,229,63,127,10,149,63,53,94,90,190,205,108,8,191,127,10,149,63,154,49,7,64,0,0,0,0};
	.const .align 4 .b8 k709YCbCr_To_601YCbCr[36] = {0,0,128,63,23,100,203,61,1,77,68,62,0,0,0,0,18,103,125,63,10,158,226,189,0,0,0,0,61,98,148,189,249,191,123,63};
	.const .align 4 .b8 k601YCbCr_To_709YCbCr[36] = {0,0,128,63,122,165,236,189,179,237,84,190,0,0,0,0,204,98,130,63,216,188,234,61,0,0,0,0,74,179,153,61,234,61,131,63};
	.const .align 4 .b8 kYCbCrOffset[12] = {0,0,128,65,0,0,0,67,0,0,0,67};
	.const .align 4 .b8 kYCbCrFullRangeOffset[12] = {0,0,0,0,0,0,0,67,0,0,0,67};
	.const .align 4 .b8 kRGB32f_To_YIQ[36] = {135,22,153,62,162,69,22,63,213,120,233,61,216,128,24,63,27,133,140,190,149,124,164,190,236,135,88,62,134,200,5,191,22,77,159,62};
	.const .align 4 .b8 kYIQ_To_RGB32f[36] = {0,0,128,63,20,208,116,63,219,249,30,63,0,0,128,63,177,80,139,190,2,188,37,191,0,0,128,63,45,178,141,191,85,48,218,63};

