1005 lines
30 KiB
Plaintext
1005 lines
30 KiB
Plaintext
//
|
|
// Generated by NVIDIA NVVM Compiler
|
|
//
|
|
// Compiler Build ID: CL-23083092
|
|
// Cuda compilation tools, release 9.1, V9.1.85
|
|
// Based on LLVM 3.4svn
|
|
//
|
|
|
|
.version 6.1
|
|
.target sm_30
|
|
.address_size 64
|
|
|
|
// .globl _Z6oxMainv
|
|
.global .align 8 .b8 pixelID[8];
|
|
.global .align 8 .b8 resolution[8];
|
|
.global .align 4 .b8 normal[12];
|
|
.global .align 4 .b8 camPos[12];
|
|
.global .align 4 .b8 root[4];
|
|
.global .align 4 .u32 imageEnabled;
|
|
.global .texref lightmap;
|
|
.global .align 16 .b8 tileInfo[16];
|
|
.global .align 4 .u32 additive;
|
|
.global .align 1 .b8 inputImageL0[1];
|
|
.global .align 1 .b8 inputImageL1x[1];
|
|
.global .align 1 .b8 inputImageL1y[1];
|
|
.global .align 1 .b8 inputImageL1z[1];
|
|
.global .align 1 .b8 inputImageL2a[1];
|
|
.global .align 1 .b8 inputImageL2b[1];
|
|
.global .align 1 .b8 inputImageL2c[1];
|
|
.global .align 1 .b8 inputImageL2d[1];
|
|
.global .align 1 .b8 inputImageL2e[1];
|
|
.global .align 1 .b8 outputImageL0[1];
|
|
.global .align 1 .b8 outputImageL1x[1];
|
|
.global .align 1 .b8 outputImageL1y[1];
|
|
.global .align 1 .b8 outputImageL1z[1];
|
|
.global .align 1 .b8 outputImageL2a[1];
|
|
.global .align 1 .b8 outputImageL2b[1];
|
|
.global .align 1 .b8 outputImageL2c[1];
|
|
.global .align 1 .b8 outputImageL2d[1];
|
|
.global .align 1 .b8 outputImageL2e[1];
|
|
.global .align 4 .b8 _ZN21rti_internal_typeinfo7pixelIDE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
|
|
.global .align 4 .b8 _ZN21rti_internal_typeinfo10resolutionE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
|
|
.global .align 4 .b8 _ZN21rti_internal_typeinfo6normalE[8] = {82, 97, 121, 0, 12, 0, 0, 0};
|
|
.global .align 4 .b8 _ZN21rti_internal_typeinfo6camPosE[8] = {82, 97, 121, 0, 12, 0, 0, 0};
|
|
.global .align 4 .b8 _ZN21rti_internal_typeinfo4rootE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
|
|
.global .align 4 .b8 _ZN21rti_internal_typeinfo12imageEnabledE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
|
|
.global .align 4 .b8 _ZN21rti_internal_typeinfo8tileInfoE[8] = {82, 97, 121, 0, 16, 0, 0, 0};
|
|
.global .align 4 .b8 _ZN21rti_internal_typeinfo8additiveE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
|
|
.global .align 8 .u64 _ZN21rti_internal_register20reg_bitness_detectorE;
|
|
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail0E;
|
|
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail1E;
|
|
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail2E;
|
|
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail3E;
|
|
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail4E;
|
|
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail5E;
|
|
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail6E;
|
|
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail7E;
|
|
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail8E;
|
|
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail9E;
|
|
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail0E;
|
|
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail1E;
|
|
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail2E;
|
|
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail3E;
|
|
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail4E;
|
|
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail5E;
|
|
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail6E;
|
|
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail7E;
|
|
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail8E;
|
|
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail9E;
|
|
.global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_xE;
|
|
.global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_yE;
|
|
.global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_zE;
|
|
.global .align 8 .b8 _ZN21rti_internal_typename7pixelIDE[6] = {117, 105, 110, 116, 50, 0};
|
|
.global .align 8 .b8 _ZN21rti_internal_typename10resolutionE[6] = {117, 105, 110, 116, 50, 0};
|
|
.global .align 8 .b8 _ZN21rti_internal_typename6normalE[7] = {102, 108, 111, 97, 116, 51, 0};
|
|
.global .align 8 .b8 _ZN21rti_internal_typename6camPosE[7] = {102, 108, 111, 97, 116, 51, 0};
|
|
.global .align 16 .b8 _ZN21rti_internal_typename4rootE[9] = {114, 116, 79, 98, 106, 101, 99, 116, 0};
|
|
.global .align 4 .b8 _ZN21rti_internal_typename12imageEnabledE[4] = {105, 110, 116, 0};
|
|
.global .align 8 .b8 _ZN21rti_internal_typename8tileInfoE[6] = {117, 105, 110, 116, 52, 0};
|
|
.global .align 4 .b8 _ZN21rti_internal_typename8additiveE[4] = {105, 110, 116, 0};
|
|
.global .align 4 .u32 _ZN21rti_internal_typeenum7pixelIDE = 4919;
|
|
.global .align 4 .u32 _ZN21rti_internal_typeenum10resolutionE = 4919;
|
|
.global .align 4 .u32 _ZN21rti_internal_typeenum6normalE = 4919;
|
|
.global .align 4 .u32 _ZN21rti_internal_typeenum6camPosE = 4919;
|
|
.global .align 4 .u32 _ZN21rti_internal_typeenum4rootE = 4919;
|
|
.global .align 4 .u32 _ZN21rti_internal_typeenum12imageEnabledE = 4919;
|
|
.global .align 4 .u32 _ZN21rti_internal_typeenum8tileInfoE = 4919;
|
|
.global .align 4 .u32 _ZN21rti_internal_typeenum8additiveE = 4919;
|
|
.global .align 16 .b8 _ZN21rti_internal_semantic7pixelIDE[14] = {114, 116, 76, 97, 117, 110, 99, 104, 73, 110, 100, 101, 120, 0};
|
|
.global .align 16 .b8 _ZN21rti_internal_semantic10resolutionE[12] = {114, 116, 76, 97, 117, 110, 99, 104, 68, 105, 109, 0};
|
|
.global .align 16 .b8 _ZN21rti_internal_semantic6normalE[17] = {97, 116, 116, 114, 105, 98, 117, 116, 101, 32, 110, 111, 114, 109, 97, 108, 0};
|
|
.global .align 1 .b8 _ZN21rti_internal_semantic6camPosE[1];
|
|
.global .align 1 .b8 _ZN21rti_internal_semantic4rootE[1];
|
|
.global .align 1 .b8 _ZN21rti_internal_semantic12imageEnabledE[1];
|
|
.global .align 1 .b8 _ZN21rti_internal_semantic8tileInfoE[1];
|
|
.global .align 1 .b8 _ZN21rti_internal_semantic8additiveE[1];
|
|
.global .align 1 .b8 _ZN23rti_internal_annotation7pixelIDE[1];
|
|
.global .align 1 .b8 _ZN23rti_internal_annotation10resolutionE[1];
|
|
.global .align 1 .b8 _ZN23rti_internal_annotation6normalE[1];
|
|
.global .align 1 .b8 _ZN23rti_internal_annotation6camPosE[1];
|
|
.global .align 1 .b8 _ZN23rti_internal_annotation4rootE[1];
|
|
.global .align 1 .b8 _ZN23rti_internal_annotation12imageEnabledE[1];
|
|
.global .align 1 .b8 _ZN23rti_internal_annotation8tileInfoE[1];
|
|
.global .align 1 .b8 _ZN23rti_internal_annotation8additiveE[1];
|
|
|
|
.visible .entry _Z6oxMainv(
|
|
|
|
)
|
|
{
|
|
.reg .b16 %rs<85>;
|
|
.reg .f32 %f<112>;
|
|
.reg .b32 %r<391>;
|
|
.reg .b64 %rd<409>;
|
|
|
|
|
|
ld.global.v2.u32 {%r131, %r132}, [pixelID];
|
|
cvt.u64.u32 %rd3, %r131;
|
|
cvt.u64.u32 %rd4, %r132;
|
|
mov.u64 %rd391, inputImageL0;
|
|
cvta.global.u64 %rd2, %rd391;
|
|
mov.u32 %r129, 2;
|
|
mov.u32 %r130, 8;
|
|
mov.u64 %rd390, 0;
|
|
// inline asm
|
|
call (%rd1), _rt_buffer_get_64, (%rd2, %r129, %r130, %rd3, %rd4, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs1, [%rd1];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f1, %rs1;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r135, %r136}, [pixelID];
|
|
cvt.u64.u32 %rd9, %r135;
|
|
cvt.u64.u32 %rd10, %r136;
|
|
// inline asm
|
|
call (%rd7), _rt_buffer_get_64, (%rd2, %r129, %r130, %rd9, %rd10, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs2, [%rd7+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f2, %rs2;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r139, %r140}, [pixelID];
|
|
cvt.u64.u32 %rd15, %r139;
|
|
cvt.u64.u32 %rd16, %r140;
|
|
// inline asm
|
|
call (%rd13), _rt_buffer_get_64, (%rd2, %r129, %r130, %rd15, %rd16, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs3, [%rd13+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f3, %rs3;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r143, %r144}, [pixelID];
|
|
cvt.u64.u32 %rd21, %r143;
|
|
cvt.u64.u32 %rd22, %r144;
|
|
mov.u64 %rd392, inputImageL1x;
|
|
cvta.global.u64 %rd20, %rd392;
|
|
// inline asm
|
|
call (%rd19), _rt_buffer_get_64, (%rd20, %r129, %r130, %rd21, %rd22, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs4, [%rd19];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f4, %rs4;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r147, %r148}, [pixelID];
|
|
cvt.u64.u32 %rd27, %r147;
|
|
cvt.u64.u32 %rd28, %r148;
|
|
// inline asm
|
|
call (%rd25), _rt_buffer_get_64, (%rd20, %r129, %r130, %rd27, %rd28, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs5, [%rd25+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f5, %rs5;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r151, %r152}, [pixelID];
|
|
cvt.u64.u32 %rd33, %r151;
|
|
cvt.u64.u32 %rd34, %r152;
|
|
// inline asm
|
|
call (%rd31), _rt_buffer_get_64, (%rd20, %r129, %r130, %rd33, %rd34, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs6, [%rd31+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f6, %rs6;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r155, %r156}, [pixelID];
|
|
cvt.u64.u32 %rd39, %r155;
|
|
cvt.u64.u32 %rd40, %r156;
|
|
mov.u64 %rd393, inputImageL1y;
|
|
cvta.global.u64 %rd38, %rd393;
|
|
// inline asm
|
|
call (%rd37), _rt_buffer_get_64, (%rd38, %r129, %r130, %rd39, %rd40, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs7, [%rd37];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f7, %rs7;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r159, %r160}, [pixelID];
|
|
cvt.u64.u32 %rd45, %r159;
|
|
cvt.u64.u32 %rd46, %r160;
|
|
// inline asm
|
|
call (%rd43), _rt_buffer_get_64, (%rd38, %r129, %r130, %rd45, %rd46, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs8, [%rd43+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f8, %rs8;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r163, %r164}, [pixelID];
|
|
cvt.u64.u32 %rd51, %r163;
|
|
cvt.u64.u32 %rd52, %r164;
|
|
// inline asm
|
|
call (%rd49), _rt_buffer_get_64, (%rd38, %r129, %r130, %rd51, %rd52, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs9, [%rd49+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f9, %rs9;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r167, %r168}, [pixelID];
|
|
cvt.u64.u32 %rd57, %r167;
|
|
cvt.u64.u32 %rd58, %r168;
|
|
mov.u64 %rd394, inputImageL1z;
|
|
cvta.global.u64 %rd56, %rd394;
|
|
// inline asm
|
|
call (%rd55), _rt_buffer_get_64, (%rd56, %r129, %r130, %rd57, %rd58, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs10, [%rd55];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f10, %rs10;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r171, %r172}, [pixelID];
|
|
cvt.u64.u32 %rd63, %r171;
|
|
cvt.u64.u32 %rd64, %r172;
|
|
// inline asm
|
|
call (%rd61), _rt_buffer_get_64, (%rd56, %r129, %r130, %rd63, %rd64, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs11, [%rd61+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f11, %rs11;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r175, %r176}, [pixelID];
|
|
cvt.u64.u32 %rd69, %r175;
|
|
cvt.u64.u32 %rd70, %r176;
|
|
// inline asm
|
|
call (%rd67), _rt_buffer_get_64, (%rd56, %r129, %r130, %rd69, %rd70, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs12, [%rd67+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f12, %rs12;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r179, %r180}, [pixelID];
|
|
cvt.u64.u32 %rd75, %r179;
|
|
cvt.u64.u32 %rd76, %r180;
|
|
mov.u64 %rd395, inputImageL2a;
|
|
cvta.global.u64 %rd74, %rd395;
|
|
// inline asm
|
|
call (%rd73), _rt_buffer_get_64, (%rd74, %r129, %r130, %rd75, %rd76, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs13, [%rd73];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f13, %rs13;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r183, %r184}, [pixelID];
|
|
cvt.u64.u32 %rd81, %r183;
|
|
cvt.u64.u32 %rd82, %r184;
|
|
// inline asm
|
|
call (%rd79), _rt_buffer_get_64, (%rd74, %r129, %r130, %rd81, %rd82, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs14, [%rd79+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f14, %rs14;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r187, %r188}, [pixelID];
|
|
cvt.u64.u32 %rd87, %r187;
|
|
cvt.u64.u32 %rd88, %r188;
|
|
// inline asm
|
|
call (%rd85), _rt_buffer_get_64, (%rd74, %r129, %r130, %rd87, %rd88, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs15, [%rd85+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f15, %rs15;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r191, %r192}, [pixelID];
|
|
cvt.u64.u32 %rd93, %r191;
|
|
cvt.u64.u32 %rd94, %r192;
|
|
mov.u64 %rd396, inputImageL2b;
|
|
cvta.global.u64 %rd92, %rd396;
|
|
// inline asm
|
|
call (%rd91), _rt_buffer_get_64, (%rd92, %r129, %r130, %rd93, %rd94, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs16, [%rd91];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f16, %rs16;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r195, %r196}, [pixelID];
|
|
cvt.u64.u32 %rd99, %r195;
|
|
cvt.u64.u32 %rd100, %r196;
|
|
// inline asm
|
|
call (%rd97), _rt_buffer_get_64, (%rd92, %r129, %r130, %rd99, %rd100, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs17, [%rd97+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f17, %rs17;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r199, %r200}, [pixelID];
|
|
cvt.u64.u32 %rd105, %r199;
|
|
cvt.u64.u32 %rd106, %r200;
|
|
// inline asm
|
|
call (%rd103), _rt_buffer_get_64, (%rd92, %r129, %r130, %rd105, %rd106, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs18, [%rd103+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f18, %rs18;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r203, %r204}, [pixelID];
|
|
cvt.u64.u32 %rd111, %r203;
|
|
cvt.u64.u32 %rd112, %r204;
|
|
mov.u64 %rd397, inputImageL2c;
|
|
cvta.global.u64 %rd110, %rd397;
|
|
// inline asm
|
|
call (%rd109), _rt_buffer_get_64, (%rd110, %r129, %r130, %rd111, %rd112, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs19, [%rd109];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f19, %rs19;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r207, %r208}, [pixelID];
|
|
cvt.u64.u32 %rd117, %r207;
|
|
cvt.u64.u32 %rd118, %r208;
|
|
// inline asm
|
|
call (%rd115), _rt_buffer_get_64, (%rd110, %r129, %r130, %rd117, %rd118, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs20, [%rd115+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f20, %rs20;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r211, %r212}, [pixelID];
|
|
cvt.u64.u32 %rd123, %r211;
|
|
cvt.u64.u32 %rd124, %r212;
|
|
// inline asm
|
|
call (%rd121), _rt_buffer_get_64, (%rd110, %r129, %r130, %rd123, %rd124, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs21, [%rd121+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f21, %rs21;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r215, %r216}, [pixelID];
|
|
cvt.u64.u32 %rd129, %r215;
|
|
cvt.u64.u32 %rd130, %r216;
|
|
mov.u64 %rd398, inputImageL2d;
|
|
cvta.global.u64 %rd128, %rd398;
|
|
// inline asm
|
|
call (%rd127), _rt_buffer_get_64, (%rd128, %r129, %r130, %rd129, %rd130, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs22, [%rd127];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f22, %rs22;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r219, %r220}, [pixelID];
|
|
cvt.u64.u32 %rd135, %r219;
|
|
cvt.u64.u32 %rd136, %r220;
|
|
// inline asm
|
|
call (%rd133), _rt_buffer_get_64, (%rd128, %r129, %r130, %rd135, %rd136, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs23, [%rd133+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f23, %rs23;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r223, %r224}, [pixelID];
|
|
cvt.u64.u32 %rd141, %r223;
|
|
cvt.u64.u32 %rd142, %r224;
|
|
// inline asm
|
|
call (%rd139), _rt_buffer_get_64, (%rd128, %r129, %r130, %rd141, %rd142, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs24, [%rd139+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f24, %rs24;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r227, %r228}, [pixelID];
|
|
cvt.u64.u32 %rd147, %r227;
|
|
cvt.u64.u32 %rd148, %r228;
|
|
mov.u64 %rd399, inputImageL2e;
|
|
cvta.global.u64 %rd146, %rd399;
|
|
// inline asm
|
|
call (%rd145), _rt_buffer_get_64, (%rd146, %r129, %r130, %rd147, %rd148, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs25, [%rd145];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f25, %rs25;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r231, %r232}, [pixelID];
|
|
cvt.u64.u32 %rd153, %r231;
|
|
cvt.u64.u32 %rd154, %r232;
|
|
// inline asm
|
|
call (%rd151), _rt_buffer_get_64, (%rd146, %r129, %r130, %rd153, %rd154, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs26, [%rd151+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f26, %rs26;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r235, %r236}, [pixelID];
|
|
cvt.u64.u32 %rd159, %r235;
|
|
cvt.u64.u32 %rd160, %r236;
|
|
// inline asm
|
|
call (%rd157), _rt_buffer_get_64, (%rd146, %r129, %r130, %rd159, %rd160, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs27, [%rd157+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f27, %rs27;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r239, %r240}, [pixelID];
|
|
cvt.u64.u32 %rd165, %r239;
|
|
cvt.u64.u32 %rd166, %r240;
|
|
mov.u64 %rd400, outputImageL0;
|
|
cvta.global.u64 %rd164, %rd400;
|
|
// inline asm
|
|
call (%rd163), _rt_buffer_get_64, (%rd164, %r129, %r130, %rd165, %rd166, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs28, [%rd163];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f28, %rs28;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r243, %r244}, [pixelID];
|
|
cvt.u64.u32 %rd171, %r243;
|
|
cvt.u64.u32 %rd172, %r244;
|
|
// inline asm
|
|
call (%rd169), _rt_buffer_get_64, (%rd164, %r129, %r130, %rd171, %rd172, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs29, [%rd169+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f29, %rs29;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r247, %r248}, [pixelID];
|
|
cvt.u64.u32 %rd177, %r247;
|
|
cvt.u64.u32 %rd178, %r248;
|
|
// inline asm
|
|
call (%rd175), _rt_buffer_get_64, (%rd164, %r129, %r130, %rd177, %rd178, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs30, [%rd175+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f30, %rs30;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r251, %r252}, [pixelID];
|
|
cvt.u64.u32 %rd183, %r251;
|
|
cvt.u64.u32 %rd184, %r252;
|
|
mov.u64 %rd401, outputImageL1x;
|
|
cvta.global.u64 %rd182, %rd401;
|
|
// inline asm
|
|
call (%rd181), _rt_buffer_get_64, (%rd182, %r129, %r130, %rd183, %rd184, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs31, [%rd181];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f31, %rs31;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r255, %r256}, [pixelID];
|
|
cvt.u64.u32 %rd189, %r255;
|
|
cvt.u64.u32 %rd190, %r256;
|
|
// inline asm
|
|
call (%rd187), _rt_buffer_get_64, (%rd182, %r129, %r130, %rd189, %rd190, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs32, [%rd187+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f32, %rs32;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r259, %r260}, [pixelID];
|
|
cvt.u64.u32 %rd195, %r259;
|
|
cvt.u64.u32 %rd196, %r260;
|
|
// inline asm
|
|
call (%rd193), _rt_buffer_get_64, (%rd182, %r129, %r130, %rd195, %rd196, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs33, [%rd193+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f33, %rs33;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r263, %r264}, [pixelID];
|
|
cvt.u64.u32 %rd201, %r263;
|
|
cvt.u64.u32 %rd202, %r264;
|
|
mov.u64 %rd402, outputImageL1y;
|
|
cvta.global.u64 %rd200, %rd402;
|
|
// inline asm
|
|
call (%rd199), _rt_buffer_get_64, (%rd200, %r129, %r130, %rd201, %rd202, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs34, [%rd199];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f34, %rs34;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r267, %r268}, [pixelID];
|
|
cvt.u64.u32 %rd207, %r267;
|
|
cvt.u64.u32 %rd208, %r268;
|
|
// inline asm
|
|
call (%rd205), _rt_buffer_get_64, (%rd200, %r129, %r130, %rd207, %rd208, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs35, [%rd205+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f35, %rs35;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r271, %r272}, [pixelID];
|
|
cvt.u64.u32 %rd213, %r271;
|
|
cvt.u64.u32 %rd214, %r272;
|
|
// inline asm
|
|
call (%rd211), _rt_buffer_get_64, (%rd200, %r129, %r130, %rd213, %rd214, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs36, [%rd211+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f36, %rs36;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r275, %r276}, [pixelID];
|
|
cvt.u64.u32 %rd219, %r275;
|
|
cvt.u64.u32 %rd220, %r276;
|
|
mov.u64 %rd403, outputImageL1z;
|
|
cvta.global.u64 %rd218, %rd403;
|
|
// inline asm
|
|
call (%rd217), _rt_buffer_get_64, (%rd218, %r129, %r130, %rd219, %rd220, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs37, [%rd217];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f37, %rs37;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r279, %r280}, [pixelID];
|
|
cvt.u64.u32 %rd225, %r279;
|
|
cvt.u64.u32 %rd226, %r280;
|
|
// inline asm
|
|
call (%rd223), _rt_buffer_get_64, (%rd218, %r129, %r130, %rd225, %rd226, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs38, [%rd223+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f38, %rs38;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r283, %r284}, [pixelID];
|
|
cvt.u64.u32 %rd231, %r283;
|
|
cvt.u64.u32 %rd232, %r284;
|
|
// inline asm
|
|
call (%rd229), _rt_buffer_get_64, (%rd218, %r129, %r130, %rd231, %rd232, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs39, [%rd229+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f39, %rs39;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r287, %r288}, [pixelID];
|
|
cvt.u64.u32 %rd237, %r287;
|
|
cvt.u64.u32 %rd238, %r288;
|
|
mov.u64 %rd404, outputImageL2a;
|
|
cvta.global.u64 %rd236, %rd404;
|
|
// inline asm
|
|
call (%rd235), _rt_buffer_get_64, (%rd236, %r129, %r130, %rd237, %rd238, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs40, [%rd235];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f40, %rs40;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r291, %r292}, [pixelID];
|
|
cvt.u64.u32 %rd243, %r291;
|
|
cvt.u64.u32 %rd244, %r292;
|
|
// inline asm
|
|
call (%rd241), _rt_buffer_get_64, (%rd236, %r129, %r130, %rd243, %rd244, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs41, [%rd241+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f41, %rs41;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r295, %r296}, [pixelID];
|
|
cvt.u64.u32 %rd249, %r295;
|
|
cvt.u64.u32 %rd250, %r296;
|
|
// inline asm
|
|
call (%rd247), _rt_buffer_get_64, (%rd236, %r129, %r130, %rd249, %rd250, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs42, [%rd247+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f42, %rs42;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r299, %r300}, [pixelID];
|
|
cvt.u64.u32 %rd255, %r299;
|
|
cvt.u64.u32 %rd256, %r300;
|
|
mov.u64 %rd405, outputImageL2b;
|
|
cvta.global.u64 %rd254, %rd405;
|
|
// inline asm
|
|
call (%rd253), _rt_buffer_get_64, (%rd254, %r129, %r130, %rd255, %rd256, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs43, [%rd253];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f43, %rs43;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r303, %r304}, [pixelID];
|
|
cvt.u64.u32 %rd261, %r303;
|
|
cvt.u64.u32 %rd262, %r304;
|
|
// inline asm
|
|
call (%rd259), _rt_buffer_get_64, (%rd254, %r129, %r130, %rd261, %rd262, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs44, [%rd259+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f44, %rs44;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r307, %r308}, [pixelID];
|
|
cvt.u64.u32 %rd267, %r307;
|
|
cvt.u64.u32 %rd268, %r308;
|
|
// inline asm
|
|
call (%rd265), _rt_buffer_get_64, (%rd254, %r129, %r130, %rd267, %rd268, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs45, [%rd265+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f45, %rs45;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r311, %r312}, [pixelID];
|
|
cvt.u64.u32 %rd273, %r311;
|
|
cvt.u64.u32 %rd274, %r312;
|
|
mov.u64 %rd406, outputImageL2c;
|
|
cvta.global.u64 %rd272, %rd406;
|
|
// inline asm
|
|
call (%rd271), _rt_buffer_get_64, (%rd272, %r129, %r130, %rd273, %rd274, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs46, [%rd271];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f46, %rs46;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r315, %r316}, [pixelID];
|
|
cvt.u64.u32 %rd279, %r315;
|
|
cvt.u64.u32 %rd280, %r316;
|
|
// inline asm
|
|
call (%rd277), _rt_buffer_get_64, (%rd272, %r129, %r130, %rd279, %rd280, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs47, [%rd277+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f47, %rs47;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r319, %r320}, [pixelID];
|
|
cvt.u64.u32 %rd285, %r319;
|
|
cvt.u64.u32 %rd286, %r320;
|
|
// inline asm
|
|
call (%rd283), _rt_buffer_get_64, (%rd272, %r129, %r130, %rd285, %rd286, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs48, [%rd283+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f48, %rs48;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r323, %r324}, [pixelID];
|
|
cvt.u64.u32 %rd291, %r323;
|
|
cvt.u64.u32 %rd292, %r324;
|
|
mov.u64 %rd407, outputImageL2d;
|
|
cvta.global.u64 %rd290, %rd407;
|
|
// inline asm
|
|
call (%rd289), _rt_buffer_get_64, (%rd290, %r129, %r130, %rd291, %rd292, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs49, [%rd289];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f49, %rs49;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r327, %r328}, [pixelID];
|
|
cvt.u64.u32 %rd297, %r327;
|
|
cvt.u64.u32 %rd298, %r328;
|
|
// inline asm
|
|
call (%rd295), _rt_buffer_get_64, (%rd290, %r129, %r130, %rd297, %rd298, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs50, [%rd295+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f50, %rs50;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r331, %r332}, [pixelID];
|
|
cvt.u64.u32 %rd303, %r331;
|
|
cvt.u64.u32 %rd304, %r332;
|
|
// inline asm
|
|
call (%rd301), _rt_buffer_get_64, (%rd290, %r129, %r130, %rd303, %rd304, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs51, [%rd301+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f51, %rs51;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r335, %r336}, [pixelID];
|
|
cvt.u64.u32 %rd309, %r335;
|
|
cvt.u64.u32 %rd310, %r336;
|
|
mov.u64 %rd408, outputImageL2e;
|
|
cvta.global.u64 %rd308, %rd408;
|
|
// inline asm
|
|
call (%rd307), _rt_buffer_get_64, (%rd308, %r129, %r130, %rd309, %rd310, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs52, [%rd307];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f52, %rs52;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r339, %r340}, [pixelID];
|
|
cvt.u64.u32 %rd315, %r339;
|
|
cvt.u64.u32 %rd316, %r340;
|
|
// inline asm
|
|
call (%rd313), _rt_buffer_get_64, (%rd308, %r129, %r130, %rd315, %rd316, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs53, [%rd313+2];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f53, %rs53;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r343, %r344}, [pixelID];
|
|
cvt.u64.u32 %rd321, %r343;
|
|
cvt.u64.u32 %rd322, %r344;
|
|
// inline asm
|
|
call (%rd319), _rt_buffer_get_64, (%rd308, %r129, %r130, %rd321, %rd322, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs54, [%rd319+4];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f54, %rs54;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r347, %r348}, [pixelID];
|
|
cvt.u64.u32 %rd327, %r347;
|
|
cvt.u64.u32 %rd328, %r348;
|
|
// inline asm
|
|
call (%rd325), _rt_buffer_get_64, (%rd2, %r129, %r130, %rd327, %rd328, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs55, [%rd325+6];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f55, %rs55;}
|
|
|
|
// inline asm
|
|
ld.global.v2.u32 {%r351, %r352}, [pixelID];
|
|
cvt.u64.u32 %rd333, %r351;
|
|
cvt.u64.u32 %rd334, %r352;
|
|
// inline asm
|
|
call (%rd331), _rt_buffer_get_64, (%rd164, %r129, %r130, %rd333, %rd334, %rd390, %rd390);
|
|
// inline asm
|
|
ld.u16 %rs56, [%rd331+6];
|
|
// inline asm
|
|
{ cvt.f32.f16 %f56, %rs56;}
|
|
|
|
// inline asm
|
|
min.f32 %f60, %f55, %f56;
|
|
add.f32 %f85, %f1, %f28;
|
|
add.f32 %f86, %f2, %f29;
|
|
add.f32 %f87, %f3, %f30;
|
|
add.f32 %f88, %f4, %f31;
|
|
add.f32 %f89, %f5, %f32;
|
|
add.f32 %f90, %f6, %f33;
|
|
add.f32 %f91, %f7, %f34;
|
|
add.f32 %f92, %f8, %f35;
|
|
add.f32 %f93, %f9, %f36;
|
|
add.f32 %f94, %f10, %f37;
|
|
add.f32 %f95, %f11, %f38;
|
|
add.f32 %f96, %f12, %f39;
|
|
add.f32 %f97, %f13, %f40;
|
|
add.f32 %f98, %f14, %f41;
|
|
add.f32 %f99, %f15, %f42;
|
|
add.f32 %f100, %f16, %f43;
|
|
add.f32 %f101, %f17, %f44;
|
|
add.f32 %f102, %f18, %f45;
|
|
add.f32 %f103, %f19, %f46;
|
|
add.f32 %f104, %f20, %f47;
|
|
add.f32 %f105, %f21, %f48;
|
|
add.f32 %f106, %f22, %f49;
|
|
add.f32 %f107, %f23, %f50;
|
|
add.f32 %f108, %f24, %f51;
|
|
add.f32 %f109, %f25, %f52;
|
|
add.f32 %f110, %f26, %f53;
|
|
add.f32 %f111, %f27, %f54;
|
|
mul.f32 %f57, %f85, %f60;
|
|
mul.f32 %f58, %f86, %f60;
|
|
mul.f32 %f59, %f87, %f60;
|
|
mul.f32 %f61, %f88, %f60;
|
|
mul.f32 %f62, %f89, %f60;
|
|
mul.f32 %f63, %f90, %f60;
|
|
mul.f32 %f64, %f91, %f60;
|
|
mul.f32 %f65, %f92, %f60;
|
|
mul.f32 %f66, %f93, %f60;
|
|
mul.f32 %f67, %f94, %f60;
|
|
mul.f32 %f68, %f95, %f60;
|
|
mul.f32 %f69, %f96, %f60;
|
|
mul.f32 %f70, %f97, %f60;
|
|
mul.f32 %f71, %f98, %f60;
|
|
mul.f32 %f72, %f99, %f60;
|
|
mul.f32 %f73, %f100, %f60;
|
|
mul.f32 %f74, %f101, %f60;
|
|
mul.f32 %f75, %f102, %f60;
|
|
mul.f32 %f76, %f103, %f60;
|
|
mul.f32 %f77, %f104, %f60;
|
|
mul.f32 %f78, %f105, %f60;
|
|
mul.f32 %f79, %f106, %f60;
|
|
mul.f32 %f80, %f107, %f60;
|
|
mul.f32 %f81, %f108, %f60;
|
|
mul.f32 %f82, %f109, %f60;
|
|
mul.f32 %f83, %f110, %f60;
|
|
mul.f32 %f84, %f111, %f60;
|
|
ld.global.v2.u32 {%r355, %r356}, [pixelID];
|
|
cvt.u64.u32 %rd339, %r355;
|
|
cvt.u64.u32 %rd340, %r356;
|
|
// inline asm
|
|
call (%rd337), _rt_buffer_get_64, (%rd164, %r129, %r130, %rd339, %rd340, %rd390, %rd390);
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs59, %f59;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs58, %f58;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs57, %f57;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs60, %f60;}
|
|
|
|
// inline asm
|
|
st.v4.u16 [%rd337], {%rs57, %rs58, %rs59, %rs60};
|
|
ld.global.v2.u32 {%r359, %r360}, [pixelID];
|
|
cvt.u64.u32 %rd345, %r359;
|
|
cvt.u64.u32 %rd346, %r360;
|
|
// inline asm
|
|
call (%rd343), _rt_buffer_get_64, (%rd182, %r129, %r130, %rd345, %rd346, %rd390, %rd390);
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs63, %f63;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs62, %f62;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs61, %f61;}
|
|
|
|
// inline asm
|
|
st.v4.u16 [%rd343], {%rs61, %rs62, %rs63, %rs60};
|
|
ld.global.v2.u32 {%r363, %r364}, [pixelID];
|
|
cvt.u64.u32 %rd351, %r363;
|
|
cvt.u64.u32 %rd352, %r364;
|
|
// inline asm
|
|
call (%rd349), _rt_buffer_get_64, (%rd200, %r129, %r130, %rd351, %rd352, %rd390, %rd390);
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs66, %f66;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs65, %f65;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs64, %f64;}
|
|
|
|
// inline asm
|
|
st.v4.u16 [%rd349], {%rs64, %rs65, %rs66, %rs60};
|
|
ld.global.v2.u32 {%r367, %r368}, [pixelID];
|
|
cvt.u64.u32 %rd357, %r367;
|
|
cvt.u64.u32 %rd358, %r368;
|
|
// inline asm
|
|
call (%rd355), _rt_buffer_get_64, (%rd218, %r129, %r130, %rd357, %rd358, %rd390, %rd390);
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs69, %f69;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs68, %f68;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs67, %f67;}
|
|
|
|
// inline asm
|
|
st.v4.u16 [%rd355], {%rs67, %rs68, %rs69, %rs60};
|
|
ld.global.v2.u32 {%r371, %r372}, [pixelID];
|
|
cvt.u64.u32 %rd363, %r371;
|
|
cvt.u64.u32 %rd364, %r372;
|
|
// inline asm
|
|
call (%rd361), _rt_buffer_get_64, (%rd236, %r129, %r130, %rd363, %rd364, %rd390, %rd390);
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs72, %f72;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs71, %f71;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs70, %f70;}
|
|
|
|
// inline asm
|
|
st.v4.u16 [%rd361], {%rs70, %rs71, %rs72, %rs60};
|
|
ld.global.v2.u32 {%r375, %r376}, [pixelID];
|
|
cvt.u64.u32 %rd369, %r375;
|
|
cvt.u64.u32 %rd370, %r376;
|
|
// inline asm
|
|
call (%rd367), _rt_buffer_get_64, (%rd254, %r129, %r130, %rd369, %rd370, %rd390, %rd390);
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs75, %f75;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs74, %f74;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs73, %f73;}
|
|
|
|
// inline asm
|
|
st.v4.u16 [%rd367], {%rs73, %rs74, %rs75, %rs60};
|
|
ld.global.v2.u32 {%r379, %r380}, [pixelID];
|
|
cvt.u64.u32 %rd375, %r379;
|
|
cvt.u64.u32 %rd376, %r380;
|
|
// inline asm
|
|
call (%rd373), _rt_buffer_get_64, (%rd272, %r129, %r130, %rd375, %rd376, %rd390, %rd390);
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs78, %f78;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs77, %f77;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs76, %f76;}
|
|
|
|
// inline asm
|
|
st.v4.u16 [%rd373], {%rs76, %rs77, %rs78, %rs60};
|
|
ld.global.v2.u32 {%r383, %r384}, [pixelID];
|
|
cvt.u64.u32 %rd381, %r383;
|
|
cvt.u64.u32 %rd382, %r384;
|
|
// inline asm
|
|
call (%rd379), _rt_buffer_get_64, (%rd290, %r129, %r130, %rd381, %rd382, %rd390, %rd390);
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs81, %f81;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs80, %f80;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs79, %f79;}
|
|
|
|
// inline asm
|
|
st.v4.u16 [%rd379], {%rs79, %rs80, %rs81, %rs60};
|
|
ld.global.v2.u32 {%r387, %r388}, [pixelID];
|
|
cvt.u64.u32 %rd387, %r387;
|
|
cvt.u64.u32 %rd388, %r388;
|
|
// inline asm
|
|
call (%rd385), _rt_buffer_get_64, (%rd308, %r129, %r130, %rd387, %rd388, %rd390, %rd390);
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs84, %f84;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs83, %f83;}
|
|
|
|
// inline asm
|
|
// inline asm
|
|
{ cvt.rn.f16.f32 %rs82, %f82;}
|
|
|
|
// inline asm
|
|
st.v4.u16 [%rd385], {%rs82, %rs83, %rs84, %rs60};
|
|
ret;
|
|
}
|
|
|
|
|