// #include <ATen/ATen.h>

#include <cuda.h>
#include <cuda_runtime.h>
#include <cstdio>
#include <cmath>
#include <iostream>

namespace {

template <typename scalar_t>
__device__ __forceinline__ void single_mul(
    scalar_t x_re,
    scalar_t x_im,
    scalar_t y_re,
    scalar_t y_im,
    scalar_t* out_re,
    scalar_t* out_im) {

    scalar_t uavc = x_re * (y_re + y_im);
    *out_re += uavc - (x_re + x_im) * y_im;
    *out_im += (x_im - x_re) * y_re + uavc;
}

template <typename scalar_t>
__device__ __forceinline__ void single_add(
    scalar_t x_re,
    scalar_t x_im,
    scalar_t y_re,
    scalar_t y_im,
    scalar_t* out_re,
    scalar_t* out_im) {

    *out_re += x_re + y_re;
    *out_im += x_im + y_im;
}

/**
Complex multiplication of tensors using shared memory and barrier
synchronization.

Compute the element wise complex multiplication for each thread in the block and
write the result to the shared memory. Then synchronize the threads and in the
log based fashion sum up the results for each output pixel through its channels,
if they are present in the cache. The stride is the number of threads per block
times the I (the two float representation of the complex numbers).
*/
template <typename scalar_t>
__global__ void complex_mul_cuda_kernel(
    const scalar_t* __restrict__ x,
    const scalar_t* __restrict__ y,
    scalar_t* __restrict__ out,
    const int N, const int F, const int C, const int H, const int W) {
    // The size of the shared memory cache should be twice the number of threads
    // per block as we store the real and imaginary part of the result.
    extern __shared__ float cache[];   // cache for the result of the complex multiplication

    const int I = 2; // the last dimension for the complex number
    const int plane_size = H * W;
    const int channel_size = plane_size * I;
    const int image_size = C * channel_size;  // size of the image from the batch
    // number of complex values in the input that we iterate through
    const int nr_values = C * H * W;

    const int n = blockIdx.x; // current index of an image/input map in the batch
    const int f = blockIdx.y; // current index of a filter from the filter bank
    const int block_size = blockDim.x;
    const int thread_nr = threadIdx.x;

    // stride for the H*W map is equal to the number of threads declared in a block
    const int stride = block_size * I; // we need H*W threads per plane, each deals with I numbers

    const int n_idx = n * image_size;  // start index in the batch for this input map
    const int f_idx = f * image_size;  // start index in the bank for this filter

    // find index for the output
    const int no_idx = n * (F * channel_size); // output index for the batch data point
    const int fo_idx = f * channel_size;       // output index for the filter/channel

    // Each H*W plane contains H*W*I elements in depth.
    // We linearize it and start from 0, move by #threads*I steps in outer loop.
    const int start_idx = threadIdx.x*I;

    // index in the input map
    int N_idx = n_idx + start_idx; // index across the first channel plane (in the input map n).
    const int last_N_idx = n_idx + image_size;  // last index for the starting position to compute the sum through each channel for this pixel

    // To prevent us from a deadlock, we have to always execute __syncthreads();
    // for all the threads in the block. Each thread has to do the same number of
    // iterations for any loop. To ensure that, we keep all threads running,
    // even though, some of them are really idle. We keep the loop running to
    // the multiple of the block size that is greater than the number of values
    // in the input map in total: C*H*W - this is a number of complex cells in the
    // input map.
    const int num_blocks = (nr_values + block_size - 1) / block_size;
    const int last_block_idx = n_idx + num_blocks * block_size * I;

    // index in the filter
    int F_idx = f_idx + start_idx;

    // index in the output, we compute cells on a flat plane (no channels)
    int base_O_idx = no_idx + fo_idx;
    int run_O_idx = (start_idx % channel_size);

    int thread_cidx = thread_nr * I;

    printf("N_idx:%d, last_block_idx:%d, last_N_idx:%d\n", N_idx, last_block_idx, last_N_idx);

    while (N_idx < last_block_idx)  {

        // Zero out caches.
        cache[thread_cidx] = 0;
        cache[thread_cidx + 1] = 0;

        if (N_idx < last_N_idx - 1) {
            scalar_t out_re = 0;
            scalar_t out_im = 0;

            scalar_t x_re = x[N_idx];
            scalar_t x_im = x[N_idx + 1];
            scalar_t y_re = y[F_idx];
            scalar_t y_im = y[F_idx + 1];
            single_mul(x_re, x_im, y_re, y_im, &out_re, &out_im);

            cache[thread_cidx] = out_re;
            cache[thread_cidx + 1] = out_im;
        }

        __syncthreads();  // Make the results visible to all threads.

        // It is not O(logN) but O(N) as of now. For each element in the output
        // map we have a dedicated thread. The thread goes through all the
        // channels present in the cache.
        if (thread_nr < plane_size) {
            for (int cache_index = thread_nr + plane_size;
                     cache_index < block_size;
                     cache_index += plane_size) {
                cache[thread_cidx] += cache[cache_index*I];
                cache[thread_cidx + 1] += cache[cache_index*I + 1];
            }
            // Move the summed values (across the channels) for each pixel to
            // the output.
            const int O_idx = base_O_idx + run_O_idx;
            out[O_idx] += cache[thread_cidx];
            out[O_idx + 1] = cache[thread_cidx + 1];
        }

        N_idx += stride;
        F_idx += stride;
        run_O_idx = (run_O_idx + stride) % channel_size;

        // Make sure that all cache cells are zeroed out before moving on.
        // We need this as in the second part we access cache cells that do not
        // belong only to this thread.
        __syncthreads();
    }
}


} // namespace

//void complex_mul_stride_no_permute_cuda(
//    at::Tensor x,
//    at::Tensor y,
//    at::Tensor out,
//    int threads = 1024) {
//
//    const auto N = x.size(0);  // batch_size
//    const auto F = y.size(0);  // filter_bank_size
//    const auto C = x.size(1);  // number of channels
//    const auto H = x.size(2);  // height of the matrix
//    const auto W = x.size(3);  // width of the matrix
//
//    const auto x_blocks = N;
//    const auto y_blocks = F;
//    const dim3 blocks(x_blocks, y_blocks);
//
//    AT_DISPATCH_FLOATING_TYPES(x.type(), "complex_mul_cuda",
//    ([&] {
//        complex_mul_cuda_kernel<scalar_t><<<blocks, threads>>>(
//        x.data<scalar_t>(), y.data<scalar_t>(), out.data<scalar_t>(),
//        N, F, C, H, W);
//    }));
//}

//template <typename scalar_t>
//void complex_mul_stride_no_permute_cuda_pure(
//    at::Tensor x,
//    at::Tensor y,
//    at::Tensor out,
//    int threads = 1024) {
//
//    const auto N = x.size(0);  // batch_size
//    const auto F = y.size(0);  // filter_bank_size
//    const auto C = x.size(1);  // number of channels
//    const auto H = x.size(2);  // height of the matrix
//    const auto W = x.size(3);  // width of the matrix
//
//    const auto x_blocks = N;
//    const auto y_blocks = F;
//    const dim3 blocks(x_blocks, y_blocks);
//
//    // Run kernel on the GPU
//    complex_mul_cuda_kernel<scalar_t><<<blocks, 1024>>>(
//        x.data<scalar_t>(), y.data<scalar_t>(), out.data<scalar_t>(),
//        N, F, C, H, W);
//}

/**
Uncomment the pytorch related stuff.

Compile:
ady@skr-compute1:/tmp/pycharm_project_154/cnns/nnlib/pytorch_cuda/complex_mul_cuda$ nvcc complex_mul_kernel_stride_no_permute.cu -o complex_mul_profile.out
ady@skr-compute1:/tmp/pycharm_project_154/cnns/nnlib/pytorch_cuda/complex_mul_cuda$ nvprof ./complex_mul_profile.out

nvidia

/usr/local/cuda/bin/nvcc -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include/torch/csrc/api/include -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include/TH -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include/THC -I/usr/local/cuda/include -I/local/ady/anaconda3/include/python3.6m -c complex_mul_kernel.cu -o complex_mul_kernel_stride_no_permute.out -std=c++11
nvcc -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include/torch/csrc/api/include -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include/TH -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include/THC -I/usr/local/cuda/include -I/local/ady/anaconda3/include/python3.6m complex_mul_kernel_stride_no_permute.cu -o complex_mul_kernel_stride_no_permute.out -std=c++11
Segmentation fault
*/
int main(void)
{
    int N = 1;
    int F = 1;
    int C = 4;
    int H = 3;
    int W = 2;
    int I = 2;
    int size_input = N * C * H * W * I;
    int size_filter = F * C * H * W * I;
    int size_output = N * F * H * W * I;
    int cuda_block_threads = 16;

    // auto dims = {128, 32, 16, 8, 2};
    //    at::Tensor x = at::randn({128, 32, 16, 8, 2});
    //    at::Tensor y = at::randn({128, 32, 16, 8, 2});
    //    at::Tensor out = at::zeros({128, 32, 16, 8, 2});
    float *x, *y, * out;

    // Allocate unified memory - accessible from cpu or gpu
    cudaMallocManaged(&x, size_input*sizeof(float));
    cudaMallocManaged(&y, size_filter*sizeof(float));
    cudaMallocManaged(&out, size_output*sizeof(float));

    for (int j=0; j<H; ++j) {
        for (int i=0; i<W; ++i) {
            x[(j*W+i)*2] = 3;
            x[(j*W+i)*2 + 1] = 1;
            y[(j*W+i)*2] = 4;
            y[(j*W+i)*2 + 1] = 2;
        }
    }

    for (int i=0; i<H*W*2; i+=2) {
        printf("%p %d: %f, %f, %f, %f\n", x, i, x[i], x[i+1], y[i], y[i+1]);
    }

//    float *dz;  // device z
//    cudaMalloc(&dz, 9*sizeof(float));
//    cudaMemcpy(dz, hz, 9*sizeof(float), cudaMemcpyHostToDevice);

    const dim3 blocks(N, F);

    complex_mul_cuda_kernel<float><<<blocks, cuda_block_threads,
        cuda_block_threads*2>>>(x, y, out, N, F, C, H, W);

    for (int i=0; i<H*W*C; i+=2) {
        printf("%d: %f, %f\n", i, out[i], out[i+1]);
    }

    cudaFree(x);
    cudaFree(y);
    cudaFree(out);
//    cudaFree(dz);

    printf("finished computation\n");

    return 0;
}
	.file	"tmpxft_00384533_00000000-6_complex_mul_shared_mem.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3642:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3642:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_ZL119__device_stub__ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiiiPKfS0_Pfiiiii, @function
_ZL119__device_stub__ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiiiPKfS0_Pfiiiii:
.LFB3664:
	.cfi_startproc
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movl	%ecx, 12(%rsp)
	leaq	48(%rsp), %rcx
	movl	%r8d, 8(%rsp)
	movl	%r9d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	16(%rsp), %rax
	movq	%rdi, 16(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	192(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	200(%rsp), %rax
	movq	%rsi, 24(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 32(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movl	$1, 64(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	48(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 208
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L2:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3664:
	.size	_ZL119__device_stub__ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiiiPKfS0_Pfiiiii, .-_ZL119__device_stub__ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiiiPKfS0_Pfiiiii
	.section	.text._ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii,"axG",@progbits,_ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii,comdat
	.weak	_ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
	.hidden	_ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
	.type	_ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii, @function
_ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii:
.LFB3968:
	.cfi_startproc
	endbr64
	jmp	_ZL119__device_stub__ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiiiPKfS0_Pfiiiii
	.cfi_endproc
.LFE3968:
	.size	_ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii, .-_ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC4:
	.string	"%p %d: %f, %f, %f, %f\n"
.LC5:
	.string	"%d: %f, %f\n"
.LC6:
	.string	"finished computation\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3639:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	$1, %edx
	movl	$192, %esi
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$64, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	call	cudaMallocManaged@PLT
	leaq	16(%rsp), %rdi
	movl	$1, %edx
	movl	$192, %esi
	call	cudaMallocManaged@PLT
	movl	$1, %edx
	leaq	24(%rsp), %rdi
	movl	$48, %esi
	call	cudaMallocManaged@PLT
	movq	8(%rsp), %rcx
	xorl	%eax, %eax
	movq	16(%rsp), %rdx
	movss	.LC0(%rip), %xmm3
	movss	.LC1(%rip), %xmm2
	movss	.LC2(%rip), %xmm1
	movss	.LC3(%rip), %xmm0
.L9:
	movss	%xmm3, (%rcx,%rax)
	movss	%xmm2, 4(%rcx,%rax)
	movss	%xmm1, (%rdx,%rax)
	movss	%xmm0, 4(%rdx,%rax)
	movss	%xmm3, 8(%rcx,%rax)
	movss	%xmm2, 12(%rcx,%rax)
	movss	%xmm1, 8(%rdx,%rax)
	movss	%xmm0, 12(%rdx,%rax)
	addq	$16, %rax
	cmpq	$48, %rax
	jne	.L9
	movl	$4, %ebx
	xorl	%ebp, %ebp
	leaq	.LC4(%rip), %r12
.L10:
	movq	16(%rsp), %rax
	movq	8(%rsp), %rdx
	movl	%ebp, %ecx
	movq	%r12, %rsi
	movl	$2, %edi
	addl	$2, %ebp
	cvtss2sd	(%rax,%rbx), %xmm3
	cvtss2sd	-4(%rax,%rbx), %xmm2
	movb	$4, %al
	cvtss2sd	-4(%rdx,%rbx), %xmm0
	cvtss2sd	(%rdx,%rbx), %xmm1
	call	__printf_chk@PLT
	addq	$8, %rbx
	cmpl	$12, %ebp
	jne	.L10
	movl	$268435457, %edx
	xorl	%r9d, %r9d
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	salq	$4, %rdx
	movl	$32, %r8d
	movl	$1, %esi
	movq	%rdi, 32(%rsp)
	movl	$1, 40(%rsp)
	movq	%rdx, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L11
	pushq	$2
	.cfi_def_cfa_offset 104
	movl	$4, %r9d
	movl	$1, %r8d
	movl	$1, %ecx
	pushq	$3
	.cfi_def_cfa_offset 112
	movq	40(%rsp), %rdx
	movq	32(%rsp), %rsi
	movq	24(%rsp), %rdi
	call	_ZL119__device_stub__ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiiiPKfS0_Pfiiiii
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L11:
	movl	$4, %ebp
	xorl	%ebx, %ebx
	leaq	.LC5(%rip), %r12
.L12:
	movq	24(%rsp), %rax
	movl	%ebx, %edx
	movq	%r12, %rsi
	movl	$2, %edi
	addl	$2, %ebx
	cvtss2sd	-4(%rax,%rbp), %xmm0
	cvtss2sd	(%rax,%rbp), %xmm1
	movb	$2, %al
	addq	$8, %rbp
	call	__printf_chk@PLT
	cmpl	$24, %ebx
	jne	.L12
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	xorl	%eax, %eax
	movl	$2, %edi
	leaq	.LC6(%rip), %rsi
	call	__printf_chk@PLT
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L13
	call	__stack_chk_fail@PLT
.L13:
	addq	$64, %rsp
	.cfi_def_cfa_offset 32
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3639:
	.size	main, .-main
	.section	.rodata.str1.1
.LC7:
	.string	"_ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3667:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC7(%rip), %rdx
	movq	%rax, %rdi
	leaq	_ZN54_GLOBAL__N__d33534c5_25_complex_mul_shared_mem_cu_main23complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3667:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC0:
	.long	1077936128
	.align 4
.LC1:
	.long	1065353216
	.align 4
.LC2:
	.long	1082130432
	.align 4
.LC3:
	.long	1073741824
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

// #include <ATen/ATen.h>

#include <hip/hip_runtime.h>

#include <cstdio>
#include <cmath>
#include <iostream>

namespace {

template <typename scalar_t>
__device__ __forceinline__ void single_mul(
    scalar_t x_re,
    scalar_t x_im,
    scalar_t y_re,
    scalar_t y_im,
    scalar_t* out_re,
    scalar_t* out_im) {

    scalar_t uavc = x_re * (y_re + y_im);
    *out_re += uavc - (x_re + x_im) * y_im;
    *out_im += (x_im - x_re) * y_re + uavc;
}

template <typename scalar_t>
__device__ __forceinline__ void single_add(
    scalar_t x_re,
    scalar_t x_im,
    scalar_t y_re,
    scalar_t y_im,
    scalar_t* out_re,
    scalar_t* out_im) {

    *out_re += x_re + y_re;
    *out_im += x_im + y_im;
}

/**
Complex multiplication of tensors using shared memory and barrier
synchronization.

Compute the element wise complex multiplication for each thread in the block and
write the result to the shared memory. Then synchronize the threads and in the
log based fashion sum up the results for each output pixel through its channels,
if they are present in the cache. The stride is the number of threads per block
times the I (the two float representation of the complex numbers).
*/
template <typename scalar_t>
__global__ void complex_mul_cuda_kernel(
    const scalar_t* __restrict__ x,
    const scalar_t* __restrict__ y,
    scalar_t* __restrict__ out,
    const int N, const int F, const int C, const int H, const int W) {
    // The size of the shared memory cache should be twice the number of threads
    // per block as we store the real and imaginary part of the result.
    extern __shared__ float cache[];   // cache for the result of the complex multiplication

    const int I = 2; // the last dimension for the complex number
    const int plane_size = H * W;
    const int channel_size = plane_size * I;
    const int image_size = C * channel_size;  // size of the image from the batch
    // number of complex values in the input that we iterate through
    const int nr_values = C * H * W;

    const int n = blockIdx.x; // current index of an image/input map in the batch
    const int f = blockIdx.y; // current index of a filter from the filter bank
    const int block_size = blockDim.x;
    const int thread_nr = threadIdx.x;

    // stride for the H*W map is equal to the number of threads declared in a block
    const int stride = block_size * I; // we need H*W threads per plane, each deals with I numbers

    const int n_idx = n * image_size;  // start index in the batch for this input map
    const int f_idx = f * image_size;  // start index in the bank for this filter

    // find index for the output
    const int no_idx = n * (F * channel_size); // output index for the batch data point
    const int fo_idx = f * channel_size;       // output index for the filter/channel

    // Each H*W plane contains H*W*I elements in depth.
    // We linearize it and start from 0, move by #threads*I steps in outer loop.
    const int start_idx = threadIdx.x*I;

    // index in the input map
    int N_idx = n_idx + start_idx; // index across the first channel plane (in the input map n).
    const int last_N_idx = n_idx + image_size;  // last index for the starting position to compute the sum through each channel for this pixel

    // To prevent us from a deadlock, we have to always execute __syncthreads();
    // for all the threads in the block. Each thread has to do the same number of
    // iterations for any loop. To ensure that, we keep all threads running,
    // even though, some of them are really idle. We keep the loop running to
    // the multiple of the block size that is greater than the number of values
    // in the input map in total: C*H*W - this is a number of complex cells in the
    // input map.
    const int num_blocks = (nr_values + block_size - 1) / block_size;
    const int last_block_idx = n_idx + num_blocks * block_size * I;

    // index in the filter
    int F_idx = f_idx + start_idx;

    // index in the output, we compute cells on a flat plane (no channels)
    int base_O_idx = no_idx + fo_idx;
    int run_O_idx = (start_idx % channel_size);

    int thread_cidx = thread_nr * I;

    printf("N_idx:%d, last_block_idx:%d, last_N_idx:%d\n", N_idx, last_block_idx, last_N_idx);

    while (N_idx < last_block_idx)  {

        // Zero out caches.
        cache[thread_cidx] = 0;
        cache[thread_cidx + 1] = 0;

        if (N_idx < last_N_idx - 1) {
            scalar_t out_re = 0;
            scalar_t out_im = 0;

            scalar_t x_re = x[N_idx];
            scalar_t x_im = x[N_idx + 1];
            scalar_t y_re = y[F_idx];
            scalar_t y_im = y[F_idx + 1];
            single_mul(x_re, x_im, y_re, y_im, &out_re, &out_im);

            cache[thread_cidx] = out_re;
            cache[thread_cidx + 1] = out_im;
        }

        __syncthreads();  // Make the results visible to all threads.

        // It is not O(logN) but O(N) as of now. For each element in the output
        // map we have a dedicated thread. The thread goes through all the
        // channels present in the cache.
        if (thread_nr < plane_size) {
            for (int cache_index = thread_nr + plane_size;
                     cache_index < block_size;
                     cache_index += plane_size) {
                cache[thread_cidx] += cache[cache_index*I];
                cache[thread_cidx + 1] += cache[cache_index*I + 1];
            }
            // Move the summed values (across the channels) for each pixel to
            // the output.
            const int O_idx = base_O_idx + run_O_idx;
            out[O_idx] += cache[thread_cidx];
            out[O_idx + 1] = cache[thread_cidx + 1];
        }

        N_idx += stride;
        F_idx += stride;
        run_O_idx = (run_O_idx + stride) % channel_size;

        // Make sure that all cache cells are zeroed out before moving on.
        // We need this as in the second part we access cache cells that do not
        // belong only to this thread.
        __syncthreads();
    }
}


} // namespace

//void complex_mul_stride_no_permute_cuda(
//    at::Tensor x,
//    at::Tensor y,
//    at::Tensor out,
//    int threads = 1024) {
//
//    const auto N = x.size(0);  // batch_size
//    const auto F = y.size(0);  // filter_bank_size
//    const auto C = x.size(1);  // number of channels
//    const auto H = x.size(2);  // height of the matrix
//    const auto W = x.size(3);  // width of the matrix
//
//    const auto x_blocks = N;
//    const auto y_blocks = F;
//    const dim3 blocks(x_blocks, y_blocks);
//
//    AT_DISPATCH_FLOATING_TYPES(x.type(), "complex_mul_cuda",
//    ([&] {
//        complex_mul_cuda_kernel<scalar_t><<<blocks, threads>>>(
//        x.data<scalar_t>(), y.data<scalar_t>(), out.data<scalar_t>(),
//        N, F, C, H, W);
//    }));
//}

//template <typename scalar_t>
//void complex_mul_stride_no_permute_cuda_pure(
//    at::Tensor x,
//    at::Tensor y,
//    at::Tensor out,
//    int threads = 1024) {
//
//    const auto N = x.size(0);  // batch_size
//    const auto F = y.size(0);  // filter_bank_size
//    const auto C = x.size(1);  // number of channels
//    const auto H = x.size(2);  // height of the matrix
//    const auto W = x.size(3);  // width of the matrix
//
//    const auto x_blocks = N;
//    const auto y_blocks = F;
//    const dim3 blocks(x_blocks, y_blocks);
//
//    // Run kernel on the GPU
//    complex_mul_cuda_kernel<scalar_t><<<blocks, 1024>>>(
//        x.data<scalar_t>(), y.data<scalar_t>(), out.data<scalar_t>(),
//        N, F, C, H, W);
//}

/**
Uncomment the pytorch related stuff.

Compile:
ady@skr-compute1:/tmp/pycharm_project_154/cnns/nnlib/pytorch_cuda/complex_mul_cuda$ nvcc complex_mul_kernel_stride_no_permute.cu -o complex_mul_profile.out
ady@skr-compute1:/tmp/pycharm_project_154/cnns/nnlib/pytorch_cuda/complex_mul_cuda$ nvprof ./complex_mul_profile.out

nvidia

/usr/local/cuda/bin/nvcc -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include/torch/csrc/api/include -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include/TH -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include/THC -I/usr/local/cuda/include -I/local/ady/anaconda3/include/python3.6m -c complex_mul_kernel.cu -o complex_mul_kernel_stride_no_permute.out -std=c++11
nvcc -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include/torch/csrc/api/include -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include/TH -I/local/ady/anaconda3/lib/python3.6/site-packages/torch/lib/include/THC -I/usr/local/cuda/include -I/local/ady/anaconda3/include/python3.6m complex_mul_kernel_stride_no_permute.cu -o complex_mul_kernel_stride_no_permute.out -std=c++11
Segmentation fault
*/
int main(void)
{
    int N = 1;
    int F = 1;
    int C = 4;
    int H = 3;
    int W = 2;
    int I = 2;
    int size_input = N * C * H * W * I;
    int size_filter = F * C * H * W * I;
    int size_output = N * F * H * W * I;
    int cuda_block_threads = 16;

    // auto dims = {128, 32, 16, 8, 2};
    //    at::Tensor x = at::randn({128, 32, 16, 8, 2});
    //    at::Tensor y = at::randn({128, 32, 16, 8, 2});
    //    at::Tensor out = at::zeros({128, 32, 16, 8, 2});
    float *x, *y, * out;

    // Allocate unified memory - accessible from cpu or gpu
    hipMallocManaged(&x, size_input*sizeof(float));
    hipMallocManaged(&y, size_filter*sizeof(float));
    hipMallocManaged(&out, size_output*sizeof(float));

    for (int j=0; j<H; ++j) {
        for (int i=0; i<W; ++i) {
            x[(j*W+i)*2] = 3;
            x[(j*W+i)*2 + 1] = 1;
            y[(j*W+i)*2] = 4;
            y[(j*W+i)*2 + 1] = 2;
        }
    }

    for (int i=0; i<H*W*2; i+=2) {
        printf("%p %d: %f, %f, %f, %f\n", x, i, x[i], x[i+1], y[i], y[i+1]);
    }

//    float *dz;  // device z
//    cudaMalloc(&dz, 9*sizeof(float));
//    cudaMemcpy(dz, hz, 9*sizeof(float), cudaMemcpyHostToDevice);

    const dim3 blocks(N, F);

    complex_mul_cuda_kernel<float><<<blocks, cuda_block_threads,
        cuda_block_threads*2>>>(x, y, out, N, F, C, H, W);

    for (int i=0; i<H*W*C; i+=2) {
        printf("%d: %f, %f\n", i, out[i], out[i+1]);
    }

    hipFree(x);
    hipFree(y);
    hipFree(out);
//    cudaFree(dz);

    printf("finished computation\n");

    return 0;
}
	.text
	.file	"complex_mul_shared_mem.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	leaq	16(%rsp), %rbx
	movl	$192, %esi
	movq	%rbx, %rdi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	8(%rsp), %r14
	movl	$192, %esi
	movq	%r14, %rdi
	movl	$1, %edx
	callq	hipMallocManaged
	movq	%rsp, %rdi
	movl	$48, %esi
	movl	$1, %edx
	callq	hipMallocManaged
	movq	(%rbx), %rax
	movl	$12, %ecx
	movq	(%r14), %rdx
	movabsq	$4575657222486360064, %rsi      # imm = 0x3F80000040400000
	movabsq	$4611686019509518336, %rdi      # imm = 0x4000000040800000
.LBB0_1:                                # %.preheader71
                                        # =>This Inner Loop Header: Depth=1
	movq	%rsi, -12(%rax,%rcx)
	movq	%rdi, -12(%rdx,%rcx)
	movq	%rsi, -4(%rax,%rcx)
	movq	%rdi, -4(%rdx,%rcx)
	addq	$16, %rcx
	cmpq	$60, %rcx
	jne	.LBB0_1
# %bb.2:                                # %.preheader.preheader
	xorl	%ebx, %ebx
.LBB0_3:                                # %.preheader
                                        # =>This Inner Loop Header: Depth=1
	movq	16(%rsp), %rsi
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rsi,%rbx,4), %xmm0
	xorps	%xmm1, %xmm1
	cvtss2sd	4(%rsi,%rbx,4), %xmm1
	movq	8(%rsp), %rax
	xorps	%xmm2, %xmm2
	cvtss2sd	(%rax,%rbx,4), %xmm2
	xorps	%xmm3, %xmm3
	cvtss2sd	4(%rax,%rbx,4), %xmm3
	movl	$.L.str, %edi
	movl	%ebx, %edx
	movb	$4, %al
	callq	printf
	leaq	2(%rbx), %rax
	cmpq	$10, %rbx
	movq	%rax, %rbx
	jb	.LBB0_3
# %bb.4:
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	15(%rdi), %rdx
	movl	$32, %r8d
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_5
# %bb.8:
	movq	16(%rsp), %rdi
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdx
	movl	$1, %ecx
	movl	$1, %r8d
	movl	$4, %r9d
	pushq	$2
	.cfi_adjust_cfa_offset 8
	pushq	$3
	.cfi_adjust_cfa_offset 8
	callq	_ZN12_GLOBAL__N_138__device_stub__complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
.LBB0_5:                                # %.preheader86
	xorl	%ebx, %ebx
.LBB0_6:                                # =>This Inner Loop Header: Depth=1
	movq	(%rsp), %rax
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rax,%rbx,4), %xmm0
	xorps	%xmm1, %xmm1
	cvtss2sd	4(%rax,%rbx,4), %xmm1
	movl	$.L.str.1, %edi
	movl	%ebx, %esi
	movb	$2, %al
	callq	printf
	leaq	2(%rbx), %rax
	cmpq	$22, %rbx
	movq	%rax, %rbx
	jb	.LBB0_6
# %bb.7:
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	movl	$.Lstr, %edi
	callq	puts@PLT
	xorl	%eax, %eax
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.type	_ZN12_GLOBAL__N_138__device_stub__complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii,@function # -- Begin function _ZN12_GLOBAL__N_138__device_stub__complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
_ZN12_GLOBAL__N_138__device_stub__complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii: # @_ZN12_GLOBAL__N_138__device_stub__complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$160, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	20(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	16(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	12(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	216(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_ZN12_GLOBAL__N_138__device_stub__complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii, .Lfunc_end1-_ZN12_GLOBAL__N_138__device_stub__complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"%p %d: %f, %f, %f, %f\n"
	.size	.L.str, 23

	.type	_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii,@object # @_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
	.section	.rodata,"a",@progbits
	.p2align	3, 0x0
_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii:
	.quad	_ZN12_GLOBAL__N_138__device_stub__complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
	.size	_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii, 8

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"%d: %f, %f\n"
	.size	.L.str.1, 12

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii"
	.size	.L__unnamed_1, 64

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"finished computation"
	.size	.Lstr, 21

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _ZN12_GLOBAL__N_138__device_stub__complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.section	.text._ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii,"axG",@progbits,_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii,comdat
	.globl	_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii ; -- Begin function _ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
	.p2align	8
	.type	_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii,@function
_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii: ; @_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
; %bb.0:
	s_clause 0x1
	s_load_b32 s24, s[0:1], 0x3c
	s_load_b64 s[16:17], s[0:1], 0x80
	v_mbcnt_lo_u32_b32 v3, -1, 0
	v_mov_b32_e32 v1, 0
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v8, v3
	;;#ASMSTART
	;;#ASMEND
	v_readfirstlane_b32 s2, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_eq_u32_e64 s2, s2, v8
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_6
; %bb.1:
	v_mov_b32_e32 v4, 0
	s_mov_b32 s4, exec_lo
	s_waitcnt lgkmcnt(0)
	global_load_b64 v[11:12], v4, s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_clause 0x1
	global_load_b64 v[1:2], v4, s[16:17] offset:40
	global_load_b64 v[5:6], v4, s[16:17]
	s_waitcnt vmcnt(1)
	v_and_b32_e32 v1, v1, v11
	v_and_b32_e32 v2, v2, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v7, v1, 24
	v_mul_lo_u32 v2, v2, 24
	v_mul_lo_u32 v1, v1, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v2, v7, v2
	s_waitcnt vmcnt(0)
	v_add_co_u32 v1, vcc_lo, v5, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, v6, v2, vcc_lo
	global_load_b64 v[9:10], v[1:2], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[1:2], v4, v[9:12], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmpx_ne_u64_e64 v[1:2], v[11:12]
	s_cbranch_execz .LBB0_5
; %bb.2:                                ; %.preheader3.i.i.i.preheader
	s_mov_b32 s5, 0
.LBB0_3:                                ; %.preheader3.i.i.i
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	s_clause 0x1
	global_load_b64 v[5:6], v4, s[16:17] offset:40
	global_load_b64 v[9:10], v4, s[16:17]
	v_dual_mov_b32 v12, v2 :: v_dual_mov_b32 v11, v1
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v5, v5, v11
	v_and_b32_e32 v7, v6, v12
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, v5, 24, v[9:10]
	v_mad_u64_u32 v[5:6], null, v7, 24, v[2:3]
	s_delay_alu instid0(VALU_DEP_1)
	v_mov_b32_e32 v2, v5
	global_load_b64 v[9:10], v[1:2], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[1:2], v4, v[9:12], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmp_eq_u64_e32 vcc_lo, v[1:2], v[11:12]
	s_or_b32 s5, vcc_lo, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s5
	s_cbranch_execnz .LBB0_3
; %bb.4:                                ; %Flow556
	s_or_b32 exec_lo, exec_lo, s5
.LBB0_5:                                ; %Flow558
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s4
.LBB0_6:                                ; %.loopexit4.i.i.i
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s3
	v_mov_b32_e32 v9, 0
	v_readfirstlane_b32 s8, v1
	v_readfirstlane_b32 s9, v2
	s_mov_b32 s6, exec_lo
	s_waitcnt lgkmcnt(0)
	s_clause 0x1
	global_load_b64 v[10:11], v9, s[16:17] offset:40
	global_load_b128 v[4:7], v9, s[16:17]
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s4, v10
	v_readfirstlane_b32 s5, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[4:5], s[8:9], s[4:5]
	s_mul_i32 s3, s5, 24
	s_mul_hi_u32 s10, s4, 24
	s_mul_i32 s11, s4, 24
	s_and_saveexec_b32 s7, s2
	s_cbranch_execz .LBB0_8
; %bb.7:
	v_dual_mov_b32 v10, s6 :: v_dual_mov_b32 v11, v9
	s_add_i32 s6, s10, s3
	s_waitcnt vmcnt(0)
	v_add_co_u32 v1, vcc_lo, v4, s11
	v_add_co_ci_u32_e32 v2, vcc_lo, s6, v5, vcc_lo
	v_dual_mov_b32 v12, 2 :: v_dual_mov_b32 v13, 1
	global_store_b128 v[1:2], v[10:13], off offset:8
.LBB0_8:
	s_or_b32 exec_lo, exec_lo, s7
	s_lshl_b64 s[4:5], s[4:5], 12
	v_lshlrev_b64 v[1:2], 6, v[8:9]
	s_waitcnt vmcnt(0)
	v_add_co_u32 v6, vcc_lo, v6, s4
	v_add_co_ci_u32_e32 v7, vcc_lo, s5, v7, vcc_lo
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[12:13], s[0:1], 0x10
	s_mov_b32 s20, 0
	v_add_co_u32 v1, vcc_lo, v6, v1
	s_mov_b32 s21, s20
	s_mov_b32 s22, s20
	s_mov_b32 s23, s20
	v_add_co_ci_u32_e32 v2, vcc_lo, v7, v2, vcc_lo
	v_dual_mov_b32 v8, 33 :: v_dual_mov_b32 v11, v9
	v_mov_b32_e32 v10, v9
	v_dual_mov_b32 v12, s20 :: v_dual_mov_b32 v15, s23
	v_dual_mov_b32 v13, s21 :: v_dual_mov_b32 v14, s22
	s_clause 0x3
	global_store_b128 v[1:2], v[8:11], off
	global_store_b128 v[1:2], v[12:15], off offset:16
	global_store_b128 v[1:2], v[12:15], off offset:32
	global_store_b128 v[1:2], v[12:15], off offset:48
	s_and_saveexec_b32 s18, s2
	s_cbranch_execz .LBB0_15
; %bb.9:
	v_mov_b32_e32 v12, 0
	s_mov_b32 s19, exec_lo
	s_clause 0x1
	global_load_b64 v[15:16], v12, s[16:17] offset:32 glc
	global_load_b64 v[6:7], v12, s[16:17] offset:40
	v_dual_mov_b32 v13, s8 :: v_dual_mov_b32 v14, s9
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v7, s9, v7
	v_and_b32_e32 v6, s8, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v7, v7, 24
	v_mul_hi_u32 v8, v6, 24
	v_mul_lo_u32 v6, v6, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v7, v8, v7
	v_add_co_u32 v10, vcc_lo, v4, v6
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v11, vcc_lo, v5, v7, vcc_lo
	global_store_b64 v[10:11], v[15:16], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[8:9], v12, v[13:16], s[16:17] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u64_e64 v[8:9], v[15:16]
	s_cbranch_execz .LBB0_11
.LBB0_10:                               ; %.preheader1.i.i.i
                                        ; =>This Inner Loop Header: Depth=1
	v_dual_mov_b32 v6, s8 :: v_dual_mov_b32 v7, s9
	s_sleep 1
	global_store_b64 v[10:11], v[8:9], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[6:7], v12, v[6:9], s[16:17] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[6:7], v[8:9]
	v_dual_mov_b32 v9, v7 :: v_dual_mov_b32 v8, v6
	s_or_b32 s20, vcc_lo, s20
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s20
	s_cbranch_execnz .LBB0_10
.LBB0_11:                               ; %Flow554
	s_or_b32 exec_lo, exec_lo, s19
	v_mov_b32_e32 v9, 0
	s_mov_b32 s20, exec_lo
	s_mov_b32 s19, exec_lo
	v_mbcnt_lo_u32_b32 v8, s20, 0
	global_load_b64 v[6:7], v9, s[16:17] offset:16
	v_cmpx_eq_u32_e32 0, v8
	s_cbranch_execz .LBB0_13
; %bb.12:
	s_bcnt1_i32_b32 s20, s20
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v8, s20
	s_waitcnt vmcnt(0)
	global_atomic_add_u64 v[6:7], v[8:9], off offset:8
.LBB0_13:
	s_or_b32 exec_lo, exec_lo, s19
	s_waitcnt vmcnt(0)
	global_load_b64 v[8:9], v[6:7], off offset:16
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[8:9]
	s_cbranch_vccnz .LBB0_15
; %bb.14:
	global_load_b32 v6, v[6:7], off offset:24
	v_mov_b32_e32 v7, 0
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s19, v6
	s_waitcnt_vscnt null, 0x0
	global_store_b64 v[8:9], v[6:7], off
	s_and_b32 m0, s19, 0xff
	s_sendmsg sendmsg(MSG_INTERRUPT)
.LBB0_15:                               ; %Flow555
	s_or_b32 exec_lo, exec_lo, s18
	s_add_i32 s10, s10, s3
	v_add_co_u32 v4, vcc_lo, v4, s11
	v_add_co_ci_u32_e32 v5, vcc_lo, s10, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, v4, 20
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
.LBB0_16:                               ; =>This Inner Loop Header: Depth=1
	v_mov_b32_e32 v6, 1
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_18
; %bb.17:                               ;   in Loop: Header=BB0_16 Depth=1
	global_load_b32 v6, v[4:5], off glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_and_b32_e32 v6, 1, v6
.LBB0_18:                               ;   in Loop: Header=BB0_16 Depth=1
	s_or_b32 exec_lo, exec_lo, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s3, v6
	s_cmp_eq_u32 s3, 0
	s_cbranch_scc1 .LBB0_20
; %bb.19:                               ;   in Loop: Header=BB0_16 Depth=1
	s_mov_b32 s3, 0
	s_sleep 1
	s_branch .LBB0_21
.LBB0_20:                               ;   in Loop: Header=BB0_16 Depth=1
	s_mov_b32 s3, -1
.LBB0_21:                               ; %Flow549
                                        ;   in Loop: Header=BB0_16 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s3
	s_cbranch_vccnz .LBB0_16
; %bb.22:
	global_load_b64 v[8:9], v[1:2], off
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_26
; %bb.23:
	v_mov_b32_e32 v10, 0
	s_clause 0x2
	global_load_b64 v[1:2], v10, s[16:17] offset:40
	global_load_b64 v[11:12], v10, s[16:17] offset:24 glc
	global_load_b64 v[6:7], v10, s[16:17]
	s_waitcnt vmcnt(2)
	v_add_co_u32 v13, vcc_lo, v1, 1
	v_add_co_ci_u32_e32 v14, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, v13, s8
	v_add_co_ci_u32_e32 v5, vcc_lo, s9, v14, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[4:5]
	v_dual_cndmask_b32 v5, v5, v14 :: v_dual_cndmask_b32 v4, v4, v13
	v_and_b32_e32 v2, v5, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v1, v4, v1
	v_mul_lo_u32 v2, v2, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v13, v1, 24
	v_mul_lo_u32 v1, v1, 24
	v_add_nc_u32_e32 v2, v13, v2
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_u32 v1, vcc_lo, v6, v1
	v_mov_b32_e32 v6, v11
	v_add_co_ci_u32_e32 v2, vcc_lo, v7, v2, vcc_lo
	v_mov_b32_e32 v7, v12
	global_store_b64 v[1:2], v[11:12], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[6:7], v10, v[4:7], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, v[6:7], v[11:12]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_26
; %bb.24:                               ; %.preheader.i.i.i.preheader
	s_mov_b32 s2, 0
.LBB0_25:                               ; %.preheader.i.i.i
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	global_store_b64 v[1:2], v[6:7], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[11:12], v10, v[4:7], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[11:12], v[6:7]
	v_dual_mov_b32 v6, v11 :: v_dual_mov_b32 v7, v12
	s_or_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB0_25
.LBB0_26:                               ; %__ockl_printf_begin.exit
	s_or_b32 exec_lo, exec_lo, s3
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, .str@rel32@lo+4
	s_addc_u32 s9, s9, .str@rel32@hi+12
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u64 s[8:9], 0
	s_cbranch_scc0 .LBB0_112
; %bb.27:
	s_waitcnt vmcnt(0)
	v_dual_mov_b32 v5, v9 :: v_dual_and_b32 v30, 2, v8
	v_dual_mov_b32 v2, 0 :: v_dual_mov_b32 v13, 1
	v_and_b32_e32 v4, -3, v8
	v_mov_b32_e32 v12, 2
	s_mov_b64 s[10:11], 44
.LBB0_28:                               ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_31 Depth 2
                                        ;     Child Loop BB0_38 Depth 2
                                        ;     Child Loop BB0_46 Depth 2
                                        ;     Child Loop BB0_54 Depth 2
                                        ;     Child Loop BB0_62 Depth 2
                                        ;     Child Loop BB0_70 Depth 2
                                        ;     Child Loop BB0_78 Depth 2
                                        ;     Child Loop BB0_86 Depth 2
                                        ;     Child Loop BB0_94 Depth 2
                                        ;     Child Loop BB0_100 Depth 2
                                        ;     Child Loop BB0_109 Depth 2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_u64_e64 s2, s[10:11], 56
                                        ; implicit-def: $vgpr6_vgpr7
                                        ; implicit-def: $sgpr26
	s_and_b32 s2, s2, exec_lo
	s_cselect_b32 s18, s10, 56
	s_cselect_b32 s19, s11, 0
	s_cmp_gt_u32 s18, 7
	s_mov_b32 s2, -1
	s_cbranch_scc1 .LBB0_33
; %bb.29:                               ;   in Loop: Header=BB0_28 Depth=1
	s_waitcnt vmcnt(0)
	v_mov_b32_e32 v6, 0
	v_mov_b32_e32 v7, 0
	s_cmp_eq_u32 s18, 0
	s_cbranch_scc1 .LBB0_32
; %bb.30:                               ; %.preheader31.i.preheader
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_lshl_b64 s[2:3], s[18:19], 3
	s_mov_b64 s[20:21], 0
	s_mov_b64 s[22:23], s[8:9]
.LBB0_31:                               ; %.preheader31.i
                                        ;   Parent Loop BB0_28 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_u8 v1, v2, s[22:23]
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v1, 0xffff, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_lshlrev_b64 v[10:11], s20, v[1:2]
	s_add_u32 s20, s20, 8
	s_addc_u32 s21, s21, 0
	s_add_u32 s22, s22, 1
	s_addc_u32 s23, s23, 0
	s_cmp_lg_u32 s2, s20
	v_or_b32_e32 v6, v10, v6
	v_or_b32_e32 v7, v11, v7
	s_cbranch_scc1 .LBB0_31
.LBB0_32:                               ; %Flow525
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s2, 0
	s_mov_b32 s26, 0
.LBB0_33:                               ; %Flow527
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_and_not1_b32 vcc_lo, exec_lo, s2
	s_mov_b64 s[2:3], s[8:9]
	s_cbranch_vccnz .LBB0_35
; %bb.34:                               ;   in Loop: Header=BB0_28 Depth=1
	global_load_b64 v[6:7], v2, s[8:9]
	s_add_i32 s26, s18, -8
	s_add_u32 s2, s8, 8
	s_addc_u32 s3, s9, 0
.LBB0_35:                               ; %.loopexit32.i
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_cmp_gt_u32 s26, 7
	s_cbranch_scc1 .LBB0_40
; %bb.36:                               ;   in Loop: Header=BB0_28 Depth=1
	v_mov_b32_e32 v14, 0
	v_mov_b32_e32 v15, 0
	s_cmp_eq_u32 s26, 0
	s_cbranch_scc1 .LBB0_39
; %bb.37:                               ; %.preheader29.i.preheader
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b64 s[20:21], 0
	s_mov_b64 s[22:23], 0
.LBB0_38:                               ; %.preheader29.i
                                        ;   Parent Loop BB0_28 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s28, s2, s22
	s_addc_u32 s29, s3, s23
	s_add_u32 s22, s22, 1
	global_load_u8 v1, v2, s[28:29]
	s_addc_u32 s23, s23, 0
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v1, 0xffff, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[10:11], s20, v[1:2]
	s_add_u32 s20, s20, 8
	s_addc_u32 s21, s21, 0
	s_cmp_lg_u32 s26, s22
	v_or_b32_e32 v14, v10, v14
	s_delay_alu instid0(VALU_DEP_2)
	v_or_b32_e32 v15, v11, v15
	s_cbranch_scc1 .LBB0_38
.LBB0_39:                               ; %Flow520
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s20, 0
	s_mov_b32 s25, 0
	s_branch .LBB0_41
.LBB0_40:                               ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s20, -1
                                        ; implicit-def: $vgpr14_vgpr15
                                        ; implicit-def: $sgpr25
.LBB0_41:                               ; %Flow522
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s20
	s_cbranch_vccnz .LBB0_43
; %bb.42:                               ;   in Loop: Header=BB0_28 Depth=1
	global_load_b64 v[14:15], v2, s[2:3]
	s_add_i32 s25, s26, -8
	s_add_u32 s2, s2, 8
	s_addc_u32 s3, s3, 0
.LBB0_43:                               ; %.loopexit30.i
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_cmp_gt_u32 s25, 7
	s_cbranch_scc1 .LBB0_48
; %bb.44:                               ;   in Loop: Header=BB0_28 Depth=1
	v_mov_b32_e32 v16, 0
	v_mov_b32_e32 v17, 0
	s_cmp_eq_u32 s25, 0
	s_cbranch_scc1 .LBB0_47
; %bb.45:                               ; %.preheader27.i.preheader
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b64 s[20:21], 0
	s_mov_b64 s[22:23], 0
.LBB0_46:                               ; %.preheader27.i
                                        ;   Parent Loop BB0_28 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s26, s2, s22
	s_addc_u32 s27, s3, s23
	s_add_u32 s22, s22, 1
	global_load_u8 v1, v2, s[26:27]
	s_addc_u32 s23, s23, 0
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v1, 0xffff, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[10:11], s20, v[1:2]
	s_add_u32 s20, s20, 8
	s_addc_u32 s21, s21, 0
	s_cmp_lg_u32 s25, s22
	v_or_b32_e32 v16, v10, v16
	s_delay_alu instid0(VALU_DEP_2)
	v_or_b32_e32 v17, v11, v17
	s_cbranch_scc1 .LBB0_46
.LBB0_47:                               ; %Flow515
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s20, 0
	s_mov_b32 s26, 0
	s_branch .LBB0_49
.LBB0_48:                               ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s20, -1
                                        ; implicit-def: $sgpr26
.LBB0_49:                               ; %Flow517
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s20
	s_cbranch_vccnz .LBB0_51
; %bb.50:                               ;   in Loop: Header=BB0_28 Depth=1
	global_load_b64 v[16:17], v2, s[2:3]
	s_add_i32 s26, s25, -8
	s_add_u32 s2, s2, 8
	s_addc_u32 s3, s3, 0
.LBB0_51:                               ; %.loopexit28.i
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_cmp_gt_u32 s26, 7
	s_cbranch_scc1 .LBB0_56
; %bb.52:                               ;   in Loop: Header=BB0_28 Depth=1
	v_mov_b32_e32 v18, 0
	v_mov_b32_e32 v19, 0
	s_cmp_eq_u32 s26, 0
	s_cbranch_scc1 .LBB0_55
; %bb.53:                               ; %.preheader25.i.preheader
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b64 s[20:21], 0
	s_mov_b64 s[22:23], 0
.LBB0_54:                               ; %.preheader25.i
                                        ;   Parent Loop BB0_28 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s28, s2, s22
	s_addc_u32 s29, s3, s23
	s_add_u32 s22, s22, 1
	global_load_u8 v1, v2, s[28:29]
	s_addc_u32 s23, s23, 0
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v1, 0xffff, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[10:11], s20, v[1:2]
	s_add_u32 s20, s20, 8
	s_addc_u32 s21, s21, 0
	s_cmp_lg_u32 s26, s22
	v_or_b32_e32 v18, v10, v18
	s_delay_alu instid0(VALU_DEP_2)
	v_or_b32_e32 v19, v11, v19
	s_cbranch_scc1 .LBB0_54
.LBB0_55:                               ; %Flow510
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s20, 0
	s_mov_b32 s25, 0
	s_branch .LBB0_57
.LBB0_56:                               ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s20, -1
                                        ; implicit-def: $vgpr18_vgpr19
                                        ; implicit-def: $sgpr25
.LBB0_57:                               ; %Flow512
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s20
	s_cbranch_vccnz .LBB0_59
; %bb.58:                               ;   in Loop: Header=BB0_28 Depth=1
	global_load_b64 v[18:19], v2, s[2:3]
	s_add_i32 s25, s26, -8
	s_add_u32 s2, s2, 8
	s_addc_u32 s3, s3, 0
.LBB0_59:                               ; %.loopexit26.i
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_cmp_gt_u32 s25, 7
	s_cbranch_scc1 .LBB0_64
; %bb.60:                               ;   in Loop: Header=BB0_28 Depth=1
	v_mov_b32_e32 v20, 0
	v_mov_b32_e32 v21, 0
	s_cmp_eq_u32 s25, 0
	s_cbranch_scc1 .LBB0_63
; %bb.61:                               ; %.preheader23.i.preheader
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b64 s[20:21], 0
	s_mov_b64 s[22:23], 0
.LBB0_62:                               ; %.preheader23.i
                                        ;   Parent Loop BB0_28 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s26, s2, s22
	s_addc_u32 s27, s3, s23
	s_add_u32 s22, s22, 1
	global_load_u8 v1, v2, s[26:27]
	s_addc_u32 s23, s23, 0
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v1, 0xffff, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[10:11], s20, v[1:2]
	s_add_u32 s20, s20, 8
	s_addc_u32 s21, s21, 0
	s_cmp_lg_u32 s25, s22
	v_or_b32_e32 v20, v10, v20
	s_delay_alu instid0(VALU_DEP_2)
	v_or_b32_e32 v21, v11, v21
	s_cbranch_scc1 .LBB0_62
.LBB0_63:                               ; %Flow505
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s20, 0
	s_mov_b32 s26, 0
	s_branch .LBB0_65
.LBB0_64:                               ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s20, -1
                                        ; implicit-def: $sgpr26
.LBB0_65:                               ; %Flow507
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s20
	s_cbranch_vccnz .LBB0_67
; %bb.66:                               ;   in Loop: Header=BB0_28 Depth=1
	global_load_b64 v[20:21], v2, s[2:3]
	s_add_i32 s26, s25, -8
	s_add_u32 s2, s2, 8
	s_addc_u32 s3, s3, 0
.LBB0_67:                               ; %.loopexit24.i
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_cmp_gt_u32 s26, 7
	s_cbranch_scc1 .LBB0_72
; %bb.68:                               ;   in Loop: Header=BB0_28 Depth=1
	v_mov_b32_e32 v22, 0
	v_mov_b32_e32 v23, 0
	s_cmp_eq_u32 s26, 0
	s_cbranch_scc1 .LBB0_71
; %bb.69:                               ; %.preheader21.i.preheader
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b64 s[20:21], 0
	s_mov_b64 s[22:23], 0
.LBB0_70:                               ; %.preheader21.i
                                        ;   Parent Loop BB0_28 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s28, s2, s22
	s_addc_u32 s29, s3, s23
	s_add_u32 s22, s22, 1
	global_load_u8 v1, v2, s[28:29]
	s_addc_u32 s23, s23, 0
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v1, 0xffff, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[10:11], s20, v[1:2]
	s_add_u32 s20, s20, 8
	s_addc_u32 s21, s21, 0
	s_cmp_lg_u32 s26, s22
	v_or_b32_e32 v22, v10, v22
	s_delay_alu instid0(VALU_DEP_2)
	v_or_b32_e32 v23, v11, v23
	s_cbranch_scc1 .LBB0_70
.LBB0_71:                               ; %Flow500
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s20, 0
	s_mov_b32 s25, 0
	s_branch .LBB0_73
.LBB0_72:                               ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s20, -1
                                        ; implicit-def: $vgpr22_vgpr23
                                        ; implicit-def: $sgpr25
.LBB0_73:                               ; %Flow502
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s20
	s_cbranch_vccnz .LBB0_75
; %bb.74:                               ;   in Loop: Header=BB0_28 Depth=1
	global_load_b64 v[22:23], v2, s[2:3]
	s_add_i32 s25, s26, -8
	s_add_u32 s2, s2, 8
	s_addc_u32 s3, s3, 0
.LBB0_75:                               ; %.loopexit22.i
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_cmp_gt_u32 s25, 7
	s_cbranch_scc1 .LBB0_80
; %bb.76:                               ;   in Loop: Header=BB0_28 Depth=1
	v_mov_b32_e32 v24, 0
	v_mov_b32_e32 v25, 0
	s_cmp_eq_u32 s25, 0
	s_cbranch_scc1 .LBB0_79
; %bb.77:                               ; %.preheader.i.preheader
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b64 s[20:21], 0
	s_mov_b64 s[22:23], s[2:3]
.LBB0_78:                               ; %.preheader.i
                                        ;   Parent Loop BB0_28 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_u8 v1, v2, s[22:23]
	s_add_i32 s25, s25, -1
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v1, 0xffff, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_lshlrev_b64 v[10:11], s20, v[1:2]
	s_add_u32 s20, s20, 8
	s_addc_u32 s21, s21, 0
	s_add_u32 s22, s22, 1
	s_addc_u32 s23, s23, 0
	s_cmp_lg_u32 s25, 0
	v_or_b32_e32 v24, v10, v24
	v_or_b32_e32 v25, v11, v25
	s_cbranch_scc1 .LBB0_78
.LBB0_79:                               ; %Flow495
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s20, 0
	s_branch .LBB0_81
.LBB0_80:                               ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s20, -1
.LBB0_81:                               ; %Flow497
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s20
	s_cbranch_vccnz .LBB0_83
; %bb.82:                               ;   in Loop: Header=BB0_28 Depth=1
	global_load_b64 v[24:25], v2, s[2:3]
.LBB0_83:                               ; %.loopexit.i
                                        ;   in Loop: Header=BB0_28 Depth=1
	v_dual_mov_b32 v1, v3 :: v_dual_mov_b32 v10, 0
	;;#ASMSTART
	;;#ASMEND
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s2, v1
	v_mov_b32_e32 v11, 0
	v_cmp_eq_u32_e64 s2, s2, v1
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_89
; %bb.84:                               ;   in Loop: Header=BB0_28 Depth=1
	global_load_b64 v[28:29], v2, s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_clause 0x1
	global_load_b64 v[10:11], v2, s[16:17] offset:40
	global_load_b64 v[26:27], v2, s[16:17]
	s_mov_b32 s20, exec_lo
	s_waitcnt vmcnt(1)
	v_and_b32_e32 v11, v11, v29
	v_and_b32_e32 v10, v10, v28
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v11, v11, 24
	v_mul_hi_u32 v31, v10, 24
	v_mul_lo_u32 v10, v10, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v11, v31, v11
	s_waitcnt vmcnt(0)
	v_add_co_u32 v10, vcc_lo, v26, v10
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v11, vcc_lo, v27, v11, vcc_lo
	global_load_b64 v[26:27], v[10:11], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[10:11], v2, v[26:29], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmpx_ne_u64_e64 v[10:11], v[28:29]
	s_cbranch_execz .LBB0_88
; %bb.85:                               ; %.preheader3.i.i19.i.preheader
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s21, 0
.LBB0_86:                               ; %.preheader3.i.i19.i
                                        ;   Parent Loop BB0_28 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_sleep 1
	s_clause 0x1
	global_load_b64 v[26:27], v2, s[16:17] offset:40
	global_load_b64 v[31:32], v2, s[16:17]
	v_dual_mov_b32 v29, v11 :: v_dual_mov_b32 v28, v10
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_and_b32_e32 v26, v26, v28
	s_waitcnt vmcnt(0)
	v_mad_u64_u32 v[10:11], null, v26, 24, v[31:32]
	v_and_b32_e32 v31, v27, v29
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[26:27], null, v31, 24, v[11:12]
	v_mov_b32_e32 v11, v26
	global_load_b64 v[26:27], v[10:11], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[10:11], v2, v[26:29], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmp_eq_u64_e32 vcc_lo, v[10:11], v[28:29]
	s_or_b32 s21, vcc_lo, s21
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s21
	s_cbranch_execnz .LBB0_86
; %bb.87:                               ; %Flow490
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_or_b32 exec_lo, exec_lo, s21
.LBB0_88:                               ; %Flow492
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s20
.LBB0_89:                               ; %.loopexit4.i.i14.i
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s3
	s_clause 0x1
	global_load_b64 v[31:32], v2, s[16:17] offset:40
	global_load_b128 v[26:29], v2, s[16:17]
	v_readfirstlane_b32 s20, v10
	v_readfirstlane_b32 s21, v11
	s_mov_b32 s27, exec_lo
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s22, v31
	v_readfirstlane_b32 s23, v32
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[22:23], s[20:21], s[22:23]
	s_mul_i32 s3, s23, 24
	s_mul_hi_u32 s25, s22, 24
	s_mul_i32 s26, s22, 24
	s_and_saveexec_b32 s28, s2
	s_cbranch_execz .LBB0_91
; %bb.90:                               ;   in Loop: Header=BB0_28 Depth=1
	v_dual_mov_b32 v10, s27 :: v_dual_mov_b32 v11, v2
	s_add_i32 s27, s25, s3
	s_waitcnt vmcnt(0)
	v_add_co_u32 v31, vcc_lo, v26, s26
	v_add_co_ci_u32_e32 v32, vcc_lo, s27, v27, vcc_lo
	global_store_b128 v[31:32], v[10:13], off offset:8
.LBB0_91:                               ;   in Loop: Header=BB0_28 Depth=1
	s_or_b32 exec_lo, exec_lo, s28
	v_cmp_gt_u64_e64 vcc_lo, s[10:11], 56
	v_or_b32_e32 v10, 0, v5
	v_or_b32_e32 v11, v4, v30
	s_lshl_b64 s[22:23], s[22:23], 12
	s_lshl_b32 s27, s18, 2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_add_i32 s27, s27, 28
	v_dual_cndmask_b32 v5, v10, v5 :: v_dual_cndmask_b32 v4, v11, v4
	v_lshlrev_b64 v[10:11], 6, v[1:2]
	s_waitcnt vmcnt(0)
	v_add_co_u32 v1, vcc_lo, v28, s22
	v_add_co_ci_u32_e32 v28, vcc_lo, s23, v29, vcc_lo
	s_and_b32 s27, s27, 0x1e0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_u32 v10, vcc_lo, v1, v10
	v_and_or_b32 v4, 0xffffff1f, v4, s27
	v_add_co_ci_u32_e32 v11, vcc_lo, v28, v11, vcc_lo
	s_clause 0x3
	global_store_b128 v[10:11], v[4:7], off
	global_store_b128 v[10:11], v[14:17], off offset:16
	global_store_b128 v[10:11], v[18:21], off offset:32
	global_store_b128 v[10:11], v[22:25], off offset:48
	s_and_saveexec_b32 s22, s2
	s_cbranch_execz .LBB0_99
; %bb.92:                               ;   in Loop: Header=BB0_28 Depth=1
	s_clause 0x1
	global_load_b64 v[18:19], v2, s[16:17] offset:32 glc
	global_load_b64 v[4:5], v2, s[16:17] offset:40
	v_dual_mov_b32 v16, s20 :: v_dual_mov_b32 v17, s21
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s28, v4
	v_readfirstlane_b32 s29, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[28:29], s[28:29], s[20:21]
	s_mul_i32 s23, s29, 24
	s_mul_hi_u32 s27, s28, 24
	s_mul_i32 s28, s28, 24
	s_add_i32 s27, s27, s23
	v_add_co_u32 v14, vcc_lo, v26, s28
	v_add_co_ci_u32_e32 v15, vcc_lo, s27, v27, vcc_lo
	s_mov_b32 s23, exec_lo
	global_store_b64 v[14:15], v[18:19], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[6:7], v2, v[16:19], s[16:17] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u64_e64 v[6:7], v[18:19]
	s_cbranch_execz .LBB0_95
; %bb.93:                               ; %.preheader1.i.i17.i.preheader
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s27, 0
.LBB0_94:                               ; %.preheader1.i.i17.i
                                        ;   Parent Loop BB0_28 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	v_dual_mov_b32 v4, s20 :: v_dual_mov_b32 v5, s21
	s_sleep 1
	global_store_b64 v[14:15], v[6:7], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[4:5], v2, v[4:7], s[16:17] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[4:5], v[6:7]
	v_dual_mov_b32 v7, v5 :: v_dual_mov_b32 v6, v4
	s_or_b32 s27, vcc_lo, s27
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s27
	s_cbranch_execnz .LBB0_94
.LBB0_95:                               ; %Flow488
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_or_b32 exec_lo, exec_lo, s23
	global_load_b64 v[4:5], v2, s[16:17] offset:16
	s_mov_b32 s27, exec_lo
	s_mov_b32 s23, exec_lo
	v_mbcnt_lo_u32_b32 v1, s27, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v1
	s_cbranch_execz .LBB0_97
; %bb.96:                               ;   in Loop: Header=BB0_28 Depth=1
	s_bcnt1_i32_b32 s27, s27
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v1, s27
	s_waitcnt vmcnt(0)
	global_atomic_add_u64 v[4:5], v[1:2], off offset:8
.LBB0_97:                               ;   in Loop: Header=BB0_28 Depth=1
	s_or_b32 exec_lo, exec_lo, s23
	s_waitcnt vmcnt(0)
	global_load_b64 v[6:7], v[4:5], off offset:16
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[6:7]
	s_cbranch_vccnz .LBB0_99
; %bb.98:                               ;   in Loop: Header=BB0_28 Depth=1
	global_load_b32 v1, v[4:5], off offset:24
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s23, v1
	s_waitcnt_vscnt null, 0x0
	global_store_b64 v[6:7], v[1:2], off
	s_and_b32 m0, s23, 0xff
	s_sendmsg sendmsg(MSG_INTERRUPT)
.LBB0_99:                               ; %Flow489
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_or_b32 exec_lo, exec_lo, s22
	s_add_i32 s25, s25, s3
	v_add_co_u32 v1, vcc_lo, v26, s26
	v_add_co_ci_u32_e32 v5, vcc_lo, s25, v27, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, v1, 20
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
.LBB0_100:                              ;   Parent Loop BB0_28 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	v_mov_b32_e32 v1, 1
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_102
; %bb.101:                              ;   in Loop: Header=BB0_100 Depth=2
	global_load_b32 v1, v[4:5], off glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_and_b32_e32 v1, 1, v1
.LBB0_102:                              ;   in Loop: Header=BB0_100 Depth=2
	s_or_b32 exec_lo, exec_lo, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s3, v1
	s_cmp_eq_u32 s3, 0
	s_cbranch_scc1 .LBB0_104
; %bb.103:                              ;   in Loop: Header=BB0_100 Depth=2
	s_mov_b32 s3, 0
	s_sleep 1
	s_branch .LBB0_105
.LBB0_104:                              ;   in Loop: Header=BB0_100 Depth=2
	s_mov_b32 s3, -1
.LBB0_105:                              ; %Flow483
                                        ;   in Loop: Header=BB0_100 Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s3
	s_cbranch_vccnz .LBB0_100
; %bb.106:                              ;   in Loop: Header=BB0_28 Depth=1
	global_load_b128 v[4:7], v[10:11], off
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_110
; %bb.107:                              ;   in Loop: Header=BB0_28 Depth=1
	s_clause 0x2
	global_load_b64 v[6:7], v2, s[16:17] offset:40
	global_load_b64 v[10:11], v2, s[16:17] offset:24 glc
	global_load_b64 v[16:17], v2, s[16:17]
	s_waitcnt vmcnt(2)
	v_add_co_u32 v1, vcc_lo, v6, 1
	v_add_co_ci_u32_e32 v18, vcc_lo, 0, v7, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v14, vcc_lo, v1, s20
	v_add_co_ci_u32_e32 v15, vcc_lo, s21, v18, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[14:15]
	v_dual_cndmask_b32 v15, v15, v18 :: v_dual_cndmask_b32 v14, v14, v1
	v_and_b32_e32 v1, v15, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v6, v14, v6
	v_mul_hi_u32 v7, v6, 24
	v_mul_lo_u32 v6, v6, 24
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_co_u32 v6, vcc_lo, v16, v6
	v_mov_b32_e32 v16, v10
	v_mul_lo_u32 v1, v1, 24
	v_add_nc_u32_e32 v1, v7, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e32 v7, vcc_lo, v17, v1, vcc_lo
	v_mov_b32_e32 v17, v11
	global_store_b64 v[6:7], v[10:11], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[16:17], v2, v[14:17], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, v[16:17], v[10:11]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_110
; %bb.108:                              ; %.preheader.i.i16.i.preheader
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_mov_b32 s2, 0
.LBB0_109:                              ; %.preheader.i.i16.i
                                        ;   Parent Loop BB0_28 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_sleep 1
	global_store_b64 v[6:7], v[16:17], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[10:11], v2, v[14:17], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[10:11], v[16:17]
	v_dual_mov_b32 v17, v11 :: v_dual_mov_b32 v16, v10
	s_or_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB0_109
.LBB0_110:                              ; %__ockl_hostcall_preview.exit20.i
                                        ;   in Loop: Header=BB0_28 Depth=1
	s_or_b32 exec_lo, exec_lo, s3
	s_sub_u32 s10, s10, s18
	s_subb_u32 s11, s11, s19
	s_add_u32 s8, s8, s18
	s_addc_u32 s9, s9, s19
	s_cmp_lg_u64 s[10:11], 0
	s_cbranch_scc1 .LBB0_28
; %bb.111:                              ; %Flow528
	s_mov_b32 s2, 0
	s_branch .LBB0_113
.LBB0_112:
	s_mov_b32 s2, -1
                                        ; implicit-def: $vgpr4_vgpr5
.LBB0_113:                              ; %Flow543
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 vcc_lo, exec_lo, s2
	s_cbranch_vccz .LBB0_142
; %bb.114:
	v_dual_mov_b32 v10, v3 :: v_dual_mov_b32 v1, 0
	;;#ASMSTART
	;;#ASMEND
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s2, v10
	v_mov_b32_e32 v2, 0
	v_cmp_eq_u32_e64 s2, s2, v10
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_120
; %bb.115:
	s_waitcnt vmcnt(0)
	v_mov_b32_e32 v4, 0
	s_mov_b32 s8, exec_lo
	global_load_b64 v[13:14], v4, s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_clause 0x1
	global_load_b64 v[1:2], v4, s[16:17] offset:40
	global_load_b64 v[5:6], v4, s[16:17]
	s_waitcnt vmcnt(1)
	v_and_b32_e32 v1, v1, v13
	v_and_b32_e32 v2, v2, v14
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v7, v1, 24
	v_mul_lo_u32 v2, v2, 24
	v_mul_lo_u32 v1, v1, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v2, v7, v2
	s_waitcnt vmcnt(0)
	v_add_co_u32 v1, vcc_lo, v5, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, v6, v2, vcc_lo
	global_load_b64 v[11:12], v[1:2], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[1:2], v4, v[11:14], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmpx_ne_u64_e64 v[1:2], v[13:14]
	s_cbranch_execz .LBB0_119
; %bb.116:                              ; %.preheader3.i.i.i100.preheader
	s_mov_b32 s9, 0
.LBB0_117:                              ; %.preheader3.i.i.i100
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	s_clause 0x1
	global_load_b64 v[5:6], v4, s[16:17] offset:40
	global_load_b64 v[11:12], v4, s[16:17]
	v_dual_mov_b32 v14, v2 :: v_dual_mov_b32 v13, v1
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v5, v5, v13
	v_and_b32_e32 v7, v6, v14
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, v5, 24, v[11:12]
	v_mad_u64_u32 v[5:6], null, v7, 24, v[2:3]
	s_delay_alu instid0(VALU_DEP_1)
	v_mov_b32_e32 v2, v5
	global_load_b64 v[11:12], v[1:2], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[1:2], v4, v[11:14], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmp_eq_u64_e32 vcc_lo, v[1:2], v[13:14]
	s_or_b32 s9, vcc_lo, s9
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s9
	s_cbranch_execnz .LBB0_117
; %bb.118:                              ; %Flow540
	s_or_b32 exec_lo, exec_lo, s9
.LBB0_119:                              ; %Flow542
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s8
.LBB0_120:                              ; %.loopexit4.i.i.i95
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s3
	v_mov_b32_e32 v11, 0
	v_readfirstlane_b32 s8, v1
	v_readfirstlane_b32 s9, v2
	s_mov_b32 s20, exec_lo
	s_clause 0x1
	global_load_b64 v[12:13], v11, s[16:17] offset:40
	global_load_b128 v[4:7], v11, s[16:17]
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s10, v12
	v_readfirstlane_b32 s11, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[10:11], s[8:9], s[10:11]
	s_mul_i32 s3, s11, 24
	s_mul_hi_u32 s18, s10, 24
	s_mul_i32 s19, s10, 24
	s_and_saveexec_b32 s21, s2
	s_cbranch_execz .LBB0_122
; %bb.121:
	v_dual_mov_b32 v12, s20 :: v_dual_mov_b32 v13, v11
	s_add_i32 s20, s18, s3
	s_waitcnt vmcnt(0)
	v_add_co_u32 v1, vcc_lo, v4, s19
	v_add_co_ci_u32_e32 v2, vcc_lo, s20, v5, vcc_lo
	v_dual_mov_b32 v14, 2 :: v_dual_mov_b32 v15, 1
	global_store_b128 v[1:2], v[12:15], off offset:8
.LBB0_122:
	s_or_b32 exec_lo, exec_lo, s21
	s_lshl_b64 s[10:11], s[10:11], 12
	v_lshlrev_b64 v[1:2], 6, v[10:11]
	s_waitcnt vmcnt(0)
	v_add_co_u32 v6, vcc_lo, v6, s10
	v_add_co_ci_u32_e32 v7, vcc_lo, s11, v7, vcc_lo
	s_mov_b32 s20, 0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, v6, v1
	s_mov_b32 s21, s20
	s_mov_b32 s22, s20
	s_mov_b32 s23, s20
	v_and_or_b32 v8, 0xffffff1f, v8, 32
	v_add_co_ci_u32_e32 v2, vcc_lo, v7, v2, vcc_lo
	v_mov_b32_e32 v10, v11
	v_dual_mov_b32 v12, s20 :: v_dual_mov_b32 v15, s23
	v_dual_mov_b32 v13, s21 :: v_dual_mov_b32 v14, s22
	s_clause 0x3
	global_store_b128 v[1:2], v[8:11], off
	global_store_b128 v[1:2], v[12:15], off offset:16
	global_store_b128 v[1:2], v[12:15], off offset:32
	global_store_b128 v[1:2], v[12:15], off offset:48
	s_and_saveexec_b32 s10, s2
	s_cbranch_execz .LBB0_130
; %bb.123:
	v_dual_mov_b32 v12, 0 :: v_dual_mov_b32 v13, s8
	v_mov_b32_e32 v14, s9
	s_clause 0x1
	global_load_b64 v[15:16], v12, s[16:17] offset:32 glc
	global_load_b64 v[6:7], v12, s[16:17] offset:40
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s20, v6
	v_readfirstlane_b32 s21, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[20:21], s[20:21], s[8:9]
	s_mul_i32 s11, s21, 24
	s_mul_hi_u32 s21, s20, 24
	s_mul_i32 s20, s20, 24
	s_add_i32 s21, s21, s11
	v_add_co_u32 v10, vcc_lo, v4, s20
	v_add_co_ci_u32_e32 v11, vcc_lo, s21, v5, vcc_lo
	s_mov_b32 s11, exec_lo
	global_store_b64 v[10:11], v[15:16], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[8:9], v12, v[13:16], s[16:17] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u64_e64 v[8:9], v[15:16]
	s_cbranch_execz .LBB0_126
; %bb.124:                              ; %.preheader1.i.i.i98.preheader
	s_mov_b32 s20, 0
.LBB0_125:                              ; %.preheader1.i.i.i98
                                        ; =>This Inner Loop Header: Depth=1
	v_dual_mov_b32 v6, s8 :: v_dual_mov_b32 v7, s9
	s_sleep 1
	global_store_b64 v[10:11], v[8:9], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[6:7], v12, v[6:9], s[16:17] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[6:7], v[8:9]
	v_dual_mov_b32 v9, v7 :: v_dual_mov_b32 v8, v6
	s_or_b32 s20, vcc_lo, s20
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s20
	s_cbranch_execnz .LBB0_125
.LBB0_126:                              ; %Flow538
	s_or_b32 exec_lo, exec_lo, s11
	v_mov_b32_e32 v9, 0
	s_mov_b32 s20, exec_lo
	s_mov_b32 s11, exec_lo
	v_mbcnt_lo_u32_b32 v8, s20, 0
	global_load_b64 v[6:7], v9, s[16:17] offset:16
	v_cmpx_eq_u32_e32 0, v8
	s_cbranch_execz .LBB0_128
; %bb.127:
	s_bcnt1_i32_b32 s20, s20
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v8, s20
	s_waitcnt vmcnt(0)
	global_atomic_add_u64 v[6:7], v[8:9], off offset:8
.LBB0_128:
	s_or_b32 exec_lo, exec_lo, s11
	s_waitcnt vmcnt(0)
	global_load_b64 v[8:9], v[6:7], off offset:16
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[8:9]
	s_cbranch_vccnz .LBB0_130
; %bb.129:
	global_load_b32 v6, v[6:7], off offset:24
	v_mov_b32_e32 v7, 0
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s11, v6
	s_waitcnt_vscnt null, 0x0
	global_store_b64 v[8:9], v[6:7], off
	s_and_b32 m0, s11, 0xff
	s_sendmsg sendmsg(MSG_INTERRUPT)
.LBB0_130:                              ; %Flow539
	s_or_b32 exec_lo, exec_lo, s10
	s_add_i32 s18, s18, s3
	v_add_co_u32 v4, vcc_lo, v4, s19
	v_add_co_ci_u32_e32 v5, vcc_lo, s18, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, v4, 20
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
.LBB0_131:                              ; =>This Inner Loop Header: Depth=1
	v_mov_b32_e32 v6, 1
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_133
; %bb.132:                              ;   in Loop: Header=BB0_131 Depth=1
	global_load_b32 v6, v[4:5], off glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_and_b32_e32 v6, 1, v6
.LBB0_133:                              ;   in Loop: Header=BB0_131 Depth=1
	s_or_b32 exec_lo, exec_lo, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s3, v6
	s_cmp_eq_u32 s3, 0
	s_cbranch_scc1 .LBB0_135
; %bb.134:                              ;   in Loop: Header=BB0_131 Depth=1
	s_mov_b32 s3, 0
	s_sleep 1
	s_branch .LBB0_136
.LBB0_135:                              ;   in Loop: Header=BB0_131 Depth=1
	s_mov_b32 s3, -1
.LBB0_136:                              ; %Flow533
                                        ;   in Loop: Header=BB0_131 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s3
	s_cbranch_vccnz .LBB0_131
; %bb.137:
	global_load_b128 v[4:7], v[1:2], off
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_141
; %bb.138:
	v_mov_b32_e32 v10, 0
	s_clause 0x2
	global_load_b64 v[1:2], v10, s[16:17] offset:40
	global_load_b64 v[11:12], v10, s[16:17] offset:24 glc
	global_load_b64 v[8:9], v10, s[16:17]
	s_waitcnt vmcnt(2)
	v_add_co_u32 v13, vcc_lo, v1, 1
	v_add_co_ci_u32_e32 v14, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v6, vcc_lo, v13, s8
	v_add_co_ci_u32_e32 v7, vcc_lo, s9, v14, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[6:7]
	v_dual_cndmask_b32 v7, v7, v14 :: v_dual_cndmask_b32 v6, v6, v13
	v_and_b32_e32 v2, v7, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v1, v6, v1
	v_mul_lo_u32 v2, v2, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v13, v1, 24
	v_mul_lo_u32 v1, v1, 24
	v_add_nc_u32_e32 v2, v13, v2
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_u32 v1, vcc_lo, v8, v1
	v_mov_b32_e32 v8, v11
	v_add_co_ci_u32_e32 v2, vcc_lo, v9, v2, vcc_lo
	v_mov_b32_e32 v9, v12
	global_store_b64 v[1:2], v[11:12], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[8:9], v10, v[6:9], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, v[8:9], v[11:12]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_141
; %bb.139:                              ; %.preheader.i.i.i97.preheader
	s_mov_b32 s2, 0
.LBB0_140:                              ; %.preheader.i.i.i97
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	global_store_b64 v[1:2], v[8:9], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[11:12], v10, v[6:9], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[11:12], v[8:9]
	v_dual_mov_b32 v8, v11 :: v_dual_mov_b32 v9, v12
	s_or_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB0_140
.LBB0_141:                              ; %__ockl_hostcall_preview.exit.i
	s_or_b32 exec_lo, exec_lo, s3
.LBB0_142:                              ; %__ockl_printf_append_string_n.exit
	s_waitcnt vmcnt(0)
	v_dual_mov_b32 v1, v3 :: v_dual_mov_b32 v6, 0
	;;#ASMSTART
	;;#ASMEND
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s2, v1
	v_mov_b32_e32 v7, 0
	v_cmp_eq_u32_e64 s2, s2, v1
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_148
; %bb.143:
	v_mov_b32_e32 v2, 0
	s_mov_b32 s8, exec_lo
	global_load_b64 v[8:9], v2, s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_clause 0x1
	global_load_b64 v[6:7], v2, s[16:17] offset:40
	global_load_b64 v[10:11], v2, s[16:17]
	s_waitcnt vmcnt(1)
	v_and_b32_e32 v7, v7, v9
	v_and_b32_e32 v6, v6, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v7, v7, 24
	v_mul_hi_u32 v12, v6, 24
	v_mul_lo_u32 v6, v6, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v7, v12, v7
	s_waitcnt vmcnt(0)
	v_add_co_u32 v6, vcc_lo, v10, v6
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v7, vcc_lo, v11, v7, vcc_lo
	global_load_b64 v[6:7], v[6:7], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[6:7], v2, v[6:9], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmpx_ne_u64_e64 v[6:7], v[8:9]
	s_cbranch_execz .LBB0_147
; %bb.144:                              ; %.preheader3.i.i.i107.preheader
	s_mov_b32 s9, 0
.LBB0_145:                              ; %.preheader3.i.i.i107
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	s_clause 0x1
	global_load_b64 v[10:11], v2, s[16:17] offset:40
	global_load_b64 v[12:13], v2, s[16:17]
	v_dual_mov_b32 v9, v7 :: v_dual_mov_b32 v8, v6
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_and_b32_e32 v10, v10, v8
	s_waitcnt vmcnt(0)
	v_mad_u64_u32 v[6:7], null, v10, 24, v[12:13]
	v_and_b32_e32 v12, v11, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[10:11], null, v12, 24, v[7:8]
	v_mov_b32_e32 v7, v10
	global_load_b64 v[6:7], v[6:7], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[6:7], v2, v[6:9], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmp_eq_u64_e32 vcc_lo, v[6:7], v[8:9]
	s_or_b32 s9, vcc_lo, s9
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s9
	s_cbranch_execnz .LBB0_145
; %bb.146:                              ; %Flow476
	s_or_b32 exec_lo, exec_lo, s9
.LBB0_147:                              ; %Flow478
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s8
.LBB0_148:                              ; %.loopexit4.i.i.i101
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s3
	v_mov_b32_e32 v2, 0
	v_readfirstlane_b32 s18, v6
	v_readfirstlane_b32 s19, v7
	s_mov_b32 s8, exec_lo
	s_clause 0x1
	global_load_b64 v[12:13], v2, s[16:17] offset:40
	global_load_b128 v[8:11], v2, s[16:17]
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s10, v12
	v_readfirstlane_b32 s11, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[20:21], s[18:19], s[10:11]
	s_mul_i32 s3, s21, 24
	s_mul_hi_u32 s23, s20, 24
	s_mul_i32 s25, s20, 24
	s_and_saveexec_b32 s9, s2
	s_cbranch_execz .LBB0_150
; %bb.149:
	v_dual_mov_b32 v12, s8 :: v_dual_mov_b32 v13, v2
	s_add_i32 s8, s23, s3
	s_waitcnt vmcnt(0)
	v_add_co_u32 v6, vcc_lo, v8, s25
	v_add_co_ci_u32_e32 v7, vcc_lo, s8, v9, vcc_lo
	v_dual_mov_b32 v14, 2 :: v_dual_mov_b32 v15, 1
	global_store_b128 v[6:7], v[12:15], off offset:8
.LBB0_150:
	s_or_b32 exec_lo, exec_lo, s9
	s_load_b128 s[8:11], s[0:1], 0x1c
	s_lshl_b64 s[0:1], s[20:21], 12
	s_mov_b32 s28, 0
	v_lshlrev_b64 v[12:13], 6, v[1:2]
	s_waitcnt vmcnt(0)
	v_add_co_u32 v1, vcc_lo, v10, s0
	s_mov_b32 s29, s28
	s_mov_b32 s30, s28
	s_mov_b32 s31, s28
	v_lshlrev_b32_e32 v15, 1, v0
	v_mov_b32_e32 v7, v2
	v_add_co_ci_u32_e32 v2, vcc_lo, s1, v11, vcc_lo
	v_add_co_u32 v1, vcc_lo, v1, v12
	v_and_or_b32 v4, 0xffffff1f, v4, 32
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v2, vcc_lo, v2, v13, vcc_lo
	v_dual_mov_b32 v10, s28 :: v_dual_mov_b32 v11, s29
	v_mov_b32_e32 v12, s30
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s20, s11, s10
	v_mov_b32_e32 v13, s31
	s_lshl_b32 s1, s20, 1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s21, s1, s9
	s_mul_i32 s22, s21, s14
	s_delay_alu instid0(SALU_CYCLE_1)
	v_add_nc_u32_e32 v6, s22, v15
	s_clause 0x3
	global_store_b128 v[1:2], v[4:7], off
	global_store_b128 v[1:2], v[10:13], off offset:16
	global_store_b128 v[1:2], v[10:13], off offset:32
	global_store_b128 v[1:2], v[10:13], off offset:48
	s_and_saveexec_b32 s0, s2
	s_cbranch_execz .LBB0_158
; %bb.151:
	v_dual_mov_b32 v7, 0 :: v_dual_mov_b32 v16, s18
	v_mov_b32_e32 v17, s19
	s_clause 0x1
	global_load_b64 v[18:19], v7, s[16:17] offset:32 glc
	global_load_b64 v[4:5], v7, s[16:17] offset:40
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s10, v4
	v_readfirstlane_b32 s11, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[10:11], s[10:11], s[18:19]
	s_mul_i32 s11, s11, 24
	s_mul_hi_u32 s26, s10, 24
	s_mul_i32 s10, s10, 24
	s_add_i32 s26, s26, s11
	v_add_co_u32 v4, vcc_lo, v8, s10
	v_add_co_ci_u32_e32 v5, vcc_lo, s26, v9, vcc_lo
	s_mov_b32 s10, exec_lo
	global_store_b64 v[4:5], v[18:19], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[12:13], v7, v[16:19], s[16:17] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u64_e64 v[12:13], v[18:19]
	s_cbranch_execz .LBB0_154
; %bb.152:                              ; %.preheader1.i.i.i105.preheader
	s_mov_b32 s11, 0
.LBB0_153:                              ; %.preheader1.i.i.i105
                                        ; =>This Inner Loop Header: Depth=1
	v_dual_mov_b32 v10, s18 :: v_dual_mov_b32 v11, s19
	s_sleep 1
	global_store_b64 v[4:5], v[12:13], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[10:11], v7, v[10:13], s[16:17] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[10:11], v[12:13]
	v_dual_mov_b32 v13, v11 :: v_dual_mov_b32 v12, v10
	s_or_b32 s11, vcc_lo, s11
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s11
	s_cbranch_execnz .LBB0_153
.LBB0_154:                              ; %Flow474
	s_or_b32 exec_lo, exec_lo, s10
	v_mov_b32_e32 v11, 0
	s_mov_b32 s11, exec_lo
	s_mov_b32 s10, exec_lo
	v_mbcnt_lo_u32_b32 v7, s11, 0
	global_load_b64 v[4:5], v11, s[16:17] offset:16
	v_cmpx_eq_u32_e32 0, v7
	s_cbranch_execz .LBB0_156
; %bb.155:
	s_bcnt1_i32_b32 s11, s11
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v10, s11
	s_waitcnt vmcnt(0)
	global_atomic_add_u64 v[4:5], v[10:11], off offset:8
.LBB0_156:
	s_or_b32 exec_lo, exec_lo, s10
	s_waitcnt vmcnt(0)
	global_load_b64 v[10:11], v[4:5], off offset:16
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[10:11]
	s_cbranch_vccnz .LBB0_158
; %bb.157:
	global_load_b32 v4, v[4:5], off offset:24
	v_mov_b32_e32 v5, 0
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s10, v4
	s_waitcnt_vscnt null, 0x0
	global_store_b64 v[10:11], v[4:5], off
	s_and_b32 m0, s10, 0xff
	s_sendmsg sendmsg(MSG_INTERRUPT)
.LBB0_158:                              ; %Flow475
	s_or_b32 exec_lo, exec_lo, s0
	s_add_i32 s23, s23, s3
	v_add_co_u32 v4, vcc_lo, v8, s25
	v_add_co_ci_u32_e32 v5, vcc_lo, s23, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, v4, 20
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
.LBB0_159:                              ; =>This Inner Loop Header: Depth=1
	v_mov_b32_e32 v7, 1
	s_and_saveexec_b32 s0, s2
	s_cbranch_execz .LBB0_161
; %bb.160:                              ;   in Loop: Header=BB0_159 Depth=1
	global_load_b32 v7, v[4:5], off glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_and_b32_e32 v7, 1, v7
.LBB0_161:                              ;   in Loop: Header=BB0_159 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s0, v7
	s_cmp_eq_u32 s0, 0
	s_cbranch_scc1 .LBB0_163
; %bb.162:                              ;   in Loop: Header=BB0_159 Depth=1
	s_mov_b32 s0, 0
	s_sleep 1
	s_branch .LBB0_164
.LBB0_163:                              ;   in Loop: Header=BB0_159 Depth=1
	s_mov_b32 s0, -1
.LBB0_164:                              ; %Flow469
                                        ;   in Loop: Header=BB0_159 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s0
	s_cbranch_vccnz .LBB0_159
; %bb.165:
	global_load_b64 v[7:8], v[1:2], off
	s_and_saveexec_b32 s0, s2
	s_cbranch_execz .LBB0_169
; %bb.166:
	v_mov_b32_e32 v4, 0
	s_clause 0x2
	global_load_b64 v[1:2], v4, s[16:17] offset:40
	global_load_b64 v[13:14], v4, s[16:17] offset:24 glc
	global_load_b64 v[11:12], v4, s[16:17]
	s_waitcnt vmcnt(2)
	v_add_co_u32 v5, vcc_lo, v1, 1
	v_add_co_ci_u32_e32 v16, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v9, vcc_lo, v5, s18
	v_add_co_ci_u32_e32 v10, vcc_lo, s19, v16, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[9:10]
	v_dual_cndmask_b32 v9, v9, v5 :: v_dual_cndmask_b32 v10, v10, v16
	v_and_b32_e32 v1, v9, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v2, v10, v2
	v_mul_hi_u32 v5, v1, 24
	v_mul_lo_u32 v1, v1, 24
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_co_u32 v1, vcc_lo, v11, v1
	v_mov_b32_e32 v11, v13
	v_mul_lo_u32 v2, v2, 24
	v_add_nc_u32_e32 v2, v5, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e32 v2, vcc_lo, v12, v2, vcc_lo
	v_mov_b32_e32 v12, v14
	global_store_b64 v[1:2], v[13:14], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[11:12], v4, v[9:12], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, v[11:12], v[13:14]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_169
; %bb.167:                              ; %.preheader.i.i.i104.preheader
	s_mov_b32 s2, 0
.LBB0_168:                              ; %.preheader.i.i.i104
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	global_store_b64 v[1:2], v[11:12], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[13:14], v4, v[9:12], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[13:14], v[11:12]
	v_dual_mov_b32 v11, v13 :: v_dual_mov_b32 v12, v14
	s_or_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB0_168
.LBB0_169:                              ; %__ockl_printf_append_args.exit
	s_or_b32 exec_lo, exec_lo, s0
	v_mov_b32_e32 v9, v3
	v_mov_b32_e32 v1, 0
	v_mov_b32_e32 v2, 0
	;;#ASMSTART
	;;#ASMEND
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s0, v9
	v_cmp_eq_u32_e64 s0, s0, v9
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s2, s0
	s_cbranch_execz .LBB0_175
; %bb.170:
	v_mov_b32_e32 v4, 0
	s_mov_b32 s3, exec_lo
	global_load_b64 v[12:13], v4, s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_clause 0x1
	global_load_b64 v[1:2], v4, s[16:17] offset:40
	global_load_b64 v[10:11], v4, s[16:17]
	s_waitcnt vmcnt(1)
	v_and_b32_e32 v1, v1, v12
	v_and_b32_e32 v2, v2, v13
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v5, v1, 24
	v_mul_lo_u32 v2, v2, 24
	v_mul_lo_u32 v1, v1, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v2, v5, v2
	s_waitcnt vmcnt(0)
	v_add_co_u32 v1, vcc_lo, v10, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, v11, v2, vcc_lo
	global_load_b64 v[10:11], v[1:2], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[1:2], v4, v[10:13], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmpx_ne_u64_e64 v[1:2], v[12:13]
	s_cbranch_execz .LBB0_174
; %bb.171:                              ; %.preheader3.i.i.i114.preheader
	s_mov_b32 s10, 0
.LBB0_172:                              ; %.preheader3.i.i.i114
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	s_clause 0x1
	global_load_b64 v[10:11], v4, s[16:17] offset:40
	global_load_b64 v[16:17], v4, s[16:17]
	v_dual_mov_b32 v13, v2 :: v_dual_mov_b32 v12, v1
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_and_b32_e32 v5, v10, v12
	s_waitcnt vmcnt(0)
	v_mad_u64_u32 v[1:2], null, v5, 24, v[16:17]
	v_and_b32_e32 v5, v11, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[10:11], null, v5, 24, v[2:3]
	v_mov_b32_e32 v2, v10
	global_load_b64 v[10:11], v[1:2], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[1:2], v4, v[10:13], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmp_eq_u64_e32 vcc_lo, v[1:2], v[12:13]
	s_or_b32 s10, vcc_lo, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s10
	s_cbranch_execnz .LBB0_172
; %bb.173:                              ; %Flow462
	s_or_b32 exec_lo, exec_lo, s10
.LBB0_174:                              ; %Flow464
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s3
.LBB0_175:                              ; %.loopexit4.i.i.i108
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s2
	v_mov_b32_e32 v10, 0
	v_readfirstlane_b32 s2, v1
	v_readfirstlane_b32 s3, v2
	s_mov_b32 s18, exec_lo
	s_clause 0x1
	global_load_b64 v[4:5], v10, s[16:17] offset:40
	global_load_b128 v[11:14], v10, s[16:17]
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s10, v4
	v_readfirstlane_b32 s11, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[10:11], s[2:3], s[10:11]
	s_mul_i32 s23, s11, 24
	s_mul_hi_u32 s25, s10, 24
	s_mul_i32 s26, s10, 24
	s_and_saveexec_b32 s19, s0
	s_cbranch_execz .LBB0_177
; %bb.176:
	v_dual_mov_b32 v16, s18 :: v_dual_mov_b32 v17, v10
	s_add_i32 s18, s25, s23
	s_waitcnt vmcnt(0)
	v_add_co_u32 v1, vcc_lo, v11, s26
	v_add_co_ci_u32_e32 v2, vcc_lo, s18, v12, vcc_lo
	v_dual_mov_b32 v18, 2 :: v_dual_mov_b32 v19, 1
	global_store_b128 v[1:2], v[16:19], off offset:8
.LBB0_177:
	s_or_b32 exec_lo, exec_lo, s19
	s_and_b32 s18, 0xffff, s24
	s_mul_i32 s9, s20, s9
	v_cvt_f32_u32_e32 v1, s18
	s_add_i32 s9, s9, s18
	s_sub_i32 s24, 0, s18
	s_add_i32 s9, s9, -1
	s_lshl_b64 s[10:11], s[10:11], 12
	v_rcp_iflag_f32_e32 v1, v1
	s_ashr_i32 s27, s9, 31
	s_waitcnt vmcnt(0)
	v_add_co_u32 v4, vcc_lo, v13, s10
	s_add_i32 s28, s9, s27
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v14, vcc_lo
	s_xor_b32 s28, s28, s27
	v_and_or_b32 v7, 0xffffff1f, v7, 32
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v1, v1
	v_readfirstlane_b32 s19, v1
	v_lshlrev_b64 v[1:2], 6, v[9:10]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_mul_i32 s24, s24, s19
	v_add_co_u32 v1, vcc_lo, v4, v1
	s_mul_hi_u32 s24, s19, s24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_add_co_ci_u32_e32 v2, vcc_lo, v5, v2, vcc_lo
	s_add_i32 s19, s19, s24
	s_mul_hi_u32 s19, s28, s19
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s19, s19, s18
	s_sub_i32 s19, s28, s19
	s_mov_b32 s28, 0
	s_sub_i32 s24, s19, s18
	s_cmp_ge_u32 s19, s18
	s_mov_b32 s29, s28
	s_cselect_b32 s19, s24, s19
	s_mov_b32 s30, s28
	s_sub_i32 s24, s19, s18
	s_cmp_ge_u32 s19, s18
	s_mov_b32 s31, s28
	v_mov_b32_e32 v16, s28
	s_cselect_b32 s19, s24, s19
	v_mov_b32_e32 v17, s29
	s_xor_b32 s19, s19, s27
	v_mov_b32_e32 v19, s31
	s_sub_i32 s10, s27, s19
	v_mov_b32_e32 v18, s30
	s_add_i32 s9, s9, s10
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b32 s19, s9, 1
	s_add_i32 s9, s19, s22
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v9, s9
	s_clause 0x3
	global_store_b128 v[1:2], v[7:10], off
	global_store_b128 v[1:2], v[16:19], off offset:16
	global_store_b128 v[1:2], v[16:19], off offset:32
	global_store_b128 v[1:2], v[16:19], off offset:48
	s_and_saveexec_b32 s10, s0
	s_cbranch_execz .LBB0_185
; %bb.178:
	v_dual_mov_b32 v13, 0 :: v_dual_mov_b32 v16, s2
	v_mov_b32_e32 v17, s3
	s_clause 0x1
	global_load_b64 v[18:19], v13, s[16:17] offset:32 glc
	global_load_b64 v[4:5], v13, s[16:17] offset:40
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s28, v4
	v_readfirstlane_b32 s29, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[28:29], s[28:29], s[2:3]
	s_mul_i32 s11, s29, 24
	s_mul_hi_u32 s24, s28, 24
	s_mul_i32 s27, s28, 24
	s_add_i32 s24, s24, s11
	v_add_co_u32 v4, vcc_lo, v11, s27
	v_add_co_ci_u32_e32 v5, vcc_lo, s24, v12, vcc_lo
	s_mov_b32 s11, exec_lo
	global_store_b64 v[4:5], v[18:19], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[9:10], v13, v[16:19], s[16:17] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u64_e64 v[9:10], v[18:19]
	s_cbranch_execz .LBB0_181
; %bb.179:                              ; %.preheader1.i.i.i112.preheader
	s_mov_b32 s24, 0
.LBB0_180:                              ; %.preheader1.i.i.i112
                                        ; =>This Inner Loop Header: Depth=1
	v_dual_mov_b32 v7, s2 :: v_dual_mov_b32 v8, s3
	s_sleep 1
	global_store_b64 v[4:5], v[9:10], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[7:8], v13, v[7:10], s[16:17] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[7:8], v[9:10]
	v_dual_mov_b32 v10, v8 :: v_dual_mov_b32 v9, v7
	s_or_b32 s24, vcc_lo, s24
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s24
	s_cbranch_execnz .LBB0_180
.LBB0_181:                              ; %Flow460
	s_or_b32 exec_lo, exec_lo, s11
	v_mov_b32_e32 v8, 0
	s_mov_b32 s24, exec_lo
	s_mov_b32 s11, exec_lo
	v_mbcnt_lo_u32_b32 v7, s24, 0
	global_load_b64 v[4:5], v8, s[16:17] offset:16
	v_cmpx_eq_u32_e32 0, v7
	s_cbranch_execz .LBB0_183
; %bb.182:
	s_bcnt1_i32_b32 s24, s24
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v7, s24
	s_waitcnt vmcnt(0)
	global_atomic_add_u64 v[4:5], v[7:8], off offset:8
.LBB0_183:
	s_or_b32 exec_lo, exec_lo, s11
	s_waitcnt vmcnt(0)
	global_load_b64 v[7:8], v[4:5], off offset:16
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[7:8]
	s_cbranch_vccnz .LBB0_185
; %bb.184:
	global_load_b32 v4, v[4:5], off offset:24
	v_mov_b32_e32 v5, 0
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s11, v4
	s_waitcnt_vscnt null, 0x0
	global_store_b64 v[7:8], v[4:5], off
	s_and_b32 m0, s11, 0xff
	s_sendmsg sendmsg(MSG_INTERRUPT)
.LBB0_185:                              ; %Flow461
	s_or_b32 exec_lo, exec_lo, s10
	s_add_i32 s25, s25, s23
	v_add_co_u32 v4, vcc_lo, v11, s26
	v_add_co_ci_u32_e32 v5, vcc_lo, s25, v12, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, v4, 20
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
.LBB0_186:                              ; =>This Inner Loop Header: Depth=1
	v_mov_b32_e32 v7, 1
	s_and_saveexec_b32 s10, s0
	s_cbranch_execz .LBB0_188
; %bb.187:                              ;   in Loop: Header=BB0_186 Depth=1
	global_load_b32 v7, v[4:5], off glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_and_b32_e32 v7, 1, v7
.LBB0_188:                              ;   in Loop: Header=BB0_186 Depth=1
	s_or_b32 exec_lo, exec_lo, s10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s10, v7
	s_cmp_eq_u32 s10, 0
	s_cbranch_scc1 .LBB0_190
; %bb.189:                              ;   in Loop: Header=BB0_186 Depth=1
	s_mov_b32 s10, 0
	s_sleep 1
	s_branch .LBB0_191
.LBB0_190:                              ;   in Loop: Header=BB0_186 Depth=1
	s_mov_b32 s10, -1
.LBB0_191:                              ; %Flow455
                                        ;   in Loop: Header=BB0_186 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s10
	s_cbranch_vccnz .LBB0_186
; %bb.192:
	global_load_b64 v[1:2], v[1:2], off
	s_and_saveexec_b32 s10, s0
	s_cbranch_execz .LBB0_196
; %bb.193:
	v_mov_b32_e32 v11, 0
	s_clause 0x2
	global_load_b64 v[4:5], v11, s[16:17] offset:40
	global_load_b64 v[12:13], v11, s[16:17] offset:24 glc
	global_load_b64 v[9:10], v11, s[16:17]
	s_waitcnt vmcnt(2)
	v_add_co_u32 v14, vcc_lo, v4, 1
	v_add_co_ci_u32_e32 v16, vcc_lo, 0, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v7, vcc_lo, v14, s2
	v_add_co_ci_u32_e32 v8, vcc_lo, s3, v16, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[7:8]
	v_dual_cndmask_b32 v8, v8, v16 :: v_dual_cndmask_b32 v7, v7, v14
	v_and_b32_e32 v5, v8, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v4, v7, v4
	v_mul_lo_u32 v5, v5, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v14, v4, 24
	v_mul_lo_u32 v4, v4, 24
	v_add_nc_u32_e32 v5, v14, v5
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_u32 v4, vcc_lo, v9, v4
	v_mov_b32_e32 v9, v12
	v_add_co_ci_u32_e32 v5, vcc_lo, v10, v5, vcc_lo
	v_mov_b32_e32 v10, v13
	global_store_b64 v[4:5], v[12:13], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[9:10], v11, v[7:10], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, v[9:10], v[12:13]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_196
; %bb.194:                              ; %.preheader.i.i.i111.preheader
	s_mov_b32 s0, 0
.LBB0_195:                              ; %.preheader.i.i.i111
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	global_store_b64 v[4:5], v[9:10], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[12:13], v11, v[7:10], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[12:13], v[9:10]
	v_dual_mov_b32 v9, v12 :: v_dual_mov_b32 v10, v13
	s_or_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s0
	s_cbranch_execnz .LBB0_195
.LBB0_196:                              ; %__ockl_printf_append_args.exit115
	s_or_b32 exec_lo, exec_lo, s10
	;;#ASMSTART
	;;#ASMEND
	v_readfirstlane_b32 s0, v3
	v_mov_b32_e32 v11, 0
	v_mov_b32_e32 v12, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_eq_u32_e64 s0, s0, v3
	s_and_saveexec_b32 s2, s0
	s_cbranch_execz .LBB0_202
; %bb.197:
	v_mov_b32_e32 v4, 0
	s_mov_b32 s3, exec_lo
	global_load_b64 v[9:10], v4, s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_clause 0x1
	global_load_b64 v[7:8], v4, s[16:17] offset:40
	global_load_b64 v[11:12], v4, s[16:17]
	s_waitcnt vmcnt(1)
	v_and_b32_e32 v5, v8, v10
	v_and_b32_e32 v7, v7, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v5, v5, 24
	v_mul_hi_u32 v8, v7, 24
	v_mul_lo_u32 v7, v7, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v5, v8, v5
	s_waitcnt vmcnt(0)
	v_add_co_u32 v7, vcc_lo, v11, v7
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v8, vcc_lo, v12, v5, vcc_lo
	global_load_b64 v[7:8], v[7:8], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[11:12], v4, v[7:10], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmpx_ne_u64_e64 v[11:12], v[9:10]
	s_cbranch_execz .LBB0_201
; %bb.198:                              ; %.preheader3.i.i.i122.preheader
	s_mov_b32 s10, 0
.LBB0_199:                              ; %.preheader3.i.i.i122
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	s_clause 0x1
	global_load_b64 v[7:8], v4, s[16:17] offset:40
	global_load_b64 v[13:14], v4, s[16:17]
	v_dual_mov_b32 v9, v11 :: v_dual_mov_b32 v10, v12
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_and_b32_e32 v5, v7, v9
	s_waitcnt vmcnt(0)
	v_mad_u64_u32 v[11:12], null, v5, 24, v[13:14]
	v_and_b32_e32 v13, v8, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v5, v12
	v_mad_u64_u32 v[7:8], null, v13, 24, v[5:6]
	s_delay_alu instid0(VALU_DEP_1)
	v_mov_b32_e32 v12, v7
	global_load_b64 v[7:8], v[11:12], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[11:12], v4, v[7:10], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmp_eq_u64_e32 vcc_lo, v[11:12], v[9:10]
	s_or_b32 s10, vcc_lo, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s10
	s_cbranch_execnz .LBB0_199
; %bb.200:                              ; %Flow448
	s_or_b32 exec_lo, exec_lo, s10
.LBB0_201:                              ; %Flow450
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s3
.LBB0_202:                              ; %.loopexit4.i.i.i116
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s2
	v_mov_b32_e32 v4, 0
	v_readfirstlane_b32 s2, v11
	v_readfirstlane_b32 s3, v12
	s_mov_b32 s26, exec_lo
	s_clause 0x1
	global_load_b64 v[13:14], v4, s[16:17] offset:40
	global_load_b128 v[7:10], v4, s[16:17]
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s10, v13
	v_readfirstlane_b32 s11, v14
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[10:11], s[2:3], s[10:11]
	s_mul_i32 s23, s11, 24
	s_mul_hi_u32 s24, s10, 24
	s_mul_i32 s25, s10, 24
	s_and_saveexec_b32 s27, s0
	s_cbranch_execz .LBB0_204
; %bb.203:
	v_dual_mov_b32 v11, s26 :: v_dual_mov_b32 v12, v4
	s_add_i32 s26, s24, s23
	s_waitcnt vmcnt(0)
	v_add_co_u32 v16, vcc_lo, v7, s25
	v_add_co_ci_u32_e32 v17, vcc_lo, s26, v8, vcc_lo
	v_dual_mov_b32 v13, 2 :: v_dual_mov_b32 v14, 1
	global_store_b128 v[16:17], v[11:14], off offset:8
.LBB0_204:
	s_or_b32 exec_lo, exec_lo, s27
	s_lshl_b64 s[10:11], s[10:11], 12
	v_lshlrev_b64 v[11:12], 6, v[3:4]
	s_waitcnt vmcnt(0)
	v_add_co_u32 v3, vcc_lo, v9, s10
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v10, vcc_lo
	s_mov_b32 s28, 0
	s_add_i32 s10, s22, s21
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_u32 v13, vcc_lo, v3, v11
	s_mov_b32 s29, s28
	s_mov_b32 s30, s28
	s_mov_b32 s31, s28
	v_and_or_b32 v1, 0xffffff1d, v1, 34
	v_add_co_ci_u32_e32 v14, vcc_lo, v5, v12, vcc_lo
	v_mov_b32_e32 v3, s10
	v_dual_mov_b32 v9, s28 :: v_dual_mov_b32 v12, s31
	v_dual_mov_b32 v10, s29 :: v_dual_mov_b32 v11, s30
	s_clause 0x3
	global_store_b128 v[13:14], v[1:4], off
	global_store_b128 v[13:14], v[9:12], off offset:16
	global_store_b128 v[13:14], v[9:12], off offset:32
	global_store_b128 v[13:14], v[9:12], off offset:48
	s_and_saveexec_b32 s11, s0
	s_cbranch_execz .LBB0_212
; %bb.205:
	v_dual_mov_b32 v5, 0 :: v_dual_mov_b32 v12, s3
	v_mov_b32_e32 v11, s2
	s_clause 0x1
	global_load_b64 v[13:14], v5, s[16:17] offset:32 glc
	global_load_b64 v[1:2], v5, s[16:17] offset:40
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s26, v1
	v_readfirstlane_b32 s27, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[26:27], s[26:27], s[2:3]
	s_mul_i32 s22, s27, 24
	s_mul_hi_u32 s27, s26, 24
	s_mul_i32 s26, s26, 24
	s_add_i32 s27, s27, s22
	v_add_co_u32 v9, vcc_lo, v7, s26
	v_add_co_ci_u32_e32 v10, vcc_lo, s27, v8, vcc_lo
	s_mov_b32 s22, exec_lo
	global_store_b64 v[9:10], v[13:14], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[3:4], v5, v[11:14], s[16:17] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u64_e64 v[3:4], v[13:14]
	s_cbranch_execz .LBB0_208
; %bb.206:                              ; %.preheader1.i.i.i120.preheader
	s_mov_b32 s26, 0
.LBB0_207:                              ; %.preheader1.i.i.i120
                                        ; =>This Inner Loop Header: Depth=1
	v_dual_mov_b32 v1, s2 :: v_dual_mov_b32 v2, s3
	s_sleep 1
	global_store_b64 v[9:10], v[3:4], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[1:2], v5, v[1:4], s[16:17] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[1:2], v[3:4]
	v_dual_mov_b32 v4, v2 :: v_dual_mov_b32 v3, v1
	s_or_b32 s26, vcc_lo, s26
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s26
	s_cbranch_execnz .LBB0_207
.LBB0_208:                              ; %Flow446
	s_or_b32 exec_lo, exec_lo, s22
	v_mov_b32_e32 v4, 0
	s_mov_b32 s26, exec_lo
	s_mov_b32 s22, exec_lo
	v_mbcnt_lo_u32_b32 v3, s26, 0
	global_load_b64 v[1:2], v4, s[16:17] offset:16
	v_cmpx_eq_u32_e32 0, v3
	s_cbranch_execz .LBB0_210
; %bb.209:
	s_bcnt1_i32_b32 s26, s26
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v3, s26
	s_waitcnt vmcnt(0)
	global_atomic_add_u64 v[1:2], v[3:4], off offset:8
.LBB0_210:
	s_or_b32 exec_lo, exec_lo, s22
	s_waitcnt vmcnt(0)
	global_load_b64 v[3:4], v[1:2], off offset:16
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[3:4]
	s_cbranch_vccnz .LBB0_212
; %bb.211:
	global_load_b32 v1, v[1:2], off offset:24
	v_mov_b32_e32 v2, 0
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s22, v1
	s_waitcnt_vscnt null, 0x0
	global_store_b64 v[3:4], v[1:2], off
	s_and_b32 m0, s22, 0xff
	s_sendmsg sendmsg(MSG_INTERRUPT)
.LBB0_212:                              ; %Flow447
	s_or_b32 exec_lo, exec_lo, s11
	s_add_i32 s24, s24, s23
	v_add_co_u32 v1, vcc_lo, v7, s25
	v_add_co_ci_u32_e32 v2, vcc_lo, s24, v8, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, v1, 20
	v_add_co_ci_u32_e32 v2, vcc_lo, 0, v2, vcc_lo
.LBB0_213:                              ; =>This Inner Loop Header: Depth=1
	v_mov_b32_e32 v3, 1
	s_and_saveexec_b32 s11, s0
	s_cbranch_execz .LBB0_215
; %bb.214:                              ;   in Loop: Header=BB0_213 Depth=1
	global_load_b32 v3, v[1:2], off glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_and_b32_e32 v3, 1, v3
.LBB0_215:                              ;   in Loop: Header=BB0_213 Depth=1
	s_or_b32 exec_lo, exec_lo, s11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s11, v3
	s_cmp_eq_u32 s11, 0
	s_cbranch_scc1 .LBB0_217
; %bb.216:                              ;   in Loop: Header=BB0_213 Depth=1
	s_mov_b32 s11, 0
	s_sleep 1
	s_branch .LBB0_218
.LBB0_217:                              ;   in Loop: Header=BB0_213 Depth=1
	s_mov_b32 s11, -1
.LBB0_218:                              ; %Flow441
                                        ;   in Loop: Header=BB0_213 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s11
	s_cbranch_vccnz .LBB0_213
; %bb.219:
	s_and_saveexec_b32 s11, s0
	s_cbranch_execz .LBB0_223
; %bb.220:
	v_mov_b32_e32 v5, 0
	s_clause 0x2
	global_load_b64 v[3:4], v5, s[16:17] offset:40
	global_load_b64 v[9:10], v5, s[16:17] offset:24 glc
	global_load_b64 v[7:8], v5, s[16:17]
	s_waitcnt vmcnt(2)
	v_add_co_u32 v11, vcc_lo, v3, 1
	v_add_co_ci_u32_e32 v12, vcc_lo, 0, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, v11, s2
	v_add_co_ci_u32_e32 v2, vcc_lo, s3, v12, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[1:2]
	v_dual_cndmask_b32 v2, v2, v12 :: v_dual_cndmask_b32 v1, v1, v11
	v_and_b32_e32 v4, v2, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v3, v1, v3
	v_mul_lo_u32 v4, v4, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v11, v3, 24
	v_mul_lo_u32 v3, v3, 24
	v_add_nc_u32_e32 v4, v11, v4
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_u32 v7, vcc_lo, v7, v3
	v_mov_b32_e32 v3, v9
	v_add_co_ci_u32_e32 v8, vcc_lo, v8, v4, vcc_lo
	v_mov_b32_e32 v4, v10
	global_store_b64 v[7:8], v[9:10], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[3:4], v5, v[1:4], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, v[3:4], v[9:10]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_223
; %bb.221:                              ; %.preheader.i.i.i119.preheader
	s_mov_b32 s0, 0
.LBB0_222:                              ; %.preheader.i.i.i119
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	global_store_b64 v[7:8], v[3:4], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[9:10], v5, v[1:4], s[16:17] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[9:10], v[3:4]
	v_dual_mov_b32 v3, v9 :: v_dual_mov_b32 v4, v10
	s_or_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s0
	s_cbranch_execnz .LBB0_222
.LBB0_223:                              ; %__ockl_printf_append_args.exit123
	s_or_b32 exec_lo, exec_lo, s11
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s0, exec_lo
	v_cmpx_gt_i32_e64 s19, v15
	s_cbranch_execz .LBB0_234
; %bb.224:                              ; %.lr.ph149
	s_ashr_i32 s0, s1, 31
	v_lshl_add_u32 v5, v15, 2, 0
	s_add_i32 s2, s1, s0
	v_cmp_gt_i32_e32 vcc_lo, s20, v0
	s_xor_b32 s2, s2, s0
	s_lshl_b32 s3, s18, 1
	v_cvt_f32_u32_e32 v1, s2
	s_sub_i32 s0, 0, s2
	s_add_i32 s10, s10, -1
	v_mov_b32_e32 v9, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v3, v1
	v_mad_u64_u32 v[1:2], null, s21, s15, v[15:16]
	s_delay_alu instid0(VALU_DEP_2)
	v_mul_lo_u32 v4, s0, v3
	s_mul_i32 s0, s14, s8
	s_mov_b32 s8, 0
	s_add_i32 s11, s0, s15
	s_lshl_b32 s14, s20, 3
	s_mul_i32 s11, s11, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v3, v4
	v_add_nc_u32_e32 v4, s20, v0
	v_cmp_gt_i32_e64 s0, s18, v4
	v_lshl_add_u32 v0, v4, 3, 0
	s_delay_alu instid0(VALU_DEP_4)
	v_add_nc_u32_e32 v8, v3, v2
.LBB0_225:                              ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_230 Depth 2
	s_mov_b32 s15, exec_lo
	ds_store_2addr_b32 v5, v9, v9 offset1:1
	v_cmpx_gt_i32_e64 s10, v6
	s_cbranch_execz .LBB0_227
; %bb.226:                              ;   in Loop: Header=BB0_225 Depth=1
	v_ashrrev_i32_e32 v2, 31, v1
	v_ashrrev_i32_e32 v7, 31, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	v_lshlrev_b64 v[10:11], 2, v[6:7]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v2, s1, s6, v2
	v_add_co_ci_u32_e64 v3, s1, s7, v3, s1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v10, s1, s4, v10
	v_add_co_ci_u32_e64 v11, s1, s5, v11, s1
	global_load_b64 v[2:3], v[2:3], off
	global_load_b64 v[10:11], v[10:11], off
	s_waitcnt vmcnt(1)
	v_add_f32_e32 v7, v2, v3
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v12, v10, v11
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v7, v10, v7 :: v_dual_sub_f32 v10, v11, v10
	v_fma_f32 v3, -v12, v3, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v7, v10, v2 :: v_dual_add_f32 v2, 0, v3
	v_add_f32_e32 v3, 0, v7
	ds_store_2addr_b32 v5, v2, v3 offset1:1
.LBB0_227:                              ;   in Loop: Header=BB0_225 Depth=1
	s_or_b32 exec_lo, exec_lo, s15
	v_ashrrev_i32_e32 v2, 31, v15
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	v_add_nc_u32_e32 v3, v15, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v3, v3, v2
	v_mul_hi_u32 v7, v3, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v7, v7, s2
	v_sub_nc_u32_e32 v3, v3, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_subrev_nc_u32_e32 v7, s2, v3
	v_cmp_le_u32_e64 s1, s2, v3
	v_cndmask_b32_e64 v3, v3, v7, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_subrev_nc_u32_e32 v7, s2, v3
	v_cmp_le_u32_e64 s1, s2, v3
	v_cndmask_b32_e64 v3, v3, v7, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v3, v3, v2
	v_sub_nc_u32_e32 v7, v3, v2
	s_and_saveexec_b32 s15, vcc_lo
	s_cbranch_execz .LBB0_233
; %bb.228:                              ; %.preheader
                                        ;   in Loop: Header=BB0_225 Depth=1
	ds_load_2addr_b32 v[2:3], v5 offset1:1
	s_and_saveexec_b32 s16, s0
	s_cbranch_execz .LBB0_232
; %bb.229:                              ; %.lr.ph.preheader
                                        ;   in Loop: Header=BB0_225 Depth=1
	v_mov_b32_e32 v10, v0
	v_mov_b32_e32 v11, v4
	s_mov_b32 s17, 0
.LBB0_230:                              ; %.lr.ph
                                        ;   Parent Loop BB0_225 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	ds_load_b32 v12, v10
	v_add_nc_u32_e32 v11, s20, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_le_i32_e64 s1, s18, v11
	s_or_b32 s17, s1, s17
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v2, v12, v2
	ds_store_b32 v5, v2
	ds_load_b32 v12, v10 offset:4
	s_waitcnt lgkmcnt(0)
	v_dual_add_f32 v3, v12, v3 :: v_dual_add_nc_u32 v10, s14, v10
	ds_store_b32 v5, v3 offset:4
	s_and_not1_b32 exec_lo, exec_lo, s17
	s_cbranch_execnz .LBB0_230
; %bb.231:                              ; %Flow
                                        ;   in Loop: Header=BB0_225 Depth=1
	s_or_b32 exec_lo, exec_lo, s17
.LBB0_232:                              ; %Flow434
                                        ;   in Loop: Header=BB0_225 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	s_or_b32 exec_lo, exec_lo, s16
	v_add_nc_u32_e32 v10, s11, v7
	v_ashrrev_i32_e32 v11, 31, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[10:11], 2, v[10:11]
	v_add_co_u32 v10, s1, s12, v10
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v11, s1, s13, v11, s1
	global_load_b32 v12, v[10:11], off
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_add_f32_e32 v2, v2, v12
	global_store_b64 v[10:11], v[2:3], off
.LBB0_233:                              ; %Flow435
                                        ;   in Loop: Header=BB0_225 Depth=1
	s_or_b32 exec_lo, exec_lo, s15
	v_add_nc_u32_e32 v6, s3, v6
	v_add_nc_u32_e32 v1, s3, v1
	v_add_nc_u32_e32 v15, s3, v7
	s_waitcnt_vscnt null, 0x0
	s_barrier
	v_cmp_le_i32_e64 s1, s9, v6
	buffer_gl0_inv
	s_or_b32 s8, s1, s8
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s8
	s_cbranch_execnz .LBB0_225
.LBB0_234:                              ; %._crit_edge150
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 304
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 33
		.amdhsa_next_free_sgpr 32
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.section	.text._ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii,"axG",@progbits,_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii,comdat
.Lfunc_end0:
	.size	_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii, .Lfunc_end0-_ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 10320
; NumSgprs: 34
; NumVgprs: 33
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 4
; VGPRBlocks: 4
; NumSGPRsForWavesPerEU: 34
; NumVGPRsForWavesPerEU: 33
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	.str,@object                    ; @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.str:
	.asciz	"N_idx:%d, last_block_idx:%d, last_N_idx:%d\n"
	.size	.str, 44

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .actual_access:  read_only
        .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .actual_access:  read_only
        .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         52
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         60
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         62
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         64
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         66
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         68
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         70
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         112
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         128
        .size:           8
        .value_kind:     hidden_hostcall_buffer
      - .offset:         168
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 304
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii
    .private_segment_fixed_size: 0
    .sgpr_count:     34
    .sgpr_spill_count: 0
    .symbol:         _ZN12_GLOBAL__N_123complex_mul_cuda_kernelIfEEvPKT_S3_PS1_iiiii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     33
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
