#include <stdio.h>
#include <stdlib.h>
#include <string.h> /* memcpy */
#include <math.h>
#include <stdint.h>

void *cuda_upload_var(void *host_var, int size)
{
	void *cuda_var;
	cudaMalloc(&cuda_var, 4);
	cudaMemcpy(cuda_var, host_var, size, cudaMemcpyHostToDevice);
	return cuda_var;
}
void cuda_download_var(void *cuda_var, void *host_var, int size)
{
	cudaMemcpy(host_var, cuda_var, size, cudaMemcpyDeviceToHost);
	cudaFree(cuda_var);
}

typedef struct intfield5
{
    int *m;
    int size[5];
    int is_device_field;
} intfield5;

void memcpy_field_intfield5(intfield5 dst, intfield5 src)
{
    if (dst.is_device_field == 0 && src.is_device_field == 0) {
        cudaMemcpy(dst.m, src.m, (sizeof(*dst.m))*dst.size[0]*dst.size[1]*dst.size[2]*dst.size[3]*dst.size[4], cudaMemcpyHostToHost);
    }
    if (dst.is_device_field == 1 && src.is_device_field == 0) {
        cudaMemcpy(dst.m, src.m, (sizeof(*dst.m))*dst.size[0]*dst.size[1]*dst.size[2]*dst.size[3]*dst.size[4], cudaMemcpyHostToDevice);
    }
    if (dst.is_device_field == 0 && src.is_device_field == 1) {
        cudaMemcpy(dst.m, src.m, (sizeof(*dst.m))*dst.size[0]*dst.size[1]*dst.size[2]*dst.size[3]*dst.size[4], cudaMemcpyDeviceToHost);
    }
    if (dst.is_device_field == 1 && src.is_device_field == 1) {
        cudaMemcpy(dst.m, src.m, (sizeof(*dst.m))*dst.size[0]*dst.size[1]*dst.size[2]*dst.size[3]*dst.size[4], cudaMemcpyDeviceToDevice);
    }
}

int size_intfield5(intfield5 field, int index)
{
    return field.size[index];
}

intfield5 alloc_device_field_intfield5(int size_0, int size_1, int size_2, int size_3, int size_4)
{
    intfield5 field;
    cudaMalloc((void**)&field.m, (sizeof(*field.m))*size_0*size_1*size_2*size_3*size_4);
    field.size[0] = size_0;
    field.size[1] = size_1;
    field.size[2] = size_2;
    field.size[3] = size_3;
    field.size[4] = size_4;
    field.is_device_field = 1;
    return field;
}

void free_device_field_intfield5(intfield5 field)
{
    cudaFree(field.m);
}

typedef struct intmat5
{
    int m[5];
} intmat5;

/* Adapted from: */

/* Z_2 lattice gauge simulation */

/* Michael Creutz <creutz@bnl.gov>     */

/* http://thy.phy.bnl.gov/~creutz/z2.c */





/* the lattice is of dimensions SIZE**4  */

/* @todo Move field decl to main */

__constant__ const int SIZE = 10;

typedef intfield5 Links; /* Last index is link direction */

intfield5 link;

__constant__ const int RAND_DATA_COUNT = 128;

/* Poor man's random generator */

__constant__ const float rand_data[128] = {
    0.765778,
    0.380508,
    0.976123,
    0.047972,
    0.027949,
    0.493132,
    0.145068,
    0.937659,
    0.688443,
    0.317046,
    0.803646,
    0.917738,
    0.513913,
    0.363706,
    0.137274,
    0.666660,
    0.250019,
    0.622242,
    0.021247,
    0.406825,
    0.707708,
    0.856293,
    0.947693,
    0.207796,
    0.362935,
    0.902242,
    0.427960,
    0.704711,
    0.613763,
    0.660261,
    0.378255,
    0.654958,
    0.936904,
    0.683342,
    0.891384,
    0.299881,
    0.064560,
    0.300503,
    0.572774,
    0.132678,
    0.132292,
    0.438706,
    0.594546,
    0.837315,
    0.180435,
    0.215016,
    0.726831,
    0.767127,
    0.556461,
    0.860724,
    0.132273,
    0.288679,
    0.001132,
    0.946316,
    0.740891,
    0.502307,
    0.189147,
    0.609733,
    0.716687,
    0.098146,
    0.650990,
    0.476326,
    0.958396,
    0.458836,
    0.834419,
    0.876043,
    0.820873,
    0.433127,
    0.800544,
    0.939788,
    0.741833,
    0.905454,
    0.796914,
    0.567545,
    0.054171,
    0.333496,
    0.247967,
    0.880176,
    0.760589,
    0.769755,
    0.011049,
    0.361483,
    0.829162,
    0.228125,
    0.572835,
    0.854979,
    0.070170,
    0.759810,
    0.022272,
    0.477994,
    0.014528,
    0.991334,
    0.314297,
    0.940028,
    0.235618,
    0.840691,
    0.882266,
    0.840194,
    0.985364,
    0.713334,
    0.697650,
    0.090573,
    0.262273,
    0.534600,
    0.761973,
    0.146971,
    0.667842,
    0.069159,
    0.102225,
    0.982492,
    0.933260,
    0.441284,
    0.149844,
    0.039490,
    0.520590,
    0.071531,
    0.141776,
    0.701622,
    0.213773,
    0.717888,
    0.621524,
    0.285984,
    0.442431,
    0.471437,
    0.197912,
    0.314655,
    0.496274,
    0.896794
};

__host__ __device__ void moveup(intmat5 *x, int d)
{
    x->m[1*d] += 1;

    if (x->m[1*d] >= SIZE) {
        x->m[1*d] -= SIZE;
    }
}

__host__ __device__ void movedown(intmat5 *x, int d)
{
    x->m[1*d] -= 1;

    if (x->m[1*d] < 0) {
        x->m[1*d] += SIZE;
    }
}
__global__ void kernel_0(intfield5 link)
{
    if (threadIdx.x + blockIdx.x*blockDim.x >= link.size[0]*link.size[1]*link.size[2]*link.size[3]*link.size[4]) {
        return;
    }
    intmat5 id;
    id.m[1*0] = (threadIdx.x + blockIdx.x*blockDim.x) % link.size[0]/1;
    id.m[1*1] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1])/link.size[0];
    id.m[1*2] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1]*link.size[2])/(link.size[0]*link.size[1]);
    id.m[1*3] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1]*link.size[2]*link.size[3])/(link.size[0]*link.size[1]*link.size[2]);
    id.m[1*4] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1]*link.size[2]*link.size[3]*link.size[4])/(link.size[0]*link.size[1]*link.size[2]*link.size[3]);
    link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*id.m[1*4]] = 1;
}


void coldstart()
{

    {
        dim3 dim_grid(link.size[0]*link.size[1]*link.size[2]*link.size[3]*link.size[4]/128 + 1, 1, 1);
        dim3 dim_block(128, 1, 1);
        kernel_0<<<dim_grid, dim_block>>>(link);
    }
}
__global__ void kernel_1(intfield5 link, double beta, int iter, float *cuda_action, int oddeven_phase)
{
    if (threadIdx.x + blockIdx.x*blockDim.x >= link.size[0]*link.size[1]*link.size[2]*link.size[3]*link.size[4]) {
        return;
    }
    intmat5 id;
    id.m[1*0] = (threadIdx.x + blockIdx.x*blockDim.x) % link.size[0]/1;
    id.m[1*1] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1])/link.size[0];
    id.m[1*2] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1]*link.size[2])/(link.size[0]*link.size[1]);
    id.m[1*3] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1]*link.size[2]*link.size[3])/(link.size[0]*link.size[1]*link.size[2]);
    id.m[1*4] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1]*link.size[2]*link.size[3]*link.size[4])/(link.size[0]*link.size[1]*link.size[2]*link.size[3]);
    if ((id.m[1*0] + id.m[1*1] + id.m[1*2] + id.m[1*3]) % 2 == oddeven_phase % 2 || id.m[1*4] != oddeven_phase/2) {
        return;
    }

    int dperp;

    float staplesum = 0;

    int staple;

    float bplus;

    float bminus;

    int d = id.m[1*4];

    for (dperp = 0; dperp < 4; dperp += 1) {

        if (dperp != d) {
            movedown(&id, dperp);

            int v1 = link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*dperp];

            int v2 = link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*d];
            staple = v1*v2;
            moveup(&id, d);
            staple *= link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*dperp];
            moveup(&id, dperp);
            staplesum += staple;
            staple = link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*dperp];
            moveup(&id, dperp);
            movedown(&id, d);
            staple *= link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*d];
            movedown(&id, dperp);
            staple *= link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*dperp];
            staplesum += staple;
        }
    }
    bplus = exp(beta*staplesum);
    bminus = 1/bplus;
    bplus = bplus/(bplus + bminus);

    int rand_ix = id.m[1*0] + id.m[1*1]*SIZE + id.m[1*3]*SIZE*SIZE + id.m[1*4]*SIZE*SIZE*SIZE + iter*SIZE*SIZE*SIZE*SIZE;

    if (rand_data[rand_ix % RAND_DATA_COUNT] < bplus) {
        link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*d] = 1;
        atomicAdd(cuda_action, staplesum);
    } else {
        link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*d] = -1;
        atomicAdd(cuda_action, -staplesum);
    }
}


double update(double beta, int iter)
{

    float action = 0.000000;

    {
        float *cuda_action = (float*)cuda_upload_var(&action, sizeof(action));
        dim3 dim_grid(link.size[0]*link.size[1]*link.size[2]*link.size[3]*link.size[4]/128 + 1, 1, 1);
        dim3 dim_block(128, 1, 1);
        for (int cuda_phase = 0; cuda_phase < 8; ++cuda_phase) {
            kernel_1<<<dim_grid, dim_block>>>(link, beta, iter, cuda_action, cuda_phase);
        }
        cuda_download_var(cuda_action, &action, sizeof(action));
    }
    action /= SIZE*SIZE*SIZE*SIZE*4*6;

    return 1.000000 - action;
}

int main()
{
    link = alloc_device_field_intfield5(SIZE, SIZE, SIZE, SIZE, 4);

    double beta;

    double action;

    double dbeta = 0.010000;
    coldstart();

    int iter = 0;

    for (beta = 1; beta > 0.000000; beta -= dbeta) {
        action = update(beta, iter);
        printf("%g\t%g\n", beta, action);

        ++iter;
    }
    printf("\n\n");

    for (beta = 0; beta < 1.000000; beta += dbeta) {
        action = update(beta, iter);
        printf("%g\t%g\n", beta, action);

        ++iter;
    }
    free_device_field_intfield5(link);

    return 0;
}

	.file	"tmpxft_002e809f_00000000-6_z2.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2040:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2040:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z15cuda_upload_varPvi
	.type	_Z15cuda_upload_varPvi, @function
_Z15cuda_upload_varPvi:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%esi, %ebx
	movl	$4, %esi
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %rdi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movslq	%ebx, %rdx
	movq	%rbp, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	(%rsp), %rax
	movq	8(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L3
	call	__stack_chk_fail@PLT
.L3:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z15cuda_upload_varPvi, .-_Z15cuda_upload_varPvi
	.globl	_Z17cuda_download_varPvS_i
	.type	_Z17cuda_download_varPvS_i, @function
_Z17cuda_download_varPvS_i:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	movslq	%edx, %rdx
	movq	%rsi, %rdi
	movl	$2, %ecx
	movq	%rbx, %rsi
	call	cudaMemcpy@PLT
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	cudaFree@PLT
	.cfi_endproc
.LFE2028:
	.size	_Z17cuda_download_varPvS_i, .-_Z17cuda_download_varPvS_i
	.globl	_Z22memcpy_field_intfield59intfield5S_
	.type	_Z22memcpy_field_intfield59intfield5S_, @function
_Z22memcpy_field_intfield59intfield5S_:
.LFB2029:
	.cfi_startproc
	endbr64
	movl	36(%rsp), %r11d
	movq	8(%rsp), %rdi
	movl	16(%rsp), %eax
	movslq	20(%rsp), %r9
	movslq	24(%rsp), %r10
	movslq	28(%rsp), %r8
	movslq	32(%rsp), %rcx
	movq	40(%rsp), %rsi
	movl	68(%rsp), %edx
	testl	%r11d, %r11d
	jne	.L9
	testl	%edx, %edx
	jne	.L10
	movslq	%eax, %rdx
	imulq	%r9, %rdx
	imulq	%r10, %rdx
	imulq	%r8, %rdx
	imulq	%rcx, %rdx
	xorl	%ecx, %ecx
	salq	$2, %rdx
	jmp	.L15
.L9:
	decl	%r11d
	jne	.L8
	testl	%edx, %edx
	jne	.L12
	movslq	%eax, %rdx
	imulq	%r9, %rdx
	imulq	%r10, %rdx
	imulq	%r8, %rdx
	imulq	%rcx, %rdx
	movl	$1, %ecx
	salq	$2, %rdx
	jmp	.L15
.L12:
	decl	%edx
	jne	.L8
	movslq	%eax, %rdx
	imulq	%r9, %rdx
	imulq	%r10, %rdx
	imulq	%r8, %rdx
	imulq	%rcx, %rdx
	movl	$3, %ecx
	salq	$2, %rdx
.L15:
	jmp	cudaMemcpy@PLT
.L10:
	decl	%edx
	jne	.L8
	movslq	%eax, %rdx
	imulq	%r9, %rdx
	imulq	%r10, %rdx
	imulq	%r8, %rdx
	imulq	%rcx, %rdx
	movl	$2, %ecx
	salq	$2, %rdx
	jmp	.L15
.L8:
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z22memcpy_field_intfield59intfield5S_, .-_Z22memcpy_field_intfield59intfield5S_
	.globl	_Z14size_intfield59intfield5i
	.type	_Z14size_intfield59intfield5i, @function
_Z14size_intfield59intfield5i:
.LFB2030:
	.cfi_startproc
	endbr64
	movslq	%edi, %rdi
	movl	16(%rsp,%rdi,4), %eax
	ret
	.cfi_endproc
.LFE2030:
	.size	_Z14size_intfield59intfield5i, .-_Z14size_intfield59intfield5i
	.globl	_Z28alloc_device_field_intfield5iiiii
	.type	_Z28alloc_device_field_intfield5iiiii, @function
_Z28alloc_device_field_intfield5iiiii:
.LFB2031:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%ecx, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%edx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%esi, %ebp
	movslq	%r8d, %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rsi, %r14
	movq	%rdi, %rbx
	pushq	%rax
	.cfi_def_cfa_offset 64
	movslq	%r9d, %rax
	imulq	%rax, %rsi
	movq	%rax, %r15
	movslq	%ecx, %rax
	imulq	%rax, %rsi
	movslq	%edx, %rax
	imulq	%rax, %rsi
	movslq	%ebp, %rax
	imulq	%rax, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	%ebp, 8(%rbx)
	movq	%rbx, %rax
	movl	%r12d, 12(%rbx)
	movl	%r13d, 16(%rbx)
	movl	%r14d, 20(%rbx)
	movl	%r15d, 24(%rbx)
	movl	$1, 28(%rbx)
	popq	%rdx
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	ret
	.cfi_endproc
.LFE2031:
	.size	_Z28alloc_device_field_intfield5iiiii, .-_Z28alloc_device_field_intfield5iiiii
	.globl	_Z27free_device_field_intfield59intfield5
	.type	_Z27free_device_field_intfield59intfield5, @function
_Z27free_device_field_intfield59intfield5:
.LFB2032:
	.cfi_startproc
	endbr64
	movq	8(%rsp), %rdi
	jmp	cudaFree@PLT
	.cfi_endproc
.LFE2032:
	.size	_Z27free_device_field_intfield59intfield5, .-_Z27free_device_field_intfield59intfield5
	.globl	_Z6moveupP7intmat5i
	.type	_Z6moveupP7intmat5i, @function
_Z6moveupP7intmat5i:
.LFB2033:
	.cfi_startproc
	endbr64
	movslq	%esi, %rsi
	movl	(%rdi,%rsi,4), %edx
	leal	1(%rdx), %eax
	cmpl	$9, %eax
	jle	.L21
	leal	-9(%rdx), %eax
.L21:
	movl	%eax, (%rdi,%rsi,4)
	ret
	.cfi_endproc
.LFE2033:
	.size	_Z6moveupP7intmat5i, .-_Z6moveupP7intmat5i
	.globl	_Z8movedownP7intmat5i
	.type	_Z8movedownP7intmat5i, @function
_Z8movedownP7intmat5i:
.LFB2034:
	.cfi_startproc
	endbr64
	movslq	%esi, %rsi
	movl	(%rdi,%rsi,4), %eax
	movl	%eax, %edx
	decl	%edx
	jns	.L23
	leal	9(%rax), %edx
.L23:
	movl	%edx, (%rdi,%rsi,4)
	ret
	.cfi_endproc
.LFE2034:
	.size	_Z8movedownP7intmat5i, .-_Z8movedownP7intmat5i
	.globl	_Z35__device_stub__Z8kernel_09intfield5R9intfield5
	.type	_Z35__device_stub__Z8kernel_09intfield5R9intfield5, @function
_Z35__device_stub__Z8kernel_09intfield5R9intfield5:
.LFB2062:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rcx
	movq	%rdi, 64(%rsp)
	movq	%rsp, %rdx
	movabsq	$4294967297, %rax
	leaq	28(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, 24(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 28(%rsp)
	movl	$1, 36(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L24
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z8kernel_09intfield5(%rip), %rdi
	pushq	8(%rsp)
	.cfi_def_cfa_offset 112
	movq	44(%rsp), %rcx
	movl	52(%rsp), %r8d
	movq	32(%rsp), %rsi
	movl	40(%rsp), %edx
	leaq	80(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L24:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L26
	call	__stack_chk_fail@PLT
.L26:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2062:
	.size	_Z35__device_stub__Z8kernel_09intfield5R9intfield5, .-_Z35__device_stub__Z8kernel_09intfield5R9intfield5
	.globl	_Z8kernel_09intfield5
	.type	_Z8kernel_09intfield5, @function
_Z8kernel_09intfield5:
.LFB2063:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	16(%rsp), %rdi
	call	_Z35__device_stub__Z8kernel_09intfield5R9intfield5
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2063:
	.size	_Z8kernel_09intfield5, .-_Z8kernel_09intfield5
	.globl	_Z9coldstartv
	.type	_Z9coldstartv, @function
_Z9coldstartv:
.LFB2035:
	.cfi_startproc
	endbr64
	subq	$72, %rsp
	.cfi_def_cfa_offset 80
	movl	$128, %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	movl	8+link(%rip), %eax
	imull	12+link(%rip), %eax
	imull	16+link(%rip), %eax
	imull	20+link(%rip), %eax
	imull	24+link(%rip), %eax
	cltd
	idivl	%ecx
	movl	$33554433, %edx
	movl	$1, %ecx
	salq	$7, %rdx
	incl	%eax
	movl	%eax, (%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 4(%rsp)
	movq	(%rsp), %rdi
	movl	8(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L30
	movaps	link(%rip), %xmm0
	movaps	16+link(%rip), %xmm1
	leaq	24(%rsp), %rdi
	movups	%xmm0, 24(%rsp)
	movups	%xmm1, 40(%rsp)
	call	_Z35__device_stub__Z8kernel_09intfield5R9intfield5
.L30:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L32
	call	__stack_chk_fail@PLT
.L32:
	addq	$72, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2035:
	.size	_Z9coldstartv, .-_Z9coldstartv
	.globl	_Z40__device_stub__Z8kernel_19intfield5diPfiR9intfield5diPfi
	.type	_Z40__device_stub__Z8kernel_19intfield5diPfiR9intfield5diPfi, @function
_Z40__device_stub__Z8kernel_19intfield5diPfiR9intfield5diPfi:
.LFB2064:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%esi, 20(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 16(%rsp)
	leaq	40(%rsp), %rcx
	movsd	%xmm0, 24(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movq	%rdi, 96(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rax, 104(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movl	$1, 56(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L34
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z8kernel_19intfield5diPfi(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L34:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L36
	call	__stack_chk_fail@PLT
.L36:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2064:
	.size	_Z40__device_stub__Z8kernel_19intfield5diPfiR9intfield5diPfi, .-_Z40__device_stub__Z8kernel_19intfield5diPfiR9intfield5diPfi
	.globl	_Z8kernel_19intfield5diPfi
	.type	_Z8kernel_19intfield5diPfi, @function
_Z8kernel_19intfield5diPfi:
.LFB2065:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	movl	%edx, %ecx
	movq	%rsi, %rdx
	movl	%edi, %esi
	leaq	16(%rsp), %rdi
	call	_Z40__device_stub__Z8kernel_19intfield5diPfiR9intfield5diPfi
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2065:
	.size	_Z8kernel_19intfield5diPfi, .-_Z8kernel_19intfield5diPfi
	.globl	_Z6updatedi
	.type	_Z6updatedi, @function
_Z6updatedi:
.LFB2036:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$4, %esi
	movq	%xmm0, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	xorl	%r13d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%edi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %r12
	movl	$0x00000000, 12(%rsp)
	leaq	40(%rsp), %r14
	movq	%r12, %rdi
	call	_Z15cuda_upload_varPvi
	movl	$128, %ecx
	movl	$1, 36(%rsp)
	movq	%rax, %rbx
	movl	8+link(%rip), %eax
	imull	12+link(%rip), %eax
	imull	16+link(%rip), %eax
	imull	20+link(%rip), %eax
	imull	24+link(%rip), %eax
	cltd
	idivl	%ecx
	incl	%eax
	movl	%eax, 16(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 20(%rsp)
	addq	$127, %rax
	movq	%rax, 28(%rsp)
.L42:
	movl	36(%rsp), %ecx
	movq	28(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	16(%rsp), %rdi
	movl	24(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movaps	link(%rip), %xmm2
	movl	%r13d, %ecx
	movq	%rbx, %rdx
	movl	%ebp, %esi
	movaps	16+link(%rip), %xmm3
	movq	%r15, %xmm0
	movq	%r14, %rdi
	movups	%xmm2, 40(%rsp)
	movups	%xmm3, 56(%rsp)
	call	_Z40__device_stub__Z8kernel_19intfield5diPfiR9intfield5diPfi
.L41:
	incl	%r13d
	cmpl	$8, %r13d
	jne	.L42
	movl	$4, %edx
	movq	%r12, %rsi
	movq	%rbx, %rdi
	call	_Z17cuda_download_varPvS_i
	movsd	.LC2(%rip), %xmm0
	movss	12(%rsp), %xmm1
	divss	.LC1(%rip), %xmm1
	cvtss2sd	%xmm1, %xmm1
	subsd	%xmm1, %xmm0
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L43
	call	__stack_chk_fail@PLT
.L43:
	addq	$88, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	ret
	.cfi_endproc
.LFE2036:
	.size	_Z6updatedi, .-_Z6updatedi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC4:
	.string	"%g\t%g\n"
.LC6:
	.string	"\n\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2037:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movl	$160000, %esi
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	leaq	.LC4(%rip), %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	xorl	%ebx, %ebx
	subq	$48, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	call	cudaMalloc@PLT
	movabsq	$42949672970, %rax
	movq	%rax, 16(%rsp)
	movups	8(%rsp), %xmm4
	movq	%rax, 24(%rsp)
	movl	$1073741825, %eax
	salq	$2, %rax
	movaps	%xmm4, link(%rip)
	movq	%rax, 32(%rsp)
	movups	24(%rsp), %xmm5
	movaps	%xmm5, 16+link(%rip)
	call	_Z9coldstartv
	movsd	.LC2(%rip), %xmm6
	movq	%xmm6, %r14
.L47:
	movl	%ebx, %edi
	movq	%r14, %xmm0
	incl	%ebx
	call	_Z6updatedi
	movq	%rbp, %rsi
	movl	$2, %edi
	movb	$2, %al
	movaps	%xmm0, %xmm1
	movq	%r14, %xmm0
	call	__printf_chk@PLT
	movq	%r14, %xmm2
	subsd	.LC5(%rip), %xmm2
	movq	%xmm2, %r14
	cmpl	$100, %ebx
	jne	.L47
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%r14d, %r14d
	call	__printf_chk@PLT
.L48:
	movl	%ebx, %edi
	movq	%r14, %xmm0
	incl	%ebx
	call	_Z6updatedi
	movq	%rbp, %rsi
	movl	$2, %edi
	movb	$2, %al
	movaps	%xmm0, %xmm1
	movq	%r14, %xmm0
	call	__printf_chk@PLT
	movsd	.LC5(%rip), %xmm3
	movq	%r14, %xmm7
	addsd	%xmm7, %xmm3
	movq	%xmm3, %r14
	cmpl	$200, %ebx
	jne	.L48
	movq	link(%rip), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L49
	call	__stack_chk_fail@PLT
.L49:
	addq	$48, %rsp
	.cfi_def_cfa_offset 32
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2037:
	.size	main, .-main
	.section	.rodata.str1.1
.LC7:
	.string	"_Z8kernel_19intfield5diPfi"
.LC8:
	.string	"_Z8kernel_09intfield5"
.LC9:
	.string	"SIZE"
.LC10:
	.string	"RAND_DATA_COUNT"
.LC11:
	.string	"rand_data"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2067:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -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
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z8kernel_19intfield5diPfi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC8(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z8kernel_09intfield5(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC9(%rip), %rdx
	movl	$4, %r9d
	leaq	_ZL4SIZE(%rip), %rsi
	pushq	$1
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC10(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$4, %r9d
	leaq	_ZL15RAND_DATA_COUNT(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC11(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$512, %r9d
	movq	%rdx, %rcx
	leaq	_ZL9rand_data(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%r8
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2067:
	.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
	.align 32
	.type	_ZL9rand_data, @object
	.size	_ZL9rand_data, 512
_ZL9rand_data:
	.long	1061423623
	.long	1052955122
	.long	1064952626
	.long	1027898954
	.long	1021637965
	.long	1056734156
	.long	1041534133
	.long	1064307308
	.long	1060126157
	.long	1050825690
	.long	1062058943
	.long	1063973089
	.long	1057198029
	.long	1052391340
	.long	1041011086
	.long	1059760699
	.long	1048576638
	.long	1059015488
	.long	1018039856
	.long	1053838174
	.long	1060449370
	.long	1062942213
	.long	1064475650
	.long	1045743738
	.long	1052365470
	.long	1063713109
	.long	1054547347
	.long	1060399089
	.long	1058873234
	.long	1059653341
	.long	1052879524
	.long	1059564372
	.long	1064294641
	.long	1060040576
	.long	1063530942
	.long	1050249729
	.long	1032075273
	.long	1050270599
	.long	1058185553
	.long	1040702654
	.long	1040676750
	.long	1054907923
	.long	1058550827
	.long	1062623815
	.long	1043907572
	.long	1046228264
	.long	1060770201
	.long	1061446255
	.long	1057911866
	.long	1063016552
	.long	1040675475
	.long	1049873852
	.long	982802334
	.long	1064452548
	.long	1061006088
	.long	1057003313
	.long	1044492224
	.long	1058805622
	.long	1060600013
	.long	1036583109
	.long	1059497800
	.long	1056170240
	.long	1064655217
	.long	1055583373
	.long	1062575228
	.long	1063273563
	.long	1062347964
	.long	1054720722
	.long	1062006900
	.long	1064343026
	.long	1061021892
	.long	1063766997
	.long	1061945998
	.long	1058097825
	.long	1029562985
	.long	1051377661
	.long	1048439568
	.long	1063342903
	.long	1061336566
	.long	1061490346
	.long	1010108125
	.long	1052316749
	.long	1062487030
	.long	1047107994
	.long	1058186577
	.long	1062920167
	.long	1032828234
	.long	1061323496
	.long	1018590149
	.long	1056226209
	.long	1013843673
	.long	1065207825
	.long	1050733449
	.long	1064347053
	.long	1047610840
	.long	1062680454
	.long	1063377967
	.long	1062672116
	.long	1065107665
	.long	1060543759
	.long	1060280625
	.long	1035566678
	.long	1048987814
	.long	1057545100
	.long	1061359786
	.long	1041661841
	.long	1059780529
	.long	1032692540
	.long	1037130583
	.long	1065059481
	.long	1064233505
	.long	1054994426
	.long	1041854645
	.long	1025622084
	.long	1057310051
	.long	1033010904
	.long	1041313210
	.long	1060347264
	.long	1046144847
	.long	1060620162
	.long	1059003442
	.long	1049783423
	.long	1055032913
	.long	1056006193
	.long	1045080433
	.long	1050745462
	.long	1056839584
	.long	1063621707
	.align 4
	.type	_ZL15RAND_DATA_COUNT, @object
	.size	_ZL15RAND_DATA_COUNT, 4
_ZL15RAND_DATA_COUNT:
	.long	128
	.globl	link
	.bss
	.align 32
	.type	link, @object
	.size	link, 32
link:
	.zero	32
	.section	.rodata
	.align 4
	.type	_ZL4SIZE, @object
	.size	_ZL4SIZE, 4
_ZL4SIZE:
	.long	10
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC1:
	.long	1214930944
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC2:
	.long	0
	.long	1072693248
	.align 8
.LC5:
	.long	1202590843
	.long	1065646817
	.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 <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h> /* memcpy */
#include <math.h>
#include <stdint.h>

void *cuda_upload_var(void *host_var, int size)
{
	void *cuda_var;
	hipMalloc(&cuda_var, 4);
	hipMemcpy(cuda_var, host_var, size, hipMemcpyHostToDevice);
	return cuda_var;
}
void cuda_download_var(void *cuda_var, void *host_var, int size)
{
	hipMemcpy(host_var, cuda_var, size, hipMemcpyDeviceToHost);
	hipFree(cuda_var);
}

typedef struct intfield5
{
    int *m;
    int size[5];
    int is_device_field;
} intfield5;

void memcpy_field_intfield5(intfield5 dst, intfield5 src)
{
    if (dst.is_device_field == 0 && src.is_device_field == 0) {
        hipMemcpy(dst.m, src.m, (sizeof(*dst.m))*dst.size[0]*dst.size[1]*dst.size[2]*dst.size[3]*dst.size[4], hipMemcpyHostToHost);
    }
    if (dst.is_device_field == 1 && src.is_device_field == 0) {
        hipMemcpy(dst.m, src.m, (sizeof(*dst.m))*dst.size[0]*dst.size[1]*dst.size[2]*dst.size[3]*dst.size[4], hipMemcpyHostToDevice);
    }
    if (dst.is_device_field == 0 && src.is_device_field == 1) {
        hipMemcpy(dst.m, src.m, (sizeof(*dst.m))*dst.size[0]*dst.size[1]*dst.size[2]*dst.size[3]*dst.size[4], hipMemcpyDeviceToHost);
    }
    if (dst.is_device_field == 1 && src.is_device_field == 1) {
        hipMemcpy(dst.m, src.m, (sizeof(*dst.m))*dst.size[0]*dst.size[1]*dst.size[2]*dst.size[3]*dst.size[4], hipMemcpyDeviceToDevice);
    }
}

int size_intfield5(intfield5 field, int index)
{
    return field.size[index];
}

intfield5 alloc_device_field_intfield5(int size_0, int size_1, int size_2, int size_3, int size_4)
{
    intfield5 field;
    hipMalloc((void**)&field.m, (sizeof(*field.m))*size_0*size_1*size_2*size_3*size_4);
    field.size[0] = size_0;
    field.size[1] = size_1;
    field.size[2] = size_2;
    field.size[3] = size_3;
    field.size[4] = size_4;
    field.is_device_field = 1;
    return field;
}

void free_device_field_intfield5(intfield5 field)
{
    hipFree(field.m);
}

typedef struct intmat5
{
    int m[5];
} intmat5;

/* Adapted from: */

/* Z_2 lattice gauge simulation */

/* Michael Creutz <creutz@bnl.gov>     */

/* http://thy.phy.bnl.gov/~creutz/z2.c */





/* the lattice is of dimensions SIZE**4  */

/* @todo Move field decl to main */

__constant__ const int SIZE = 10;

typedef intfield5 Links; /* Last index is link direction */

intfield5 link;

__constant__ const int RAND_DATA_COUNT = 128;

/* Poor man's random generator */

__constant__ const float rand_data[128] = {
    0.765778,
    0.380508,
    0.976123,
    0.047972,
    0.027949,
    0.493132,
    0.145068,
    0.937659,
    0.688443,
    0.317046,
    0.803646,
    0.917738,
    0.513913,
    0.363706,
    0.137274,
    0.666660,
    0.250019,
    0.622242,
    0.021247,
    0.406825,
    0.707708,
    0.856293,
    0.947693,
    0.207796,
    0.362935,
    0.902242,
    0.427960,
    0.704711,
    0.613763,
    0.660261,
    0.378255,
    0.654958,
    0.936904,
    0.683342,
    0.891384,
    0.299881,
    0.064560,
    0.300503,
    0.572774,
    0.132678,
    0.132292,
    0.438706,
    0.594546,
    0.837315,
    0.180435,
    0.215016,
    0.726831,
    0.767127,
    0.556461,
    0.860724,
    0.132273,
    0.288679,
    0.001132,
    0.946316,
    0.740891,
    0.502307,
    0.189147,
    0.609733,
    0.716687,
    0.098146,
    0.650990,
    0.476326,
    0.958396,
    0.458836,
    0.834419,
    0.876043,
    0.820873,
    0.433127,
    0.800544,
    0.939788,
    0.741833,
    0.905454,
    0.796914,
    0.567545,
    0.054171,
    0.333496,
    0.247967,
    0.880176,
    0.760589,
    0.769755,
    0.011049,
    0.361483,
    0.829162,
    0.228125,
    0.572835,
    0.854979,
    0.070170,
    0.759810,
    0.022272,
    0.477994,
    0.014528,
    0.991334,
    0.314297,
    0.940028,
    0.235618,
    0.840691,
    0.882266,
    0.840194,
    0.985364,
    0.713334,
    0.697650,
    0.090573,
    0.262273,
    0.534600,
    0.761973,
    0.146971,
    0.667842,
    0.069159,
    0.102225,
    0.982492,
    0.933260,
    0.441284,
    0.149844,
    0.039490,
    0.520590,
    0.071531,
    0.141776,
    0.701622,
    0.213773,
    0.717888,
    0.621524,
    0.285984,
    0.442431,
    0.471437,
    0.197912,
    0.314655,
    0.496274,
    0.896794
};

__host__ __device__ void moveup(intmat5 *x, int d)
{
    x->m[1*d] += 1;

    if (x->m[1*d] >= SIZE) {
        x->m[1*d] -= SIZE;
    }
}

__host__ __device__ void movedown(intmat5 *x, int d)
{
    x->m[1*d] -= 1;

    if (x->m[1*d] < 0) {
        x->m[1*d] += SIZE;
    }
}
__global__ void kernel_0(intfield5 link)
{
    if (threadIdx.x + blockIdx.x*blockDim.x >= link.size[0]*link.size[1]*link.size[2]*link.size[3]*link.size[4]) {
        return;
    }
    intmat5 id;
    id.m[1*0] = (threadIdx.x + blockIdx.x*blockDim.x) % link.size[0]/1;
    id.m[1*1] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1])/link.size[0];
    id.m[1*2] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1]*link.size[2])/(link.size[0]*link.size[1]);
    id.m[1*3] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1]*link.size[2]*link.size[3])/(link.size[0]*link.size[1]*link.size[2]);
    id.m[1*4] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1]*link.size[2]*link.size[3]*link.size[4])/(link.size[0]*link.size[1]*link.size[2]*link.size[3]);
    link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*id.m[1*4]] = 1;
}


void coldstart()
{

    {
        dim3 dim_grid(link.size[0]*link.size[1]*link.size[2]*link.size[3]*link.size[4]/128 + 1, 1, 1);
        dim3 dim_block(128, 1, 1);
        kernel_0<<<dim_grid, dim_block>>>(link);
    }
}
__global__ void kernel_1(intfield5 link, double beta, int iter, float *cuda_action, int oddeven_phase)
{
    if (threadIdx.x + blockIdx.x*blockDim.x >= link.size[0]*link.size[1]*link.size[2]*link.size[3]*link.size[4]) {
        return;
    }
    intmat5 id;
    id.m[1*0] = (threadIdx.x + blockIdx.x*blockDim.x) % link.size[0]/1;
    id.m[1*1] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1])/link.size[0];
    id.m[1*2] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1]*link.size[2])/(link.size[0]*link.size[1]);
    id.m[1*3] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1]*link.size[2]*link.size[3])/(link.size[0]*link.size[1]*link.size[2]);
    id.m[1*4] = (threadIdx.x + blockIdx.x*blockDim.x) % (link.size[0]*link.size[1]*link.size[2]*link.size[3]*link.size[4])/(link.size[0]*link.size[1]*link.size[2]*link.size[3]);
    if ((id.m[1*0] + id.m[1*1] + id.m[1*2] + id.m[1*3]) % 2 == oddeven_phase % 2 || id.m[1*4] != oddeven_phase/2) {
        return;
    }

    int dperp;

    float staplesum = 0;

    int staple;

    float bplus;

    float bminus;

    int d = id.m[1*4];

    for (dperp = 0; dperp < 4; dperp += 1) {

        if (dperp != d) {
            movedown(&id, dperp);

            int v1 = link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*dperp];

            int v2 = link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*d];
            staple = v1*v2;
            moveup(&id, d);
            staple *= link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*dperp];
            moveup(&id, dperp);
            staplesum += staple;
            staple = link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*dperp];
            moveup(&id, dperp);
            movedown(&id, d);
            staple *= link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*d];
            movedown(&id, dperp);
            staple *= link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*dperp];
            staplesum += staple;
        }
    }
    bplus = exp(beta*staplesum);
    bminus = 1/bplus;
    bplus = bplus/(bplus + bminus);

    int rand_ix = id.m[1*0] + id.m[1*1]*SIZE + id.m[1*3]*SIZE*SIZE + id.m[1*4]*SIZE*SIZE*SIZE + iter*SIZE*SIZE*SIZE*SIZE;

    if (rand_data[rand_ix % RAND_DATA_COUNT] < bplus) {
        link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*d] = 1;
        atomicAdd(cuda_action, staplesum);
    } else {
        link.m[link.size[1]*link.size[2]*link.size[3]*link.size[4]*id.m[1*0] + link.size[2]*link.size[3]*link.size[4]*id.m[1*1] + link.size[3]*link.size[4]*id.m[1*2] + link.size[4]*id.m[1*3] + 1*d] = -1;
        atomicAdd(cuda_action, -staplesum);
    }
}


double update(double beta, int iter)
{

    float action = 0.000000;

    {
        float *cuda_action = (float*)cuda_upload_var(&action, sizeof(action));
        dim3 dim_grid(link.size[0]*link.size[1]*link.size[2]*link.size[3]*link.size[4]/128 + 1, 1, 1);
        dim3 dim_block(128, 1, 1);
        for (int cuda_phase = 0; cuda_phase < 8; ++cuda_phase) {
            kernel_1<<<dim_grid, dim_block>>>(link, beta, iter, cuda_action, cuda_phase);
        }
        cuda_download_var(cuda_action, &action, sizeof(action));
    }
    action /= SIZE*SIZE*SIZE*SIZE*4*6;

    return 1.000000 - action;
}

int main()
{
    link = alloc_device_field_intfield5(SIZE, SIZE, SIZE, SIZE, 4);

    double beta;

    double action;

    double dbeta = 0.010000;
    coldstart();

    int iter = 0;

    for (beta = 1; beta > 0.000000; beta -= dbeta) {
        action = update(beta, iter);
        printf("%g\t%g\n", beta, action);

        ++iter;
    }
    printf("\n\n");

    for (beta = 0; beta < 1.000000; beta += dbeta) {
        action = update(beta, iter);
        printf("%g\t%g\n", beta, action);

        ++iter;
    }
    free_device_field_intfield5(link);

    return 0;
}

	.text
	.file	"z2.hip"
	.globl	_Z15cuda_upload_varPvi          # -- Begin function _Z15cuda_upload_varPvi
	.type	_Z15cuda_upload_varPvi,@function
_Z15cuda_upload_varPvi:                 # @_Z15cuda_upload_varPvi
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$16, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	%esi, %ebx
	movq	%rdi, %r14
	leaq	8(%rsp), %r15
	movl	$4, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movslq	%ebx, %rdx
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r15), %rax
	addq	$16, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z15cuda_upload_varPvi, .Lfunc_end0-_Z15cuda_upload_varPvi
	.cfi_endproc
                                        # -- End function
	.globl	_Z17cuda_download_varPvS_i      # -- Begin function _Z17cuda_download_varPvS_i
	.type	_Z17cuda_download_varPvS_i,@function
_Z17cuda_download_varPvS_i:             # @_Z17cuda_download_varPvS_i
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	%rdi, %rbx
	movslq	%edx, %rdx
	movq	%rsi, %rdi
	movq	%rbx, %rsi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	hipFree                         # TAILCALL
.Lfunc_end1:
	.size	_Z17cuda_download_varPvS_i, .Lfunc_end1-_Z17cuda_download_varPvS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z22memcpy_field_intfield59intfield5S_ # -- Begin function _Z22memcpy_field_intfield59intfield5S_
	.type	_Z22memcpy_field_intfield59intfield5S_,@function
_Z22memcpy_field_intfield59intfield5S_: # @_Z22memcpy_field_intfield59intfield5S_
	.cfi_startproc
# %bb.0:
	movl	36(%rsp), %edx
	movl	68(%rsp), %eax
	xorl	%ecx, %ecx
	movl	%edx, %esi
	orl	%eax, %esi
	je	.LBB2_8
# %bb.1:
	cmpl	$1, %edx
	jne	.LBB2_3
# %bb.2:
	movl	$1, %ecx
	testl	%eax, %eax
	je	.LBB2_8
.LBB2_3:
	testl	%edx, %edx
	jne	.LBB2_5
# %bb.4:
	movl	$2, %ecx
	cmpl	$1, %eax
	je	.LBB2_8
.LBB2_5:                                # %.thread13
	cmpl	$1, %edx
	jne	.LBB2_7
# %bb.6:                                # %.thread13
	movl	$3, %ecx
	cmpl	$1, %eax
	jne	.LBB2_7
.LBB2_8:                                # %.thread13.thread.sink.split
	leaq	40(%rsp), %rax
	leaq	8(%rsp), %rdx
	movq	(%rdx), %rdi
	movq	(%rax), %rsi
	movslq	8(%rdx), %rax
	movslq	12(%rdx), %r8
	imulq	%rax, %r8
	movslq	16(%rdx), %rax
	movslq	20(%rdx), %r9
	imulq	%rax, %r9
	imulq	%r8, %r9
	movslq	24(%rdx), %rdx
	imulq	%r9, %rdx
	shlq	$2, %rdx
	jmp	hipMemcpy                       # TAILCALL
.LBB2_7:                                # %.thread13.thread
	retq
.Lfunc_end2:
	.size	_Z22memcpy_field_intfield59intfield5S_, .Lfunc_end2-_Z22memcpy_field_intfield59intfield5S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z14size_intfield59intfield5i   # -- Begin function _Z14size_intfield59intfield5i
	.type	_Z14size_intfield59intfield5i,@function
_Z14size_intfield59intfield5i:          # @_Z14size_intfield59intfield5i
	.cfi_startproc
# %bb.0:
	movslq	%edi, %rax
	movl	16(%rsp,%rax,4), %eax
	retq
.Lfunc_end3:
	.size	_Z14size_intfield59intfield5i, .Lfunc_end3-_Z14size_intfield59intfield5i
	.cfi_endproc
                                        # -- End function
	.globl	_Z28alloc_device_field_intfield5iiiii # -- Begin function _Z28alloc_device_field_intfield5iiiii
	.type	_Z28alloc_device_field_intfield5iiiii,@function
_Z28alloc_device_field_intfield5iiiii:  # @_Z28alloc_device_field_intfield5iiiii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdi, %rbx
	movslq	%esi, %r14
	movslq	%edx, %r15
	movq	%r14, %rax
	imulq	%r15, %rax
	movslq	%ecx, %r12
	movslq	%r8d, %r13
	movq	%r12, %rsi
	imulq	%r13, %rsi
	imulq	%rax, %rsi
	movslq	%r9d, %rbp
	imulq	%rbp, %rsi
	shlq	$2, %rsi
	callq	hipMalloc
	movl	%r14d, 8(%rbx)
	movl	%r15d, 12(%rbx)
	movl	%r12d, 16(%rbx)
	movl	%r13d, 20(%rbx)
	movl	%ebp, 24(%rbx)
	movl	$1, 28(%rbx)
	movq	%rbx, %rax
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end4:
	.size	_Z28alloc_device_field_intfield5iiiii, .Lfunc_end4-_Z28alloc_device_field_intfield5iiiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z27free_device_field_intfield59intfield5 # -- Begin function _Z27free_device_field_intfield59intfield5
	.type	_Z27free_device_field_intfield59intfield5,@function
_Z27free_device_field_intfield59intfield5: # @_Z27free_device_field_intfield59intfield5
	.cfi_startproc
# %bb.0:
	movq	8(%rsp), %rdi
	jmp	hipFree                         # TAILCALL
.Lfunc_end5:
	.size	_Z27free_device_field_intfield59intfield5, .Lfunc_end5-_Z27free_device_field_intfield59intfield5
	.cfi_endproc
                                        # -- End function
	.globl	_Z6moveupP7intmat5i             # -- Begin function _Z6moveupP7intmat5i
	.type	_Z6moveupP7intmat5i,@function
_Z6moveupP7intmat5i:                    # @_Z6moveupP7intmat5i
	.cfi_startproc
# %bb.0:
	movslq	%esi, %rax
	movl	(%rdi,%rax,4), %ecx
	leal	1(%rcx), %edx
	leal	-9(%rcx), %esi
	cmpl	$9, %ecx
	cmovll	%edx, %esi
	movl	%esi, (%rdi,%rax,4)
	retq
.Lfunc_end6:
	.size	_Z6moveupP7intmat5i, .Lfunc_end6-_Z6moveupP7intmat5i
	.cfi_endproc
                                        # -- End function
	.globl	_Z8movedownP7intmat5i           # -- Begin function _Z8movedownP7intmat5i
	.type	_Z8movedownP7intmat5i,@function
_Z8movedownP7intmat5i:                  # @_Z8movedownP7intmat5i
	.cfi_startproc
# %bb.0:
	movslq	%esi, %rax
	movl	(%rdi,%rax,4), %ecx
	xorl	%edx, %edx
	testl	%ecx, %ecx
	setg	%dl
	negl	%edx
	orl	$9, %edx
	addl	%ecx, %edx
	movl	%edx, (%rdi,%rax,4)
	retq
.Lfunc_end7:
	.size	_Z8movedownP7intmat5i, .Lfunc_end7-_Z8movedownP7intmat5i
	.cfi_endproc
                                        # -- End function
	.globl	_Z23__device_stub__kernel_09intfield5 # -- Begin function _Z23__device_stub__kernel_09intfield5
	.type	_Z23__device_stub__kernel_09intfield5,@function
_Z23__device_stub__kernel_09intfield5:  # @_Z23__device_stub__kernel_09intfield5
	.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	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	112(%rsp), %rax
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z8kernel_09intfield5, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	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_end8:
	.size	_Z23__device_stub__kernel_09intfield5, .Lfunc_end8-_Z23__device_stub__kernel_09intfield5
	.cfi_endproc
                                        # -- End function
	.globl	_Z9coldstartv                   # -- Begin function _Z9coldstartv
	.type	_Z9coldstartv,@function
_Z9coldstartv:                          # @_Z9coldstartv
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movl	link+12(%rip), %eax
	imull	link+8(%rip), %eax
	imull	link+16(%rip), %eax
	imull	link+20(%rip), %eax
	imull	link+24(%rip), %eax
	leal	127(%rax), %edi
	testl	%eax, %eax
	cmovnsl	%eax, %edi
	sarl	$7, %edi
	incl	%edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$128, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB9_2
# %bb.1:
	movups	link+16(%rip), %xmm0
	movups	%xmm0, 16(%rsp)
	movups	link(%rip), %xmm0
	movups	%xmm0, (%rsp)
	callq	_Z23__device_stub__kernel_09intfield5
.LBB9_2:
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end9:
	.size	_Z9coldstartv, .Lfunc_end9-_Z9coldstartv
	.cfi_endproc
                                        # -- End function
	.globl	_Z23__device_stub__kernel_19intfield5diPfi # -- Begin function _Z23__device_stub__kernel_19intfield5diPfi
	.type	_Z23__device_stub__kernel_19intfield5diPfi,@function
_Z23__device_stub__kernel_19intfield5diPfi: # @_Z23__device_stub__kernel_19intfield5diPfi
	.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	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movsd	%xmm0, (%rax)
	leaq	12(%rsp), %rcx
	movl	%edi, (%rcx)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	8(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	176(%rsp), %rdx
	leaq	80(%rsp), %rbx
	movq	%rdx, (%rbx)
	movq	%rax, 8(%rbx)
	movq	%rcx, 16(%rbx)
	movq	%rdi, 24(%rbx)
	movq	%rsi, 32(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z8kernel_19intfield5diPfi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	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_end10:
	.size	_Z23__device_stub__kernel_19intfield5diPfi, .Lfunc_end10-_Z23__device_stub__kernel_19intfield5diPfi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z6updatedi
.LCPI11_0:
	.long	0xc86a6000                      # float -2.4E+5
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI11_1:
	.quad	0x3ff0000000000000              # double 1
	.text
	.globl	_Z6updatedi
	.type	_Z6updatedi,@function
_Z6updatedi:                            # @_Z6updatedi
	.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	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	%edi, %ebx
	movsd	%xmm0, 48(%rsp)                 # 8-byte Spill
	movabsq	$4294967296, %r14               # imm = 0x100000000
	leaq	44(%rsp), %r15
	movl	$0, (%r15)
	leaq	56(%rsp), %r12
	movl	$4, %esi
	movq	%r12, %rdi
	callq	hipMalloc
	movq	(%r12), %rdi
	movl	$4, %edx
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	link+12(%rip), %eax
	imull	link+8(%rip), %eax
	imull	link+16(%rip), %eax
	imull	link+20(%rip), %eax
	movq	(%r12), %r15
	imull	link+24(%rip), %eax
	leal	127(%rax), %r12d
	testl	%eax, %eax
	cmovnsl	%eax, %r12d
	sarl	$7, %r12d
	incl	%r12d
	btsq	$32, %r12
	xorl	%r13d, %r13d
	subq	$-128, %r14
.LBB11_1:                               # =>This Inner Loop Header: Depth=1
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r14, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB11_3
# %bb.2:                                #   in Loop: Header=BB11_1 Depth=1
	movups	link+16(%rip), %xmm0
	movups	%xmm0, 16(%rsp)
	movups	link(%rip), %xmm0
	movups	%xmm0, (%rsp)
	movsd	48(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movl	%ebx, %edi
	movq	%r15, %rsi
	movl	%r13d, %edx
	callq	_Z23__device_stub__kernel_19intfield5diPfi
.LBB11_3:                               #   in Loop: Header=BB11_1 Depth=1
	incl	%r13d
	cmpl	$8, %r13d
	jne	.LBB11_1
# %bb.4:
	leaq	44(%rsp), %rbx
	movl	$4, %edx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	%r15, %rdi
	callq	hipFree
	movss	(%rbx), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	divss	.LCPI11_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	addsd	.LCPI11_1(%rip), %xmm0
	addq	$64, %rsp
	.cfi_def_cfa_offset 48
	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_end11:
	.size	_Z6updatedi, .Lfunc_end11-_Z6updatedi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI12_0:
	.long	10                              # 0xa
	.long	10                              # 0xa
	.long	4                               # 0x4
	.long	1                               # 0x1
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI12_1:
	.quad	0x3ff0000000000000              # double 1
.LCPI12_2:
	.quad	0xbf847ae147ae147b              # double -0.01
.LCPI12_3:
	.quad	0x3f847ae147ae147b              # double 0.01
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$48, %rsp
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -16
	leaq	16(%rsp), %rbx
	movl	$160000, %esi                   # imm = 0x27100
	movq	%rbx, %rdi
	callq	hipMalloc
	movabsq	$42949672970, %rax              # imm = 0xA0000000A
	movq	%rax, 8(%rbx)
	movaps	.LCPI12_0(%rip), %xmm0          # xmm0 = [10,10,4,1]
	movups	%xmm0, 16(%rbx)
	movups	(%rbx), %xmm0
	movups	%xmm0, link(%rip)
	movups	16(%rbx), %xmm0
	movups	%xmm0, link+16(%rip)
	callq	_Z9coldstartv
	movsd	.LCPI12_1(%rip), %xmm0          # xmm0 = mem[0],zero
	xorl	%ebx, %ebx
.LBB12_1:                               # =>This Inner Loop Header: Depth=1
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movl	%ebx, %edi
	callq	_Z6updatedi
	movaps	%xmm0, %xmm1
	movl	$.L.str, %edi
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movb	$2, %al
	callq	printf
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	incl	%ebx
	addsd	.LCPI12_2(%rip), %xmm0
	cmpl	$100, %ebx
	jne	.LBB12_1
# %bb.2:
	movl	$.Lstr, %edi
	callq	puts@PLT
	xorpd	%xmm0, %xmm0
	movl	$100, %ebx
.LBB12_3:                               # =>This Inner Loop Header: Depth=1
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movl	%ebx, %edi
	callq	_Z6updatedi
	movaps	%xmm0, %xmm1
	movl	$.L.str, %edi
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movb	$2, %al
	callq	printf
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	incl	%ebx
	addsd	.LCPI12_3(%rip), %xmm0
	cmpl	$200, %ebx
	jne	.LBB12_3
# %bb.4:
	movq	link(%rip), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$48, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end12:
	.size	main, .Lfunc_end12-main
	.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:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB13_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB13_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z8kernel_09intfield5, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z8kernel_19intfield5diPfi, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end13:
	.size	__hip_module_ctor, .Lfunc_end13-__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	.LBB14_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
.LBB14_2:
	retq
.Lfunc_end14:
	.size	__hip_module_dtor, .Lfunc_end14-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	link,@object                    # @link
	.bss
	.globl	link
	.p2align	3, 0x0
link:
	.zero	32
	.size	link, 32

	.type	_Z8kernel_09intfield5,@object   # @_Z8kernel_09intfield5
	.section	.rodata,"a",@progbits
	.globl	_Z8kernel_09intfield5
	.p2align	3, 0x0
_Z8kernel_09intfield5:
	.quad	_Z23__device_stub__kernel_09intfield5
	.size	_Z8kernel_09intfield5, 8

	.type	_Z8kernel_19intfield5diPfi,@object # @_Z8kernel_19intfield5diPfi
	.globl	_Z8kernel_19intfield5diPfi
	.p2align	3, 0x0
_Z8kernel_19intfield5diPfi:
	.quad	_Z23__device_stub__kernel_19intfield5diPfi
	.size	_Z8kernel_19intfield5diPfi, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"%g\t%g\n"
	.size	.L.str, 7

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z8kernel_09intfield5"
	.size	.L__unnamed_1, 22

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z8kernel_19intfield5diPfi"
	.size	.L__unnamed_2, 27

	.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	"\n"
	.size	.Lstr, 2

	.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 _Z23__device_stub__kernel_09intfield5
	.addrsig_sym _Z23__device_stub__kernel_19intfield5diPfi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z8kernel_09intfield5
	.addrsig_sym _Z8kernel_19intfield5diPfi
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z8kernel_09intfield5   ; -- Begin function _Z8kernel_09intfield5
	.globl	_Z8kernel_09intfield5
	.p2align	8
	.type	_Z8kernel_09intfield5,@function
_Z8kernel_09intfield5:                  ; @_Z8kernel_09intfield5
; %bb.0:
	s_clause 0x2
	s_load_b32 s2, s[0:1], 0x2c
	s_load_b128 s[4:7], s[0:1], 0x8
	s_load_b32 s8, s[0:1], 0x18
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_mul_i32 s10, s5, s4
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mul_i32 s9, s10, s6
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s3, s9, s7
	s_mul_i32 s2, s3, s8
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_u32_e32 vcc_lo, s2, v1
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB0_2
; %bb.1:
	v_cvt_f32_u32_e32 v0, s10
	v_cvt_f32_u32_e32 v3, s9
	v_cvt_f32_u32_e32 v2, s3
	s_sub_i32 s2, 0, s10
	v_cvt_f32_u32_e32 v4, s4
	v_rcp_iflag_f32_e32 v0, v0
	v_rcp_iflag_f32_e32 v3, v3
	v_rcp_iflag_f32_e32 v2, v2
	s_sub_i32 s11, 0, s3
	v_rcp_iflag_f32_e32 v4, v4
	s_load_b64 s[0:1], s[0:1], 0x0
	s_delay_alu instid0(TRANS32_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v0, 0x4f7ffffe, v0 :: v_dual_mul_f32 v3, 0x4f7ffffe, v3
	v_cvt_u32_f32_e32 v0, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_u32_f32_e32 v3, v3
	v_mul_lo_u32 v5, s2, v0
	s_sub_i32 s2, 0, s9
	s_delay_alu instid0(VALU_DEP_2) | instid1(SALU_CYCLE_1)
	v_mul_lo_u32 v7, s2, v3
	s_sub_i32 s2, 0, s4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v5, v0, v5
	v_mul_f32_e32 v2, 0x4f7ffffe, v2
	v_mul_hi_u32 v7, v3, v7
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_dual_mul_f32 v4, 0x4f7ffffe, v4 :: v_dual_add_nc_u32 v5, v0, v5
	v_cvt_u32_f32_e32 v2, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v3, v3, v7
	v_cvt_u32_f32_e32 v4, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_lo_u32 v6, s11, v2
	v_mul_hi_u32 v7, v1, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_lo_u32 v8, s2, v4
	v_mul_hi_u32 v6, v2, v6
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v7, v7, s9
	v_sub_nc_u32_e32 v7, v1, v7
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v0, v2, v6
	v_mul_hi_u32 v2, v4, v8
	v_mul_hi_u32 v6, v1, v5
	v_mul_hi_u32 v0, v1, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v2, v4, v2
	v_mul_lo_u32 v4, v6, s10
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_lo_u32 v6, v0, s3
	v_mul_hi_u32 v8, v1, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_sub_nc_u32_e32 v4, v1, v4
	v_sub_nc_u32_e32 v6, v1, v6
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_lo_u32 v8, v8, s4
	v_subrev_nc_u32_e32 v10, s10, v4
	v_cmp_le_u32_e64 s2, s10, v4
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_subrev_nc_u32_e32 v11, s3, v6
	v_cmp_le_u32_e32 vcc_lo, s3, v6
	v_add_nc_u32_e32 v9, 1, v0
	v_cndmask_b32_e64 v4, v4, v10, s2
	v_sub_nc_u32_e32 v8, v1, v8
	s_mul_i32 s2, s5, s6
	v_cndmask_b32_e32 v6, v6, v11, vcc_lo
	v_cndmask_b32_e32 v0, v0, v9, vcc_lo
	v_subrev_nc_u32_e32 v9, s9, v7
	v_cmp_le_u32_e32 vcc_lo, s9, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_dual_cndmask_b32 v7, v7, v9 :: v_dual_add_nc_u32 v10, 1, v0
	v_cmp_le_u32_e32 vcc_lo, s3, v6
	v_subrev_nc_u32_e32 v9, s10, v4
	v_subrev_nc_u32_e32 v6, s9, v7
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_cndmask_b32_e32 v0, v0, v10, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s10, v4
	v_cndmask_b32_e32 v4, v4, v9, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s9, v7
	v_subrev_nc_u32_e32 v9, s4, v8
	v_cndmask_b32_e32 v6, v7, v6, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s4, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_mul_hi_u32 v5, v6, v5
	v_cndmask_b32_e32 v7, v8, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_subrev_nc_u32_e32 v8, s4, v7
	v_cmp_le_u32_e32 vcc_lo, s4, v7
	v_add_nc_u32_e32 v11, 1, v5
	v_mul_lo_u32 v10, v0, s3
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_cndmask_b32_e32 v7, v7, v8, vcc_lo
	v_mul_hi_u32 v2, v4, v2
	v_mul_lo_u32 v8, v5, s10
	v_sub_nc_u32_e32 v1, v1, v10
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_mul_lo_u32 v7, s2, v7
	v_mul_lo_u32 v9, v2, s4
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_mul_hi_u32 v3, v1, v3
	v_add_nc_u32_e32 v10, 1, v2
	v_sub_nc_u32_e32 v6, v6, v8
	v_sub_nc_u32_e32 v4, v4, v9
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_lo_u32 v9, v3, s9
	v_cmp_le_u32_e64 s2, s10, v6
	v_add_nc_u32_e32 v8, 1, v3
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_subrev_nc_u32_e32 v12, s4, v4
	v_cmp_le_u32_e32 vcc_lo, s4, v4
	v_cndmask_b32_e64 v5, v5, v11, s2
	v_sub_nc_u32_e32 v1, v1, v9
	v_cndmask_b32_e32 v2, v2, v10, vcc_lo
	v_cndmask_b32_e32 v4, v4, v12, vcc_lo
	v_subrev_nc_u32_e32 v10, s10, v6
	s_delay_alu instid0(VALU_DEP_4)
	v_cmp_le_u32_e32 vcc_lo, s9, v1
	v_subrev_nc_u32_e32 v11, s9, v1
	v_add_nc_u32_e32 v9, 1, v2
	v_cmp_le_u32_e64 s3, s4, v4
	v_cndmask_b32_e64 v4, v6, v10, s2
	v_dual_cndmask_b32 v3, v3, v8 :: v_dual_add_nc_u32 v6, 1, v5
	v_cndmask_b32_e32 v1, v1, v11, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v2, v2, v9, s3
	v_cmp_le_u32_e32 vcc_lo, s10, v4
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v8, 1, v3
	v_mul_lo_u32 v2, v2, s6
	v_cndmask_b32_e32 v4, v5, v6, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s9, v1
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v1, v3, v8, vcc_lo
	v_add3_u32 v4, v2, v7, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[2:3], null, v4, s7, v[1:2]
	v_mad_u64_u32 v[3:4], null, v2, s8, v[0:1]
	v_mov_b32_e32 v2, 1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v4, 31, v3
	v_lshlrev_b64 v[0:1], 2, v[3:4]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b32 v[0:1], v2, off
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8kernel_09intfield5
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 13
		.amdhsa_next_free_sgpr 16
		.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
	.text
.Lfunc_end0:
	.size	_Z8kernel_09intfield5, .Lfunc_end0-_Z8kernel_09intfield5
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 864
; NumSgprs: 18
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 13
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z8kernel_19intfield5diPfi ; -- Begin function _Z8kernel_19intfield5diPfi
	.globl	_Z8kernel_19intfield5diPfi
	.p2align	8
	.type	_Z8kernel_19intfield5diPfi,@function
_Z8kernel_19intfield5diPfi:             ; @_Z8kernel_19intfield5diPfi
; %bb.0:
	s_clause 0x2
	s_load_b32 s4, s[2:3], 0x4c
	s_load_b128 s[8:11], s[2:3], 0x8
	s_load_b32 s6, s[2:3], 0x18
	v_and_b32_e32 v6, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_mul_i32 s7, s9, s8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[2:3], null, s15, s4, v[6:7]
	s_mul_i32 s5, s7, s10
	s_mul_i32 s12, s5, s11
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mul_i32 s4, s12, s6
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_u32_e32 vcc_lo, s4, v2
	s_and_saveexec_b32 s4, vcc_lo
	s_cbranch_execz .LBB1_15
; %bb.1:
	v_cvt_f32_u32_e32 v1, s7
	v_cvt_f32_u32_e32 v4, s12
	v_cvt_f32_u32_e32 v5, s8
	s_sub_i32 s4, 0, s7
	s_sub_i32 s14, 0, s12
	v_rcp_iflag_f32_e32 v1, v1
	v_rcp_iflag_f32_e32 v4, v4
	v_rcp_iflag_f32_e32 v5, v5
	v_cvt_f32_u32_e32 v3, s5
	s_sub_i32 s13, 0, s5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(TRANS32_DEP_3)
	v_rcp_iflag_f32_e32 v3, v3
	v_dual_mul_f32 v1, 0x4f7ffffe, v1 :: v_dual_mul_f32 v4, 0x4f7ffffe, v4
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v5, 0x4f7ffffe, v5
	v_cvt_u32_f32_e32 v1, v1
	v_cvt_u32_f32_e32 v4, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cvt_u32_f32_e32 v5, v5
	v_mul_lo_u32 v7, s4, v1
	s_delay_alu instid0(VALU_DEP_3)
	v_mul_lo_u32 v9, s14, v4
	s_sub_i32 s4, 0, s8
	s_delay_alu instid0(VALU_DEP_3) | instid1(SALU_CYCLE_1)
	v_mul_lo_u32 v10, s4, v5
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v7, v1, v7
	v_mul_hi_u32 v9, v4, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v7, v1, v7
	v_add_nc_u32_e32 v1, v4, v9
	v_mul_hi_u32 v4, v5, v10
	v_mul_f32_e32 v3, 0x4f7ffffe, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v1, v2, v1
	v_add_nc_u32_e32 v4, v5, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cvt_u32_f32_e32 v3, v3
	v_add_nc_u32_e32 v11, 1, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v10, v2, v4
	v_mul_lo_u32 v8, s13, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v10, v10, s8
	v_mul_hi_u32 v8, v3, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v10, v2, v10
	v_add_nc_u32_e32 v3, v3, v8
	v_mul_hi_u32 v8, v2, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v9, v2, v3
	v_mul_lo_u32 v5, v8, s7
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_mul_lo_u32 v8, v9, s5
	v_mul_lo_u32 v9, v1, s12
	v_sub_nc_u32_e32 v5, v2, v5
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_sub_nc_u32_e32 v8, v2, v8
	v_sub_nc_u32_e32 v9, v2, v9
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_subrev_nc_u32_e32 v12, s7, v5
	v_cmp_le_u32_e64 s4, s7, v5
	v_subrev_nc_u32_e32 v13, s5, v8
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_le_u32_e32 vcc_lo, s12, v9
	v_cndmask_b32_e64 v5, v5, v12, s4
	v_cmp_le_u32_e64 s4, s5, v8
	v_cndmask_b32_e32 v1, v1, v11, vcc_lo
	v_subrev_nc_u32_e32 v11, s12, v9
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_subrev_nc_u32_e32 v12, s7, v5
	v_cndmask_b32_e64 v8, v8, v13, s4
	s_load_b32 s4, s[2:3], 0x38
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v9, v9, v11, vcc_lo
	v_add_nc_u32_e32 v11, 1, v1
	v_subrev_nc_u32_e32 v13, s5, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_le_u32_e32 vcc_lo, s12, v9
	v_cndmask_b32_e32 v1, v1, v11, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s7, v5
	v_subrev_nc_u32_e32 v11, s8, v10
	s_delay_alu instid0(VALU_DEP_3)
	v_mul_lo_u32 v9, v1, s12
	v_cndmask_b32_e32 v5, v5, v12, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s5, v8
	s_load_b64 s[12:13], s[0:1], 0x4
	v_cndmask_b32_e32 v8, v8, v13, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s8, v10
	v_sub_nc_u32_e32 v9, v2, v9
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v7, v8, v7
	v_cndmask_b32_e32 v10, v10, v11, vcc_lo
	v_subrev_nc_u32_e32 v2, s8, v10
	v_cmp_le_u32_e32 vcc_lo, s8, v10
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_lo_u32 v12, v7, s7
	v_cndmask_b32_e32 v2, v10, v2, vcc_lo
	v_mul_hi_u32 v3, v9, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v8, v8, v12
	v_add_nc_u32_e32 v12, 1, v3
	v_mul_hi_u32 v4, v5, v4
	v_mul_lo_u32 v13, v3, s5
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_subrev_nc_u32_e32 v14, s7, v8
	v_cmp_le_u32_e64 s0, s7, v8
	v_mul_lo_u32 v11, v4, s8
	v_add_nc_u32_e32 v10, 1, v4
	v_sub_nc_u32_e32 v9, v9, v13
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v8, v8, v14, s0
	v_sub_nc_u32_e32 v5, v5, v11
	v_add_nc_u32_e32 v11, 1, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s8, v5
	v_cndmask_b32_e64 v7, v7, v11, s0
	v_subrev_nc_u32_e32 v11, s5, v9
	v_cndmask_b32_e32 v4, v4, v10, vcc_lo
	v_subrev_nc_u32_e32 v10, s8, v5
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v13, 1, v7
	v_cndmask_b32_e32 v5, v5, v10, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s5, v9
	v_add_nc_u32_e32 v10, 1, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_cmp_le_u32_e64 s0, s8, v5
	v_dual_cndmask_b32 v12, v3, v12 :: v_dual_cndmask_b32 v5, v9, v11
	v_cmp_le_u32_e32 vcc_lo, s7, v8
	v_cndmask_b32_e64 v3, v4, v10, s0
	s_delay_alu instid0(VALU_DEP_3)
	v_add_nc_u32_e32 v9, 1, v12
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s0, s12, 16
	v_cndmask_b32_e32 v4, v7, v13, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s5, v5
	v_add_nc_u32_e32 v7, v3, v2
	s_mul_i32 s0, s0, s13
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(VALU_DEP_3)
	v_mul_lo_u32 v6, s0, v6
	v_cndmask_b32_e32 v5, v12, v9, vcc_lo
	v_bfe_u32 v9, v0, 10, 10
	v_bfe_u32 v0, v0, 20, 10
	s_lshr_b32 s0, s4, 31
	v_add3_u32 v7, v7, v4, v5
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_mul_u32_u24_e32 v9, s13, v9
	s_add_i32 s0, s4, s0
	s_and_b32 s1, s0, -2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshrrev_b32_e32 v8, 31, v7
	v_add3_u32 v0, v6, v9, v0
	s_ashr_i32 s0, s0, 1
	s_sub_i32 s1, s4, s1
	v_cmp_eq_u32_e32 vcc_lo, s0, v1
	v_add_nc_u32_e32 v8, v7, v8
	v_mul_lo_u32 v0, v0, 20
	ds_store_2addr_b32 v0, v2, v3 offset1:1
	ds_store_2addr_b32 v0, v4, v5 offset0:2 offset1:3
	ds_store_b32 v0, v1 offset:16
	v_and_b32_e32 v8, -2, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v6, v7, v8
	v_cmp_ne_u32_e64 s0, s1, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s0, vcc_lo
	s_and_b32 exec_lo, exec_lo, s0
	s_cbranch_execz .LBB1_15
; %bb.2:                                ; %.preheader
	s_load_b64 s[4:5], s[2:3], 0x0
	s_mul_i32 s7, s6, s11
	v_lshl_add_u32 v7, v1, 2, v0
	v_mov_b32_e32 v6, 0
	v_mov_b32_e32 v8, v0
	s_mul_i32 s8, s7, s10
	s_mov_b64 s[0:1], 0
	s_mul_i32 s12, s8, s9
.LBB1_3:                                ; =>This Inner Loop Header: Depth=1
	s_mov_b32 s13, exec_lo
	v_cmpx_ne_u32_e64 s0, v1
	s_cbranch_execz .LBB1_5
; %bb.4:                                ;   in Loop: Header=BB1_3 Depth=1
	ds_load_b32 v2, v8
	s_waitcnt lgkmcnt(0)
	v_cmp_gt_i32_e32 vcc_lo, 1, v2
	v_cndmask_b32_e64 v3, -1, 9, vcc_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v2, v3, v2
	ds_store_b32 v8, v2
	ds_load_b32 v2, v7
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, 8, v2
	v_cndmask_b32_e64 v3, 1, -9, vcc_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v9, v3, v2
	ds_load_2addr_b32 v[2:3], v0 offset1:1
	ds_load_2addr_b32 v[4:5], v0 offset0:2 offset1:3
	ds_store_b32 v7, v9
	ds_load_b32 v9, v8
	s_waitcnt lgkmcnt(3)
	v_mul_lo_u32 v2, v2, s12
	s_waitcnt lgkmcnt(2)
	v_mul_lo_u32 v19, v5, s6
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, 8, v9
	v_cndmask_b32_e64 v10, 1, -9, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v13, v10, v9
	v_cmp_lt_i32_e32 vcc_lo, 8, v13
	v_cndmask_b32_e64 v9, 1, -9, vcc_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v17, v9, v13
	ds_load_2addr_b32 v[9:10], v0 offset1:1
	ds_load_2addr_b32 v[11:12], v0 offset0:2 offset1:3
	ds_store_b32 v8, v13
	ds_load_2addr_b32 v[13:14], v0 offset1:1
	ds_load_2addr_b32 v[15:16], v0 offset0:2 offset1:3
	ds_store_b32 v8, v17
	ds_load_b32 v17, v7
	s_waitcnt lgkmcnt(0)
	v_cmp_gt_i32_e32 vcc_lo, 1, v17
	v_cndmask_b32_e64 v18, -1, 9, vcc_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v17, v18, v17
	ds_store_b32 v7, v17
	ds_load_b32 v18, v8
	v_mul_lo_u32 v17, v4, s7
	v_mad_u64_u32 v[4:5], null, v3, s8, v[2:3]
	v_mov_b32_e32 v2, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v21, v4, v17, v19
	v_ashrrev_i32_e32 v22, 31, v21
	v_add_nc_u32_e32 v17, v21, v1
	s_waitcnt lgkmcnt(0)
	v_cmp_gt_i32_e32 vcc_lo, 1, v18
	v_cndmask_b32_e64 v3, -1, 9, vcc_lo
	v_add_co_u32 v21, vcc_lo, s0, v21
	v_add_co_ci_u32_e32 v22, vcc_lo, s1, v22, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v3, v3, v18
	v_ashrrev_i32_e32 v18, 31, v17
	v_mad_u64_u32 v[4:5], null, s9, v9, v[2:3]
	v_mov_b32_e32 v5, v14
	ds_load_2addr_b32 v[9:10], v0 offset1:1
	ds_load_2addr_b32 v[19:20], v0 offset0:2 offset1:3
	ds_store_b32 v8, v3
	ds_load_2addr_b32 v[2:3], v0 offset1:1
	v_lshlrev_b64 v[17:18], 2, v[17:18]
	v_mad_u64_u32 v[23:24], null, s9, v13, v[5:6]
	v_lshlrev_b64 v[13:14], 2, v[21:22]
	v_mad_u64_u32 v[21:22], null, s10, v4, v[11:12]
	ds_load_2addr_b32 v[4:5], v0 offset0:2 offset1:3
	v_mov_b32_e32 v11, v12
	v_mov_b32_e32 v12, v16
	v_mad_u64_u32 v[24:25], null, s10, v23, v[15:16]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_mad_u64_u32 v[15:16], null, s11, v21, v[11:12]
	s_waitcnt lgkmcnt(3)
	v_mul_lo_u32 v23, v20, s6
	v_mad_u64_u32 v[21:22], null, s11, v24, v[12:13]
	v_mul_lo_u32 v12, v9, s12
	s_waitcnt lgkmcnt(1)
	v_mov_b32_e32 v11, v3
	v_mul_lo_u32 v22, v10, s8
	s_delay_alu instid0(VALU_DEP_2)
	v_mad_u64_u32 v[9:10], null, s9, v2, v[11:12]
	v_mad_u64_u32 v[10:11], null, s6, v15, s[0:1]
	v_mad_u64_u32 v[15:16], null, s6, v21, s[0:1]
	v_mul_lo_u32 v21, v19, s7
	v_add3_u32 v22, v12, v1, v22
	s_waitcnt lgkmcnt(0)
	v_mad_u64_u32 v[19:20], null, s10, v9, v[4:5]
	v_mov_b32_e32 v9, v5
	v_add_co_u32 v12, vcc_lo, s4, v13
	v_add_co_ci_u32_e32 v13, vcc_lo, s5, v14, vcc_lo
	v_add3_u32 v14, v22, v21, v23
	s_delay_alu instid0(VALU_DEP_4)
	v_mad_u64_u32 v[20:21], null, s11, v19, v[9:10]
	v_ashrrev_i32_e32 v11, 31, v10
	v_ashrrev_i32_e32 v16, 31, v15
	v_add_co_u32 v17, vcc_lo, s4, v17
	v_add_co_ci_u32_e32 v18, vcc_lo, s5, v18, vcc_lo
	v_mad_u64_u32 v[23:24], null, s6, v20, s[0:1]
	v_lshlrev_b64 v[9:10], 2, v[10:11]
	v_lshlrev_b64 v[21:22], 2, v[15:16]
	v_ashrrev_i32_e32 v15, 31, v14
	s_clause 0x1
	global_load_b32 v13, v[12:13], off
	global_load_b32 v16, v[17:18], off
	v_add_co_u32 v9, vcc_lo, s4, v9
	v_lshlrev_b64 v[14:15], 2, v[14:15]
	v_ashrrev_i32_e32 v24, 31, v23
	v_add_co_ci_u32_e32 v10, vcc_lo, s5, v10, vcc_lo
	v_add_co_u32 v19, vcc_lo, s4, v21
	v_add_co_ci_u32_e32 v20, vcc_lo, s5, v22, vcc_lo
	v_add_co_u32 v14, vcc_lo, s4, v14
	v_lshlrev_b64 v[11:12], 2, v[23:24]
	v_add_co_ci_u32_e32 v15, vcc_lo, s5, v15, vcc_lo
	s_clause 0x2
	global_load_b32 v17, v[9:10], off
	global_load_b32 v18, v[19:20], off
	global_load_b32 v14, v[14:15], off
	v_add_co_u32 v9, vcc_lo, s4, v11
	v_add_co_ci_u32_e32 v10, vcc_lo, s5, v12, vcc_lo
	global_load_b32 v9, v[9:10], off
	s_waitcnt vmcnt(4)
	v_mul_lo_u32 v10, v16, v13
	s_waitcnt vmcnt(3)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v10, v10, v17
	s_waitcnt vmcnt(1)
	v_mul_lo_u32 v11, v14, v18
	v_cvt_f32_i32_e32 v10, v10
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v9, v11, v9
	v_add_f32_e32 v6, v6, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_f32_i32_e32 v9, v9
	v_add_f32_e32 v6, v6, v9
.LBB1_5:                                ;   in Loop: Header=BB1_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s13
	v_add_nc_u32_e32 v8, 4, v8
	s_add_u32 s0, s0, 1
	s_addc_u32 s1, s1, 0
	s_cmp_lg_u32 s0, 4
	s_cbranch_scc1 .LBB1_3
; %bb.6:
	v_cvt_f64_f32_e32 v[7:8], v6
	s_clause 0x1
	s_load_b64 s[0:1], s[2:3], 0x20
	s_load_b32 s7, s[2:3], 0x28
	s_mov_b32 s12, 0x6a5dcb37
	s_mov_b32 s13, 0x3e5ade15
	ds_load_b32 v0, v0 offset:16
	v_mul_lo_u32 v15, v3, 10
	v_mul_lo_u32 v16, 0x64, v5
	v_mul_lo_u32 v3, v3, s10
	s_waitcnt lgkmcnt(0)
	s_mulk_i32 s7, 0x2710
	s_delay_alu instid0(VALU_DEP_3) | instid1(SALU_CYCLE_1)
	v_add3_u32 v15, v2, s7, v15
	v_mul_lo_u32 v0, 0x3e8, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v0, v15, v16, v0
	v_ashrrev_i32_e32 v15, 31, v0
	v_mul_f64 v[7:8], v[7:8], s[0:1]
	s_mov_b32 s0, 0x652b82fe
	s_mov_b32 s1, 0x3ff71547
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v15, 25, v15
	v_add_nc_u32_e32 v15, v0, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v15, 0xffffff80, v15
	v_sub_nc_u32_e32 v15, v0, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v16, 31, v15
	v_lshlrev_b64 v[15:16], 2, v[15:16]
	v_mul_f64 v[9:10], v[7:8], s[0:1]
	s_mov_b32 s0, 0xfefa39ef
	s_mov_b32 s1, 0xbfe62e42
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rndne_f64_e32 v[9:10], v[9:10]
	v_fma_f64 v[11:12], v[9:10], s[0:1], v[7:8]
	s_mov_b32 s0, 0x3b39803f
	s_mov_b32 s1, 0xbc7abc9e
	v_cvt_i32_f64_e32 v0, v[9:10]
	s_delay_alu instid0(VALU_DEP_2)
	v_fma_f64 v[11:12], v[9:10], s[0:1], v[11:12]
	s_mov_b32 s0, 0xfca7ab0c
	s_mov_b32 s1, 0x3e928af3
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[13:14], v[11:12], s[12:13], s[0:1]
	s_mov_b32 s0, 0x623fde64
	s_mov_b32 s1, 0x3ec71dee
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[13:14], v[11:12], v[13:14], s[0:1]
	s_mov_b32 s0, 0x7c89e6b0
	s_mov_b32 s1, 0x3efa0199
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[13:14], v[11:12], v[13:14], s[0:1]
	s_mov_b32 s0, 0x14761f6e
	s_mov_b32 s1, 0x3f2a01a0
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[13:14], v[11:12], v[13:14], s[0:1]
	s_mov_b32 s0, 0x1852b7b0
	s_mov_b32 s1, 0x3f56c16c
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[13:14], v[11:12], v[13:14], s[0:1]
	s_mov_b32 s0, 0x11122322
	s_mov_b32 s1, 0x3f811111
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[13:14], v[11:12], v[13:14], s[0:1]
	s_mov_b32 s0, 0x555502a1
	s_mov_b32 s1, 0x3fa55555
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[13:14], v[11:12], v[13:14], s[0:1]
	s_mov_b32 s0, 0x55555511
	s_mov_b32 s1, 0x3fc55555
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[13:14], v[11:12], v[13:14], s[0:1]
	s_mov_b32 s0, 11
	s_mov_b32 s1, 0x3fe00000
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[13:14], v[11:12], v[13:14], s[0:1]
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, _ZL9rand_data@rel32@lo+4
	s_addc_u32 s1, s1, _ZL9rand_data@rel32@hi+12
	v_add_co_u32 v15, vcc_lo, v15, s0
	v_add_co_ci_u32_e32 v16, vcc_lo, s1, v16, vcc_lo
	v_cmp_nlt_f64_e32 vcc_lo, 0x40900000, v[7:8]
	v_cmp_ngt_f64_e64 s0, 0xc090cc00, v[7:8]
	global_load_b32 v15, v[15:16], off
	v_fma_f64 v[13:14], v[11:12], v[13:14], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[9:10], v[11:12], v[13:14], 1.0
	v_ldexp_f64 v[9:10], v[9:10], v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v0, 0x7ff00000, v10, vcc_lo
	s_and_b32 vcc_lo, s0, vcc_lo
	v_cndmask_b32_e32 v7, 0, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cndmask_b32_e64 v8, 0, v0, s0
	s_mul_i32 s0, s9, s10
	v_mul_lo_u32 v2, s0, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_f32_f64_e32 v7, v[7:8]
	v_add3_u32 v4, v3, v2, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f32 v0, null, v7, v7, 1.0
	v_div_scale_f32 v10, vcc_lo, 1.0, v7, 1.0
	v_rcp_f32_e32 v8, v0
	s_waitcnt_depctr 0xfff
	v_fma_f32 v9, -v0, v8, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v8, v9, v8
	v_mul_f32_e32 v9, v10, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v11, -v0, v9, v10
	v_fmac_f32_e32 v9, v11, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v0, -v0, v9, v10
	v_div_fmas_f32 v0, v0, v8, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fixup_f32 v0, v0, v7, 1.0
	v_add_f32_e32 v8, v0, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f32 v9, null, v8, v8, v7
	v_div_scale_f32 v11, vcc_lo, v7, v8, v7
	v_rcp_f32_e32 v10, v9
	s_waitcnt_depctr 0xfff
	v_fma_f32 v0, -v9, v10, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v10, v0, v10
	v_dual_mov_b32 v0, v5 :: v_dual_mul_f32 v5, v11, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[2:3], null, v4, s11, v[0:1]
	v_fma_f32 v0, -v9, v5, v11
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[3:4], null, v2, s6, v[1:2]
	v_fmac_f32_e32 v5, v0, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v0, -v9, v5, v11
	v_ashrrev_i32_e32 v4, 31, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v0, v0, v10, v5
	v_div_fixup_f32 v2, v0, v8, v7
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[3:4]
                                        ; implicit-def: $vgpr3
	v_add_co_u32 v0, s0, s4, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v1, s0, s5, v1, s0
	s_mov_b32 s0, exec_lo
	s_waitcnt vmcnt(0)
	v_cmpx_nlt_f32_e32 v15, v2
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB1_8
; %bb.7:
	v_mov_b32_e32 v2, -1
	v_xor_b32_e32 v3, 0x80000000, v6
	global_store_b32 v[0:1], v2, off
                                        ; implicit-def: $vgpr0_vgpr1
.LBB1_8:                                ; %Flow222
	s_and_not1_saveexec_b32 s0, s0
	s_cbranch_execz .LBB1_10
; %bb.9:
	v_dual_mov_b32 v2, 1 :: v_dual_mov_b32 v3, v6
	global_store_b32 v[0:1], v2, off
.LBB1_10:                               ; %.sink.split
	s_or_b32 exec_lo, exec_lo, s0
	v_bfrev_b32_e32 v2, 1
	s_mov_b32 s0, exec_lo
.LBB1_11:                               ; %ComputeLoop
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_ctz_i32_b32 s1, s0
	v_readlane_b32 s4, v3, s1
	s_lshl_b32 s1, 1, s1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s0, s0, s1
	s_cmp_lg_u32 s0, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v2, s4, v2
	s_cbranch_scc1 .LBB1_11
; %bb.12:                               ; %ComputeEnd
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s4, 0
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB1_15
; %bb.13:
	s_load_b64 s[0:1], s[2:3], 0x30
	v_mov_b32_e32 v3, 0
	s_waitcnt lgkmcnt(0)
	global_load_b32 v1, v3, s[0:1]
.LBB1_14:                               ; %atomicrmw.start
                                        ; =>This Inner Loop Header: Depth=1
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v0, v1, v2
	global_atomic_cmpswap_b32 v0, v3, v[0:1], s[0:1] glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u32_e32 vcc_lo, v0, v1
	v_mov_b32_e32 v1, v0
	s_or_b32 s4, vcc_lo, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s4
	s_cbranch_execnz .LBB1_14
.LBB1_15:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8kernel_19intfield5diPfi
		.amdhsa_group_segment_fixed_size 20480
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 320
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 1
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 26
		.amdhsa_next_free_sgpr 16
		.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
	.text
.Lfunc_end1:
	.size	_Z8kernel_19intfield5diPfi, .Lfunc_end1-_Z8kernel_19intfield5diPfi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 3036
; NumSgprs: 18
; NumVgprs: 26
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 20480 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 3
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 26
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	_ZL9rand_data,@object           ; @_ZL9rand_data
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
_ZL9rand_data:
	.long	0x3f440a07                      ; float 0.765778005
	.long	0x3ec2d1f2                      ; float 0.380508006
	.long	0x3f79e332                      ; float 0.976122975
	.long	0x3d447e4a                      ; float 0.0479720011
	.long	0x3ce4f54d                      ; float 0.0279489998
	.long	0x3efc7bcc                      ; float 0.493131995
	.long	0x3e148cb5                      ; float 0.145068005
	.long	0x3f700a6c                      ; float 0.937659025
	.long	0x3f303dcd                      ; float 0.688443005
	.long	0x3ea253da                      ; float 0.317045987
	.long	0x3f4dbbbf                      ; float 0.803646028
	.long	0x3f6af0e1                      ; float 0.91773802
	.long	0x3f038fcd                      ; float 0.513912976
	.long	0x3eba37ac                      ; float 0.363705993
	.long	0x3e0c918e                      ; float 0.137273997
	.long	0x3f2aaa3b                      ; float 0.66666001
	.long	0x3e80027e                      ; float 0.250019014
	.long	0x3f1f4b40                      ; float 0.622241973
	.long	0x3cae0e30                      ; float 0.0212469995
	.long	0x3ed04b5e                      ; float 0.406825006
	.long	0x3f352c5a                      ; float 0.707708001
	.long	0x3f5b3605                      ; float 0.856293022
	.long	0x3f729c02                      ; float 0.94769299
	.long	0x3e54c87a                      ; float 0.207796007
	.long	0x3eb9d29e                      ; float 0.362935007
	.long	0x3f66f955                      ; float 0.902242004
	.long	0x3edb1d93                      ; float 0.427960008
	.long	0x3f3467f1                      ; float 0.704711019
	.long	0x3f1d1f92                      ; float 0.613762975
	.long	0x3f2906dd                      ; float 0.660260975
	.long	0x3ec1aaa4                      ; float 0.37825501
	.long	0x3f27ab54                      ; float 0.65495801
	.long	0x3f6fd8f1                      ; float 0.936904013
	.long	0x3f2eef80                      ; float 0.68334198
	.long	0x3f6431be                      ; float 0.891384005
	.long	0x3e998a01                      ; float 0.299881011
	.long	0x3d843809                      ; float 0.0645600036
	.long	0x3e99db87                      ; float 0.300502986
	.long	0x3f12a151                      ; float 0.572773993
	.long	0x3e07dcbe                      ; float 0.132678002
	.long	0x3e07778e                      ; float 0.132292002
	.long	0x3ee09e13                      ; float 0.438706011
	.long	0x3f18342b                      ; float 0.59454602
	.long	0x3f565a47                      ; float 0.837315022
	.long	0x3e38c3f4                      ; float 0.180435002
	.long	0x3e5c2d28                      ; float 0.215016007
	.long	0x3f3a1199                      ; float 0.726831018
	.long	0x3f44626f                      ; float 0.767126977
	.long	0x3f0e743a                      ; float 0.556460977
	.long	0x3f5c5868                      ; float 0.860723972
	.long	0x3e077293                      ; float 0.132273003
	.long	0x3e93cdbc                      ; float 0.288679004
	.long	0x3a945f9e                      ; float 0.001132
	.long	0x3f7241c4                      ; float 0.946316003
	.long	0x3f3dab08                      ; float 0.740890979
	.long	0x3f009731                      ; float 0.502306998
	.long	0x3e41afc0                      ; float 0.189146996
	.long	0x3f1c1776                      ; float 0.609732985
	.long	0x3f3778cd                      ; float 0.716687023
	.long	0x3dc900c5                      ; float 0.098145999
	.long	0x3f26a748                      ; float 0.650990009
	.long	0x3ef3e100                      ; float 0.476325989
	.long	0x3f755971                      ; float 0.958396017
	.long	0x3eeaec8d                      ; float 0.458835989
	.long	0x3f559c7c                      ; float 0.834419012
	.long	0x3f60445b                      ; float 0.876043021
	.long	0x3f5224bc                      ; float 0.820873022
	.long	0x3eddc2d2                      ; float 0.433126986
	.long	0x3f4cf074                      ; float 0.800544024
	.long	0x3f7095f2                      ; float 0.939787983
	.long	0x3f3de8c4                      ; float 0.741832972
	.long	0x3f67cbd5                      ; float 0.905453979
	.long	0x3f4c028e                      ; float 0.796913981
	.long	0x3f114aa1                      ; float 0.567544997
	.long	0x3d5de269                      ; float 0.0541709997
	.long	0x3eaabffd                      ; float 0.333496004
	.long	0x3e7deb10                      ; float 0.247967005
	.long	0x3f615337                      ; float 0.880176007
	.long	0x3f42b5f6                      ; float 0.760589003
	.long	0x3f450eaa                      ; float 0.769755005
	.long	0x3c3506dd                      ; float 0.0110489996
	.long	0x3eb9144d                      ; float 0.361483008
	.long	0x3f5443f6                      ; float 0.829162001
	.long	0x3e69999a                      ; float 0.228125006
	.long	0x3f12a551                      ; float 0.572835028
	.long	0x3f5adfe7                      ; float 0.854978978
	.long	0x3d8fb54a                      ; float 0.0701700002
	.long	0x3f4282e8                      ; float 0.75980997
	.long	0x3cb673c5                      ; float 0.0222720001
	.long	0x3ef4bba1                      ; float 0.477993995
	.long	0x3c6e06d9                      ; float 0.0145279998
	.long	0x3f7dc811                      ; float 0.991334021
	.long	0x3ea0eb89                      ; float 0.314296991
	.long	0x3f70a5ad                      ; float 0.940028011
	.long	0x3e7145d8                      ; float 0.235617995
	.long	0x3f573786                      ; float 0.84069097
	.long	0x3f61dc2f                      ; float 0.882265985
	.long	0x3f5716f4                      ; float 0.840193987
	.long	0x3f7c40d1                      ; float 0.985364019
	.long	0x3f369d0f                      ; float 0.713334023
	.long	0x3f329931                      ; float 0.697650015
	.long	0x3db97e56                      ; float 0.0905729979
	.long	0x3e8648a6                      ; float 0.262273014
	.long	0x3f08db8c                      ; float 0.534600019
	.long	0x3f4310aa                      ; float 0.761973023
	.long	0x3e167f91                      ; float 0.146971002
	.long	0x3f2af7b1                      ; float 0.66784197
	.long	0x3d8da33c                      ; float 0.0691590011
	.long	0x3dd15b57                      ; float 0.102224998
	.long	0x3f7b8499                      ; float 0.982492029
	.long	0x3f6eea21                      ; float 0.933260023
	.long	0x3ee1effa                      ; float 0.441284001
	.long	0x3e1970b5                      ; float 0.149844006
	.long	0x3d21c044                      ; float 0.0394899994
	.long	0x3f054563                      ; float 0.520590007
	.long	0x3d927ed8                      ; float 0.0715309978
	.long	0x3e112dba                      ; float 0.141775995
	.long	0x3f339d80                      ; float 0.701622009
	.long	0x3e5ae74f                      ; float 0.213772997
	.long	0x3f37c782                      ; float 0.717887998
	.long	0x3f1f1c32                      ; float 0.621523976
	.long	0x3e926c7f                      ; float 0.28598401
	.long	0x3ee28651                      ; float 0.442431003
	.long	0x3ef16031                      ; float 0.471437007
	.long	0x3e4aa971                      ; float 0.197911993
	.long	0x3ea11a76                      ; float 0.314655006
	.long	0x3efe17a0                      ; float 0.496273994
	.long	0x3f65944b                      ; float 0.896794021
	.size	_ZL9rand_data, 512

	.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:
      - .offset:         0
        .size:           32
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8kernel_09intfield5
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z8kernel_09intfield5.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           32
        .value_kind:     by_value
      - .offset:         32
        .size:           8
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .offset:         56
        .size:           4
        .value_kind:     by_value
      - .offset:         64
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         68
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         72
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         76
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         78
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         80
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         82
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         84
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         86
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         128
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 20480
    .kernarg_segment_align: 8
    .kernarg_segment_size: 320
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8kernel_19intfield5diPfi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z8kernel_19intfield5diPfi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     26
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
