/* CUDA exercise to convert a simple serial code for a brute force
   largest prime number search into CUDA (32-bit, int version). This
   initial code is serial, but it is written as CUDA code for your
   convenience, so should be compiled with nvcc (see below). Your task
   is to convert the serial computation to a kernel computation. In
   the simplest case, use atomicMax to find the globally largest prime
   number.

   All prime numbers can be expressed as 6*k-1 or 6*k+1, k being an
   integer. We provide the range of k to probe as macro parameters
   KMIN and KMAX (see below).

   You should get a speedup ~20x with atomicMax.

Hints:

* It's very convenient to use a two-dimensional grid of blocks,
  defined as "dim3 Nblocks (NBLOCKS, 2, 1);". The second grid
  dimension is used to derive the two values of j=(-1; 1) inside the
  kernel: "int j = 2*blockIdx.y - 1;". This way, there will be only
  one loop inside the kernel - for y.

* When you get a failure (not a prime) inside the y loop, you can exit
  the thread with "return" (no need to use "break").




To compile:

nvcc -arch=sm_20 -O2 primes.cu -o primes

*/

#include <sys/time.h>
#include <ctype.h>
#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#include <unistd.h>

// Range of k-numbers for primes search:
#define KMIN 100000000
// Should be smaller than 357,913,941 (because we are using signed int)
#define KMAX 100100000

// Number of threads in one block (possible range is 32...1024):
#define BLOCK_SIZE 256

// Number of blocks to run:
#define NBLOCKS (KMAX-KMIN+BLOCK_SIZE)/BLOCK_SIZE

/* Subtract the `struct timeval' values X and Y,
   storing the result in RESULT.
   Return 1 if the difference is negative, otherwise 0.  */

// It messes up with y!

int
timeval_subtract (double *result, struct timeval *x, struct timeval *y)
{
  struct timeval result0;

  /* Perform the carry for the later subtraction by updating y. */
  if (x->tv_usec < y->tv_usec) {
    int nsec = (y->tv_usec - x->tv_usec) / 1000000 + 1;
    y->tv_usec -= 1000000 * nsec;
    y->tv_sec += nsec;
  }
  if (x->tv_usec - y->tv_usec > 1000000) {
    int nsec = (y->tv_usec - x->tv_usec) / 1000000;
    y->tv_usec += 1000000 * nsec;
    y->tv_sec -= nsec;
  }

  /* Compute the time remaining to wait.
     tv_usec is certainly positive. */
  result0.tv_sec = x->tv_sec - y->tv_sec;
  result0.tv_usec = x->tv_usec - y->tv_usec;
  *result = ((double)result0.tv_usec)/1e6 + (double)result0.tv_sec;

  /* Return 1 if result is negative. */
  return x->tv_sec < y->tv_sec;
}


//++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++


// Kernel(s) should go here:




int main (int argc,char **argv)
{
  struct timeval  tdr0, tdr1, tdr;
  double restime;
  int devid, devcount, error, success;
  int xmax, ymax, x, y;

  if (BLOCK_SIZE>1024)
    {
      printf ("Bad BLOCK_SIZE: %d\n", BLOCK_SIZE);
      exit (1);
    }

  /* find number of device in current "context" */
  cudaGetDevice(&devid);
  /* find how many devices are available */
  if (cudaGetDeviceCount(&devcount) || devcount==0)
    {
      printf ("No CUDA devices!\n");
      exit (1);
    }
  else
    {
      cudaDeviceProp deviceProp; 
      cudaGetDeviceProperties (&deviceProp, devid);
      printf ("Device count, devid: %d %d\n", devcount, devid);
      printf ("Device: %s\n", deviceProp.name);
      printf("[deviceProp.major.deviceProp.minor] = [%d.%d]\n\n", deviceProp.major, deviceProp.minor);
    }

  //--------------------------------------------------------------------------------
  if (error = cudaDeviceSynchronize())
    {
      printf ("Error %d\n", error);
      exit (error);
    }
  gettimeofday (&tdr0, NULL);


  // This serial computation will have to be replaced by calls to kernel(s):
  xmax = 0;
  for (int k=KMIN; k<=KMAX; k++)
    {
      // testing "-1" and "+1" cases:
      for (int j=-1; j<2; j=j+2)
	{
	  // Prime candidate:
	  x = 6*k + j;
	  // We should be dividing by numbers up to sqrt(x):
	  ymax = (int)ceil(sqrt((double)x));

	  // Primality test:
	  for (y=3; y<=ymax; y=y+2)
	    {
	      // Tpo be a success, the modulus should not be equal to zero:
	      success = x % y;
	      if (!success)
		break;
	    }

	  if (success && x > xmax)
	    {
	      xmax = x;
	    }
	}
    }

  if (error = cudaDeviceSynchronize())
    {
      printf ("Error %d\n", error);
      exit (error);
    }
  gettimeofday (&tdr1, NULL);
  tdr = tdr0;
  timeval_subtract (&restime, &tdr1, &tdr);
  printf ("%d\n", xmax);
  printf ("Time: %e\n", restime);
  //--------------------------------------------------------------------------------



  return 0;

}

	.file	"tmpxft_00353424_00000000-6_primes.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2044:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2044:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z16timeval_subtractPdP7timevalS1_
	.type	_Z16timeval_subtractPdP7timevalS1_, @function
_Z16timeval_subtractPdP7timevalS1_:
.LFB2040:
	.cfi_startproc
	endbr64
	movq	%rdx, %rcx
	movq	8(%rsi), %rdx
	movq	8(%rcx), %r8
	cmpq	%r8, %rdx
	jge	.L3
	movq	%r8, %rax
	movl	$1000000, %r9d
	subq	%rdx, %rax
	cqto
	idivq	%r9
	incl	%eax
	imull	$1000000, %eax, %edx
	cltq
	addq	%rax, (%rcx)
	movslq	%edx, %rdx
	subq	%rdx, %r8
	movq	%r8, 8(%rcx)
.L3:
	movq	8(%rsi), %rdx
	movq	8(%rcx), %r8
	movq	%rdx, %rax
	subq	%r8, %rax
	cmpq	$1000000, %rax
	jle	.L4
	movq	%r8, %rax
	movl	$1000000, %r9d
	subq	%rdx, %rax
	cqto
	idivq	%r9
	imull	$1000000, %eax, %edx
	cltq
	subq	%rax, (%rcx)
	movslq	%edx, %rdx
	addq	%r8, %rdx
	movq	%rdx, 8(%rcx)
.L4:
	movq	(%rsi), %r8
	movq	(%rcx), %rdx
	movq	8(%rsi), %rax
	subq	8(%rcx), %rax
	cvtsi2sdq	%rax, %xmm0
	movq	%r8, %rax
	divsd	.LC0(%rip), %xmm0
	subq	%rdx, %rax
	cvtsi2sdq	%rax, %xmm1
	xorl	%eax, %eax
	cmpq	%rdx, %r8
	setl	%al
	addsd	%xmm1, %xmm0
	movsd	%xmm0, (%rdi)
	ret
	.cfi_endproc
.LFE2040:
	.size	_Z16timeval_subtractPdP7timevalS1_, .-_Z16timeval_subtractPdP7timevalS1_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"Device count, devid: %d %d\n"
.LC2:
	.string	"Device: %s\n"
.LC3:
	.string	"[deviceProp.major.deviceProp.minor] = [%d.%d]\n\n"
.LC4:
	.string	"No CUDA devices!\n"
.LC5:
	.string	"Error %d\n"
.LC6:
	.string	"%d\n"
.LC7:
	.string	"Time: %e\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2041:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$1104, %rsp
	.cfi_def_cfa_offset 1152
	movq	%fs:40, %rax
	movq	%rax, 1096(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %rdi
	call	cudaGetDevice@PLT
	leaq	4(%rsp), %rdi
	call	cudaGetDeviceCount@PLT
	testl	%eax, %eax
	jne	.L7
	cmpl	$0, 4(%rsp)
	je	.L7
	leaq	64(%rsp), %rbx
	movl	(%rsp), %esi
	movq	%rbx, %rdi
	call	cudaGetDeviceProperties_v2@PLT
	movl	(%rsp), %ecx
	movl	4(%rsp), %edx
	xorl	%eax, %eax
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	%rbx, %rdx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC2(%rip), %rsi
	call	__printf_chk@PLT
	movl	428(%rsp), %ecx
	movl	$2, %edi
	xorl	%eax, %eax
	movl	424(%rsp), %edx
	leaq	.LC3(%rip), %rsi
	call	__printf_chk@PLT
	call	cudaDeviceSynchronize@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L27
	movl	%eax, %edx
	movl	$2, %edi
	leaq	.LC5(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	%ebx, %edi
	jmp	.L29
.L7:
	movl	$2, %edi
	leaq	.LC4(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1, %edi
	jmp	.L29
.L31:
	addl	$6, %r12d
	cmpl	$600600005, %r12d
	jne	.L10
	call	cudaDeviceSynchronize@PLT
	movl	%eax, %ebp
	testl	%eax, %eax
	je	.L17
	movl	%eax, %edx
	movl	$2, %edi
	leaq	.LC5(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	%ebp, %edi
.L29:
	call	exit@PLT
.L27:
	leaq	16(%rsp), %rdi
	xorl	%esi, %esi
	movl	$599999999, %r12d
	call	gettimeofday@PLT
.L10:
	movl	%r12d, %ebp
	orl	$-1, %r14d
.L15:
	cvtsi2sdl	%ebp, %xmm0
	call	sqrt@PLT
	call	ceil@PLT
	movl	$3, %ecx
	cvttsd2sil	%xmm0, %esi
.L11:
	cmpl	%esi, %ecx
	jg	.L30
	movl	%ebp, %eax
	cltd
	idivl	%ecx
	movl	%edx, %r13d
	testl	%edx, %edx
	je	.L12
	addl	$2, %ecx
	jmp	.L11
.L30:
	testl	%r13d, %r13d
	je	.L12
	cmpl	%ebx, %ebp
	cmovge	%ebp, %ebx
.L12:
	addl	$2, %ebp
	decl	%r14d
	je	.L31
	movl	$1, %r14d
	jmp	.L15
.L17:
	leaq	32(%rsp), %rbp
	xorl	%esi, %esi
	movq	%rbp, %rdi
	call	gettimeofday@PLT
	movups	16(%rsp), %xmm1
	leaq	48(%rsp), %rdx
	movq	%rbp, %rsi
	leaq	8(%rsp), %rdi
	movups	%xmm1, 48(%rsp)
	call	_Z16timeval_subtractPdP7timevalS1_
	leaq	.LC6(%rip), %rsi
	xorl	%eax, %eax
	movl	%ebx, %edx
	movl	$2, %edi
	call	__printf_chk@PLT
	movsd	8(%rsp), %xmm0
	leaq	.LC7(%rip), %rsi
	movb	$1, %al
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	1096(%rsp), %rax
	subq	%fs:40, %rax
	je	.L18
	call	__stack_chk_fail@PLT
.L18:
	addq	$1104, %rsp
	.cfi_def_cfa_offset 48
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2041:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2067:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2067:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC0:
	.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:

/* CUDA exercise to convert a simple serial code for a brute force
   largest prime number search into CUDA (32-bit, int version). This
   initial code is serial, but it is written as CUDA code for your
   convenience, so should be compiled with nvcc (see below). Your task
   is to convert the serial computation to a kernel computation. In
   the simplest case, use atomicMax to find the globally largest prime
   number.

   All prime numbers can be expressed as 6*k-1 or 6*k+1, k being an
   integer. We provide the range of k to probe as macro parameters
   KMIN and KMAX (see below).

   You should get a speedup ~20x with atomicMax.

Hints:

* It's very convenient to use a two-dimensional grid of blocks,
  defined as "dim3 Nblocks (NBLOCKS, 2, 1);". The second grid
  dimension is used to derive the two values of j=(-1; 1) inside the
  kernel: "int j = 2*blockIdx.y - 1;". This way, there will be only
  one loop inside the kernel - for y.

* When you get a failure (not a prime) inside the y loop, you can exit
  the thread with "return" (no need to use "break").




To compile:

nvcc -arch=sm_20 -O2 primes.cu -o primes

*/


#include <hip/hip_runtime.h>
#include <sys/time.h>
#include <ctype.h>
#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#include <unistd.h>

// Range of k-numbers for primes search:
#define KMIN 100000000
// Should be smaller than 357,913,941 (because we are using signed int)
#define KMAX 100100000

// Number of threads in one block (possible range is 32...1024):
#define BLOCK_SIZE 256

// Number of blocks to run:
#define NBLOCKS (KMAX-KMIN+BLOCK_SIZE)/BLOCK_SIZE

/* Subtract the `struct timeval' values X and Y,
   storing the result in RESULT.
   Return 1 if the difference is negative, otherwise 0.  */

// It messes up with y!

int
timeval_subtract (double *result, struct timeval *x, struct timeval *y)
{
  struct timeval result0;

  /* Perform the carry for the later subtraction by updating y. */
  if (x->tv_usec < y->tv_usec) {
    int nsec = (y->tv_usec - x->tv_usec) / 1000000 + 1;
    y->tv_usec -= 1000000 * nsec;
    y->tv_sec += nsec;
  }
  if (x->tv_usec - y->tv_usec > 1000000) {
    int nsec = (y->tv_usec - x->tv_usec) / 1000000;
    y->tv_usec += 1000000 * nsec;
    y->tv_sec -= nsec;
  }

  /* Compute the time remaining to wait.
     tv_usec is certainly positive. */
  result0.tv_sec = x->tv_sec - y->tv_sec;
  result0.tv_usec = x->tv_usec - y->tv_usec;
  *result = ((double)result0.tv_usec)/1e6 + (double)result0.tv_sec;

  /* Return 1 if result is negative. */
  return x->tv_sec < y->tv_sec;
}


//++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++


// Kernel(s) should go here:




int main (int argc,char **argv)
{
  struct timeval  tdr0, tdr1, tdr;
  double restime;
  int devid, devcount, error, success;
  int xmax, ymax, x, y;

  if (BLOCK_SIZE>1024)
    {
      printf ("Bad BLOCK_SIZE: %d\n", BLOCK_SIZE);
      exit (1);
    }

  /* find number of device in current "context" */
  hipGetDevice(&devid);
  /* find how many devices are available */
  if (hipGetDeviceCount(&devcount) || devcount==0)
    {
      printf ("No CUDA devices!\n");
      exit (1);
    }
  else
    {
      hipDeviceProp_t deviceProp; 
      hipGetDeviceProperties (&deviceProp, devid);
      printf ("Device count, devid: %d %d\n", devcount, devid);
      printf ("Device: %s\n", deviceProp.name);
      printf("[deviceProp.major.deviceProp.minor] = [%d.%d]\n\n", deviceProp.major, deviceProp.minor);
    }

  //--------------------------------------------------------------------------------
  if (error = hipDeviceSynchronize())
    {
      printf ("Error %d\n", error);
      exit (error);
    }
  gettimeofday (&tdr0, NULL);


  // This serial computation will have to be replaced by calls to kernel(s):
  xmax = 0;
  for (int k=KMIN; k<=KMAX; k++)
    {
      // testing "-1" and "+1" cases:
      for (int j=-1; j<2; j=j+2)
	{
	  // Prime candidate:
	  x = 6*k + j;
	  // We should be dividing by numbers up to sqrt(x):
	  ymax = (int)ceil(sqrt((double)x));

	  // Primality test:
	  for (y=3; y<=ymax; y=y+2)
	    {
	      // Tpo be a success, the modulus should not be equal to zero:
	      success = x % y;
	      if (!success)
		break;
	    }

	  if (success && x > xmax)
	    {
	      xmax = x;
	    }
	}
    }

  if (error = hipDeviceSynchronize())
    {
      printf ("Error %d\n", error);
      exit (error);
    }
  gettimeofday (&tdr1, NULL);
  tdr = tdr0;
  timeval_subtract (&restime, &tdr1, &tdr);
  printf ("%d\n", xmax);
  printf ("Time: %e\n", restime);
  //--------------------------------------------------------------------------------



  return 0;

}

	.text
	.file	"primes.hip"
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z16timeval_subtractPdP7timevalS1_
.LCPI0_0:
	.quad	0x412e848000000000              # double 1.0E+6
	.text
	.globl	_Z16timeval_subtractPdP7timevalS1_
	.type	_Z16timeval_subtractPdP7timevalS1_,@function
_Z16timeval_subtractPdP7timevalS1_:     # @_Z16timeval_subtractPdP7timevalS1_
	.cfi_startproc
# %bb.0:
	movq	%rdx, %rcx
	movq	8(%rsi), %rdx
	movq	(%rcx), %r8
	movq	8(%rcx), %r9
	movq	%r9, %rax
	movabsq	$4835703278458516699, %r10      # imm = 0x431BDE82D7B634DB
	subq	%rdx, %rax
	jle	.LBB0_2
# %bb.1:
	imulq	%r10
	movq	%rdx, %rax
	shrq	$18, %rax
	shrq	$63, %rdx
	addl	%edx, %eax
	incl	%eax
	imull	$1000000, %eax, %edx            # imm = 0xF4240
	movslq	%edx, %rdx
	subq	%rdx, %r9
	movq	%r9, 8(%rcx)
	cltq
	addq	%rax, %r8
	movq	%r8, (%rcx)
	movq	8(%rsi), %rdx
.LBB0_2:
	movq	%rdx, %rax
	subq	%r9, %rax
	cmpq	$1000001, %rax                  # imm = 0xF4241
	jl	.LBB0_4
# %bb.3:
	movq	%r9, %rax
	subq	%rdx, %rax
	imulq	%r10
	movq	%rdx, %rax
	shrq	$63, %rax
	sarq	$18, %rdx
	addq	%rax, %rdx
	movabsq	$4294967296000000, %r10         # imm = 0xF424000000000
	imulq	%rdx, %r10
	sarq	$32, %r10
	addq	%r9, %r10
	movq	%r10, 8(%rcx)
	movslq	%edx, %rax
	subq	%rax, %r8
	movq	%r8, (%rcx)
	movq	8(%rsi), %rax
	subq	%r10, %rax
.LBB0_4:
	movq	(%rsi), %rcx
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI0_0(%rip), %xmm0
	xorl	%eax, %eax
	subq	%r8, %rcx
	cvtsi2sd	%rcx, %xmm1
	addsd	%xmm0, %xmm1
	movsd	%xmm1, (%rdi)
	setl	%al
	retq
.Lfunc_end0:
	.size	_Z16timeval_subtractPdP7timevalS1_, .Lfunc_end0-_Z16timeval_subtractPdP7timevalS1_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI1_0:
	.quad	0x0000000000000000              # double 0
	.text
	.globl	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	$1528, %rsp                     # imm = 0x5F8
	.cfi_def_cfa_offset 1584
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	leaq	4(%rsp), %rdi
	callq	hipGetDevice
	movq	%rsp, %rdi
	callq	hipGetDeviceCount
	testl	%eax, %eax
	jne	.LBB1_18
# %bb.1:
	cmpl	$0, (%rsp)
	je	.LBB1_18
# %bb.2:
	movl	4(%rsp), %esi
	leaq	56(%rsp), %r14
	movq	%r14, %rdi
	callq	hipGetDevicePropertiesR0600
	movl	(%rsp), %esi
	movl	4(%rsp), %edx
	xorl	%ebx, %ebx
	movl	$.L.str.1, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.2, %edi
	movq	%r14, %rsi
	xorl	%eax, %eax
	callq	printf
	movl	360(%r14), %esi
	movl	364(%r14), %edx
	movl	$.L.str.3, %edi
	xorl	%eax, %eax
	callq	printf
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB1_19
# %bb.3:
	leaq	56(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movl	$100000000, %r14d               # imm = 0x5F5E100
	xorl	%r12d, %r12d
.LBB1_4:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_5 Depth 2
                                        #       Child Loop BB1_13 Depth 3
	leal	(%r14,%r14), %eax
	leal	(%rax,%rax,2), %r15d
	movl	$-1, %r13d
.LBB1_5:                                #   Parent Loop BB1_4 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB1_13 Depth 3
	leal	(%r15,%r13), %ebp
	xorps	%xmm0, %xmm0
	cvtsi2sd	%ebp, %xmm0
	ucomisd	.LCPI1_0(%rip), %xmm0
	jb	.LBB1_10
# %bb.6:                                #   in Loop: Header=BB1_5 Depth=2
	sqrtsd	%xmm0, %xmm0
	jmp	.LBB1_11
.LBB1_10:                               # %call.sqrt
                                        #   in Loop: Header=BB1_5 Depth=2
	callq	sqrt
.LBB1_11:                               # %.split
                                        #   in Loop: Header=BB1_5 Depth=2
	callq	ceil@PLT
	cvttsd2si	%xmm0, %ecx
	cmpl	$3, %ecx
	jge	.LBB1_12
# %bb.16:                               # %._crit_edge
                                        #   in Loop: Header=BB1_5 Depth=2
	cmpl	%ebx, %ebp
	cmovlel	%ebx, %ebp
	testl	%r12d, %r12d
	cmovnel	%ebp, %ebx
	jmp	.LBB1_17
.LBB1_12:                               # %.lr.ph.preheader
                                        #   in Loop: Header=BB1_5 Depth=2
	movl	$3, %esi
.LBB1_13:                               # %.lr.ph
                                        #   Parent Loop BB1_4 Depth=1
                                        #     Parent Loop BB1_5 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	xorl	%r12d, %r12d
	movl	%ebp, %eax
	xorl	%edx, %edx
	divl	%esi
	testl	%edx, %edx
	je	.LBB1_17
# %bb.14:                               #   in Loop: Header=BB1_13 Depth=3
	addl	$2, %esi
	cmpl	%ecx, %esi
	jle	.LBB1_13
# %bb.15:                               # %._crit_edge.thread55
                                        #   in Loop: Header=BB1_5 Depth=2
	cmpl	%ebx, %ebp
	cmovgl	%ebp, %ebx
	movl	%edx, %r12d
.LBB1_17:                               # %._crit_edge.thread
                                        #   in Loop: Header=BB1_5 Depth=2
	leal	2(%r13), %eax
	testl	%r13d, %r13d
	movl	%eax, %r13d
	js	.LBB1_5
# %bb.7:                                #   in Loop: Header=BB1_4 Depth=1
	incl	%r14d
	cmpl	$100100001, %r14d               # imm = 0x5F767A1
	jne	.LBB1_4
# %bb.8:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB1_19
# %bb.9:
	leaq	40(%rsp), %r14
	movq	%r14, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movups	56(%rsp), %xmm0
	leaq	16(%rsp), %rdx
	movaps	%xmm0, (%rdx)
	leaq	8(%rsp), %r15
	movq	%r15, %rdi
	movq	%r14, %rsi
	callq	_Z16timeval_subtractPdP7timevalS1_
	movl	$.L.str.5, %edi
	movl	%ebx, %esi
	xorl	%eax, %eax
	callq	printf
	movsd	(%r15), %xmm0                   # xmm0 = mem[0],zero
	movl	$.L.str.6, %edi
	movb	$1, %al
	callq	printf
	xorl	%eax, %eax
	addq	$1528, %rsp                     # imm = 0x5F8
	.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
.LBB1_19:
	.cfi_def_cfa_offset 1584
	movl	%eax, %ebp
	movl	$.L.str.4, %edi
	movl	%eax, %esi
	xorl	%eax, %eax
	callq	printf
	movl	%ebp, %edi
	callq	exit
.LBB1_18:
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	$1, %edi
	callq	exit
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"Device count, devid: %d %d\n"
	.size	.L.str.1, 28

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Device: %s\n"
	.size	.L.str.2, 12

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"[deviceProp.major.deviceProp.minor] = [%d.%d]\n\n"
	.size	.L.str.3, 48

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Error %d\n"
	.size	.L.str.4, 10

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"%d\n"
	.size	.L.str.5, 4

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Time: %e\n"
	.size	.L.str.6, 10

	.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	"No CUDA devices!"
	.size	.Lstr, 17

	.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_

	.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_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
