Skip to content

Instantly share code, notes, and snippets.

@drisspg

drisspg/README.md

Last active July 24, 2026 14:30
Show Gist options

  • Select an option

    Learn more about clone URLs
  • Save drisspg/4ced47465f083119405a28e6a5833872 to your computer and use it in GitHub Desktop.

Select an option

Learn more about clone URLs
Save drisspg/4ced47465f083119405a28e6a5833872 to your computer and use it in GitHub Desktop.
Download ZIP
SM100 ptxas VIMNMX3 integer-abs miscompile reproducer
Raw
README.md

SM100 ptxas integer-abs miscompile reproducer

This package contains both the trimmed Python/Triton reproducer and a CUDA Driver API + PTX variant that does not require Python, PyTorch, Triton, or Inductor at runtime.

Both variants require an SM100 GPU and a compatible CUDA driver. They were verified with bundled ptxas-blackwell 13.3 (V13.3.33) and driver 580.126.20 on GB200.

Trimmed Python/Triton reproducer

Run repro.py with an affected Triton/ptxas installation:

TRITON_ALWAYS_COMPILE=1 TRITON_CACHE_DIR=/tmp/sm100-abs-repro python repro.py

Expected result before the workaround:

expected key count: 169
maximum key-count error: 819.0
AssertionError

The same script reports zero error with the Triton PTX canonicalization.

Standalone C++/PTX reproducer

This path requires g++, CUDA headers, and ptxas with PTX 9.3 / sm_100a support:

PTXAS=/path/to/ptxas-blackwell CUDA_HOME=/usr/local/cuda-13.1 bash run.sh

run.sh:

  1. Builds repro.cpp against the CUDA Driver API.
  2. Compiles the same repro.ptx with normal optimized ptxas.
  3. Compiles it again with --opt-level=0.
  4. Launches both cubins and compares every row against the expected count of 169.

Expected output on an affected compiler:

expected key count: 169
maximum key-count error: 819
mismatched rows: 20 / 128
expected key count: 169
maximum key-count error: 0
mismatched rows: 0 / 128
optimized ptxas status: 1 (expected 1 on affected versions)
ptxas O0 status: 0 (expected 0)

Compiler error

The PTX computes an integer absolute value as max(x, -x) and then combines it with another distance. Optimized ptxas can fuse this into a VIMNMX3 whose two inputs both reference x, dropping the register holding -x. O0 preserves the correct calculation.

The PTX SHA-256 is:

c152929836c75f23dfaea2e66f9b2ca7c530f3ba1665b36218a7f75abee9b969
Raw
repro.cpp
#include <cuda.h>
#include <algorithm>
#include <array>
#include <cmath>
#include <cstdlib>
#include <iostream>
void checkCuda(CUresult result, const char* expression) {
if (result == CUDA_SUCCESS) {
return;
}
const char* name = nullptr;
const char* description = nullptr;
cuGetErrorName(result, &name);
cuGetErrorString(result, &description);
std::cerr << expression << " failed: " << (name ? name : "unknown") << ": "
<< (description ? description : "unknown") << '\n';
std::exit(2);
}
#define CUDA_CHECK(expression) checkCuda((expression), #expression)
int main(int argc, char** argv) {
if (argc != 2) {
std::cerr << "usage: " << argv[0] << " <cubin>\n";
return 2;
}
CUDA_CHECK(cuInit(0));
CUdevice device;
CUDA_CHECK(cuDeviceGet(&device, 0));
CUcontext context;
CUDA_CHECK(cuDevicePrimaryCtxRetain(&context, device));
CUDA_CHECK(cuCtxSetCurrent(context));
CUmodule module;
CUDA_CHECK(cuModuleLoad(&module, argv[1]));
CUfunction function;
CUDA_CHECK(cuModuleGetFunction(&function, module, "integer_abs_ptxas_kernel"));
constexpr size_t blockM = 128;
constexpr size_t blockN = 64;
constexpr size_t headDim = 64;
constexpr size_t iterations = 26;
constexpr size_t expectedCount = 13 * 13;
constexpr unsigned int sharedMemoryBytes = 41488;
CUdeviceptr q;
CUdeviceptr k;
CUdeviceptr output;
CUDA_CHECK(cuMemAlloc(&q, blockM * headDim * sizeof(unsigned short)));
CUDA_CHECK(cuMemAlloc(&k, iterations * blockN * headDim * sizeof(unsigned short)));
CUDA_CHECK(cuMemAlloc(&output, blockM * sizeof(float)));
CUDA_CHECK(cuMemsetD8(q, 0, blockM * headDim * sizeof(unsigned short)));
CUDA_CHECK(cuMemsetD8(k, 0, iterations * blockN * headDim * sizeof(unsigned short)));
CUDA_CHECK(cuMemsetD8(output, 0, blockM * sizeof(float)));
CUdeviceptr globalScratch = 0;
CUdeviceptr profileScratch = 0;
void* arguments[] = {&q, &k, &output, &globalScratch, &profileScratch};
CUDA_CHECK(cuLaunchKernel(function, 1, 1, 1, 128, 1, 1, sharedMemoryBytes, nullptr,
arguments, nullptr));
CUDA_CHECK(cuCtxSynchronize());
std::array<float, blockM> result;
CUDA_CHECK(cuMemcpyDtoH(result.data(), output, result.size() * sizeof(float)));
float maximumError = 0.0f;
size_t mismatches = 0;
for (float value : result) {
maximumError = std::max(maximumError, std::abs(value - expectedCount));
mismatches += value != expectedCount;
}
std::cout << "expected key count: " << expectedCount << '\n';
std::cout << "maximum key-count error: " << maximumError << '\n';
std::cout << "mismatched rows: " << mismatches << " / " << blockM << '\n';
CUDA_CHECK(cuMemFree(output));
CUDA_CHECK(cuMemFree(k));
CUDA_CHECK(cuMemFree(q));
CUDA_CHECK(cuModuleUnload(module));
CUDA_CHECK(cuCtxSetCurrent(nullptr));
CUDA_CHECK(cuDevicePrimaryCtxRelease(device));
return maximumError == 0.0f ? 0 : 1;
}
Raw
repro.ptx
//
// Generated by LLVM NVPTX Back-End
//
.version 9.3
.target sm_100a
.address_size 64
// .globl integer_abs_ptxas_kernel // -- Begin function integer_abs_ptxas_kernel
.extern .shared .align 16 .b8 global_smem[];
// @integer_abs_ptxas_kernel
.visible .entry integer_abs_ptxas_kernel(
.param .u64 .ptr .global .align 1 integer_abs_ptxas_kernel_param_0,
.param .u64 .ptr .global .align 1 integer_abs_ptxas_kernel_param_1,
.param .u64 .ptr .global .align 1 integer_abs_ptxas_kernel_param_2,
.param .u64 .ptr .global .align 1 integer_abs_ptxas_kernel_param_3,
.param .u64 .ptr .global .align 1 integer_abs_ptxas_kernel_param_4
)
.reqntid 128
{
.reg .pred %p<173>;
.reg .b16 %rs<107>;
.reg .b32 %r<2461>;
.reg .b64 %rd<584>;
.loc 1 1176 0 // test_core.py:1176:0
$L__func_begin0:
.loc 1 1176 0 // test_core.py:1176:0
// %bb.0:
ld.param.b64 %rd18, [integer_abs_ptxas_kernel_param_1];
ld.param.b64 %rd36, [integer_abs_ptxas_kernel_param_0];
$L__tmp0:
.loc 1 1176 1 // test_core.py:1176:1
mov.u32 %r1, %tid.x;
setp.lt.u32 %p1, %r1, 32;
mov.b32 %r39, global_smem;
// begin inline asm
@%p1 tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%r39], 64;
// end inline asm
bar.sync 0;
ld.shared.b32 %r1390, [global_smem];
bar.sync 0;
// begin inline asm
@%p1 tcgen05.relinquish_alloc_permit.cta_group::1.sync.aligned;
// end inline asm
.loc 1 1181 16 // test_core.py:1181:16
shr.u32 %r2, %r1, 5;
and.b32 %r3, %r1, 127;
shr.u32 %r82, %r1, 3;
bfe.u32 %r4, %r1, 3, 4;
and.b32 %r5, %r1, 64;
bfe.u32 %r83, %r1, 6, 1;
.loc 1 1182 9 // test_core.py:1182:9
or.b32 %r84, %r83, 58;
or.b32 %r85, %r83, 60;
or.b32 %r86, %r83, 62;
or.b32 %r87, %r83, 64;
or.b32 %r88, %r83, 66;
or.b32 %r89, %r83, 68;
or.b32 %r90, %r83, 70;
or.b32 %r91, %r83, 72;
or.b32 %r92, %r83, 74;
or.b32 %r93, %r83, 76;
or.b32 %r94, %r83, 78;
or.b32 %r95, %r83, 80;
or.b32 %r96, %r83, 82;
or.b32 %r97, %r83, 84;
or.b32 %r98, %r83, 86;
or.b32 %r99, %r83, 88;
or.b32 %r100, %r83, 90;
or.b32 %r101, %r83, 92;
or.b32 %r102, %r83, 94;
or.b32 %r103, %r83, 96;
or.b32 %r104, %r83, 98;
or.b32 %r105, %r83, 100;
or.b32 %r106, %r83, 102;
or.b32 %r107, %r83, 104;
or.b32 %r108, %r83, 106;
or.b32 %r109, %r83, 108;
or.b32 %r110, %r83, 110;
or.b32 %r111, %r83, 112;
or.b32 %r112, %r83, 114;
or.b32 %r113, %r83, 116;
or.b32 %r114, %r83, 118;
or.b32 %r115, %r83, 120;
.loc 1 1183 9 // test_core.py:1183:9
shl.b32 %r116, %r1, 3;
and.b32 %r117, %r116, 56;
and.b32 %r118, %r1, 63;
.loc 1 1184 25 // test_core.py:1184:25
shl.b32 %r119, %r4, 6;
shl.b32 %r120, %r82, 6;
or.b32 %r121, %r120, 7168;
.loc 1 1184 17 // test_core.py:1184:17
mul.wide.u32 %rd37, %r119, 2;
add.s64 %rd38, %rd36, %rd37;
mul.wide.u32 %rd39, %r121, 2;
add.s64 %rd40, %rd36, %rd39;
mul.wide.u32 %rd41, %r117, 2;
add.s64 %rd20, %rd38, %rd41;
add.s64 %rd21, %rd20, 2048;
add.s64 %rd22, %rd20, 4096;
add.s64 %rd23, %rd20, 6144;
add.s64 %rd24, %rd20, 8192;
add.s64 %rd25, %rd20, 10240;
add.s64 %rd26, %rd20, 12288;
add.s64 %rd27, %rd40, %rd41;
.loc 1 1184 9 // test_core.py:1184:9
// begin inline asm
.loc 1 1184 9 // test_core.py:1184:9
mov.u32 %r40, 0x0;
.loc 1 1185 9 // test_core.py:1185:9
mov.u32 %r41, 0x0;
.loc 1 1186 9 // test_core.py:1186:9
mov.u32 %r42, 0x0;
.loc 1 1187 9 // test_core.py:1187:9
mov.u32 %r43, 0x0;
.loc 1 1188 9 // test_core.py:1188:9
ld.global.v4.b32 { %r40, %r41, %r42, %r43 }, [ %rd20 + 0 ];
// end inline asm
// begin inline asm
.loc 1 1184 9 // test_core.py:1184:9
mov.u32 %r44, 0x0;
.loc 1 1185 9 // test_core.py:1185:9
mov.u32 %r45, 0x0;
.loc 1 1186 9 // test_core.py:1186:9
mov.u32 %r46, 0x0;
.loc 1 1187 9 // test_core.py:1187:9
mov.u32 %r47, 0x0;
.loc 1 1188 9 // test_core.py:1188:9
ld.global.v4.b32 { %r44, %r45, %r46, %r47 }, [ %rd21 + 0 ];
// end inline asm
// begin inline asm
.loc 1 1184 9 // test_core.py:1184:9
mov.u32 %r48, 0x0;
.loc 1 1185 9 // test_core.py:1185:9
mov.u32 %r49, 0x0;
.loc 1 1186 9 // test_core.py:1186:9
mov.u32 %r50, 0x0;
.loc 1 1187 9 // test_core.py:1187:9
mov.u32 %r51, 0x0;
.loc 1 1188 9 // test_core.py:1188:9
ld.global.v4.b32 { %r48, %r49, %r50, %r51 }, [ %rd22 + 0 ];
// end inline asm
// begin inline asm
.loc 1 1184 9 // test_core.py:1184:9
mov.u32 %r52, 0x0;
.loc 1 1185 9 // test_core.py:1185:9
mov.u32 %r53, 0x0;
.loc 1 1186 9 // test_core.py:1186:9
mov.u32 %r54, 0x0;
.loc 1 1187 9 // test_core.py:1187:9
mov.u32 %r55, 0x0;
.loc 1 1188 9 // test_core.py:1188:9
ld.global.v4.b32 { %r52, %r53, %r54, %r55 }, [ %rd23 + 0 ];
// end inline asm
// begin inline asm
.loc 1 1184 9 // test_core.py:1184:9
mov.u32 %r56, 0x0;
.loc 1 1185 9 // test_core.py:1185:9
mov.u32 %r57, 0x0;
.loc 1 1186 9 // test_core.py:1186:9
mov.u32 %r58, 0x0;
.loc 1 1187 9 // test_core.py:1187:9
mov.u32 %r59, 0x0;
.loc 1 1188 9 // test_core.py:1188:9
ld.global.v4.b32 { %r56, %r57, %r58, %r59 }, [ %rd24 + 0 ];
// end inline asm
// begin inline asm
.loc 1 1184 9 // test_core.py:1184:9
mov.u32 %r60, 0x0;
.loc 1 1185 9 // test_core.py:1185:9
mov.u32 %r61, 0x0;
.loc 1 1186 9 // test_core.py:1186:9
mov.u32 %r62, 0x0;
.loc 1 1187 9 // test_core.py:1187:9
mov.u32 %r63, 0x0;
.loc 1 1188 9 // test_core.py:1188:9
ld.global.v4.b32 { %r60, %r61, %r62, %r63 }, [ %rd25 + 0 ];
// end inline asm
// begin inline asm
.loc 1 1184 9 // test_core.py:1184:9
mov.u32 %r64, 0x0;
.loc 1 1185 9 // test_core.py:1185:9
mov.u32 %r65, 0x0;
.loc 1 1186 9 // test_core.py:1186:9
mov.u32 %r66, 0x0;
.loc 1 1187 9 // test_core.py:1187:9
mov.u32 %r67, 0x0;
.loc 1 1188 9 // test_core.py:1188:9
ld.global.v4.b32 { %r64, %r65, %r66, %r67 }, [ %rd26 + 0 ];
// end inline asm
// begin inline asm
.loc 1 1184 9 // test_core.py:1184:9
mov.u32 %r68, 0x0;
.loc 1 1185 9 // test_core.py:1185:9
mov.u32 %r69, 0x0;
.loc 1 1186 9 // test_core.py:1186:9
mov.u32 %r70, 0x0;
.loc 1 1187 9 // test_core.py:1187:9
mov.u32 %r71, 0x0;
.loc 1 1188 9 // test_core.py:1188:9
ld.global.v4.b32 { %r68, %r69, %r70, %r71 }, [ %rd27 + 0 ];
// end inline asm
shl.b32 %r122, %r3, 4;
shl.b32 %r32, %r1, 1;
and.b32 %r123, %r32, 112;
xor.b32 %r124, %r122, %r123;
add.s32 %r73, %r39, %r124;
st.shared.v4.b32 [%r73+24576], {%r40, %r41, %r42, %r43};
st.shared.v4.b32 [%r73+26624], {%r44, %r45, %r46, %r47};
st.shared.v4.b32 [%r73+28672], {%r48, %r49, %r50, %r51};
st.shared.v4.b32 [%r73+30720], {%r52, %r53, %r54, %r55};
st.shared.v4.b32 [%r73+32768], {%r56, %r57, %r58, %r59};
st.shared.v4.b32 [%r73+34816], {%r60, %r61, %r62, %r63};
st.shared.v4.b32 [%r73+36864], {%r64, %r65, %r66, %r67};
st.shared.v4.b32 [%r73+38912], {%r68, %r69, %r70, %r71};
.loc 1 1189 5 // test_core.py:1189:5
setp.eq.b32 %p4, %r2, 0;
elect.sync %r125|%p5, -1;
and.pred %p2, %p4, %p5;
add.s32 %r2452, %r39, 41472;
// begin inline asm
@%p2 mbarrier.init.shared::cta.b64 [%r2452], 1;
// end inline asm
elect.sync %r126|%p6, -1;
and.pred %p3, %p4, %p6;
add.s32 %r72, %r39, 41480;
// begin inline asm
@%p3 mbarrier.init.shared::cta.b64 [%r72], 1;
// end inline asm
.loc 1 1191 21 // test_core.py:1191:21
add.s64 %rd42, %rd18, %rd37;
add.s64 %rd28, %rd42, %rd41;
add.s64 %rd29, %rd28, 2048;
add.s64 %rd30, %rd28, 4096;
add.s64 %rd31, %rd28, 6144;
mov.b32 %r74, 16;
.loc 1 1191 13 // test_core.py:1191:13
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r73 + 0 ], [ %rd28 + 0 ], 0x10, %r74;
// end inline asm
add.s32 %r75, %r73, 2048;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r75 + 0 ], [ %rd29 + 0 ], 0x10, %r74;
// end inline asm
add.s32 %r76, %r73, 4096;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r76 + 0 ], [ %rd30 + 0 ], 0x10, %r74;
// end inline asm
add.s32 %r77, %r73, 6144;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r77 + 0 ], [ %rd31 + 0 ], 0x10, %r74;
// end inline asm
cp.async.commit_group;
.loc 1 1191 21 // test_core.py:1191:21
add.s64 %rd43, %rd18, %rd39;
add.s64 %rd32, %rd28, 8192;
add.s64 %rd33, %rd28, 10240;
add.s64 %rd34, %rd28, 12288;
add.s64 %rd35, %rd43, %rd41;
.loc 1 1191 13 // test_core.py:1191:13
bar.sync 0;
add.s32 %r78, %r73, 8192;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r78 + 0 ], [ %rd32 + 0 ], 0x10, %r74;
// end inline asm
add.s32 %r79, %r73, 10240;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r79 + 0 ], [ %rd33 + 0 ], 0x10, %r74;
// end inline asm
add.s32 %r80, %r73, 12288;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r80 + 0 ], [ %rd34 + 0 ], 0x10, %r74;
// end inline asm
add.s32 %r81, %r73, 14336;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r81 + 0 ], [ %rd35 + 0 ], 0x10, %r74;
// end inline asm
cp.async.commit_group;
cp.async.wait_group 1;
bar.sync 0;
.loc 1 1192 14 // test_core.py:1192:14
fence.proxy.async.shared::cta;
shfl.sync.idx.b32 %r33, %r2, 0, 31, -1;
setp.eq.b32 %p7, %r33, 0;
bar.sync 0;
@%p7 bra $L__BB0_2;
// %bb.1: // %._crit_edge
.loc 1 0 14 // test_core.py:0:14
add.s32 %r127, %r39, 24576;
bfe.u32 %r128, %r127, 4, 15;
cvt.u64.u32 %rd44, %r128;
or.b64 %rd582, %rd44, 4611756662049472512;
add.s64 %rd581, %rd44, 4611756662049472514;
add.s64 %rd580, %rd44, 4611756662049472516;
add.s64 %rd579, %rd44, 4611756662049472518;
.loc 1 1192 14 // test_core.py:1192:14
bra.uni $L__BB0_3;
$L__BB0_2:
elect.sync %r131|%p9, -1;
add.s32 %r132, %r39, 24576;
bfe.u32 %r133, %r132, 4, 15;
cvt.u64.u32 %rd49, %r133;
bfe.u32 %r134, %r39, 4, 15;
cvt.u64.u32 %rd50, %r134;
or.b64 %rd582, %rd49, 4611756662049472512;
or.b64 %rd45, %rd50, 4611756662049472512;
mov.b32 %r129, 135266320;
mov.pred %p8, 0;
// begin inline asm
@%p9 tcgen05.mma.cta_group::1.kind::f16 [ %r1390 + 0 ], %rd582, %rd45, %r129, %p8;
// end inline asm
add.s64 %rd581, %rd49, 4611756662049472514;
add.s64 %rd46, %rd50, 4611756662049472514;
mov.pred %p10, -1;
// begin inline asm
@%p9 tcgen05.mma.cta_group::1.kind::f16 [ %r1390 + 0 ], %rd581, %rd46, %r129, %p10;
// end inline asm
add.s64 %rd580, %rd49, 4611756662049472516;
add.s64 %rd47, %rd50, 4611756662049472516;
// begin inline asm
@%p9 tcgen05.mma.cta_group::1.kind::f16 [ %r1390 + 0 ], %rd580, %rd47, %r129, %p10;
// end inline asm
add.s64 %rd579, %rd49, 4611756662049472518;
add.s64 %rd48, %rd50, 4611756662049472518;
// begin inline asm
@%p9 tcgen05.mma.cta_group::1.kind::f16 [ %r1390 + 0 ], %rd579, %rd48, %r129, %p10;
// end inline asm
add.s32 %r130, %r39, 41472;
// begin inline asm
@%p9 tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%r130];
// end inline asm
$L__BB0_3:
.loc 1 0 14 // test_core.py:0:14
ld.param.b64 %rd19, [integer_abs_ptxas_kernel_param_2];
or.b32 %r6, %r83, 6;
or.b32 %r7, %r83, 8;
or.b32 %r8, %r83, 10;
or.b32 %r9, %r83, 12;
or.b32 %r10, %r83, 14;
or.b32 %r11, %r83, 16;
or.b32 %r12, %r83, 18;
or.b32 %r13, %r83, 20;
or.b32 %r14, %r83, 22;
or.b32 %r15, %r83, 24;
or.b32 %r16, %r83, 26;
or.b32 %r17, %r83, 28;
or.b32 %r18, %r83, 30;
or.b32 %r19, %r83, 32;
or.b32 %r20, %r83, 34;
or.b32 %r21, %r83, 36;
or.b32 %r22, %r83, 38;
or.b32 %r23, %r83, 40;
or.b32 %r24, %r83, 42;
or.b32 %r25, %r83, 44;
or.b32 %r26, %r83, 46;
or.b32 %r27, %r83, 48;
or.b32 %r28, %r83, 50;
or.b32 %r29, %r83, 52;
or.b32 %r30, %r83, 54;
or.b32 %r31, %r83, 56;
mov.b64 %rd16, {%r85, %r84};
mov.b64 %rd15, {%r87, %r86};
mov.b64 %rd14, {%r89, %r88};
mov.b64 %rd13, {%r91, %r90};
mov.b64 %rd12, {%r93, %r92};
mov.b64 %rd11, {%r95, %r94};
mov.b64 %rd10, {%r97, %r96};
mov.b64 %rd9, {%r99, %r98};
mov.b64 %rd8, {%r101, %r100};
mov.b64 %rd7, {%r103, %r102};
mov.b64 %rd6, {%r105, %r104};
mov.b64 %rd5, {%r107, %r106};
mov.b64 %rd4, {%r109, %r108};
mov.b64 %rd3, {%r111, %r110};
mov.b64 %rd2, {%r113, %r112};
mov.b64 %rd1, {%r115, %r114};
or.b32 %r2457, %r118, 10624;
.loc 1 1191 21 // test_core.py:1191:21
add.s64 %rd51, %rd28, 16384;
add.s64 %rd52, %rd28, 18432;
add.s64 %rd53, %rd28, 20480;
add.s64 %rd54, %rd28, 22528;
.loc 1 1191 13 // test_core.py:1191:13
bar.sync 0;
add.s32 %r135, %r73, 16384;
mov.b32 %r136, 16;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r135 + 0 ], [ %rd51 + 0 ], 0x10, %r136;
// end inline asm
add.s32 %r137, %r73, 18432;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r137 + 0 ], [ %rd52 + 0 ], 0x10, %r136;
// end inline asm
add.s32 %r138, %r73, 20480;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r138 + 0 ], [ %rd53 + 0 ], 0x10, %r136;
// end inline asm
add.s32 %r139, %r73, 22528;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r139 + 0 ], [ %rd54 + 0 ], 0x10, %r136;
// end inline asm
cp.async.commit_group;
.loc 1 1192 14 // test_core.py:1192:14
fence.proxy.async.shared::cta;
and.b32 %r140, %r1, 96;
shl.b32 %r141, %r140, 1;
add.s32 %r34, %r39, %r141;
shl.b32 %r142, %r1, 6;
and.b32 %r143, %r142, 64;
shl.b32 %r144, %r5, 1;
add.s32 %r145, %r39, %r143;
add.s32 %r35, %r145, %r144;
add.s32 %r36, %r39, %r5;
and.b32 %r146, %r32, 60;
add.s32 %r147, %r34, %r141;
add.s32 %r148, %r147, %r146;
add.s32 %r37, %r148, %r143;
shl.b32 %r149, %r33, 21;
and.b32 %r150, %r149, 6291456;
add.s32 %r1387, %r150, %r1390;
.loc 1 1189 5 // test_core.py:1189:5
mul.wide.u32 %rd55, %r4, 128;
and.b32 %r151, %r1, 7;
mul.wide.u32 %rd56, %r151, 16;
or.b64 %rd57, %rd55, %rd56;
add.s64 %rd17, %rd18, %rd57;
mov.b32 %r2458, 0f00000000;
mov.b32 %r2455, 1;
mov.b32 %r2454, 2;
mov.b32 %r2451, 0;
mov.b64 %rd583, 0;
mov.b32 %r2453, %r2451;
mov.b32 %r2456, %r2451;
mov.b32 %r2459, %r2458;
mov.b32 %r2460, %r2451;
bra.uni $L__BB0_4;
$L__BB0_6: // in Loop: Header=BB0_4 Depth=1
.loc 1 0 0 // test_core.py:0:0
add.f32 %r2459, %r2459, %r1300;
add.f32 %r2458, %r2458, %r1304;
.loc 1 1189 5 // test_core.py:1189:5
setp.lt.u32 %p90, %r2460, 23;
.loc 1 1192 14 // test_core.py:1192:14
add.s32 %r1317, %r2455, 1;
setp.gt.s32 %p91, %r1317, 1;
selp.b32 %r2455, 0, %r1317, %p91;
selp.b32 %r1318, 1, 0, %p91;
xor.b32 %r38, %r2456, %r1318;
.loc 1 1189 5 // test_core.py:1189:5
add.s32 %r1319, %r2454, 1;
setp.gt.s32 %p92, %r1319, 2;
selp.b32 %r2454, 0, %r1319, %p92;
.loc 1 1191 21 // test_core.py:1191:21
add.s64 %rd322, %rd17, %rd583;
add.s64 %rd318, %rd322, 24576;
add.s64 %rd319, %rd322, 26624;
add.s64 %rd320, %rd322, 28672;
.loc 1 1191 13 // test_core.py:1191:13
add.s64 %rd321, %rd322, 30720;
shl.b32 %r1320, %r2454, 13;
bar.sync 0;
add.s32 %r1312, %r73, %r1320;
selp.b32 %r1313, 16, 0, %p90;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r1312 + 0 ], [ %rd318 + 0 ], 0x10, %r1313;
// end inline asm
add.s32 %r1314, %r1312, 2048;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r1314 + 0 ], [ %rd319 + 0 ], 0x10, %r1313;
// end inline asm
add.s32 %r1315, %r1312, 4096;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r1315 + 0 ], [ %rd320 + 0 ], 0x10, %r1313;
// end inline asm
add.s32 %r1316, %r1312, 6144;
// begin inline asm
.loc 1 1191 13 // test_core.py:1191:13
cp.async.cg.shared.global [ %r1316 + 0 ], [ %rd321 + 0 ], 0x10, %r1313;
// end inline asm
cp.async.commit_group;
.loc 1 1189 5 // test_core.py:1189:5
add.s32 %r2460, %r2460, 1;
add.s64 %rd583, %rd583, 8192;
cvt.u32.u64 %r1321, %rd583;
.loc 1 1195 9 // test_core.py:1195:9
add.s32 %r2457, %r2457, 64;
.loc 1 1189 5 // test_core.py:1189:5
setp.ne.b32 %p93, %r1321, 204800;
mov.b32 %r2451, %r2456;
mov.b32 %r2456, %r38;
@!%p93 bra $L__BB0_7;
$L__BB0_4: // =>This Inner Loop Header: Depth=1
$L__tmp1:
.loc 1 1193 25 // test_core.py:1193:25
.loc 1 1170 40, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1170:40
shr.u32 %r216, %r2457, 7;
.loc 1 1170 27, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1170:27
sub.s32 %r217, 89, %r216;
.loc 1 1170 20, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1170:20
abs.s32 %r218, %r217;
.loc 1 1171 40, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1171:40
and.b32 %r219, %r2457, 127;
.loc 1 1171 27, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1171:27
sub.s32 %r220, 6, %r219;
sub.s32 %r221, %r6, %r219;
sub.s32 %r222, %r7, %r219;
sub.s32 %r223, %r8, %r219;
sub.s32 %r224, %r9, %r219;
sub.s32 %r225, %r10, %r219;
sub.s32 %r226, %r11, %r219;
sub.s32 %r227, %r12, %r219;
sub.s32 %r228, %r13, %r219;
sub.s32 %r229, %r14, %r219;
sub.s32 %r230, %r15, %r219;
sub.s32 %r231, %r16, %r219;
sub.s32 %r232, %r17, %r219;
sub.s32 %r233, %r18, %r219;
sub.s32 %r234, %r19, %r219;
sub.s32 %r235, %r20, %r219;
sub.s32 %r236, %r21, %r219;
sub.s32 %r237, %r22, %r219;
sub.s32 %r238, %r23, %r219;
sub.s32 %r239, %r24, %r219;
sub.s32 %r240, %r25, %r219;
sub.s32 %r241, %r26, %r219;
sub.s32 %r242, %r27, %r219;
sub.s32 %r243, %r28, %r219;
sub.s32 %r244, %r29, %r219;
sub.s32 %r245, %r30, %r219;
sub.s32 %r246, %r31, %r219;
sub.s32 %r247, 121, %r219;
.loc 1 1171 20, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1171:20
abs.s32 %r248, %r220;
neg.s32 %r249, %r221;
max.s32 %r250, %r221, %r249;
neg.s32 %r251, %r222;
max.s32 %r252, %r222, %r251;
neg.s32 %r253, %r223;
max.s32 %r254, %r223, %r253;
neg.s32 %r255, %r224;
max.s32 %r256, %r224, %r255;
neg.s32 %r257, %r225;
max.s32 %r258, %r225, %r257;
neg.s32 %r259, %r226;
max.s32 %r260, %r226, %r259;
neg.s32 %r261, %r227;
max.s32 %r262, %r227, %r261;
neg.s32 %r263, %r228;
max.s32 %r264, %r228, %r263;
neg.s32 %r265, %r229;
max.s32 %r266, %r229, %r265;
neg.s32 %r267, %r230;
max.s32 %r268, %r230, %r267;
neg.s32 %r269, %r231;
max.s32 %r270, %r231, %r269;
neg.s32 %r271, %r232;
max.s32 %r272, %r232, %r271;
neg.s32 %r273, %r233;
max.s32 %r274, %r233, %r273;
neg.s32 %r275, %r234;
max.s32 %r276, %r234, %r275;
neg.s32 %r277, %r235;
max.s32 %r278, %r235, %r277;
neg.s32 %r279, %r236;
max.s32 %r280, %r236, %r279;
neg.s32 %r281, %r237;
max.s32 %r282, %r237, %r281;
neg.s32 %r283, %r238;
max.s32 %r284, %r238, %r283;
neg.s32 %r285, %r239;
max.s32 %r286, %r239, %r285;
neg.s32 %r287, %r240;
max.s32 %r288, %r240, %r287;
neg.s32 %r289, %r241;
max.s32 %r290, %r241, %r289;
neg.s32 %r291, %r242;
max.s32 %r292, %r242, %r291;
neg.s32 %r293, %r243;
max.s32 %r294, %r243, %r293;
neg.s32 %r295, %r244;
max.s32 %r296, %r244, %r295;
neg.s32 %r297, %r245;
max.s32 %r298, %r245, %r297;
neg.s32 %r299, %r246;
max.s32 %r300, %r246, %r299;
abs.s32 %r301, %r247;
.loc 1 1172 12, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1172:12
max.u32 %r302, %r218, %r248;
max.s32 %r303, %r218, %r250;
max.s32 %r304, %r218, %r252;
max.s32 %r305, %r218, %r254;
max.s32 %r306, %r218, %r256;
max.s32 %r307, %r218, %r258;
max.s32 %r308, %r218, %r260;
max.s32 %r309, %r218, %r262;
max.s32 %r310, %r218, %r264;
max.s32 %r311, %r218, %r266;
max.s32 %r312, %r218, %r268;
max.s32 %r313, %r218, %r270;
max.s32 %r314, %r218, %r272;
max.s32 %r315, %r218, %r274;
max.s32 %r316, %r218, %r276;
max.s32 %r317, %r218, %r278;
max.s32 %r318, %r218, %r280;
max.s32 %r319, %r218, %r282;
max.s32 %r320, %r218, %r284;
max.s32 %r321, %r218, %r286;
max.s32 %r322, %r218, %r288;
max.s32 %r323, %r218, %r290;
max.s32 %r324, %r218, %r292;
max.s32 %r325, %r218, %r294;
max.s32 %r326, %r218, %r296;
max.s32 %r327, %r218, %r298;
max.s32 %r328, %r218, %r300;
max.u32 %r329, %r218, %r301;
setp.lt.u32 %p11, %r302, 7;
setp.lt.u32 %p12, %r303, 7;
setp.lt.u32 %p13, %r304, 7;
setp.lt.u32 %p14, %r305, 7;
setp.lt.u32 %p15, %r306, 7;
setp.lt.u32 %p16, %r307, 7;
setp.lt.u32 %p17, %r308, 7;
setp.lt.u32 %p18, %r309, 7;
setp.lt.u32 %p19, %r310, 7;
setp.lt.u32 %p20, %r311, 7;
setp.lt.u32 %p21, %r312, 7;
setp.lt.u32 %p22, %r313, 7;
setp.lt.u32 %p23, %r314, 7;
setp.lt.u32 %p24, %r315, 7;
setp.lt.u32 %p25, %r316, 7;
setp.lt.u32 %p26, %r317, 7;
setp.lt.u32 %p27, %r318, 7;
setp.lt.u32 %p28, %r319, 7;
setp.lt.u32 %p29, %r320, 7;
setp.lt.u32 %p30, %r321, 7;
setp.lt.u32 %p31, %r322, 7;
setp.lt.u32 %p32, %r323, 7;
setp.lt.u32 %p33, %r324, 7;
setp.lt.u32 %p34, %r325, 7;
setp.lt.u32 %p35, %r326, 7;
setp.lt.u32 %p36, %r327, 7;
setp.lt.u32 %p37, %r328, 7;
.loc 1 1171 27, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1171:27
mov.b64 {%r330, %r331}, %rd16;
sub.s32 %r332, %r331, %r219;
sub.s32 %r333, %r330, %r219;
mov.b64 {%r334, %r335}, %rd15;
sub.s32 %r336, %r335, %r219;
sub.s32 %r337, %r334, %r219;
mov.b64 {%r338, %r339}, %rd14;
sub.s32 %r340, %r339, %r219;
sub.s32 %r341, %r338, %r219;
mov.b64 {%r342, %r343}, %rd13;
sub.s32 %r344, %r343, %r219;
sub.s32 %r345, %r342, %r219;
mov.b64 {%r346, %r347}, %rd12;
sub.s32 %r348, %r347, %r219;
sub.s32 %r349, %r346, %r219;
mov.b64 {%r350, %r351}, %rd11;
sub.s32 %r352, %r351, %r219;
sub.s32 %r353, %r350, %r219;
mov.b64 {%r354, %r355}, %rd10;
sub.s32 %r356, %r355, %r219;
sub.s32 %r357, %r354, %r219;
mov.b64 {%r358, %r359}, %rd9;
sub.s32 %r360, %r359, %r219;
sub.s32 %r361, %r358, %r219;
mov.b64 {%r362, %r363}, %rd3;
sub.s32 %r364, %r363, %r219;
sub.s32 %r365, %r362, %r219;
mov.b64 {%r366, %r367}, %rd4;
sub.s32 %r368, %r367, %r219;
sub.s32 %r369, %r366, %r219;
mov.b64 {%r370, %r371}, %rd1;
sub.s32 %r372, %r371, %r219;
sub.s32 %r373, %r370, %r219;
mov.b64 {%r374, %r375}, %rd2;
sub.s32 %r376, %r375, %r219;
sub.s32 %r377, %r374, %r219;
mov.b64 {%r378, %r379}, %rd7;
sub.s32 %r380, %r379, %r219;
sub.s32 %r381, %r378, %r219;
mov.b64 {%r382, %r383}, %rd8;
sub.s32 %r384, %r383, %r219;
sub.s32 %r385, %r382, %r219;
mov.b64 {%r386, %r387}, %rd5;
sub.s32 %r388, %r387, %r219;
sub.s32 %r389, %r386, %r219;
mov.b64 {%r390, %r391}, %rd6;
sub.s32 %r392, %r391, %r219;
sub.s32 %r393, %r390, %r219;
.loc 1 1171 20, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1171:20
neg.s32 %r394, %r393;
max.s32 %r395, %r393, %r394;
neg.s32 %r396, %r392;
max.s32 %r397, %r392, %r396;
neg.s32 %r398, %r389;
max.s32 %r399, %r389, %r398;
neg.s32 %r400, %r388;
max.s32 %r401, %r388, %r400;
neg.s32 %r402, %r385;
max.s32 %r403, %r385, %r402;
neg.s32 %r404, %r384;
max.s32 %r405, %r384, %r404;
neg.s32 %r406, %r381;
max.s32 %r407, %r381, %r406;
neg.s32 %r408, %r380;
max.s32 %r409, %r380, %r408;
neg.s32 %r410, %r377;
max.s32 %r411, %r377, %r410;
neg.s32 %r412, %r376;
max.s32 %r413, %r376, %r412;
neg.s32 %r414, %r373;
max.s32 %r415, %r373, %r414;
neg.s32 %r416, %r372;
max.s32 %r417, %r372, %r416;
neg.s32 %r418, %r369;
max.s32 %r419, %r369, %r418;
neg.s32 %r420, %r368;
max.s32 %r421, %r368, %r420;
neg.s32 %r422, %r365;
max.s32 %r423, %r365, %r422;
neg.s32 %r424, %r364;
max.s32 %r425, %r364, %r424;
neg.s32 %r426, %r361;
max.s32 %r427, %r361, %r426;
neg.s32 %r428, %r360;
max.s32 %r429, %r360, %r428;
neg.s32 %r430, %r357;
max.s32 %r431, %r357, %r430;
neg.s32 %r432, %r356;
max.s32 %r433, %r356, %r432;
neg.s32 %r434, %r353;
max.s32 %r435, %r353, %r434;
neg.s32 %r436, %r352;
max.s32 %r437, %r352, %r436;
neg.s32 %r438, %r349;
max.s32 %r439, %r349, %r438;
neg.s32 %r440, %r348;
max.s32 %r441, %r348, %r440;
neg.s32 %r442, %r345;
max.s32 %r443, %r345, %r442;
neg.s32 %r444, %r344;
max.s32 %r445, %r344, %r444;
neg.s32 %r446, %r341;
max.s32 %r447, %r341, %r446;
neg.s32 %r448, %r340;
max.s32 %r449, %r340, %r448;
neg.s32 %r450, %r337;
max.s32 %r451, %r337, %r450;
neg.s32 %r452, %r336;
max.s32 %r453, %r336, %r452;
neg.s32 %r454, %r333;
max.s32 %r455, %r333, %r454;
neg.s32 %r456, %r332;
max.s32 %r457, %r332, %r456;
.loc 1 1172 12, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1172:12
max.s32 %r458, %r218, %r457;
max.s32 %r459, %r218, %r455;
max.s32 %r460, %r218, %r453;
max.s32 %r461, %r218, %r451;
max.s32 %r462, %r218, %r449;
max.s32 %r463, %r218, %r447;
max.s32 %r464, %r218, %r445;
max.s32 %r465, %r218, %r443;
max.s32 %r466, %r218, %r441;
max.s32 %r467, %r218, %r439;
max.s32 %r468, %r218, %r437;
max.s32 %r469, %r218, %r435;
max.s32 %r470, %r218, %r433;
max.s32 %r471, %r218, %r431;
max.s32 %r472, %r218, %r429;
max.s32 %r473, %r218, %r427;
max.s32 %r474, %r218, %r425;
max.s32 %r475, %r218, %r423;
max.s32 %r476, %r218, %r421;
max.s32 %r477, %r218, %r419;
max.s32 %r478, %r218, %r417;
max.s32 %r479, %r218, %r415;
max.s32 %r480, %r218, %r413;
max.s32 %r481, %r218, %r411;
max.s32 %r482, %r218, %r409;
max.s32 %r483, %r218, %r407;
max.s32 %r484, %r218, %r405;
max.s32 %r485, %r218, %r403;
max.s32 %r486, %r218, %r401;
max.s32 %r487, %r218, %r399;
max.s32 %r488, %r218, %r397;
max.s32 %r489, %r218, %r395;
setp.lt.u32 %p38, %r489, 7;
selp.b16 %rs1, 1, 0, %p38;
shl.b16 %rs2, %rs1, 2;
setp.lt.u32 %p39, %r488, 7;
selp.b16 %rs3, -1, 0, %p39;
shl.b16 %rs4, %rs3, 3;
or.b16 %rs5, %rs4, %rs2;
setp.lt.u32 %p40, %r487, 7;
selp.b16 %rs6, 1, 0, %p40;
setp.lt.u32 %p41, %r486, 7;
selp.b16 %rs7, -1, 0, %p41;
shl.b16 %rs8, %rs7, 1;
or.b16 %rs9, %rs6, %rs8;
and.b16 %rs10, %rs9, 3;
or.b16 %rs11, %rs10, %rs5;
and.b16 %rs12, %rs11, 15;
shl.b16 %rs13, %rs12, 8;
setp.lt.u32 %p42, %r485, 7;
selp.b16 %rs14, 1, 0, %p42;
shl.b16 %rs15, %rs14, 2;
setp.lt.u32 %p43, %r484, 7;
selp.b16 %rs16, -1, 0, %p43;
shl.b16 %rs17, %rs16, 3;
or.b16 %rs18, %rs17, %rs15;
setp.lt.u32 %p44, %r483, 7;
selp.b16 %rs19, 1, 0, %p44;
setp.lt.u32 %p45, %r482, 7;
selp.b16 %rs20, -1, 0, %p45;
shl.b16 %rs21, %rs20, 1;
or.b16 %rs22, %rs19, %rs21;
and.b16 %rs23, %rs22, 3;
or.b16 %rs24, %rs23, %rs18;
shl.b16 %rs25, %rs24, 12;
or.b16 %rs26, %rs25, %rs13;
setp.lt.u32 %p46, %r481, 7;
selp.b16 %rs27, 1, 0, %p46;
shl.b16 %rs28, %rs27, 2;
setp.lt.u32 %p47, %r480, 7;
selp.b16 %rs29, -1, 0, %p47;
shl.b16 %rs30, %rs29, 3;
or.b16 %rs31, %rs30, %rs28;
setp.lt.u32 %p48, %r479, 7;
selp.b16 %rs32, 1, 0, %p48;
setp.lt.u32 %p49, %r478, 7;
selp.b16 %rs33, -1, 0, %p49;
shl.b16 %rs34, %rs33, 1;
or.b16 %rs35, %rs32, %rs34;
and.b16 %rs36, %rs35, 3;
or.b16 %rs37, %rs36, %rs31;
and.b16 %rs38, %rs37, 15;
setp.lt.u32 %p50, %r477, 7;
selp.b16 %rs39, 1, 0, %p50;
shl.b16 %rs40, %rs39, 2;
setp.lt.u32 %p51, %r476, 7;
selp.b16 %rs41, -1, 0, %p51;
shl.b16 %rs42, %rs41, 3;
or.b16 %rs43, %rs42, %rs40;
setp.lt.u32 %p52, %r475, 7;
selp.b16 %rs44, 1, 0, %p52;
setp.lt.u32 %p53, %r474, 7;
selp.b16 %rs45, -1, 0, %p53;
shl.b16 %rs46, %rs45, 1;
or.b16 %rs47, %rs44, %rs46;
and.b16 %rs48, %rs47, 3;
or.b16 %rs49, %rs48, %rs43;
shl.b16 %rs50, %rs49, 4;
or.b16 %rs51, %rs38, %rs50;
and.b16 %rs52, %rs51, 255;
or.b16 %rs53, %rs52, %rs26;
cvt.u32.u16 %r490, %rs53;
setp.lt.u32 %p54, %r473, 7;
setp.lt.u32 %p55, %r472, 7;
setp.lt.u32 %p56, %r471, 7;
setp.lt.u32 %p57, %r470, 7;
setp.lt.u32 %p58, %r469, 7;
setp.lt.u32 %p59, %r468, 7;
setp.lt.u32 %p60, %r467, 7;
setp.lt.u32 %p61, %r466, 7;
setp.lt.u32 %p62, %r465, 7;
setp.lt.u32 %p63, %r464, 7;
setp.lt.u32 %p64, %r463, 7;
setp.lt.u32 %p65, %r462, 7;
setp.lt.u32 %p66, %r461, 7;
setp.lt.u32 %p67, %r460, 7;
setp.lt.u32 %p68, %r459, 7;
setp.lt.u32 %p69, %r458, 7;
setp.lt.u32 %p70, %r329, 7;
$L__tmp2:
.loc 1 1193 25 // test_core.py:1193:25
selp.f32 %r491, 0f3F800000, 0f00000000, %p11;
selp.f32 %r492, 0f3F800000, 0f00000000, %p12;
selp.f32 %r493, 0f3F800000, 0f00000000, %p13;
selp.f32 %r494, 0f3F800000, 0f00000000, %p14;
selp.f32 %r495, 0f3F800000, 0f00000000, %p15;
selp.f32 %r496, 0f3F800000, 0f00000000, %p16;
selp.f32 %r497, 0f3F800000, 0f00000000, %p17;
selp.f32 %r498, 0f3F800000, 0f00000000, %p18;
selp.f32 %r499, 0f3F800000, 0f00000000, %p19;
selp.f32 %r500, 0f3F800000, 0f00000000, %p20;
selp.f32 %r501, 0f3F800000, 0f00000000, %p21;
selp.f32 %r502, 0f3F800000, 0f00000000, %p22;
selp.f32 %r503, 0f3F800000, 0f00000000, %p23;
selp.f32 %r504, 0f3F800000, 0f00000000, %p24;
selp.f32 %r505, 0f3F800000, 0f00000000, %p25;
selp.f32 %r506, 0f3F800000, 0f00000000, %p26;
selp.f32 %r507, 0f3F800000, 0f00000000, %p27;
selp.f32 %r508, 0f3F800000, 0f00000000, %p28;
selp.f32 %r509, 0f3F800000, 0f00000000, %p29;
selp.f32 %r510, 0f3F800000, 0f00000000, %p30;
selp.f32 %r511, 0f3F800000, 0f00000000, %p31;
selp.f32 %r512, 0f3F800000, 0f00000000, %p32;
selp.f32 %r513, 0f3F800000, 0f00000000, %p33;
selp.f32 %r514, 0f3F800000, 0f00000000, %p34;
selp.f32 %r515, 0f3F800000, 0f00000000, %p35;
selp.f32 %r516, 0f3F800000, 0f00000000, %p36;
selp.f32 %r517, 0f3F800000, 0f00000000, %p37;
selp.f32 %r518, 0f3F800000, 0f00000000, %p69;
selp.f32 %r519, 0f3F800000, 0f00000000, %p68;
selp.f32 %r520, 0f3F800000, 0f00000000, %p67;
selp.f32 %r521, 0f3F800000, 0f00000000, %p66;
selp.f32 %r522, 0f3F800000, 0f00000000, %p65;
selp.f32 %r523, 0f3F800000, 0f00000000, %p64;
selp.f32 %r524, 0f3F800000, 0f00000000, %p63;
selp.f32 %r525, 0f3F800000, 0f00000000, %p62;
selp.f32 %r526, 0f3F800000, 0f00000000, %p61;
selp.f32 %r527, 0f3F800000, 0f00000000, %p60;
selp.f32 %r528, 0f3F800000, 0f00000000, %p59;
selp.f32 %r529, 0f3F800000, 0f00000000, %p58;
selp.f32 %r530, 0f3F800000, 0f00000000, %p57;
selp.f32 %r531, 0f3F800000, 0f00000000, %p56;
selp.f32 %r532, 0f3F800000, 0f00000000, %p55;
selp.f32 %r533, 0f3F800000, 0f00000000, %p54;
shr.u32 %r534, %r490, 15;
and.b32 %r535, %r534, 1;
setp.ne.b32 %p71, %r535, 0;
selp.b32 %r536, 1, 0, %p71;
cvt.rn.f32.u32 %r537, %r536;
shr.u32 %r538, %r490, 14;
and.b32 %r539, %r538, 1;
setp.ne.b32 %p72, %r539, 0;
selp.b32 %r540, 1, 0, %p72;
cvt.rn.f32.u32 %r541, %r540;
shr.u32 %r542, %r490, 13;
and.b32 %r543, %r542, 1;
setp.ne.b32 %p73, %r543, 0;
selp.b32 %r544, 1, 0, %p73;
cvt.rn.f32.u32 %r545, %r544;
shr.u32 %r546, %r490, 12;
and.b32 %r547, %r546, 1;
setp.ne.b32 %p74, %r547, 0;
selp.b32 %r548, 1, 0, %p74;
cvt.rn.f32.u32 %r549, %r548;
shr.u32 %r550, %r490, 11;
and.b32 %r551, %r550, 1;
setp.ne.b32 %p75, %r551, 0;
selp.b32 %r552, 1, 0, %p75;
cvt.rn.f32.u32 %r553, %r552;
shr.u32 %r554, %r490, 10;
and.b32 %r555, %r554, 1;
setp.ne.b32 %p76, %r555, 0;
selp.b32 %r556, 1, 0, %p76;
cvt.rn.f32.u32 %r557, %r556;
shr.u32 %r558, %r490, 9;
and.b32 %r559, %r558, 1;
setp.ne.b32 %p77, %r559, 0;
selp.b32 %r560, 1, 0, %p77;
cvt.rn.f32.u32 %r561, %r560;
shr.u32 %r562, %r490, 8;
and.b32 %r563, %r562, 1;
setp.ne.b32 %p78, %r563, 0;
selp.b32 %r564, 1, 0, %p78;
cvt.rn.f32.u32 %r565, %r564;
shr.u32 %r566, %r490, 7;
and.b32 %r567, %r566, 1;
setp.ne.b32 %p79, %r567, 0;
selp.b32 %r568, 1, 0, %p79;
cvt.rn.f32.u32 %r569, %r568;
shr.u32 %r570, %r490, 6;
and.b32 %r571, %r570, 1;
setp.ne.b32 %p80, %r571, 0;
selp.b32 %r572, 1, 0, %p80;
cvt.rn.f32.u32 %r573, %r572;
shr.u32 %r574, %r490, 5;
and.b32 %r575, %r574, 1;
setp.ne.b32 %p81, %r575, 0;
selp.b32 %r576, 1, 0, %p81;
cvt.rn.f32.u32 %r577, %r576;
shr.u32 %r578, %r490, 4;
and.b32 %r579, %r578, 1;
setp.ne.b32 %p82, %r579, 0;
selp.b32 %r580, 1, 0, %p82;
cvt.rn.f32.u32 %r581, %r580;
shr.u32 %r582, %r490, 3;
and.b32 %r583, %r582, 1;
setp.ne.b32 %p83, %r583, 0;
selp.b32 %r584, 1, 0, %p83;
cvt.rn.f32.u32 %r585, %r584;
shr.u32 %r586, %r490, 2;
and.b32 %r587, %r586, 1;
setp.ne.b32 %p84, %r587, 0;
selp.b32 %r588, 1, 0, %p84;
cvt.rn.f32.u32 %r589, %r588;
shr.u32 %r590, %r490, 1;
and.b32 %r591, %r590, 1;
setp.ne.b32 %p85, %r591, 0;
selp.b32 %r592, 1, 0, %p85;
cvt.rn.f32.u32 %r593, %r592;
selp.f32 %r594, 0f3F800000, 0f00000000, %p48;
selp.f32 %r595, 0f3F800000, 0f00000000, %p70;
$L__tmp3:
.loc 1 1193 18 // test_core.py:1193:18
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r596, %r491, 16, 31, -1;
$L__tmp4:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r597, %r491, %r596;
$L__tmp5:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r598, %r597, 8, 31, -1;
$L__tmp6:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r599, %r597, %r598;
$L__tmp7:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r600, %r599, 4, 31, -1;
$L__tmp8:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r601, %r599, %r600;
$L__tmp9:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r602, %r601, 2, 31, -1;
$L__tmp10:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r603, %r601, %r602;
$L__tmp11:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r604, %r603, 1, 31, -1;
shfl.sync.bfly.b32 %r605, %r491, 16, 31, -1;
$L__tmp12:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r606, %r491, %r605;
$L__tmp13:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r607, %r606, 8, 31, -1;
$L__tmp14:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r608, %r606, %r607;
$L__tmp15:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r609, %r608, 4, 31, -1;
$L__tmp16:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r610, %r608, %r609;
$L__tmp17:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r611, %r610, 2, 31, -1;
$L__tmp18:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r612, %r610, %r611;
$L__tmp19:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r613, %r612, 1, 31, -1;
shfl.sync.bfly.b32 %r614, %r491, 16, 31, -1;
$L__tmp20:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r615, %r491, %r614;
$L__tmp21:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r616, %r615, 8, 31, -1;
$L__tmp22:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r617, %r615, %r616;
$L__tmp23:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r618, %r617, 4, 31, -1;
$L__tmp24:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r619, %r617, %r618;
$L__tmp25:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r620, %r619, 2, 31, -1;
$L__tmp26:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r621, %r619, %r620;
$L__tmp27:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r622, %r621, 1, 31, -1;
shfl.sync.bfly.b32 %r623, %r492, 16, 31, -1;
$L__tmp28:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r624, %r492, %r623;
$L__tmp29:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r625, %r624, 8, 31, -1;
$L__tmp30:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r626, %r624, %r625;
$L__tmp31:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r627, %r626, 4, 31, -1;
$L__tmp32:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r628, %r626, %r627;
$L__tmp33:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r629, %r628, 2, 31, -1;
$L__tmp34:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r630, %r628, %r629;
$L__tmp35:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r631, %r630, 1, 31, -1;
mov.b64 %rd58, {%r622, %r631};
mov.b64 %rd59, {%r604, %r613};
$L__tmp36:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd60, {%r621, %r630};
mov.b64 %rd61, {%r603, %r612};
add.f32x2 %rd62, %rd61, %rd59;
add.f32x2 %rd63, %rd60, %rd58;
$L__tmp37:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r632, %r493, 16, 31, -1;
$L__tmp38:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r633, %r493, %r632;
$L__tmp39:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r634, %r633, 8, 31, -1;
$L__tmp40:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r635, %r633, %r634;
$L__tmp41:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r636, %r635, 4, 31, -1;
$L__tmp42:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r637, %r635, %r636;
$L__tmp43:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r638, %r637, 2, 31, -1;
$L__tmp44:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r639, %r637, %r638;
$L__tmp45:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r640, %r639, 1, 31, -1;
shfl.sync.bfly.b32 %r641, %r494, 16, 31, -1;
$L__tmp46:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r642, %r494, %r641;
$L__tmp47:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r643, %r642, 8, 31, -1;
$L__tmp48:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r644, %r642, %r643;
$L__tmp49:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r645, %r644, 4, 31, -1;
$L__tmp50:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r646, %r644, %r645;
$L__tmp51:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r647, %r646, 2, 31, -1;
$L__tmp52:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r648, %r646, %r647;
$L__tmp53:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r649, %r648, 1, 31, -1;
shfl.sync.bfly.b32 %r650, %r495, 16, 31, -1;
$L__tmp54:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r651, %r495, %r650;
$L__tmp55:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r652, %r651, 8, 31, -1;
$L__tmp56:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r653, %r651, %r652;
$L__tmp57:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r654, %r653, 4, 31, -1;
$L__tmp58:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r655, %r653, %r654;
$L__tmp59:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r656, %r655, 2, 31, -1;
$L__tmp60:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r657, %r655, %r656;
$L__tmp61:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r658, %r657, 1, 31, -1;
shfl.sync.bfly.b32 %r659, %r496, 16, 31, -1;
$L__tmp62:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r660, %r496, %r659;
$L__tmp63:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r661, %r660, 8, 31, -1;
$L__tmp64:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r662, %r660, %r661;
$L__tmp65:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r663, %r662, 4, 31, -1;
$L__tmp66:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r664, %r662, %r663;
$L__tmp67:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r665, %r664, 2, 31, -1;
$L__tmp68:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r666, %r664, %r665;
$L__tmp69:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r667, %r666, 1, 31, -1;
mov.b64 %rd64, {%r658, %r667};
mov.b64 %rd65, {%r640, %r649};
$L__tmp70:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd66, {%r657, %r666};
mov.b64 %rd67, {%r639, %r648};
add.f32x2 %rd68, %rd67, %rd65;
add.f32x2 %rd69, %rd66, %rd64;
$L__tmp71:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r668, %r497, 16, 31, -1;
$L__tmp72:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r669, %r497, %r668;
$L__tmp73:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r670, %r669, 8, 31, -1;
$L__tmp74:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r671, %r669, %r670;
$L__tmp75:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r672, %r671, 4, 31, -1;
$L__tmp76:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r673, %r671, %r672;
$L__tmp77:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r674, %r673, 2, 31, -1;
$L__tmp78:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r675, %r673, %r674;
$L__tmp79:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r676, %r675, 1, 31, -1;
shfl.sync.bfly.b32 %r677, %r498, 16, 31, -1;
$L__tmp80:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r678, %r498, %r677;
$L__tmp81:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r679, %r678, 8, 31, -1;
$L__tmp82:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r680, %r678, %r679;
$L__tmp83:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r681, %r680, 4, 31, -1;
$L__tmp84:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r682, %r680, %r681;
$L__tmp85:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r683, %r682, 2, 31, -1;
$L__tmp86:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r684, %r682, %r683;
$L__tmp87:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r685, %r684, 1, 31, -1;
shfl.sync.bfly.b32 %r686, %r499, 16, 31, -1;
$L__tmp88:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r687, %r499, %r686;
$L__tmp89:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r688, %r687, 8, 31, -1;
$L__tmp90:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r689, %r687, %r688;
$L__tmp91:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r690, %r689, 4, 31, -1;
$L__tmp92:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r691, %r689, %r690;
$L__tmp93:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r692, %r691, 2, 31, -1;
$L__tmp94:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r693, %r691, %r692;
$L__tmp95:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r694, %r693, 1, 31, -1;
shfl.sync.bfly.b32 %r695, %r500, 16, 31, -1;
$L__tmp96:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r696, %r500, %r695;
$L__tmp97:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r697, %r696, 8, 31, -1;
$L__tmp98:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r698, %r696, %r697;
$L__tmp99:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r699, %r698, 4, 31, -1;
$L__tmp100:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r700, %r698, %r699;
$L__tmp101:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r701, %r700, 2, 31, -1;
$L__tmp102:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r702, %r700, %r701;
$L__tmp103:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r703, %r702, 1, 31, -1;
mov.b64 %rd70, {%r694, %r703};
mov.b64 %rd71, {%r676, %r685};
$L__tmp104:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd72, {%r693, %r702};
mov.b64 %rd73, {%r675, %r684};
add.f32x2 %rd74, %rd73, %rd71;
add.f32x2 %rd75, %rd72, %rd70;
$L__tmp105:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r704, %r501, 16, 31, -1;
$L__tmp106:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r705, %r501, %r704;
$L__tmp107:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r706, %r705, 8, 31, -1;
$L__tmp108:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r707, %r705, %r706;
$L__tmp109:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r708, %r707, 4, 31, -1;
$L__tmp110:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r709, %r707, %r708;
$L__tmp111:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r710, %r709, 2, 31, -1;
$L__tmp112:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r711, %r709, %r710;
$L__tmp113:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r712, %r711, 1, 31, -1;
shfl.sync.bfly.b32 %r713, %r502, 16, 31, -1;
$L__tmp114:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r714, %r502, %r713;
$L__tmp115:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r715, %r714, 8, 31, -1;
$L__tmp116:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r716, %r714, %r715;
$L__tmp117:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r717, %r716, 4, 31, -1;
$L__tmp118:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r718, %r716, %r717;
$L__tmp119:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r719, %r718, 2, 31, -1;
$L__tmp120:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r720, %r718, %r719;
$L__tmp121:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r721, %r720, 1, 31, -1;
shfl.sync.bfly.b32 %r722, %r503, 16, 31, -1;
$L__tmp122:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r723, %r503, %r722;
$L__tmp123:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r724, %r723, 8, 31, -1;
$L__tmp124:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r725, %r723, %r724;
$L__tmp125:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r726, %r725, 4, 31, -1;
$L__tmp126:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r727, %r725, %r726;
$L__tmp127:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r728, %r727, 2, 31, -1;
$L__tmp128:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r729, %r727, %r728;
$L__tmp129:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r730, %r729, 1, 31, -1;
shfl.sync.bfly.b32 %r731, %r504, 16, 31, -1;
$L__tmp130:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r732, %r504, %r731;
$L__tmp131:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r733, %r732, 8, 31, -1;
$L__tmp132:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r734, %r732, %r733;
$L__tmp133:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r735, %r734, 4, 31, -1;
$L__tmp134:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r736, %r734, %r735;
$L__tmp135:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r737, %r736, 2, 31, -1;
$L__tmp136:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r738, %r736, %r737;
$L__tmp137:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r739, %r738, 1, 31, -1;
mov.b64 %rd76, {%r730, %r739};
mov.b64 %rd77, {%r712, %r721};
$L__tmp138:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd78, {%r729, %r738};
mov.b64 %rd79, {%r711, %r720};
add.f32x2 %rd80, %rd79, %rd77;
add.f32x2 %rd81, %rd78, %rd76;
$L__tmp139:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r740, %r505, 16, 31, -1;
$L__tmp140:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r741, %r505, %r740;
$L__tmp141:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r742, %r741, 8, 31, -1;
$L__tmp142:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r743, %r741, %r742;
$L__tmp143:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r744, %r743, 4, 31, -1;
$L__tmp144:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r745, %r743, %r744;
$L__tmp145:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r746, %r745, 2, 31, -1;
$L__tmp146:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r747, %r745, %r746;
$L__tmp147:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r748, %r747, 1, 31, -1;
shfl.sync.bfly.b32 %r749, %r506, 16, 31, -1;
$L__tmp148:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r750, %r506, %r749;
$L__tmp149:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r751, %r750, 8, 31, -1;
$L__tmp150:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r752, %r750, %r751;
$L__tmp151:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r753, %r752, 4, 31, -1;
$L__tmp152:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r754, %r752, %r753;
$L__tmp153:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r755, %r754, 2, 31, -1;
$L__tmp154:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r756, %r754, %r755;
$L__tmp155:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r757, %r756, 1, 31, -1;
shfl.sync.bfly.b32 %r758, %r507, 16, 31, -1;
$L__tmp156:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r759, %r507, %r758;
$L__tmp157:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r760, %r759, 8, 31, -1;
$L__tmp158:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r761, %r759, %r760;
$L__tmp159:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r762, %r761, 4, 31, -1;
$L__tmp160:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r763, %r761, %r762;
$L__tmp161:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r764, %r763, 2, 31, -1;
$L__tmp162:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r765, %r763, %r764;
$L__tmp163:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r766, %r765, 1, 31, -1;
shfl.sync.bfly.b32 %r767, %r508, 16, 31, -1;
$L__tmp164:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r768, %r508, %r767;
$L__tmp165:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r769, %r768, 8, 31, -1;
$L__tmp166:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r770, %r768, %r769;
$L__tmp167:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r771, %r770, 4, 31, -1;
$L__tmp168:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r772, %r770, %r771;
$L__tmp169:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r773, %r772, 2, 31, -1;
$L__tmp170:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r774, %r772, %r773;
$L__tmp171:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r775, %r774, 1, 31, -1;
mov.b64 %rd82, {%r766, %r775};
mov.b64 %rd83, {%r748, %r757};
$L__tmp172:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd84, {%r765, %r774};
mov.b64 %rd85, {%r747, %r756};
add.f32x2 %rd86, %rd85, %rd83;
add.f32x2 %rd87, %rd84, %rd82;
$L__tmp173:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r776, %r509, 16, 31, -1;
$L__tmp174:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r777, %r509, %r776;
$L__tmp175:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r778, %r777, 8, 31, -1;
$L__tmp176:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r779, %r777, %r778;
$L__tmp177:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r780, %r779, 4, 31, -1;
$L__tmp178:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r781, %r779, %r780;
$L__tmp179:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r782, %r781, 2, 31, -1;
$L__tmp180:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r783, %r781, %r782;
$L__tmp181:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r784, %r783, 1, 31, -1;
shfl.sync.bfly.b32 %r785, %r510, 16, 31, -1;
$L__tmp182:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r786, %r510, %r785;
$L__tmp183:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r787, %r786, 8, 31, -1;
$L__tmp184:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r788, %r786, %r787;
$L__tmp185:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r789, %r788, 4, 31, -1;
$L__tmp186:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r790, %r788, %r789;
$L__tmp187:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r791, %r790, 2, 31, -1;
$L__tmp188:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r792, %r790, %r791;
$L__tmp189:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r793, %r792, 1, 31, -1;
shfl.sync.bfly.b32 %r794, %r511, 16, 31, -1;
$L__tmp190:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r795, %r511, %r794;
$L__tmp191:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r796, %r795, 8, 31, -1;
$L__tmp192:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r797, %r795, %r796;
$L__tmp193:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r798, %r797, 4, 31, -1;
$L__tmp194:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r799, %r797, %r798;
$L__tmp195:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r800, %r799, 2, 31, -1;
$L__tmp196:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r801, %r799, %r800;
$L__tmp197:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r802, %r801, 1, 31, -1;
shfl.sync.bfly.b32 %r803, %r512, 16, 31, -1;
$L__tmp198:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r804, %r512, %r803;
$L__tmp199:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r805, %r804, 8, 31, -1;
$L__tmp200:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r806, %r804, %r805;
$L__tmp201:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r807, %r806, 4, 31, -1;
$L__tmp202:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r808, %r806, %r807;
$L__tmp203:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r809, %r808, 2, 31, -1;
$L__tmp204:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r810, %r808, %r809;
$L__tmp205:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r811, %r810, 1, 31, -1;
mov.b64 %rd88, {%r802, %r811};
mov.b64 %rd89, {%r784, %r793};
$L__tmp206:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd90, {%r801, %r810};
mov.b64 %rd91, {%r783, %r792};
add.f32x2 %rd92, %rd91, %rd89;
add.f32x2 %rd93, %rd90, %rd88;
$L__tmp207:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r812, %r513, 16, 31, -1;
$L__tmp208:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r813, %r513, %r812;
$L__tmp209:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r814, %r813, 8, 31, -1;
$L__tmp210:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r815, %r813, %r814;
$L__tmp211:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r816, %r815, 4, 31, -1;
$L__tmp212:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r817, %r815, %r816;
$L__tmp213:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r818, %r817, 2, 31, -1;
$L__tmp214:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r819, %r817, %r818;
$L__tmp215:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r820, %r819, 1, 31, -1;
shfl.sync.bfly.b32 %r821, %r514, 16, 31, -1;
$L__tmp216:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r822, %r514, %r821;
$L__tmp217:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r823, %r822, 8, 31, -1;
$L__tmp218:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r824, %r822, %r823;
$L__tmp219:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r825, %r824, 4, 31, -1;
$L__tmp220:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r826, %r824, %r825;
$L__tmp221:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r827, %r826, 2, 31, -1;
$L__tmp222:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r828, %r826, %r827;
$L__tmp223:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r829, %r828, 1, 31, -1;
shfl.sync.bfly.b32 %r830, %r515, 16, 31, -1;
$L__tmp224:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r831, %r515, %r830;
$L__tmp225:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r832, %r831, 8, 31, -1;
$L__tmp226:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r833, %r831, %r832;
$L__tmp227:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r834, %r833, 4, 31, -1;
$L__tmp228:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r835, %r833, %r834;
$L__tmp229:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r836, %r835, 2, 31, -1;
$L__tmp230:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r837, %r835, %r836;
$L__tmp231:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r838, %r837, 1, 31, -1;
shfl.sync.bfly.b32 %r839, %r516, 16, 31, -1;
$L__tmp232:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r840, %r516, %r839;
$L__tmp233:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r841, %r840, 8, 31, -1;
$L__tmp234:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r842, %r840, %r841;
$L__tmp235:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r843, %r842, 4, 31, -1;
$L__tmp236:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r844, %r842, %r843;
$L__tmp237:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r845, %r844, 2, 31, -1;
$L__tmp238:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r846, %r844, %r845;
$L__tmp239:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r847, %r846, 1, 31, -1;
mov.b64 %rd94, {%r838, %r847};
mov.b64 %rd95, {%r820, %r829};
$L__tmp240:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd96, {%r837, %r846};
mov.b64 %rd97, {%r819, %r828};
add.f32x2 %rd98, %rd97, %rd95;
add.f32x2 %rd99, %rd96, %rd94;
$L__tmp241:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r848, %r517, 16, 31, -1;
$L__tmp242:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r849, %r517, %r848;
$L__tmp243:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r850, %r849, 8, 31, -1;
$L__tmp244:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r851, %r849, %r850;
$L__tmp245:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r852, %r851, 4, 31, -1;
$L__tmp246:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r853, %r851, %r852;
$L__tmp247:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r854, %r853, 2, 31, -1;
$L__tmp248:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r855, %r853, %r854;
$L__tmp249:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r856, %r855, 1, 31, -1;
shfl.sync.bfly.b32 %r857, %r518, 16, 31, -1;
$L__tmp250:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r858, %r518, %r857;
$L__tmp251:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r859, %r858, 8, 31, -1;
$L__tmp252:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r860, %r858, %r859;
$L__tmp253:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r861, %r860, 4, 31, -1;
$L__tmp254:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r862, %r860, %r861;
$L__tmp255:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r863, %r862, 2, 31, -1;
$L__tmp256:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r864, %r862, %r863;
$L__tmp257:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r865, %r864, 1, 31, -1;
shfl.sync.bfly.b32 %r866, %r519, 16, 31, -1;
$L__tmp258:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r867, %r519, %r866;
$L__tmp259:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r868, %r867, 8, 31, -1;
$L__tmp260:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r869, %r867, %r868;
$L__tmp261:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r870, %r869, 4, 31, -1;
$L__tmp262:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r871, %r869, %r870;
$L__tmp263:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r872, %r871, 2, 31, -1;
$L__tmp264:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r873, %r871, %r872;
$L__tmp265:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r874, %r873, 1, 31, -1;
shfl.sync.bfly.b32 %r875, %r520, 16, 31, -1;
$L__tmp266:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r876, %r520, %r875;
$L__tmp267:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r877, %r876, 8, 31, -1;
$L__tmp268:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r878, %r876, %r877;
$L__tmp269:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r879, %r878, 4, 31, -1;
$L__tmp270:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r880, %r878, %r879;
$L__tmp271:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r881, %r880, 2, 31, -1;
$L__tmp272:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r882, %r880, %r881;
$L__tmp273:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r883, %r882, 1, 31, -1;
mov.b64 %rd100, {%r874, %r883};
mov.b64 %rd101, {%r856, %r865};
$L__tmp274:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd102, {%r873, %r882};
mov.b64 %rd103, {%r855, %r864};
add.f32x2 %rd104, %rd103, %rd101;
add.f32x2 %rd105, %rd102, %rd100;
$L__tmp275:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r884, %r521, 16, 31, -1;
$L__tmp276:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r885, %r521, %r884;
$L__tmp277:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r886, %r885, 8, 31, -1;
$L__tmp278:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r887, %r885, %r886;
$L__tmp279:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r888, %r887, 4, 31, -1;
$L__tmp280:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r889, %r887, %r888;
$L__tmp281:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r890, %r889, 2, 31, -1;
$L__tmp282:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r891, %r889, %r890;
$L__tmp283:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r892, %r891, 1, 31, -1;
shfl.sync.bfly.b32 %r893, %r522, 16, 31, -1;
$L__tmp284:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r894, %r522, %r893;
$L__tmp285:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r895, %r894, 8, 31, -1;
$L__tmp286:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r896, %r894, %r895;
$L__tmp287:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r897, %r896, 4, 31, -1;
$L__tmp288:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r898, %r896, %r897;
$L__tmp289:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r899, %r898, 2, 31, -1;
$L__tmp290:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r900, %r898, %r899;
$L__tmp291:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r901, %r900, 1, 31, -1;
shfl.sync.bfly.b32 %r902, %r523, 16, 31, -1;
$L__tmp292:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r903, %r523, %r902;
$L__tmp293:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r904, %r903, 8, 31, -1;
$L__tmp294:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r905, %r903, %r904;
$L__tmp295:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r906, %r905, 4, 31, -1;
$L__tmp296:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r907, %r905, %r906;
$L__tmp297:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r908, %r907, 2, 31, -1;
$L__tmp298:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r909, %r907, %r908;
$L__tmp299:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r910, %r909, 1, 31, -1;
shfl.sync.bfly.b32 %r911, %r524, 16, 31, -1;
$L__tmp300:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r912, %r524, %r911;
$L__tmp301:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r913, %r912, 8, 31, -1;
$L__tmp302:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r914, %r912, %r913;
$L__tmp303:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r915, %r914, 4, 31, -1;
$L__tmp304:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r916, %r914, %r915;
$L__tmp305:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r917, %r916, 2, 31, -1;
$L__tmp306:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r918, %r916, %r917;
$L__tmp307:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r919, %r918, 1, 31, -1;
mov.b64 %rd106, {%r910, %r919};
mov.b64 %rd107, {%r892, %r901};
$L__tmp308:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd108, {%r909, %r918};
mov.b64 %rd109, {%r891, %r900};
add.f32x2 %rd110, %rd109, %rd107;
add.f32x2 %rd111, %rd108, %rd106;
$L__tmp309:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r920, %r525, 16, 31, -1;
$L__tmp310:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r921, %r525, %r920;
$L__tmp311:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r922, %r921, 8, 31, -1;
$L__tmp312:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r923, %r921, %r922;
$L__tmp313:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r924, %r923, 4, 31, -1;
$L__tmp314:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r925, %r923, %r924;
$L__tmp315:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r926, %r925, 2, 31, -1;
$L__tmp316:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r927, %r925, %r926;
$L__tmp317:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r928, %r927, 1, 31, -1;
shfl.sync.bfly.b32 %r929, %r526, 16, 31, -1;
$L__tmp318:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r930, %r526, %r929;
$L__tmp319:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r931, %r930, 8, 31, -1;
$L__tmp320:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r932, %r930, %r931;
$L__tmp321:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r933, %r932, 4, 31, -1;
$L__tmp322:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r934, %r932, %r933;
$L__tmp323:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r935, %r934, 2, 31, -1;
$L__tmp324:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r936, %r934, %r935;
$L__tmp325:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r937, %r936, 1, 31, -1;
shfl.sync.bfly.b32 %r938, %r527, 16, 31, -1;
$L__tmp326:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r939, %r527, %r938;
$L__tmp327:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r940, %r939, 8, 31, -1;
$L__tmp328:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r941, %r939, %r940;
$L__tmp329:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r942, %r941, 4, 31, -1;
$L__tmp330:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r943, %r941, %r942;
$L__tmp331:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r944, %r943, 2, 31, -1;
$L__tmp332:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r945, %r943, %r944;
$L__tmp333:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r946, %r945, 1, 31, -1;
shfl.sync.bfly.b32 %r947, %r528, 16, 31, -1;
$L__tmp334:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r948, %r528, %r947;
$L__tmp335:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r949, %r948, 8, 31, -1;
$L__tmp336:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r950, %r948, %r949;
$L__tmp337:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r951, %r950, 4, 31, -1;
$L__tmp338:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r952, %r950, %r951;
$L__tmp339:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r953, %r952, 2, 31, -1;
$L__tmp340:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r954, %r952, %r953;
$L__tmp341:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r955, %r954, 1, 31, -1;
mov.b64 %rd112, {%r946, %r955};
mov.b64 %rd113, {%r928, %r937};
$L__tmp342:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd114, {%r945, %r954};
mov.b64 %rd115, {%r927, %r936};
add.f32x2 %rd116, %rd115, %rd113;
add.f32x2 %rd117, %rd114, %rd112;
$L__tmp343:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r956, %r529, 16, 31, -1;
$L__tmp344:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r957, %r529, %r956;
$L__tmp345:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r958, %r957, 8, 31, -1;
$L__tmp346:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r959, %r957, %r958;
$L__tmp347:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r960, %r959, 4, 31, -1;
$L__tmp348:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r961, %r959, %r960;
$L__tmp349:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r962, %r961, 2, 31, -1;
$L__tmp350:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r963, %r961, %r962;
$L__tmp351:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r964, %r963, 1, 31, -1;
shfl.sync.bfly.b32 %r965, %r530, 16, 31, -1;
$L__tmp352:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r966, %r530, %r965;
$L__tmp353:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r967, %r966, 8, 31, -1;
$L__tmp354:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r968, %r966, %r967;
$L__tmp355:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r969, %r968, 4, 31, -1;
$L__tmp356:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r970, %r968, %r969;
$L__tmp357:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r971, %r970, 2, 31, -1;
$L__tmp358:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r972, %r970, %r971;
$L__tmp359:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r973, %r972, 1, 31, -1;
shfl.sync.bfly.b32 %r974, %r531, 16, 31, -1;
$L__tmp360:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r975, %r531, %r974;
$L__tmp361:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r976, %r975, 8, 31, -1;
$L__tmp362:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r977, %r975, %r976;
$L__tmp363:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r978, %r977, 4, 31, -1;
$L__tmp364:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r979, %r977, %r978;
$L__tmp365:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r980, %r979, 2, 31, -1;
$L__tmp366:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r981, %r979, %r980;
$L__tmp367:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r982, %r981, 1, 31, -1;
shfl.sync.bfly.b32 %r983, %r532, 16, 31, -1;
$L__tmp368:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r984, %r532, %r983;
$L__tmp369:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r985, %r984, 8, 31, -1;
$L__tmp370:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r986, %r984, %r985;
$L__tmp371:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r987, %r986, 4, 31, -1;
$L__tmp372:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r988, %r986, %r987;
$L__tmp373:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r989, %r988, 2, 31, -1;
$L__tmp374:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r990, %r988, %r989;
$L__tmp375:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r991, %r990, 1, 31, -1;
mov.b64 %rd118, {%r982, %r991};
mov.b64 %rd119, {%r964, %r973};
$L__tmp376:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd120, {%r981, %r990};
mov.b64 %rd121, {%r963, %r972};
add.f32x2 %rd122, %rd121, %rd119;
add.f32x2 %rd123, %rd120, %rd118;
$L__tmp377:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r992, %r533, 16, 31, -1;
$L__tmp378:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r993, %r533, %r992;
$L__tmp379:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r994, %r993, 8, 31, -1;
$L__tmp380:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r995, %r993, %r994;
$L__tmp381:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r996, %r995, 4, 31, -1;
$L__tmp382:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r997, %r995, %r996;
$L__tmp383:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r998, %r997, 2, 31, -1;
$L__tmp384:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r999, %r997, %r998;
$L__tmp385:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1000, %r999, 1, 31, -1;
shfl.sync.bfly.b32 %r1001, %r537, 16, 31, -1;
$L__tmp386:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1002, %r537, %r1001;
$L__tmp387:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1003, %r1002, 8, 31, -1;
$L__tmp388:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1004, %r1002, %r1003;
$L__tmp389:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1005, %r1004, 4, 31, -1;
$L__tmp390:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1006, %r1004, %r1005;
$L__tmp391:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1007, %r1006, 2, 31, -1;
$L__tmp392:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1008, %r1006, %r1007;
$L__tmp393:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1009, %r1008, 1, 31, -1;
shfl.sync.bfly.b32 %r1010, %r541, 16, 31, -1;
$L__tmp394:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1011, %r541, %r1010;
$L__tmp395:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1012, %r1011, 8, 31, -1;
$L__tmp396:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1013, %r1011, %r1012;
$L__tmp397:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1014, %r1013, 4, 31, -1;
$L__tmp398:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1015, %r1013, %r1014;
$L__tmp399:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1016, %r1015, 2, 31, -1;
$L__tmp400:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1017, %r1015, %r1016;
$L__tmp401:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1018, %r1017, 1, 31, -1;
shfl.sync.bfly.b32 %r1019, %r545, 16, 31, -1;
$L__tmp402:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1020, %r545, %r1019;
$L__tmp403:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1021, %r1020, 8, 31, -1;
$L__tmp404:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1022, %r1020, %r1021;
$L__tmp405:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1023, %r1022, 4, 31, -1;
$L__tmp406:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1024, %r1022, %r1023;
$L__tmp407:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1025, %r1024, 2, 31, -1;
$L__tmp408:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1026, %r1024, %r1025;
$L__tmp409:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1027, %r1026, 1, 31, -1;
mov.b64 %rd124, {%r1018, %r1027};
mov.b64 %rd125, {%r1000, %r1009};
$L__tmp410:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd126, {%r1017, %r1026};
mov.b64 %rd127, {%r999, %r1008};
add.f32x2 %rd128, %rd126, %rd124;
add.f32x2 %rd129, %rd127, %rd125;
$L__tmp411:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1028, %r549, 16, 31, -1;
$L__tmp412:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1029, %r549, %r1028;
$L__tmp413:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1030, %r1029, 8, 31, -1;
$L__tmp414:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1031, %r1029, %r1030;
$L__tmp415:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1032, %r1031, 4, 31, -1;
$L__tmp416:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1033, %r1031, %r1032;
$L__tmp417:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1034, %r1033, 2, 31, -1;
$L__tmp418:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1035, %r1033, %r1034;
$L__tmp419:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1036, %r1035, 1, 31, -1;
shfl.sync.bfly.b32 %r1037, %r553, 16, 31, -1;
$L__tmp420:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1038, %r553, %r1037;
$L__tmp421:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1039, %r1038, 8, 31, -1;
$L__tmp422:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1040, %r1038, %r1039;
$L__tmp423:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1041, %r1040, 4, 31, -1;
$L__tmp424:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1042, %r1040, %r1041;
$L__tmp425:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1043, %r1042, 2, 31, -1;
$L__tmp426:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1044, %r1042, %r1043;
$L__tmp427:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1045, %r1044, 1, 31, -1;
shfl.sync.bfly.b32 %r1046, %r557, 16, 31, -1;
$L__tmp428:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1047, %r557, %r1046;
$L__tmp429:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1048, %r1047, 8, 31, -1;
$L__tmp430:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1049, %r1047, %r1048;
$L__tmp431:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1050, %r1049, 4, 31, -1;
$L__tmp432:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1051, %r1049, %r1050;
$L__tmp433:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1052, %r1051, 2, 31, -1;
$L__tmp434:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1053, %r1051, %r1052;
$L__tmp435:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1054, %r1053, 1, 31, -1;
shfl.sync.bfly.b32 %r1055, %r561, 16, 31, -1;
$L__tmp436:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1056, %r561, %r1055;
$L__tmp437:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1057, %r1056, 8, 31, -1;
$L__tmp438:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1058, %r1056, %r1057;
$L__tmp439:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1059, %r1058, 4, 31, -1;
$L__tmp440:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1060, %r1058, %r1059;
$L__tmp441:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1061, %r1060, 2, 31, -1;
$L__tmp442:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1062, %r1060, %r1061;
$L__tmp443:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1063, %r1062, 1, 31, -1;
mov.b64 %rd130, {%r1054, %r1063};
mov.b64 %rd131, {%r1036, %r1045};
$L__tmp444:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd132, {%r1053, %r1062};
mov.b64 %rd133, {%r1035, %r1044};
add.f32x2 %rd134, %rd133, %rd131;
add.f32x2 %rd135, %rd132, %rd130;
$L__tmp445:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1064, %r565, 16, 31, -1;
$L__tmp446:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1065, %r565, %r1064;
$L__tmp447:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1066, %r1065, 8, 31, -1;
$L__tmp448:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1067, %r1065, %r1066;
$L__tmp449:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1068, %r1067, 4, 31, -1;
$L__tmp450:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1069, %r1067, %r1068;
$L__tmp451:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1070, %r1069, 2, 31, -1;
$L__tmp452:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1071, %r1069, %r1070;
$L__tmp453:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1072, %r1071, 1, 31, -1;
shfl.sync.bfly.b32 %r1073, %r569, 16, 31, -1;
$L__tmp454:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1074, %r569, %r1073;
$L__tmp455:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1075, %r1074, 8, 31, -1;
$L__tmp456:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1076, %r1074, %r1075;
$L__tmp457:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1077, %r1076, 4, 31, -1;
$L__tmp458:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1078, %r1076, %r1077;
$L__tmp459:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1079, %r1078, 2, 31, -1;
$L__tmp460:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1080, %r1078, %r1079;
$L__tmp461:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1081, %r1080, 1, 31, -1;
shfl.sync.bfly.b32 %r1082, %r573, 16, 31, -1;
$L__tmp462:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1083, %r573, %r1082;
$L__tmp463:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1084, %r1083, 8, 31, -1;
$L__tmp464:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1085, %r1083, %r1084;
$L__tmp465:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1086, %r1085, 4, 31, -1;
$L__tmp466:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1087, %r1085, %r1086;
$L__tmp467:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1088, %r1087, 2, 31, -1;
$L__tmp468:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1089, %r1087, %r1088;
$L__tmp469:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1090, %r1089, 1, 31, -1;
shfl.sync.bfly.b32 %r1091, %r577, 16, 31, -1;
$L__tmp470:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1092, %r577, %r1091;
$L__tmp471:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1093, %r1092, 8, 31, -1;
$L__tmp472:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1094, %r1092, %r1093;
$L__tmp473:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1095, %r1094, 4, 31, -1;
$L__tmp474:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1096, %r1094, %r1095;
$L__tmp475:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1097, %r1096, 2, 31, -1;
$L__tmp476:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1098, %r1096, %r1097;
$L__tmp477:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1099, %r1098, 1, 31, -1;
mov.b64 %rd136, {%r1090, %r1099};
mov.b64 %rd137, {%r1072, %r1081};
$L__tmp478:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd138, {%r1089, %r1098};
mov.b64 %rd139, {%r1071, %r1080};
add.f32x2 %rd140, %rd139, %rd137;
add.f32x2 %rd141, %rd138, %rd136;
$L__tmp479:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1100, %r581, 16, 31, -1;
$L__tmp480:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1101, %r581, %r1100;
$L__tmp481:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1102, %r1101, 8, 31, -1;
$L__tmp482:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1103, %r1101, %r1102;
$L__tmp483:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1104, %r1103, 4, 31, -1;
$L__tmp484:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1105, %r1103, %r1104;
$L__tmp485:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1106, %r1105, 2, 31, -1;
$L__tmp486:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1107, %r1105, %r1106;
$L__tmp487:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1108, %r1107, 1, 31, -1;
shfl.sync.bfly.b32 %r1109, %r585, 16, 31, -1;
$L__tmp488:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1110, %r585, %r1109;
$L__tmp489:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1111, %r1110, 8, 31, -1;
$L__tmp490:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1112, %r1110, %r1111;
$L__tmp491:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1113, %r1112, 4, 31, -1;
$L__tmp492:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1114, %r1112, %r1113;
$L__tmp493:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1115, %r1114, 2, 31, -1;
$L__tmp494:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1116, %r1114, %r1115;
$L__tmp495:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1117, %r1116, 1, 31, -1;
shfl.sync.bfly.b32 %r1118, %r589, 16, 31, -1;
$L__tmp496:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1119, %r589, %r1118;
$L__tmp497:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1120, %r1119, 8, 31, -1;
$L__tmp498:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1121, %r1119, %r1120;
$L__tmp499:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1122, %r1121, 4, 31, -1;
$L__tmp500:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1123, %r1121, %r1122;
$L__tmp501:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1124, %r1123, 2, 31, -1;
$L__tmp502:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1125, %r1123, %r1124;
$L__tmp503:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1126, %r1125, 1, 31, -1;
shfl.sync.bfly.b32 %r1127, %r593, 16, 31, -1;
$L__tmp504:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1128, %r593, %r1127;
$L__tmp505:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1129, %r1128, 8, 31, -1;
$L__tmp506:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1130, %r1128, %r1129;
$L__tmp507:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1131, %r1130, 4, 31, -1;
$L__tmp508:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1132, %r1130, %r1131;
$L__tmp509:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1133, %r1132, 2, 31, -1;
$L__tmp510:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1134, %r1132, %r1133;
$L__tmp511:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1135, %r1134, 1, 31, -1;
mov.b64 %rd142, {%r1126, %r1135};
mov.b64 %rd143, {%r1108, %r1117};
$L__tmp512:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd144, {%r1125, %r1134};
mov.b64 %rd145, {%r1107, %r1116};
add.f32x2 %rd146, %rd145, %rd143;
add.f32x2 %rd147, %rd144, %rd142;
$L__tmp513:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1136, %r594, 16, 31, -1;
$L__tmp514:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1137, %r594, %r1136;
$L__tmp515:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1138, %r1137, 8, 31, -1;
$L__tmp516:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1139, %r1137, %r1138;
$L__tmp517:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1140, %r1139, 4, 31, -1;
$L__tmp518:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1141, %r1139, %r1140;
$L__tmp519:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1142, %r1141, 2, 31, -1;
$L__tmp520:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1143, %r1141, %r1142;
$L__tmp521:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1144, %r1143, 1, 31, -1;
shfl.sync.bfly.b32 %r1145, %r595, 16, 31, -1;
$L__tmp522:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1146, %r595, %r1145;
$L__tmp523:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1147, %r1146, 8, 31, -1;
$L__tmp524:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1148, %r1146, %r1147;
$L__tmp525:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1149, %r1148, 4, 31, -1;
$L__tmp526:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1150, %r1148, %r1149;
$L__tmp527:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1151, %r1150, 2, 31, -1;
$L__tmp528:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1152, %r1150, %r1151;
$L__tmp529:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1153, %r1152, 1, 31, -1;
shfl.sync.bfly.b32 %r1154, %r595, 16, 31, -1;
$L__tmp530:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1155, %r595, %r1154;
$L__tmp531:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1156, %r1155, 8, 31, -1;
$L__tmp532:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1157, %r1155, %r1156;
$L__tmp533:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1158, %r1157, 4, 31, -1;
$L__tmp534:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1159, %r1157, %r1158;
$L__tmp535:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1160, %r1159, 2, 31, -1;
$L__tmp536:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1161, %r1159, %r1160;
$L__tmp537:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1162, %r1161, 1, 31, -1;
shfl.sync.bfly.b32 %r1163, %r595, 16, 31, -1;
$L__tmp538:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1164, %r595, %r1163;
$L__tmp539:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1165, %r1164, 8, 31, -1;
$L__tmp540:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1166, %r1164, %r1165;
$L__tmp541:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1167, %r1166, 4, 31, -1;
$L__tmp542:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1168, %r1166, %r1167;
$L__tmp543:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1169, %r1168, 2, 31, -1;
$L__tmp544:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1170, %r1168, %r1169;
$L__tmp545:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1171, %r1170, 1, 31, -1;
mov.b64 %rd148, {%r1162, %r1171};
mov.b64 %rd149, {%r1144, %r1153};
$L__tmp546:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd150, {%r1161, %r1170};
mov.b64 %rd151, {%r1143, %r1152};
add.f32x2 %rd152, %rd151, %rd149;
add.f32x2 %rd153, %rd150, %rd148;
$L__tmp547:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
st.shared.v2.b64 [%r34+40960], {%rd62, %rd63};
st.shared.v2.b64 [%r34+40976], {%rd86, %rd87};
st.shared.v2.b64 [%r34+40992], {%rd110, %rd111};
st.shared.v2.b64 [%r34+41008], {%rd134, %rd135};
bar.sync 0;
ld.shared.v2.b64 {%rd154, %rd155}, [%r35+40960];
mov.b64 {%r1172, %r1173}, %rd155;
mov.b64 {%r1174, %r1175}, %rd154;
ld.shared.v2.b64 {%rd156, %rd157}, [%r35+40976];
mov.b64 {%r1176, %r1177}, %rd157;
mov.b64 {%r1178, %r1179}, %rd156;
ld.shared.v2.b64 {%rd158, %rd159}, [%r35+40992];
mov.b64 {%r1180, %r1181}, %rd159;
mov.b64 {%r1182, %r1183}, %rd158;
ld.shared.v2.b64 {%rd160, %rd161}, [%r35+41008];
mov.b64 {%r1184, %r1185}, %rd161;
mov.b64 {%r1186, %r1187}, %rd160;
bar.sync 0;
st.shared.v2.b64 [%r34+40960], {%rd68, %rd69};
st.shared.v2.b64 [%r34+40976], {%rd92, %rd93};
st.shared.v2.b64 [%r34+40992], {%rd116, %rd117};
st.shared.v2.b64 [%r34+41008], {%rd140, %rd141};
bar.sync 0;
ld.shared.v2.b64 {%rd162, %rd163}, [%r35+40960];
mov.b64 {%r1188, %r1189}, %rd163;
mov.b64 {%r1190, %r1191}, %rd162;
ld.shared.v2.b64 {%rd164, %rd165}, [%r35+40976];
mov.b64 {%r1192, %r1193}, %rd165;
mov.b64 {%r1194, %r1195}, %rd164;
ld.shared.v2.b64 {%rd166, %rd167}, [%r35+40992];
mov.b64 {%r1196, %r1197}, %rd167;
mov.b64 {%r1198, %r1199}, %rd166;
ld.shared.v2.b64 {%rd168, %rd169}, [%r35+41008];
mov.b64 {%r1200, %r1201}, %rd169;
mov.b64 {%r1202, %r1203}, %rd168;
bar.sync 0;
st.shared.v2.b64 [%r34+40960], {%rd74, %rd75};
st.shared.v2.b64 [%r34+40976], {%rd98, %rd99};
st.shared.v2.b64 [%r34+40992], {%rd122, %rd123};
st.shared.v2.b64 [%r34+41008], {%rd146, %rd147};
bar.sync 0;
ld.shared.v2.b64 {%rd170, %rd171}, [%r35+40960];
mov.b64 {%r1204, %r1205}, %rd171;
mov.b64 {%r1206, %r1207}, %rd170;
ld.shared.v2.b64 {%rd172, %rd173}, [%r35+40976];
mov.b64 {%r1208, %r1209}, %rd173;
mov.b64 {%r1210, %r1211}, %rd172;
ld.shared.v2.b64 {%rd174, %rd175}, [%r35+40992];
mov.b64 {%r1212, %r1213}, %rd175;
mov.b64 {%r1214, %r1215}, %rd174;
ld.shared.v2.b64 {%rd176, %rd177}, [%r35+41008];
mov.b64 {%r1216, %r1217}, %rd177;
mov.b64 {%r1218, %r1219}, %rd176;
bar.sync 0;
st.shared.v2.b64 [%r34+40960], {%rd80, %rd81};
st.shared.v2.b64 [%r34+40976], {%rd104, %rd105};
st.shared.v2.b64 [%r34+40992], {%rd129, %rd128};
st.shared.v2.b64 [%r34+41008], {%rd152, %rd153};
bar.sync 0;
shfl.sync.bfly.b32 %r1220, %r1174, 1, 31, -1;
shfl.sync.bfly.b32 %r1221, %r1175, 1, 31, -1;
shfl.sync.bfly.b32 %r1222, %r1172, 1, 31, -1;
shfl.sync.bfly.b32 %r1223, %r1173, 1, 31, -1;
mov.b64 %rd178, {%r1222, %r1223};
mov.b64 %rd179, {%r1220, %r1221};
$L__tmp548:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd180, %rd154, %rd179;
add.f32x2 %rd181, %rd155, %rd178;
$L__tmp549:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1224, %r1190, 1, 31, -1;
shfl.sync.bfly.b32 %r1225, %r1191, 1, 31, -1;
shfl.sync.bfly.b32 %r1226, %r1188, 1, 31, -1;
shfl.sync.bfly.b32 %r1227, %r1189, 1, 31, -1;
mov.b64 %rd182, {%r1226, %r1227};
mov.b64 %rd183, {%r1224, %r1225};
$L__tmp550:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd184, %rd162, %rd183;
add.f32x2 %rd185, %rd163, %rd182;
$L__tmp551:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1228, %r1206, 1, 31, -1;
shfl.sync.bfly.b32 %r1229, %r1207, 1, 31, -1;
shfl.sync.bfly.b32 %r1230, %r1204, 1, 31, -1;
shfl.sync.bfly.b32 %r1231, %r1205, 1, 31, -1;
mov.b64 %rd186, {%r1230, %r1231};
mov.b64 %rd187, {%r1228, %r1229};
$L__tmp552:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd188, %rd170, %rd187;
add.f32x2 %rd189, %rd171, %rd186;
$L__tmp553:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
ld.shared.v2.b64 {%rd190, %rd191}, [%r35+40960];
mov.b64 {%r1232, %r1233}, %rd191;
mov.b64 {%r1234, %r1235}, %rd190;
shfl.sync.bfly.b32 %r1236, %r1234, 1, 31, -1;
shfl.sync.bfly.b32 %r1237, %r1235, 1, 31, -1;
shfl.sync.bfly.b32 %r1238, %r1232, 1, 31, -1;
shfl.sync.bfly.b32 %r1239, %r1233, 1, 31, -1;
mov.b64 %rd192, {%r1238, %r1239};
mov.b64 %rd193, {%r1236, %r1237};
$L__tmp554:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd194, %rd190, %rd193;
add.f32x2 %rd195, %rd191, %rd192;
$L__tmp555:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1240, %r1178, 1, 31, -1;
shfl.sync.bfly.b32 %r1241, %r1179, 1, 31, -1;
shfl.sync.bfly.b32 %r1242, %r1176, 1, 31, -1;
shfl.sync.bfly.b32 %r1243, %r1177, 1, 31, -1;
mov.b64 %rd196, {%r1242, %r1243};
mov.b64 %rd197, {%r1240, %r1241};
$L__tmp556:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd198, %rd156, %rd197;
add.f32x2 %rd199, %rd157, %rd196;
$L__tmp557:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1244, %r1194, 1, 31, -1;
shfl.sync.bfly.b32 %r1245, %r1195, 1, 31, -1;
shfl.sync.bfly.b32 %r1246, %r1192, 1, 31, -1;
shfl.sync.bfly.b32 %r1247, %r1193, 1, 31, -1;
mov.b64 %rd200, {%r1246, %r1247};
mov.b64 %rd201, {%r1244, %r1245};
$L__tmp558:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd202, %rd164, %rd201;
add.f32x2 %rd203, %rd165, %rd200;
$L__tmp559:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1248, %r1210, 1, 31, -1;
shfl.sync.bfly.b32 %r1249, %r1211, 1, 31, -1;
shfl.sync.bfly.b32 %r1250, %r1208, 1, 31, -1;
shfl.sync.bfly.b32 %r1251, %r1209, 1, 31, -1;
mov.b64 %rd204, {%r1250, %r1251};
mov.b64 %rd205, {%r1248, %r1249};
$L__tmp560:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd206, %rd172, %rd205;
add.f32x2 %rd207, %rd173, %rd204;
$L__tmp561:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
ld.shared.v2.b64 {%rd208, %rd209}, [%r35+40976];
mov.b64 {%r1252, %r1253}, %rd209;
mov.b64 {%r1254, %r1255}, %rd208;
shfl.sync.bfly.b32 %r1256, %r1254, 1, 31, -1;
shfl.sync.bfly.b32 %r1257, %r1255, 1, 31, -1;
shfl.sync.bfly.b32 %r1258, %r1252, 1, 31, -1;
shfl.sync.bfly.b32 %r1259, %r1253, 1, 31, -1;
mov.b64 %rd210, {%r1258, %r1259};
mov.b64 %rd211, {%r1256, %r1257};
$L__tmp562:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd212, %rd208, %rd211;
add.f32x2 %rd213, %rd209, %rd210;
$L__tmp563:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1260, %r1182, 1, 31, -1;
shfl.sync.bfly.b32 %r1261, %r1183, 1, 31, -1;
shfl.sync.bfly.b32 %r1262, %r1180, 1, 31, -1;
shfl.sync.bfly.b32 %r1263, %r1181, 1, 31, -1;
mov.b64 %rd214, {%r1262, %r1263};
mov.b64 %rd215, {%r1260, %r1261};
$L__tmp564:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd216, %rd158, %rd215;
add.f32x2 %rd217, %rd159, %rd214;
$L__tmp565:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1264, %r1198, 1, 31, -1;
shfl.sync.bfly.b32 %r1265, %r1199, 1, 31, -1;
shfl.sync.bfly.b32 %r1266, %r1196, 1, 31, -1;
shfl.sync.bfly.b32 %r1267, %r1197, 1, 31, -1;
mov.b64 %rd218, {%r1266, %r1267};
mov.b64 %rd219, {%r1264, %r1265};
$L__tmp566:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd220, %rd166, %rd219;
add.f32x2 %rd221, %rd167, %rd218;
$L__tmp567:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1268, %r1214, 1, 31, -1;
shfl.sync.bfly.b32 %r1269, %r1215, 1, 31, -1;
shfl.sync.bfly.b32 %r1270, %r1212, 1, 31, -1;
shfl.sync.bfly.b32 %r1271, %r1213, 1, 31, -1;
mov.b64 %rd222, {%r1270, %r1271};
mov.b64 %rd223, {%r1268, %r1269};
$L__tmp568:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd224, %rd174, %rd223;
add.f32x2 %rd225, %rd175, %rd222;
$L__tmp569:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
ld.shared.v2.b64 {%rd226, %rd227}, [%r35+40992];
mov.b64 {%r1272, %r1273}, %rd227;
mov.b64 {%r1274, %r1275}, %rd226;
shfl.sync.bfly.b32 %r1276, %r1274, 1, 31, -1;
shfl.sync.bfly.b32 %r1277, %r1275, 1, 31, -1;
shfl.sync.bfly.b32 %r1278, %r1272, 1, 31, -1;
shfl.sync.bfly.b32 %r1279, %r1273, 1, 31, -1;
mov.b64 %rd228, {%r1278, %r1279};
mov.b64 %rd229, {%r1276, %r1277};
$L__tmp570:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd230, %rd226, %rd229;
add.f32x2 %rd231, %rd227, %rd228;
$L__tmp571:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1280, %r1186, 1, 31, -1;
shfl.sync.bfly.b32 %r1281, %r1187, 1, 31, -1;
shfl.sync.bfly.b32 %r1282, %r1184, 1, 31, -1;
shfl.sync.bfly.b32 %r1283, %r1185, 1, 31, -1;
mov.b64 %rd232, {%r1282, %r1283};
mov.b64 %rd233, {%r1280, %r1281};
$L__tmp572:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd234, %rd160, %rd233;
add.f32x2 %rd235, %rd161, %rd232;
$L__tmp573:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1284, %r1202, 1, 31, -1;
shfl.sync.bfly.b32 %r1285, %r1203, 1, 31, -1;
shfl.sync.bfly.b32 %r1286, %r1200, 1, 31, -1;
shfl.sync.bfly.b32 %r1287, %r1201, 1, 31, -1;
mov.b64 %rd236, {%r1286, %r1287};
mov.b64 %rd237, {%r1284, %r1285};
$L__tmp574:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd238, %rd168, %rd237;
add.f32x2 %rd239, %rd169, %rd236;
$L__tmp575:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1288, %r1218, 1, 31, -1;
shfl.sync.bfly.b32 %r1289, %r1219, 1, 31, -1;
shfl.sync.bfly.b32 %r1290, %r1216, 1, 31, -1;
shfl.sync.bfly.b32 %r1291, %r1217, 1, 31, -1;
mov.b64 %rd240, {%r1290, %r1291};
mov.b64 %rd241, {%r1288, %r1289};
$L__tmp576:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd242, %rd176, %rd241;
add.f32x2 %rd243, %rd177, %rd240;
$L__tmp577:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
ld.shared.v2.b64 {%rd244, %rd245}, [%r35+41008];
mov.b64 {%r1292, %r1293}, %rd245;
mov.b64 {%r1294, %r1295}, %rd244;
shfl.sync.bfly.b32 %r1296, %r1294, 1, 31, -1;
shfl.sync.bfly.b32 %r1297, %r1295, 1, 31, -1;
shfl.sync.bfly.b32 %r1298, %r1292, 1, 31, -1;
shfl.sync.bfly.b32 %r1299, %r1293, 1, 31, -1;
mov.b64 %rd246, {%r1298, %r1299};
mov.b64 %rd247, {%r1296, %r1297};
$L__tmp578:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd248, %rd244, %rd247;
add.f32x2 %rd249, %rd245, %rd246;
$L__tmp579:
.loc 1 1193 9 // test_core.py:1193:9
bar.sync 0;
st.shared.v2.b64 [%r36+40960], {%rd180, %rd181};
st.shared.v2.b64 [%r36+40976], {%rd184, %rd185};
st.shared.v2.b64 [%r36+40992], {%rd188, %rd189};
st.shared.v2.b64 [%r36+41008], {%rd194, %rd195};
st.shared.v2.b64 [%r36+41088], {%rd198, %rd199};
st.shared.v2.b64 [%r36+41104], {%rd202, %rd203};
st.shared.v2.b64 [%r36+41120], {%rd206, %rd207};
st.shared.v2.b64 [%r36+41136], {%rd212, %rd213};
st.shared.v2.b64 [%r36+41216], {%rd216, %rd217};
st.shared.v2.b64 [%r36+41232], {%rd220, %rd221};
st.shared.v2.b64 [%r36+41248], {%rd224, %rd225};
st.shared.v2.b64 [%r36+41264], {%rd230, %rd231};
st.shared.v2.b64 [%r36+41344], {%rd234, %rd235};
st.shared.v2.b64 [%r36+41360], {%rd238, %rd239};
st.shared.v2.b64 [%r36+41376], {%rd242, %rd243};
st.shared.v2.b64 [%r36+41392], {%rd248, %rd249};
bar.sync 0;
ld.shared.b32 %r1300, [%r37+40960];
.loc 1 1192 14 // test_core.py:1192:14
// begin inline asm
{
.reg .pred complete;
.loc 1 1195 14 // test_core.py:1195:14
waitLoop:
.loc 1 1196 14 // test_core.py:1196:14
mbarrier.try_wait.parity.shared::cta.b64 complete, [%r2452], %r2451;
@!complete bra.uni waitLoop;
}
// end inline asm
.loc 1 1189 5 // test_core.py:1189:5
add.s32 %r1301, %r2453, 1;
setp.gt.s32 %p86, %r1301, 2;
selp.b32 %r2453, 0, %r1301, %p86;
.loc 1 1192 14 // test_core.py:1192:14
// begin inline asm
.loc 1 1192 14 // test_core.py:1192:14
tcgen05.ld.sync.aligned.32x32b.x64.b32 {%r152, %r153, %r154, %r155, %r156, %r157, %r158, %r159, %r160, %r161, %r162, %r163, %r164, %r165, %r166, %r167, %r168, %r169, %r170, %r171, %r172, %r173, %r174, %r175, %r176, %r177, %r178, %r179, %r180, %r181, %r182, %r183, %r184, %r185, %r186, %r187, %r188, %r189, %r190, %r191, %r192, %r193, %r194, %r195, %r196, %r197, %r198, %r199, %r200, %r201, %r202, %r203, %r204, %r205, %r206, %r207, %r208, %r209, %r210, %r211, %r212, %r213, %r214, %r215}, [%r1387 + 0];
// end inline asm
tcgen05.wait::ld.sync.aligned;
$L__tmp580:
.loc 1 1194 20 // test_core.py:1194:20
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1194 20 // standard.py:293:12
mov.b64 %rd250, {%r152, %r153};
mov.b64 %rd251, {%r154, %r155};
mov.b64 %rd252, {%r156, %r157};
mov.b64 %rd253, {%r158, %r159};
mov.b64 %rd254, {%r160, %r161};
mov.b64 %rd255, {%r162, %r163};
mov.b64 %rd256, {%r164, %r165};
mov.b64 %rd257, {%r166, %r167};
mov.b64 %rd258, {%r168, %r169};
mov.b64 %rd259, {%r170, %r171};
mov.b64 %rd260, {%r172, %r173};
mov.b64 %rd261, {%r174, %r175};
mov.b64 %rd262, {%r176, %r177};
mov.b64 %rd263, {%r178, %r179};
mov.b64 %rd264, {%r180, %r181};
mov.b64 %rd265, {%r182, %r183};
mov.b64 %rd266, {%r184, %r185};
mov.b64 %rd267, {%r186, %r187};
mov.b64 %rd268, {%r188, %r189};
mov.b64 %rd269, {%r190, %r191};
mov.b64 %rd270, {%r192, %r193};
mov.b64 %rd271, {%r194, %r195};
mov.b64 %rd272, {%r196, %r197};
mov.b64 %rd273, {%r198, %r199};
mov.b64 %rd274, {%r200, %r201};
mov.b64 %rd275, {%r202, %r203};
mov.b64 %rd276, {%r204, %r205};
mov.b64 %rd277, {%r206, %r207};
mov.b64 %rd278, {%r208, %r209};
mov.b64 %rd279, {%r210, %r211};
mov.b64 %rd280, {%r212, %r213};
mov.b64 %rd281, {%r214, %r215};
add.f32x2 %rd282, %rd250, %rd251;
add.f32x2 %rd283, %rd252, %rd253;
add.f32x2 %rd284, %rd254, %rd255;
add.f32x2 %rd285, %rd256, %rd257;
add.f32x2 %rd286, %rd258, %rd259;
add.f32x2 %rd287, %rd260, %rd261;
add.f32x2 %rd288, %rd262, %rd263;
add.f32x2 %rd289, %rd264, %rd265;
add.f32x2 %rd290, %rd266, %rd267;
add.f32x2 %rd291, %rd268, %rd269;
add.f32x2 %rd292, %rd270, %rd271;
add.f32x2 %rd293, %rd272, %rd273;
add.f32x2 %rd294, %rd274, %rd275;
add.f32x2 %rd295, %rd276, %rd277;
add.f32x2 %rd296, %rd278, %rd279;
add.f32x2 %rd297, %rd280, %rd281;
add.f32x2 %rd298, %rd282, %rd283;
add.f32x2 %rd299, %rd284, %rd285;
add.f32x2 %rd300, %rd286, %rd287;
add.f32x2 %rd301, %rd288, %rd289;
add.f32x2 %rd302, %rd290, %rd291;
add.f32x2 %rd303, %rd292, %rd293;
add.f32x2 %rd304, %rd294, %rd295;
add.f32x2 %rd305, %rd296, %rd297;
add.f32x2 %rd306, %rd298, %rd299;
add.f32x2 %rd307, %rd300, %rd301;
add.f32x2 %rd308, %rd302, %rd303;
add.f32x2 %rd309, %rd304, %rd305;
add.f32x2 %rd310, %rd306, %rd307;
add.f32x2 %rd311, %rd308, %rd309;
add.f32x2 %rd312, %rd310, %rd311;
$L__tmp581:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 {%r1302, %r1303}, %rd312;
add.f32 %r1304, %r1302, %r1303;
$L__tmp582:
.loc 1 1191 13 // test_core.py:1191:13
cp.async.wait_group 1;
bar.sync 0;
.loc 1 1189 5 // test_core.py:1189:5
shl.b32 %r1305, %r2455, 3;
add.s32 %r1306, %r39, %r1305;
add.s32 %r2452, %r1306, 41472;
.loc 1 1192 14 // test_core.py:1192:14
fence.proxy.async.shared::cta;
bar.sync 0;
bar.sync 0;
@!%p7 bra $L__BB0_6;
// %bb.5: // in Loop: Header=BB0_4 Depth=1
.loc 1 1191 13 // test_core.py:1191:13
shl.b32 %r1308, %r2453, 13;
add.s32 %r1309, %r39, %r1308;
.loc 1 1192 14 // test_core.py:1192:14
elect.sync %r1310|%p88, -1;
bfe.u32 %r1311, %r1309, 4, 15;
cvt.u64.u32 %rd317, %r1311;
or.b64 %rd313, %rd317, 4611756662049472512;
mov.b32 %r1307, 135266320;
mov.pred %p87, 0;
// begin inline asm
@%p88 tcgen05.mma.cta_group::1.kind::f16 [ %r1390 + 0 ], %rd582, %rd313, %r1307, %p87;
// end inline asm
add.s64 %rd314, %rd317, 4611756662049472514;
mov.pred %p89, -1;
// begin inline asm
@%p88 tcgen05.mma.cta_group::1.kind::f16 [ %r1390 + 0 ], %rd581, %rd314, %r1307, %p89;
// end inline asm
add.s64 %rd315, %rd317, 4611756662049472516;
// begin inline asm
@%p88 tcgen05.mma.cta_group::1.kind::f16 [ %r1390 + 0 ], %rd580, %rd315, %r1307, %p89;
// end inline asm
add.s64 %rd316, %rd317, 4611756662049472518;
// begin inline asm
@%p88 tcgen05.mma.cta_group::1.kind::f16 [ %r1390 + 0 ], %rd579, %rd316, %r1307, %p89;
// end inline asm
// begin inline asm
@%p88 tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%r2452];
// end inline asm
bra.uni $L__BB0_6;
$L__BB0_7:
$L__tmp583:
.loc 1 1170 40, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1170:40
shr.u32 %r1391, %r2457, 7;
.loc 1 1170 27, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1170:27
sub.s32 %r1392, 89, %r1391;
.loc 1 1170 20, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1170:20
abs.s32 %r1393, %r1392;
.loc 1 1171 40, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1171:40
and.b32 %r1394, %r2457, 127;
.loc 1 1171 27, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1171:27
sub.s32 %r1395, 6, %r1394;
sub.s32 %r1396, %r6, %r1394;
sub.s32 %r1397, %r7, %r1394;
sub.s32 %r1398, %r8, %r1394;
sub.s32 %r1399, %r9, %r1394;
sub.s32 %r1400, %r10, %r1394;
sub.s32 %r1401, %r11, %r1394;
sub.s32 %r1402, %r12, %r1394;
sub.s32 %r1403, %r13, %r1394;
sub.s32 %r1404, %r14, %r1394;
sub.s32 %r1405, %r15, %r1394;
sub.s32 %r1406, %r16, %r1394;
sub.s32 %r1407, %r17, %r1394;
sub.s32 %r1408, %r18, %r1394;
sub.s32 %r1409, %r19, %r1394;
sub.s32 %r1410, %r20, %r1394;
sub.s32 %r1411, %r21, %r1394;
sub.s32 %r1412, %r22, %r1394;
sub.s32 %r1413, %r23, %r1394;
sub.s32 %r1414, %r24, %r1394;
sub.s32 %r1415, %r25, %r1394;
sub.s32 %r1416, %r26, %r1394;
sub.s32 %r1417, %r27, %r1394;
sub.s32 %r1418, %r28, %r1394;
sub.s32 %r1419, %r29, %r1394;
sub.s32 %r1420, %r30, %r1394;
sub.s32 %r1421, %r31, %r1394;
sub.s32 %r1422, 121, %r1394;
.loc 1 1171 20, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1171:20
abs.s32 %r1423, %r1395;
neg.s32 %r1424, %r1396;
max.s32 %r1425, %r1396, %r1424;
neg.s32 %r1426, %r1397;
max.s32 %r1427, %r1397, %r1426;
neg.s32 %r1428, %r1398;
max.s32 %r1429, %r1398, %r1428;
neg.s32 %r1430, %r1399;
max.s32 %r1431, %r1399, %r1430;
neg.s32 %r1432, %r1400;
max.s32 %r1433, %r1400, %r1432;
neg.s32 %r1434, %r1401;
max.s32 %r1435, %r1401, %r1434;
neg.s32 %r1436, %r1402;
max.s32 %r1437, %r1402, %r1436;
neg.s32 %r1438, %r1403;
max.s32 %r1439, %r1403, %r1438;
neg.s32 %r1440, %r1404;
max.s32 %r1441, %r1404, %r1440;
neg.s32 %r1442, %r1405;
max.s32 %r1443, %r1405, %r1442;
neg.s32 %r1444, %r1406;
max.s32 %r1445, %r1406, %r1444;
neg.s32 %r1446, %r1407;
max.s32 %r1447, %r1407, %r1446;
neg.s32 %r1448, %r1408;
max.s32 %r1449, %r1408, %r1448;
neg.s32 %r1450, %r1409;
max.s32 %r1451, %r1409, %r1450;
neg.s32 %r1452, %r1410;
max.s32 %r1453, %r1410, %r1452;
neg.s32 %r1454, %r1411;
max.s32 %r1455, %r1411, %r1454;
neg.s32 %r1456, %r1412;
max.s32 %r1457, %r1412, %r1456;
neg.s32 %r1458, %r1413;
max.s32 %r1459, %r1413, %r1458;
neg.s32 %r1460, %r1414;
max.s32 %r1461, %r1414, %r1460;
neg.s32 %r1462, %r1415;
max.s32 %r1463, %r1415, %r1462;
neg.s32 %r1464, %r1416;
max.s32 %r1465, %r1416, %r1464;
neg.s32 %r1466, %r1417;
max.s32 %r1467, %r1417, %r1466;
neg.s32 %r1468, %r1418;
max.s32 %r1469, %r1418, %r1468;
neg.s32 %r1470, %r1419;
max.s32 %r1471, %r1419, %r1470;
neg.s32 %r1472, %r1420;
max.s32 %r1473, %r1420, %r1472;
neg.s32 %r1474, %r1421;
max.s32 %r1475, %r1421, %r1474;
abs.s32 %r1476, %r1422;
.loc 1 1172 12, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1172:12
max.u32 %r1477, %r1393, %r1423;
max.s32 %r1478, %r1393, %r1425;
max.s32 %r1479, %r1393, %r1427;
max.s32 %r1480, %r1393, %r1429;
max.s32 %r1481, %r1393, %r1431;
max.s32 %r1482, %r1393, %r1433;
max.s32 %r1483, %r1393, %r1435;
max.s32 %r1484, %r1393, %r1437;
max.s32 %r1485, %r1393, %r1439;
max.s32 %r1486, %r1393, %r1441;
max.s32 %r1487, %r1393, %r1443;
max.s32 %r1488, %r1393, %r1445;
max.s32 %r1489, %r1393, %r1447;
max.s32 %r1490, %r1393, %r1449;
max.s32 %r1491, %r1393, %r1451;
max.s32 %r1492, %r1393, %r1453;
max.s32 %r1493, %r1393, %r1455;
max.s32 %r1494, %r1393, %r1457;
max.s32 %r1495, %r1393, %r1459;
max.s32 %r1496, %r1393, %r1461;
max.s32 %r1497, %r1393, %r1463;
max.s32 %r1498, %r1393, %r1465;
max.s32 %r1499, %r1393, %r1467;
max.s32 %r1500, %r1393, %r1469;
max.s32 %r1501, %r1393, %r1471;
max.s32 %r1502, %r1393, %r1473;
max.s32 %r1503, %r1393, %r1475;
max.u32 %r1504, %r1393, %r1476;
setp.lt.u32 %p96, %r1477, 7;
setp.lt.u32 %p97, %r1478, 7;
setp.lt.u32 %p98, %r1479, 7;
setp.lt.u32 %p99, %r1480, 7;
setp.lt.u32 %p100, %r1481, 7;
setp.lt.u32 %p101, %r1482, 7;
setp.lt.u32 %p102, %r1483, 7;
setp.lt.u32 %p103, %r1484, 7;
setp.lt.u32 %p104, %r1485, 7;
setp.lt.u32 %p105, %r1486, 7;
setp.lt.u32 %p106, %r1487, 7;
setp.lt.u32 %p107, %r1488, 7;
setp.lt.u32 %p108, %r1489, 7;
setp.lt.u32 %p109, %r1490, 7;
setp.lt.u32 %p110, %r1491, 7;
setp.lt.u32 %p111, %r1492, 7;
setp.lt.u32 %p112, %r1493, 7;
setp.lt.u32 %p113, %r1494, 7;
setp.lt.u32 %p114, %r1495, 7;
setp.lt.u32 %p115, %r1496, 7;
setp.lt.u32 %p116, %r1497, 7;
setp.lt.u32 %p117, %r1498, 7;
setp.lt.u32 %p118, %r1499, 7;
setp.lt.u32 %p119, %r1500, 7;
setp.lt.u32 %p120, %r1501, 7;
setp.lt.u32 %p121, %r1502, 7;
setp.lt.u32 %p122, %r1503, 7;
.loc 1 1171 27, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1171:27
sub.s32 %r1505, %r331, %r1394;
sub.s32 %r1506, %r330, %r1394;
sub.s32 %r1507, %r335, %r1394;
sub.s32 %r1508, %r334, %r1394;
sub.s32 %r1509, %r339, %r1394;
sub.s32 %r1510, %r338, %r1394;
sub.s32 %r1511, %r343, %r1394;
sub.s32 %r1512, %r342, %r1394;
sub.s32 %r1513, %r347, %r1394;
sub.s32 %r1514, %r346, %r1394;
sub.s32 %r1515, %r351, %r1394;
sub.s32 %r1516, %r350, %r1394;
sub.s32 %r1517, %r355, %r1394;
sub.s32 %r1518, %r354, %r1394;
sub.s32 %r1519, %r359, %r1394;
sub.s32 %r1520, %r358, %r1394;
sub.s32 %r1521, %r363, %r1394;
sub.s32 %r1522, %r362, %r1394;
sub.s32 %r1523, %r367, %r1394;
sub.s32 %r1524, %r366, %r1394;
sub.s32 %r1525, %r371, %r1394;
sub.s32 %r1526, %r370, %r1394;
sub.s32 %r1527, %r375, %r1394;
sub.s32 %r1528, %r374, %r1394;
sub.s32 %r1529, %r379, %r1394;
sub.s32 %r1530, %r378, %r1394;
sub.s32 %r1531, %r383, %r1394;
sub.s32 %r1532, %r382, %r1394;
sub.s32 %r1533, %r387, %r1394;
sub.s32 %r1534, %r386, %r1394;
sub.s32 %r1535, %r391, %r1394;
sub.s32 %r1536, %r390, %r1394;
.loc 1 1171 20, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1171:20
neg.s32 %r1537, %r1536;
max.s32 %r1538, %r1536, %r1537;
neg.s32 %r1539, %r1535;
max.s32 %r1540, %r1535, %r1539;
neg.s32 %r1541, %r1534;
max.s32 %r1542, %r1534, %r1541;
neg.s32 %r1543, %r1533;
max.s32 %r1544, %r1533, %r1543;
neg.s32 %r1545, %r1532;
max.s32 %r1546, %r1532, %r1545;
neg.s32 %r1547, %r1531;
max.s32 %r1548, %r1531, %r1547;
neg.s32 %r1549, %r1530;
max.s32 %r1550, %r1530, %r1549;
neg.s32 %r1551, %r1529;
max.s32 %r1552, %r1529, %r1551;
neg.s32 %r1553, %r1528;
max.s32 %r1554, %r1528, %r1553;
neg.s32 %r1555, %r1527;
max.s32 %r1556, %r1527, %r1555;
neg.s32 %r1557, %r1526;
max.s32 %r1558, %r1526, %r1557;
neg.s32 %r1559, %r1525;
max.s32 %r1560, %r1525, %r1559;
neg.s32 %r1561, %r1524;
max.s32 %r1562, %r1524, %r1561;
neg.s32 %r1563, %r1523;
max.s32 %r1564, %r1523, %r1563;
neg.s32 %r1565, %r1522;
max.s32 %r1566, %r1522, %r1565;
neg.s32 %r1567, %r1521;
max.s32 %r1568, %r1521, %r1567;
neg.s32 %r1569, %r1520;
max.s32 %r1570, %r1520, %r1569;
neg.s32 %r1571, %r1519;
max.s32 %r1572, %r1519, %r1571;
neg.s32 %r1573, %r1518;
max.s32 %r1574, %r1518, %r1573;
neg.s32 %r1575, %r1517;
max.s32 %r1576, %r1517, %r1575;
neg.s32 %r1577, %r1516;
max.s32 %r1578, %r1516, %r1577;
neg.s32 %r1579, %r1515;
max.s32 %r1580, %r1515, %r1579;
neg.s32 %r1581, %r1514;
max.s32 %r1582, %r1514, %r1581;
neg.s32 %r1583, %r1513;
max.s32 %r1584, %r1513, %r1583;
neg.s32 %r1585, %r1512;
max.s32 %r1586, %r1512, %r1585;
neg.s32 %r1587, %r1511;
max.s32 %r1588, %r1511, %r1587;
neg.s32 %r1589, %r1510;
max.s32 %r1590, %r1510, %r1589;
neg.s32 %r1591, %r1509;
max.s32 %r1592, %r1509, %r1591;
neg.s32 %r1593, %r1508;
max.s32 %r1594, %r1508, %r1593;
neg.s32 %r1595, %r1507;
max.s32 %r1596, %r1507, %r1595;
neg.s32 %r1597, %r1506;
max.s32 %r1598, %r1506, %r1597;
neg.s32 %r1599, %r1505;
max.s32 %r1600, %r1505, %r1599;
.loc 1 1172 12, function_name $L__info_string0, inlined_at 1 1193 25 // test_core.py:1172:12
max.s32 %r1601, %r1393, %r1600;
max.s32 %r1602, %r1393, %r1598;
max.s32 %r1603, %r1393, %r1596;
max.s32 %r1604, %r1393, %r1594;
max.s32 %r1605, %r1393, %r1592;
max.s32 %r1606, %r1393, %r1590;
max.s32 %r1607, %r1393, %r1588;
max.s32 %r1608, %r1393, %r1586;
max.s32 %r1609, %r1393, %r1584;
max.s32 %r1610, %r1393, %r1582;
max.s32 %r1611, %r1393, %r1580;
max.s32 %r1612, %r1393, %r1578;
max.s32 %r1613, %r1393, %r1576;
max.s32 %r1614, %r1393, %r1574;
max.s32 %r1615, %r1393, %r1572;
max.s32 %r1616, %r1393, %r1570;
max.s32 %r1617, %r1393, %r1568;
max.s32 %r1618, %r1393, %r1566;
max.s32 %r1619, %r1393, %r1564;
max.s32 %r1620, %r1393, %r1562;
max.s32 %r1621, %r1393, %r1560;
max.s32 %r1622, %r1393, %r1558;
max.s32 %r1623, %r1393, %r1556;
max.s32 %r1624, %r1393, %r1554;
max.s32 %r1625, %r1393, %r1552;
max.s32 %r1626, %r1393, %r1550;
max.s32 %r1627, %r1393, %r1548;
max.s32 %r1628, %r1393, %r1546;
max.s32 %r1629, %r1393, %r1544;
max.s32 %r1630, %r1393, %r1542;
max.s32 %r1631, %r1393, %r1540;
max.s32 %r1632, %r1393, %r1538;
setp.lt.u32 %p123, %r1632, 7;
selp.b16 %rs54, 1, 0, %p123;
shl.b16 %rs55, %rs54, 2;
setp.lt.u32 %p124, %r1631, 7;
selp.b16 %rs56, -1, 0, %p124;
shl.b16 %rs57, %rs56, 3;
or.b16 %rs58, %rs57, %rs55;
setp.lt.u32 %p125, %r1630, 7;
selp.b16 %rs59, 1, 0, %p125;
setp.lt.u32 %p126, %r1629, 7;
selp.b16 %rs60, -1, 0, %p126;
shl.b16 %rs61, %rs60, 1;
or.b16 %rs62, %rs59, %rs61;
and.b16 %rs63, %rs62, 3;
or.b16 %rs64, %rs63, %rs58;
and.b16 %rs65, %rs64, 15;
shl.b16 %rs66, %rs65, 8;
setp.lt.u32 %p127, %r1628, 7;
selp.b16 %rs67, 1, 0, %p127;
shl.b16 %rs68, %rs67, 2;
setp.lt.u32 %p128, %r1627, 7;
selp.b16 %rs69, -1, 0, %p128;
shl.b16 %rs70, %rs69, 3;
or.b16 %rs71, %rs70, %rs68;
setp.lt.u32 %p129, %r1626, 7;
selp.b16 %rs72, 1, 0, %p129;
setp.lt.u32 %p130, %r1625, 7;
selp.b16 %rs73, -1, 0, %p130;
shl.b16 %rs74, %rs73, 1;
or.b16 %rs75, %rs72, %rs74;
and.b16 %rs76, %rs75, 3;
or.b16 %rs77, %rs76, %rs71;
shl.b16 %rs78, %rs77, 12;
or.b16 %rs79, %rs78, %rs66;
setp.lt.u32 %p131, %r1624, 7;
selp.b16 %rs80, 1, 0, %p131;
shl.b16 %rs81, %rs80, 2;
setp.lt.u32 %p132, %r1623, 7;
selp.b16 %rs82, -1, 0, %p132;
shl.b16 %rs83, %rs82, 3;
or.b16 %rs84, %rs83, %rs81;
setp.lt.u32 %p133, %r1622, 7;
selp.b16 %rs85, 1, 0, %p133;
setp.lt.u32 %p134, %r1621, 7;
selp.b16 %rs86, -1, 0, %p134;
shl.b16 %rs87, %rs86, 1;
or.b16 %rs88, %rs85, %rs87;
and.b16 %rs89, %rs88, 3;
or.b16 %rs90, %rs89, %rs84;
and.b16 %rs91, %rs90, 15;
setp.lt.u32 %p135, %r1620, 7;
selp.b16 %rs92, 1, 0, %p135;
shl.b16 %rs93, %rs92, 2;
setp.lt.u32 %p136, %r1619, 7;
selp.b16 %rs94, -1, 0, %p136;
shl.b16 %rs95, %rs94, 3;
or.b16 %rs96, %rs95, %rs93;
setp.lt.u32 %p137, %r1618, 7;
selp.b16 %rs97, 1, 0, %p137;
setp.lt.u32 %p138, %r1617, 7;
selp.b16 %rs98, -1, 0, %p138;
shl.b16 %rs99, %rs98, 1;
or.b16 %rs100, %rs97, %rs99;
and.b16 %rs101, %rs100, 3;
or.b16 %rs102, %rs101, %rs96;
shl.b16 %rs103, %rs102, 4;
or.b16 %rs104, %rs91, %rs103;
and.b16 %rs105, %rs104, 255;
or.b16 %rs106, %rs105, %rs79;
cvt.u32.u16 %r1633, %rs106;
setp.lt.u32 %p139, %r1616, 7;
setp.lt.u32 %p140, %r1615, 7;
setp.lt.u32 %p141, %r1614, 7;
setp.lt.u32 %p142, %r1613, 7;
setp.lt.u32 %p143, %r1612, 7;
setp.lt.u32 %p144, %r1611, 7;
setp.lt.u32 %p145, %r1610, 7;
setp.lt.u32 %p146, %r1609, 7;
setp.lt.u32 %p147, %r1608, 7;
setp.lt.u32 %p148, %r1607, 7;
setp.lt.u32 %p149, %r1606, 7;
setp.lt.u32 %p150, %r1605, 7;
setp.lt.u32 %p151, %r1604, 7;
setp.lt.u32 %p152, %r1603, 7;
setp.lt.u32 %p153, %r1602, 7;
setp.lt.u32 %p154, %r1601, 7;
setp.lt.u32 %p155, %r1504, 7;
$L__tmp584:
.loc 1 1193 25 // test_core.py:1193:25
selp.f32 %r1634, 0f3F800000, 0f00000000, %p96;
selp.f32 %r1635, 0f3F800000, 0f00000000, %p97;
selp.f32 %r1636, 0f3F800000, 0f00000000, %p98;
selp.f32 %r1637, 0f3F800000, 0f00000000, %p99;
selp.f32 %r1638, 0f3F800000, 0f00000000, %p100;
selp.f32 %r1639, 0f3F800000, 0f00000000, %p101;
selp.f32 %r1640, 0f3F800000, 0f00000000, %p102;
selp.f32 %r1641, 0f3F800000, 0f00000000, %p103;
selp.f32 %r1642, 0f3F800000, 0f00000000, %p104;
selp.f32 %r1643, 0f3F800000, 0f00000000, %p105;
selp.f32 %r1644, 0f3F800000, 0f00000000, %p106;
selp.f32 %r1645, 0f3F800000, 0f00000000, %p107;
selp.f32 %r1646, 0f3F800000, 0f00000000, %p108;
selp.f32 %r1647, 0f3F800000, 0f00000000, %p109;
selp.f32 %r1648, 0f3F800000, 0f00000000, %p110;
selp.f32 %r1649, 0f3F800000, 0f00000000, %p111;
selp.f32 %r1650, 0f3F800000, 0f00000000, %p112;
selp.f32 %r1651, 0f3F800000, 0f00000000, %p113;
selp.f32 %r1652, 0f3F800000, 0f00000000, %p114;
selp.f32 %r1653, 0f3F800000, 0f00000000, %p115;
selp.f32 %r1654, 0f3F800000, 0f00000000, %p116;
selp.f32 %r1655, 0f3F800000, 0f00000000, %p117;
selp.f32 %r1656, 0f3F800000, 0f00000000, %p118;
selp.f32 %r1657, 0f3F800000, 0f00000000, %p119;
selp.f32 %r1658, 0f3F800000, 0f00000000, %p120;
selp.f32 %r1659, 0f3F800000, 0f00000000, %p121;
selp.f32 %r1660, 0f3F800000, 0f00000000, %p122;
selp.f32 %r1661, 0f3F800000, 0f00000000, %p154;
selp.f32 %r1662, 0f3F800000, 0f00000000, %p153;
selp.f32 %r1663, 0f3F800000, 0f00000000, %p152;
selp.f32 %r1664, 0f3F800000, 0f00000000, %p151;
selp.f32 %r1665, 0f3F800000, 0f00000000, %p150;
selp.f32 %r1666, 0f3F800000, 0f00000000, %p149;
selp.f32 %r1667, 0f3F800000, 0f00000000, %p148;
selp.f32 %r1668, 0f3F800000, 0f00000000, %p147;
selp.f32 %r1669, 0f3F800000, 0f00000000, %p146;
selp.f32 %r1670, 0f3F800000, 0f00000000, %p145;
selp.f32 %r1671, 0f3F800000, 0f00000000, %p144;
selp.f32 %r1672, 0f3F800000, 0f00000000, %p143;
selp.f32 %r1673, 0f3F800000, 0f00000000, %p142;
selp.f32 %r1674, 0f3F800000, 0f00000000, %p141;
selp.f32 %r1675, 0f3F800000, 0f00000000, %p140;
selp.f32 %r1676, 0f3F800000, 0f00000000, %p139;
shr.u32 %r1677, %r1633, 15;
and.b32 %r1678, %r1677, 1;
setp.ne.b32 %p156, %r1678, 0;
selp.b32 %r1679, 1, 0, %p156;
cvt.rn.f32.u32 %r1680, %r1679;
shr.u32 %r1681, %r1633, 14;
and.b32 %r1682, %r1681, 1;
setp.ne.b32 %p157, %r1682, 0;
selp.b32 %r1683, 1, 0, %p157;
cvt.rn.f32.u32 %r1684, %r1683;
shr.u32 %r1685, %r1633, 13;
and.b32 %r1686, %r1685, 1;
setp.ne.b32 %p158, %r1686, 0;
selp.b32 %r1687, 1, 0, %p158;
cvt.rn.f32.u32 %r1688, %r1687;
shr.u32 %r1689, %r1633, 12;
and.b32 %r1690, %r1689, 1;
setp.ne.b32 %p159, %r1690, 0;
selp.b32 %r1691, 1, 0, %p159;
cvt.rn.f32.u32 %r1692, %r1691;
shr.u32 %r1693, %r1633, 11;
and.b32 %r1694, %r1693, 1;
setp.ne.b32 %p160, %r1694, 0;
selp.b32 %r1695, 1, 0, %p160;
cvt.rn.f32.u32 %r1696, %r1695;
shr.u32 %r1697, %r1633, 10;
and.b32 %r1698, %r1697, 1;
setp.ne.b32 %p161, %r1698, 0;
selp.b32 %r1699, 1, 0, %p161;
cvt.rn.f32.u32 %r1700, %r1699;
shr.u32 %r1701, %r1633, 9;
and.b32 %r1702, %r1701, 1;
setp.ne.b32 %p162, %r1702, 0;
selp.b32 %r1703, 1, 0, %p162;
cvt.rn.f32.u32 %r1704, %r1703;
shr.u32 %r1705, %r1633, 8;
and.b32 %r1706, %r1705, 1;
setp.ne.b32 %p163, %r1706, 0;
selp.b32 %r1707, 1, 0, %p163;
cvt.rn.f32.u32 %r1708, %r1707;
shr.u32 %r1709, %r1633, 7;
and.b32 %r1710, %r1709, 1;
setp.ne.b32 %p164, %r1710, 0;
selp.b32 %r1711, 1, 0, %p164;
cvt.rn.f32.u32 %r1712, %r1711;
shr.u32 %r1713, %r1633, 6;
and.b32 %r1714, %r1713, 1;
setp.ne.b32 %p165, %r1714, 0;
selp.b32 %r1715, 1, 0, %p165;
cvt.rn.f32.u32 %r1716, %r1715;
shr.u32 %r1717, %r1633, 5;
and.b32 %r1718, %r1717, 1;
setp.ne.b32 %p166, %r1718, 0;
selp.b32 %r1719, 1, 0, %p166;
cvt.rn.f32.u32 %r1720, %r1719;
shr.u32 %r1721, %r1633, 4;
and.b32 %r1722, %r1721, 1;
setp.ne.b32 %p167, %r1722, 0;
selp.b32 %r1723, 1, 0, %p167;
cvt.rn.f32.u32 %r1724, %r1723;
shr.u32 %r1725, %r1633, 3;
and.b32 %r1726, %r1725, 1;
setp.ne.b32 %p168, %r1726, 0;
selp.b32 %r1727, 1, 0, %p168;
cvt.rn.f32.u32 %r1728, %r1727;
shr.u32 %r1729, %r1633, 2;
and.b32 %r1730, %r1729, 1;
setp.ne.b32 %p169, %r1730, 0;
selp.b32 %r1731, 1, 0, %p169;
cvt.rn.f32.u32 %r1732, %r1731;
shr.u32 %r1733, %r1633, 1;
and.b32 %r1734, %r1733, 1;
setp.ne.b32 %p170, %r1734, 0;
selp.b32 %r1735, 1, 0, %p170;
cvt.rn.f32.u32 %r1736, %r1735;
selp.f32 %r1737, 0f3F800000, 0f00000000, %p133;
selp.f32 %r1738, 0f3F800000, 0f00000000, %p155;
$L__tmp585:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1739, %r1634, 16, 31, -1;
$L__tmp586:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1740, %r1634, %r1739;
$L__tmp587:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1741, %r1740, 8, 31, -1;
$L__tmp588:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1742, %r1740, %r1741;
$L__tmp589:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1743, %r1742, 4, 31, -1;
$L__tmp590:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1744, %r1742, %r1743;
$L__tmp591:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1745, %r1744, 2, 31, -1;
$L__tmp592:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1746, %r1744, %r1745;
$L__tmp593:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1747, %r1746, 1, 31, -1;
shfl.sync.bfly.b32 %r1748, %r1634, 16, 31, -1;
$L__tmp594:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1749, %r1634, %r1748;
$L__tmp595:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1750, %r1749, 8, 31, -1;
$L__tmp596:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1751, %r1749, %r1750;
$L__tmp597:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1752, %r1751, 4, 31, -1;
$L__tmp598:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1753, %r1751, %r1752;
$L__tmp599:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1754, %r1753, 2, 31, -1;
$L__tmp600:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1755, %r1753, %r1754;
$L__tmp601:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1756, %r1755, 1, 31, -1;
shfl.sync.bfly.b32 %r1757, %r1634, 16, 31, -1;
$L__tmp602:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1758, %r1634, %r1757;
$L__tmp603:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1759, %r1758, 8, 31, -1;
$L__tmp604:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1760, %r1758, %r1759;
$L__tmp605:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1761, %r1760, 4, 31, -1;
$L__tmp606:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1762, %r1760, %r1761;
$L__tmp607:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1763, %r1762, 2, 31, -1;
$L__tmp608:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1764, %r1762, %r1763;
$L__tmp609:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1765, %r1764, 1, 31, -1;
shfl.sync.bfly.b32 %r1766, %r1635, 16, 31, -1;
$L__tmp610:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1767, %r1635, %r1766;
$L__tmp611:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1768, %r1767, 8, 31, -1;
$L__tmp612:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1769, %r1767, %r1768;
$L__tmp613:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1770, %r1769, 4, 31, -1;
$L__tmp614:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1771, %r1769, %r1770;
$L__tmp615:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1772, %r1771, 2, 31, -1;
$L__tmp616:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1773, %r1771, %r1772;
$L__tmp617:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1774, %r1773, 1, 31, -1;
mov.b64 %rd324, {%r1765, %r1774};
mov.b64 %rd325, {%r1747, %r1756};
$L__tmp618:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd326, {%r1764, %r1773};
mov.b64 %rd327, {%r1746, %r1755};
add.f32x2 %rd328, %rd327, %rd325;
add.f32x2 %rd329, %rd326, %rd324;
$L__tmp619:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1775, %r1636, 16, 31, -1;
$L__tmp620:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1776, %r1636, %r1775;
$L__tmp621:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1777, %r1776, 8, 31, -1;
$L__tmp622:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1778, %r1776, %r1777;
$L__tmp623:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1779, %r1778, 4, 31, -1;
$L__tmp624:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1780, %r1778, %r1779;
$L__tmp625:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1781, %r1780, 2, 31, -1;
$L__tmp626:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1782, %r1780, %r1781;
$L__tmp627:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1783, %r1782, 1, 31, -1;
shfl.sync.bfly.b32 %r1784, %r1637, 16, 31, -1;
$L__tmp628:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1785, %r1637, %r1784;
$L__tmp629:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1786, %r1785, 8, 31, -1;
$L__tmp630:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1787, %r1785, %r1786;
$L__tmp631:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1788, %r1787, 4, 31, -1;
$L__tmp632:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1789, %r1787, %r1788;
$L__tmp633:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1790, %r1789, 2, 31, -1;
$L__tmp634:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1791, %r1789, %r1790;
$L__tmp635:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1792, %r1791, 1, 31, -1;
shfl.sync.bfly.b32 %r1793, %r1638, 16, 31, -1;
$L__tmp636:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1794, %r1638, %r1793;
$L__tmp637:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1795, %r1794, 8, 31, -1;
$L__tmp638:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1796, %r1794, %r1795;
$L__tmp639:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1797, %r1796, 4, 31, -1;
$L__tmp640:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1798, %r1796, %r1797;
$L__tmp641:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1799, %r1798, 2, 31, -1;
$L__tmp642:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1800, %r1798, %r1799;
$L__tmp643:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1801, %r1800, 1, 31, -1;
shfl.sync.bfly.b32 %r1802, %r1639, 16, 31, -1;
$L__tmp644:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1803, %r1639, %r1802;
$L__tmp645:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1804, %r1803, 8, 31, -1;
$L__tmp646:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1805, %r1803, %r1804;
$L__tmp647:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1806, %r1805, 4, 31, -1;
$L__tmp648:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1807, %r1805, %r1806;
$L__tmp649:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1808, %r1807, 2, 31, -1;
$L__tmp650:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1809, %r1807, %r1808;
$L__tmp651:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1810, %r1809, 1, 31, -1;
mov.b64 %rd330, {%r1801, %r1810};
mov.b64 %rd331, {%r1783, %r1792};
$L__tmp652:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd332, {%r1800, %r1809};
mov.b64 %rd333, {%r1782, %r1791};
add.f32x2 %rd334, %rd333, %rd331;
add.f32x2 %rd335, %rd332, %rd330;
$L__tmp653:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1811, %r1640, 16, 31, -1;
$L__tmp654:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1812, %r1640, %r1811;
$L__tmp655:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1813, %r1812, 8, 31, -1;
$L__tmp656:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1814, %r1812, %r1813;
$L__tmp657:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1815, %r1814, 4, 31, -1;
$L__tmp658:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1816, %r1814, %r1815;
$L__tmp659:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1817, %r1816, 2, 31, -1;
$L__tmp660:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1818, %r1816, %r1817;
$L__tmp661:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1819, %r1818, 1, 31, -1;
shfl.sync.bfly.b32 %r1820, %r1641, 16, 31, -1;
$L__tmp662:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1821, %r1641, %r1820;
$L__tmp663:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1822, %r1821, 8, 31, -1;
$L__tmp664:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1823, %r1821, %r1822;
$L__tmp665:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1824, %r1823, 4, 31, -1;
$L__tmp666:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1825, %r1823, %r1824;
$L__tmp667:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1826, %r1825, 2, 31, -1;
$L__tmp668:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1827, %r1825, %r1826;
$L__tmp669:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1828, %r1827, 1, 31, -1;
shfl.sync.bfly.b32 %r1829, %r1642, 16, 31, -1;
$L__tmp670:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1830, %r1642, %r1829;
$L__tmp671:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1831, %r1830, 8, 31, -1;
$L__tmp672:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1832, %r1830, %r1831;
$L__tmp673:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1833, %r1832, 4, 31, -1;
$L__tmp674:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1834, %r1832, %r1833;
$L__tmp675:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1835, %r1834, 2, 31, -1;
$L__tmp676:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1836, %r1834, %r1835;
$L__tmp677:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1837, %r1836, 1, 31, -1;
shfl.sync.bfly.b32 %r1838, %r1643, 16, 31, -1;
$L__tmp678:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1839, %r1643, %r1838;
$L__tmp679:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1840, %r1839, 8, 31, -1;
$L__tmp680:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1841, %r1839, %r1840;
$L__tmp681:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1842, %r1841, 4, 31, -1;
$L__tmp682:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1843, %r1841, %r1842;
$L__tmp683:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1844, %r1843, 2, 31, -1;
$L__tmp684:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1845, %r1843, %r1844;
$L__tmp685:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1846, %r1845, 1, 31, -1;
mov.b64 %rd336, {%r1837, %r1846};
mov.b64 %rd337, {%r1819, %r1828};
$L__tmp686:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd338, {%r1836, %r1845};
mov.b64 %rd339, {%r1818, %r1827};
add.f32x2 %rd340, %rd339, %rd337;
add.f32x2 %rd341, %rd338, %rd336;
$L__tmp687:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1847, %r1644, 16, 31, -1;
$L__tmp688:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1848, %r1644, %r1847;
$L__tmp689:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1849, %r1848, 8, 31, -1;
$L__tmp690:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1850, %r1848, %r1849;
$L__tmp691:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1851, %r1850, 4, 31, -1;
$L__tmp692:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1852, %r1850, %r1851;
$L__tmp693:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1853, %r1852, 2, 31, -1;
$L__tmp694:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1854, %r1852, %r1853;
$L__tmp695:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1855, %r1854, 1, 31, -1;
shfl.sync.bfly.b32 %r1856, %r1645, 16, 31, -1;
$L__tmp696:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1857, %r1645, %r1856;
$L__tmp697:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1858, %r1857, 8, 31, -1;
$L__tmp698:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1859, %r1857, %r1858;
$L__tmp699:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1860, %r1859, 4, 31, -1;
$L__tmp700:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1861, %r1859, %r1860;
$L__tmp701:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1862, %r1861, 2, 31, -1;
$L__tmp702:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1863, %r1861, %r1862;
$L__tmp703:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1864, %r1863, 1, 31, -1;
shfl.sync.bfly.b32 %r1865, %r1646, 16, 31, -1;
$L__tmp704:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1866, %r1646, %r1865;
$L__tmp705:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1867, %r1866, 8, 31, -1;
$L__tmp706:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1868, %r1866, %r1867;
$L__tmp707:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1869, %r1868, 4, 31, -1;
$L__tmp708:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1870, %r1868, %r1869;
$L__tmp709:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1871, %r1870, 2, 31, -1;
$L__tmp710:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1872, %r1870, %r1871;
$L__tmp711:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1873, %r1872, 1, 31, -1;
shfl.sync.bfly.b32 %r1874, %r1647, 16, 31, -1;
$L__tmp712:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1875, %r1647, %r1874;
$L__tmp713:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1876, %r1875, 8, 31, -1;
$L__tmp714:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1877, %r1875, %r1876;
$L__tmp715:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1878, %r1877, 4, 31, -1;
$L__tmp716:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1879, %r1877, %r1878;
$L__tmp717:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1880, %r1879, 2, 31, -1;
$L__tmp718:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1881, %r1879, %r1880;
$L__tmp719:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1882, %r1881, 1, 31, -1;
mov.b64 %rd342, {%r1873, %r1882};
mov.b64 %rd343, {%r1855, %r1864};
$L__tmp720:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd344, {%r1872, %r1881};
mov.b64 %rd345, {%r1854, %r1863};
add.f32x2 %rd346, %rd345, %rd343;
add.f32x2 %rd347, %rd344, %rd342;
$L__tmp721:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1883, %r1648, 16, 31, -1;
$L__tmp722:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1884, %r1648, %r1883;
$L__tmp723:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1885, %r1884, 8, 31, -1;
$L__tmp724:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1886, %r1884, %r1885;
$L__tmp725:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1887, %r1886, 4, 31, -1;
$L__tmp726:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1888, %r1886, %r1887;
$L__tmp727:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1889, %r1888, 2, 31, -1;
$L__tmp728:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1890, %r1888, %r1889;
$L__tmp729:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1891, %r1890, 1, 31, -1;
shfl.sync.bfly.b32 %r1892, %r1649, 16, 31, -1;
$L__tmp730:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1893, %r1649, %r1892;
$L__tmp731:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1894, %r1893, 8, 31, -1;
$L__tmp732:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1895, %r1893, %r1894;
$L__tmp733:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1896, %r1895, 4, 31, -1;
$L__tmp734:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1897, %r1895, %r1896;
$L__tmp735:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1898, %r1897, 2, 31, -1;
$L__tmp736:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1899, %r1897, %r1898;
$L__tmp737:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1900, %r1899, 1, 31, -1;
shfl.sync.bfly.b32 %r1901, %r1650, 16, 31, -1;
$L__tmp738:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1902, %r1650, %r1901;
$L__tmp739:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1903, %r1902, 8, 31, -1;
$L__tmp740:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1904, %r1902, %r1903;
$L__tmp741:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1905, %r1904, 4, 31, -1;
$L__tmp742:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1906, %r1904, %r1905;
$L__tmp743:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1907, %r1906, 2, 31, -1;
$L__tmp744:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1908, %r1906, %r1907;
$L__tmp745:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1909, %r1908, 1, 31, -1;
shfl.sync.bfly.b32 %r1910, %r1651, 16, 31, -1;
$L__tmp746:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1911, %r1651, %r1910;
$L__tmp747:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1912, %r1911, 8, 31, -1;
$L__tmp748:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1913, %r1911, %r1912;
$L__tmp749:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1914, %r1913, 4, 31, -1;
$L__tmp750:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1915, %r1913, %r1914;
$L__tmp751:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1916, %r1915, 2, 31, -1;
$L__tmp752:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1917, %r1915, %r1916;
$L__tmp753:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1918, %r1917, 1, 31, -1;
mov.b64 %rd348, {%r1909, %r1918};
mov.b64 %rd349, {%r1891, %r1900};
$L__tmp754:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd350, {%r1908, %r1917};
mov.b64 %rd351, {%r1890, %r1899};
add.f32x2 %rd352, %rd351, %rd349;
add.f32x2 %rd353, %rd350, %rd348;
$L__tmp755:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1919, %r1652, 16, 31, -1;
$L__tmp756:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1920, %r1652, %r1919;
$L__tmp757:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1921, %r1920, 8, 31, -1;
$L__tmp758:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1922, %r1920, %r1921;
$L__tmp759:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1923, %r1922, 4, 31, -1;
$L__tmp760:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1924, %r1922, %r1923;
$L__tmp761:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1925, %r1924, 2, 31, -1;
$L__tmp762:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1926, %r1924, %r1925;
$L__tmp763:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1927, %r1926, 1, 31, -1;
shfl.sync.bfly.b32 %r1928, %r1653, 16, 31, -1;
$L__tmp764:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1929, %r1653, %r1928;
$L__tmp765:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1930, %r1929, 8, 31, -1;
$L__tmp766:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1931, %r1929, %r1930;
$L__tmp767:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1932, %r1931, 4, 31, -1;
$L__tmp768:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1933, %r1931, %r1932;
$L__tmp769:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1934, %r1933, 2, 31, -1;
$L__tmp770:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1935, %r1933, %r1934;
$L__tmp771:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1936, %r1935, 1, 31, -1;
shfl.sync.bfly.b32 %r1937, %r1654, 16, 31, -1;
$L__tmp772:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1938, %r1654, %r1937;
$L__tmp773:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1939, %r1938, 8, 31, -1;
$L__tmp774:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1940, %r1938, %r1939;
$L__tmp775:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1941, %r1940, 4, 31, -1;
$L__tmp776:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1942, %r1940, %r1941;
$L__tmp777:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1943, %r1942, 2, 31, -1;
$L__tmp778:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1944, %r1942, %r1943;
$L__tmp779:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1945, %r1944, 1, 31, -1;
shfl.sync.bfly.b32 %r1946, %r1655, 16, 31, -1;
$L__tmp780:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1947, %r1655, %r1946;
$L__tmp781:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1948, %r1947, 8, 31, -1;
$L__tmp782:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1949, %r1947, %r1948;
$L__tmp783:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1950, %r1949, 4, 31, -1;
$L__tmp784:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1951, %r1949, %r1950;
$L__tmp785:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1952, %r1951, 2, 31, -1;
$L__tmp786:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1953, %r1951, %r1952;
$L__tmp787:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1954, %r1953, 1, 31, -1;
mov.b64 %rd354, {%r1945, %r1954};
mov.b64 %rd355, {%r1927, %r1936};
$L__tmp788:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd356, {%r1944, %r1953};
mov.b64 %rd357, {%r1926, %r1935};
add.f32x2 %rd358, %rd357, %rd355;
add.f32x2 %rd359, %rd356, %rd354;
$L__tmp789:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1955, %r1656, 16, 31, -1;
$L__tmp790:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1956, %r1656, %r1955;
$L__tmp791:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1957, %r1956, 8, 31, -1;
$L__tmp792:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1958, %r1956, %r1957;
$L__tmp793:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1959, %r1958, 4, 31, -1;
$L__tmp794:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1960, %r1958, %r1959;
$L__tmp795:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1961, %r1960, 2, 31, -1;
$L__tmp796:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1962, %r1960, %r1961;
$L__tmp797:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1963, %r1962, 1, 31, -1;
shfl.sync.bfly.b32 %r1964, %r1657, 16, 31, -1;
$L__tmp798:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1965, %r1657, %r1964;
$L__tmp799:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1966, %r1965, 8, 31, -1;
$L__tmp800:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1967, %r1965, %r1966;
$L__tmp801:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1968, %r1967, 4, 31, -1;
$L__tmp802:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1969, %r1967, %r1968;
$L__tmp803:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1970, %r1969, 2, 31, -1;
$L__tmp804:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1971, %r1969, %r1970;
$L__tmp805:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1972, %r1971, 1, 31, -1;
shfl.sync.bfly.b32 %r1973, %r1658, 16, 31, -1;
$L__tmp806:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1974, %r1658, %r1973;
$L__tmp807:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1975, %r1974, 8, 31, -1;
$L__tmp808:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1976, %r1974, %r1975;
$L__tmp809:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1977, %r1976, 4, 31, -1;
$L__tmp810:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1978, %r1976, %r1977;
$L__tmp811:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1979, %r1978, 2, 31, -1;
$L__tmp812:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1980, %r1978, %r1979;
$L__tmp813:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1981, %r1980, 1, 31, -1;
shfl.sync.bfly.b32 %r1982, %r1659, 16, 31, -1;
$L__tmp814:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1983, %r1659, %r1982;
$L__tmp815:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1984, %r1983, 8, 31, -1;
$L__tmp816:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1985, %r1983, %r1984;
$L__tmp817:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1986, %r1985, 4, 31, -1;
$L__tmp818:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1987, %r1985, %r1986;
$L__tmp819:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1988, %r1987, 2, 31, -1;
$L__tmp820:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1989, %r1987, %r1988;
$L__tmp821:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1990, %r1989, 1, 31, -1;
mov.b64 %rd360, {%r1981, %r1990};
mov.b64 %rd361, {%r1963, %r1972};
$L__tmp822:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd362, {%r1980, %r1989};
mov.b64 %rd363, {%r1962, %r1971};
add.f32x2 %rd364, %rd363, %rd361;
add.f32x2 %rd365, %rd362, %rd360;
$L__tmp823:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1991, %r1660, 16, 31, -1;
$L__tmp824:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1992, %r1660, %r1991;
$L__tmp825:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1993, %r1992, 8, 31, -1;
$L__tmp826:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1994, %r1992, %r1993;
$L__tmp827:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1995, %r1994, 4, 31, -1;
$L__tmp828:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1996, %r1994, %r1995;
$L__tmp829:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1997, %r1996, 2, 31, -1;
$L__tmp830:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r1998, %r1996, %r1997;
$L__tmp831:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r1999, %r1998, 1, 31, -1;
shfl.sync.bfly.b32 %r2000, %r1661, 16, 31, -1;
$L__tmp832:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2001, %r1661, %r2000;
$L__tmp833:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2002, %r2001, 8, 31, -1;
$L__tmp834:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2003, %r2001, %r2002;
$L__tmp835:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2004, %r2003, 4, 31, -1;
$L__tmp836:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2005, %r2003, %r2004;
$L__tmp837:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2006, %r2005, 2, 31, -1;
$L__tmp838:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2007, %r2005, %r2006;
$L__tmp839:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2008, %r2007, 1, 31, -1;
shfl.sync.bfly.b32 %r2009, %r1662, 16, 31, -1;
$L__tmp840:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2010, %r1662, %r2009;
$L__tmp841:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2011, %r2010, 8, 31, -1;
$L__tmp842:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2012, %r2010, %r2011;
$L__tmp843:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2013, %r2012, 4, 31, -1;
$L__tmp844:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2014, %r2012, %r2013;
$L__tmp845:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2015, %r2014, 2, 31, -1;
$L__tmp846:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2016, %r2014, %r2015;
$L__tmp847:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2017, %r2016, 1, 31, -1;
shfl.sync.bfly.b32 %r2018, %r1663, 16, 31, -1;
$L__tmp848:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2019, %r1663, %r2018;
$L__tmp849:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2020, %r2019, 8, 31, -1;
$L__tmp850:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2021, %r2019, %r2020;
$L__tmp851:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2022, %r2021, 4, 31, -1;
$L__tmp852:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2023, %r2021, %r2022;
$L__tmp853:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2024, %r2023, 2, 31, -1;
$L__tmp854:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2025, %r2023, %r2024;
$L__tmp855:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2026, %r2025, 1, 31, -1;
mov.b64 %rd366, {%r2017, %r2026};
mov.b64 %rd367, {%r1999, %r2008};
$L__tmp856:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd368, {%r2016, %r2025};
mov.b64 %rd369, {%r1998, %r2007};
add.f32x2 %rd370, %rd369, %rd367;
add.f32x2 %rd371, %rd368, %rd366;
$L__tmp857:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2027, %r1664, 16, 31, -1;
$L__tmp858:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2028, %r1664, %r2027;
$L__tmp859:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2029, %r2028, 8, 31, -1;
$L__tmp860:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2030, %r2028, %r2029;
$L__tmp861:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2031, %r2030, 4, 31, -1;
$L__tmp862:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2032, %r2030, %r2031;
$L__tmp863:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2033, %r2032, 2, 31, -1;
$L__tmp864:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2034, %r2032, %r2033;
$L__tmp865:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2035, %r2034, 1, 31, -1;
shfl.sync.bfly.b32 %r2036, %r1665, 16, 31, -1;
$L__tmp866:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2037, %r1665, %r2036;
$L__tmp867:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2038, %r2037, 8, 31, -1;
$L__tmp868:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2039, %r2037, %r2038;
$L__tmp869:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2040, %r2039, 4, 31, -1;
$L__tmp870:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2041, %r2039, %r2040;
$L__tmp871:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2042, %r2041, 2, 31, -1;
$L__tmp872:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2043, %r2041, %r2042;
$L__tmp873:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2044, %r2043, 1, 31, -1;
shfl.sync.bfly.b32 %r2045, %r1666, 16, 31, -1;
$L__tmp874:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2046, %r1666, %r2045;
$L__tmp875:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2047, %r2046, 8, 31, -1;
$L__tmp876:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2048, %r2046, %r2047;
$L__tmp877:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2049, %r2048, 4, 31, -1;
$L__tmp878:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2050, %r2048, %r2049;
$L__tmp879:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2051, %r2050, 2, 31, -1;
$L__tmp880:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2052, %r2050, %r2051;
$L__tmp881:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2053, %r2052, 1, 31, -1;
shfl.sync.bfly.b32 %r2054, %r1667, 16, 31, -1;
$L__tmp882:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2055, %r1667, %r2054;
$L__tmp883:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2056, %r2055, 8, 31, -1;
$L__tmp884:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2057, %r2055, %r2056;
$L__tmp885:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2058, %r2057, 4, 31, -1;
$L__tmp886:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2059, %r2057, %r2058;
$L__tmp887:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2060, %r2059, 2, 31, -1;
$L__tmp888:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2061, %r2059, %r2060;
$L__tmp889:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2062, %r2061, 1, 31, -1;
mov.b64 %rd372, {%r2053, %r2062};
mov.b64 %rd373, {%r2035, %r2044};
$L__tmp890:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd374, {%r2052, %r2061};
mov.b64 %rd375, {%r2034, %r2043};
add.f32x2 %rd376, %rd375, %rd373;
add.f32x2 %rd377, %rd374, %rd372;
$L__tmp891:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2063, %r1668, 16, 31, -1;
$L__tmp892:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2064, %r1668, %r2063;
$L__tmp893:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2065, %r2064, 8, 31, -1;
$L__tmp894:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2066, %r2064, %r2065;
$L__tmp895:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2067, %r2066, 4, 31, -1;
$L__tmp896:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2068, %r2066, %r2067;
$L__tmp897:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2069, %r2068, 2, 31, -1;
$L__tmp898:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2070, %r2068, %r2069;
$L__tmp899:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2071, %r2070, 1, 31, -1;
shfl.sync.bfly.b32 %r2072, %r1669, 16, 31, -1;
$L__tmp900:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2073, %r1669, %r2072;
$L__tmp901:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2074, %r2073, 8, 31, -1;
$L__tmp902:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2075, %r2073, %r2074;
$L__tmp903:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2076, %r2075, 4, 31, -1;
$L__tmp904:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2077, %r2075, %r2076;
$L__tmp905:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2078, %r2077, 2, 31, -1;
$L__tmp906:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2079, %r2077, %r2078;
$L__tmp907:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2080, %r2079, 1, 31, -1;
shfl.sync.bfly.b32 %r2081, %r1670, 16, 31, -1;
$L__tmp908:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2082, %r1670, %r2081;
$L__tmp909:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2083, %r2082, 8, 31, -1;
$L__tmp910:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2084, %r2082, %r2083;
$L__tmp911:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2085, %r2084, 4, 31, -1;
$L__tmp912:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2086, %r2084, %r2085;
$L__tmp913:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2087, %r2086, 2, 31, -1;
$L__tmp914:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2088, %r2086, %r2087;
$L__tmp915:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2089, %r2088, 1, 31, -1;
shfl.sync.bfly.b32 %r2090, %r1671, 16, 31, -1;
$L__tmp916:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2091, %r1671, %r2090;
$L__tmp917:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2092, %r2091, 8, 31, -1;
$L__tmp918:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2093, %r2091, %r2092;
$L__tmp919:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2094, %r2093, 4, 31, -1;
$L__tmp920:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2095, %r2093, %r2094;
$L__tmp921:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2096, %r2095, 2, 31, -1;
$L__tmp922:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2097, %r2095, %r2096;
$L__tmp923:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2098, %r2097, 1, 31, -1;
mov.b64 %rd378, {%r2089, %r2098};
mov.b64 %rd379, {%r2071, %r2080};
$L__tmp924:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd380, {%r2088, %r2097};
mov.b64 %rd381, {%r2070, %r2079};
add.f32x2 %rd382, %rd381, %rd379;
add.f32x2 %rd383, %rd380, %rd378;
$L__tmp925:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2099, %r1672, 16, 31, -1;
$L__tmp926:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2100, %r1672, %r2099;
$L__tmp927:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2101, %r2100, 8, 31, -1;
$L__tmp928:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2102, %r2100, %r2101;
$L__tmp929:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2103, %r2102, 4, 31, -1;
$L__tmp930:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2104, %r2102, %r2103;
$L__tmp931:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2105, %r2104, 2, 31, -1;
$L__tmp932:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2106, %r2104, %r2105;
$L__tmp933:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2107, %r2106, 1, 31, -1;
shfl.sync.bfly.b32 %r2108, %r1673, 16, 31, -1;
$L__tmp934:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2109, %r1673, %r2108;
$L__tmp935:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2110, %r2109, 8, 31, -1;
$L__tmp936:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2111, %r2109, %r2110;
$L__tmp937:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2112, %r2111, 4, 31, -1;
$L__tmp938:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2113, %r2111, %r2112;
$L__tmp939:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2114, %r2113, 2, 31, -1;
$L__tmp940:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2115, %r2113, %r2114;
$L__tmp941:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2116, %r2115, 1, 31, -1;
shfl.sync.bfly.b32 %r2117, %r1674, 16, 31, -1;
$L__tmp942:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2118, %r1674, %r2117;
$L__tmp943:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2119, %r2118, 8, 31, -1;
$L__tmp944:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2120, %r2118, %r2119;
$L__tmp945:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2121, %r2120, 4, 31, -1;
$L__tmp946:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2122, %r2120, %r2121;
$L__tmp947:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2123, %r2122, 2, 31, -1;
$L__tmp948:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2124, %r2122, %r2123;
$L__tmp949:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2125, %r2124, 1, 31, -1;
shfl.sync.bfly.b32 %r2126, %r1675, 16, 31, -1;
$L__tmp950:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2127, %r1675, %r2126;
$L__tmp951:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2128, %r2127, 8, 31, -1;
$L__tmp952:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2129, %r2127, %r2128;
$L__tmp953:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2130, %r2129, 4, 31, -1;
$L__tmp954:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2131, %r2129, %r2130;
$L__tmp955:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2132, %r2131, 2, 31, -1;
$L__tmp956:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2133, %r2131, %r2132;
$L__tmp957:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2134, %r2133, 1, 31, -1;
mov.b64 %rd384, {%r2125, %r2134};
mov.b64 %rd385, {%r2107, %r2116};
$L__tmp958:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd386, {%r2124, %r2133};
mov.b64 %rd387, {%r2106, %r2115};
add.f32x2 %rd388, %rd387, %rd385;
add.f32x2 %rd389, %rd386, %rd384;
$L__tmp959:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2135, %r1676, 16, 31, -1;
$L__tmp960:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2136, %r1676, %r2135;
$L__tmp961:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2137, %r2136, 8, 31, -1;
$L__tmp962:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2138, %r2136, %r2137;
$L__tmp963:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2139, %r2138, 4, 31, -1;
$L__tmp964:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2140, %r2138, %r2139;
$L__tmp965:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2141, %r2140, 2, 31, -1;
$L__tmp966:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2142, %r2140, %r2141;
$L__tmp967:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2143, %r2142, 1, 31, -1;
shfl.sync.bfly.b32 %r2144, %r1680, 16, 31, -1;
$L__tmp968:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2145, %r1680, %r2144;
$L__tmp969:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2146, %r2145, 8, 31, -1;
$L__tmp970:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2147, %r2145, %r2146;
$L__tmp971:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2148, %r2147, 4, 31, -1;
$L__tmp972:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2149, %r2147, %r2148;
$L__tmp973:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2150, %r2149, 2, 31, -1;
$L__tmp974:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2151, %r2149, %r2150;
$L__tmp975:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2152, %r2151, 1, 31, -1;
shfl.sync.bfly.b32 %r2153, %r1684, 16, 31, -1;
$L__tmp976:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2154, %r1684, %r2153;
$L__tmp977:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2155, %r2154, 8, 31, -1;
$L__tmp978:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2156, %r2154, %r2155;
$L__tmp979:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2157, %r2156, 4, 31, -1;
$L__tmp980:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2158, %r2156, %r2157;
$L__tmp981:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2159, %r2158, 2, 31, -1;
$L__tmp982:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2160, %r2158, %r2159;
$L__tmp983:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2161, %r2160, 1, 31, -1;
shfl.sync.bfly.b32 %r2162, %r1688, 16, 31, -1;
$L__tmp984:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2163, %r1688, %r2162;
$L__tmp985:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2164, %r2163, 8, 31, -1;
$L__tmp986:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2165, %r2163, %r2164;
$L__tmp987:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2166, %r2165, 4, 31, -1;
$L__tmp988:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2167, %r2165, %r2166;
$L__tmp989:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2168, %r2167, 2, 31, -1;
$L__tmp990:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2169, %r2167, %r2168;
$L__tmp991:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2170, %r2169, 1, 31, -1;
mov.b64 %rd390, {%r2161, %r2170};
mov.b64 %rd391, {%r2143, %r2152};
$L__tmp992:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd392, {%r2160, %r2169};
mov.b64 %rd393, {%r2142, %r2151};
add.f32x2 %rd394, %rd392, %rd390;
add.f32x2 %rd395, %rd393, %rd391;
$L__tmp993:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2171, %r1692, 16, 31, -1;
$L__tmp994:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2172, %r1692, %r2171;
$L__tmp995:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2173, %r2172, 8, 31, -1;
$L__tmp996:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2174, %r2172, %r2173;
$L__tmp997:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2175, %r2174, 4, 31, -1;
$L__tmp998:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2176, %r2174, %r2175;
$L__tmp999:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2177, %r2176, 2, 31, -1;
$L__tmp1000:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2178, %r2176, %r2177;
$L__tmp1001:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2179, %r2178, 1, 31, -1;
shfl.sync.bfly.b32 %r2180, %r1696, 16, 31, -1;
$L__tmp1002:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2181, %r1696, %r2180;
$L__tmp1003:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2182, %r2181, 8, 31, -1;
$L__tmp1004:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2183, %r2181, %r2182;
$L__tmp1005:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2184, %r2183, 4, 31, -1;
$L__tmp1006:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2185, %r2183, %r2184;
$L__tmp1007:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2186, %r2185, 2, 31, -1;
$L__tmp1008:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2187, %r2185, %r2186;
$L__tmp1009:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2188, %r2187, 1, 31, -1;
shfl.sync.bfly.b32 %r2189, %r1700, 16, 31, -1;
$L__tmp1010:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2190, %r1700, %r2189;
$L__tmp1011:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2191, %r2190, 8, 31, -1;
$L__tmp1012:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2192, %r2190, %r2191;
$L__tmp1013:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2193, %r2192, 4, 31, -1;
$L__tmp1014:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2194, %r2192, %r2193;
$L__tmp1015:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2195, %r2194, 2, 31, -1;
$L__tmp1016:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2196, %r2194, %r2195;
$L__tmp1017:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2197, %r2196, 1, 31, -1;
shfl.sync.bfly.b32 %r2198, %r1704, 16, 31, -1;
$L__tmp1018:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2199, %r1704, %r2198;
$L__tmp1019:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2200, %r2199, 8, 31, -1;
$L__tmp1020:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2201, %r2199, %r2200;
$L__tmp1021:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2202, %r2201, 4, 31, -1;
$L__tmp1022:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2203, %r2201, %r2202;
$L__tmp1023:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2204, %r2203, 2, 31, -1;
$L__tmp1024:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2205, %r2203, %r2204;
$L__tmp1025:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2206, %r2205, 1, 31, -1;
mov.b64 %rd396, {%r2197, %r2206};
mov.b64 %rd397, {%r2179, %r2188};
$L__tmp1026:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd398, {%r2196, %r2205};
mov.b64 %rd399, {%r2178, %r2187};
add.f32x2 %rd400, %rd399, %rd397;
add.f32x2 %rd401, %rd398, %rd396;
$L__tmp1027:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2207, %r1708, 16, 31, -1;
$L__tmp1028:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2208, %r1708, %r2207;
$L__tmp1029:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2209, %r2208, 8, 31, -1;
$L__tmp1030:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2210, %r2208, %r2209;
$L__tmp1031:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2211, %r2210, 4, 31, -1;
$L__tmp1032:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2212, %r2210, %r2211;
$L__tmp1033:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2213, %r2212, 2, 31, -1;
$L__tmp1034:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2214, %r2212, %r2213;
$L__tmp1035:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2215, %r2214, 1, 31, -1;
shfl.sync.bfly.b32 %r2216, %r1712, 16, 31, -1;
$L__tmp1036:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2217, %r1712, %r2216;
$L__tmp1037:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2218, %r2217, 8, 31, -1;
$L__tmp1038:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2219, %r2217, %r2218;
$L__tmp1039:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2220, %r2219, 4, 31, -1;
$L__tmp1040:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2221, %r2219, %r2220;
$L__tmp1041:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2222, %r2221, 2, 31, -1;
$L__tmp1042:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2223, %r2221, %r2222;
$L__tmp1043:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2224, %r2223, 1, 31, -1;
shfl.sync.bfly.b32 %r2225, %r1716, 16, 31, -1;
$L__tmp1044:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2226, %r1716, %r2225;
$L__tmp1045:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2227, %r2226, 8, 31, -1;
$L__tmp1046:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2228, %r2226, %r2227;
$L__tmp1047:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2229, %r2228, 4, 31, -1;
$L__tmp1048:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2230, %r2228, %r2229;
$L__tmp1049:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2231, %r2230, 2, 31, -1;
$L__tmp1050:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2232, %r2230, %r2231;
$L__tmp1051:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2233, %r2232, 1, 31, -1;
shfl.sync.bfly.b32 %r2234, %r1720, 16, 31, -1;
$L__tmp1052:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2235, %r1720, %r2234;
$L__tmp1053:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2236, %r2235, 8, 31, -1;
$L__tmp1054:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2237, %r2235, %r2236;
$L__tmp1055:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2238, %r2237, 4, 31, -1;
$L__tmp1056:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2239, %r2237, %r2238;
$L__tmp1057:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2240, %r2239, 2, 31, -1;
$L__tmp1058:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2241, %r2239, %r2240;
$L__tmp1059:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2242, %r2241, 1, 31, -1;
mov.b64 %rd402, {%r2233, %r2242};
mov.b64 %rd403, {%r2215, %r2224};
$L__tmp1060:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd404, {%r2232, %r2241};
mov.b64 %rd405, {%r2214, %r2223};
add.f32x2 %rd406, %rd405, %rd403;
add.f32x2 %rd407, %rd404, %rd402;
$L__tmp1061:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2243, %r1724, 16, 31, -1;
$L__tmp1062:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2244, %r1724, %r2243;
$L__tmp1063:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2245, %r2244, 8, 31, -1;
$L__tmp1064:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2246, %r2244, %r2245;
$L__tmp1065:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2247, %r2246, 4, 31, -1;
$L__tmp1066:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2248, %r2246, %r2247;
$L__tmp1067:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2249, %r2248, 2, 31, -1;
$L__tmp1068:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2250, %r2248, %r2249;
$L__tmp1069:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2251, %r2250, 1, 31, -1;
shfl.sync.bfly.b32 %r2252, %r1728, 16, 31, -1;
$L__tmp1070:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2253, %r1728, %r2252;
$L__tmp1071:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2254, %r2253, 8, 31, -1;
$L__tmp1072:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2255, %r2253, %r2254;
$L__tmp1073:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2256, %r2255, 4, 31, -1;
$L__tmp1074:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2257, %r2255, %r2256;
$L__tmp1075:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2258, %r2257, 2, 31, -1;
$L__tmp1076:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2259, %r2257, %r2258;
$L__tmp1077:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2260, %r2259, 1, 31, -1;
shfl.sync.bfly.b32 %r2261, %r1732, 16, 31, -1;
$L__tmp1078:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2262, %r1732, %r2261;
$L__tmp1079:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2263, %r2262, 8, 31, -1;
$L__tmp1080:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2264, %r2262, %r2263;
$L__tmp1081:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2265, %r2264, 4, 31, -1;
$L__tmp1082:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2266, %r2264, %r2265;
$L__tmp1083:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2267, %r2266, 2, 31, -1;
$L__tmp1084:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2268, %r2266, %r2267;
$L__tmp1085:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2269, %r2268, 1, 31, -1;
shfl.sync.bfly.b32 %r2270, %r1736, 16, 31, -1;
$L__tmp1086:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2271, %r1736, %r2270;
$L__tmp1087:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2272, %r2271, 8, 31, -1;
$L__tmp1088:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2273, %r2271, %r2272;
$L__tmp1089:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2274, %r2273, 4, 31, -1;
$L__tmp1090:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2275, %r2273, %r2274;
$L__tmp1091:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2276, %r2275, 2, 31, -1;
$L__tmp1092:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2277, %r2275, %r2276;
$L__tmp1093:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2278, %r2277, 1, 31, -1;
mov.b64 %rd408, {%r2269, %r2278};
mov.b64 %rd409, {%r2251, %r2260};
$L__tmp1094:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd410, {%r2268, %r2277};
mov.b64 %rd411, {%r2250, %r2259};
add.f32x2 %rd412, %rd411, %rd409;
add.f32x2 %rd413, %rd410, %rd408;
$L__tmp1095:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2279, %r1737, 16, 31, -1;
$L__tmp1096:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2280, %r1737, %r2279;
$L__tmp1097:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2281, %r2280, 8, 31, -1;
$L__tmp1098:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2282, %r2280, %r2281;
$L__tmp1099:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2283, %r2282, 4, 31, -1;
$L__tmp1100:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2284, %r2282, %r2283;
$L__tmp1101:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2285, %r2284, 2, 31, -1;
$L__tmp1102:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2286, %r2284, %r2285;
$L__tmp1103:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2287, %r2286, 1, 31, -1;
shfl.sync.bfly.b32 %r2288, %r1738, 16, 31, -1;
$L__tmp1104:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2289, %r1738, %r2288;
$L__tmp1105:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2290, %r2289, 8, 31, -1;
$L__tmp1106:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2291, %r2289, %r2290;
$L__tmp1107:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2292, %r2291, 4, 31, -1;
$L__tmp1108:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2293, %r2291, %r2292;
$L__tmp1109:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2294, %r2293, 2, 31, -1;
$L__tmp1110:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2295, %r2293, %r2294;
$L__tmp1111:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2296, %r2295, 1, 31, -1;
shfl.sync.bfly.b32 %r2297, %r1738, 16, 31, -1;
$L__tmp1112:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2298, %r1738, %r2297;
$L__tmp1113:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2299, %r2298, 8, 31, -1;
$L__tmp1114:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2300, %r2298, %r2299;
$L__tmp1115:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2301, %r2300, 4, 31, -1;
$L__tmp1116:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2302, %r2300, %r2301;
$L__tmp1117:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2303, %r2302, 2, 31, -1;
$L__tmp1118:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2304, %r2302, %r2303;
$L__tmp1119:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2305, %r2304, 1, 31, -1;
shfl.sync.bfly.b32 %r2306, %r1738, 16, 31, -1;
$L__tmp1120:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2307, %r1738, %r2306;
$L__tmp1121:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2308, %r2307, 8, 31, -1;
$L__tmp1122:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2309, %r2307, %r2308;
$L__tmp1123:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2310, %r2309, 4, 31, -1;
$L__tmp1124:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2311, %r2309, %r2310;
$L__tmp1125:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2312, %r2311, 2, 31, -1;
$L__tmp1126:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32 %r2313, %r2311, %r2312;
$L__tmp1127:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2314, %r2313, 1, 31, -1;
mov.b64 %rd414, {%r2305, %r2314};
mov.b64 %rd415, {%r2287, %r2296};
$L__tmp1128:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 %rd416, {%r2304, %r2313};
mov.b64 %rd417, {%r2286, %r2295};
add.f32x2 %rd418, %rd417, %rd415;
add.f32x2 %rd419, %rd416, %rd414;
$L__tmp1129:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
st.shared.v2.b64 [%r34+40960], {%rd328, %rd329};
st.shared.v2.b64 [%r34+40976], {%rd352, %rd353};
st.shared.v2.b64 [%r34+40992], {%rd376, %rd377};
st.shared.v2.b64 [%r34+41008], {%rd400, %rd401};
bar.sync 0;
ld.shared.v2.b64 {%rd420, %rd421}, [%r35+40960];
mov.b64 {%r2315, %r2316}, %rd421;
mov.b64 {%r2317, %r2318}, %rd420;
ld.shared.v2.b64 {%rd422, %rd423}, [%r35+40976];
mov.b64 {%r2319, %r2320}, %rd423;
mov.b64 {%r2321, %r2322}, %rd422;
ld.shared.v2.b64 {%rd424, %rd425}, [%r35+40992];
mov.b64 {%r2323, %r2324}, %rd425;
mov.b64 {%r2325, %r2326}, %rd424;
ld.shared.v2.b64 {%rd426, %rd427}, [%r35+41008];
mov.b64 {%r2327, %r2328}, %rd427;
mov.b64 {%r2329, %r2330}, %rd426;
bar.sync 0;
st.shared.v2.b64 [%r34+40960], {%rd334, %rd335};
st.shared.v2.b64 [%r34+40976], {%rd358, %rd359};
st.shared.v2.b64 [%r34+40992], {%rd382, %rd383};
st.shared.v2.b64 [%r34+41008], {%rd406, %rd407};
bar.sync 0;
ld.shared.v2.b64 {%rd428, %rd429}, [%r35+40960];
mov.b64 {%r2331, %r2332}, %rd429;
mov.b64 {%r2333, %r2334}, %rd428;
ld.shared.v2.b64 {%rd430, %rd431}, [%r35+40976];
mov.b64 {%r2335, %r2336}, %rd431;
mov.b64 {%r2337, %r2338}, %rd430;
ld.shared.v2.b64 {%rd432, %rd433}, [%r35+40992];
mov.b64 {%r2339, %r2340}, %rd433;
mov.b64 {%r2341, %r2342}, %rd432;
ld.shared.v2.b64 {%rd434, %rd435}, [%r35+41008];
mov.b64 {%r2343, %r2344}, %rd435;
mov.b64 {%r2345, %r2346}, %rd434;
bar.sync 0;
st.shared.v2.b64 [%r34+40960], {%rd340, %rd341};
st.shared.v2.b64 [%r34+40976], {%rd364, %rd365};
st.shared.v2.b64 [%r34+40992], {%rd388, %rd389};
st.shared.v2.b64 [%r34+41008], {%rd412, %rd413};
bar.sync 0;
ld.shared.v2.b64 {%rd436, %rd437}, [%r35+40960];
mov.b64 {%r2347, %r2348}, %rd437;
mov.b64 {%r2349, %r2350}, %rd436;
ld.shared.v2.b64 {%rd438, %rd439}, [%r35+40976];
mov.b64 {%r2351, %r2352}, %rd439;
mov.b64 {%r2353, %r2354}, %rd438;
ld.shared.v2.b64 {%rd440, %rd441}, [%r35+40992];
mov.b64 {%r2355, %r2356}, %rd441;
mov.b64 {%r2357, %r2358}, %rd440;
ld.shared.v2.b64 {%rd442, %rd443}, [%r35+41008];
mov.b64 {%r2359, %r2360}, %rd443;
mov.b64 {%r2361, %r2362}, %rd442;
bar.sync 0;
st.shared.v2.b64 [%r34+40960], {%rd346, %rd347};
st.shared.v2.b64 [%r34+40976], {%rd370, %rd371};
st.shared.v2.b64 [%r34+40992], {%rd395, %rd394};
st.shared.v2.b64 [%r34+41008], {%rd418, %rd419};
bar.sync 0;
shfl.sync.bfly.b32 %r2363, %r2317, 1, 31, -1;
shfl.sync.bfly.b32 %r2364, %r2318, 1, 31, -1;
shfl.sync.bfly.b32 %r2365, %r2315, 1, 31, -1;
shfl.sync.bfly.b32 %r2366, %r2316, 1, 31, -1;
mov.b64 %rd444, {%r2365, %r2366};
mov.b64 %rd445, {%r2363, %r2364};
$L__tmp1130:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd446, %rd420, %rd445;
add.f32x2 %rd447, %rd421, %rd444;
$L__tmp1131:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2367, %r2333, 1, 31, -1;
shfl.sync.bfly.b32 %r2368, %r2334, 1, 31, -1;
shfl.sync.bfly.b32 %r2369, %r2331, 1, 31, -1;
shfl.sync.bfly.b32 %r2370, %r2332, 1, 31, -1;
mov.b64 %rd448, {%r2369, %r2370};
mov.b64 %rd449, {%r2367, %r2368};
$L__tmp1132:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd450, %rd428, %rd449;
add.f32x2 %rd451, %rd429, %rd448;
$L__tmp1133:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2371, %r2349, 1, 31, -1;
shfl.sync.bfly.b32 %r2372, %r2350, 1, 31, -1;
shfl.sync.bfly.b32 %r2373, %r2347, 1, 31, -1;
shfl.sync.bfly.b32 %r2374, %r2348, 1, 31, -1;
mov.b64 %rd452, {%r2373, %r2374};
mov.b64 %rd453, {%r2371, %r2372};
$L__tmp1134:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd454, %rd436, %rd453;
add.f32x2 %rd455, %rd437, %rd452;
$L__tmp1135:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
ld.shared.v2.b64 {%rd456, %rd457}, [%r35+40960];
mov.b64 {%r2375, %r2376}, %rd457;
mov.b64 {%r2377, %r2378}, %rd456;
shfl.sync.bfly.b32 %r2379, %r2377, 1, 31, -1;
shfl.sync.bfly.b32 %r2380, %r2378, 1, 31, -1;
shfl.sync.bfly.b32 %r2381, %r2375, 1, 31, -1;
shfl.sync.bfly.b32 %r2382, %r2376, 1, 31, -1;
mov.b64 %rd458, {%r2381, %r2382};
mov.b64 %rd459, {%r2379, %r2380};
$L__tmp1136:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd460, %rd456, %rd459;
add.f32x2 %rd461, %rd457, %rd458;
$L__tmp1137:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2383, %r2321, 1, 31, -1;
shfl.sync.bfly.b32 %r2384, %r2322, 1, 31, -1;
shfl.sync.bfly.b32 %r2385, %r2319, 1, 31, -1;
shfl.sync.bfly.b32 %r2386, %r2320, 1, 31, -1;
mov.b64 %rd462, {%r2385, %r2386};
mov.b64 %rd463, {%r2383, %r2384};
$L__tmp1138:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd464, %rd422, %rd463;
add.f32x2 %rd465, %rd423, %rd462;
$L__tmp1139:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2387, %r2337, 1, 31, -1;
shfl.sync.bfly.b32 %r2388, %r2338, 1, 31, -1;
shfl.sync.bfly.b32 %r2389, %r2335, 1, 31, -1;
shfl.sync.bfly.b32 %r2390, %r2336, 1, 31, -1;
mov.b64 %rd466, {%r2389, %r2390};
mov.b64 %rd467, {%r2387, %r2388};
$L__tmp1140:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd468, %rd430, %rd467;
add.f32x2 %rd469, %rd431, %rd466;
$L__tmp1141:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2391, %r2353, 1, 31, -1;
shfl.sync.bfly.b32 %r2392, %r2354, 1, 31, -1;
shfl.sync.bfly.b32 %r2393, %r2351, 1, 31, -1;
shfl.sync.bfly.b32 %r2394, %r2352, 1, 31, -1;
mov.b64 %rd470, {%r2393, %r2394};
mov.b64 %rd471, {%r2391, %r2392};
$L__tmp1142:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd472, %rd438, %rd471;
add.f32x2 %rd473, %rd439, %rd470;
$L__tmp1143:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
ld.shared.v2.b64 {%rd474, %rd475}, [%r35+40976];
mov.b64 {%r2395, %r2396}, %rd475;
mov.b64 {%r2397, %r2398}, %rd474;
shfl.sync.bfly.b32 %r2399, %r2397, 1, 31, -1;
shfl.sync.bfly.b32 %r2400, %r2398, 1, 31, -1;
shfl.sync.bfly.b32 %r2401, %r2395, 1, 31, -1;
shfl.sync.bfly.b32 %r2402, %r2396, 1, 31, -1;
mov.b64 %rd476, {%r2401, %r2402};
mov.b64 %rd477, {%r2399, %r2400};
$L__tmp1144:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd478, %rd474, %rd477;
add.f32x2 %rd479, %rd475, %rd476;
$L__tmp1145:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2403, %r2325, 1, 31, -1;
shfl.sync.bfly.b32 %r2404, %r2326, 1, 31, -1;
shfl.sync.bfly.b32 %r2405, %r2323, 1, 31, -1;
shfl.sync.bfly.b32 %r2406, %r2324, 1, 31, -1;
mov.b64 %rd480, {%r2405, %r2406};
mov.b64 %rd481, {%r2403, %r2404};
$L__tmp1146:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd482, %rd424, %rd481;
add.f32x2 %rd483, %rd425, %rd480;
$L__tmp1147:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2407, %r2341, 1, 31, -1;
shfl.sync.bfly.b32 %r2408, %r2342, 1, 31, -1;
shfl.sync.bfly.b32 %r2409, %r2339, 1, 31, -1;
shfl.sync.bfly.b32 %r2410, %r2340, 1, 31, -1;
mov.b64 %rd484, {%r2409, %r2410};
mov.b64 %rd485, {%r2407, %r2408};
$L__tmp1148:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd486, %rd432, %rd485;
add.f32x2 %rd487, %rd433, %rd484;
$L__tmp1149:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2411, %r2357, 1, 31, -1;
shfl.sync.bfly.b32 %r2412, %r2358, 1, 31, -1;
shfl.sync.bfly.b32 %r2413, %r2355, 1, 31, -1;
shfl.sync.bfly.b32 %r2414, %r2356, 1, 31, -1;
mov.b64 %rd488, {%r2413, %r2414};
mov.b64 %rd489, {%r2411, %r2412};
$L__tmp1150:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd490, %rd440, %rd489;
add.f32x2 %rd491, %rd441, %rd488;
$L__tmp1151:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
ld.shared.v2.b64 {%rd492, %rd493}, [%r35+40992];
mov.b64 {%r2415, %r2416}, %rd493;
mov.b64 {%r2417, %r2418}, %rd492;
shfl.sync.bfly.b32 %r2419, %r2417, 1, 31, -1;
shfl.sync.bfly.b32 %r2420, %r2418, 1, 31, -1;
shfl.sync.bfly.b32 %r2421, %r2415, 1, 31, -1;
shfl.sync.bfly.b32 %r2422, %r2416, 1, 31, -1;
mov.b64 %rd494, {%r2421, %r2422};
mov.b64 %rd495, {%r2419, %r2420};
$L__tmp1152:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd496, %rd492, %rd495;
add.f32x2 %rd497, %rd493, %rd494;
$L__tmp1153:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2423, %r2329, 1, 31, -1;
shfl.sync.bfly.b32 %r2424, %r2330, 1, 31, -1;
shfl.sync.bfly.b32 %r2425, %r2327, 1, 31, -1;
shfl.sync.bfly.b32 %r2426, %r2328, 1, 31, -1;
mov.b64 %rd498, {%r2425, %r2426};
mov.b64 %rd499, {%r2423, %r2424};
$L__tmp1154:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd500, %rd426, %rd499;
add.f32x2 %rd501, %rd427, %rd498;
$L__tmp1155:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2427, %r2345, 1, 31, -1;
shfl.sync.bfly.b32 %r2428, %r2346, 1, 31, -1;
shfl.sync.bfly.b32 %r2429, %r2343, 1, 31, -1;
shfl.sync.bfly.b32 %r2430, %r2344, 1, 31, -1;
mov.b64 %rd502, {%r2429, %r2430};
mov.b64 %rd503, {%r2427, %r2428};
$L__tmp1156:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd504, %rd434, %rd503;
add.f32x2 %rd505, %rd435, %rd502;
$L__tmp1157:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
shfl.sync.bfly.b32 %r2431, %r2361, 1, 31, -1;
shfl.sync.bfly.b32 %r2432, %r2362, 1, 31, -1;
shfl.sync.bfly.b32 %r2433, %r2359, 1, 31, -1;
shfl.sync.bfly.b32 %r2434, %r2360, 1, 31, -1;
mov.b64 %rd506, {%r2433, %r2434};
mov.b64 %rd507, {%r2431, %r2432};
$L__tmp1158:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd508, %rd442, %rd507;
add.f32x2 %rd509, %rd443, %rd506;
$L__tmp1159:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1193 18 // standard.py:293:12
ld.shared.v2.b64 {%rd510, %rd511}, [%r35+41008];
mov.b64 {%r2435, %r2436}, %rd511;
mov.b64 {%r2437, %r2438}, %rd510;
shfl.sync.bfly.b32 %r2439, %r2437, 1, 31, -1;
shfl.sync.bfly.b32 %r2440, %r2438, 1, 31, -1;
shfl.sync.bfly.b32 %r2441, %r2435, 1, 31, -1;
shfl.sync.bfly.b32 %r2442, %r2436, 1, 31, -1;
mov.b64 %rd512, {%r2441, %r2442};
mov.b64 %rd513, {%r2439, %r2440};
$L__tmp1160:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
add.f32x2 %rd514, %rd510, %rd513;
add.f32x2 %rd515, %rd511, %rd512;
$L__tmp1161:
.loc 1 1193 9 // test_core.py:1193:9
bar.sync 0;
st.shared.v2.b64 [%r36+40960], {%rd446, %rd447};
st.shared.v2.b64 [%r36+40976], {%rd450, %rd451};
st.shared.v2.b64 [%r36+40992], {%rd454, %rd455};
st.shared.v2.b64 [%r36+41008], {%rd460, %rd461};
st.shared.v2.b64 [%r36+41088], {%rd464, %rd465};
st.shared.v2.b64 [%r36+41104], {%rd468, %rd469};
st.shared.v2.b64 [%r36+41120], {%rd472, %rd473};
st.shared.v2.b64 [%r36+41136], {%rd478, %rd479};
st.shared.v2.b64 [%r36+41216], {%rd482, %rd483};
st.shared.v2.b64 [%r36+41232], {%rd486, %rd487};
st.shared.v2.b64 [%r36+41248], {%rd490, %rd491};
st.shared.v2.b64 [%r36+41264], {%rd496, %rd497};
st.shared.v2.b64 [%r36+41344], {%rd500, %rd501};
st.shared.v2.b64 [%r36+41360], {%rd504, %rd505};
st.shared.v2.b64 [%r36+41376], {%rd508, %rd509};
st.shared.v2.b64 [%r36+41392], {%rd514, %rd515};
bar.sync 0;
ld.shared.b32 %r2443, [%r37+40960];
add.f32 %r2444, %r2459, %r2443;
mov.b32 %r1322, 0;
.loc 1 1192 14 // test_core.py:1192:14
// begin inline asm
{
.reg .pred complete;
.loc 1 1195 14 // test_core.py:1195:14
waitLoop:
.loc 1 1196 14 // test_core.py:1196:14
mbarrier.try_wait.parity.shared::cta.b64 complete, [%r2452], %r1322;
@!complete bra.uni waitLoop;
}
// end inline asm
// begin inline asm
.loc 1 1192 14 // test_core.py:1192:14
tcgen05.ld.sync.aligned.32x32b.x64.b32 {%r1323, %r1324, %r1325, %r1326, %r1327, %r1328, %r1329, %r1330, %r1331, %r1332, %r1333, %r1334, %r1335, %r1336, %r1337, %r1338, %r1339, %r1340, %r1341, %r1342, %r1343, %r1344, %r1345, %r1346, %r1347, %r1348, %r1349, %r1350, %r1351, %r1352, %r1353, %r1354, %r1355, %r1356, %r1357, %r1358, %r1359, %r1360, %r1361, %r1362, %r1363, %r1364, %r1365, %r1366, %r1367, %r1368, %r1369, %r1370, %r1371, %r1372, %r1373, %r1374, %r1375, %r1376, %r1377, %r1378, %r1379, %r1380, %r1381, %r1382, %r1383, %r1384, %r1385, %r1386}, [%r1387 + 0];
// end inline asm
tcgen05.wait::ld.sync.aligned;
$L__tmp1162:
.loc 2 293 12, function_name $L__info_string0, inlined_at 1 1194 20 // standard.py:293:12
mov.b64 %rd516, {%r1323, %r1324};
mov.b64 %rd517, {%r1325, %r1326};
mov.b64 %rd518, {%r1327, %r1328};
mov.b64 %rd519, {%r1329, %r1330};
mov.b64 %rd520, {%r1331, %r1332};
mov.b64 %rd521, {%r1333, %r1334};
mov.b64 %rd522, {%r1335, %r1336};
mov.b64 %rd523, {%r1337, %r1338};
mov.b64 %rd524, {%r1339, %r1340};
mov.b64 %rd525, {%r1341, %r1342};
mov.b64 %rd526, {%r1343, %r1344};
mov.b64 %rd527, {%r1345, %r1346};
mov.b64 %rd528, {%r1347, %r1348};
mov.b64 %rd529, {%r1349, %r1350};
mov.b64 %rd530, {%r1351, %r1352};
mov.b64 %rd531, {%r1353, %r1354};
mov.b64 %rd532, {%r1355, %r1356};
mov.b64 %rd533, {%r1357, %r1358};
mov.b64 %rd534, {%r1359, %r1360};
mov.b64 %rd535, {%r1361, %r1362};
mov.b64 %rd536, {%r1363, %r1364};
mov.b64 %rd537, {%r1365, %r1366};
mov.b64 %rd538, {%r1367, %r1368};
mov.b64 %rd539, {%r1369, %r1370};
mov.b64 %rd540, {%r1371, %r1372};
mov.b64 %rd541, {%r1373, %r1374};
mov.b64 %rd542, {%r1375, %r1376};
mov.b64 %rd543, {%r1377, %r1378};
mov.b64 %rd544, {%r1379, %r1380};
mov.b64 %rd545, {%r1381, %r1382};
mov.b64 %rd546, {%r1383, %r1384};
mov.b64 %rd547, {%r1385, %r1386};
add.f32x2 %rd548, %rd516, %rd517;
add.f32x2 %rd549, %rd518, %rd519;
add.f32x2 %rd550, %rd520, %rd521;
add.f32x2 %rd551, %rd522, %rd523;
add.f32x2 %rd552, %rd524, %rd525;
add.f32x2 %rd553, %rd526, %rd527;
add.f32x2 %rd554, %rd528, %rd529;
add.f32x2 %rd555, %rd530, %rd531;
add.f32x2 %rd556, %rd532, %rd533;
add.f32x2 %rd557, %rd534, %rd535;
add.f32x2 %rd558, %rd536, %rd537;
add.f32x2 %rd559, %rd538, %rd539;
add.f32x2 %rd560, %rd540, %rd541;
add.f32x2 %rd561, %rd542, %rd543;
add.f32x2 %rd562, %rd544, %rd545;
add.f32x2 %rd563, %rd546, %rd547;
add.f32x2 %rd564, %rd548, %rd549;
add.f32x2 %rd565, %rd550, %rd551;
add.f32x2 %rd566, %rd552, %rd553;
add.f32x2 %rd567, %rd554, %rd555;
add.f32x2 %rd568, %rd556, %rd557;
add.f32x2 %rd569, %rd558, %rd559;
add.f32x2 %rd570, %rd560, %rd561;
add.f32x2 %rd571, %rd562, %rd563;
add.f32x2 %rd572, %rd564, %rd565;
add.f32x2 %rd573, %rd566, %rd567;
add.f32x2 %rd574, %rd568, %rd569;
add.f32x2 %rd575, %rd570, %rd571;
add.f32x2 %rd576, %rd572, %rd573;
add.f32x2 %rd577, %rd574, %rd575;
add.f32x2 %rd578, %rd576, %rd577;
$L__tmp1163:
.loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
mov.b64 {%r2445, %r2446}, %rd578;
add.f32 %r2447, %r2445, %r2446;
$L__tmp1164:
.loc 1 1194 9 // test_core.py:1194:9
add.f32 %r2448, %r2458, %r2447;
.loc 1 1189 5 // test_core.py:1189:5
cp.async.wait_group 0;
bar.sync 0;
elect.sync %r2449|%p171, -1;
and.pred %p94, %p4, %p171;
add.s32 %r1388, %r39, 41472;
// begin inline asm
@%p94 mbarrier.inval.shared::cta.b64 [%r1388];
// end inline asm
elect.sync %r2450|%p172, -1;
and.pred %p95, %p4, %p172;
// begin inline asm
@%p95 mbarrier.inval.shared::cta.b64 [%r72];
// end inline asm
.loc 1 1197 14 // test_core.py:1197:14
mad.wide.u32 %rd323, %r3, 4, %rd19;
.loc 1 1197 47 // test_core.py:1197:47
add.f32 %r1389, %r2444, %r2448;
.loc 1 1197 5 // test_core.py:1197:5
// begin inline asm
.loc 1 1197 5 // test_core.py:1197:5
st.global.b32 [ %rd323 + 0 ], { %r1389 };
// end inline asm
.loc 1 1176 1 // test_core.py:1176:1
bar.sync 0;
// begin inline asm
@%p1 tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r1390, 64;
// end inline asm
ret;
$L__tmp1165:
$L__func_end0:
// -- End function
}
.file 1 "/home/drisspg/meta/triton/python/test/unit/language/test_core.py"
.file 2 "/home/drisspg/meta/triton/python/triton/language/standard.py"
.section .debug_abbrev
{
.b8 1 // Abbreviation Code
.b8 17 // DW_TAG_compile_unit
.b8 1 // DW_CHILDREN_yes
.b8 37 // DW_AT_producer
.b8 8 // DW_FORM_string
.b8 19 // DW_AT_language
.b8 5 // DW_FORM_data2
.b8 3 // DW_AT_name
.b8 8 // DW_FORM_string
.b8 16 // DW_AT_stmt_list
.b8 6 // DW_FORM_data4
.b8 27 // DW_AT_comp_dir
.b8 8 // DW_FORM_string
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 2 // Abbreviation Code
.b8 46 // DW_TAG_subprogram
.b8 0 // DW_CHILDREN_no
.b8 3 // DW_AT_name
.b8 8 // DW_FORM_string
.b8 32 // DW_AT_inline
.b8 11 // DW_FORM_data1
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 3 // Abbreviation Code
.b8 46 // DW_TAG_subprogram
.b8 1 // DW_CHILDREN_yes
.b8 17 // DW_AT_low_pc
.b8 1 // DW_FORM_addr
.b8 18 // DW_AT_high_pc
.b8 1 // DW_FORM_addr
.b8 49 // DW_AT_abstract_origin
.b8 19 // DW_FORM_ref4
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 4 // Abbreviation Code
.b8 29 // DW_TAG_inlined_subroutine
.b8 0 // DW_CHILDREN_no
.b8 49 // DW_AT_abstract_origin
.b8 19 // DW_FORM_ref4
.b8 17 // DW_AT_low_pc
.b8 1 // DW_FORM_addr
.b8 18 // DW_AT_high_pc
.b8 1 // DW_FORM_addr
.b8 88 // DW_AT_call_file
.b8 11 // DW_FORM_data1
.b8 89 // DW_AT_call_line
.b8 5 // DW_FORM_data2
.b8 87 // DW_AT_call_column
.b8 11 // DW_FORM_data1
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 5 // Abbreviation Code
.b8 29 // DW_TAG_inlined_subroutine
.b8 1 // DW_CHILDREN_yes
.b8 49 // DW_AT_abstract_origin
.b8 19 // DW_FORM_ref4
.b8 17 // DW_AT_low_pc
.b8 1 // DW_FORM_addr
.b8 18 // DW_AT_high_pc
.b8 1 // DW_FORM_addr
.b8 88 // DW_AT_call_file
.b8 11 // DW_FORM_data1
.b8 89 // DW_AT_call_line
.b8 5 // DW_FORM_data2
.b8 87 // DW_AT_call_column
.b8 11 // DW_FORM_data1
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 0 // EOM(3)
}
.section .debug_info
{
.b32 263 // Length of Unit
.b8 2 // DWARF version number
.b8 0
.b32 .debug_abbrev // Offset Into Abbrev. Section
.b8 8 // Address Size (in bytes)
.b8 1 // Abbrev [1] 0xb:0x100 DW_TAG_compile_unit
.b8 116 // DW_AT_producer
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 0
.b8 2 // DW_AT_language
.b8 0
.b8 116 // DW_AT_name
.b8 101
.b8 115
.b8 116
.b8 95
.b8 99
.b8 111
.b8 114
.b8 101
.b8 46
.b8 112
.b8 121
.b8 0
.b32 .debug_line // DW_AT_stmt_list
.b8 47 // DW_AT_comp_dir
.b8 104
.b8 111
.b8 109
.b8 101
.b8 47
.b8 100
.b8 114
.b8 105
.b8 115
.b8 115
.b8 112
.b8 103
.b8 47
.b8 109
.b8 101
.b8 116
.b8 97
.b8 47
.b8 116
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 47
.b8 112
.b8 121
.b8 116
.b8 104
.b8 111
.b8 110
.b8 47
.b8 116
.b8 101
.b8 115
.b8 116
.b8 47
.b8 117
.b8 110
.b8 105
.b8 116
.b8 47
.b8 108
.b8 97
.b8 110
.b8 103
.b8 117
.b8 97
.b8 103
.b8 101
.b8 0
.b8 2 // Abbrev [2] 0x5a:0x1b DW_TAG_subprogram
.b8 105 // DW_AT_name
.b8 110
.b8 116
.b8 101
.b8 103
.b8 101
.b8 114
.b8 95
.b8 97
.b8 98
.b8 115
.b8 95
.b8 112
.b8 116
.b8 120
.b8 97
.b8 115
.b8 95
.b8 107
.b8 101
.b8 114
.b8 110
.b8 101
.b8 108
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x75:0x95 DW_TAG_subprogram
.b64 $L__func_begin0 // DW_AT_low_pc
.b64 $L__func_end0 // DW_AT_high_pc
.b32 90 // DW_AT_abstract_origin
.b8 4 // Abbrev [4] 0x8a:0x19 DW_TAG_inlined_subroutine
.b32 90 // DW_AT_abstract_origin
.b64 $L__tmp1 // DW_AT_low_pc
.b64 $L__tmp584 // DW_AT_high_pc
.b8 1 // DW_AT_call_file
.b8 169 // DW_AT_call_line
.b8 4
.b8 25 // DW_AT_call_column
.b8 5 // Abbrev [5] 0xa3:0x33 DW_TAG_inlined_subroutine
.b32 90 // DW_AT_abstract_origin
.b64 $L__tmp3 // DW_AT_low_pc
.b64 $L__tmp1161 // DW_AT_high_pc
.b8 1 // DW_AT_call_file
.b8 169 // DW_AT_call_line
.b8 4
.b8 18 // DW_AT_call_column
.b8 4 // Abbrev [4] 0xbc:0x19 DW_TAG_inlined_subroutine
.b32 90 // DW_AT_abstract_origin
.b64 $L__tmp4 // DW_AT_low_pc
.b64 $L__tmp1161 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 37 // DW_AT_call_line
.b8 1
.b8 12 // DW_AT_call_column
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0xd6:0x33 DW_TAG_inlined_subroutine
.b32 90 // DW_AT_abstract_origin
.b64 $L__tmp580 // DW_AT_low_pc
.b64 $L__tmp1164 // DW_AT_high_pc
.b8 1 // DW_AT_call_file
.b8 170 // DW_AT_call_line
.b8 4
.b8 20 // DW_AT_call_column
.b8 4 // Abbrev [4] 0xef:0x19 DW_TAG_inlined_subroutine
.b32 90 // DW_AT_abstract_origin
.b64 $L__tmp581 // DW_AT_low_pc
.b64 $L__tmp1164 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 37 // DW_AT_call_line
.b8 1
.b8 12 // DW_AT_call_column
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
}
.section .debug_str
{
$L__info_string0:
.b8 105 // string offset=0 ; integer_abs_ptxas_kernel
.b8 110
.b8 116
.b8 101
.b8 103
.b8 101
.b8 114
.b8 95
.b8 97
.b8 98
.b8 115
.b8 95
.b8 112
.b8 116
.b8 120
.b8 97
.b8 115
.b8 95
.b8 107
.b8 101
.b8 114
.b8 110
.b8 101
.b8 108
.b8 0
}
.section .debug_macinfo { }
Raw
repro.py
import torch
import triton
import triton.language as tl
@triton.jit
def neighborhood_mask(m, n):
center_row = tl.maximum(tl.minimum(m // 128, 121), 6)
center_col = tl.maximum(tl.minimum(m % 128, 121), 6)
row_distance = tl.abs(center_row - n // 128)
col_distance = tl.abs(center_col - n % 128)
return tl.maximum(row_distance, col_distance) <= 6
@triton.jit
def repro(q_ptr, k_ptr, out_ptr):
block_m: tl.constexpr = 128
block_n: tl.constexpr = 64
head_dim: tl.constexpr = 64
m_offset = tl.arange(0, block_m)[:, None]
m = 89 * 128 + m_offset
d = tl.arange(0, head_dim)[None, :]
q = tl.load(q_ptr + m_offset * head_dim + d)
count = tl.zeros([block_m], tl.float32)
dot_sum = tl.zeros([block_m], tl.float32)
n = 83 * 128 + tl.arange(0, block_n)[None, :]
kv_offset = 0
for _ in range(26):
n_offset = tl.arange(0, block_n)[:, None]
k = tl.load(k_ptr + (kv_offset + n_offset) * head_dim + d)
qk = tl.dot(q, tl.trans(k), input_precision="ieee")
count += tl.sum(neighborhood_mask(m, n).to(tl.float32), 1)
dot_sum += tl.sum(qk, 1)
n += block_n
kv_offset += block_n
tl.store(out_ptr + tl.arange(0, block_m), count + dot_sum)
q = torch.zeros((128, 64), device="cuda", dtype=torch.float16)
k = torch.zeros((26 * 64, 64), device="cuda", dtype=torch.float16)
result = torch.empty(128, device="cuda")
repro[(1,)](q, k, result, num_warps=4, num_stages=3)
torch.cuda.synchronize()
expected_count = 13 * 13
maximum_error = (result - expected_count).abs().max().item()
print(f"expected key count: {expected_count}")
print(f"maximum key-count error: {maximum_error}")
assert maximum_error == 0
Raw
run.sh
#!/usr/bin/env bash
set -euo pipefail
root="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
cd "$root"
cuda_home="${CUDA_HOME:-/usr/local/cuda-13.1}"
ptxas="${PTXAS:-$(command -v ptxas)}"
cxx="${CXX:-g++}"
"$cxx" -std=c++17 -O2 -I"$cuda_home/include" repro.cpp -lcuda -o repro
common_ptxas_args=(--gpu-name=sm_100a --regAllocOptLevel=2 -lineinfo repro.ptx)
"$ptxas" "${common_ptxas_args[@]}" -o repro_o3.cubin
"$ptxas" --opt-level=0 "${common_ptxas_args[@]}" -o repro_o0.cubin
set +e
./repro repro_o3.cubin
o3_status=$?
./repro repro_o0.cubin
o0_status=$?
set -e
printf 'optimized ptxas status: %d (expected 1 on affected versions)\n' "$o3_status"
printf 'ptxas O0 status: %d (expected 0)\n' "$o0_status"
if [[ "$o3_status" -ne 1 || "$o0_status" -ne 0 ]]; then
exit 1
fi
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment