#include <cuda_runtime.h>
#include<cassert>
#include<sys/time.h>
#include<time.h>
#include<stdio.h>
#include<string>
#include<sstream>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
       fprintf(stderr,"GPUassert (%d): %s (%s, line %d)\n", code, cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

int timeval_subtract(struct timeval *result, struct timeval *t2, struct timeval *t1)
{
    unsigned int resolution=1000000;
    long int diff = (t2->tv_usec + resolution * t2->tv_sec) - (t1->tv_usec + resolution * t1->tv_sec);
    result->tv_sec = diff / resolution;
    result->tv_usec = diff % resolution;
    return (diff<0);
}

class MyInt8Tuple {
    unsigned int x0; unsigned int x1; unsigned int x2; unsigned int x3;
    unsigned int x4; unsigned int x5; unsigned int x6; unsigned int x7;
    
public:
    
    __device__ __host__ inline MyInt8Tuple()
        : x0(0), x1(0), x2(0), x3(0), x4(0), x5(0), x6(0), x7(0) {
    }
    
    __device__ __host__ inline unsigned int& operator[](const unsigned int i) {
        assert(i < 8);
        if (i == 0) return x0;
        else if (i == 1) return x1;
        else if (i == 2) return x2;
        else if (i == 3) return x3;
        else if (i == 4) return x4;
        else if (i == 5) return x5;
        else if (i == 6) return x6;
        else return x7; // i == 7
    }
};

class MyInt8Array {
    unsigned int arr[8];
    
public:
    
    __device__ __host__ inline MyInt8Array()
        : arr() {
    }
    
    __device__ __host__ unsigned int& operator[](const unsigned int i) {
        assert(i < 8);
        return arr[i];
    }
};

void populateIntArray(const unsigned int  num_elems,
                      unsigned int*       in_array
    ) {
    for(unsigned int i = 0; i < num_elems; i++) {
        in_array[i] = std::rand() % 20;
    }
}

template<typename T>
void printIntCollection(unsigned int length, std::string title, T collection) {
    fprintf(stderr, "%-12s [", (title + ":").c_str());
    bool first = true;
    for(unsigned int i = 0; i < length; i++) {
        if (first) {
            fprintf(stderr, "%2d", collection[i]);
            first = false;
        }
        else {
            fprintf(stderr, ", %2d", collection[i]);
        }
    }
    fprintf(stderr, "]\n");
}

#define NUM_ITERATIONS 50

template<typename IntType>
__global__ void testKernel(
    // IntType*            in_array,
    IntType*            out_array,
    unsigned int*       indices,
    const unsigned int  num_elems
    ) {
    const unsigned int gid = blockIdx.x*blockDim.x + threadIdx.x;
    if (gid < num_elems) {
        const unsigned int lane8 = threadIdx.x & 7;
        const unsigned int index = indices[lane8];
        
        IntType elem;
        for (unsigned int i = 0; i < NUM_ITERATIONS; i++) {
            elem[index]++;
        }
        out_array[gid] = elem;
    }
}

#define BLOCK_SIZE 512

template<typename IntType>
void test(const unsigned int num_elems,
          const unsigned int i0,
          const unsigned int i1,
          const unsigned int i2,
          const unsigned int i3,
          const unsigned int i4,
          const unsigned int i5,
          const unsigned int i6,
          const unsigned int i7
    ) {
    // fprintf(stderr, "sizeof(IntType) = %d\n", sizeof(IntType));
    
    struct timeval t_start, t_end, t_diff;
    unsigned long int elapsed;
    
    // IntType* in_array = (IntType*) malloc(num_elems * sizeof(IntType));
    IntType* out_array = (IntType*) malloc(num_elems * sizeof(IntType));
    
    // memset(in_array, 0, num_elems * sizeof(IntType));
    
    unsigned int indices[8] = {i0,i1,i2,i3,i4,i5,i6,i7};
    
    // IntType *in_array_d;
    IntType *out_array_d;
    unsigned int *indices_d;
    // cudaMalloc((void**)&in_array_d, num_elems * sizeof(IntType));
    gpuErrchk( cudaMalloc((void**)&out_array_d, num_elems * sizeof(IntType)) );
    gpuErrchk( cudaMalloc((void**)&indices_d, 8 * sizeof(unsigned int)) );
    
    // cudaMemcpy(in_array_d, in_array, num_elems * sizeof(IntType), cudaMemcpyHostToDevice);
    gpuErrchk( cudaMemset(out_array_d, 0, num_elems * sizeof(IntType)) );
    gpuErrchk( cudaMemcpy(indices_d, indices, 8 * sizeof(unsigned int), cudaMemcpyHostToDevice) );
    gpuErrchk( cudaThreadSynchronize() );
    
    const unsigned int num_blocks = (num_elems + BLOCK_SIZE - 1) / BLOCK_SIZE;
    // fprintf(stderr, "num_blocks = %d\n", num_blocks);
    
    gettimeofday(&t_start, NULL);
    
    testKernel<IntType><<<num_blocks, BLOCK_SIZE>>>(out_array_d, indices_d, num_elems);
    
    gpuErrchk( cudaPeekAtLastError() );
    gpuErrchk( cudaThreadSynchronize() );
    gettimeofday(&t_end, NULL);
    
    gpuErrchk( cudaMemcpy(out_array, out_array_d, num_elems * sizeof(IntType), cudaMemcpyDeviceToHost) );
    
    // cudaFree(in_array_d);
    gpuErrchk( cudaFree(out_array_d) );
    gpuErrchk( cudaFree(indices_d) );
    
    timeval_subtract(&t_diff, &t_end, &t_start);
    elapsed = (t_diff.tv_sec*1e6+t_diff.tv_usec);
    fprintf(stderr, "Total runtime:  %6lu microsecs\n", elapsed);
    
    std::srand(time(NULL));
    unsigned int index = std::rand() % num_elems;
    std::stringstream sstm;
    sstm << "Random element (" << index << ")";
    std::string s = sstm.str();
#ifdef PRINT
    printIntCollection(8, s, out_array[index]);
#endif
    
    // free(in_array);
    free(out_array);
}

int main(int argc, char** argv) {
    if (argc != 10) {
        printf("The program takes <num_elems> and eight indices as arguments!\n");
        return EXIT_FAILURE;
    }
    const unsigned int num_elems = strtoul(argv[1], NULL, 10);
    const unsigned int i0 = strtoul(argv[2], NULL, 10);
    const unsigned int i1 = strtoul(argv[3], NULL, 10);
    const unsigned int i2 = strtoul(argv[4], NULL, 10);
    const unsigned int i3 = strtoul(argv[5], NULL, 10);
    const unsigned int i4 = strtoul(argv[6], NULL, 10);
    const unsigned int i5 = strtoul(argv[7], NULL, 10);
    const unsigned int i6 = strtoul(argv[8], NULL, 10);
    const unsigned int i7 = strtoul(argv[9], NULL, 10);
    
    fprintf(stderr, "MyInt8Tuple:\n");
    test<MyInt8Tuple>(num_elems,i0,i1,i2,i3,i4,i5,i6,i7);
    fprintf(stderr, "MyInt8Array:\n");
    test<MyInt8Array>(num_elems,i0,i1,i2,i3,i4,i5,i6,i7);
    
    return EXIT_SUCCESS;
}

	.file	"tmpxft_00398af7_00000000-6_replayTest.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3711:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3711:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_ZL50__device_stub__Z10testKernelI11MyInt8ArrayEvPT_PjjP11MyInt8ArrayPjj, @function
_ZL50__device_stub__Z10testKernelI11MyInt8ArrayEvPT_PjjP11MyInt8ArrayPjj:
.LFB3735:
	.cfi_startproc
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z10testKernelI11MyInt8ArrayEvPT_Pjj(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	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 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L2:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3735:
	.size	_ZL50__device_stub__Z10testKernelI11MyInt8ArrayEvPT_PjjP11MyInt8ArrayPjj, .-_ZL50__device_stub__Z10testKernelI11MyInt8ArrayEvPT_PjjP11MyInt8ArrayPjj
	.type	_ZL50__device_stub__Z10testKernelI11MyInt8TupleEvPT_PjjP11MyInt8TuplePjj, @function
_ZL50__device_stub__Z10testKernelI11MyInt8TupleEvPT_PjjP11MyInt8TuplePjj:
.LFB3733:
	.cfi_startproc
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L7
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z10testKernelI11MyInt8TupleEvPT_Pjj(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	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 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L7:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L9
	call	__stack_chk_fail@PLT
.L9:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3733:
	.size	_ZL50__device_stub__Z10testKernelI11MyInt8TupleEvPT_PjjP11MyInt8TuplePjj, .-_ZL50__device_stub__Z10testKernelI11MyInt8TupleEvPT_PjjP11MyInt8TuplePjj
	.section	.text._Z10testKernelI11MyInt8TupleEvPT_Pjj,"axG",@progbits,_Z10testKernelI11MyInt8TupleEvPT_Pjj,comdat
	.weak	_Z10testKernelI11MyInt8TupleEvPT_Pjj
	.type	_Z10testKernelI11MyInt8TupleEvPT_Pjj, @function
_Z10testKernelI11MyInt8TupleEvPT_Pjj:
.LFB4040:
	.cfi_startproc
	endbr64
	jmp	_ZL50__device_stub__Z10testKernelI11MyInt8TupleEvPT_PjjP11MyInt8TuplePjj
	.cfi_endproc
.LFE4040:
	.size	_Z10testKernelI11MyInt8TupleEvPT_Pjj, .-_Z10testKernelI11MyInt8TupleEvPT_Pjj
	.section	.text._Z10testKernelI11MyInt8ArrayEvPT_Pjj,"axG",@progbits,_Z10testKernelI11MyInt8ArrayEvPT_Pjj,comdat
	.weak	_Z10testKernelI11MyInt8ArrayEvPT_Pjj
	.type	_Z10testKernelI11MyInt8ArrayEvPT_Pjj, @function
_Z10testKernelI11MyInt8ArrayEvPT_Pjj:
.LFB4042:
	.cfi_startproc
	endbr64
	jmp	_ZL50__device_stub__Z10testKernelI11MyInt8ArrayEvPT_PjjP11MyInt8ArrayPjj
	.cfi_endproc
.LFE4042:
	.size	_Z10testKernelI11MyInt8ArrayEvPT_Pjj, .-_Z10testKernelI11MyInt8ArrayEvPT_Pjj
	.section	.rodata._Z9gpuAssert9cudaErrorPKcib.str1.1,"aMS",@progbits,1
.LC0:
	.string	"GPUassert (%d): %s (%s, line %d)\n"
	.section	.text._Z9gpuAssert9cudaErrorPKcib,"axG",@progbits,_Z9gpuAssert9cudaErrorPKcib,comdat
	.weak	_Z9gpuAssert9cudaErrorPKcib
	.type	_Z9gpuAssert9cudaErrorPKcib, @function
_Z9gpuAssert9cudaErrorPKcib:
.LFB3693:
	.cfi_startproc
	endbr64
	testl	%edi, %edi
	je	.L19
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	%edx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movl	%ecx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movl	%edi, %ebx
	subq	$16, %rsp
	.cfi_def_cfa_offset 48
	movq	%rsi, 8(%rsp)
	call	cudaGetErrorString@PLT
	movq	stderr(%rip), %rdi
	movl	%ebx, %ecx
	leaq	.LC0(%rip), %rdx
	pushq	%r9
	.cfi_def_cfa_offset 56
	movq	%rax, %r8
	movl	$2, %esi
	xorl	%eax, %eax
	pushq	%r12
	.cfi_def_cfa_offset 64
	movq	24(%rsp), %r9
	call	__fprintf_chk@PLT
	popq	%r10
	.cfi_def_cfa_offset 56
	popq	%r11
	.cfi_def_cfa_offset 48
	testb	%bpl, %bpl
	je	.L13
	movl	%ebx, %edi
	call	exit@PLT
.L13:
	addq	$16, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
.L19:
	.cfi_restore 3
	.cfi_restore 6
	.cfi_restore 12
	ret
	.cfi_endproc
.LFE3693:
	.size	_Z9gpuAssert9cudaErrorPKcib, .-_Z9gpuAssert9cudaErrorPKcib
	.text
	.globl	_Z16timeval_subtractP7timevalS0_S0_
	.type	_Z16timeval_subtractP7timevalS0_S0_, @function
_Z16timeval_subtractP7timevalS0_S0_:
.LFB3694:
	.cfi_startproc
	endbr64
	imulq	$1000000, (%rsi), %rcx
	addq	8(%rsi), %rcx
	movl	$1000000, %esi
	imulq	$1000000, (%rdx), %rax
	addq	8(%rdx), %rax
	subq	%rax, %rcx
	movq	%rcx, %rax
	shrq	$63, %rcx
	cqto
	idivq	%rsi
	movq	%rax, (%rdi)
	movq	%rcx, %rax
	movq	%rdx, 8(%rdi)
	ret
	.cfi_endproc
.LFE3694:
	.size	_Z16timeval_subtractP7timevalS0_S0_, .-_Z16timeval_subtractP7timevalS0_S0_
	.globl	_Z16populateIntArrayjPj
	.type	_Z16populateIntArrayjPj, @function
_Z16populateIntArrayjPj:
.LFB3703:
	.cfi_startproc
	endbr64
	movl	%edi, %edi
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	leaq	(%rsi,%rdi,4), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movl	$20, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rsi, %rbx
.L24:
	cmpq	%r12, %rbx
	je	.L27
	call	rand@PLT
	addq	$4, %rbx
	cltd
	idivl	%ebp
	movl	%edx, -4(%rbx)
	jmp	.L24
.L27:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3703:
	.size	_Z16populateIntArrayjPj, .-_Z16populateIntArrayjPj
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"_Z10testKernelI11MyInt8ArrayEvPT_Pjj"
.LC2:
	.string	"_Z10testKernelI11MyInt8TupleEvPT_Pjj"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3738:
	.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	.LC1(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z10testKernelI11MyInt8ArrayEvPT_Pjj(%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	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z10testKernelI11MyInt8TupleEvPT_Pjj(%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
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.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
.LFE3738:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata._Z4testI11MyInt8TupleEvjjjjjjjjj.str1.1,"aMS",@progbits,1
.LC3:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/asbjornvg/mdiscriminator/master/cuda_code/replayTest.cu"
.LC6:
	.string	"Total runtime:  %6lu microsecs\n"
.LC7:
	.string	"Random element ("
.LC8:
	.string	")"
	.section	.text._Z4testI11MyInt8TupleEvjjjjjjjjj,"axG",@progbits,_Z4testI11MyInt8TupleEvjjjjjjjjj,comdat
	.weak	_Z4testI11MyInt8TupleEvjjjjjjjjj
	.type	_Z4testI11MyInt8TupleEvjjjjjjjjj, @function
_Z4testI11MyInt8TupleEvjjjjjjjjj:
.LFB4038:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4038
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	%ecx, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%r8d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%edi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%r13, %rbp
	salq	$5, %r13
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%r13, %rdi
	movl	%r9d, %ebx
	subq	$552, %rsp
	.cfi_def_cfa_offset 608
	movl	%esi, 12(%rsp)
	movl	%edx, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 536(%rsp)
	xorl	%eax, %eax
	call	malloc@PLT
	movl	12(%rsp), %esi
	movl	8(%rsp), %edx
	leaq	16(%rsp), %rdi
	movq	%rax, %r12
	movl	608(%rsp), %eax
	movl	%r15d, 120(%rsp)
	leaq	32(%rsp), %r15
	movl	%esi, 112(%rsp)
	movq	%r13, %rsi
	movl	%eax, 132(%rsp)
	movl	616(%rsp), %eax
	movl	%edx, 116(%rsp)
	movl	%eax, 136(%rsp)
	movl	624(%rsp), %eax
	movl	%ebx, 128(%rsp)
	leaq	.LC3(%rip), %rbx
	movl	%r14d, 124(%rsp)
	movl	%eax, 140(%rsp)
.LEHB0:
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$141, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	leaq	24(%rsp), %rdi
	movl	$32, %esi
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$142, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	movq	%r13, %rdx
	call	cudaMemset@PLT
	movl	$1, %ecx
	movl	$145, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	24(%rsp), %rdi
	leaq	112(%rsp), %rsi
	movl	$1, %ecx
	movl	$32, %edx
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$146, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	call	cudaThreadSynchronize@PLT
	movl	$1, %ecx
	movl	$147, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	xorl	%esi, %esi
	movq	%r15, %rdi
	call	gettimeofday@PLT
	movl	$8388609, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	leal	511(%rbp), %eax
	salq	$9, %rdx
	movl	$1, %ecx
	movl	$1, 72(%rsp)
	shrl	$9, %eax
	movq	%rdx, 64(%rsp)
	movl	%eax, 48(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 52(%rsp)
	movq	48(%rsp), %rdi
	movl	56(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L31
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	%ebp, %edx
	call	_ZL50__device_stub__Z10testKernelI11MyInt8TupleEvPT_PjjP11MyInt8TuplePjj
.L31:
	call	cudaPeekAtLastError@PLT
	movl	$1, %ecx
	movl	$156, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	leaq	48(%rsp), %r14
	call	_Z9gpuAssert9cudaErrorPKcib
	call	cudaThreadSynchronize@PLT
	movl	$1, %ecx
	movl	$157, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	xorl	%esi, %esi
	movq	%r14, %rdi
	call	gettimeofday@PLT
	movq	16(%rsp), %rsi
	movq	%r13, %rdx
	movq	%r12, %rdi
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$160, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$163, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$164, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	leaq	64(%rsp), %rdi
	movq	%r15, %rdx
	movq	%r14, %rsi
	call	_Z16timeval_subtractP7timevalS0_S0_
	cvtsi2sdq	64(%rsp), %xmm0
	mulsd	.LC4(%rip), %xmm0
	cvtsi2sdq	72(%rsp), %xmm1
	addsd	%xmm1, %xmm0
	movsd	.LC5(%rip), %xmm1
	comisd	%xmm1, %xmm0
	jnb	.L32
	cvttsd2siq	%xmm0, %rcx
	jmp	.L33
.L32:
	subsd	%xmm1, %xmm0
	cvttsd2siq	%xmm0, %rcx
	btcq	$63, %rcx
.L33:
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	leaq	.LC6(%rip), %rdx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	xorl	%edi, %edi
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
	call	rand@PLT
	xorl	%edx, %edx
	divl	%ebp
	leaq	144(%rsp), %rbp
	movq	%rbp, %rdi
	movl	%edx, %ebx
	call	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEEC1Ev@PLT
.LEHE0:
	leaq	160(%rsp), %rdi
	leaq	.LC7(%rip), %rsi
.LEHB1:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movl	%ebx, %esi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movq	%rax, %rdi
	leaq	.LC8(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	80(%rsp), %rbx
	leaq	168(%rsp), %rsi
	movq	%rbx, %rdi
	call	_ZNKSt7__cxx1115basic_stringbufIcSt11char_traitsIcESaIcEE3strEv@PLT
.LEHE1:
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEED1Ev@PLT
	movq	536(%rsp), %rax
	subq	%fs:40, %rax
	je	.L36
	jmp	.L38
.L37:
	endbr64
	movq	%rax, %rbx
.L34:
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEED1Ev@PLT
	movq	536(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L38
	movq	%rbx, %rdi
.LEHB2:
	call	_Unwind_Resume@PLT
.LEHE2:
.L38:
	call	__stack_chk_fail@PLT
.L36:
	addq	$552, %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
.LFE4038:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table._Z4testI11MyInt8TupleEvjjjjjjjjj,"aG",@progbits,_Z4testI11MyInt8TupleEvjjjjjjjjj,comdat
.LLSDA4038:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4038-.LLSDACSB4038
.LLSDACSB4038:
	.uleb128 .LEHB0-.LFB4038
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB4038
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L37-.LFB4038
	.uleb128 0
	.uleb128 .LEHB2-.LFB4038
	.uleb128 .LEHE2-.LEHB2
	.uleb128 0
	.uleb128 0
.LLSDACSE4038:
	.section	.text._Z4testI11MyInt8TupleEvjjjjjjjjj,"axG",@progbits,_Z4testI11MyInt8TupleEvjjjjjjjjj,comdat
	.size	_Z4testI11MyInt8TupleEvjjjjjjjjj, .-_Z4testI11MyInt8TupleEvjjjjjjjjj
	.section	.text._Z4testI11MyInt8ArrayEvjjjjjjjjj,"axG",@progbits,_Z4testI11MyInt8ArrayEvjjjjjjjjj,comdat
	.weak	_Z4testI11MyInt8ArrayEvjjjjjjjjj
	.type	_Z4testI11MyInt8ArrayEvjjjjjjjjj, @function
_Z4testI11MyInt8ArrayEvjjjjjjjjj:
.LFB4039:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4039
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	%ecx, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%r8d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%edi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%r13, %rbp
	salq	$5, %r13
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%r13, %rdi
	movl	%r9d, %ebx
	subq	$552, %rsp
	.cfi_def_cfa_offset 608
	movl	%esi, 12(%rsp)
	movl	%edx, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 536(%rsp)
	xorl	%eax, %eax
	call	malloc@PLT
	movl	12(%rsp), %esi
	movl	8(%rsp), %edx
	leaq	16(%rsp), %rdi
	movq	%rax, %r12
	movl	608(%rsp), %eax
	movl	%r15d, 120(%rsp)
	leaq	32(%rsp), %r15
	movl	%esi, 112(%rsp)
	movq	%r13, %rsi
	movl	%eax, 132(%rsp)
	movl	616(%rsp), %eax
	movl	%edx, 116(%rsp)
	movl	%eax, 136(%rsp)
	movl	624(%rsp), %eax
	movl	%ebx, 128(%rsp)
	leaq	.LC3(%rip), %rbx
	movl	%r14d, 124(%rsp)
	movl	%eax, 140(%rsp)
.LEHB3:
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$141, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	leaq	24(%rsp), %rdi
	movl	$32, %esi
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$142, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	movq	%r13, %rdx
	call	cudaMemset@PLT
	movl	$1, %ecx
	movl	$145, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	24(%rsp), %rdi
	leaq	112(%rsp), %rsi
	movl	$1, %ecx
	movl	$32, %edx
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$146, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	call	cudaThreadSynchronize@PLT
	movl	$1, %ecx
	movl	$147, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	xorl	%esi, %esi
	movq	%r15, %rdi
	call	gettimeofday@PLT
	movl	$8388609, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	leal	511(%rbp), %eax
	salq	$9, %rdx
	movl	$1, %ecx
	movl	$1, 72(%rsp)
	shrl	$9, %eax
	movq	%rdx, 64(%rsp)
	movl	%eax, 48(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 52(%rsp)
	movq	48(%rsp), %rdi
	movl	56(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	%ebp, %edx
	call	_ZL50__device_stub__Z10testKernelI11MyInt8ArrayEvPT_PjjP11MyInt8ArrayPjj
.L41:
	call	cudaPeekAtLastError@PLT
	movl	$1, %ecx
	movl	$156, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	leaq	48(%rsp), %r14
	call	_Z9gpuAssert9cudaErrorPKcib
	call	cudaThreadSynchronize@PLT
	movl	$1, %ecx
	movl	$157, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	xorl	%esi, %esi
	movq	%r14, %rdi
	call	gettimeofday@PLT
	movq	16(%rsp), %rsi
	movq	%r13, %rdx
	movq	%r12, %rdi
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$160, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$163, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$164, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	leaq	64(%rsp), %rdi
	movq	%r15, %rdx
	movq	%r14, %rsi
	call	_Z16timeval_subtractP7timevalS0_S0_
	cvtsi2sdq	64(%rsp), %xmm0
	mulsd	.LC4(%rip), %xmm0
	cvtsi2sdq	72(%rsp), %xmm1
	addsd	%xmm1, %xmm0
	movsd	.LC5(%rip), %xmm1
	comisd	%xmm1, %xmm0
	jnb	.L42
	cvttsd2siq	%xmm0, %rcx
	jmp	.L43
.L42:
	subsd	%xmm1, %xmm0
	cvttsd2siq	%xmm0, %rcx
	btcq	$63, %rcx
.L43:
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	leaq	.LC6(%rip), %rdx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	xorl	%edi, %edi
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
	call	rand@PLT
	xorl	%edx, %edx
	divl	%ebp
	leaq	144(%rsp), %rbp
	movq	%rbp, %rdi
	movl	%edx, %ebx
	call	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEEC1Ev@PLT
.LEHE3:
	leaq	160(%rsp), %rdi
	leaq	.LC7(%rip), %rsi
.LEHB4:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movl	%ebx, %esi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movq	%rax, %rdi
	leaq	.LC8(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	80(%rsp), %rbx
	leaq	168(%rsp), %rsi
	movq	%rbx, %rdi
	call	_ZNKSt7__cxx1115basic_stringbufIcSt11char_traitsIcESaIcEE3strEv@PLT
.LEHE4:
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEED1Ev@PLT
	movq	536(%rsp), %rax
	subq	%fs:40, %rax
	je	.L46
	jmp	.L48
.L47:
	endbr64
	movq	%rax, %rbx
.L44:
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEED1Ev@PLT
	movq	536(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L48
	movq	%rbx, %rdi
.LEHB5:
	call	_Unwind_Resume@PLT
.LEHE5:
.L48:
	call	__stack_chk_fail@PLT
.L46:
	addq	$552, %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
.LFE4039:
	.section	.gcc_except_table._Z4testI11MyInt8ArrayEvjjjjjjjjj,"aG",@progbits,_Z4testI11MyInt8ArrayEvjjjjjjjjj,comdat
.LLSDA4039:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4039-.LLSDACSB4039
.LLSDACSB4039:
	.uleb128 .LEHB3-.LFB4039
	.uleb128 .LEHE3-.LEHB3
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB4-.LFB4039
	.uleb128 .LEHE4-.LEHB4
	.uleb128 .L47-.LFB4039
	.uleb128 0
	.uleb128 .LEHB5-.LFB4039
	.uleb128 .LEHE5-.LEHB5
	.uleb128 0
	.uleb128 0
.LLSDACSE4039:
	.section	.text._Z4testI11MyInt8ArrayEvjjjjjjjjj,"axG",@progbits,_Z4testI11MyInt8ArrayEvjjjjjjjjj,comdat
	.size	_Z4testI11MyInt8ArrayEvjjjjjjjjj, .-_Z4testI11MyInt8ArrayEvjjjjjjjjj
	.section	.rodata.str1.1
.LC9:
	.string	"The program takes <num_elems> and eight indices as arguments!\n"
.LC10:
	.string	"MyInt8Tuple:\n"
.LC11:
	.string	"MyInt8Array:\n"
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB3708:
	.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
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	cmpl	$10, %edi
	je	.L51
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1, %eax
	jmp	.L50
.L51:
	movq	8(%rsi), %rdi
	movq	%rsi, %r15
	movl	$10, %edx
	xorl	%esi, %esi
	call	__isoc23_strtoul@PLT
	movq	16(%r15), %rdi
	movl	$10, %edx
	xorl	%esi, %esi
	movl	%eax, 4(%rsp)
	call	__isoc23_strtoul@PLT
	movq	24(%r15), %rdi
	movl	$10, %edx
	xorl	%esi, %esi
	movl	%eax, 8(%rsp)
	call	__isoc23_strtoul@PLT
	movq	32(%r15), %rdi
	movl	$10, %edx
	xorl	%esi, %esi
	movl	%eax, 12(%rsp)
	call	__isoc23_strtoul@PLT
	movq	40(%r15), %rdi
	movl	$10, %edx
	xorl	%esi, %esi
	movl	%eax, %ebx
	call	__isoc23_strtoul@PLT
	movq	48(%r15), %rdi
	movl	$10, %edx
	xorl	%esi, %esi
	movq	%rax, %rbp
	call	__isoc23_strtoul@PLT
	movq	56(%r15), %rdi
	movl	$10, %edx
	xorl	%esi, %esi
	movq	%rax, %r12
	call	__isoc23_strtoul@PLT
	movq	64(%r15), %rdi
	movl	$10, %edx
	xorl	%esi, %esi
	movq	%rax, %r13
	call	__isoc23_strtoul@PLT
	movq	72(%r15), %rdi
	movl	$10, %edx
	xorl	%esi, %esi
	movq	%rax, %r14
	call	__isoc23_strtoul@PLT
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	leaq	.LC10(%rip), %rdx
	movq	%rax, %r15
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movl	%r12d, %r9d
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	pushq	%rax
	.cfi_def_cfa_offset 88
	pushq	%r15
	.cfi_def_cfa_offset 96
	pushq	%r14
	.cfi_def_cfa_offset 104
	pushq	%r13
	.cfi_def_cfa_offset 112
	movl	44(%rsp), %edx
	movl	40(%rsp), %esi
	movl	36(%rsp), %edi
	call	_Z4testI11MyInt8TupleEvjjjjjjjjj
	addq	$32, %rsp
	.cfi_def_cfa_offset 80
	movl	$2, %esi
	xorl	%eax, %eax
	movq	stderr(%rip), %rdi
	leaq	.LC11(%rip), %rdx
	call	__fprintf_chk@PLT
	movl	%r12d, %r9d
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	pushq	%rdx
	.cfi_def_cfa_offset 88
	pushq	%r15
	.cfi_def_cfa_offset 96
	pushq	%r14
	.cfi_def_cfa_offset 104
	pushq	%r13
	.cfi_def_cfa_offset 112
	movl	44(%rsp), %edx
	movl	40(%rsp), %esi
	movl	36(%rsp), %edi
	call	_Z4testI11MyInt8ArrayEvjjjjjjjjj
	addq	$32, %rsp
	.cfi_def_cfa_offset 80
	xorl	%eax, %eax
.L50:
	addq	$24, %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
.LFE3708:
	.size	main, .-main
	.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.cst8,"aM",@progbits,8
	.align 8
.LC4:
	.long	0
	.long	1093567616
	.align 8
.LC5:
	.long	0
	.long	1138753536
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.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<cassert>
#include<sys/time.h>
#include<time.h>
#include<stdio.h>
#include<string>
#include<sstream>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(hipError_t code, const char *file, int line, bool abort=true)
{
   if (code != hipSuccess) 
   {
       fprintf(stderr,"GPUassert (%d): %s (%s, line %d)\n", code, hipGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

int timeval_subtract(struct timeval *result, struct timeval *t2, struct timeval *t1)
{
    unsigned int resolution=1000000;
    long int diff = (t2->tv_usec + resolution * t2->tv_sec) - (t1->tv_usec + resolution * t1->tv_sec);
    result->tv_sec = diff / resolution;
    result->tv_usec = diff % resolution;
    return (diff<0);
}

class MyInt8Tuple {
    unsigned int x0; unsigned int x1; unsigned int x2; unsigned int x3;
    unsigned int x4; unsigned int x5; unsigned int x6; unsigned int x7;
    
public:
    
    __device__ __host__ inline MyInt8Tuple()
        : x0(0), x1(0), x2(0), x3(0), x4(0), x5(0), x6(0), x7(0) {
    }
    
    __device__ __host__ inline unsigned int& operator[](const unsigned int i) {
        assert(i < 8);
        if (i == 0) return x0;
        else if (i == 1) return x1;
        else if (i == 2) return x2;
        else if (i == 3) return x3;
        else if (i == 4) return x4;
        else if (i == 5) return x5;
        else if (i == 6) return x6;
        else return x7; // i == 7
    }
};

class MyInt8Array {
    unsigned int arr[8];
    
public:
    
    __device__ __host__ inline MyInt8Array()
        : arr() {
    }
    
    __device__ __host__ unsigned int& operator[](const unsigned int i) {
        assert(i < 8);
        return arr[i];
    }
};

void populateIntArray(const unsigned int  num_elems,
                      unsigned int*       in_array
    ) {
    for(unsigned int i = 0; i < num_elems; i++) {
        in_array[i] = std::rand() % 20;
    }
}

template<typename T>
void printIntCollection(unsigned int length, std::string title, T collection) {
    fprintf(stderr, "%-12s [", (title + ":").c_str());
    bool first = true;
    for(unsigned int i = 0; i < length; i++) {
        if (first) {
            fprintf(stderr, "%2d", collection[i]);
            first = false;
        }
        else {
            fprintf(stderr, ", %2d", collection[i]);
        }
    }
    fprintf(stderr, "]\n");
}

#define NUM_ITERATIONS 50

template<typename IntType>
__global__ void testKernel(
    // IntType*            in_array,
    IntType*            out_array,
    unsigned int*       indices,
    const unsigned int  num_elems
    ) {
    const unsigned int gid = blockIdx.x*blockDim.x + threadIdx.x;
    if (gid < num_elems) {
        const unsigned int lane8 = threadIdx.x & 7;
        const unsigned int index = indices[lane8];
        
        IntType elem;
        for (unsigned int i = 0; i < NUM_ITERATIONS; i++) {
            elem[index]++;
        }
        out_array[gid] = elem;
    }
}

#define BLOCK_SIZE 512

template<typename IntType>
void test(const unsigned int num_elems,
          const unsigned int i0,
          const unsigned int i1,
          const unsigned int i2,
          const unsigned int i3,
          const unsigned int i4,
          const unsigned int i5,
          const unsigned int i6,
          const unsigned int i7
    ) {
    // fprintf(stderr, "sizeof(IntType) = %d\n", sizeof(IntType));
    
    struct timeval t_start, t_end, t_diff;
    unsigned long int elapsed;
    
    // IntType* in_array = (IntType*) malloc(num_elems * sizeof(IntType));
    IntType* out_array = (IntType*) malloc(num_elems * sizeof(IntType));
    
    // memset(in_array, 0, num_elems * sizeof(IntType));
    
    unsigned int indices[8] = {i0,i1,i2,i3,i4,i5,i6,i7};
    
    // IntType *in_array_d;
    IntType *out_array_d;
    unsigned int *indices_d;
    // cudaMalloc((void**)&in_array_d, num_elems * sizeof(IntType));
    gpuErrchk( hipMalloc((void**)&out_array_d, num_elems * sizeof(IntType)) );
    gpuErrchk( hipMalloc((void**)&indices_d, 8 * sizeof(unsigned int)) );
    
    // cudaMemcpy(in_array_d, in_array, num_elems * sizeof(IntType), cudaMemcpyHostToDevice);
    gpuErrchk( hipMemset(out_array_d, 0, num_elems * sizeof(IntType)) );
    gpuErrchk( hipMemcpy(indices_d, indices, 8 * sizeof(unsigned int), hipMemcpyHostToDevice) );
    gpuErrchk( hipDeviceSynchronize() );
    
    const unsigned int num_blocks = (num_elems + BLOCK_SIZE - 1) / BLOCK_SIZE;
    // fprintf(stderr, "num_blocks = %d\n", num_blocks);
    
    gettimeofday(&t_start, NULL);
    
    testKernel<IntType><<<num_blocks, BLOCK_SIZE>>>(out_array_d, indices_d, num_elems);
    
    gpuErrchk( hipPeekAtLastError() );
    gpuErrchk( hipDeviceSynchronize() );
    gettimeofday(&t_end, NULL);
    
    gpuErrchk( hipMemcpy(out_array, out_array_d, num_elems * sizeof(IntType), hipMemcpyDeviceToHost) );
    
    // cudaFree(in_array_d);
    gpuErrchk( hipFree(out_array_d) );
    gpuErrchk( hipFree(indices_d) );
    
    timeval_subtract(&t_diff, &t_end, &t_start);
    elapsed = (t_diff.tv_sec*1e6+t_diff.tv_usec);
    fprintf(stderr, "Total runtime:  %6lu microsecs\n", elapsed);
    
    std::srand(time(NULL));
    unsigned int index = std::rand() % num_elems;
    std::stringstream sstm;
    sstm << "Random element (" << index << ")";
    std::string s = sstm.str();
#ifdef PRINT
    printIntCollection(8, s, out_array[index]);
#endif
    
    // free(in_array);
    free(out_array);
}

int main(int argc, char** argv) {
    if (argc != 10) {
        printf("The program takes <num_elems> and eight indices as arguments!\n");
        return EXIT_FAILURE;
    }
    const unsigned int num_elems = strtoul(argv[1], NULL, 10);
    const unsigned int i0 = strtoul(argv[2], NULL, 10);
    const unsigned int i1 = strtoul(argv[3], NULL, 10);
    const unsigned int i2 = strtoul(argv[4], NULL, 10);
    const unsigned int i3 = strtoul(argv[5], NULL, 10);
    const unsigned int i4 = strtoul(argv[6], NULL, 10);
    const unsigned int i5 = strtoul(argv[7], NULL, 10);
    const unsigned int i6 = strtoul(argv[8], NULL, 10);
    const unsigned int i7 = strtoul(argv[9], NULL, 10);
    
    fprintf(stderr, "MyInt8Tuple:\n");
    test<MyInt8Tuple>(num_elems,i0,i1,i2,i3,i4,i5,i6,i7);
    fprintf(stderr, "MyInt8Array:\n");
    test<MyInt8Array>(num_elems,i0,i1,i2,i3,i4,i5,i6,i7);
    
    return EXIT_SUCCESS;
}

	.text
	.file	"replayTest.hip"
	.globl	_Z16timeval_subtractP7timevalS0_S0_ # -- Begin function _Z16timeval_subtractP7timevalS0_S0_
	.type	_Z16timeval_subtractP7timevalS0_S0_,@function
_Z16timeval_subtractP7timevalS0_S0_:    # @_Z16timeval_subtractP7timevalS0_S0_
	.cfi_startproc
# %bb.0:
	movq	(%rsi), %rax
	movq	8(%rsi), %rcx
	subq	(%rdx), %rax
	imulq	$1000000, %rax, %rax            # imm = 0xF4240
	subq	8(%rdx), %rcx
	addq	%rax, %rcx
	movabsq	$4835703278458516699, %rdx      # imm = 0x431BDE82D7B634DB
	movq	%rcx, %rax
	imulq	%rdx
	movq	%rdx, %rax
	shrq	$63, %rax
	sarq	$18, %rdx
	addq	%rax, %rdx
	movq	%rdx, (%rdi)
	imulq	$1000000, %rdx, %rax            # imm = 0xF4240
	movq	%rcx, %rdx
	subq	%rax, %rdx
	movq	%rdx, 8(%rdi)
	shrq	$63, %rcx
	movl	%ecx, %eax
	retq
.Lfunc_end0:
	.size	_Z16timeval_subtractP7timevalS0_S0_, .Lfunc_end0-_Z16timeval_subtractP7timevalS0_S0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z16populateIntArrayjPj         # -- Begin function _Z16populateIntArrayjPj
	.type	_Z16populateIntArrayjPj,@function
_Z16populateIntArrayjPj:                # @_Z16populateIntArrayjPj
	.cfi_startproc
# %bb.0:
	testl	%edi, %edi
	je	.LBB1_4
# %bb.1:                                # %.lr.ph.preheader
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %rbx
	movl	%edi, %r14d
	xorl	%r15d, %r15d
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	cltq
	imulq	$1717986919, %rax, %rcx         # imm = 0x66666667
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$35, %rcx
	addl	%edx, %ecx
	shll	$2, %ecx
	leal	(%rcx,%rcx,4), %ecx
	subl	%ecx, %eax
	movl	%eax, (%rbx,%r15,4)
	incq	%r15
	cmpq	%r15, %r14
	jne	.LBB1_2
# %bb.3:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r14
	.cfi_restore %r15
.LBB1_4:                                # %._crit_edge
	retq
.Lfunc_end1:
	.size	_Z16populateIntArrayjPj, .Lfunc_end1-_Z16populateIntArrayjPj
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.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
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	cmpl	$10, %edi
	je	.LBB2_2
# %bb.1:
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	$1, %ebx
.LBB2_3:
	movl	%ebx, %eax
	addq	$40, %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
.LBB2_2:
	.cfi_def_cfa_offset 96
	movq	%rsi, %rbp
	movq	8(%rsi), %rdi
	xorl	%ebx, %ebx
	xorl	%esi, %esi
	movl	$10, %edx
	callq	__isoc23_strtoul
	movq	%rax, (%rsp)                    # 8-byte Spill
	movq	16(%rbp), %rdi
	xorl	%esi, %esi
	movl	$10, %edx
	callq	__isoc23_strtoul
	movq	%rax, %r15
	movq	24(%rbp), %rdi
	xorl	%esi, %esi
	movl	$10, %edx
	callq	__isoc23_strtoul
	movq	%rax, %r12
	movq	32(%rbp), %rdi
	xorl	%esi, %esi
	movl	$10, %edx
	callq	__isoc23_strtoul
	movq	%rax, %r13
	movq	40(%rbp), %rdi
	xorl	%esi, %esi
	movl	$10, %edx
	callq	__isoc23_strtoul
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movq	48(%rbp), %rdi
	xorl	%esi, %esi
	movl	$10, %edx
	callq	__isoc23_strtoul
	movq	%rax, 24(%rsp)                  # 8-byte Spill
	movq	56(%rbp), %rdi
	xorl	%esi, %esi
	movl	$10, %edx
	callq	__isoc23_strtoul
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movq	64(%rbp), %rdi
	xorl	%esi, %esi
	movl	$10, %edx
	callq	__isoc23_strtoul
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	movq	72(%rbp), %rdi
	xorl	%esi, %esi
	movl	$10, %edx
	callq	__isoc23_strtoul
	movq	%rax, %r14
	movq	stderr(%rip), %rcx
	movl	$.L.str.1, %edi
	movl	$13, %esi
	movl	$1, %edx
	callq	fwrite@PLT
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movq	8(%rsp), %rdi                   # 8-byte Reload
                                        # kill: def $edi killed $edi killed $rdi
	movl	%r15d, %esi
	movl	%r12d, %edx
	movl	%r13d, %ecx
	movq	40(%rsp), %rbp                  # 8-byte Reload
	movl	%ebp, %r8d
	movq	32(%rsp), %r9                   # 8-byte Reload
                                        # kill: def $r9d killed $r9d killed $r9
	pushq	%r14
	.cfi_adjust_cfa_offset 8
	pushq	24(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	40(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	callq	_Z4testI11MyInt8TupleEvjjjjjjjjj
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	movq	stderr(%rip), %rcx
	movl	$.L.str.2, %edi
	movl	$13, %esi
	movl	$1, %edx
	callq	fwrite@PLT
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movq	8(%rsp), %rdi                   # 8-byte Reload
                                        # kill: def $edi killed $edi killed $rdi
	movl	%r15d, %esi
	movl	%r12d, %edx
	movl	%r13d, %ecx
	movl	%ebp, %r8d
	movq	32(%rsp), %r9                   # 8-byte Reload
                                        # kill: def $r9d killed $r9d killed $r9
	pushq	%r14
	.cfi_adjust_cfa_offset 8
	pushq	24(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	40(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	callq	_Z4testI11MyInt8ArrayEvjjjjjjjjj
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	jmp	.LBB2_3
.Lfunc_end2:
	.size	main, .Lfunc_end2-main
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z4testI11MyInt8TupleEvjjjjjjjjj
.LCPI3_0:
	.quad	0x412e848000000000              # double 1.0E+6
.LCPI3_1:
	.quad	0x43e0000000000000              # double 9.2233720368547758E+18
	.section	.text._Z4testI11MyInt8TupleEvjjjjjjjjj,"axG",@progbits,_Z4testI11MyInt8TupleEvjjjjjjjjj,comdat
	.weak	_Z4testI11MyInt8TupleEvjjjjjjjjj
	.type	_Z4testI11MyInt8TupleEvjjjjjjjjj,@function
_Z4testI11MyInt8TupleEvjjjjjjjjj:       # @_Z4testI11MyInt8TupleEvjjjjjjjjj
.Lfunc_begin0:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception0
# %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
	subq	$520, %rsp                      # imm = 0x208
	.cfi_def_cfa_offset 576
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%r9d, %ebp
	movl	%r8d, %r12d
	movl	%ecx, %r13d
	movl	%edx, %ebx
	movl	%esi, %r14d
                                        # kill: def $edi killed $edi def $rdi
	movq	%rdi, 56(%rsp)                  # 8-byte Spill
	movl	%edi, %r15d
	shlq	$5, %r15
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	movl	%r14d, 16(%rsp)
	movl	%ebx, 20(%rsp)
	movl	%r13d, 24(%rsp)
	movl	%r12d, 28(%rsp)
	movl	%ebp, 32(%rsp)
	movl	576(%rsp), %eax
	movl	%eax, 36(%rsp)
	movl	584(%rsp), %eax
	movl	%eax, 40(%rsp)
	movl	592(%rsp), %eax
	movl	%eax, 44(%rsp)
	movq	%rsp, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB3_1
# %bb.3:                                # %_Z9gpuAssert10hipError_tPKcib.exit
	leaq	8(%rsp), %rdi
	movl	$32, %esi
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB3_4
# %bb.5:                                # %_Z9gpuAssert10hipError_tPKcib.exit26
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	movq	%r15, %rdx
	callq	hipMemset
	testl	%eax, %eax
	movq	56(%rsp), %r14                  # 8-byte Reload
	jne	.LBB3_6
# %bb.7:                                # %_Z9gpuAssert10hipError_tPKcib.exit28
	movq	8(%rsp), %rdi
	leaq	16(%rsp), %rsi
	movl	$32, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB3_8
# %bb.9:                                # %_Z9gpuAssert10hipError_tPKcib.exit30
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_10
# %bb.11:                               # %_Z9gpuAssert10hipError_tPKcib.exit32
	leal	511(%r14), %ebx
	shrl	$9, %ebx
	leaq	80(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	btsq	$32, %rbx
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$512, %rdx                      # imm = 0x200
	movq	%rbx, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_13
# %bb.12:
	movq	(%rsp), %rdi
	movq	8(%rsp), %rsi
	movl	%r14d, %edx
	callq	_Z25__device_stub__testKernelI11MyInt8TupleEvPT_Pjj
.LBB3_13:
	callq	hipPeekAtLastError
	testl	%eax, %eax
	movq	48(%rsp), %rbx                  # 8-byte Reload
	jne	.LBB3_14
# %bb.15:                               # %_Z9gpuAssert10hipError_tPKcib.exit34
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_16
# %bb.17:                               # %_Z9gpuAssert10hipError_tPKcib.exit36
	leaq	64(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%rsp), %rsi
	movq	%rbx, %rdi
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB3_18
# %bb.19:                               # %_Z9gpuAssert10hipError_tPKcib.exit38
	movq	(%rsp), %rdi
	callq	hipFree
	testl	%eax, %eax
	jne	.LBB3_20
# %bb.21:                               # %_Z9gpuAssert10hipError_tPKcib.exit40
	movq	8(%rsp), %rdi
	callq	hipFree
	testl	%eax, %eax
	jne	.LBB3_22
# %bb.23:                               # %_Z9gpuAssert10hipError_tPKcib.exit42
	movq	64(%rsp), %rax
	movq	72(%rsp), %rcx
	subq	80(%rsp), %rax
	imulq	$1000000, %rax, %rax            # imm = 0xF4240
	subq	88(%rsp), %rcx
	addq	%rax, %rcx
	movabsq	$4835703278458516699, %rdx      # imm = 0x431BDE82D7B634DB
	movq	%rcx, %rax
	imulq	%rdx
	movq	%rdx, %rax
	shrq	$63, %rax
	sarq	$18, %rdx
	addq	%rax, %rdx
	imulq	$1000000, %rdx, %rax            # imm = 0xF4240
	subq	%rax, %rcx
	cvtsi2sd	%rdx, %xmm0
	mulsd	.LCPI3_0(%rip), %xmm0
	cvtsi2sd	%rcx, %xmm1
	addsd	%xmm0, %xmm1
	cvttsd2si	%xmm1, %rax
	movq	%rax, %rcx
	subsd	.LCPI3_1(%rip), %xmm1
	cvttsd2si	%xmm1, %rdx
	sarq	$63, %rcx
	andq	%rcx, %rdx
	orq	%rax, %rdx
	movq	stderr(%rip), %rdi
	movl	$.L.str.4, %esi
	xorl	%eax, %eax
	callq	fprintf
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	callq	rand
	movl	%eax, %ebp
	leaq	128(%rsp), %rdi
	callq	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEEC1Ev
	leaq	144(%rsp), %r15
.Ltmp0:
	movl	$.L.str.5, %esi
	movl	$16, %edx
	movq	%r15, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp1:
# %bb.24:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit
	movl	%ebp, %eax
	xorl	%edx, %edx
	divl	%r14d
                                        # kill: def $edx killed $edx def $rdx
.Ltmp2:
	movq	%r15, %rdi
	movq	%rdx, %rsi
	callq	_ZNSo9_M_insertImEERSoT_
.Ltmp3:
# %bb.25:                               # %_ZNSolsEj.exit
.Ltmp4:
	movl	$.L.str.6, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp5:
# %bb.26:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit43
	leaq	152(%rsp), %rsi
.Ltmp7:
	leaq	96(%rsp), %rdi
	callq	_ZNKSt7__cxx1115basic_stringbufIcSt11char_traitsIcESaIcEE3strEv
.Ltmp8:
# %bb.27:                               # %_ZNKSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEE3strEv.exit
	movq	%rbx, %rdi
	callq	free
	leaq	112(%rsp), %rax
	movq	-16(%rax), %rdi
	cmpq	%rax, %rdi
	je	.LBB3_29
# %bb.28:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i
	callq	_ZdlPv
.LBB3_29:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit
	leaq	128(%rsp), %rdi
	movl	$_ZTTNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEEE, %esi
	callq	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEED2Ev
	leaq	256(%rsp), %rdi
	callq	_ZNSt8ios_baseD2Ev
	addq	$520, %rsp                      # imm = 0x208
	.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
.LBB3_1:
	.cfi_def_cfa_offset 576
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$141, %r9d
	jmp	.LBB3_2
.LBB3_4:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$142, %r9d
	jmp	.LBB3_2
.LBB3_6:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$145, %r9d
	jmp	.LBB3_2
.LBB3_8:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$146, %r9d
	jmp	.LBB3_2
.LBB3_10:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$147, %r9d
	jmp	.LBB3_2
.LBB3_14:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$156, %r9d
	jmp	.LBB3_2
.LBB3_16:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$157, %r9d
	jmp	.LBB3_2
.LBB3_18:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$160, %r9d
	jmp	.LBB3_2
.LBB3_20:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$163, %r9d
	jmp	.LBB3_2
.LBB3_22:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$164, %r9d
.LBB3_2:
	xorl	%eax, %eax
	callq	fprintf
	movl	%ebp, %edi
	callq	exit
.LBB3_31:
.Ltmp9:
	jmp	.LBB3_32
.LBB3_30:
.Ltmp6:
.LBB3_32:
	movq	%rax, %rbx
	leaq	128(%rsp), %rdi
	movl	$_ZTTNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEEE, %esi
	callq	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEED2Ev
	leaq	256(%rsp), %rdi
	callq	_ZNSt8ios_baseD2Ev
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end3:
	.size	_Z4testI11MyInt8TupleEvjjjjjjjjj, .Lfunc_end3-_Z4testI11MyInt8TupleEvjjjjjjjjj
	.cfi_endproc
	.section	.gcc_except_table._Z4testI11MyInt8TupleEvjjjjjjjjj,"aG",@progbits,_Z4testI11MyInt8TupleEvjjjjjjjjj,comdat
	.p2align	2, 0x0
GCC_except_table3:
.Lexception0:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end0-.Lcst_begin0
.Lcst_begin0:
	.uleb128 .Lfunc_begin0-.Lfunc_begin0    # >> Call Site 1 <<
	.uleb128 .Ltmp0-.Lfunc_begin0           #   Call between .Lfunc_begin0 and .Ltmp0
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp0-.Lfunc_begin0           # >> Call Site 2 <<
	.uleb128 .Ltmp5-.Ltmp0                  #   Call between .Ltmp0 and .Ltmp5
	.uleb128 .Ltmp6-.Lfunc_begin0           #     jumps to .Ltmp6
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp7-.Lfunc_begin0           # >> Call Site 3 <<
	.uleb128 .Ltmp8-.Ltmp7                  #   Call between .Ltmp7 and .Ltmp8
	.uleb128 .Ltmp9-.Lfunc_begin0           #     jumps to .Ltmp9
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp8-.Lfunc_begin0           # >> Call Site 4 <<
	.uleb128 .Lfunc_end3-.Ltmp8             #   Call between .Ltmp8 and .Lfunc_end3
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end0:
	.p2align	2, 0x0
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z4testI11MyInt8ArrayEvjjjjjjjjj
.LCPI4_0:
	.quad	0x412e848000000000              # double 1.0E+6
.LCPI4_1:
	.quad	0x43e0000000000000              # double 9.2233720368547758E+18
	.section	.text._Z4testI11MyInt8ArrayEvjjjjjjjjj,"axG",@progbits,_Z4testI11MyInt8ArrayEvjjjjjjjjj,comdat
	.weak	_Z4testI11MyInt8ArrayEvjjjjjjjjj
	.type	_Z4testI11MyInt8ArrayEvjjjjjjjjj,@function
_Z4testI11MyInt8ArrayEvjjjjjjjjj:       # @_Z4testI11MyInt8ArrayEvjjjjjjjjj
.Lfunc_begin1:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception1
# %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
	subq	$520, %rsp                      # imm = 0x208
	.cfi_def_cfa_offset 576
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%r9d, %ebp
	movl	%r8d, %r12d
	movl	%ecx, %r13d
	movl	%edx, %ebx
	movl	%esi, %r14d
                                        # kill: def $edi killed $edi def $rdi
	movq	%rdi, 56(%rsp)                  # 8-byte Spill
	movl	%edi, %r15d
	shlq	$5, %r15
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	movl	%r14d, 16(%rsp)
	movl	%ebx, 20(%rsp)
	movl	%r13d, 24(%rsp)
	movl	%r12d, 28(%rsp)
	movl	%ebp, 32(%rsp)
	movl	576(%rsp), %eax
	movl	%eax, 36(%rsp)
	movl	584(%rsp), %eax
	movl	%eax, 40(%rsp)
	movl	592(%rsp), %eax
	movl	%eax, 44(%rsp)
	movq	%rsp, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB4_1
# %bb.3:                                # %_Z9gpuAssert10hipError_tPKcib.exit
	leaq	8(%rsp), %rdi
	movl	$32, %esi
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB4_4
# %bb.5:                                # %_Z9gpuAssert10hipError_tPKcib.exit26
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	movq	%r15, %rdx
	callq	hipMemset
	testl	%eax, %eax
	movq	56(%rsp), %r14                  # 8-byte Reload
	jne	.LBB4_6
# %bb.7:                                # %_Z9gpuAssert10hipError_tPKcib.exit28
	movq	8(%rsp), %rdi
	leaq	16(%rsp), %rsi
	movl	$32, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB4_8
# %bb.9:                                # %_Z9gpuAssert10hipError_tPKcib.exit30
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB4_10
# %bb.11:                               # %_Z9gpuAssert10hipError_tPKcib.exit32
	leal	511(%r14), %ebx
	shrl	$9, %ebx
	leaq	80(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	btsq	$32, %rbx
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$512, %rdx                      # imm = 0x200
	movq	%rbx, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_13
# %bb.12:
	movq	(%rsp), %rdi
	movq	8(%rsp), %rsi
	movl	%r14d, %edx
	callq	_Z25__device_stub__testKernelI11MyInt8ArrayEvPT_Pjj
.LBB4_13:
	callq	hipPeekAtLastError
	testl	%eax, %eax
	movq	48(%rsp), %rbx                  # 8-byte Reload
	jne	.LBB4_14
# %bb.15:                               # %_Z9gpuAssert10hipError_tPKcib.exit34
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB4_16
# %bb.17:                               # %_Z9gpuAssert10hipError_tPKcib.exit36
	leaq	64(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%rsp), %rsi
	movq	%rbx, %rdi
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB4_18
# %bb.19:                               # %_Z9gpuAssert10hipError_tPKcib.exit38
	movq	(%rsp), %rdi
	callq	hipFree
	testl	%eax, %eax
	jne	.LBB4_20
# %bb.21:                               # %_Z9gpuAssert10hipError_tPKcib.exit40
	movq	8(%rsp), %rdi
	callq	hipFree
	testl	%eax, %eax
	jne	.LBB4_22
# %bb.23:                               # %_Z9gpuAssert10hipError_tPKcib.exit42
	movq	64(%rsp), %rax
	movq	72(%rsp), %rcx
	subq	80(%rsp), %rax
	imulq	$1000000, %rax, %rax            # imm = 0xF4240
	subq	88(%rsp), %rcx
	addq	%rax, %rcx
	movabsq	$4835703278458516699, %rdx      # imm = 0x431BDE82D7B634DB
	movq	%rcx, %rax
	imulq	%rdx
	movq	%rdx, %rax
	shrq	$63, %rax
	sarq	$18, %rdx
	addq	%rax, %rdx
	imulq	$1000000, %rdx, %rax            # imm = 0xF4240
	subq	%rax, %rcx
	cvtsi2sd	%rdx, %xmm0
	mulsd	.LCPI4_0(%rip), %xmm0
	cvtsi2sd	%rcx, %xmm1
	addsd	%xmm0, %xmm1
	cvttsd2si	%xmm1, %rax
	movq	%rax, %rcx
	subsd	.LCPI4_1(%rip), %xmm1
	cvttsd2si	%xmm1, %rdx
	sarq	$63, %rcx
	andq	%rcx, %rdx
	orq	%rax, %rdx
	movq	stderr(%rip), %rdi
	movl	$.L.str.4, %esi
	xorl	%eax, %eax
	callq	fprintf
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	callq	rand
	movl	%eax, %ebp
	leaq	128(%rsp), %rdi
	callq	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEEC1Ev
	leaq	144(%rsp), %r15
.Ltmp10:
	movl	$.L.str.5, %esi
	movl	$16, %edx
	movq	%r15, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp11:
# %bb.24:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit
	movl	%ebp, %eax
	xorl	%edx, %edx
	divl	%r14d
                                        # kill: def $edx killed $edx def $rdx
.Ltmp12:
	movq	%r15, %rdi
	movq	%rdx, %rsi
	callq	_ZNSo9_M_insertImEERSoT_
.Ltmp13:
# %bb.25:                               # %_ZNSolsEj.exit
.Ltmp14:
	movl	$.L.str.6, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp15:
# %bb.26:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit43
	leaq	152(%rsp), %rsi
.Ltmp17:
	leaq	96(%rsp), %rdi
	callq	_ZNKSt7__cxx1115basic_stringbufIcSt11char_traitsIcESaIcEE3strEv
.Ltmp18:
# %bb.27:                               # %_ZNKSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEE3strEv.exit
	movq	%rbx, %rdi
	callq	free
	leaq	112(%rsp), %rax
	movq	-16(%rax), %rdi
	cmpq	%rax, %rdi
	je	.LBB4_29
# %bb.28:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i
	callq	_ZdlPv
.LBB4_29:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit
	leaq	128(%rsp), %rdi
	movl	$_ZTTNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEEE, %esi
	callq	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEED2Ev
	leaq	256(%rsp), %rdi
	callq	_ZNSt8ios_baseD2Ev
	addq	$520, %rsp                      # imm = 0x208
	.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
.LBB4_1:
	.cfi_def_cfa_offset 576
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$141, %r9d
	jmp	.LBB4_2
.LBB4_4:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$142, %r9d
	jmp	.LBB4_2
.LBB4_6:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$145, %r9d
	jmp	.LBB4_2
.LBB4_8:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$146, %r9d
	jmp	.LBB4_2
.LBB4_10:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$147, %r9d
	jmp	.LBB4_2
.LBB4_14:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$156, %r9d
	jmp	.LBB4_2
.LBB4_16:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$157, %r9d
	jmp	.LBB4_2
.LBB4_18:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$160, %r9d
	jmp	.LBB4_2
.LBB4_20:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$163, %r9d
	jmp	.LBB4_2
.LBB4_22:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.3, %r8d
	movq	%rbx, %rdi
	movl	%ebp, %edx
	movq	%rax, %rcx
	movl	$164, %r9d
.LBB4_2:
	xorl	%eax, %eax
	callq	fprintf
	movl	%ebp, %edi
	callq	exit
.LBB4_31:
.Ltmp19:
	jmp	.LBB4_32
.LBB4_30:
.Ltmp16:
.LBB4_32:
	movq	%rax, %rbx
	leaq	128(%rsp), %rdi
	movl	$_ZTTNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEEE, %esi
	callq	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEED2Ev
	leaq	256(%rsp), %rdi
	callq	_ZNSt8ios_baseD2Ev
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end4:
	.size	_Z4testI11MyInt8ArrayEvjjjjjjjjj, .Lfunc_end4-_Z4testI11MyInt8ArrayEvjjjjjjjjj
	.cfi_endproc
	.section	.gcc_except_table._Z4testI11MyInt8ArrayEvjjjjjjjjj,"aG",@progbits,_Z4testI11MyInt8ArrayEvjjjjjjjjj,comdat
	.p2align	2, 0x0
GCC_except_table4:
.Lexception1:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end1-.Lcst_begin1
.Lcst_begin1:
	.uleb128 .Lfunc_begin1-.Lfunc_begin1    # >> Call Site 1 <<
	.uleb128 .Ltmp10-.Lfunc_begin1          #   Call between .Lfunc_begin1 and .Ltmp10
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp10-.Lfunc_begin1          # >> Call Site 2 <<
	.uleb128 .Ltmp15-.Ltmp10                #   Call between .Ltmp10 and .Ltmp15
	.uleb128 .Ltmp16-.Lfunc_begin1          #     jumps to .Ltmp16
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp17-.Lfunc_begin1          # >> Call Site 3 <<
	.uleb128 .Ltmp18-.Ltmp17                #   Call between .Ltmp17 and .Ltmp18
	.uleb128 .Ltmp19-.Lfunc_begin1          #     jumps to .Ltmp19
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp18-.Lfunc_begin1          # >> Call Site 4 <<
	.uleb128 .Lfunc_end4-.Ltmp18            #   Call between .Ltmp18 and .Lfunc_end4
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end1:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text._Z25__device_stub__testKernelI11MyInt8TupleEvPT_Pjj,"axG",@progbits,_Z25__device_stub__testKernelI11MyInt8TupleEvPT_Pjj,comdat
	.weak	_Z25__device_stub__testKernelI11MyInt8TupleEvPT_Pjj # -- Begin function _Z25__device_stub__testKernelI11MyInt8TupleEvPT_Pjj
	.type	_Z25__device_stub__testKernelI11MyInt8TupleEvPT_Pjj,@function
_Z25__device_stub__testKernelI11MyInt8TupleEvPT_Pjj: # @_Z25__device_stub__testKernelI11MyInt8TupleEvPT_Pjj
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%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	$_Z10testKernelI11MyInt8TupleEvPT_Pjj, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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_end5:
	.size	_Z25__device_stub__testKernelI11MyInt8TupleEvPT_Pjj, .Lfunc_end5-_Z25__device_stub__testKernelI11MyInt8TupleEvPT_Pjj
	.cfi_endproc
                                        # -- End function
	.section	.text._Z25__device_stub__testKernelI11MyInt8ArrayEvPT_Pjj,"axG",@progbits,_Z25__device_stub__testKernelI11MyInt8ArrayEvPT_Pjj,comdat
	.weak	_Z25__device_stub__testKernelI11MyInt8ArrayEvPT_Pjj # -- Begin function _Z25__device_stub__testKernelI11MyInt8ArrayEvPT_Pjj
	.type	_Z25__device_stub__testKernelI11MyInt8ArrayEvPT_Pjj,@function
_Z25__device_stub__testKernelI11MyInt8ArrayEvPT_Pjj: # @_Z25__device_stub__testKernelI11MyInt8ArrayEvPT_Pjj
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%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	$_Z10testKernelI11MyInt8ArrayEvPT_Pjj, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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_end6:
	.size	_Z25__device_stub__testKernelI11MyInt8ArrayEvPT_Pjj, .Lfunc_end6-_Z25__device_stub__testKernelI11MyInt8ArrayEvPT_Pjj
	.cfi_endproc
                                        # -- End function
	.text
	.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	.LBB7_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB7_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10testKernelI11MyInt8TupleEvPT_Pjj, %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	$_Z10testKernelI11MyInt8ArrayEvPT_Pjj, %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_end7:
	.size	__hip_module_ctor, .Lfunc_end7-__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	.LBB8_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
.LBB8_2:
	retq
.Lfunc_end8:
	.size	__hip_module_dtor, .Lfunc_end8-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"MyInt8Tuple:\n"
	.size	.L.str.1, 14

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"MyInt8Array:\n"
	.size	.L.str.2, 14

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"/home/ubuntu/Datasets/stackv2/train-structured-repos-hip-sm89/asbjornvg/mdiscriminator/master/cuda_code/replayTest.hip"
	.size	.L.str.3, 119

	.type	_Z10testKernelI11MyInt8TupleEvPT_Pjj,@object # @_Z10testKernelI11MyInt8TupleEvPT_Pjj
	.section	.rodata._Z10testKernelI11MyInt8TupleEvPT_Pjj,"aG",@progbits,_Z10testKernelI11MyInt8TupleEvPT_Pjj,comdat
	.weak	_Z10testKernelI11MyInt8TupleEvPT_Pjj
	.p2align	3, 0x0
_Z10testKernelI11MyInt8TupleEvPT_Pjj:
	.quad	_Z25__device_stub__testKernelI11MyInt8TupleEvPT_Pjj
	.size	_Z10testKernelI11MyInt8TupleEvPT_Pjj, 8

	.type	.L.str.4,@object                # @.str.4
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.4:
	.asciz	"Total runtime:  %6lu microsecs\n"
	.size	.L.str.4, 32

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"Random element ("
	.size	.L.str.5, 17

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	")"
	.size	.L.str.6, 2

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"GPUassert (%d): %s (%s, line %d)\n"
	.size	.L.str.7, 34

	.type	_Z10testKernelI11MyInt8ArrayEvPT_Pjj,@object # @_Z10testKernelI11MyInt8ArrayEvPT_Pjj
	.section	.rodata._Z10testKernelI11MyInt8ArrayEvPT_Pjj,"aG",@progbits,_Z10testKernelI11MyInt8ArrayEvPT_Pjj,comdat
	.weak	_Z10testKernelI11MyInt8ArrayEvPT_Pjj
	.p2align	3, 0x0
_Z10testKernelI11MyInt8ArrayEvPT_Pjj:
	.quad	_Z25__device_stub__testKernelI11MyInt8ArrayEvPT_Pjj
	.size	_Z10testKernelI11MyInt8ArrayEvPT_Pjj, 8

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z10testKernelI11MyInt8TupleEvPT_Pjj"
	.size	.L__unnamed_1, 37

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z10testKernelI11MyInt8ArrayEvPT_Pjj"
	.size	.L__unnamed_2, 37

	.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	"The program takes <num_elems> and eight indices as arguments!"
	.size	.Lstr, 62

	.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 _Z25__device_stub__testKernelI11MyInt8TupleEvPT_Pjj
	.addrsig_sym __gxx_personality_v0
	.addrsig_sym _Z25__device_stub__testKernelI11MyInt8ArrayEvPT_Pjj
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Unwind_Resume
	.addrsig_sym _Z10testKernelI11MyInt8TupleEvPT_Pjj
	.addrsig_sym _Z10testKernelI11MyInt8ArrayEvPT_Pjj
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.section	.text._Z10testKernelI11MyInt8TupleEvPT_Pjj,"axG",@progbits,_Z10testKernelI11MyInt8TupleEvPT_Pjj,comdat
	.protected	_Z10testKernelI11MyInt8TupleEvPT_Pjj ; -- Begin function _Z10testKernelI11MyInt8TupleEvPT_Pjj
	.globl	_Z10testKernelI11MyInt8TupleEvPT_Pjj
	.p2align	8
	.type	_Z10testKernelI11MyInt8TupleEvPT_Pjj,@function
_Z10testKernelI11MyInt8TupleEvPT_Pjj:   ; @_Z10testKernelI11MyInt8TupleEvPT_Pjj
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x24
	s_load_b32 s3, s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_u32_e64 s3, v1
	s_cbranch_execz .LBB0_32
; %bb.1:
	s_load_b128 s[0:3], s[0:1], 0x0
	v_and_b32_e32 v0, 7, v0
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_2)
	v_lshlrev_b32_e32 v0, 2, v0
	s_waitcnt lgkmcnt(0)
	global_load_b32 v0, v0, s[2:3]
	s_mov_b32 s2, 50
	s_clause 0x7
	scratch_store_b32 off, v2, off offset:4
	scratch_store_b32 off, v2, off offset:8
	scratch_store_b32 off, v2, off offset:12
	scratch_store_b32 off, v2, off offset:16
	scratch_store_b32 off, v2, off offset:20
	scratch_store_b32 off, v2, off offset:24
	scratch_store_b32 off, v2, off offset:28
	scratch_store_b32 off, v2, off offset:32
.LBB0_2:                                ; %NodeBlock29
                                        ; =>This Inner Loop Header: Depth=1
	s_mov_b32 s4, 0
	s_mov_b32 s3, exec_lo
                                        ; implicit-def: $vgpr2
	s_waitcnt vmcnt(0)
	v_cmpx_lt_i32_e32 2, v0
	s_xor_b32 s3, exec_lo, s3
	s_cbranch_execz .LBB0_18
; %bb.3:                                ; %NodeBlock27
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s5, 0
	s_mov_b32 s4, exec_lo
                                        ; implicit-def: $vgpr2
	v_cmpx_lt_i32_e32 4, v0
	s_xor_b32 s4, exec_lo, s4
	s_cbranch_execz .LBB0_11
; %bb.4:                                ; %NodeBlock25
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s6, 0
	s_mov_b32 s5, exec_lo
                                        ; implicit-def: $vgpr2
	v_cmpx_lt_i32_e32 5, v0
	s_xor_b32 s5, exec_lo, s5
	s_cbranch_execz .LBB0_8
; %bb.5:                                ; %LeafBlock23
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s6, -1
	s_mov_b32 s7, exec_lo
                                        ; implicit-def: $vgpr2
	v_cmpx_eq_u32_e32 6, v0
; %bb.6:                                ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v2, 28
	s_xor_b32 s6, exec_lo, -1
; %bb.7:                                ; %Flow37
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s7
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s6, s6, exec_lo
.LBB0_8:                                ; %Flow36
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_and_not1_saveexec_b32 s5, s5
; %bb.9:                                ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v2, 24
; %bb.10:                               ; %Flow38
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s5, s6, exec_lo
.LBB0_11:                               ; %Flow35
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_and_not1_saveexec_b32 s4, s4
	s_cbranch_execz .LBB0_17
; %bb.12:                               ; %NodeBlock21
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s6, exec_lo
                                        ; implicit-def: $vgpr2
	v_cmpx_lt_i32_e32 3, v0
	s_xor_b32 s6, exec_lo, s6
; %bb.13:                               ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v2, 20
; %bb.14:                               ; %Flow
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_and_not1_saveexec_b32 s6, s6
; %bb.15:                               ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v2, 16
; %bb.16:                               ; %Flow31
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s6
.LBB0_17:                               ; %Flow39
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s4
	s_and_b32 s4, s5, exec_lo
.LBB0_18:                               ; %Flow34
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_and_not1_saveexec_b32 s3, s3
	s_cbranch_execz .LBB0_28
; %bb.19:                               ; %NodeBlock19
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s5, exec_lo
                                        ; implicit-def: $vgpr2
	v_cmpx_lt_i32_e32 0, v0
	s_xor_b32 s5, exec_lo, s5
	s_cbranch_execz .LBB0_25
; %bb.20:                               ; %NodeBlock
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s6, exec_lo
                                        ; implicit-def: $vgpr2
	v_cmpx_lt_i32_e32 1, v0
	s_xor_b32 s6, exec_lo, s6
; %bb.21:                               ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v2, 12
; %bb.22:                               ; %Flow32
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_and_not1_saveexec_b32 s6, s6
; %bb.23:                               ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v2, 8
; %bb.24:                               ; %Flow33
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s6
.LBB0_25:                               ; %Flow41
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_saveexec_b32 s5, s5
	s_mov_b32 s6, s4
	s_xor_b32 exec_lo, exec_lo, s5
; %bb.26:                               ; %LeafBlock
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_cmp_ne_u32_e32 vcc_lo, 0, v0
	v_mov_b32_e32 v2, 4
	s_and_not1_b32 s6, s4, exec_lo
	s_and_b32 s7, vcc_lo, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s6, s6, s7
; %bb.27:                               ; %Flow42
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s4, s4, exec_lo
	s_and_b32 s5, s6, exec_lo
	s_or_b32 s4, s4, s5
.LBB0_28:                               ; %Flow40
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s3
	s_and_saveexec_b32 s3, s4
; %bb.29:                               ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v2, 32
; %bb.30:                               ; %_ZN11MyInt8TupleixEj.exit
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s3
	scratch_load_b32 v3, v2, off
	s_add_i32 s2, s2, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s2, 0
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v3, 1, v3
	scratch_store_b32 v2, v3, off
	s_cbranch_scc0 .LBB0_2
; %bb.31:
	s_clause 0x7
	scratch_load_b32 v3, off, off offset:4
	scratch_load_b32 v4, off, off offset:8
	scratch_load_b32 v5, off, off offset:12
	scratch_load_b32 v6, off, off offset:16
	scratch_load_b32 v7, off, off offset:20
	scratch_load_b32 v8, off, off offset:24
	scratch_load_b32 v9, off, off offset:28
	scratch_load_b32 v10, off, off offset:32
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 5, v[1:2]
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	s_waitcnt vmcnt(4)
	global_store_b128 v[0:1], v[3:6], off
	s_waitcnt vmcnt(0)
	global_store_b128 v[0:1], v[7:10], off offset:16
.LBB0_32:
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10testKernelI11MyInt8TupleEvPT_Pjj
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 36
		.amdhsa_kernarg_size 280
		.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 1
		.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 11
		.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
	.section	.text._Z10testKernelI11MyInt8TupleEvPT_Pjj,"axG",@progbits,_Z10testKernelI11MyInt8TupleEvPT_Pjj,comdat
.Lfunc_end0:
	.size	_Z10testKernelI11MyInt8TupleEvPT_Pjj, .Lfunc_end0-_Z10testKernelI11MyInt8TupleEvPT_Pjj
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 624
; NumSgprs: 18
; NumVgprs: 11
; ScratchSize: 36
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 11
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 1
; 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
	.section	.text._Z10testKernelI11MyInt8ArrayEvPT_Pjj,"axG",@progbits,_Z10testKernelI11MyInt8ArrayEvPT_Pjj,comdat
	.protected	_Z10testKernelI11MyInt8ArrayEvPT_Pjj ; -- Begin function _Z10testKernelI11MyInt8ArrayEvPT_Pjj
	.globl	_Z10testKernelI11MyInt8ArrayEvPT_Pjj
	.p2align	8
	.type	_Z10testKernelI11MyInt8ArrayEvPT_Pjj,@function
_Z10testKernelI11MyInt8ArrayEvPT_Pjj:   ; @_Z10testKernelI11MyInt8ArrayEvPT_Pjj
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[2:3], 0x24
	s_load_b32 s5, s[2:3], 0x10
	v_and_b32_e32 v3, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[1:2], null, s15, s4, v[3:4]
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_u32_e64 s5, v1
	s_cbranch_execz .LBB1_2
; %bb.1:
	s_load_b128 s[4:7], s[2:3], 0x0
	v_and_b32_e32 v2, 7, v3
	s_load_b64 s[0:1], s[0:1], 0x4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	v_lshlrev_b32_e32 v2, 2, v2
	s_waitcnt lgkmcnt(0)
	global_load_b32 v4, v2, s[6:7]
	s_lshr_b32 s0, s0, 16
	s_mul_i32 s0, s0, s1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v2, s0, v3
	v_bfe_u32 v3, v0, 10, 10
	v_bfe_u32 v0, v0, 20, 10
	v_mad_u32_u24 v2, v3, s1, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_lshl_u32 v0, v2, v0, 5
	v_mov_b32_e32 v2, 0
	ds_store_2addr_b32 v0, v2, v2 offset0:2 offset1:3
	ds_store_2addr_b32 v0, v2, v2 offset1:1
	ds_store_2addr_b32 v0, v2, v2 offset0:6 offset1:7
	ds_store_2addr_b32 v0, v2, v2 offset0:4 offset1:5
	s_waitcnt vmcnt(0)
	v_lshl_add_u32 v3, v4, 2, v0
	ds_load_b32 v4, v3
	s_waitcnt lgkmcnt(0)
	v_add_nc_u32_e32 v4, 50, v4
	ds_store_b32 v3, v4
	ds_load_2addr_b32 v[3:4], v0 offset0:4 offset1:5
	ds_load_2addr_b32 v[9:10], v0 offset0:2 offset1:3
	ds_load_2addr_b32 v[7:8], v0 offset1:1
	ds_load_2addr_b32 v[5:6], v0 offset0:6 offset1:7
	v_lshlrev_b64 v[0:1], 5, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s4, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s5, v1, vcc_lo
	s_waitcnt lgkmcnt(1)
	global_store_b128 v[0:1], v[7:10], off
	s_waitcnt lgkmcnt(0)
	global_store_b128 v[0:1], v[3:6], off offset:16
.LBB1_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10testKernelI11MyInt8ArrayEvPT_Pjj
		.amdhsa_group_segment_fixed_size 32768
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.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 11
		.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
	.section	.text._Z10testKernelI11MyInt8ArrayEvPT_Pjj,"axG",@progbits,_Z10testKernelI11MyInt8ArrayEvPT_Pjj,comdat
.Lfunc_end1:
	.size	_Z10testKernelI11MyInt8ArrayEvPT_Pjj, .Lfunc_end1-_Z10testKernelI11MyInt8ArrayEvPT_Pjj
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 332
; NumSgprs: 18
; NumVgprs: 11
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 32768 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 11
; 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	__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:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10testKernelI11MyInt8TupleEvPT_Pjj
    .private_segment_fixed_size: 36
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10testKernelI11MyInt8TupleEvPT_Pjj.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     11
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 32768
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10testKernelI11MyInt8ArrayEvPT_Pjj
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10testKernelI11MyInt8ArrayEvPT_Pjj.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     11
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
