/*
   A multithreaded C-program for MT19937.
   Original single threaded C reference coded by Takuji Nishimurar
   and Makoto Matsumoto, with initialization improved 2002/1/26.
   Multithreaded C implementation coded by Eric Mills.

   Before using, initialize the state by using mt19937gi(seed)
   or mt19937gai(init_key, key_length) for the global memory versions or
   mt19937si(seed) or mt19937sai(init_key, key_length) for all shared
   memory versions.

   Copyright (C) 1997 - 2002, Makoto Matsumoto and Takuji Nishimura,
   All rights reserved.
   Multithreaded implementation Copyright (C) 2007, Eric Mills.
   All rights reserved.

   Redistribution and use in source and binary forms, with or without
   modification, are permitted provided that the following conditions
   are met:

     1. Redistributions of source code must retain the above copyright
        notice, this list of conditions and the following disclaimer.

     2. Redistributions in binary form must reproduce the above copyright
        notice, this list of conditions and the following disclaimer in the
        documentation and/or other materials provided with the distribution.

     3. The names of its contributors may not be used to endorse or promote
        products derived from this software without specific prior written
        permission.

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


   Any feedback is very welcome.
   http://www.math.sci.hiroshima-u.ac.jp/~m-mat/MT/emt.html
   email: m-mat @ math.sci.hiroshima-u.ac.jp (remove space)
*/

#define	NVG80				/* For Nvidia G80 achitecture where mod is VERY slow */

#ifdef NVG80
#define	mod(x, y)	((x) < (y) ? (x) : (x) - (y))	/* Short mod - known input range */
#else
#define	mod(x, y)	((x) % (y))
#endif

#ifdef _WIN32
typedef unsigned int uint;
#endif

#define N		624
#define M		397
#define INIT_MULT	1812433253	/* See Knuth TAOCP Vol2. 3rd Ed. P.106 for multiplier. */
#define	ARRAY_SEED	19650218	/* Seed for initial setup before incorp array seed */
#define MATRIX_A	0x9908b0df	/* Constant vector a */
#define UPPER_MASK	0x80000000	/* Most significant w-r bits */
#define LOWER_MASK	0x7fffffff	/* Least significant r bits */
#define	TEMPER1		0x9d2c5680
#define	TEMPER2		0xefc60000

/* First a global memory implementation that uses 2 global reads and 1 global
 * write per result and keeps only 2 words of state in permanent shared memory. */

#define	MAX_THREADS	227	/* Set to minimise shared memory allocation (max blockDim.x) */
#define	MAX_BLOCKS	256	/* Set to minimise global memory allocation (max gridDim.x) */

__shared__ int	mtNext;		/* Start of next block of seeds */
__shared__ uint	mtNexti;	/* Indirect on above to save global read cycle */
__device__ uint g_seeds[MAX_BLOCKS][N];
__constant__ uint mag01[2] = {0, MATRIX_A};	/* 2 way bus conflict for each read */

/* Init by single seed - single threaded as only used once */
__device__ static void
mt19937gi(uint seed)
{
    int			i;

    mtNext = 0;
    if (threadIdx.x == 0)
    {
	g_seeds[blockIdx.x][0] = mtNexti = seed;
	for (i = 1; i < N; i++)
	{
	    seed = (INIT_MULT * (seed ^ (seed >> 30)) + i); 
	    g_seeds[blockIdx.x][i] = seed;
	}
    }
    return;
}

/* Init by array - single threaded as only used once, opt to reduce global refs */
__device__ static void
mt19937gai(uint* seeds, uint length)
{
    mt19937gi(ARRAY_SEED);
    if (threadIdx.x == 0)
    {
	int	i = 1;
	int	j = 0;
	int	k;
	uint	mti;				/* g_seeds[i] */
	uint	mtj;				/* g_seeds[i - 1] */

	mti = g_seeds[blockIdx.x][0];
	for (k = N > length ? N : length; k != 0; k--)
	{
	    mtj = mti;
	    mti = g_seeds[blockIdx.x][i];
	    mti = (mti ^ ((mtj ^ (mtj >> 30)) * 1664525)) + seeds[j] + j;
	    g_seeds[blockIdx.x][i] = mti;
	    if (++i >= N)
	    {
		g_seeds[blockIdx.x][0] = mti;
		i = 1;
	    }
	    if (++j >= length)
	    {
		j = 0;
	    }
	}
	for (k = N - 1; k != 0; k--)
	{
	    mtj = mti;
	    mti = g_seeds[blockIdx.x][i];
	    mti = (mti ^ ((mtj ^ (mtj >> 30)) * 1566083941)) - i;
	    g_seeds[blockIdx.x][i] = mti;
	    if (++i >= N)
	    {
		g_seeds[blockIdx.x][0] = mti;
		i = 1;
	    }
	}
	g_seeds[blockIdx.x][0] = mtNexti = 0x80000000;	/* MSB is 1; assuring non-zero initial array */ 
    }
    return;
}

/* Return next MT random by increasing thread ID, Good for 1 - 227 threads.
 * Note you should wind back MAX_THREADS to your max requirement
 * to keep auto allocation of shared mem to a minimum.
 * Best as a general purpose library routine. */
__device__ static uint
mt19937g(void)
{
    int			kk;
    uint		y;
    const int		tid = threadIdx.x;
    __shared__ uint	seed[MAX_THREADS + 1];

    kk = mod(mtNext + tid, N);
    __syncthreads();				/* Finish with mtNext & g_seeds ready from last call & init */

    seed[tid + 1] = g_seeds[blockIdx.x][mod(kk + 1, N)];	/* Sequential but not aligned */
    if (tid == blockDim.x - 1)
    {
	mtNext = kk + 1;
	seed[0] = mtNexti;
	mtNexti = seed[blockDim.x];
    }
    __syncthreads();				/* seed[] ready */

    y = (seed[tid] & UPPER_MASK) | (seed[tid + 1] & LOWER_MASK);
    y = g_seeds[blockIdx.x][kk < N - M ? kk + M : kk + (M - N)] ^ (y >> 1) ^ mag01[y & 1];
    g_seeds[blockIdx.x][kk] = y;		/* Does not overlap above reads */
    y ^= (y >> 11);				/* Tempering */
    y ^= (y <<  7) & TEMPER1;
    y ^= (y << 15) & TEMPER2;
    y ^= (y >> 18);
    return y;
}

/* Generalised global memory version for any number of threads.
 * Note only runs up to 227 at a time, rest loop and block till all done.
 * Runs fractional warps at each end so not perfect utilisation.
 * Uses 228 words of auto allocated shared mem. */
__device__ static uint
mt19937gl(void)
{
    int			jj;
    int			kk;
    uint		y;
    int			tid;			/* Offset thread ID */
    __shared__ uint	seed[N - M + 1];

    kk = mod(mtNext + threadIdx.x, N);		/* G80 limited to 512 threads */
    __syncthreads();				/* Finish with mtNext & g_seeds set from init */

    if (threadIdx.x == blockDim.x - 1)
    {
	mtNext = kk + 1;			/* Modded next call */
    }
    jj = 0;
    do
    {
	__syncthreads();			/* g_seeds set from last loop */

	tid = threadIdx.x - jj;
	if (0 <= tid && tid < N - M)
	{
	    seed[tid + 1] = g_seeds[blockIdx.x][mod(kk + 1, N)];	/* Sequential but not aligned */
	    y = min(N - M, blockDim.x - jj);
	    if (tid == y - 1)			/* Last thread this loop */
	    {
		seed[0] = mtNexti;
		mtNexti = seed[y];
	    }
	}
	__syncthreads();			/* seed[] ready */

	if (0 <= tid && tid < N - M)
	{
	    y = (seed[tid] & UPPER_MASK) | (seed[tid + 1] & LOWER_MASK);
	    y = g_seeds[blockIdx.x][kk < N - M ? kk + M : kk + (M - N)] ^ (y >> 1) ^ mag01[y & 1];
	    g_seeds[blockIdx.x][kk] = y;	/* Does not overlap reads above */
	}
    } while ((jj += N - M) < blockDim.x);

    y ^= (y >> 11);				/* Tempering */
    y ^= (y <<  7) & TEMPER1;
    y ^= (y << 15) & TEMPER2;
    y ^= (y >> 18);
    return y;
}

/*************************************************************************************
 * This is a shared memory implementation that keeps the full 626 words of state
 * in shared memory. Faster for heavy random work where you can afford shared mem. */

__shared__ int	mtNexts;			/* Start of next block of seeds */
__shared__ uint	s_seeds[N + 1];

/* Init by single seed - single threaded as only used once */
__device__ static void
mt19937si(uint seed)
{
    int		i;

    if (threadIdx.x == 0)
    {
	mtNexts = 0;
	s_seeds[0] = seed;
	for (i = 1; i < N; i++)
	{
	    seed = (INIT_MULT * (seed ^ (seed >> 30)) + i); 
	    s_seeds[i] = seed;
	}
    }
    __syncthreads();				/* Ensure mtNexts set & needed for mt19937w() */
    return;
}

/* Init by array - single threaded as only used once */
__device__ static void
mt19937sai(uint* seeds, uint length)
{
    mt19937si(ARRAY_SEED);
    if (threadIdx.x == 0)
    {
	int	i = 1;
	int	j = 0;
	int	k;

	for (k = N > length ? N : length; k != 0; k--)
	{
	    s_seeds[i] = (s_seeds[i] ^ ((s_seeds[i - 1] ^ (s_seeds[i - 1] >> 30)) * 1664525)) + seeds[j] + j;
	    if (++i >= N)
	    {
		s_seeds[0] = s_seeds[N - 1];
		i = 1;
	    }
	    if (++j >= length)
	    {
		j = 0;
	    }
	}
	for (k = N - 1; k != 0; k--)
	{
	    s_seeds[i] = (s_seeds[i] ^ ((s_seeds[i - 1] ^ (s_seeds[i - 1] >> 30)) * 1566083941)) - i;
	    if (++i >= N)
	    {
		s_seeds[0] = s_seeds[N - 1];
		i = 1;
	    }
	}
	s_seeds[0] = 0x80000000;		/* MSB is 1; assuring non-zero initial array */ 
    }
    __syncthreads();				/* Needed for mt19937w() */
    return;
}

/* Return next MT random by increasing thread ID for 1-227 threads. */
__device__ static uint
mt19937s(void)
{
    int		kk;
    uint	y;
    const int	tid = threadIdx.x;

    kk = mod(mtNexts + tid, N);
    __syncthreads();				/* Finished with mtNexts & s_seed[] ready from last run */

    if (tid == blockDim.x - 1)
    {
	mtNexts = kk + 1;			/* Will get modded on next call */
    }
    y = (s_seeds[kk] & UPPER_MASK) | (s_seeds[kk + 1] & LOWER_MASK);
    y = s_seeds[kk < N - M ? kk + M : kk + (M - N)] ^ (y >> 1) ^ mag01[y & 1];
    //y = s_seeds[kk < N - M ? kk + M : kk + (M - N)] ^ (y >> 1) ^ (y & 1 ? MATRIX_A : 0);	// Same speed
    __syncthreads();				/* All done before we update */

    s_seeds[kk] = y;
    if (kk == 0)				/* Copy up for next round */
    {
	s_seeds[N] = y;
    }
    y ^= (y >> 11);				/* Tempering */
    y ^= (y <<  7) & TEMPER1;
    y ^= (y << 15) & TEMPER2;
    y ^= (y >> 18);
    return y;
}

/* General shared memory version for any number of threads.
 * Note only up to 227 threads are run at any one time,
 * the rest loop and block till all are done. */
__device__ static uint
mt19937sl(void)
{
    int		jj;
    int		kk;
    uint	y;
    int		tid;				/* Offset thread ID */

    kk = mod(mtNexts + threadIdx.x, N);		/* G80 limited to 512 threads */
    __syncthreads();				/* Finished with mtNexts & s_seed[] ready from init */

    if (threadIdx.x == blockDim.x - 1)
    {
	mtNexts = kk + 1;			/* Will get modded on next call */
    }
    jj = 0;
    do
    {
	__syncthreads();			/* s_seeds[] ready from last loop */

	tid = threadIdx.x - jj;
	if (0 <= tid && tid < N - M)
	{
	    y = (s_seeds[kk] & UPPER_MASK) | (s_seeds[kk + 1] & LOWER_MASK);
	    y = s_seeds[kk < N - M ? kk + M : kk + (M - N)] ^ (y >> 1) ^ mag01[y & 1];
	}
	__syncthreads();			/* All done before we update */

	if (0 <= tid && tid < N - M)
	{
	    s_seeds[kk] = y;
	    if (kk == 0)
	    {
		s_seeds[N] = y;
	    }
	}
    } while ((jj += N - M) < blockDim.x);

    y ^= (y >> 11);				/* Tempering */
    y ^= (y <<  7) & TEMPER1;
    y ^= (y << 15) & TEMPER2;
    y ^= (y >> 18);
    return y;
}

/***************************************************************************************
 * This is an implementation of a full step in 1 call - all 624 results returned at once
 * in pairs - 64 bit version. It may be run with 227-312 threads and will drop numbers
 * from the sequence if < 312 (not incorrect).
 * Original idea for this version was first presented by Brian Budge. */

#define B2      224             /* Size of second block */

__device__ static uint2
mt19937w(const int tid)
{
   int   kk;
   uint  y;
   uint2 ret;

   kk = tid;

   /* First 227 */
   if (kk < N-M) {
      y = (s_seeds[kk]&UPPER_MASK)|(s_seeds[kk+1]&LOWER_MASK);
      y = s_seeds[kk+M] ^ (y >> 1) ^ mag01[y & 1];
   }
   __syncthreads();

   if (kk < N-M) {
       s_seeds[kk] = y;
       if (kk == 0)
       {
	   s_seeds[N] = y;
       }
   }
   kk += N-M;
   __syncthreads();

   /* Next 224 */
   if (kk < N-M + B2) {
      y = (s_seeds[kk]&UPPER_MASK)|(s_seeds[kk+1]&LOWER_MASK);
      y = s_seeds[kk+(M-N)] ^ (y >> 1) ^ mag01[y & 1];
   }
   __syncthreads();

   if (kk < N-M + B2) {
       s_seeds[kk] = y;
   }
   kk += B2;
   __syncthreads();

   /* Last 173 */
   if (kk < N) {
       y = (s_seeds[kk]&UPPER_MASK)|(s_seeds[kk+1]&LOWER_MASK);
       y = s_seeds[kk+(M-N)] ^ (y >> 1) ^ mag01[y & 1];
   }
   __syncthreads();

   if (kk < N) {
       s_seeds[kk] = y;
   }
   __syncthreads();

   ret.x = s_seeds[2*tid];
   ret.x ^= (ret.x >> 11);				/* Tempering */
   ret.x ^= (ret.x <<  7) & TEMPER1;
   ret.x ^= (ret.x << 15) & TEMPER2;
   ret.x ^= (ret.x >> 18);
   
   ret.y = s_seeds[2*tid+1];
   ret.y ^= (ret.y >> 11);
   ret.y ^= (ret.y <<  7) & TEMPER1;
   ret.y ^= (ret.y << 15) & TEMPER2;
   ret.y ^= (ret.y >> 18);

   return ret;
}
/*******************************************************************************
 * For reference this is the original C single threaded source: */
#if 0
static unsigned long mt[N]; /* the array for the state vector  */
static int mti=N+1; /* mti==N+1 means mt[N] is not initialized */

/* initializes mt[N] with a seed */
void init_genrand(unsigned long s)
{
    mt[0]= s & 0xffffffffUL;
    for (mti=1; mti<N; mti++) {
        mt[mti] = 
	    (1812433253UL * (mt[mti-1] ^ (mt[mti-1] >> 30)) + mti); 
        /* See Knuth TAOCP Vol2. 3rd Ed. P.106 for multiplier. */
        /* In the previous versions, MSBs of the seed affect   */
        /* only MSBs of the array mt[].                        */
        /* 2002/01/09 modified by Makoto Matsumoto             */
        mt[mti] &= 0xffffffffUL;
        /* for >32 bit machines */
    }
}

/* initialize by an array with array-length */
/* init_key is the array for initializing keys */
/* key_length is its length */
/* slight change for C++, 2004/2/26 */
void init_by_array(unsigned long init_key[], int key_length)
{
    int i, j, k;
    init_genrand(19650218UL);
    i=1; j=0;
    k = (N>key_length ? N : key_length);
    for (; k; k--) {
        mt[i] = (mt[i] ^ ((mt[i-1] ^ (mt[i-1] >> 30)) * 1664525UL))
          + init_key[j] + j; /* non linear */
        mt[i] &= 0xffffffffUL; /* for WORDSIZE > 32 machines */
        i++; j++;
        if (i>=N) { mt[0] = mt[N-1]; i=1; }
        if (j>=key_length) j=0;
    }
    for (k=N-1; k; k--) {
        mt[i] = (mt[i] ^ ((mt[i-1] ^ (mt[i-1] >> 30)) * 1566083941UL))
          - i; /* non linear */
        mt[i] &= 0xffffffffUL; /* for WORDSIZE > 32 machines */
        i++;
        if (i>=N) { mt[0] = mt[N-1]; i=1; }
    }

    mt[0] = 0x80000000UL; /* MSB is 1; assuring non-zero initial array */ 
}

/* generates a random number on [0,0xffffffff]-interval */
unsigned long genrand_int32(void)
{
    unsigned long y;
    static unsigned long mag01[2]={0x0UL, MATRIX_A};
    /* mag01[x] = x * MATRIX_A  for x=0,1 */

    if (mti >= N) { /* generate N words at one time */
        int kk;

        if (mti == N+1)   /* if init_genrand() has not been called, */
            init_genrand(5489UL); /* a default initial seed is used */

        for (kk=0;kk<N-M;kk++) {
            y = (mt[kk]&UPPER_MASK)|(mt[kk+1]&LOWER_MASK);
            mt[kk] = mt[kk+M] ^ (y >> 1) ^ mag01[y & 0x1UL];
        }
        for (;kk<N-1;kk++) {
            y = (mt[kk]&UPPER_MASK)|(mt[kk+1]&LOWER_MASK);
            mt[kk] = mt[kk+(M-N)] ^ (y >> 1) ^ mag01[y & 0x1UL];
        }
        y = (mt[N-1]&UPPER_MASK)|(mt[0]&LOWER_MASK);
        mt[N-1] = mt[M-1] ^ (y >> 1) ^ mag01[y & 0x1UL];

        mti = 0;
    }
  
    y = mt[mti++];

    /* Tempering */
    y ^= (y >> 11);
    y ^= (y << 7) & 0x9d2c5680UL;
    y ^= (y << 15) & 0xefc60000UL;
    y ^= (y >> 18);

    return y;
}
#endif


	.file	"tmpxft_002564ac_00000000-6_mt19937_ref.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2019:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2019:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"g_seeds"
.LC1:
	.string	"mag01"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2042:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	movl	$638976, %r9d
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	movq	%rdx, %rcx
	leaq	_ZL7g_seeds(%rip), %rsi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC1(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL5mag01(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2042:
	.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	_ZL5mag01
	.comm	_ZL5mag01,8,8
	.local	_ZL7g_seeds
	.comm	_ZL7g_seeds,638976,32
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
/*
   A multithreaded C-program for MT19937.
   Original single threaded C reference coded by Takuji Nishimurar
   and Makoto Matsumoto, with initialization improved 2002/1/26.
   Multithreaded C implementation coded by Eric Mills.

   Before using, initialize the state by using mt19937gi(seed)
   or mt19937gai(init_key, key_length) for the global memory versions or
   mt19937si(seed) or mt19937sai(init_key, key_length) for all shared
   memory versions.

   Copyright (C) 1997 - 2002, Makoto Matsumoto and Takuji Nishimura,
   All rights reserved.
   Multithreaded implementation Copyright (C) 2007, Eric Mills.
   All rights reserved.

   Redistribution and use in source and binary forms, with or without
   modification, are permitted provided that the following conditions
   are met:

     1. Redistributions of source code must retain the above copyright
        notice, this list of conditions and the following disclaimer.

     2. Redistributions in binary form must reproduce the above copyright
        notice, this list of conditions and the following disclaimer in the
        documentation and/or other materials provided with the distribution.

     3. The names of its contributors may not be used to endorse or promote
        products derived from this software without specific prior written
        permission.

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


   Any feedback is very welcome.
   http://www.math.sci.hiroshima-u.ac.jp/~m-mat/MT/emt.html
   email: m-mat @ math.sci.hiroshima-u.ac.jp (remove space)
*/

#define	NVG80				/* For Nvidia G80 achitecture where mod is VERY slow */

#ifdef NVG80
#define	mod(x, y)	((x) < (y) ? (x) : (x) - (y))	/* Short mod - known input range */
#else
#define	mod(x, y)	((x) % (y))
#endif

#ifdef _WIN32
typedef unsigned int uint;
#endif

#define N		624
#define M		397
#define INIT_MULT	1812433253	/* See Knuth TAOCP Vol2. 3rd Ed. P.106 for multiplier. */
#define	ARRAY_SEED	19650218	/* Seed for initial setup before incorp array seed */
#define MATRIX_A	0x9908b0df	/* Constant vector a */
#define UPPER_MASK	0x80000000	/* Most significant w-r bits */
#define LOWER_MASK	0x7fffffff	/* Least significant r bits */
#define	TEMPER1		0x9d2c5680
#define	TEMPER2		0xefc60000

/* First a global memory implementation that uses 2 global reads and 1 global
 * write per result and keeps only 2 words of state in permanent shared memory. */

#define	MAX_THREADS	227	/* Set to minimise shared memory allocation (max blockDim.x) */
#define	MAX_BLOCKS	256	/* Set to minimise global memory allocation (max gridDim.x) */

__shared__ int	mtNext;		/* Start of next block of seeds */
__shared__ uint	mtNexti;	/* Indirect on above to save global read cycle */
__device__ uint g_seeds[MAX_BLOCKS][N];
__constant__ uint mag01[2] = {0, MATRIX_A};	/* 2 way bus conflict for each read */

/* Init by single seed - single threaded as only used once */
__device__ static void
mt19937gi(uint seed)
{
    int			i;

    mtNext = 0;
    if (threadIdx.x == 0)
    {
	g_seeds[blockIdx.x][0] = mtNexti = seed;
	for (i = 1; i < N; i++)
	{
	    seed = (INIT_MULT * (seed ^ (seed >> 30)) + i); 
	    g_seeds[blockIdx.x][i] = seed;
	}
    }
    return;
}

/* Init by array - single threaded as only used once, opt to reduce global refs */
__device__ static void
mt19937gai(uint* seeds, uint length)
{
    mt19937gi(ARRAY_SEED);
    if (threadIdx.x == 0)
    {
	int	i = 1;
	int	j = 0;
	int	k;
	uint	mti;				/* g_seeds[i] */
	uint	mtj;				/* g_seeds[i - 1] */

	mti = g_seeds[blockIdx.x][0];
	for (k = N > length ? N : length; k != 0; k--)
	{
	    mtj = mti;
	    mti = g_seeds[blockIdx.x][i];
	    mti = (mti ^ ((mtj ^ (mtj >> 30)) * 1664525)) + seeds[j] + j;
	    g_seeds[blockIdx.x][i] = mti;
	    if (++i >= N)
	    {
		g_seeds[blockIdx.x][0] = mti;
		i = 1;
	    }
	    if (++j >= length)
	    {
		j = 0;
	    }
	}
	for (k = N - 1; k != 0; k--)
	{
	    mtj = mti;
	    mti = g_seeds[blockIdx.x][i];
	    mti = (mti ^ ((mtj ^ (mtj >> 30)) * 1566083941)) - i;
	    g_seeds[blockIdx.x][i] = mti;
	    if (++i >= N)
	    {
		g_seeds[blockIdx.x][0] = mti;
		i = 1;
	    }
	}
	g_seeds[blockIdx.x][0] = mtNexti = 0x80000000;	/* MSB is 1; assuring non-zero initial array */ 
    }
    return;
}

/* Return next MT random by increasing thread ID, Good for 1 - 227 threads.
 * Note you should wind back MAX_THREADS to your max requirement
 * to keep auto allocation of shared mem to a minimum.
 * Best as a general purpose library routine. */
__device__ static uint
mt19937g(void)
{
    int			kk;
    uint		y;
    const int		tid = threadIdx.x;
    __shared__ uint	seed[MAX_THREADS + 1];

    kk = mod(mtNext + tid, N);
    __syncthreads();				/* Finish with mtNext & g_seeds ready from last call & init */

    seed[tid + 1] = g_seeds[blockIdx.x][mod(kk + 1, N)];	/* Sequential but not aligned */
    if (tid == blockDim.x - 1)
    {
	mtNext = kk + 1;
	seed[0] = mtNexti;
	mtNexti = seed[blockDim.x];
    }
    __syncthreads();				/* seed[] ready */

    y = (seed[tid] & UPPER_MASK) | (seed[tid + 1] & LOWER_MASK);
    y = g_seeds[blockIdx.x][kk < N - M ? kk + M : kk + (M - N)] ^ (y >> 1) ^ mag01[y & 1];
    g_seeds[blockIdx.x][kk] = y;		/* Does not overlap above reads */
    y ^= (y >> 11);				/* Tempering */
    y ^= (y <<  7) & TEMPER1;
    y ^= (y << 15) & TEMPER2;
    y ^= (y >> 18);
    return y;
}

/* Generalised global memory version for any number of threads.
 * Note only runs up to 227 at a time, rest loop and block till all done.
 * Runs fractional warps at each end so not perfect utilisation.
 * Uses 228 words of auto allocated shared mem. */
__device__ static uint
mt19937gl(void)
{
    int			jj;
    int			kk;
    uint		y;
    int			tid;			/* Offset thread ID */
    __shared__ uint	seed[N - M + 1];

    kk = mod(mtNext + threadIdx.x, N);		/* G80 limited to 512 threads */
    __syncthreads();				/* Finish with mtNext & g_seeds set from init */

    if (threadIdx.x == blockDim.x - 1)
    {
	mtNext = kk + 1;			/* Modded next call */
    }
    jj = 0;
    do
    {
	__syncthreads();			/* g_seeds set from last loop */

	tid = threadIdx.x - jj;
	if (0 <= tid && tid < N - M)
	{
	    seed[tid + 1] = g_seeds[blockIdx.x][mod(kk + 1, N)];	/* Sequential but not aligned */
	    y = min(N - M, blockDim.x - jj);
	    if (tid == y - 1)			/* Last thread this loop */
	    {
		seed[0] = mtNexti;
		mtNexti = seed[y];
	    }
	}
	__syncthreads();			/* seed[] ready */

	if (0 <= tid && tid < N - M)
	{
	    y = (seed[tid] & UPPER_MASK) | (seed[tid + 1] & LOWER_MASK);
	    y = g_seeds[blockIdx.x][kk < N - M ? kk + M : kk + (M - N)] ^ (y >> 1) ^ mag01[y & 1];
	    g_seeds[blockIdx.x][kk] = y;	/* Does not overlap reads above */
	}
    } while ((jj += N - M) < blockDim.x);

    y ^= (y >> 11);				/* Tempering */
    y ^= (y <<  7) & TEMPER1;
    y ^= (y << 15) & TEMPER2;
    y ^= (y >> 18);
    return y;
}

/*************************************************************************************
 * This is a shared memory implementation that keeps the full 626 words of state
 * in shared memory. Faster for heavy random work where you can afford shared mem. */

__shared__ int	mtNexts;			/* Start of next block of seeds */
__shared__ uint	s_seeds[N + 1];

/* Init by single seed - single threaded as only used once */
__device__ static void
mt19937si(uint seed)
{
    int		i;

    if (threadIdx.x == 0)
    {
	mtNexts = 0;
	s_seeds[0] = seed;
	for (i = 1; i < N; i++)
	{
	    seed = (INIT_MULT * (seed ^ (seed >> 30)) + i); 
	    s_seeds[i] = seed;
	}
    }
    __syncthreads();				/* Ensure mtNexts set & needed for mt19937w() */
    return;
}

/* Init by array - single threaded as only used once */
__device__ static void
mt19937sai(uint* seeds, uint length)
{
    mt19937si(ARRAY_SEED);
    if (threadIdx.x == 0)
    {
	int	i = 1;
	int	j = 0;
	int	k;

	for (k = N > length ? N : length; k != 0; k--)
	{
	    s_seeds[i] = (s_seeds[i] ^ ((s_seeds[i - 1] ^ (s_seeds[i - 1] >> 30)) * 1664525)) + seeds[j] + j;
	    if (++i >= N)
	    {
		s_seeds[0] = s_seeds[N - 1];
		i = 1;
	    }
	    if (++j >= length)
	    {
		j = 0;
	    }
	}
	for (k = N - 1; k != 0; k--)
	{
	    s_seeds[i] = (s_seeds[i] ^ ((s_seeds[i - 1] ^ (s_seeds[i - 1] >> 30)) * 1566083941)) - i;
	    if (++i >= N)
	    {
		s_seeds[0] = s_seeds[N - 1];
		i = 1;
	    }
	}
	s_seeds[0] = 0x80000000;		/* MSB is 1; assuring non-zero initial array */ 
    }
    __syncthreads();				/* Needed for mt19937w() */
    return;
}

/* Return next MT random by increasing thread ID for 1-227 threads. */
__device__ static uint
mt19937s(void)
{
    int		kk;
    uint	y;
    const int	tid = threadIdx.x;

    kk = mod(mtNexts + tid, N);
    __syncthreads();				/* Finished with mtNexts & s_seed[] ready from last run */

    if (tid == blockDim.x - 1)
    {
	mtNexts = kk + 1;			/* Will get modded on next call */
    }
    y = (s_seeds[kk] & UPPER_MASK) | (s_seeds[kk + 1] & LOWER_MASK);
    y = s_seeds[kk < N - M ? kk + M : kk + (M - N)] ^ (y >> 1) ^ mag01[y & 1];
    //y = s_seeds[kk < N - M ? kk + M : kk + (M - N)] ^ (y >> 1) ^ (y & 1 ? MATRIX_A : 0);	// Same speed
    __syncthreads();				/* All done before we update */

    s_seeds[kk] = y;
    if (kk == 0)				/* Copy up for next round */
    {
	s_seeds[N] = y;
    }
    y ^= (y >> 11);				/* Tempering */
    y ^= (y <<  7) & TEMPER1;
    y ^= (y << 15) & TEMPER2;
    y ^= (y >> 18);
    return y;
}

/* General shared memory version for any number of threads.
 * Note only up to 227 threads are run at any one time,
 * the rest loop and block till all are done. */
__device__ static uint
mt19937sl(void)
{
    int		jj;
    int		kk;
    uint	y;
    int		tid;				/* Offset thread ID */

    kk = mod(mtNexts + threadIdx.x, N);		/* G80 limited to 512 threads */
    __syncthreads();				/* Finished with mtNexts & s_seed[] ready from init */

    if (threadIdx.x == blockDim.x - 1)
    {
	mtNexts = kk + 1;			/* Will get modded on next call */
    }
    jj = 0;
    do
    {
	__syncthreads();			/* s_seeds[] ready from last loop */

	tid = threadIdx.x - jj;
	if (0 <= tid && tid < N - M)
	{
	    y = (s_seeds[kk] & UPPER_MASK) | (s_seeds[kk + 1] & LOWER_MASK);
	    y = s_seeds[kk < N - M ? kk + M : kk + (M - N)] ^ (y >> 1) ^ mag01[y & 1];
	}
	__syncthreads();			/* All done before we update */

	if (0 <= tid && tid < N - M)
	{
	    s_seeds[kk] = y;
	    if (kk == 0)
	    {
		s_seeds[N] = y;
	    }
	}
    } while ((jj += N - M) < blockDim.x);

    y ^= (y >> 11);				/* Tempering */
    y ^= (y <<  7) & TEMPER1;
    y ^= (y << 15) & TEMPER2;
    y ^= (y >> 18);
    return y;
}

/***************************************************************************************
 * This is an implementation of a full step in 1 call - all 624 results returned at once
 * in pairs - 64 bit version. It may be run with 227-312 threads and will drop numbers
 * from the sequence if < 312 (not incorrect).
 * Original idea for this version was first presented by Brian Budge. */

#define B2      224             /* Size of second block */

__device__ static uint2
mt19937w(const int tid)
{
   int   kk;
   uint  y;
   uint2 ret;

   kk = tid;

   /* First 227 */
   if (kk < N-M) {
      y = (s_seeds[kk]&UPPER_MASK)|(s_seeds[kk+1]&LOWER_MASK);
      y = s_seeds[kk+M] ^ (y >> 1) ^ mag01[y & 1];
   }
   __syncthreads();

   if (kk < N-M) {
       s_seeds[kk] = y;
       if (kk == 0)
       {
	   s_seeds[N] = y;
       }
   }
   kk += N-M;
   __syncthreads();

   /* Next 224 */
   if (kk < N-M + B2) {
      y = (s_seeds[kk]&UPPER_MASK)|(s_seeds[kk+1]&LOWER_MASK);
      y = s_seeds[kk+(M-N)] ^ (y >> 1) ^ mag01[y & 1];
   }
   __syncthreads();

   if (kk < N-M + B2) {
       s_seeds[kk] = y;
   }
   kk += B2;
   __syncthreads();

   /* Last 173 */
   if (kk < N) {
       y = (s_seeds[kk]&UPPER_MASK)|(s_seeds[kk+1]&LOWER_MASK);
       y = s_seeds[kk+(M-N)] ^ (y >> 1) ^ mag01[y & 1];
   }
   __syncthreads();

   if (kk < N) {
       s_seeds[kk] = y;
   }
   __syncthreads();

   ret.x = s_seeds[2*tid];
   ret.x ^= (ret.x >> 11);				/* Tempering */
   ret.x ^= (ret.x <<  7) & TEMPER1;
   ret.x ^= (ret.x << 15) & TEMPER2;
   ret.x ^= (ret.x >> 18);
   
   ret.y = s_seeds[2*tid+1];
   ret.y ^= (ret.y >> 11);
   ret.y ^= (ret.y <<  7) & TEMPER1;
   ret.y ^= (ret.y << 15) & TEMPER2;
   ret.y ^= (ret.y >> 18);

   return ret;
}
/*******************************************************************************
 * For reference this is the original C single threaded source: */
#if 0
static unsigned long mt[N]; /* the array for the state vector  */
static int mti=N+1; /* mti==N+1 means mt[N] is not initialized */

/* initializes mt[N] with a seed */
void init_genrand(unsigned long s)
{
    mt[0]= s & 0xffffffffUL;
    for (mti=1; mti<N; mti++) {
        mt[mti] = 
	    (1812433253UL * (mt[mti-1] ^ (mt[mti-1] >> 30)) + mti); 
        /* See Knuth TAOCP Vol2. 3rd Ed. P.106 for multiplier. */
        /* In the previous versions, MSBs of the seed affect   */
        /* only MSBs of the array mt[].                        */
        /* 2002/01/09 modified by Makoto Matsumoto             */
        mt[mti] &= 0xffffffffUL;
        /* for >32 bit machines */
    }
}

/* initialize by an array with array-length */
/* init_key is the array for initializing keys */
/* key_length is its length */
/* slight change for C++, 2004/2/26 */
void init_by_array(unsigned long init_key[], int key_length)
{
    int i, j, k;
    init_genrand(19650218UL);
    i=1; j=0;
    k = (N>key_length ? N : key_length);
    for (; k; k--) {
        mt[i] = (mt[i] ^ ((mt[i-1] ^ (mt[i-1] >> 30)) * 1664525UL))
          + init_key[j] + j; /* non linear */
        mt[i] &= 0xffffffffUL; /* for WORDSIZE > 32 machines */
        i++; j++;
        if (i>=N) { mt[0] = mt[N-1]; i=1; }
        if (j>=key_length) j=0;
    }
    for (k=N-1; k; k--) {
        mt[i] = (mt[i] ^ ((mt[i-1] ^ (mt[i-1] >> 30)) * 1566083941UL))
          - i; /* non linear */
        mt[i] &= 0xffffffffUL; /* for WORDSIZE > 32 machines */
        i++;
        if (i>=N) { mt[0] = mt[N-1]; i=1; }
    }

    mt[0] = 0x80000000UL; /* MSB is 1; assuring non-zero initial array */ 
}

/* generates a random number on [0,0xffffffff]-interval */
unsigned long genrand_int32(void)
{
    unsigned long y;
    static unsigned long mag01[2]={0x0UL, MATRIX_A};
    /* mag01[x] = x * MATRIX_A  for x=0,1 */

    if (mti >= N) { /* generate N words at one time */
        int kk;

        if (mti == N+1)   /* if init_genrand() has not been called, */
            init_genrand(5489UL); /* a default initial seed is used */

        for (kk=0;kk<N-M;kk++) {
            y = (mt[kk]&UPPER_MASK)|(mt[kk+1]&LOWER_MASK);
            mt[kk] = mt[kk+M] ^ (y >> 1) ^ mag01[y & 0x1UL];
        }
        for (;kk<N-1;kk++) {
            y = (mt[kk]&UPPER_MASK)|(mt[kk+1]&LOWER_MASK);
            mt[kk] = mt[kk+(M-N)] ^ (y >> 1) ^ mag01[y & 0x1UL];
        }
        y = (mt[N-1]&UPPER_MASK)|(mt[0]&LOWER_MASK);
        mt[N-1] = mt[M-1] ^ (y >> 1) ^ mag01[y & 0x1UL];

        mti = 0;
    }
  
    y = mt[mti++];

    /* Tempering */
    y ^= (y >> 11);
    y ^= (y << 7) & 0x9d2c5680UL;
    y ^= (y << 15) & 0xefc60000UL;
    y ^= (y >> 18);

    return y;
}
#endif


	.text
	.file	"mt19937_ref.hip"
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB0_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB0_2:
	movl	$g_seeds, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$638976, %r9d                   # imm = 0x9C000
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$mag01, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movl	$8, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$1
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$__hip_module_dtor, %edi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end0:
	.size	__hip_module_ctor, .Lfunc_end0-__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	.LBB1_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
.LBB1_2:
	retq
.Lfunc_end1:
	.size	__hip_module_dtor, .Lfunc_end1-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	g_seeds,@object                 # @g_seeds
	.local	g_seeds
	.comm	g_seeds,638976,16
	.type	mag01,@object                   # @mag01
	.local	mag01
	.comm	mag01,8,4
	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"g_seeds"
	.size	.L__unnamed_1, 8

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"mag01"
	.size	.L__unnamed_2, 6

	.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 __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym g_seeds
	.addrsig_sym mag01
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.protected	g_seeds                 ; @g_seeds
	.type	g_seeds,@object
	.section	.bss,"aw",@nobits
	.globl	g_seeds
	.p2align	4, 0x0
g_seeds:
	.zero	638976
	.size	g_seeds, 638976

	.protected	mag01                   ; @mag01
	.type	mag01,@object
	.data
	.globl	mag01
	.p2align	2, 0x0
mag01:
	.long	0                               ; 0x0
	.long	2567483615                      ; 0x9908b0df
	.size	mag01, 8

	.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 g_seeds
	.addrsig_sym mag01
	.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
