/*
Collatz code for CS 4380 / CS 5351

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

Redistribution in source or binary form, with or without modification,
is *not* permitted. Use in source or binary form, 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

Bryan Valenzuela
Connor Steed
*/

#include <cstdio>
#include <cuda.h>
#include <algorithm>
#include <sys/time.h>

static const int ThreadsPerBlock = 512;

static __global__ void collatz(const long start, const long bound, const long step, int* const maxlen)
{
  const long i = threadIdx.x + blockIdx.x * (long)blockDim.x + start;

  if( (i - start) % step == 0)
  if (i < bound){
    long val = i;
    int len = 1;
    while (val != 1) {
      len++;
      if ((val % 2) == 0) {
        val /= 2;  // even
      } else {
        val = 3 * val + 1;  // odd
      }
    }
      atomicMax(maxlen, len);
  }

}

static void CheckCuda()
{
  cudaError_t e;
  cudaDeviceSynchronize();
  if (cudaSuccess != (e = cudaGetLastError())) {
    fprintf(stderr, "CUDA error %d: %s\n", e, cudaGetErrorString(e));
    exit(-1);
  }
}

int main(int argc, char *argv[])
{
  printf("Collatz v1.5\n");

  // check command line
  if (argc != 4) {fprintf(stderr, "USAGE: %s start bound step\n", argv[0]); exit(-1);}
  const long start = atol(argv[1]);
  const long bound = atol(argv[2]);
  const long step = atol(argv[3]);
  if (start < 1) {fprintf(stderr, "ERROR: start value must be at least 1\n"); exit(-1);}
  if (bound <= start) {fprintf(stderr, "ERROR: bound must be larger than start\n"); exit(-1);}
  if (step < 1) {fprintf(stderr, "ERROR: step size must be at least 1\n"); exit(-1);}
  printf("start value: %ld\n", start);
  printf("upper bound: %ld\n", bound);
  printf("step size: %ld\n", step);

  int maxlen = 0;
  int size = sizeof(int);

  int* d_maxlen;

  if (cudaSuccess != cudaMalloc((void **)&d_maxlen, size)) {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);}
  // start time
  timeval beg, end;
  gettimeofday(&beg, NULL);

  // execute timed code
  collatz<<<(bound + ThreadsPerBlock - 1) / ThreadsPerBlock, ThreadsPerBlock>>>(start, bound, step, d_maxlen);
  cudaDeviceSynchronize();

  // end time
  gettimeofday(&end, NULL);
  const double runtime = end.tv_sec - beg.tv_sec + (end.tv_usec - beg.tv_usec) / 1000000.0;
  printf("compute time: %.6f s\n", runtime);

   CheckCuda();
  if (cudaSuccess != cudaMemcpy(&maxlen, d_maxlen, sizeof(int), cudaMemcpyDeviceToHost)) {fprintf(stderr, "ERROR: copying from device failed\n"); exit(-1);}


  // print result
  printf("maximum sequence length: %d elements\n", maxlen);

  cudaFree(d_maxlen);
  return 0;
}

	.file	"tmpxft_003a38dd_00000000-6_collatz_cuda.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2311:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2311:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_ZL29__device_stub__Z7collatzlllPilllPi, @function
_ZL29__device_stub__Z7collatzlllPilllPi:
.LFB2333:
	.cfi_startproc
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_ZL7collatzlllPi(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L2:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2333:
	.size	_ZL29__device_stub__Z7collatzlllPilllPi, .-_ZL29__device_stub__Z7collatzlllPilllPi
	.type	_ZL7collatzlllPi, @function
_ZL7collatzlllPi:
.LFB2334:
	.cfi_startproc
	endbr64
	jmp	_ZL29__device_stub__Z7collatzlllPilllPi
	.cfi_endproc
.LFE2334:
	.size	_ZL7collatzlllPi, .-_ZL7collatzlllPi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Collatz v1.5\n"
.LC1:
	.string	"USAGE: %s start bound step\n"
.LC2:
	.string	"ERROR: start value must be at least 1\n"
.LC3:
	.string	"ERROR: bound must be larger than start\n"
.LC4:
	.string	"ERROR: step size must be at least 1\n"
.LC5:
	.string	"start value: %ld\n"
.LC6:
	.string	"upper bound: %ld\n"
.LC7:
	.string	"step size: %ld\n"
.LC8:
	.string	"ERROR: could not allocate memory\n"
.LC9:
	.string	"ERROR: copying to device failed\n"
.LC11:
	.string	"compute time: %.6f s\n"
.LC12:
	.string	"CUDA error %d: %s\n"
.LC13:
	.string	"ERROR: copying from device failed\n"
.LC14:
	.string	"maximum sequence length: %d elements\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2308:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rsi, %r12
	leaq	.LC0(%rip), %rsi
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	%edi, %ebx
	movl	$2, %edi
	subq	$88, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	cmpl	$4, %ebx
	je	.L9
	movq	(%r12), %rcx
	movq	stderr(%rip), %rdi
	leaq	.LC1(%rip), %rdx
	xorl	%eax, %eax
	movl	$2, %esi
	call	__fprintf_chk@PLT
	jmp	.L20
.L9:
	movq	8(%r12), %rdi
	call	atol@PLT
	movq	16(%r12), %rdi
	movq	%rax, %rbx
	call	atol@PLT
	movq	24(%r12), %rdi
	movq	%rax, %rbp
	call	atol@PLT
	movq	%rax, %r13
	testq	%rbx, %rbx
	jg	.L10
	leaq	.LC2(%rip), %rdx
.L21:
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L20:
	orl	$-1, %edi
	call	exit@PLT
.L10:
	leaq	.LC3(%rip), %rdx
	cmpq	%rbp, %rbx
	jge	.L21
	leaq	.LC4(%rip), %rdx
	testq	%rax, %rax
	jle	.L21
	movq	%rbx, %rdx
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbp, %rdx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC6(%rip), %rsi
	call	__printf_chk@PLT
	movq	%r13, %rdx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC7(%rip), %rsi
	call	__printf_chk@PLT
	xorl	%eax, %eax
	leaq	16(%rsp), %rdi
	movl	$4, %esi
	movl	%eax, 12(%rsp)
	call	cudaMalloc@PLT
	leaq	.LC8(%rip), %rdx
	testl	%eax, %eax
	jne	.L21
	movq	16(%rsp), %rdi
	leaq	12(%rsp), %r12
	movl	$4, %edx
	movl	$1, %ecx
	movq	%r12, %rsi
	call	cudaMemcpy@PLT
	leaq	.LC9(%rip), %rdx
	testl	%eax, %eax
	jne	.L21
	xorl	%esi, %esi
	leaq	40(%rsp), %rdi
	call	gettimeofday@PLT
	movl	$8388609, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	leaq	511(%rbp), %rax
	salq	$9, %rdx
	movl	$1, %ecx
	movl	$1, 64(%rsp)
	sarq	$9, %rax
	movq	%rdx, 56(%rsp)
	movl	%eax, 28(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	28(%rsp), %rdi
	movl	36(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L15
	movq	16(%rsp), %rcx
	movq	%r13, %rdx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_ZL29__device_stub__Z7collatzlllPilllPi
.L15:
	call	cudaDeviceSynchronize@PLT
	leaq	56(%rsp), %rdi
	xorl	%esi, %esi
	call	gettimeofday@PLT
	movq	64(%rsp), %rax
	subq	48(%rsp), %rax
	leaq	.LC11(%rip), %rsi
	cvtsi2sdq	%rax, %xmm0
	movq	56(%rsp), %rax
	subq	40(%rsp), %rax
	divsd	.LC10(%rip), %xmm0
	cvtsi2sdq	%rax, %xmm1
	movl	$2, %edi
	movb	$1, %al
	addsd	%xmm1, %xmm0
	call	__printf_chk@PLT
	call	cudaDeviceSynchronize@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L16
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	movq	stderr(%rip), %rdi
	movl	%ebx, %ecx
	leaq	.LC12(%rip), %rdx
	movq	%rax, %r8
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	jmp	.L20
.L16:
	movq	16(%rsp), %rsi
	movl	$4, %edx
	movl	$2, %ecx
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	leaq	.LC13(%rip), %rdx
	testl	%eax, %eax
	jne	.L21
	movl	12(%rsp), %edx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC14(%rip), %rsi
	call	__printf_chk@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L18
	call	__stack_chk_fail@PLT
.L18:
	addq	$88, %rsp
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2308:
	.size	main, .-main
	.section	.rodata.str1.1
.LC15:
	.string	"_Z7collatzlllPi"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2336:
	.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	.LC15(%rip), %rdx
	movq	%rax, %rdi
	leaq	_ZL7collatzlllPi(%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
.LFE2336:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC10:
	.long	0
	.long	1093567616
	.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:

/*
Collatz code for CS 4380 / CS 5351

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

Redistribution in source or binary form, with or without modification,
is *not* permitted. Use in source or binary form, 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

Bryan Valenzuela
Connor Steed
*/

#include <cstdio>
#include <hip/hip_runtime.h>
#include <algorithm>
#include <sys/time.h>

static const int ThreadsPerBlock = 512;

static __global__ void collatz(const long start, const long bound, const long step, int* const maxlen)
{
  const long i = threadIdx.x + blockIdx.x * (long)blockDim.x + start;

  if( (i - start) % step == 0)
  if (i < bound){
    long val = i;
    int len = 1;
    while (val != 1) {
      len++;
      if ((val % 2) == 0) {
        val /= 2;  // even
      } else {
        val = 3 * val + 1;  // odd
      }
    }
      atomicMax(maxlen, len);
  }

}

static void CheckCuda()
{
  hipError_t e;
  hipDeviceSynchronize();
  if (hipSuccess != (e = hipGetLastError())) {
    fprintf(stderr, "CUDA error %d: %s\n", e, hipGetErrorString(e));
    exit(-1);
  }
}

int main(int argc, char *argv[])
{
  printf("Collatz v1.5\n");

  // check command line
  if (argc != 4) {fprintf(stderr, "USAGE: %s start bound step\n", argv[0]); exit(-1);}
  const long start = atol(argv[1]);
  const long bound = atol(argv[2]);
  const long step = atol(argv[3]);
  if (start < 1) {fprintf(stderr, "ERROR: start value must be at least 1\n"); exit(-1);}
  if (bound <= start) {fprintf(stderr, "ERROR: bound must be larger than start\n"); exit(-1);}
  if (step < 1) {fprintf(stderr, "ERROR: step size must be at least 1\n"); exit(-1);}
  printf("start value: %ld\n", start);
  printf("upper bound: %ld\n", bound);
  printf("step size: %ld\n", step);

  int maxlen = 0;
  int size = sizeof(int);

  int* d_maxlen;

  if (hipSuccess != hipMalloc((void **)&d_maxlen, size)) {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);}
  // start time
  timeval beg, end;
  gettimeofday(&beg, NULL);

  // execute timed code
  collatz<<<(bound + ThreadsPerBlock - 1) / ThreadsPerBlock, ThreadsPerBlock>>>(start, bound, step, d_maxlen);
  hipDeviceSynchronize();

  // end time
  gettimeofday(&end, NULL);
  const double runtime = end.tv_sec - beg.tv_sec + (end.tv_usec - beg.tv_usec) / 1000000.0;
  printf("compute time: %.6f s\n", runtime);

   CheckCuda();
  if (hipSuccess != hipMemcpy(&maxlen, d_maxlen, sizeof(int), hipMemcpyDeviceToHost)) {fprintf(stderr, "ERROR: copying from device failed\n"); exit(-1);}


  // print result
  printf("maximum sequence length: %d elements\n", maxlen);

  hipFree(d_maxlen);
  return 0;
}

	.text
	.file	"collatz_cuda.hip"
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI0_0:
	.quad	0x412e848000000000              # double 1.0E+6
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$48, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %r15
	movl	%edi, %ebx
	movl	$.Lstr, %edi
	callq	puts@PLT
	cmpl	$4, %ebx
	jne	.LBB0_11
# %bb.1:
	movq	8(%r15), %rdi
	callq	atol
	movq	%rax, %rbx
	movq	16(%r15), %rdi
	callq	atol
	movq	%rax, %r14
	movq	24(%r15), %rdi
	callq	atol
	testq	%rbx, %rbx
	jle	.LBB0_12
# %bb.2:
	cmpq	%rbx, %r14
	jle	.LBB0_13
# %bb.3:
	movq	%rax, %r15
	testq	%rax, %rax
	jle	.LBB0_14
# %bb.4:
	movl	$.L.str.5, %edi
	movq	%rbx, %rsi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.6, %edi
	movq	%r14, %rsi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.7, %edi
	movq	%r15, %rsi
	xorl	%eax, %eax
	callq	printf
	movl	$0, 4(%rsp)
	leaq	8(%rsp), %rdi
	movl	$4, %esi
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB0_15
# %bb.5:
	movq	8(%rsp), %rdi
	leaq	4(%rsp), %rsi
	movl	$4, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB0_16
# %bb.6:
	leaq	16(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	%r14, %rax
	addq	$511, %rax                      # imm = 0x1FF
	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
	jne	.LBB0_8
# %bb.7:
	movq	8(%rsp), %rcx
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	callq	_ZL22__device_stub__collatzlllPi
.LBB0_8:
	callq	hipDeviceSynchronize
	leaq	32(%rsp), %rbx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%rbx), %rax
	movq	8(%rbx), %rcx
	subq	16(%rsp), %rax
	cvtsi2sd	%rax, %xmm1
	subq	24(%rsp), %rcx
	cvtsi2sd	%rcx, %xmm0
	divsd	.LCPI0_0(%rip), %xmm0
	addsd	%xmm1, %xmm0
	movl	$.L.str.10, %edi
	movb	$1, %al
	callq	printf
	callq	hipDeviceSynchronize
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB0_17
# %bb.9:                                # %_ZL9CheckCudav.exit
	movq	8(%rsp), %rsi
	leaq	4(%rsp), %rdi
	movl	$4, %edx
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB0_18
# %bb.10:
	movl	4(%rsp), %esi
	movl	$.L.str.12, %edi
	xorl	%eax, %eax
	callq	printf
	movq	8(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$48, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB0_11:
	.cfi_def_cfa_offset 80
	movq	stderr(%rip), %rdi
	movq	(%r15), %rdx
	movl	$.L.str.1, %esi
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB0_20
.LBB0_12:
	movq	stderr(%rip), %rcx
	movl	$.L.str.2, %edi
	movl	$38, %esi
	jmp	.LBB0_19
.LBB0_13:
	movq	stderr(%rip), %rcx
	movl	$.L.str.3, %edi
	movl	$39, %esi
	jmp	.LBB0_19
.LBB0_14:
	movq	stderr(%rip), %rcx
	movl	$.L.str.4, %edi
	movl	$36, %esi
	jmp	.LBB0_19
.LBB0_15:
	movq	stderr(%rip), %rcx
	movl	$.L.str.8, %edi
	movl	$33, %esi
	jmp	.LBB0_19
.LBB0_16:
	movq	stderr(%rip), %rcx
	movl	$.L.str.9, %edi
	movl	$32, %esi
	jmp	.LBB0_19
.LBB0_17:
	movl	%eax, %ebx
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.13, %esi
	movq	%r14, %rdi
	movl	%ebx, %edx
	movq	%rax, %rcx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB0_20
.LBB0_18:
	movq	stderr(%rip), %rcx
	movl	$.L.str.11, %edi
	movl	$34, %esi
.LBB0_19:
	movl	$1, %edx
	callq	fwrite@PLT
.LBB0_20:
	movl	$-1, %edi
	callq	exit
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.type	_ZL22__device_stub__collatzlllPi,@function # -- Begin function _ZL22__device_stub__collatzlllPi
_ZL22__device_stub__collatzlllPi:       # @_ZL22__device_stub__collatzlllPi
	.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), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	16(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_ZL7collatzlllPi, %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_end1:
	.size	_ZL22__device_stub__collatzlllPi, .Lfunc_end1-_ZL22__device_stub__collatzlllPi
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_ZL7collatzlllPi, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"USAGE: %s start bound step\n"
	.size	.L.str.1, 28

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"ERROR: start value must be at least 1\n"
	.size	.L.str.2, 39

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"ERROR: bound must be larger than start\n"
	.size	.L.str.3, 40

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"ERROR: step size must be at least 1\n"
	.size	.L.str.4, 37

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"start value: %ld\n"
	.size	.L.str.5, 18

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"upper bound: %ld\n"
	.size	.L.str.6, 18

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"step size: %ld\n"
	.size	.L.str.7, 16

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"ERROR: could not allocate memory\n"
	.size	.L.str.8, 34

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

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

	.type	.L.str.10,@object               # @.str.10
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.10:
	.asciz	"compute time: %.6f s\n"
	.size	.L.str.10, 22

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"ERROR: copying from device failed\n"
	.size	.L.str.11, 35

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"maximum sequence length: %d elements\n"
	.size	.L.str.12, 38

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"CUDA error %d: %s\n"
	.size	.L.str.13, 19

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_ZL7collatzlllPi"
	.size	.L__unnamed_1, 17

	.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	"Collatz v1.5"
	.size	.Lstr, 13

	.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__collatzlllPi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _ZL7collatzlllPi
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.section	.text._ZL7collatzlllPi,"axG",@progbits,_ZL7collatzlllPi,comdat
	.globl	_ZL7collatzlllPi                ; -- Begin function _ZL7collatzlllPi
	.p2align	8
	.type	_ZL7collatzlllPi,@function
_ZL7collatzlllPi:                       ; @_ZL7collatzlllPi
; %bb.0:
	s_clause 0x1
	s_load_b256 s[4:11], s[0:1], 0x0
	s_load_b32 s1, s[0:1], 0x2c
	s_waitcnt lgkmcnt(0)
	s_ashr_i32 s0, s9, 31
	s_and_b32 s12, s1, 0xffff
	s_add_u32 s2, s8, s0
	s_mov_b32 s1, s0
	s_addc_u32 s3, s9, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_xor_b64 s[2:3], s[2:3], s[0:1]
	v_cvt_f32_u32_e32 v1, s2
	v_cvt_f32_u32_e32 v2, s3
	s_sub_u32 s0, 0, s2
	s_subb_u32 s1, 0, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmamk_f32 v1, v2, 0x4f800000, v1
	v_rcp_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x5f7ffffc, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v2, 0x2f800000, v1
	v_trunc_f32_e32 v2, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmamk_f32 v1, v2, 0xcf800000, v1
	v_cvt_u32_f32_e32 v2, v2
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v3, s0, v2
	v_mul_hi_u32 v4, s0, v1
	v_mul_lo_u32 v5, s1, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v3, v4, v3
	v_mul_lo_u32 v4, s0, v1
	v_add_nc_u32_e32 v3, v3, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v5, v1, v4
	v_mul_lo_u32 v6, v1, v3
	v_mul_hi_u32 v7, v1, v3
	v_mul_hi_u32 v8, v2, v4
	v_mul_lo_u32 v4, v2, v4
	v_mul_hi_u32 v9, v2, v3
	v_mul_lo_u32 v3, v2, v3
	v_add_co_u32 v5, vcc_lo, v5, v6
	v_add_co_ci_u32_e32 v6, vcc_lo, 0, v7, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, v5, v4
	v_add_co_ci_u32_e32 v4, vcc_lo, v6, v8, vcc_lo
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, v4, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, 0, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v5, vcc_lo, v1, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, v2, v4, vcc_lo
	v_mov_b32_e32 v1, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_mul_hi_u32 v2, s0, v5
	v_mul_lo_u32 v6, s1, v5
	v_mul_lo_u32 v3, s0, v4
	v_mul_lo_u32 v8, s0, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v7, v2, v3
	v_mad_u64_u32 v[2:3], null, s12, s15, v[0:1]
	v_mul_hi_u32 v1, v5, v8
	v_mul_hi_u32 v10, v4, v8
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v0, v7, v6
	v_mul_lo_u32 v8, v4, v8
	v_ashrrev_i32_e32 v9, 31, v3
	v_mul_lo_u32 v6, v5, v0
	v_mul_hi_u32 v7, v5, v0
	v_mul_hi_u32 v11, v4, v0
	v_mul_lo_u32 v0, v4, v0
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v1, vcc_lo, v1, v6
	v_add_co_ci_u32_e32 v6, vcc_lo, 0, v7, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, v1, v8
	v_add_co_ci_u32_e32 v1, vcc_lo, v6, v10, vcc_lo
	v_add_co_ci_u32_e32 v6, vcc_lo, 0, v11, vcc_lo
	v_add_co_u32 v7, vcc_lo, v2, v9
	v_add_co_ci_u32_e32 v8, vcc_lo, v3, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v0, vcc_lo, v1, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, 0, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v10, v7, v9
	v_add_co_u32 v6, vcc_lo, v5, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v11, vcc_lo, v4, v1, vcc_lo
	v_xor_b32_e32 v8, v8, v9
	v_mul_hi_u32 v12, v10, v6
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mad_u64_u32 v[0:1], null, v10, v11, 0
	v_mad_u64_u32 v[4:5], null, v8, v6, 0
	v_mad_u64_u32 v[6:7], null, v8, v11, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v0, vcc_lo, v12, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, 0, v1, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, v0, v4
	v_add_co_ci_u32_e32 v0, vcc_lo, v1, v5, vcc_lo
	v_add_co_ci_u32_e32 v1, vcc_lo, 0, v7, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, v0, v6
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v1, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_mul_lo_u32 v6, s3, v4
	v_mad_u64_u32 v[0:1], null, s2, v4, 0
	v_mul_lo_u32 v4, s2, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_co_u32 v0, vcc_lo, v10, v0
	v_add3_u32 v1, v1, v4, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v4, v8, v1
	v_subrev_co_ci_u32_e64 v4, s0, s3, v4, vcc_lo
	v_sub_co_ci_u32_e32 v1, vcc_lo, v8, v1, vcc_lo
	v_sub_co_u32 v5, vcc_lo, v0, s2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_subrev_co_ci_u32_e64 v6, s0, 0, v4, vcc_lo
	v_cmp_le_u32_e64 s0, s2, v0
	v_subrev_co_ci_u32_e32 v4, vcc_lo, s3, v4, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s3, v1
	v_cndmask_b32_e64 v7, 0, -1, s0
	v_cmp_le_u32_e64 s0, s2, v5
	v_cndmask_b32_e64 v11, 0, -1, vcc_lo
	v_cmp_eq_u32_e32 vcc_lo, s3, v6
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v8, 0, -1, s0
	v_cmp_le_u32_e64 s0, s3, v6
	v_cndmask_b32_e64 v10, 0, -1, s0
	v_cmp_eq_u32_e64 s0, s3, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v8, v10, v8, vcc_lo
	v_sub_co_u32 v10, vcc_lo, v5, s2
	v_subrev_co_ci_u32_e32 v4, vcc_lo, 0, v4, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 0, v8
	v_cndmask_b32_e64 v7, v11, v7, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v4, v6, v4 :: v_dual_cndmask_b32 v5, v5, v10
	v_cmp_ne_u32_e32 vcc_lo, 0, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v0, v0, v5 :: v_dual_cndmask_b32 v1, v1, v4
	v_xor_b32_e32 v0, v0, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v1, v1, v9
	v_sub_co_u32 v4, vcc_lo, v0, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_sub_co_ci_u32_e32 v5, vcc_lo, v1, v9, vcc_lo
	v_add_co_u32 v0, vcc_lo, v2, s4
	v_add_co_ci_u32_e32 v1, vcc_lo, s5, v3, vcc_lo
	v_cmp_eq_u64_e32 vcc_lo, 0, v[4:5]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_gt_i64_e64 s0, s[6:7], v[0:1]
	s_and_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_13
; %bb.1:                                ; %.preheader
	v_mov_b32_e32 v2, 1
	s_mov_b32 s0, exec_lo
	v_cmpx_ne_u64_e32 1, v[0:1]
	s_cbranch_execz .LBB0_9
; %bb.2:                                ; %.lr.ph.preheader
	s_mov_b32 s2, 1
	s_mov_b32 s1, 0
.LBB0_3:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	v_and_b32_e32 v2, 1, v0
	s_mov_b32 s3, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 1, v2
	s_xor_b32 s3, exec_lo, s3
; %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:                                ; %Flow31
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_and_not1_saveexec_b32 s3, s3
; %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, s3
	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 s2, s2, 1
	v_mov_b32_e32 v2, s2
	s_or_b32 s1, vcc_lo, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s1
	s_cbranch_execnz .LBB0_3
; %bb.8:                                ; %Flow32
	s_or_b32 exec_lo, exec_lo, s1
.LBB0_9:                                ; %Flow33
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s0
	s_mov_b32 s1, exec_lo
	s_brev_b32 s0, 1
.LBB0_10:                               ; %ComputeLoop
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s2, s1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s3, v2, s2
	s_lshl_b32 s2, 1, s2
	s_and_not1_b32 s1, s1, s2
	s_delay_alu instid0(VALU_DEP_1)
	s_max_i32 s0, s0, s3
	s_cmp_lg_u32 s1, 0
	s_cbranch_scc1 .LBB0_10
; %bb.11:                               ; %ComputeEnd
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s1, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s1, exec_lo, s1
	s_cbranch_execz .LBB0_13
; %bb.12:
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0
	global_atomic_max_i32 v0, v1, s[10:11]
.LBB0_13:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _ZL7collatzlllPi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 13
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.section	.text._ZL7collatzlllPi,"axG",@progbits,_ZL7collatzlllPi,comdat
.Lfunc_end0:
	.size	_ZL7collatzlllPi, .Lfunc_end0-_ZL7collatzlllPi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1160
; NumSgprs: 18
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 13
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.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
      - .offset:         16
        .size:           8
        .value_kind:     by_value
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _ZL7collatzlllPi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _ZL7collatzlllPi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
