/*
Authors: Jose Garcia Kameron Bush
Collatz code for CS 4380 / CS 5351

Copyright (c) 2019 Texas State University. All rights reserved.

Redistribution in source or binary form, with or without modification,
is *not* permitted. Use in source and binary forms, with or without
modification, is only permitted for academic use in CS 4380 or CS 5351
at Texas State University.

THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR
ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
(INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON
ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

Author: Martin Burtscher
*/

#include <cstdio>
#include <cuda.h>

static const int ThreadsPerBlock = 512;

static int* d_maxlen;

static __global__ void collatz(const long start, const long stop, int* const maxlen)
{
  // todo: process odd values from start (assume start to be odd) to stop (inclusively if stop is odd) with one thread per value (based on code from previous project
  // compute sequence lengths
  //*********has to be only odds*************
  //int maxlen = 0;
  const int i = (threadIdx.x + blockIdx.x * (long)blockDim.x) * 2 + start;
  if (i <= stop) {
    long val = i;
    int len = 1;
    while (val != 1) {
      len++;
      if ((val % 2) == 0) {
        val = val / 2;  // even
      } else {
        val = 3 * val + 1;  // odd
      }
    }
     if (*maxlen < len) {
       atomicMax(maxlen, len);
     }
	  }
  }


void GPU_Init()
{
  int maxlen = 0;
  if (cudaSuccess != cudaMalloc((void **)&d_maxlen, sizeof(int))) {fprintf(stderr, "ERROR: could not allocate memory\n"); exit(-1);}
  if (cudaSuccess != cudaMemcpy(d_maxlen, &maxlen, sizeof(int), cudaMemcpyHostToDevice)) {fprintf(stderr, "ERROR: copying to device failed\n"); exit(-1);}
}

void GPU_Exec(const long start, const long stop)
{
  if (start <= stop) {
    collatz<<<((stop - start + 2) / 2 + ThreadsPerBlock - 1) / ThreadsPerBlock, ThreadsPerBlock>>>(start, stop, d_maxlen);
  }
}

int GPU_Fini()
{
  int maxlen;

  // todo: copy the result from the device to the host and free the device memory
  if (cudaSuccess != cudaMemcpy(&maxlen, d_maxlen, sizeof(int), cudaMemcpyDeviceToHost)) {
    fprintf(stderr, "ERROR: copying from device failed\n"); exit(-1);
  }
  cudaFree(d_maxlen);

  return maxlen;
}

	.file	"tmpxft_002dc536_00000000-6_collatz_hyb.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2032:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2032:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_ZL28__device_stub__Z7collatzllPillPi, @function
_ZL28__device_stub__Z7collatzllPillPi:
.LFB2054:
	.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
	movq	%rdx, 8(%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	8(%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	_ZL7collatzllPi(%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
.LFE2054:
	.size	_ZL28__device_stub__Z7collatzllPillPi, .-_ZL28__device_stub__Z7collatzllPillPi
	.type	_ZL7collatzllPi, @function
_ZL7collatzllPi:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_ZL28__device_stub__Z7collatzllPillPi
	.cfi_endproc
.LFE2055:
	.size	_ZL7collatzllPi, .-_ZL7collatzllPi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"ERROR: could not allocate memory\n"
.LC1:
	.string	"ERROR: copying to device failed\n"
	.text
	.globl	_Z8GPU_Initv
	.type	_Z8GPU_Initv, @function
_Z8GPU_Initv:
.LFB2027:
	.cfi_startproc
	endbr64
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$4, %esi
	leaq	_ZL8d_maxlen(%rip), %rdi
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	xorl	%eax, %eax
	movl	%eax, 4(%rsp)
	call	cudaMalloc@PLT
	leaq	.LC0(%rip), %rdx
	testl	%eax, %eax
	jne	.L13
	leaq	4(%rsp), %rsi
	movl	$1, %ecx
	movl	$4, %edx
	movq	_ZL8d_maxlen(%rip), %rdi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L8
	leaq	.LC1(%rip), %rdx
.L13:
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	orl	$-1, %edi
	call	exit@PLT
.L8:
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L11
	call	__stack_chk_fail@PLT
.L11:
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z8GPU_Initv, .-_Z8GPU_Initv
	.globl	_Z8GPU_Execll
	.type	_Z8GPU_Execll, @function
_Z8GPU_Execll:
.LFB2028:
	.cfi_startproc
	endbr64
	cmpq	%rsi, %rdi
	jg	.L18
	movq	%rsi, %rax
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	$8388609, %edx
	movq	%rsi, %rbp
	subq	%rdi, %rax
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	xorl	%r9d, %r9d
	addq	$2, %rax
	xorl	%r8d, %r8d
	salq	$9, %rdx
	movl	$1, %ecx
	sarq	%rax
	addq	$511, %rax
	subq	$40, %rsp
	.cfi_def_cfa_offset 64
	sarq	$9, %rax
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L14
	movq	_ZL8d_maxlen(%rip), %rdx
	addq	$40, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 24
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_restore 3
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_restore 6
	.cfi_def_cfa_offset 8
	jmp	_ZL28__device_stub__Z7collatzllPillPi
.L14:
	.cfi_restore_state
	addq	$40, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
.L18:
	.cfi_restore 3
	.cfi_restore 6
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z8GPU_Execll, .-_Z8GPU_Execll
	.section	.rodata.str1.1
.LC2:
	.string	"ERROR: copying from device failed\n"
	.text
	.globl	_Z8GPU_Finiv
	.type	_Z8GPU_Finiv, @function
_Z8GPU_Finiv:
.LFB2029:
	.cfi_startproc
	endbr64
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movq	_ZL8d_maxlen(%rip), %rsi
	movl	$2, %ecx
	movl	$4, %edx
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	leaq	4(%rsp), %rdi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L22
	movq	stderr(%rip), %rdi
	leaq	.LC2(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	orl	$-1, %edi
	call	exit@PLT
.L22:
	movq	_ZL8d_maxlen(%rip), %rdi
	call	cudaFree@PLT
	movl	4(%rsp), %eax
	movq	8(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L23
	call	__stack_chk_fail@PLT
.L23:
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z8GPU_Finiv, .-_Z8GPU_Finiv
	.section	.rodata.str1.1
.LC3:
	.string	"_Z7collatzllPi"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2057:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC3(%rip), %rdx
	movq	%rax, %rdi
	leaq	_ZL7collatzllPi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2057:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.local	_ZL8d_maxlen
	.comm	_ZL8d_maxlen,8,8
	.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:

/*
Authors: Jose Garcia Kameron Bush
Collatz code for CS 4380 / CS 5351

Copyright (c) 2019 Texas State University. All rights reserved.

Redistribution in source or binary form, with or without modification,
is *not* permitted. Use in source and binary forms, with or without
modification, is only permitted for academic use in CS 4380 or CS 5351
at Texas State University.

THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR
ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
(INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON
ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

Author: Martin Burtscher
*/

#include <cstdio>
#include <hip/hip_runtime.h>

static const int ThreadsPerBlock = 512;

static int* d_maxlen;

static __global__ void collatz(const long start, const long stop, int* const maxlen)
{
  // todo: process odd values from start (assume start to be odd) to stop (inclusively if stop is odd) with one thread per value (based on code from previous project
  // compute sequence lengths
  //*********has to be only odds*************
  //int maxlen = 0;
  const int i = (threadIdx.x + blockIdx.x * (long)blockDim.x) * 2 + start;
  if (i <= stop) {
    long val = i;
    int len = 1;
    while (val != 1) {
      len++;
      if ((val % 2) == 0) {
        val = val / 2;  // even
      } else {
        val = 3 * val + 1;  // odd
      }
    }
     if (*maxlen < len) {
       atomicMax(maxlen, len);
     }
	  }
  }


void GPU_Init()
{
  int maxlen = 0;
  if (hipSuccess != hipMalloc((void **)&d_maxlen, sizeof(int))) {fprintf(stderr, "ERROR: could not allocate memory\n"); exit(-1);}
  if (hipSuccess != hipMemcpy(d_maxlen, &maxlen, sizeof(int), hipMemcpyHostToDevice)) {fprintf(stderr, "ERROR: copying to device failed\n"); exit(-1);}
}

void GPU_Exec(const long start, const long stop)
{
  if (start <= stop) {
    collatz<<<((stop - start + 2) / 2 + ThreadsPerBlock - 1) / ThreadsPerBlock, ThreadsPerBlock>>>(start, stop, d_maxlen);
  }
}

int GPU_Fini()
{
  int maxlen;

  // todo: copy the result from the device to the host and free the device memory
  if (hipSuccess != hipMemcpy(&maxlen, d_maxlen, sizeof(int), hipMemcpyDeviceToHost)) {
    fprintf(stderr, "ERROR: copying from device failed\n"); exit(-1);
  }
  hipFree(d_maxlen);

  return maxlen;
}

	.text
	.file	"collatz_hyb.hip"
	.globl	_Z8GPU_Initv                    # -- Begin function _Z8GPU_Initv
	.type	_Z8GPU_Initv,@function
_Z8GPU_Initv:                           # @_Z8GPU_Initv
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movl	$0, 4(%rsp)
	movl	$_ZL8d_maxlen, %edi
	movl	$4, %esi
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB0_1
# %bb.3:
	movq	_ZL8d_maxlen(%rip), %rdi
	leaq	4(%rsp), %rsi
	movl	$4, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB0_4
# %bb.5:
	popq	%rax
	.cfi_def_cfa_offset 8
	retq
.LBB0_1:
	.cfi_def_cfa_offset 16
	movq	stderr(%rip), %rcx
	movl	$.L.str, %edi
	movl	$33, %esi
	jmp	.LBB0_2
.LBB0_4:
	movq	stderr(%rip), %rcx
	movl	$.L.str.1, %edi
	movl	$32, %esi
.LBB0_2:
	movl	$1, %edx
	callq	fwrite@PLT
	movl	$-1, %edi
	callq	exit
.Lfunc_end0:
	.size	_Z8GPU_Initv, .Lfunc_end0-_Z8GPU_Initv
	.cfi_endproc
                                        # -- End function
	.globl	_Z8GPU_Execll                   # -- Begin function _Z8GPU_Execll
	.type	_Z8GPU_Execll,@function
_Z8GPU_Execll:                          # @_Z8GPU_Execll
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rsi, %rax
	subq	%rdi, %rax
	jl	.LBB1_2
# %bb.1:
	movq	%rsi, %rbx
	movq	%rdi, %r14
	leaq	2(%rax), %rcx
	shrq	$63, %rcx
	addq	%rcx, %rax
	addq	$2, %rax
	sarq	%rax
	leaq	511(%rax), %rcx
	addq	$1022, %rax                     # imm = 0x3FE
	testq	%rcx, %rcx
	cmovnsq	%rcx, %rax
	shrq	$9, %rax
	movl	%eax, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$512, %rdx                      # imm = 0x200
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB1_3
.LBB1_2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.LBB1_3:
	.cfi_def_cfa_offset 32
	movq	_ZL8d_maxlen(%rip), %rdx
	movq	%r14, %rdi
	movq	%rbx, %rsi
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	_ZL22__device_stub__collatzllPi # TAILCALL
.Lfunc_end1:
	.size	_Z8GPU_Execll, .Lfunc_end1-_Z8GPU_Execll
	.cfi_endproc
                                        # -- End function
	.type	_ZL22__device_stub__collatzllPi,@function # -- Begin function _ZL22__device_stub__collatzllPi
_ZL22__device_stub__collatzllPi:        # @_ZL22__device_stub__collatzllPi
	.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	24(%rsp), %rsi
	movq	%rdx, (%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	16(%rsp), %r12
	leaq	8(%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	$_ZL7collatzllPi, %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_end2:
	.size	_ZL22__device_stub__collatzllPi, .Lfunc_end2-_ZL22__device_stub__collatzllPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z8GPU_Finiv                    # -- Begin function _Z8GPU_Finiv
	.type	_Z8GPU_Finiv,@function
_Z8GPU_Finiv:                           # @_Z8GPU_Finiv
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	_ZL8d_maxlen(%rip), %rsi
	leaq	4(%rsp), %rdi
	movl	$4, %edx
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB3_2
# %bb.1:
	movq	_ZL8d_maxlen(%rip), %rdi
	callq	hipFree
	movl	4(%rsp), %eax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.LBB3_2:
	.cfi_def_cfa_offset 16
	movq	stderr(%rip), %rcx
	movl	$.L.str.2, %edi
	movl	$34, %esi
	movl	$1, %edx
	callq	fwrite@PLT
	movl	$-1, %edi
	callq	exit
.Lfunc_end3:
	.size	_Z8GPU_Finiv, .Lfunc_end3-_Z8GPU_Finiv
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_ZL7collatzllPi, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__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	.LBB5_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
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_ZL8d_maxlen,@object            # @_ZL8d_maxlen
	.local	_ZL8d_maxlen
	.comm	_ZL8d_maxlen,8,8
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"ERROR: could not allocate memory\n"
	.size	.L.str, 34

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"ERROR: copying to device failed\n"
	.size	.L.str.1, 33

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

	.type	.L.str.2,@object                # @.str.2
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.2:
	.asciz	"ERROR: copying from device failed\n"
	.size	.L.str.2, 35

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_ZL7collatzllPi"
	.size	.L__unnamed_1, 16

	.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

	.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 _ZL22__device_stub__collatzllPi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _ZL8d_maxlen
	.addrsig_sym _ZL7collatzllPi
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.section	.text._ZL7collatzllPi,"axG",@progbits,_ZL7collatzllPi,comdat
	.globl	_ZL7collatzllPi                 ; -- Begin function _ZL7collatzllPi
	.p2align	8
	.type	_ZL7collatzllPi,@function
_ZL7collatzllPi:                        ; @_ZL7collatzllPi
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x24
	s_load_b128 s[4:7], s[0:1], 0x0
	v_add_co_u32 v2, vcc_lo, 0, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s2, s2, s15
	v_add_lshl_u32 v0, s2, v0, 1
	s_mov_b32 s2, 0
	s_mov_b32 s3, 1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v3, vcc_lo, s4, v0, vcc_lo
	s_mov_b32 s4, exec_lo
	v_ashrrev_i32_e32 v1, 31, v3
	v_mov_b32_e32 v0, v3
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_ge_i64_e64 s[6:7], v[0:1]
	s_cbranch_execz .LBB0_14
; %bb.1:                                ; %.preheader
	s_load_b64 s[0:1], s[0:1], 0x10
	v_cmp_ne_u64_e32 vcc_lo, s[2:3], v[2:3]
	v_mov_b32_e32 v2, 1
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB0_9
; %bb.2:                                ; %.lr.ph.preheader
	s_mov_b32 s4, 1
	s_mov_b32 s3, 0
.LBB0_3:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	v_and_b32_e32 v2, 1, v0
	s_mov_b32 s5, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 1, v2
	s_xor_b32 s5, exec_lo, s5
; %bb.4:                                ;   in Loop: Header=BB0_3 Depth=1
	v_mad_u64_u32 v[2:3], null, v0, 3, 1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v0, v3
	v_mad_u64_u32 v[3:4], null, v1, 3, v[0:1]
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mov_b32 v0, v2 :: v_dual_mov_b32 v1, v3
; %bb.5:                                ; %Flow29
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_and_not1_saveexec_b32 s5, s5
; %bb.6:                                ;   in Loop: Header=BB0_3 Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_ashrrev_i64 v[0:1], 1, v[0:1]
; %bb.7:                                ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_eq_u64_e32 vcc_lo, 1, v[0:1]
	s_add_i32 s4, s4, 1
	v_mov_b32_e32 v2, s4
	s_or_b32 s3, vcc_lo, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s3
	s_cbranch_execnz .LBB0_3
; %bb.8:                                ; %Flow30
	s_or_b32 exec_lo, exec_lo, s3
.LBB0_9:                                ; %Flow31
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s2
	s_waitcnt lgkmcnt(0)
	s_load_b32 s2, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, s2, v2
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_14
; %bb.10:
	s_mov_b32 s3, exec_lo
	s_brev_b32 s2, 1
.LBB0_11:                               ; %ComputeLoop
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s4, s3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s5, v2, s4
	s_lshl_b32 s4, 1, s4
	s_and_not1_b32 s3, s3, s4
	s_delay_alu instid0(VALU_DEP_1)
	s_max_i32 s2, s2, s5
	s_cmp_lg_u32 s3, 0
	s_cbranch_scc1 .LBB0_11
; %bb.12:                               ; %ComputeEnd
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s3, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s3, exec_lo, s3
	s_cbranch_execz .LBB0_14
; %bb.13:
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2
	global_atomic_max_i32 v0, v1, s[0:1]
.LBB0_14:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _ZL7collatzllPi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.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 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 5
		.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._ZL7collatzllPi,"axG",@progbits,_ZL7collatzllPi,comdat
.Lfunc_end0:
	.size	_ZL7collatzllPi, .Lfunc_end0-_ZL7collatzllPi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 384
; NumSgprs: 18
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 5
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.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:
      - .offset:         0
        .size:           8
        .value_kind:     by_value
      - .offset:         8
        .size:           8
        .value_kind:     by_value
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .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:           _ZL7collatzllPi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _ZL7collatzllPi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
