当前位置：首页 > news >正文

ptx 简介03，ldmatrix 的应用实例解析

news 2025/9/12 5:45:22

1. 实例编译运行

main.cu

//nvcc -g -lineinfo -std=c++17 -arch=native main.cu -o main#include <iostream>
#include <thrust/device_vector.h>/*
ldmatrix.sync.aligned.shape.num{.trans}{.ss}.type r, [p];.shape  = {.m8n8};
.num    = {.x1, .x2, .x4};
.ss     = {.shared{::cta}};
.type   = {.b16};
*/__device__ 
void ldmatrix_x2(unsigned int (&x)[2], const void* ptr){ asm volatile("ldmatrix.sync.aligned.m8n8.x2.shared.b16 {%0, %1}, [%2];" : "=r"(x[0]), "=r"(x[1]): "l"(__cvta_generic_to_shared(ptr))); 
}__global__
void mykernel(const int* loadOffsets, bool print){alignas(16) __shared__ half A[128 * 16];for(int i = threadIdx.x; i < 128*16; i += blockDim.x){A[i] = i;}__syncthreads();const int lane = threadIdx.x % 32;unsigned int result[2];const int offset = loadOffsets[lane];ldmatrix_x2(result, &A[offset]);half2 loaded[2];memcpy(&loaded[0], &result[0], sizeof(half2) * 2);if(print){for(int m = 0; m < 2; m++){for(int t = 0; t < 32; t++){if(lane == t){printf("%4d %4d ", int(loaded[m].x), int(loaded[m].y));if(lane % 4 == 3){printf("\n");}}__syncwarp();}if(lane == 0){printf("\n");}__syncwarp();}}
}int main(){thrust::device_vector<int> d_loadOffsets(32, 0);for(int i = 0; i < 16; i++){const int row = i % 8;const int matrix = i / 8;d_loadOffsets[i] = row * 16 + matrix * 8;}mykernel<<<1,32>>>(d_loadOffsets.data().get(), true);cudaDeviceSynchronize();// Shared Load Matrix: Requests 16.384, Wavefronts 33.393, Bank Conflicts 0for(int i = 0; i < 16; i++){const int row = i / 2;const int matrix = i % 2;d_loadOffsets[i] = row * 16 + matrix * 8;}std::cout << "offsets: ";for(int i = 0; i < 16; i++){std::cout << d_loadOffsets[i] << " ";}std::cout << "\n";mykernel<<<1024,512>>>(d_loadOffsets.data().get(), false);cudaDeviceSynchronize();// Shared Load Matrix: Requests 16.384, Wavefronts 131.674, Bank Conflicts 98.304for(int i = 0; i < 16; i++){const int row = i / 2;const int matrix = i % 2;d_loadOffsets[i] = (4*row) * 16 + matrix * 8;}std::cout << "offsets: ";for(int i = 0; i < 16; i++){std::cout << d_loadOffsets[i] << " ";}std::cout << "\n";mykernel<<<1024,512>>>(d_loadOffsets.data().get(), false);cudaDeviceSynchronize();// Shared Load Matrix: Requests 16.384, Wavefronts 66.488, Bank Conflicts 32.768for(int i = 0; i < 16; i++){const int row = i % 8;const int matrix = i / 8;d_loadOffsets[i] = row * 16 + matrix * 8;}std::cout << "offsets: ";for(int i = 0; i < 16; i++){std::cout << d_loadOffsets[i] << " ";}std::cout << "\n";mykernel<<<1024,512>>>(d_loadOffsets.data().get(), false);cudaDeviceSynchronize();// Shared Load Matrix: Requests 16.384, Wavefronts 263.070, Bank Conflicts 229.376for(int i = 0; i < 16; i++){const int row = i % 8;const int matrix = i / 8;d_loadOffsets[i] = (4*row) * 16 + matrix * 8;}std::cout << "offsets: ";for(int i = 0; i < 16; i++){std::cout << d_loadOffsets[i] << " ";}std::cout << "\n";mykernel<<<1024,512>>>(d_loadOffsets.data().get(), false);cudaDeviceSynchronize();
}

编译运行：

nvcc -g -lineinfo -std=c++17 -arch=native main.cu -o main

或者 device 代码 debug版：

$ nvcc -g -G -std=c++17 -arch=native main.cu -o main

2. 实例功能解析

详细解析这个执行了 ldmatrix 的 CUDA Device 函数。这是一个非常经典且高效的用法。

2.1. 函数签名解析

__device__ void ldmatrix_x2(unsigned int (&x)[2], const void* ptr)

__device__: 声明这是一个在 GPU 上执行的函数。

unsigned int (&x)[2]: 这是一个对包含 2 个 unsigned int 的数组的引用。使用引用 (&) 允许函数直接修改调用者传入的数组，避免了传值拷贝。这个数组的两个元素 x[0] 和 x[1] 将被用作内联汇编中的目标寄存器。

const void* ptr: 这是一个指向共享内存中某个数据的通用指针。const 表示函数不会通过这个指针修改数据，void* 提供了灵活性，可以指向任何类型的数据。

2.2. 内联汇编详解

asm volatile("ldmatrix.sync.aligned.m8n8.x2.shared.b16 {%0, %1}, [%2];": "=r"(x[0]), "=r"(x[1])  // Output operands: "l"(__cvta_generic_to_shared(ptr))); // Input operand

我们逐部分分解：

2.2.1. 汇编模板字符串 (`"ldmatrix.sync.aligned.m8n8.x2.shared.b16 {%0, %1}, [%2];"`)

这是要执行的 PTX 指令。

ldmatrix: 指令本身，加载矩阵。

.sync: Warp 级同步指令，确保 Warp 内所有活跃线程协同执行。

.aligned: 强制要求源内存地址 (ptr) 必须是 16 字节对齐的。

.m8n8: 指定从内存中加载的数据布局对应于一个 8 行 x 8 列的矩阵。

.x2: 指定矩阵中每个元素的大小是 2 字节（16 位）。这意味着它用于加载 half (FP16)、__half、uint16_t、int16_t 等类型的数据。

.shared: 明确指定源数据位于共享内存（Shared Memory） 中。

.b16: 指定内存访问模式。.b16 表示这是一次 16 字节的访问。这与共享内存的 bank 宽度和高效访问模式有关。

{%0, %1}: 这是目标操作数列表。占位符 %0 和 %1 将会被编译器替换为后面约束列表中找到的实际寄存器。这里它要求 2 个 32 位的寄存器。

为什么是 2 个？ 一个 8x8 的矩阵，每个元素 2 字节，总大小为 8 * 8 * 2 = 128 字节。一个 Warp 有 32 个线程。ldmatrix 指令将这 128 字节的数据转置后，分布到整个 Warp 的线程寄存器中。每个线程负责 128 / 32 = 4 字节的数据。一个 32 位寄存器是 4 字节，所以每个线程需要 1 个寄存器来存储它的那部分数据。那么为什么这里列表里有 2 个？实际上，这条指令是在加载2个这样的 8x8 矩阵。.m8n8.x2 加载一个矩阵需要 4 个寄存器？不，这里的关键在于指令的变体。在 SM_70+ 上，ldmatrix 可以加载 1、2 或 4 个矩阵。这里的语法暗示它正在加载 2 个 矩阵（虽然数字没有明确写出来，但目标寄存器的数量是 2，这对应于加载 2 个 .m8n8.x2 矩阵，每个线程获得 2 个寄存器，总共 8 字节的数据）。更常见的可能是加载 1 个矩阵需要 4 个寄存器？让我们澄清：对于 .m8n8.x2 加载，加载 1 个矩阵需要 4 个目标寄存器？不，标准理解是：加载 1 个 .m8n8.x2 矩阵到 Warp 中，需要 4 个寄存器（%0, %1, %2, %3）。加载 2 个需要 8 个寄存器。这个函数只用了 2 个，所以它可能是一个特例，或者用于加载更小的块？另一种可能是用于加载 4 个 .m8n8.x4 矩阵？不，这里是 .x2。可能这个函数名 x2 指的是“扩展”或其他含义，而不是元素大小。根据指令和参数数量，它很可能是用于加载2个.m8n8.x2`矩阵。每个线程获得 2 个寄存器（8 字节）的数据。

[%2]: 这是源操作数。它是一个包含共享内存地址的寄存器。%2 将被替换为输入操作数提供的值。

2.2.2. 输出操作数 (`: "=r"(x[0]), "=r"(x[1])`)

"=r": 约束修饰符。

= 表示这是一个只写的输出操作数。

r 表示要求编译器分配一个32 位通用寄存器来保存这个值。

(x[0]), (x[1]): 对应的 C++ 变量。指令执行后，目标寄存器 %0 和 %1 中的值会被写回到数组 x 的这两个元素中。

作用：告诉编译器：“请为 x[0] 和 x[1] 分配两个寄存器。执行汇编指令后，结果将在这两个寄存器中，请将它们写回 x[0] 和 x[1]。”

2.2.3. 输入操作数 (`: "l"(__cvta_generic_to_shared(ptr))`)

这是最精妙和关键的部分。

__cvta_generic_to_shared(ptr): 这是一个 CUDA 内部函数。

作用：它将一个通用指针 (ptr) 转换为其对应的共享内存空间下的地址值。

原理：在 PTX 中，不同的内存空间（全局、共享、本地等）有独立的地址空间。一个通用的 void* 指针不能直接用于 ldmatrix 的 shared 操作。这个函数执行必要的位操作，提取出专用于共享内存地址空间的地址比特位。

"l": 这是一个约束修饰符。

l 表示一个 32 位的专用寄存器，通常用于存储地址**。这与通用寄存器 r 略有不同，编译器知道这个寄存器将用于寻址。

作用：告诉编译器：“计算 __cvta_generic_to_shared(ptr) 这个表达式的值，并将其放入一个专用的地址寄存器中，然后在汇编模板中用 %2 来引用这个寄存器。”

2.2.4. `volatile` 关键字

防止编译器优化掉这条汇编指令（例如，因为它看起来没有使用输出 x），或者将其移出循环。确保指令严格按照代码中的位置和执行次数运行。

2.3. 函数功能总结

这个 ldmatrix_x2 函数的功能是：

让一个 Warp（32 个线程）协同工作，从共享内存中 ptr 所指的、16 字节对齐的地址开始，加载 2 个连续的 8x8 矩阵（每个元素 2 字节）。数据在加载过程中会被重新排列（转置）。加载完成后，每个线程会获得 8 字节（2 个 unsigned int）的数据，存储在其 x[0] 和 x[1] 中。

这些数据通常是更大矩阵乘法操作中的一个小块（Tile）。每个线程持有的 x[0] 和 x[1] 是转置后矩阵的一小部分，它们的形式非常适合直接作为输入喂给后续的 mma（矩阵乘加）指令，从而实现极其高效的矩阵计算。

注意事项：

调用约定：这个函数必须由整个 Warp 的线程同时调用，且 ptr 的值在 Warp 内必须一致（通常是通过广播获得）。
对齐：ptr 必须是 16 字节对齐的，否则行为未定义。
数据布局：共享内存中的数据必须按照 ldmatrix 指令所期望的布局进行排列，这通常由之前的数据存储步骤（例如使用 st.shared.v2.b32 之类的指令）来保证。

这个函数是手动优化 CUDA 核函数、充分发挥 Tensor Core 性能的典型代表。

3. ldmatrix 功能系统解析

CUDA PTX 中的 ldmatrix 指令是高效利用 Tensor Cores（张量核心）进行矩阵计算的关键所在。接着前面的具体实例，这里更为系统第介绍一下 ldmatrix 指令的原理用法。

3.1. 指令概述与原理

目的

ldmatrix（Load Matrix）指令用于从一个线程束（Warp）内线程协同访问的连续共享内存区域中，高效地加载一个小的、密集的矩阵块（如 8x8），并将其转置后分布到该 Warp 中多个线程的寄存器中。

核心思想

Tensor Cores 执行的是 D = A * B + C 操作，其中 A、B、C、D 都是小矩阵。然而，全局内存或共享内存中的数据通常按行主序或列主序存储。ldmatrix 指令在数据从共享内存加载到寄存器的过程中，巧妙地完成了数据重排（转置），使得数据在寄存器中的布局恰好符合 Tensor Cores 所期望的输入格式，从而避免了显式的转置操作，极大提升了效率。

工作原理

一个 Warp（32 个线程）共同协作，从共享内存中读取一片连续的数据。每个线程负责读取数据的一部分。指令会自动地将这些数据重新组织（转置），并存入指定线程的指定寄存器。最终，整个 Warp 的寄存器合在一起，就构成了一个完整的、经过转置的矩阵。

3.2. 指令语法格式

完整的 PTX 语法如下：

ldmatrix.sync.aligned.{num}{.trans}{.ss}.type [rd1, rd2, ...], [rs1, rs2]; 
// 或者更常见的格式，指定矩阵形状：
ldmatrix.sync.aligned.shape.{num}{.trans}{.ss}.rspace [rd1, rd2, ...], [rs];

3.3. 指令中各域详解

`.sync` (Synchronization)

作用

指定这是一个Warp-level 同步指令。指令的执行会涉及 Warp 中所有活跃线程的协同操作。.sync 后缀确保所有线程在逻辑上同时参与此次加载。

可选值

在较新的架构中，可以指定 .sync.syncid 以实现更细粒度的同步，但通常直接使用 .sync。

`.aligned` (Alignment)

作用

指定共享内存的源地址必须是 16 字节对齐的。这是为了满足内存子系统的高效访问要求。如果地址未对齐，执行结果将是未定义的。

注意

这是一个强制要求，不是可选项。你必须确保传入的共享内存指针是 16 字节对齐的。

`.{num}` (Number of Matrices)

作用

指定一次指令调用要加载的矩阵数量。

可选值

.1：加载 1 个矩阵;

.2：加载 2 个矩阵;.4：加载 4 个矩阵;

影响

加载的矩阵数量直接决定了目标寄存器的数量。例如，加载一个 8x8x16 的矩阵（.m8n8 + .x2）需要 4 个寄存器（8*8*2/32/1？更正：通常加载 1 个 .m8n8.x4 矩阵需要 8 个寄存器）。加载 .4 个矩阵就需要 4 倍数量的寄存器。

`.{trans}` (Transposition)

作用

指定是否对加载的矩阵进行转置。

可选值

(空)：不进行转置，按原样加载；

.trans：对加载的矩阵进行转置；

这是关键

这个功能是为了适配 Tensor Cores 的输入。例如，在计算 A * B 时，可能需要将 B 矩阵转置后再输入给 Tensor Core。使用 .trans 可以在加载时一步完成，无需后续单独的转置指令。

`.{ss}` (Element Size / Storage Spacing)

作用

指定源数据中每个矩阵元素的大小和存储间隔。

可选值

.x1：8 位元素（如 char, uint8_t）；.x2：16 位元素（如 half, __half, short）。这是用于 FP16 张量计算最常见的大小；

.x4：32 位元素（如 float, int）；

`.{type}` / `.{rspace}` (Type / Resource Space)

作用

指定源数据所在的内存空间。

可选值

.shared：源数据位于共享内存中。这是 ldmatrix 最常用、最主要的使用场景；

.global：源数据位于全局内存中。（在某些架构上支持，但不如从共享内存加载高效）；

`.[rd1, rd2, ...]` (Destination Registers)

作用

目标操作数，是一个寄存器列表，用于接收加载来的矩阵数据。

要求

寄存器的数量取决于 {num}, {ss} 和矩阵形状。例如，加载 1 个 8x8 的矩阵（.m8n8），每个元素是 32位（.x4），则需要 (8 * 8 * 4) / 32 = 8 个 32 位寄存器；

寄存器必须是 32 位宽的（例如 %r0, %f1）；

列表中的寄存器必须是连续的；

`.[rs1, rs2]` / `[rs]` (Source Address)

作用

源操作数，是包含共享内存地址的寄存器。

要求

通常是一个包含 32 位地址的寄存器（例如 %r0）；

该地址必须指向共享内存，并且必须是 16 字节对齐的（由 .aligned 保证）；

`.{shape}` (Matrix Shape - 替代方案)

作用

另一种语法是明确指定矩阵的形状，这通常更直观。

可选值

.m8n8：加载一个 8x8 的矩阵。这是最常用的形状；

.m8n8k4 等：用于更复杂的加载模式，但 .m8n8 是基础；

3.4. 用法示例与解释

假设我们要从共享内存加载一个 8x8 的 FP16 矩阵，并对其进行转置，然后分布到寄存器中。

PTX 代码：

ldmatrix.sync.aligned.m8n8.x2.trans.shared.b16 {%0, %1, %2, %3}, [%4];

分解：

.sync.aligned：Warp 同步且地址对齐；

.m8n8：加载 8x8 的矩阵；

.x2：源元素是 16 位（FP16）；

.trans：加载时进行转置；

.shared.b16：从共享内存以 16 字节的访问模式读取；

{%0, %1, %2, %3}：需要 4 个 32 位目标寄存器；

*计算：一个 8x8 FP16 矩阵总大小 = 8 * 8 * 2字节 = 128 字节。一个 Warp 有 32 个线程，每个线程负责 128 / 32 = 4 字节的数据。一个 32 位寄存器正好是 4 字节，所以每个线程需要 1 个寄存器。但为什么这里有 4 个？实际上，ldmatrix 指令的寄存器列表是每个线程持有的寄存器数量？不，更准确的说法是：这条指令为整个 Warp 指定了 4 个连续的寄存器，但每个线程看到的是这些寄存器中的不同部分。通常，加载一个 .m8n8.x2 矩阵需要 4 个目标寄存器。

[%4]：源地址寄存器，其值是一个 16 字节对齐的共享内存地址；

在 CUDA C++ 中的内联汇编用法：

__shared__ half smem_buffer[64]; // 8x8 FP16 矩阵asm volatile ("ldmatrix.sync.aligned.m8n8.x2.trans.shared.b16 {%0, %1, %2, %3}, [%4];": "=r"(reg0), "=r"(reg1), "=r"(reg2), "=r"(reg3) // 4个输出寄存器: "r"(smem_buffer) // 输入：共享内存地址// 可能还需要 clobber 列表，但有时可省略
);

3.5. 总结

ldmatrix 是一条极其强大的指令，它将数据加载和数据重排（转置） 两个耗时的操作合并为一条高效的硬件指令。它的设计完美契合了 Tensor Cores 的工作方式，是实现高性能矩阵乘法（尤其是深度学习推理和训练）的核心原语之一。理解其各个参数的含义对于在 PTX 或 CUDA 内联汇编中正确使用它至关重要。

4. 附录：PTX 代码

生成 ptx 文件：

nvcc -ptx -lineinfo -std=c++17 -arch=native main.cu -o main.ptx

或者不带源码行号

$ nvcc -ptx --gpu-architecture=sm_120 main.cu -o main_sm_120.ptx

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-35583870
// Cuda compilation tools, release 12.8, V12.8.93
// Based on NVVM 20.0.0
//.version 8.7
.target sm_120
.address_size 64// .globl	_Z8mykernelPKib
.extern .func  (.param .b32 func_retval0) vprintf
(.param .b64 vprintf_param_0,.param .b64 vprintf_param_1
)
;
// _ZZ8mykernelPKibE1A has been demoted
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__45__cpo5beginE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__45__cpo3endE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__45__cpo6cbeginE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__45__cpo4cendE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__45__cpo6rbeginE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__45__cpo4rendE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__45__cpo7crbeginE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__45__cpo5crendE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__439_GLOBAL__N__63e6e6d0_7_main_cu_0f110be86ignoreE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__419piecewise_constructE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__48in_placeE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__47nulloptE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std3__420unreachable_sentinelE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo4swapE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo9iter_moveE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo7advanceE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo5beginE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo3endE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo6cbeginE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo4cendE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo9iter_swapE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo4nextE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo4prevE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo4dataE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo5cdataE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo4sizeE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo5ssizeE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be84cuda3std6ranges3__45__cpo8distanceE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS8cuda_cub3parE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS8cuda_cub10par_nosyncE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS6system6detail10sequential3seqE[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS12placeholders2_1E[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS12placeholders2_2E[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS12placeholders2_3E[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS12placeholders2_4E[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS12placeholders2_5E[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS12placeholders2_6E[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS12placeholders2_7E[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS12placeholders2_8E[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS12placeholders2_9E[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS12placeholders3_10E[1];
.global .align 1 .b8 _ZN37_INTERNAL_63e6e6d0_7_main_cu_0f110be86thrust21THRUST_200700_1200_NS3seqE[1];
.global .align 1 .b8 $str[9] = {37, 52, 100, 32, 37, 52, 100, 32};
.global .align 1 .b8 $str$1[2] = {10};.visible .entry _Z8mykernelPKib(.param .u64 _Z8mykernelPKib_param_0,.param .u8 _Z8mykernelPKib_param_1
)
{.local .align 8 .b8 	__local_depot0[8];.reg .b64 	%SP;.reg .b64 	%SPL;.reg .pred 	%p<136>;.reg .b16 	%rs<145>;.reg .b32 	%r<451>;.reg .b64 	%rd<333>;.loc	1 24 0// demoted variable.shared .align 16 .b8 _ZZ8mykernelPKibE1A[4096];mov.u64 	%SPL, __local_depot0;cvta.local.u64 	%SP, %SPL;ld.param.u64 	%rd2, [_Z8mykernelPKib_param_0];ld.param.s8 	%rs7, [_Z8mykernelPKib_param_1];add.u64 	%rd3, %SP, 0;.loc	1 27 5add.u64 	%rd1, %SPL, 0;mov.u32 	%r1, %tid.x;mov.u32 	%r2, %ntid.x;add.s32 	%r29, %r1, %r2;cvt.u16.u32 	%rs8, %r29;mov.u16 	%rs9, 2047;sub.s16 	%rs10, %rs9, %rs8;cvt.u16.u32 	%rs11, %r2;div.u16 	%rs1, %rs10, %rs11;and.b16  	%rs2, %rs1, 3;setp.eq.s16 	%p1, %rs2, 2;mov.u32 	%r445, %r1;@%p1 bra 	$L__BB0_3;.loc	1 27 5cvt.u32.u16 	%r3, %rs2;mov.b32 	%r444, 0;mov.u32 	%r33, _ZZ8mykernelPKibE1A;mov.u32 	%r445, %r1;
$L__BB0_2:.pragma "nounroll";.loc	1 28 9add.s32 	%r32, %r445, %r445;add.s32 	%r34, %r33, %r32;.loc	1 28 9.loc	2 212 43, function_name $L__info_string0, inlined_at 1 28 9.loc	2 1017 1, function_name $L__info_string1, inlined_at 2 212 43// begin inline asmcvt.rn.f16.s32 %rs12, %r445;// end inline asm.loc	2 212 43, function_name $L__info_string0, inlined_at 1 28 9st.shared.u16 	[%r34], %rs12;.loc	1 27 42add.s32 	%r445, %r445, %r2;.loc	1 27 5add.s32 	%r444, %r444, 1;xor.b32  	%r35, %r444, %r3;setp.ne.s32 	%p2, %r35, 2;@%p2 bra 	$L__BB0_2;
$L__BB0_3:.loc	1 27 5setp.lt.u16 	%p3, %rs1, 2;@%p3 bra 	$L__BB0_6;.loc	1 27 5mad.lo.s32 	%r449, %r2, 3, %r445;shl.b32 	%r10, %r2, 2;add.s32 	%r11, %r2, %r2;add.s32 	%r448, %r445, %r11;add.s32 	%r447, %r2, %r445;mul.lo.s32 	%r14, %r2, 6;add.s32 	%r36, %r445, %r445;mov.u32 	%r37, _ZZ8mykernelPKibE1A;add.s32 	%r446, %r37, %r36;shl.b32 	%r16, %r2, 3;
$L__BB0_5:.loc	2 1017 1, function_name $L__info_string1, inlined_at 2 212 43// begin inline asmcvt.rn.f16.s32 %rs13, %r445;// end inline asm.loc	2 212 43, function_name $L__info_string0, inlined_at 1 28 9st.shared.u16 	[%r446], %rs13;.loc	1 27 42add.s32 	%r42, %r445, %r2;.loc	2 1017 1, function_name $L__info_string1, inlined_at 2 212 43add.s32 	%r43, %r446, %r11;// begin inline asmcvt.rn.f16.s32 %rs14, %r447;// end inline asm.loc	2 212 43, function_name $L__info_string0, inlined_at 1 28 9st.shared.u16 	[%r43], %rs14;.loc	1 27 42add.s32 	%r44, %r42, %r2;.loc	2 1017 1, function_name $L__info_string1, inlined_at 2 212 43add.s32 	%r45, %r446, %r10;// begin inline asmcvt.rn.f16.s32 %rs15, %r448;// end inline asm.loc	2 212 43, function_name $L__info_string0, inlined_at 1 28 9st.shared.u16 	[%r45], %rs15;.loc	1 27 42add.s32 	%r46, %r44, %r2;.loc	2 1017 1, function_name $L__info_string1, inlined_at 2 212 43add.s32 	%r47, %r446, %r14;// begin inline asmcvt.rn.f16.s32 %rs16, %r449;// end inline asm.loc	2 212 43, function_name $L__info_string0, inlined_at 1 28 9st.shared.u16 	[%r47], %rs16;.loc	1 27 42add.s32 	%r445, %r46, %r2;.loc	1 27 5add.s32 	%r449, %r449, %r10;add.s32 	%r448, %r448, %r10;add.s32 	%r447, %r447, %r10;add.s32 	%r446, %r446, %r16;setp.lt.u32 	%p4, %r445, 2048;@%p4 bra 	$L__BB0_5;
$L__BB0_6:.loc	1 30 5bar.sync 	0;.loc	1 32 5and.b32  	%r27, %r1, 31;cvta.to.global.u64 	%rd5, %rd2;.loc	1 35 5shl.b32 	%r50, %r1, 2;cvt.u64.u32 	%rd6, %r50;and.b64  	%rd7, %rd6, 124;add.s64 	%rd8, %rd5, %rd7;ld.global.u32 	%r51, [%rd8];.loc	1 36 5.loc	1 19 15, function_name $L__info_string2, inlined_at 1 36 5.loc	3 151 3, function_name $L__info_string3, inlined_at 1 19 15add.s32 	%r52, %r51, %r51;mov.u32 	%r53, _ZZ8mykernelPKibE1A;add.s32 	%r54, %r53, %r52;cvt.u64.u32 	%rd4, %r54;.loc	1 19 15, function_name $L__info_string2, inlined_at 1 36 5// begin inline asmldmatrix.sync.aligned.m8n8.x2.shared.b16 {%r48, %r49}, [%rd4];// end inline asm.loc	1 39 5cvt.u16.u32 	%rs3, %r48;{ .reg .b16 tmp; mov.b32 {tmp, %rs4}, %r48; }cvt.u16.u32 	%rs5, %r49;{ .reg .b16 tmp; mov.b32 {tmp, %rs6}, %r49; }.loc	1 41 5setp.eq.s16 	%p5, %rs7, 0;@%p5 bra 	$L__BB0_204;.loc	1 0 5and.b32  	%r28, %r1, 3;setp.ne.s32 	%p6, %r27, 0;.loc	1 44 17@%p6 bra 	$L__BB0_10;.loc	1 0 17setp.ne.s32 	%p7, %r28, 3;.loc	1 45 21.loc	2 166 35, function_name $L__info_string4, inlined_at 1 45 21.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r55, %rs3;// end inline asm.loc	1 45 21.loc	2 166 35, function_name $L__info_string4, inlined_at 1 45 21.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r56, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r55, %r56};mov.u64 	%rd9, $str;cvta.global.u64 	%rd10, %rd9;{ // callseq 0, 0.param .b64 param0;st.param.b64 	[param0+0], %rd10;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r57, [retval0+0];} // callseq 0.loc	1 46 21@%p7 bra 	$L__BB0_10;.loc	1 47 25mov.u64 	%rd12, $str$1;cvta.global.u64 	%rd13, %rd12;{ // callseq 1, 0.param .b64 param0;st.param.b64 	[param0+0], %rd13;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r59, [retval0+0];} // callseq 1
$L__BB0_10:.loc	1 50 17.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p8, %r27, 1;@%p8 bra 	$L__BB0_13;.loc	1 0 17setp.ne.s32 	%p9, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r61, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r62, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r61, %r62};mov.u64 	%rd14, $str;cvta.global.u64 	%rd15, %rd14;{ // callseq 2, 0.param .b64 param0;st.param.b64 	[param0+0], %rd15;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r63, [retval0+0];} // callseq 2.loc	1 46 21@%p9 bra 	$L__BB0_13;.loc	1 47 25mov.u64 	%rd17, $str$1;cvta.global.u64 	%rd18, %rd17;{ // callseq 3, 0.param .b64 param0;st.param.b64 	[param0+0], %rd18;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r65, [retval0+0];} // callseq 3
$L__BB0_13:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p10, %r27, 2;@%p10 bra 	$L__BB0_16;.loc	1 0 17setp.ne.s32 	%p11, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r67, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r68, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r67, %r68};mov.u64 	%rd19, $str;cvta.global.u64 	%rd20, %rd19;{ // callseq 4, 0.param .b64 param0;st.param.b64 	[param0+0], %rd20;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r69, [retval0+0];} // callseq 4.loc	1 46 21@%p11 bra 	$L__BB0_16;.loc	1 47 25mov.u64 	%rd22, $str$1;cvta.global.u64 	%rd23, %rd22;{ // callseq 5, 0.param .b64 param0;st.param.b64 	[param0+0], %rd23;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r71, [retval0+0];} // callseq 5
$L__BB0_16:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p12, %r27, 3;@%p12 bra 	$L__BB0_19;.loc	1 0 17setp.ne.s32 	%p13, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r73, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r74, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r73, %r74};mov.u64 	%rd24, $str;cvta.global.u64 	%rd25, %rd24;{ // callseq 6, 0.param .b64 param0;st.param.b64 	[param0+0], %rd25;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r75, [retval0+0];} // callseq 6.loc	1 46 21@%p13 bra 	$L__BB0_19;.loc	1 47 25mov.u64 	%rd27, $str$1;cvta.global.u64 	%rd28, %rd27;{ // callseq 7, 0.param .b64 param0;st.param.b64 	[param0+0], %rd28;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r77, [retval0+0];} // callseq 7
$L__BB0_19:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p14, %r27, 4;@%p14 bra 	$L__BB0_22;.loc	1 0 17setp.ne.s32 	%p15, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r79, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r80, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r79, %r80};mov.u64 	%rd29, $str;cvta.global.u64 	%rd30, %rd29;{ // callseq 8, 0.param .b64 param0;st.param.b64 	[param0+0], %rd30;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r81, [retval0+0];} // callseq 8.loc	1 46 21@%p15 bra 	$L__BB0_22;.loc	1 47 25mov.u64 	%rd32, $str$1;cvta.global.u64 	%rd33, %rd32;{ // callseq 9, 0.param .b64 param0;st.param.b64 	[param0+0], %rd33;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r83, [retval0+0];} // callseq 9
$L__BB0_22:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p16, %r27, 5;@%p16 bra 	$L__BB0_25;.loc	1 0 17setp.ne.s32 	%p17, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r85, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r86, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r85, %r86};mov.u64 	%rd34, $str;cvta.global.u64 	%rd35, %rd34;{ // callseq 10, 0.param .b64 param0;st.param.b64 	[param0+0], %rd35;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r87, [retval0+0];} // callseq 10.loc	1 46 21@%p17 bra 	$L__BB0_25;.loc	1 47 25mov.u64 	%rd37, $str$1;cvta.global.u64 	%rd38, %rd37;{ // callseq 11, 0.param .b64 param0;st.param.b64 	[param0+0], %rd38;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r89, [retval0+0];} // callseq 11
$L__BB0_25:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p18, %r27, 6;@%p18 bra 	$L__BB0_28;.loc	1 0 17setp.ne.s32 	%p19, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r91, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r92, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r91, %r92};mov.u64 	%rd39, $str;cvta.global.u64 	%rd40, %rd39;{ // callseq 12, 0.param .b64 param0;st.param.b64 	[param0+0], %rd40;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r93, [retval0+0];} // callseq 12.loc	1 46 21@%p19 bra 	$L__BB0_28;.loc	1 47 25mov.u64 	%rd42, $str$1;cvta.global.u64 	%rd43, %rd42;{ // callseq 13, 0.param .b64 param0;st.param.b64 	[param0+0], %rd43;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r95, [retval0+0];} // callseq 13
$L__BB0_28:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p20, %r27, 7;@%p20 bra 	$L__BB0_31;.loc	1 0 17setp.ne.s32 	%p21, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r97, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r98, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r97, %r98};mov.u64 	%rd44, $str;cvta.global.u64 	%rd45, %rd44;{ // callseq 14, 0.param .b64 param0;st.param.b64 	[param0+0], %rd45;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r99, [retval0+0];} // callseq 14.loc	1 46 21@%p21 bra 	$L__BB0_31;.loc	1 47 25mov.u64 	%rd47, $str$1;cvta.global.u64 	%rd48, %rd47;{ // callseq 15, 0.param .b64 param0;st.param.b64 	[param0+0], %rd48;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r101, [retval0+0];} // callseq 15
$L__BB0_31:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p22, %r27, 8;@%p22 bra 	$L__BB0_34;.loc	1 0 17setp.ne.s32 	%p23, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r103, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r104, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r103, %r104};mov.u64 	%rd49, $str;cvta.global.u64 	%rd50, %rd49;{ // callseq 16, 0.param .b64 param0;st.param.b64 	[param0+0], %rd50;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r105, [retval0+0];} // callseq 16.loc	1 46 21@%p23 bra 	$L__BB0_34;.loc	1 47 25mov.u64 	%rd52, $str$1;cvta.global.u64 	%rd53, %rd52;{ // callseq 17, 0.param .b64 param0;st.param.b64 	[param0+0], %rd53;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r107, [retval0+0];} // callseq 17
$L__BB0_34:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p24, %r27, 9;@%p24 bra 	$L__BB0_37;.loc	1 0 17setp.ne.s32 	%p25, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r109, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r110, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r109, %r110};mov.u64 	%rd54, $str;cvta.global.u64 	%rd55, %rd54;{ // callseq 18, 0.param .b64 param0;st.param.b64 	[param0+0], %rd55;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r111, [retval0+0];} // callseq 18.loc	1 46 21@%p25 bra 	$L__BB0_37;.loc	1 47 25mov.u64 	%rd57, $str$1;cvta.global.u64 	%rd58, %rd57;{ // callseq 19, 0.param .b64 param0;st.param.b64 	[param0+0], %rd58;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r113, [retval0+0];} // callseq 19
$L__BB0_37:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p26, %r27, 10;@%p26 bra 	$L__BB0_40;.loc	1 0 17setp.ne.s32 	%p27, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r115, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r116, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r115, %r116};mov.u64 	%rd59, $str;cvta.global.u64 	%rd60, %rd59;{ // callseq 20, 0.param .b64 param0;st.param.b64 	[param0+0], %rd60;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r117, [retval0+0];} // callseq 20.loc	1 46 21@%p27 bra 	$L__BB0_40;.loc	1 47 25mov.u64 	%rd62, $str$1;cvta.global.u64 	%rd63, %rd62;{ // callseq 21, 0.param .b64 param0;st.param.b64 	[param0+0], %rd63;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r119, [retval0+0];} // callseq 21
$L__BB0_40:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p28, %r27, 11;@%p28 bra 	$L__BB0_43;.loc	1 0 17setp.ne.s32 	%p29, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r121, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r122, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r121, %r122};mov.u64 	%rd64, $str;cvta.global.u64 	%rd65, %rd64;{ // callseq 22, 0.param .b64 param0;st.param.b64 	[param0+0], %rd65;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r123, [retval0+0];} // callseq 22.loc	1 46 21@%p29 bra 	$L__BB0_43;.loc	1 47 25mov.u64 	%rd67, $str$1;cvta.global.u64 	%rd68, %rd67;{ // callseq 23, 0.param .b64 param0;st.param.b64 	[param0+0], %rd68;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r125, [retval0+0];} // callseq 23
$L__BB0_43:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p30, %r27, 12;@%p30 bra 	$L__BB0_46;.loc	1 0 17setp.ne.s32 	%p31, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r127, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r128, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r127, %r128};mov.u64 	%rd69, $str;cvta.global.u64 	%rd70, %rd69;{ // callseq 24, 0.param .b64 param0;st.param.b64 	[param0+0], %rd70;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r129, [retval0+0];} // callseq 24.loc	1 46 21@%p31 bra 	$L__BB0_46;.loc	1 47 25mov.u64 	%rd72, $str$1;cvta.global.u64 	%rd73, %rd72;{ // callseq 25, 0.param .b64 param0;st.param.b64 	[param0+0], %rd73;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r131, [retval0+0];} // callseq 25
$L__BB0_46:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p32, %r27, 13;@%p32 bra 	$L__BB0_49;.loc	1 0 17setp.ne.s32 	%p33, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r133, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r134, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r133, %r134};mov.u64 	%rd74, $str;cvta.global.u64 	%rd75, %rd74;{ // callseq 26, 0.param .b64 param0;st.param.b64 	[param0+0], %rd75;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r135, [retval0+0];} // callseq 26.loc	1 46 21@%p33 bra 	$L__BB0_49;.loc	1 47 25mov.u64 	%rd77, $str$1;cvta.global.u64 	%rd78, %rd77;{ // callseq 27, 0.param .b64 param0;st.param.b64 	[param0+0], %rd78;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r137, [retval0+0];} // callseq 27
$L__BB0_49:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p34, %r27, 14;@%p34 bra 	$L__BB0_52;.loc	1 0 17setp.ne.s32 	%p35, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r139, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r140, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r139, %r140};mov.u64 	%rd79, $str;cvta.global.u64 	%rd80, %rd79;{ // callseq 28, 0.param .b64 param0;st.param.b64 	[param0+0], %rd80;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r141, [retval0+0];} // callseq 28.loc	1 46 21@%p35 bra 	$L__BB0_52;.loc	1 47 25mov.u64 	%rd82, $str$1;cvta.global.u64 	%rd83, %rd82;{ // callseq 29, 0.param .b64 param0;st.param.b64 	[param0+0], %rd83;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r143, [retval0+0];} // callseq 29
$L__BB0_52:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p36, %r27, 15;@%p36 bra 	$L__BB0_55;.loc	1 0 17setp.ne.s32 	%p37, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r145, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r146, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r145, %r146};mov.u64 	%rd84, $str;cvta.global.u64 	%rd85, %rd84;{ // callseq 30, 0.param .b64 param0;st.param.b64 	[param0+0], %rd85;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r147, [retval0+0];} // callseq 30.loc	1 46 21@%p37 bra 	$L__BB0_55;.loc	1 47 25mov.u64 	%rd87, $str$1;cvta.global.u64 	%rd88, %rd87;{ // callseq 31, 0.param .b64 param0;st.param.b64 	[param0+0], %rd88;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r149, [retval0+0];} // callseq 31
$L__BB0_55:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p38, %r27, 16;@%p38 bra 	$L__BB0_58;.loc	1 0 17setp.ne.s32 	%p39, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r151, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r152, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r151, %r152};mov.u64 	%rd89, $str;cvta.global.u64 	%rd90, %rd89;{ // callseq 32, 0.param .b64 param0;st.param.b64 	[param0+0], %rd90;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r153, [retval0+0];} // callseq 32.loc	1 46 21@%p39 bra 	$L__BB0_58;.loc	1 47 25mov.u64 	%rd92, $str$1;cvta.global.u64 	%rd93, %rd92;{ // callseq 33, 0.param .b64 param0;st.param.b64 	[param0+0], %rd93;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r155, [retval0+0];} // callseq 33
$L__BB0_58:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p40, %r27, 17;@%p40 bra 	$L__BB0_61;.loc	1 0 17setp.ne.s32 	%p41, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r157, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r158, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r157, %r158};mov.u64 	%rd94, $str;cvta.global.u64 	%rd95, %rd94;{ // callseq 34, 0.param .b64 param0;st.param.b64 	[param0+0], %rd95;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r159, [retval0+0];} // callseq 34.loc	1 46 21@%p41 bra 	$L__BB0_61;.loc	1 47 25mov.u64 	%rd97, $str$1;cvta.global.u64 	%rd98, %rd97;{ // callseq 35, 0.param .b64 param0;st.param.b64 	[param0+0], %rd98;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r161, [retval0+0];} // callseq 35
$L__BB0_61:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p42, %r27, 18;@%p42 bra 	$L__BB0_64;.loc	1 0 17setp.ne.s32 	%p43, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r163, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r164, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r163, %r164};mov.u64 	%rd99, $str;cvta.global.u64 	%rd100, %rd99;{ // callseq 36, 0.param .b64 param0;st.param.b64 	[param0+0], %rd100;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r165, [retval0+0];} // callseq 36.loc	1 46 21@%p43 bra 	$L__BB0_64;.loc	1 47 25mov.u64 	%rd102, $str$1;cvta.global.u64 	%rd103, %rd102;{ // callseq 37, 0.param .b64 param0;st.param.b64 	[param0+0], %rd103;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r167, [retval0+0];} // callseq 37
$L__BB0_64:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p44, %r27, 19;@%p44 bra 	$L__BB0_67;.loc	1 0 17setp.ne.s32 	%p45, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r169, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r170, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r169, %r170};mov.u64 	%rd104, $str;cvta.global.u64 	%rd105, %rd104;{ // callseq 38, 0.param .b64 param0;st.param.b64 	[param0+0], %rd105;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r171, [retval0+0];} // callseq 38.loc	1 46 21@%p45 bra 	$L__BB0_67;.loc	1 47 25mov.u64 	%rd107, $str$1;cvta.global.u64 	%rd108, %rd107;{ // callseq 39, 0.param .b64 param0;st.param.b64 	[param0+0], %rd108;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r173, [retval0+0];} // callseq 39
$L__BB0_67:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p46, %r27, 20;@%p46 bra 	$L__BB0_70;.loc	1 0 17setp.ne.s32 	%p47, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r175, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r176, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r175, %r176};mov.u64 	%rd109, $str;cvta.global.u64 	%rd110, %rd109;{ // callseq 40, 0.param .b64 param0;st.param.b64 	[param0+0], %rd110;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r177, [retval0+0];} // callseq 40.loc	1 46 21@%p47 bra 	$L__BB0_70;.loc	1 47 25mov.u64 	%rd112, $str$1;cvta.global.u64 	%rd113, %rd112;{ // callseq 41, 0.param .b64 param0;st.param.b64 	[param0+0], %rd113;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r179, [retval0+0];} // callseq 41
$L__BB0_70:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p48, %r27, 21;@%p48 bra 	$L__BB0_73;.loc	1 0 17setp.ne.s32 	%p49, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r181, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r182, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r181, %r182};mov.u64 	%rd114, $str;cvta.global.u64 	%rd115, %rd114;{ // callseq 42, 0.param .b64 param0;st.param.b64 	[param0+0], %rd115;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r183, [retval0+0];} // callseq 42.loc	1 46 21@%p49 bra 	$L__BB0_73;.loc	1 47 25mov.u64 	%rd117, $str$1;cvta.global.u64 	%rd118, %rd117;{ // callseq 43, 0.param .b64 param0;st.param.b64 	[param0+0], %rd118;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r185, [retval0+0];} // callseq 43
$L__BB0_73:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p50, %r27, 22;@%p50 bra 	$L__BB0_76;.loc	1 0 17setp.ne.s32 	%p51, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r187, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r188, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r187, %r188};mov.u64 	%rd119, $str;cvta.global.u64 	%rd120, %rd119;{ // callseq 44, 0.param .b64 param0;st.param.b64 	[param0+0], %rd120;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r189, [retval0+0];} // callseq 44.loc	1 46 21@%p51 bra 	$L__BB0_76;.loc	1 47 25mov.u64 	%rd122, $str$1;cvta.global.u64 	%rd123, %rd122;{ // callseq 45, 0.param .b64 param0;st.param.b64 	[param0+0], %rd123;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r191, [retval0+0];} // callseq 45
$L__BB0_76:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p52, %r27, 23;@%p52 bra 	$L__BB0_79;.loc	1 0 17setp.ne.s32 	%p53, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r193, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r194, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r193, %r194};mov.u64 	%rd124, $str;cvta.global.u64 	%rd125, %rd124;{ // callseq 46, 0.param .b64 param0;st.param.b64 	[param0+0], %rd125;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r195, [retval0+0];} // callseq 46.loc	1 46 21@%p53 bra 	$L__BB0_79;.loc	1 47 25mov.u64 	%rd127, $str$1;cvta.global.u64 	%rd128, %rd127;{ // callseq 47, 0.param .b64 param0;st.param.b64 	[param0+0], %rd128;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r197, [retval0+0];} // callseq 47
$L__BB0_79:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p54, %r27, 24;@%p54 bra 	$L__BB0_82;.loc	1 0 17setp.ne.s32 	%p55, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r199, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r200, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r199, %r200};mov.u64 	%rd129, $str;cvta.global.u64 	%rd130, %rd129;{ // callseq 48, 0.param .b64 param0;st.param.b64 	[param0+0], %rd130;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r201, [retval0+0];} // callseq 48.loc	1 46 21@%p55 bra 	$L__BB0_82;.loc	1 47 25mov.u64 	%rd132, $str$1;cvta.global.u64 	%rd133, %rd132;{ // callseq 49, 0.param .b64 param0;st.param.b64 	[param0+0], %rd133;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r203, [retval0+0];} // callseq 49
$L__BB0_82:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p56, %r27, 25;@%p56 bra 	$L__BB0_85;.loc	1 0 17setp.ne.s32 	%p57, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r205, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r206, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r205, %r206};mov.u64 	%rd134, $str;cvta.global.u64 	%rd135, %rd134;{ // callseq 50, 0.param .b64 param0;st.param.b64 	[param0+0], %rd135;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r207, [retval0+0];} // callseq 50.loc	1 46 21@%p57 bra 	$L__BB0_85;.loc	1 47 25mov.u64 	%rd137, $str$1;cvta.global.u64 	%rd138, %rd137;{ // callseq 51, 0.param .b64 param0;st.param.b64 	[param0+0], %rd138;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r209, [retval0+0];} // callseq 51
$L__BB0_85:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p58, %r27, 26;@%p58 bra 	$L__BB0_88;.loc	1 0 17setp.ne.s32 	%p59, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r211, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r212, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r211, %r212};mov.u64 	%rd139, $str;cvta.global.u64 	%rd140, %rd139;{ // callseq 52, 0.param .b64 param0;st.param.b64 	[param0+0], %rd140;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r213, [retval0+0];} // callseq 52.loc	1 46 21@%p59 bra 	$L__BB0_88;.loc	1 47 25mov.u64 	%rd142, $str$1;cvta.global.u64 	%rd143, %rd142;{ // callseq 53, 0.param .b64 param0;st.param.b64 	[param0+0], %rd143;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r215, [retval0+0];} // callseq 53
$L__BB0_88:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p60, %r27, 27;@%p60 bra 	$L__BB0_91;.loc	1 0 17setp.ne.s32 	%p61, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r217, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r218, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r217, %r218};mov.u64 	%rd144, $str;cvta.global.u64 	%rd145, %rd144;{ // callseq 54, 0.param .b64 param0;st.param.b64 	[param0+0], %rd145;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r219, [retval0+0];} // callseq 54.loc	1 46 21@%p61 bra 	$L__BB0_91;.loc	1 47 25mov.u64 	%rd147, $str$1;cvta.global.u64 	%rd148, %rd147;{ // callseq 55, 0.param .b64 param0;st.param.b64 	[param0+0], %rd148;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r221, [retval0+0];} // callseq 55
$L__BB0_91:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p62, %r27, 28;@%p62 bra 	$L__BB0_94;.loc	1 0 17setp.ne.s32 	%p63, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r223, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r224, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r223, %r224};mov.u64 	%rd149, $str;cvta.global.u64 	%rd150, %rd149;{ // callseq 56, 0.param .b64 param0;st.param.b64 	[param0+0], %rd150;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r225, [retval0+0];} // callseq 56.loc	1 46 21@%p63 bra 	$L__BB0_94;.loc	1 47 25mov.u64 	%rd152, $str$1;cvta.global.u64 	%rd153, %rd152;{ // callseq 57, 0.param .b64 param0;st.param.b64 	[param0+0], %rd153;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r227, [retval0+0];} // callseq 57
$L__BB0_94:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p64, %r27, 29;@%p64 bra 	$L__BB0_97;.loc	1 0 17setp.ne.s32 	%p65, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r229, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r230, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r229, %r230};mov.u64 	%rd154, $str;cvta.global.u64 	%rd155, %rd154;{ // callseq 58, 0.param .b64 param0;st.param.b64 	[param0+0], %rd155;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r231, [retval0+0];} // callseq 58.loc	1 46 21@%p65 bra 	$L__BB0_97;.loc	1 47 25mov.u64 	%rd157, $str$1;cvta.global.u64 	%rd158, %rd157;{ // callseq 59, 0.param .b64 param0;st.param.b64 	[param0+0], %rd158;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r233, [retval0+0];} // callseq 59
$L__BB0_97:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p66, %r27, 30;@%p66 bra 	$L__BB0_100;.loc	1 0 17setp.ne.s32 	%p67, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r235, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r236, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r235, %r236};mov.u64 	%rd159, $str;cvta.global.u64 	%rd160, %rd159;{ // callseq 60, 0.param .b64 param0;st.param.b64 	[param0+0], %rd160;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r237, [retval0+0];} // callseq 60.loc	1 46 21@%p67 bra 	$L__BB0_100;.loc	1 47 25mov.u64 	%rd162, $str$1;cvta.global.u64 	%rd163, %rd162;{ // callseq 61, 0.param .b64 param0;st.param.b64 	[param0+0], %rd163;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r239, [retval0+0];} // callseq 61
$L__BB0_100:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17setp.ne.s32 	%p68, %r27, 31;@%p68 bra 	$L__BB0_103;.loc	1 0 17setp.ne.s32 	%p69, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r241, %rs3;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r242, %rs4;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r241, %r242};mov.u64 	%rd164, $str;cvta.global.u64 	%rd165, %rd164;{ // callseq 62, 0.param .b64 param0;st.param.b64 	[param0+0], %rd165;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r243, [retval0+0];} // callseq 62.loc	1 46 21@%p69 bra 	$L__BB0_103;.loc	1 47 25mov.u64 	%rd167, $str$1;cvta.global.u64 	%rd168, %rd167;{ // callseq 63, 0.param .b64 param0;st.param.b64 	[param0+0], %rd168;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r245, [retval0+0];} // callseq 63
$L__BB0_103:.loc	1 0 25setp.ne.s32 	%p70, %r27, 0;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 52 13@%p70 bra 	$L__BB0_105;.loc	1 53 17mov.u64 	%rd169, $str$1;cvta.global.u64 	%rd170, %rd169;{ // callseq 64, 0.param .b64 param0;st.param.b64 	[param0+0], %rd170;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r247, [retval0+0];} // callseq 64
$L__BB0_105:.loc	1 0 17setp.ne.s32 	%p71, %r27, 0;.loc	1 55 13.loc	4 110 3, function_name $L__info_string6, inlined_at 1 55 13bar.warp.sync 	-1;.loc	1 44 17@%p71 bra 	$L__BB0_108;.loc	1 0 17setp.ne.s32 	%p72, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r249, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r250, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r249, %r250};mov.u64 	%rd171, $str;cvta.global.u64 	%rd172, %rd171;{ // callseq 65, 0.param .b64 param0;st.param.b64 	[param0+0], %rd172;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r251, [retval0+0];} // callseq 65.loc	1 46 21@%p72 bra 	$L__BB0_108;.loc	1 47 25mov.u64 	%rd174, $str$1;cvta.global.u64 	%rd175, %rd174;{ // callseq 66, 0.param .b64 param0;st.param.b64 	[param0+0], %rd175;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r253, [retval0+0];} // callseq 66
$L__BB0_108:.loc	1 44 17setp.ne.s32 	%p73, %r27, 1;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p73 bra 	$L__BB0_111;.loc	1 0 17setp.ne.s32 	%p74, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r255, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r256, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r255, %r256};mov.u64 	%rd176, $str;cvta.global.u64 	%rd177, %rd176;{ // callseq 67, 0.param .b64 param0;st.param.b64 	[param0+0], %rd177;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r257, [retval0+0];} // callseq 67.loc	1 46 21@%p74 bra 	$L__BB0_111;.loc	1 47 25mov.u64 	%rd179, $str$1;cvta.global.u64 	%rd180, %rd179;{ // callseq 68, 0.param .b64 param0;st.param.b64 	[param0+0], %rd180;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r259, [retval0+0];} // callseq 68
$L__BB0_111:.loc	1 44 17setp.ne.s32 	%p75, %r27, 2;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p75 bra 	$L__BB0_114;.loc	1 0 17setp.ne.s32 	%p76, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r261, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r262, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r261, %r262};mov.u64 	%rd181, $str;cvta.global.u64 	%rd182, %rd181;{ // callseq 69, 0.param .b64 param0;st.param.b64 	[param0+0], %rd182;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r263, [retval0+0];} // callseq 69.loc	1 46 21@%p76 bra 	$L__BB0_114;.loc	1 47 25mov.u64 	%rd184, $str$1;cvta.global.u64 	%rd185, %rd184;{ // callseq 70, 0.param .b64 param0;st.param.b64 	[param0+0], %rd185;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r265, [retval0+0];} // callseq 70
$L__BB0_114:.loc	1 44 17setp.ne.s32 	%p77, %r27, 3;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p77 bra 	$L__BB0_117;.loc	1 0 17setp.ne.s32 	%p78, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r267, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r268, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r267, %r268};mov.u64 	%rd186, $str;cvta.global.u64 	%rd187, %rd186;{ // callseq 71, 0.param .b64 param0;st.param.b64 	[param0+0], %rd187;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r269, [retval0+0];} // callseq 71.loc	1 46 21@%p78 bra 	$L__BB0_117;.loc	1 47 25mov.u64 	%rd189, $str$1;cvta.global.u64 	%rd190, %rd189;{ // callseq 72, 0.param .b64 param0;st.param.b64 	[param0+0], %rd190;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r271, [retval0+0];} // callseq 72
$L__BB0_117:.loc	1 44 17setp.ne.s32 	%p79, %r27, 4;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p79 bra 	$L__BB0_120;.loc	1 0 17setp.ne.s32 	%p80, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r273, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r274, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r273, %r274};mov.u64 	%rd191, $str;cvta.global.u64 	%rd192, %rd191;{ // callseq 73, 0.param .b64 param0;st.param.b64 	[param0+0], %rd192;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r275, [retval0+0];} // callseq 73.loc	1 46 21@%p80 bra 	$L__BB0_120;.loc	1 47 25mov.u64 	%rd194, $str$1;cvta.global.u64 	%rd195, %rd194;{ // callseq 74, 0.param .b64 param0;st.param.b64 	[param0+0], %rd195;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r277, [retval0+0];} // callseq 74
$L__BB0_120:.loc	1 44 17setp.ne.s32 	%p81, %r27, 5;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p81 bra 	$L__BB0_123;.loc	1 0 17setp.ne.s32 	%p82, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r279, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r280, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r279, %r280};mov.u64 	%rd196, $str;cvta.global.u64 	%rd197, %rd196;{ // callseq 75, 0.param .b64 param0;st.param.b64 	[param0+0], %rd197;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r281, [retval0+0];} // callseq 75.loc	1 46 21@%p82 bra 	$L__BB0_123;.loc	1 47 25mov.u64 	%rd199, $str$1;cvta.global.u64 	%rd200, %rd199;{ // callseq 76, 0.param .b64 param0;st.param.b64 	[param0+0], %rd200;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r283, [retval0+0];} // callseq 76
$L__BB0_123:.loc	1 44 17setp.ne.s32 	%p83, %r27, 6;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p83 bra 	$L__BB0_126;.loc	1 0 17setp.ne.s32 	%p84, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r285, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r286, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r285, %r286};mov.u64 	%rd201, $str;cvta.global.u64 	%rd202, %rd201;{ // callseq 77, 0.param .b64 param0;st.param.b64 	[param0+0], %rd202;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r287, [retval0+0];} // callseq 77.loc	1 46 21@%p84 bra 	$L__BB0_126;.loc	1 47 25mov.u64 	%rd204, $str$1;cvta.global.u64 	%rd205, %rd204;{ // callseq 78, 0.param .b64 param0;st.param.b64 	[param0+0], %rd205;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r289, [retval0+0];} // callseq 78
$L__BB0_126:.loc	1 44 17setp.ne.s32 	%p85, %r27, 7;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p85 bra 	$L__BB0_129;.loc	1 0 17setp.ne.s32 	%p86, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r291, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r292, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r291, %r292};mov.u64 	%rd206, $str;cvta.global.u64 	%rd207, %rd206;{ // callseq 79, 0.param .b64 param0;st.param.b64 	[param0+0], %rd207;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r293, [retval0+0];} // callseq 79.loc	1 46 21@%p86 bra 	$L__BB0_129;.loc	1 47 25mov.u64 	%rd209, $str$1;cvta.global.u64 	%rd210, %rd209;{ // callseq 80, 0.param .b64 param0;st.param.b64 	[param0+0], %rd210;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r295, [retval0+0];} // callseq 80
$L__BB0_129:.loc	1 44 17setp.ne.s32 	%p87, %r27, 8;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p87 bra 	$L__BB0_132;.loc	1 0 17setp.ne.s32 	%p88, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r297, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r298, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r297, %r298};mov.u64 	%rd211, $str;cvta.global.u64 	%rd212, %rd211;{ // callseq 81, 0.param .b64 param0;st.param.b64 	[param0+0], %rd212;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r299, [retval0+0];} // callseq 81.loc	1 46 21@%p88 bra 	$L__BB0_132;.loc	1 47 25mov.u64 	%rd214, $str$1;cvta.global.u64 	%rd215, %rd214;{ // callseq 82, 0.param .b64 param0;st.param.b64 	[param0+0], %rd215;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r301, [retval0+0];} // callseq 82
$L__BB0_132:.loc	1 44 17setp.ne.s32 	%p89, %r27, 9;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p89 bra 	$L__BB0_135;.loc	1 0 17setp.ne.s32 	%p90, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r303, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r304, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r303, %r304};mov.u64 	%rd216, $str;cvta.global.u64 	%rd217, %rd216;{ // callseq 83, 0.param .b64 param0;st.param.b64 	[param0+0], %rd217;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r305, [retval0+0];} // callseq 83.loc	1 46 21@%p90 bra 	$L__BB0_135;.loc	1 47 25mov.u64 	%rd219, $str$1;cvta.global.u64 	%rd220, %rd219;{ // callseq 84, 0.param .b64 param0;st.param.b64 	[param0+0], %rd220;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r307, [retval0+0];} // callseq 84
$L__BB0_135:.loc	1 44 17setp.ne.s32 	%p91, %r27, 10;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p91 bra 	$L__BB0_138;.loc	1 0 17setp.ne.s32 	%p92, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r309, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r310, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r309, %r310};mov.u64 	%rd221, $str;cvta.global.u64 	%rd222, %rd221;{ // callseq 85, 0.param .b64 param0;st.param.b64 	[param0+0], %rd222;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r311, [retval0+0];} // callseq 85.loc	1 46 21@%p92 bra 	$L__BB0_138;.loc	1 47 25mov.u64 	%rd224, $str$1;cvta.global.u64 	%rd225, %rd224;{ // callseq 86, 0.param .b64 param0;st.param.b64 	[param0+0], %rd225;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r313, [retval0+0];} // callseq 86
$L__BB0_138:.loc	1 44 17setp.ne.s32 	%p93, %r27, 11;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p93 bra 	$L__BB0_141;.loc	1 0 17setp.ne.s32 	%p94, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r315, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r316, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r315, %r316};mov.u64 	%rd226, $str;cvta.global.u64 	%rd227, %rd226;{ // callseq 87, 0.param .b64 param0;st.param.b64 	[param0+0], %rd227;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r317, [retval0+0];} // callseq 87.loc	1 46 21@%p94 bra 	$L__BB0_141;.loc	1 47 25mov.u64 	%rd229, $str$1;cvta.global.u64 	%rd230, %rd229;{ // callseq 88, 0.param .b64 param0;st.param.b64 	[param0+0], %rd230;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r319, [retval0+0];} // callseq 88
$L__BB0_141:.loc	1 44 17setp.ne.s32 	%p95, %r27, 12;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p95 bra 	$L__BB0_144;.loc	1 0 17setp.ne.s32 	%p96, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r321, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r322, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r321, %r322};mov.u64 	%rd231, $str;cvta.global.u64 	%rd232, %rd231;{ // callseq 89, 0.param .b64 param0;st.param.b64 	[param0+0], %rd232;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r323, [retval0+0];} // callseq 89.loc	1 46 21@%p96 bra 	$L__BB0_144;.loc	1 47 25mov.u64 	%rd234, $str$1;cvta.global.u64 	%rd235, %rd234;{ // callseq 90, 0.param .b64 param0;st.param.b64 	[param0+0], %rd235;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r325, [retval0+0];} // callseq 90
$L__BB0_144:.loc	1 44 17setp.ne.s32 	%p97, %r27, 13;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p97 bra 	$L__BB0_147;.loc	1 0 17setp.ne.s32 	%p98, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r327, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r328, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r327, %r328};mov.u64 	%rd236, $str;cvta.global.u64 	%rd237, %rd236;{ // callseq 91, 0.param .b64 param0;st.param.b64 	[param0+0], %rd237;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r329, [retval0+0];} // callseq 91.loc	1 46 21@%p98 bra 	$L__BB0_147;.loc	1 47 25mov.u64 	%rd239, $str$1;cvta.global.u64 	%rd240, %rd239;{ // callseq 92, 0.param .b64 param0;st.param.b64 	[param0+0], %rd240;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r331, [retval0+0];} // callseq 92
$L__BB0_147:.loc	1 44 17setp.ne.s32 	%p99, %r27, 14;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p99 bra 	$L__BB0_150;.loc	1 0 17setp.ne.s32 	%p100, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r333, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r334, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r333, %r334};mov.u64 	%rd241, $str;cvta.global.u64 	%rd242, %rd241;{ // callseq 93, 0.param .b64 param0;st.param.b64 	[param0+0], %rd242;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r335, [retval0+0];} // callseq 93.loc	1 46 21@%p100 bra 	$L__BB0_150;.loc	1 47 25mov.u64 	%rd244, $str$1;cvta.global.u64 	%rd245, %rd244;{ // callseq 94, 0.param .b64 param0;st.param.b64 	[param0+0], %rd245;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r337, [retval0+0];} // callseq 94
$L__BB0_150:.loc	1 44 17setp.ne.s32 	%p101, %r27, 15;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p101 bra 	$L__BB0_153;.loc	1 0 17setp.ne.s32 	%p102, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r339, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r340, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r339, %r340};mov.u64 	%rd246, $str;cvta.global.u64 	%rd247, %rd246;{ // callseq 95, 0.param .b64 param0;st.param.b64 	[param0+0], %rd247;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r341, [retval0+0];} // callseq 95.loc	1 46 21@%p102 bra 	$L__BB0_153;.loc	1 47 25mov.u64 	%rd249, $str$1;cvta.global.u64 	%rd250, %rd249;{ // callseq 96, 0.param .b64 param0;st.param.b64 	[param0+0], %rd250;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r343, [retval0+0];} // callseq 96
$L__BB0_153:.loc	1 44 17setp.ne.s32 	%p103, %r27, 16;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p103 bra 	$L__BB0_156;.loc	1 0 17setp.ne.s32 	%p104, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r345, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r346, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r345, %r346};mov.u64 	%rd251, $str;cvta.global.u64 	%rd252, %rd251;{ // callseq 97, 0.param .b64 param0;st.param.b64 	[param0+0], %rd252;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r347, [retval0+0];} // callseq 97.loc	1 46 21@%p104 bra 	$L__BB0_156;.loc	1 47 25mov.u64 	%rd254, $str$1;cvta.global.u64 	%rd255, %rd254;{ // callseq 98, 0.param .b64 param0;st.param.b64 	[param0+0], %rd255;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r349, [retval0+0];} // callseq 98
$L__BB0_156:.loc	1 44 17setp.ne.s32 	%p105, %r27, 17;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p105 bra 	$L__BB0_159;.loc	1 0 17setp.ne.s32 	%p106, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r351, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r352, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r351, %r352};mov.u64 	%rd256, $str;cvta.global.u64 	%rd257, %rd256;{ // callseq 99, 0.param .b64 param0;st.param.b64 	[param0+0], %rd257;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r353, [retval0+0];} // callseq 99.loc	1 46 21@%p106 bra 	$L__BB0_159;.loc	1 47 25mov.u64 	%rd259, $str$1;cvta.global.u64 	%rd260, %rd259;{ // callseq 100, 0.param .b64 param0;st.param.b64 	[param0+0], %rd260;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r355, [retval0+0];} // callseq 100
$L__BB0_159:.loc	1 44 17setp.ne.s32 	%p107, %r27, 18;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p107 bra 	$L__BB0_162;.loc	1 0 17setp.ne.s32 	%p108, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r357, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r358, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r357, %r358};mov.u64 	%rd261, $str;cvta.global.u64 	%rd262, %rd261;{ // callseq 101, 0.param .b64 param0;st.param.b64 	[param0+0], %rd262;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r359, [retval0+0];} // callseq 101.loc	1 46 21@%p108 bra 	$L__BB0_162;.loc	1 47 25mov.u64 	%rd264, $str$1;cvta.global.u64 	%rd265, %rd264;{ // callseq 102, 0.param .b64 param0;st.param.b64 	[param0+0], %rd265;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r361, [retval0+0];} // callseq 102
$L__BB0_162:.loc	1 44 17setp.ne.s32 	%p109, %r27, 19;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p109 bra 	$L__BB0_165;.loc	1 0 17setp.ne.s32 	%p110, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r363, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r364, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r363, %r364};mov.u64 	%rd266, $str;cvta.global.u64 	%rd267, %rd266;{ // callseq 103, 0.param .b64 param0;st.param.b64 	[param0+0], %rd267;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r365, [retval0+0];} // callseq 103.loc	1 46 21@%p110 bra 	$L__BB0_165;.loc	1 47 25mov.u64 	%rd269, $str$1;cvta.global.u64 	%rd270, %rd269;{ // callseq 104, 0.param .b64 param0;st.param.b64 	[param0+0], %rd270;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r367, [retval0+0];} // callseq 104
$L__BB0_165:.loc	1 44 17setp.ne.s32 	%p111, %r27, 20;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p111 bra 	$L__BB0_168;.loc	1 0 17setp.ne.s32 	%p112, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r369, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r370, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r369, %r370};mov.u64 	%rd271, $str;cvta.global.u64 	%rd272, %rd271;{ // callseq 105, 0.param .b64 param0;st.param.b64 	[param0+0], %rd272;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r371, [retval0+0];} // callseq 105.loc	1 46 21@%p112 bra 	$L__BB0_168;.loc	1 47 25mov.u64 	%rd274, $str$1;cvta.global.u64 	%rd275, %rd274;{ // callseq 106, 0.param .b64 param0;st.param.b64 	[param0+0], %rd275;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r373, [retval0+0];} // callseq 106
$L__BB0_168:.loc	1 44 17setp.ne.s32 	%p113, %r27, 21;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p113 bra 	$L__BB0_171;.loc	1 0 17setp.ne.s32 	%p114, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r375, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r376, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r375, %r376};mov.u64 	%rd276, $str;cvta.global.u64 	%rd277, %rd276;{ // callseq 107, 0.param .b64 param0;st.param.b64 	[param0+0], %rd277;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r377, [retval0+0];} // callseq 107.loc	1 46 21@%p114 bra 	$L__BB0_171;.loc	1 47 25mov.u64 	%rd279, $str$1;cvta.global.u64 	%rd280, %rd279;{ // callseq 108, 0.param .b64 param0;st.param.b64 	[param0+0], %rd280;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r379, [retval0+0];} // callseq 108
$L__BB0_171:.loc	1 44 17setp.ne.s32 	%p115, %r27, 22;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p115 bra 	$L__BB0_174;.loc	1 0 17setp.ne.s32 	%p116, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r381, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r382, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r381, %r382};mov.u64 	%rd281, $str;cvta.global.u64 	%rd282, %rd281;{ // callseq 109, 0.param .b64 param0;st.param.b64 	[param0+0], %rd282;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r383, [retval0+0];} // callseq 109.loc	1 46 21@%p116 bra 	$L__BB0_174;.loc	1 47 25mov.u64 	%rd284, $str$1;cvta.global.u64 	%rd285, %rd284;{ // callseq 110, 0.param .b64 param0;st.param.b64 	[param0+0], %rd285;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r385, [retval0+0];} // callseq 110
$L__BB0_174:.loc	1 44 17setp.ne.s32 	%p117, %r27, 23;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p117 bra 	$L__BB0_177;.loc	1 0 17setp.ne.s32 	%p118, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r387, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r388, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r387, %r388};mov.u64 	%rd286, $str;cvta.global.u64 	%rd287, %rd286;{ // callseq 111, 0.param .b64 param0;st.param.b64 	[param0+0], %rd287;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r389, [retval0+0];} // callseq 111.loc	1 46 21@%p118 bra 	$L__BB0_177;.loc	1 47 25mov.u64 	%rd289, $str$1;cvta.global.u64 	%rd290, %rd289;{ // callseq 112, 0.param .b64 param0;st.param.b64 	[param0+0], %rd290;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r391, [retval0+0];} // callseq 112
$L__BB0_177:.loc	1 44 17setp.ne.s32 	%p119, %r27, 24;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p119 bra 	$L__BB0_180;.loc	1 0 17setp.ne.s32 	%p120, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r393, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r394, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r393, %r394};mov.u64 	%rd291, $str;cvta.global.u64 	%rd292, %rd291;{ // callseq 113, 0.param .b64 param0;st.param.b64 	[param0+0], %rd292;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r395, [retval0+0];} // callseq 113.loc	1 46 21@%p120 bra 	$L__BB0_180;.loc	1 47 25mov.u64 	%rd294, $str$1;cvta.global.u64 	%rd295, %rd294;{ // callseq 114, 0.param .b64 param0;st.param.b64 	[param0+0], %rd295;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r397, [retval0+0];} // callseq 114
$L__BB0_180:.loc	1 44 17setp.ne.s32 	%p121, %r27, 25;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p121 bra 	$L__BB0_183;.loc	1 0 17setp.ne.s32 	%p122, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r399, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r400, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r399, %r400};mov.u64 	%rd296, $str;cvta.global.u64 	%rd297, %rd296;{ // callseq 115, 0.param .b64 param0;st.param.b64 	[param0+0], %rd297;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r401, [retval0+0];} // callseq 115.loc	1 46 21@%p122 bra 	$L__BB0_183;.loc	1 47 25mov.u64 	%rd299, $str$1;cvta.global.u64 	%rd300, %rd299;{ // callseq 116, 0.param .b64 param0;st.param.b64 	[param0+0], %rd300;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r403, [retval0+0];} // callseq 116
$L__BB0_183:.loc	1 44 17setp.ne.s32 	%p123, %r27, 26;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p123 bra 	$L__BB0_186;.loc	1 0 17setp.ne.s32 	%p124, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r405, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r406, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r405, %r406};mov.u64 	%rd301, $str;cvta.global.u64 	%rd302, %rd301;{ // callseq 117, 0.param .b64 param0;st.param.b64 	[param0+0], %rd302;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r407, [retval0+0];} // callseq 117.loc	1 46 21@%p124 bra 	$L__BB0_186;.loc	1 47 25mov.u64 	%rd304, $str$1;cvta.global.u64 	%rd305, %rd304;{ // callseq 118, 0.param .b64 param0;st.param.b64 	[param0+0], %rd305;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r409, [retval0+0];} // callseq 118
$L__BB0_186:.loc	1 44 17setp.ne.s32 	%p125, %r27, 27;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p125 bra 	$L__BB0_189;.loc	1 0 17setp.ne.s32 	%p126, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r411, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r412, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r411, %r412};mov.u64 	%rd306, $str;cvta.global.u64 	%rd307, %rd306;{ // callseq 119, 0.param .b64 param0;st.param.b64 	[param0+0], %rd307;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r413, [retval0+0];} // callseq 119.loc	1 46 21@%p126 bra 	$L__BB0_189;.loc	1 47 25mov.u64 	%rd309, $str$1;cvta.global.u64 	%rd310, %rd309;{ // callseq 120, 0.param .b64 param0;st.param.b64 	[param0+0], %rd310;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r415, [retval0+0];} // callseq 120
$L__BB0_189:.loc	1 44 17setp.ne.s32 	%p127, %r27, 28;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p127 bra 	$L__BB0_192;.loc	1 0 17setp.ne.s32 	%p128, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r417, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r418, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r417, %r418};mov.u64 	%rd311, $str;cvta.global.u64 	%rd312, %rd311;{ // callseq 121, 0.param .b64 param0;st.param.b64 	[param0+0], %rd312;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r419, [retval0+0];} // callseq 121.loc	1 46 21@%p128 bra 	$L__BB0_192;.loc	1 47 25mov.u64 	%rd314, $str$1;cvta.global.u64 	%rd315, %rd314;{ // callseq 122, 0.param .b64 param0;st.param.b64 	[param0+0], %rd315;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r421, [retval0+0];} // callseq 122
$L__BB0_192:.loc	1 44 17setp.ne.s32 	%p129, %r27, 29;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p129 bra 	$L__BB0_195;.loc	1 0 17setp.ne.s32 	%p130, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r423, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r424, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r423, %r424};mov.u64 	%rd316, $str;cvta.global.u64 	%rd317, %rd316;{ // callseq 123, 0.param .b64 param0;st.param.b64 	[param0+0], %rd317;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r425, [retval0+0];} // callseq 123.loc	1 46 21@%p130 bra 	$L__BB0_195;.loc	1 47 25mov.u64 	%rd319, $str$1;cvta.global.u64 	%rd320, %rd319;{ // callseq 124, 0.param .b64 param0;st.param.b64 	[param0+0], %rd320;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r427, [retval0+0];} // callseq 124
$L__BB0_195:.loc	1 44 17setp.ne.s32 	%p131, %r27, 30;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p131 bra 	$L__BB0_198;.loc	1 0 17setp.ne.s32 	%p132, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r429, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r430, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r429, %r430};mov.u64 	%rd321, $str;cvta.global.u64 	%rd322, %rd321;{ // callseq 125, 0.param .b64 param0;st.param.b64 	[param0+0], %rd322;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r431, [retval0+0];} // callseq 125.loc	1 46 21@%p132 bra 	$L__BB0_198;.loc	1 47 25mov.u64 	%rd324, $str$1;cvta.global.u64 	%rd325, %rd324;{ // callseq 126, 0.param .b64 param0;st.param.b64 	[param0+0], %rd325;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r433, [retval0+0];} // callseq 126
$L__BB0_198:.loc	1 44 17setp.ne.s32 	%p133, %r27, 31;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 44 17@%p133 bra 	$L__BB0_201;.loc	1 0 17setp.ne.s32 	%p134, %r28, 3;.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r435, %rs5;// end inline asm.loc	2 858 1, function_name $L__info_string5, inlined_at 2 166 35// begin inline asmcvt.rzi.s32.f16 %r436, %rs6;// end inline asm.loc	1 45 21st.local.v2.u32 	[%rd1], {%r435, %r436};mov.u64 	%rd326, $str;cvta.global.u64 	%rd327, %rd326;{ // callseq 127, 0.param .b64 param0;st.param.b64 	[param0+0], %rd327;.param .b64 param1;st.param.b64 	[param1+0], %rd3;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r437, [retval0+0];} // callseq 127.loc	1 46 21@%p134 bra 	$L__BB0_201;.loc	1 47 25mov.u64 	%rd329, $str$1;cvta.global.u64 	%rd330, %rd329;{ // callseq 128, 0.param .b64 param0;st.param.b64 	[param0+0], %rd330;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r439, [retval0+0];} // callseq 128
$L__BB0_201:.loc	1 0 25setp.ne.s32 	%p135, %r27, 0;.loc	4 110 3, function_name $L__info_string6, inlined_at 1 50 17bar.warp.sync 	-1;.loc	1 52 13@%p135 bra 	$L__BB0_203;.loc	1 53 17mov.u64 	%rd331, $str$1;cvta.global.u64 	%rd332, %rd331;{ // callseq 129, 0.param .b64 param0;st.param.b64 	[param0+0], %rd332;.param .b64 param1;st.param.b64 	[param1+0], 0;.param .b32 retval0;call.uni (retval0), vprintf, (param0, param1);ld.param.b32 	%r441, [retval0+0];} // callseq 129
$L__BB0_203:.loc	4 110 3, function_name $L__info_string6, inlined_at 1 55 13bar.warp.sync 	-1;
$L__BB0_204:.loc	1 58 1ret;}// .globl	_ZN3cub18CUB_200700_1200_NS11EmptyKernelIvEEvv
.visible .entry _ZN3cub18CUB_200700_1200_NS11EmptyKernelIvEEvv()
{.loc	5 90 0.loc	5 91 1ret;}// .globl	_ZN3cub18CUB_200700_1200_NS6detail8for_each13static_kernelINS2_12policy_hub_t12policy_350_tEmN6thrust21THRUST_200700_1200_NS8cuda_cub20__uninitialized_fill7functorINS7_10device_ptrIiEEiEEEEvT0_T1_
.visible .entry _ZN3cub18CUB_200700_1200_NS6detail8for_each13static_kernelINS2_12policy_hub_t12policy_350_tEmN6thrust21THRUST_200700_1200_NS8cuda_cub20__uninitialized_fill7functorINS7_10device_ptrIiEEiEEEEvT0_T1_(.param .u64 _ZN3cub18CUB_200700_1200_NS6detail8for_each13static_kernelINS2_12policy_hub_t12policy_350_tEmN6thrust21THRUST_200700_1200_NS8cuda_cub20__uninitialized_fill7functorINS7_10device_ptrIiEEiEEEEvT0_T1__param_0,.param .align 8 .b8 _ZN3cub18CUB_200700_1200_NS6detail8for_each13static_kernelINS2_12policy_hub_t12policy_350_tEmN6thrust21THRUST_200700_1200_NS8cuda_cub20__uninitialized_fill7functorINS7_10device_ptrIiEEiEEEEvT0_T1__param_1[16]
)
.maxntid 256, 1, 1
{.reg .pred 	%p<4>;.reg .b16 	%rs<5>;.reg .b32 	%r<12>;.reg .b64 	%rd<16>;.loc	6 129 0ld.param.u32 	%r3, [_ZN3cub18CUB_200700_1200_NS6detail8for_each13static_kernelINS2_12policy_hub_t12policy_350_tEmN6thrust21THRUST_200700_1200_NS8cuda_cub20__uninitialized_fill7functorINS7_10device_ptrIiEEiEEEEvT0_T1__param_1+8];ld.param.u64 	%rd4, [_ZN3cub18CUB_200700_1200_NS6detail8for_each13static_kernelINS2_12policy_hub_t12policy_350_tEmN6thrust21THRUST_200700_1200_NS8cuda_cub20__uninitialized_fill7functorINS7_10device_ptrIiEEiEEEEvT0_T1__param_1];ld.param.u64 	%rd5, [_ZN3cub18CUB_200700_1200_NS6detail8for_each13static_kernelINS2_12policy_hub_t12policy_350_tEmN6thrust21THRUST_200700_1200_NS8cuda_cub20__uninitialized_fill7functorINS7_10device_ptrIiEEiEEEEvT0_T1__param_0];.loc	6 137 3mov.u32 	%r9, %ctaid.x;mul.wide.u32 	%rd1, %r9, 512;.loc	6 138 3sub.s64 	%rd2, %rd5, %rd1;.loc	6 141 3setp.lt.u64 	%p1, %rd2, 512;@%p1 bra 	$L__BB2_2;.loc	6 0 3mov.u32 	%r11, %tid.x;.loc	6 143 5cvta.to.global.u64 	%rd11, %rd4;.loc	6 143 5.loc	7 71 7, function_name $L__info_string7, inlined_at 6 143 5cvt.u64.u32 	%rd12, %r11;.loc	7 75 9, function_name $L__info_string7, inlined_at 6 143 5add.s64 	%rd13, %rd1, %rd12;.loc	7 75 9, function_name $L__info_string7, inlined_at 6 143 5.loc	8 72 5, function_name $L__info_string8, inlined_at 7 75 9.loc	9 381 5, function_name $L__info_string9, inlined_at 8 72 5.loc	9 630 3, function_name $L__info_string10, inlined_at 9 381 5.loc	9 427 5, function_name $L__info_string11, inlined_at 9 630 3.loc	9 233 5, function_name $L__info_string12, inlined_at 9 427 5.loc	10 214 5, function_name $L__info_string13, inlined_at 9 233 5shl.b64 	%rd14, %rd13, 2;add.s64 	%rd15, %rd11, %rd14;.loc	8 78 5, function_name $L__info_string8, inlined_at 7 75 9st.global.u32 	[%rd15], %r3;st.global.u32 	[%rd15+1024], %r3;.loc	6 149 1bra.uni 	$L__BB2_6;
$L__BB2_2:.loc	6 147 5cvta.to.global.u64 	%rd6, %rd4;mov.u32 	%r2, %tid.x;.loc	6 147 5.loc	7 71 7, function_name $L__info_string14, inlined_at 6 147 5cvt.u64.u32 	%rd7, %r2;.loc	7 73 7, function_name $L__info_string14, inlined_at 6 147 5setp.le.u64 	%p2, %rd2, %rd7;.loc	7 75 9, function_name $L__info_string14, inlined_at 6 147 5add.s64 	%rd8, %rd1, %rd7;.loc	7 75 9, function_name $L__info_string14, inlined_at 6 147 5.loc	8 72 5, function_name $L__info_string8, inlined_at 7 75 9.loc	9 381 5, function_name $L__info_string9, inlined_at 8 72 5.loc	9 630 3, function_name $L__info_string10, inlined_at 9 381 5.loc	9 427 5, function_name $L__info_string11, inlined_at 9 630 3.loc	9 233 5, function_name $L__info_string12, inlined_at 9 427 5.loc	10 214 5, function_name $L__info_string13, inlined_at 9 233 5shl.b64 	%rd9, %rd8, 2;add.s64 	%rd3, %rd6, %rd9;.loc	7 73 7, function_name $L__info_string14, inlined_at 6 147 5@%p2 bra 	$L__BB2_4;.loc	8 78 5, function_name $L__info_string8, inlined_at 7 75 9st.global.u32 	[%rd3], %r3;
$L__BB2_4:.loc	7 71 7, function_name $L__info_string14, inlined_at 6 147 5add.s32 	%r10, %r2, 256;cvt.u64.u32 	%rd10, %r10;.loc	7 73 7, function_name $L__info_string14, inlined_at 6 147 5setp.le.u64 	%p3, %rd2, %rd10;@%p3 bra 	$L__BB2_6;.loc	8 78 5, function_name $L__info_string8, inlined_at 7 75 9st.global.u32 	[%rd3+1024], %r3;
$L__BB2_6:.loc	6 149 1ret;}.file	1 "/home/hipper/ex_ldmatrix/main.cu".file	2 "/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda_fp16.hpp".file	3 "/usr/local/cuda/bin/../targets/x86_64-linux/include/sm_20_intrinsics.hpp".file	4 "/usr/local/cuda/bin/../targets/x86_64-linux/include/sm_30_intrinsics.hpp".file	5 "/usr/local/cuda/bin/../targets/x86_64-linux/include/cub/util_device.cuh".file	6 "/usr/local/cuda/bin/../targets/x86_64-linux/include/cub/device/dispatch/kernels/for_each.cuh".file	7 "/usr/local/cuda/bin/../targets/x86_64-linux/include/cub/agent/agent_for.cuh".file	8 "/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/system/cuda/detail/uninitialized_fill.h".file	9 "/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_facade.h".file	10 "/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_adaptor.h".section	.debug_str{
$L__info_string0:
.b8 95
.b8 90
.b8 78
.b8 54
.b8 95
.b8 95
.b8 104
.b8 97
.b8 108
.b8 102
.b8 97
.b8 83
.b8 69
.b8 105
.b8 0
$L__info_string1:
.b8 95
.b8 90
.b8 78
.b8 51
.b8 55
.b8 95
.b8 73
.b8 78
.b8 84
.b8 69
.b8 82
.b8 78
.b8 65
.b8 76
.b8 95
.b8 54
.b8 51
.b8 101
.b8 54
.b8 101
.b8 54
.b8 100
.b8 48
.b8 95
.b8 55
.b8 95
.b8 109
.b8 97
.b8 105
.b8 110
.b8 95
.b8 99
.b8 117
.b8 95
.b8 48
.b8 102
.b8 49
.b8 49
.b8 48
.b8 98
.b8 101
.b8 56
.b8 49
.b8 51
.b8 95
.b8 95
.b8 105
.b8 110
.b8 116
.b8 50
.b8 104
.b8 97
.b8 108
.b8 102
.b8 95
.b8 114
.b8 110
.b8 69
.b8 105
.b8 0
$L__info_string2:
.b8 95
.b8 90
.b8 49
.b8 49
.b8 108
.b8 100
.b8 109
.b8 97
.b8 116
.b8 114
.b8 105
.b8 120
.b8 95
.b8 120
.b8 50
.b8 82
.b8 65
.b8 50
.b8 95
.b8 106
.b8 80
.b8 75
.b8 118
.b8 0
$L__info_string3:
.b8 95
.b8 90
.b8 78
.b8 51
.b8 55
.b8 95
.b8 73
.b8 78
.b8 84
.b8 69
.b8 82
.b8 78
.b8 65
.b8 76
.b8 95
.b8 54
.b8 51
.b8 101
.b8 54
.b8 101
.b8 54
.b8 100
.b8 48
.b8 95
.b8 55
.b8 95
.b8 109
.b8 97
.b8 105
.b8 110
.b8 95
.b8 99
.b8 117
.b8 95
.b8 48
.b8 102
.b8 49
.b8 49
.b8 48
.b8 98
.b8 101
.b8 56
.b8 50
.b8 52
.b8 95
.b8 95
.b8 99
.b8 118
.b8 116
.b8 97
.b8 95
.b8 103
.b8 101
.b8 110
.b8 101
.b8 114
.b8 105
.b8 99
.b8 95
.b8 116
.b8 111
.b8 95
.b8 115
.b8 104
.b8 97
.b8 114
.b8 101
.b8 100
.b8 69
.b8 80
.b8 75
.b8 118
.b8 0
$L__info_string4:
.b8 95
.b8 90
.b8 78
.b8 75
.b8 54
.b8 95
.b8 95
.b8 104
.b8 97
.b8 108
.b8 102
.b8 99
.b8 118
.b8 105
.b8 69
.b8 118
.b8 0
$L__info_string5:
.b8 95
.b8 90
.b8 78
.b8 51
.b8 55
.b8 95
.b8 73
.b8 78
.b8 84
.b8 69
.b8 82
.b8 78
.b8 65
.b8 76
.b8 95
.b8 54
.b8 51
.b8 101
.b8 54
.b8 101
.b8 54
.b8 100
.b8 48
.b8 95
.b8 55
.b8 95
.b8 109
.b8 97
.b8 105
.b8 110
.b8 95
.b8 99
.b8 117
.b8 95
.b8 48
.b8 102
.b8 49
.b8 49
.b8 48
.b8 98
.b8 101
.b8 56
.b8 49
.b8 51
.b8 95
.b8 95
.b8 104
.b8 97
.b8 108
.b8 102
.b8 50
.b8 105
.b8 110
.b8 116
.b8 95
.b8 114
.b8 122
.b8 69
.b8 54
.b8 95
.b8 95
.b8 104
.b8 97
.b8 108
.b8 102
.b8 0
$L__info_string6:
.b8 95
.b8 90
.b8 78
.b8 51
.b8 55
.b8 95
.b8 73
.b8 78
.b8 84
.b8 69
.b8 82
.b8 78
.b8 65
.b8 76
.b8 95
.b8 54
.b8 51
.b8 101
.b8 54
.b8 101
.b8 54
.b8 100
.b8 48
.b8 95
.b8 55
.b8 95
.b8 109
.b8 97
.b8 105
.b8 110
.b8 95
.b8 99
.b8 117
.b8 95
.b8 48
.b8 102
.b8 49
.b8 49
.b8 48
.b8 98
.b8 101
.b8 56
.b8 49
.b8 48
.b8 95
.b8 95
.b8 115
.b8 121
.b8 110
.b8 99
.b8 119
.b8 97
.b8 114
.b8 112
.b8 69
.b8 106
.b8 0
$L__info_string7:
.b8 95
.b8 90
.b8 78
.b8 51
.b8 99
.b8 117
.b8 98
.b8 49
.b8 56
.b8 67
.b8 85
.b8 66
.b8 95
.b8 50
.b8 48
.b8 48
.b8 55
.b8 48
.b8 48
.b8 95
.b8 49
.b8 50
.b8 48
.b8 48
.b8 95
.b8 78
.b8 83
.b8 54
.b8 100
.b8 101
.b8 116
.b8 97
.b8 105
.b8 108
.b8 56
.b8 102
.b8 111
.b8 114
.b8 95
.b8 101
.b8 97
.b8 99
.b8 104
.b8 50
.b8 49
.b8 97
.b8 103
.b8 101
.b8 110
.b8 116
.b8 95
.b8 98
.b8 108
.b8 111
.b8 99
.b8 107
.b8 95
.b8 115
.b8 116
.b8 114
.b8 105
.b8 112
.b8 101
.b8 100
.b8 95
.b8 116
.b8 73
.b8 78
.b8 83
.b8 50
.b8 95
.b8 56
.b8 112
.b8 111
.b8 108
.b8 105
.b8 99
.b8 121
.b8 95
.b8 116
.b8 73
.b8 76
.b8 105
.b8 50
.b8 53
.b8 54
.b8 69
.b8 76
.b8 105
.b8 50
.b8 69
.b8 69
.b8 69
.b8 109
.b8 78
.b8 54
.b8 116
.b8 104
.b8 114
.b8 117
.b8 115
.b8 116
.b8 50
.b8 49
.b8 84
.b8 72
.b8 82
.b8 85
.b8 83
.b8 84
.b8 95
.b8 50
.b8 48
.b8 48
.b8 55
.b8 48
.b8 48
.b8 95
.b8 49
.b8 50
.b8 48
.b8 48
.b8 95
.b8 78
.b8 83
.b8 56
.b8 99
.b8 117
.b8 100
.b8 97
.b8 95
.b8 99
.b8 117
.b8 98
.b8 50
.b8 48
.b8 95
.b8 95
.b8 117
.b8 110
.b8 105
.b8 110
.b8 105
.b8 116
.b8 105
.b8 97
.b8 108
.b8 105
.b8 122
.b8 101
.b8 100
.b8 95
.b8 102
.b8 105
.b8 108
.b8 108
.b8 55
.b8 102
.b8 117
.b8 110
.b8 99
.b8 116
.b8 111
.b8 114
.b8 73
.b8 78
.b8 83
.b8 55
.b8 95
.b8 49
.b8 48
.b8 100
.b8 101
.b8 118
.b8 105
.b8 99
.b8 101
.b8 95
.b8 112
.b8 116
.b8 114
.b8 73
.b8 105
.b8 69
.b8 69
.b8 105
.b8 69
.b8 69
.b8 69
.b8 49
.b8 50
.b8 99
.b8 111
.b8 110
.b8 115
.b8 117
.b8 109
.b8 101
.b8 95
.b8 116
.b8 105
.b8 108
.b8 101
.b8 73
.b8 76
.b8 98
.b8 49
.b8 69
.b8 69
.b8 69
.b8 118
.b8 105
.b8 105
.b8 0
$L__info_string8:
.b8 95
.b8 90
.b8 78
.b8 54
.b8 116
.b8 104
.b8 114
.b8 117
.b8 115
.b8 116
.b8 50
.b8 49
.b8 84
.b8 72
.b8 82
.b8 85
.b8 83
.b8 84
.b8 95
.b8 50
.b8 48
.b8 48
.b8 55
.b8 48
.b8 48
.b8 95
.b8 49
.b8 50
.b8 48
.b8 48
.b8 95
.b8 78
.b8 83
.b8 56
.b8 99
.b8 117
.b8 100
.b8 97
.b8 95
.b8 99
.b8 117
.b8 98
.b8 50
.b8 48
.b8 95
.b8 95
.b8 117
.b8 110
.b8 105
.b8 110
.b8 105
.b8 116
.b8 105
.b8 97
.b8 108
.b8 105
.b8 122
.b8 101
.b8 100
.b8 95
.b8 102
.b8 105
.b8 108
.b8 108
.b8 55
.b8 102
.b8 117
.b8 110
.b8 99
.b8 116
.b8 111
.b8 114
.b8 73
.b8 78
.b8 83
.b8 48
.b8 95
.b8 49
.b8 48
.b8 100
.b8 101
.b8 118
.b8 105
.b8 99
.b8 101
.b8 95
.b8 112
.b8 116
.b8 114
.b8 73
.b8 105
.b8 69
.b8 69
.b8 105
.b8 69
.b8 99
.b8 108
.b8 73
.b8 109
.b8 69
.b8 69
.b8 118
.b8 84
.b8 95
.b8 0
$L__info_string9:
.b8 95
.b8 90
.b8 78
.b8 75
.b8 54
.b8 116
.b8 104
.b8 114
.b8 117
.b8 115
.b8 116
.b8 50
.b8 49
.b8 84
.b8 72
.b8 82
.b8 85
.b8 83
.b8 84
.b8 95
.b8 50
.b8 48
.b8 48
.b8 55
.b8 48
.b8 48
.b8 95
.b8 49
.b8 50
.b8 48
.b8 48
.b8 95
.b8 78
.b8 83
.b8 49
.b8 53
.b8 105
.b8 116
.b8 101
.b8 114
.b8 97
.b8 116
.b8 111
.b8 114
.b8 95
.b8 102
.b8 97
.b8 99
.b8 97
.b8 100
.b8 101
.b8 73
.b8 78
.b8 83
.b8 48
.b8 95
.b8 49
.b8 48
.b8 100
.b8 101
.b8 118
.b8 105
.b8 99
.b8 101
.b8 95
.b8 112
.b8 116
.b8 114
.b8 73
.b8 105
.b8 69
.b8 69
.b8 105
.b8 78
.b8 83
.b8 48
.b8 95
.b8 56
.b8 99
.b8 117
.b8 100
.b8 97
.b8 95
.b8 99
.b8 117
.b8 98
.b8 51
.b8 116
.b8 97
.b8 103
.b8 69
.b8 78
.b8 83
.b8 48
.b8 95
.b8 50
.b8 55
.b8 114
.b8 97
.b8 110
.b8 100
.b8 111
.b8 109
.b8 95
.b8 97
.b8 99
.b8 99
.b8 101
.b8 115
.b8 115
.b8 95
.b8 116
.b8 114
.b8 97
.b8 118
.b8 101
.b8 114
.b8 115
.b8 97
.b8 108
.b8 95
.b8 116
.b8 97
.b8 103
.b8 69
.b8 78
.b8 83
.b8 48
.b8 95
.b8 49
.b8 54
.b8 100
.b8 101
.b8 118
.b8 105
.b8 99
.b8 101
.b8 95
.b8 114
.b8 101
.b8 102
.b8 101
.b8 114
.b8 101
.b8 110
.b8 99
.b8 101
.b8 73
.b8 105
.b8 69
.b8 69
.b8 108
.b8 69
.b8 105
.b8 120
.b8 69
.b8 108
.b8 0
$L__info_string10:
.b8 95
.b8 90
.b8 78
.b8 54
.b8 116
.b8 104
.b8 114
.b8 117
.b8 115
.b8 116
.b8 50
.b8 49
.b8 84
.b8 72
.b8 82
.b8 85
.b8 83
.b8 84
.b8 95
.b8 50
.b8 48
.b8 48
.b8 55
.b8 48
.b8 48
.b8 95
.b8 49
.b8 50
.b8 48
.b8 48
.b8 95
.b8 78
.b8 83
.b8 112
.b8 108
.b8 73
.b8 78
.b8 83
.b8 48
.b8 95
.b8 49
.b8 48
.b8 100
.b8 101
.b8 118
.b8 105
.b8 99
.b8 101
.b8 95
.b8 112
.b8 116
.b8 114
.b8 73
.b8 105
.b8 69
.b8 69
.b8 105
.b8 78
.b8 83
.b8 48
.b8 95
.b8 56
.b8 99
.b8 117
.b8 100
.b8 97
.b8 95
.b8 99
.b8 117
.b8 98
.b8 51
.b8 116
.b8 97
.b8 103
.b8 69
.b8 78
.b8 83
.b8 48
.b8 95
.b8 50
.b8 55
.b8 114
.b8 97
.b8 110
.b8 100
.b8 111
.b8 109
.b8 95
.b8 97
.b8 99
.b8 99
.b8 101
.b8 115
.b8 115
.b8 95
.b8 116
.b8 114
.b8 97
.b8 118
.b8 101
.b8 114
.b8 115
.b8 97
.b8 108
.b8 95
.b8 116
.b8 97
.b8 103
.b8 69
.b8 78
.b8 83
.b8 48
.b8 95
.b8 49
.b8 54
.b8 100
.b8 101
.b8 118
.b8 105
.b8 99
.b8 101
.b8 95
.b8 114
.b8 101
.b8 102
.b8 101
.b8 114
.b8 101
.b8 110
.b8 99
.b8 101
.b8 73
.b8 105
.b8 69
.b8 69
.b8 108
.b8 69
.b8 69
.b8 84
.b8 95
.b8 82
.b8 75
.b8 78
.b8 83
.b8 48
.b8 95
.b8 49
.b8 53
.b8 105
.b8 116
.b8 101
.b8 114
.b8 97
.b8 116
.b8 111
.b8 114
.b8 95
.b8 102
.b8 97
.b8 99
.b8 97
.b8 100
.b8 101
.b8 73
.b8 83
.b8 57
.b8 95
.b8 84
.b8 48
.b8 95
.b8 84
.b8 49
.b8 95
.b8 84
.b8 50
.b8 95
.b8 84
.b8 51
.b8 95
.b8 84
.b8 52
.b8 95
.b8 69
.b8 69
.b8 78
.b8 83
.b8 57
.b8 95
.b8 49
.b8 53
.b8 100
.b8 105
.b8 102
.b8 102
.b8 101
.b8 114
.b8 101
.b8 110
.b8 99
.b8 101
.b8 95
.b8 116
.b8 121
.b8 112
.b8 101
.b8 69
.b8 0
$L__info_string11:
.b8 95
.b8 90
.b8 78
.b8 54
.b8 116
.b8 104
.b8 114
.b8 117
.b8 115
.b8 116
.b8 50
.b8 49
.b8 84
.b8 72
.b8 82
.b8 85
.b8 83
.b8 84
.b8 95
.b8 50
.b8 48
.b8 48
.b8 55
.b8 48
.b8 48
.b8 95
.b8 49
.b8 50
.b8 48
.b8 48
.b8 95
.b8 78
.b8 83
.b8 49
.b8 53
.b8 105
.b8 116
.b8 101
.b8 114
.b8 97
.b8 116
.b8 111
.b8 114
.b8 95
.b8 102
.b8 97
.b8 99
.b8 97
.b8 100
.b8 101
.b8 73
.b8 78
.b8 83
.b8 48
.b8 95
.b8 49
.b8 48
.b8 100
.b8 101
.b8 118
.b8 105
.b8 99
.b8 101
.b8 95
.b8 112
.b8 116
.b8 114
.b8 73
.b8 105
.b8 69
.b8 69
.b8 105
.b8 78
.b8 83
.b8 48
.b8 95
.b8 56
.b8 99
.b8 117
.b8 100
.b8 97
.b8 95
.b8 99
.b8 117
.b8 98
.b8 51
.b8 116
.b8 97
.b8 103
.b8 69
.b8 78
.b8 83
.b8 48
.b8 95
.b8 50
.b8 55
.b8 114
.b8 97
.b8 110
.b8 100
.b8 111
.b8 109
.b8 95
.b8 97
.b8 99
.b8 99
.b8 101
.b8 115
.b8 115
.b8 95
.b8 116
.b8 114
.b8 97
.b8 118
.b8 101
.b8 114
.b8 115
.b8 97
.b8 108
.b8 95
.b8 116
.b8 97
.b8 103
.b8 69
.b8 78
.b8 83
.b8 48
.b8 95
.b8 49
.b8 54
.b8 100
.b8 101
.b8 118
.b8 105
.b8 99
.b8 101
.b8 95
.b8 114
.b8 101
.b8 102
.b8 101
.b8 114
.b8 101
.b8 110
.b8 99
.b8 101
.b8 73
.b8 105
.b8 69
.b8 69
.b8 108
.b8 69
.b8 112
.b8 76
.b8 69
.b8 108
.b8 0
$L__info_string12:
.b8 95
.b8 90
.b8 78
.b8 54
.b8 116
.b8 104
.b8 114
.b8 117
.b8 115
.b8 116
.b8 50
.b8 49
.b8 84
.b8 72
.b8 82
.b8 85
.b8 83
.b8 84
.b8 95
.b8 50
.b8 48
.b8 48
.b8 55
.b8 48
.b8 48
.b8 95
.b8 49
.b8 50
.b8 48
.b8 48
.b8 95
.b8 78
.b8 83
.b8 50
.b8 48
.b8 105
.b8 116
.b8 101
.b8 114
.b8 97
.b8 116
.b8 111
.b8 114
.b8 95
.b8 99
.b8 111
.b8 114
.b8 101
.b8 95
.b8 97
.b8 99
.b8 99
.b8 101
.b8 115
.b8 115
.b8 55
.b8 97
.b8 100
.b8 118
.b8 97
.b8 110
.b8 99
.b8 101
.b8 73
.b8 78
.b8 83
.b8 48
.b8 95
.b8 49
.b8 48
.b8 100
.b8 101
.b8 118
.b8 105
.b8 99
.b8 101
.b8 95
.b8 112
.b8 116
.b8 114
.b8 73
.b8 105
.b8 69
.b8 69
.b8 69
.b8 69
.b8 118
.b8 82
.b8 84
.b8 95
.b8 78
.b8 83
.b8 53
.b8 95
.b8 49
.b8 53
.b8 100
.b8 105
.b8 102
.b8 102
.b8 101
.b8 114
.b8 101
.b8 110
.b8 99
.b8 101
.b8 95
.b8 116
.b8 121
.b8 112
.b8 101
.b8 69
.b8 0
$L__info_string13:
.b8 95
.b8 90
.b8 78
.b8 54
.b8 116
.b8 104
.b8 114
.b8 117
.b8 115
.b8 116
.b8 50
.b8 49
.b8 84
.b8 72
.b8 82
.b8 85
.b8 83
.b8 84
.b8 95
.b8 50
.b8 48
.b8 48
.b8 55
.b8 48
.b8 48
.b8 95
.b8 49
.b8 50
.b8 48
.b8 48
.b8 95
.b8 78
.b8 83
.b8 49
.b8 54
.b8 105
.b8 116
.b8 101
.b8 114
.b8 97
.b8 116
.b8 111
.b8 114
.b8 95
.b8 97
.b8 100
.b8 97
.b8 112
.b8 116
.b8 111
.b8 114
.b8 73
.b8 78
.b8 83
.b8 48
.b8 95
.b8 49
.b8 48
.b8 100
.b8 101
.b8 118
.b8 105
.b8 99
.b8 101
.b8 95
.b8 112
.b8 116
.b8 114
.b8 73
.b8 105
.b8 69
.b8 69
.b8 80
.b8 105
.b8 105
.b8 78
.b8 83
.b8 48
.b8 95
.b8 56
.b8 99
.b8 117
.b8 100
.b8 97
.b8 95
.b8 99
.b8 117
.b8 98
.b8 51
.b8 116
.b8 97
.b8 103
.b8 69
.b8 78
.b8 83
.b8 48
.b8 95
.b8 50
.b8 55
.b8 114
.b8 97
.b8 110
.b8 100
.b8 111
.b8 109
.b8 95
.b8 97
.b8 99
.b8 99
.b8 101
.b8 115
.b8 115
.b8 95
.b8 116
.b8 114
.b8 97
.b8 118
.b8 101
.b8 114
.b8 115
.b8 97
.b8 108
.b8 95
.b8 116
.b8 97
.b8 103
.b8 69
.b8 78
.b8 83
.b8 48
.b8 95
.b8 49
.b8 54
.b8 100
.b8 101
.b8 118
.b8 105
.b8 99
.b8 101
.b8 95
.b8 114
.b8 101
.b8 102
.b8 101
.b8 114
.b8 101
.b8 110
.b8 99
.b8 101
.b8 73
.b8 105
.b8 69
.b8 69
.b8 108
.b8 69
.b8 55
.b8 97
.b8 100
.b8 118
.b8 97
.b8 110
.b8 99
.b8 101
.b8 69
.b8 108
.b8 0
$L__info_string14:
.b8 95
.b8 90
.b8 78
.b8 51
.b8 99
.b8 117
.b8 98
.b8 49
.b8 56
.b8 67
.b8 85
.b8 66
.b8 95
.b8 50
.b8 48
.b8 48
.b8 55
.b8 48
.b8 48
.b8 95
.b8 49
.b8 50
.b8 48
.b8 48
.b8 95
.b8 78
.b8 83
.b8 54
.b8 100
.b8 101
.b8 116
.b8 97
.b8 105
.b8 108
.b8 56
.b8 102
.b8 111
.b8 114
.b8 95
.b8 101
.b8 97
.b8 99
.b8 104
.b8 50
.b8 49
.b8 97
.b8 103
.b8 101
.b8 110
.b8 116
.b8 95
.b8 98
.b8 108
.b8 111
.b8 99
.b8 107
.b8 95
.b8 115
.b8 116
.b8 114
.b8 105
.b8 112
.b8 101
.b8 100
.b8 95
.b8 116
.b8 73
.b8 78
.b8 83
.b8 50
.b8 95
.b8 56
.b8 112
.b8 111
.b8 108
.b8 105
.b8 99
.b8 121
.b8 95
.b8 116
.b8 73
.b8 76
.b8 105
.b8 50
.b8 53
.b8 54
.b8 69
.b8 76
.b8 105
.b8 50
.b8 69
.b8 69
.b8 69
.b8 109
.b8 78
.b8 54
.b8 116
.b8 104
.b8 114
.b8 117
.b8 115
.b8 116
.b8 50
.b8 49
.b8 84
.b8 72
.b8 82
.b8 85
.b8 83
.b8 84
.b8 95
.b8 50
.b8 48
.b8 48
.b8 55
.b8 48
.b8 48
.b8 95
.b8 49
.b8 50
.b8 48
.b8 48
.b8 95
.b8 78
.b8 83
.b8 56
.b8 99
.b8 117
.b8 100
.b8 97
.b8 95
.b8 99
.b8 117
.b8 98
.b8 50
.b8 48
.b8 95
.b8 95
.b8 117
.b8 110
.b8 105
.b8 110
.b8 105
.b8 116
.b8 105
.b8 97
.b8 108
.b8 105
.b8 122
.b8 101
.b8 100
.b8 95
.b8 102
.b8 105
.b8 108
.b8 108
.b8 55
.b8 102
.b8 117
.b8 110
.b8 99
.b8 116
.b8 111
.b8 114
.b8 73
.b8 78
.b8 83
.b8 55
.b8 95
.b8 49
.b8 48
.b8 100
.b8 101
.b8 118
.b8 105
.b8 99
.b8 101
.b8 95
.b8 112
.b8 116
.b8 114
.b8 73
.b8 105
.b8 69
.b8 69
.b8 105
.b8 69
.b8 69
.b8 69
.b8 49
.b8 50
.b8 99
.b8 111
.b8 110
.b8 115
.b8 117
.b8 109
.b8 101
.b8 95
.b8 116
.b8 105
.b8 108
.b8 101
.b8 73
.b8 76
.b8 98
.b8 48
.b8 69
.b8 69
.b8 69
.b8 118
.b8 105
.b8 105
.b8 0}