__device__ static int hash[] = {208, 34,  231, 213, 32,  248, 233, 56,  161, 78,  24,  140, 71,  48,  140, 254, 245, 255, 247, 247, 40,  185, 248, 251, 245, 28,  124, 204, 204, 76,  36,  1,   107, 28,  234, 163, 202, 224, 245, 128, 167, 204, 9,
    92,  217, 54,  239, 174, 173, 102, 193, 189, 190, 121, 100, 108, 167, 44,  43,  77,  180, 204, 8,   81,  70,  223, 11,  38,  24,  254, 210, 210, 177, 32,  81,  195, 243, 125, 8,   169, 112, 32,  97,  53,  195, 13,
    203, 9,   47,  104, 125, 117, 114, 124, 165, 203, 181, 235, 193, 206, 70,  180, 174, 0,   167, 181, 41,  164, 30,  116, 127, 198, 245, 146, 87,  224, 149, 206, 57,  4,   192, 210, 65,  210, 129, 240, 178, 105, 228,
    108, 245, 148, 140, 40,  35,  195, 38,  58,  65,  207, 215, 253, 65,  85,  208, 76,  62,  3,   237, 55,  89,  232, 50,  217, 64,  244, 157, 199, 121, 252, 90,  17,  212, 203, 149, 152, 140, 187, 234, 177, 73,  174,
    193, 100, 192, 143, 97,  53,  145, 135, 19,  103, 13,  90,  135, 151, 199, 91,  239, 247, 33,  39,  145, 101, 120, 99,  3,   186, 86,  99,  41,  237, 203, 111, 79,  220, 135, 158, 42,  30,  154, 120, 67,  87,  167,
    135, 176, 183, 191, 253, 115, 184, 21,  233, 58,  129, 233, 142, 39,  128, 211, 118, 137, 139, 255, 114, 20,  218, 113, 154, 27,  127, 246, 250, 1,   8,   198, 250, 209, 92,  222, 173, 21,  88,  102, 219};

__device__ int noise2_cuda( int x, int y, int seed ) {
    int tmp = hash[ ( unsigned int )( y + seed ) % 256 ];
    return hash[ ( unsigned int )( tmp + x ) % 256 ];
}

__device__ int noise3_cuda( int x, int y, int z, int seed ) {
    int tmp = hash[ ( unsigned int )( y + seed ) % 256 ];
    int tmp2 = hash[ ( unsigned int )( tmp + z ) % 256 ];
    return hash[ ( unsigned int )( tmp2 + x ) % 256 ];
}

__device__ float lin_inter_cuda( float x, float y, float s ) {
    return x + s * ( y - x );
}

__device__ float smooth_inter_cuda( float x, float y, float s ) {
    return lin_inter_cuda( x, y, s * s * ( 3 - 2 * s ) );
}

__device__ float noise2d_cuda( float x, float y, int seed ) {
    int x_int = floor( x );
    int y_int = floor( y );
    float x_frac = x - x_int;
    float y_frac = y - y_int;
    int s = noise2_cuda( x_int, y_int, seed );
    int t = noise2_cuda( x_int + 1, y_int, seed );
    int u = noise2_cuda( x_int, y_int + 1, seed );
    int v = noise2_cuda( x_int + 1, y_int + 1, seed );
    float low = smooth_inter_cuda( s, t, x_frac );
    float high = smooth_inter_cuda( u, v, x_frac );
    return smooth_inter_cuda( low, high, y_frac );
}

__device__ float noise3d_cuda( float x, float y, float z, int seed ) {
    int x_int = floor( x );
    int y_int = floor( y );
    int z_int = floor( z );
    float x_frac = x - x_int;
    float y_frac = y - y_int;
    float z_frac = z - z_int;

    int o = noise3_cuda( x_int, y_int, z_int, seed );
    int p = noise3_cuda( x_int + 1, y_int, z_int, seed );
    int q = noise3_cuda( x_int, y_int + 1, z_int, seed );
    int r = noise3_cuda( x_int + 1, y_int + 1, z_int, seed );

    float low1 = smooth_inter_cuda( o, p, x_frac );
    float high1 = smooth_inter_cuda( q, r, x_frac );
    float down = smooth_inter_cuda( low1, high1, y_frac );

    int s = noise3_cuda( x_int, y_int, z_int + 1, seed );
    int t = noise3_cuda( x_int + 1, y_int, z_int + 1, seed );
    int u = noise3_cuda( x_int, y_int + 1, z_int + 1, seed );
    int v = noise3_cuda( x_int + 1, y_int + 1, z_int + 1, seed );

    float low2 = smooth_inter_cuda( s, t, x_frac );
    float high2 = smooth_inter_cuda( u, v, x_frac );
    float up = smooth_inter_cuda( low2, high2, y_frac );
    return smooth_inter_cuda( down, up, z_frac );
}

__device__ float perlin_noise2d_cuda( float x, float y, float freq, int depth, int seed ){
    float xa = x * freq;
    float ya = y * freq;
    float amp = 1.0;
    float fin = 0;
    float div = 0.0;
    int i;
    for ( i = 0; i < depth; i++ ) {
        div += 256 * amp;
        fin += noise2d_cuda( xa, ya, seed ) * amp;
        amp /= 2;
        xa *= 2;
        ya *= 2;
    }
    return fin / div;
}

__device__ float perlin_noise3d_cuda( float x, float y, float z, float freq, int depth, int seed ) {
    float xa = x * freq;
    float ya = y * freq;
    float za = z * freq;
    float amp = 1.0;
    float fin = 0;
    float div = 0.0;

    int i;
    for ( i = 0; i < depth; i++ ) {
        div += 256 * amp;
        fin += noise3d_cuda( xa, ya, za, seed ) * amp;
        amp /= 2;
        xa *= 2;
        ya *= 2;
        za *= 2;
    }

    return fin / div;
}
	.file	"tmpxft_00393414_00000000-6_perlin_noise.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2018:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2018:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z11noise2_cudaiii
	.type	_Z11noise2_cudaiii, @function
_Z11noise2_cudaiii:
.LFB2008:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2008:
	.size	_Z11noise2_cudaiii, .-_Z11noise2_cudaiii
	.globl	_Z11noise3_cudaiiii
	.type	_Z11noise3_cudaiiii, @function
_Z11noise3_cudaiiii:
.LFB2009:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2009:
	.size	_Z11noise3_cudaiiii, .-_Z11noise3_cudaiiii
	.globl	_Z14lin_inter_cudafff
	.type	_Z14lin_inter_cudafff, @function
_Z14lin_inter_cudafff:
.LFB2010:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2010:
	.size	_Z14lin_inter_cudafff, .-_Z14lin_inter_cudafff
	.globl	_Z17smooth_inter_cudafff
	.type	_Z17smooth_inter_cudafff, @function
_Z17smooth_inter_cudafff:
.LFB2157:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2157:
	.size	_Z17smooth_inter_cudafff, .-_Z17smooth_inter_cudafff
	.globl	_Z12noise2d_cudaffi
	.type	_Z12noise2d_cudaffi, @function
_Z12noise2d_cudaffi:
.LFB2012:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2012:
	.size	_Z12noise2d_cudaffi, .-_Z12noise2d_cudaffi
	.globl	_Z12noise3d_cudafffi
	.type	_Z12noise3d_cudafffi, @function
_Z12noise3d_cudafffi:
.LFB2013:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2013:
	.size	_Z12noise3d_cudafffi, .-_Z12noise3d_cudafffi
	.globl	_Z19perlin_noise2d_cudafffii
	.type	_Z19perlin_noise2d_cudafffii, @function
_Z19perlin_noise2d_cudafffii:
.LFB2014:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2014:
	.size	_Z19perlin_noise2d_cudafffii, .-_Z19perlin_noise2d_cudafffii
	.globl	_Z19perlin_noise3d_cudaffffii
	.type	_Z19perlin_noise3d_cudaffffii, @function
_Z19perlin_noise3d_cudaffffii:
.LFB2015:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2015:
	.size	_Z19perlin_noise3d_cudaffffii, .-_Z19perlin_noise3d_cudaffffii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"hash"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2041:
	.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
	movl	$1024, %r9d
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	leaq	_ZL4hash(%rip), %rsi
	movq	%rdx, %rcx
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 24
	popq	%rcx
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rsi
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2041:
	.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	_ZL4hash
	.comm	_ZL4hash,1024,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>
__device__ static int hash[] = {208, 34,  231, 213, 32,  248, 233, 56,  161, 78,  24,  140, 71,  48,  140, 254, 245, 255, 247, 247, 40,  185, 248, 251, 245, 28,  124, 204, 204, 76,  36,  1,   107, 28,  234, 163, 202, 224, 245, 128, 167, 204, 9,
    92,  217, 54,  239, 174, 173, 102, 193, 189, 190, 121, 100, 108, 167, 44,  43,  77,  180, 204, 8,   81,  70,  223, 11,  38,  24,  254, 210, 210, 177, 32,  81,  195, 243, 125, 8,   169, 112, 32,  97,  53,  195, 13,
    203, 9,   47,  104, 125, 117, 114, 124, 165, 203, 181, 235, 193, 206, 70,  180, 174, 0,   167, 181, 41,  164, 30,  116, 127, 198, 245, 146, 87,  224, 149, 206, 57,  4,   192, 210, 65,  210, 129, 240, 178, 105, 228,
    108, 245, 148, 140, 40,  35,  195, 38,  58,  65,  207, 215, 253, 65,  85,  208, 76,  62,  3,   237, 55,  89,  232, 50,  217, 64,  244, 157, 199, 121, 252, 90,  17,  212, 203, 149, 152, 140, 187, 234, 177, 73,  174,
    193, 100, 192, 143, 97,  53,  145, 135, 19,  103, 13,  90,  135, 151, 199, 91,  239, 247, 33,  39,  145, 101, 120, 99,  3,   186, 86,  99,  41,  237, 203, 111, 79,  220, 135, 158, 42,  30,  154, 120, 67,  87,  167,
    135, 176, 183, 191, 253, 115, 184, 21,  233, 58,  129, 233, 142, 39,  128, 211, 118, 137, 139, 255, 114, 20,  218, 113, 154, 27,  127, 246, 250, 1,   8,   198, 250, 209, 92,  222, 173, 21,  88,  102, 219};

__device__ int noise2_cuda( int x, int y, int seed ) {
    int tmp = hash[ ( unsigned int )( y + seed ) % 256 ];
    return hash[ ( unsigned int )( tmp + x ) % 256 ];
}

__device__ int noise3_cuda( int x, int y, int z, int seed ) {
    int tmp = hash[ ( unsigned int )( y + seed ) % 256 ];
    int tmp2 = hash[ ( unsigned int )( tmp + z ) % 256 ];
    return hash[ ( unsigned int )( tmp2 + x ) % 256 ];
}

__device__ float lin_inter_cuda( float x, float y, float s ) {
    return x + s * ( y - x );
}

__device__ float smooth_inter_cuda( float x, float y, float s ) {
    return lin_inter_cuda( x, y, s * s * ( 3 - 2 * s ) );
}

__device__ float noise2d_cuda( float x, float y, int seed ) {
    int x_int = floor( x );
    int y_int = floor( y );
    float x_frac = x - x_int;
    float y_frac = y - y_int;
    int s = noise2_cuda( x_int, y_int, seed );
    int t = noise2_cuda( x_int + 1, y_int, seed );
    int u = noise2_cuda( x_int, y_int + 1, seed );
    int v = noise2_cuda( x_int + 1, y_int + 1, seed );
    float low = smooth_inter_cuda( s, t, x_frac );
    float high = smooth_inter_cuda( u, v, x_frac );
    return smooth_inter_cuda( low, high, y_frac );
}

__device__ float noise3d_cuda( float x, float y, float z, int seed ) {
    int x_int = floor( x );
    int y_int = floor( y );
    int z_int = floor( z );
    float x_frac = x - x_int;
    float y_frac = y - y_int;
    float z_frac = z - z_int;

    int o = noise3_cuda( x_int, y_int, z_int, seed );
    int p = noise3_cuda( x_int + 1, y_int, z_int, seed );
    int q = noise3_cuda( x_int, y_int + 1, z_int, seed );
    int r = noise3_cuda( x_int + 1, y_int + 1, z_int, seed );

    float low1 = smooth_inter_cuda( o, p, x_frac );
    float high1 = smooth_inter_cuda( q, r, x_frac );
    float down = smooth_inter_cuda( low1, high1, y_frac );

    int s = noise3_cuda( x_int, y_int, z_int + 1, seed );
    int t = noise3_cuda( x_int + 1, y_int, z_int + 1, seed );
    int u = noise3_cuda( x_int, y_int + 1, z_int + 1, seed );
    int v = noise3_cuda( x_int + 1, y_int + 1, z_int + 1, seed );

    float low2 = smooth_inter_cuda( s, t, x_frac );
    float high2 = smooth_inter_cuda( u, v, x_frac );
    float up = smooth_inter_cuda( low2, high2, y_frac );
    return smooth_inter_cuda( down, up, z_frac );
}

__device__ float perlin_noise2d_cuda( float x, float y, float freq, int depth, int seed ){
    float xa = x * freq;
    float ya = y * freq;
    float amp = 1.0;
    float fin = 0;
    float div = 0.0;
    int i;
    for ( i = 0; i < depth; i++ ) {
        div += 256 * amp;
        fin += noise2d_cuda( xa, ya, seed ) * amp;
        amp /= 2;
        xa *= 2;
        ya *= 2;
    }
    return fin / div;
}

__device__ float perlin_noise3d_cuda( float x, float y, float z, float freq, int depth, int seed ) {
    float xa = x * freq;
    float ya = y * freq;
    float za = z * freq;
    float amp = 1.0;
    float fin = 0;
    float div = 0.0;

    int i;
    for ( i = 0; i < depth; i++ ) {
        div += 256 * amp;
        fin += noise3d_cuda( xa, ya, za, seed ) * amp;
        amp /= 2;
        xa *= 2;
        ya *= 2;
        za *= 2;
    }

    return fin / div;
}
	.text
	.file	"perlin_noise.hip"
	.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_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
