﻿#include "device_launch_parameters.h"
#include "curand_kernel.h"
#include "dev_noise.cuh"

__global__ void cudaNoiseGeneWithSoS(float *dev_cos_value, float *dev_sin_value, unsigned int length, unsigned int path_num,
	unsigned long long uniform_seed, float omega_amp, float delta_alpha, float delta_omega, float delta_t, float sum_amp){
	extern __shared__ float _sha[];
	unsigned int tidx = blockIdx.x * blockDim.x + threadIdx.x;
	unsigned int tidy = blockIdx.y * blockDim.y + threadIdx.y;

	unsigned int x = tidx % path_num;
	unsigned int y = (tidy * blockDim.x * gridDim.x + tidx) / path_num;

	float omega_n_I = omega_amp * cosf(delta_alpha * x) + delta_omega;
	float omega_n_Q = omega_amp * sinf(delta_alpha * x) + delta_omega;

	curandState_t rand_status;
	curand_init(uniform_seed, x, 0, &rand_status);
	float phi_n_I = curand_uniform(&rand_status);
	float phi_n_Q = curand_uniform(&rand_status);

	float *cos_value = _sha, *sin_value = _sha + path_num;
	unsigned int y_step = gridDim.y * blockDim.x * gridDim.x / path_num;
	for (; y < length; y += y_step){
		cos_value[x] = cosf(omega_n_I * delta_t*y + 2 * CR_CUDART_PI*phi_n_I);
		sin_value[x] = sinf(omega_n_Q * delta_t*y + 2 * CR_CUDART_PI*phi_n_Q);
		__syncthreads();

		for (path_num >>= 1; path_num > 0; path_num >>= 1){
			if (x < path_num){
				cos_value[x] += cos_value[x + path_num];
				sin_value[x] += sin_value[x + path_num];
			}
			__syncthreads();
		}

		if (x == 0){
			dev_cos_value[y] = sum_amp * cos_value[x];
			dev_sin_value[y] = sum_amp * sin_value[x];
		}
		__syncthreads();
	}
}

__global__ void cudaNoiseOmegaCulc(float *dev_omega_n_I, float *dev_omega_n_Q, unsigned int length,
	float omega_amp, float delta_alpha, float delta_omega){
	unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
	if (tid < length){
		dev_omega_n_I[tid] = omega_amp*cosf(delta_alpha*tid) + delta_omega;
		dev_omega_n_Q[tid] = omega_amp*sinf(delta_alpha*tid) + delta_omega;
	}
}

__global__ void cudaNoiseSoSCulc(float *dev_cos_value, float *dev_sin_value,
	unsigned int pitch_width, unsigned int width, unsigned int heigth, float delta_t,
	float *dev_omega_n_I, float *dev_omega_n_Q, float *dev_phi_n){
	unsigned int x = threadIdx.x,
		y = threadIdx.y,
		tidy = blockIdx.y * blockDim.y + y;
	__shared__ float sha_cos_value[BLOCK_DIM_Y][BLOCK_DIM_X],
		sha_sin_value[BLOCK_DIM_Y][BLOCK_DIM_X];
	sha_cos_value[y][x] = 0;
	sha_sin_value[y][x] = 0;
	__syncthreads();

	if (tidy < heigth){
		__shared__ float sha_omega_n_I[BLOCK_DIM_Y], sha_omega_n_Q[BLOCK_DIM_Y],
			sha_phi_n_I[BLOCK_DIM_Y], sha_phi_n_Q[BLOCK_DIM_Y];

		if (x == 0){
			sha_omega_n_I[y] = dev_omega_n_I[tidy];
			sha_omega_n_Q[y] = dev_omega_n_Q[tidy];
			sha_phi_n_I[y] = dev_phi_n[tidy];
			sha_phi_n_Q[y] = dev_phi_n[heigth + tidy];
		}
		__syncthreads();

		for (unsigned int tidx = blockIdx.x * blockDim.x + x;
			tidx < width; tidx += gridDim.x*blockDim.x){
			sha_cos_value[y][x] = cosf(sha_omega_n_I[y] * delta_t*tidx + 2 * CR_CUDART_PI*sha_phi_n_I[y]);
			sha_sin_value[y][x] = sinf(sha_omega_n_Q[y] * delta_t*tidx + 2 * CR_CUDART_PI*sha_phi_n_Q[y]);
			__syncthreads();

			for (unsigned int heigth_ii = blockDim.y / 2, extra = blockDim.y % 2;
				heigth_ii > 0; extra = heigth_ii % 2, heigth_ii /= 2){
				if (y < heigth_ii){
					sha_cos_value[y][x] += sha_cos_value[heigth_ii + extra + y][x];
					sha_sin_value[y][x] += sha_sin_value[heigth_ii + extra + y][x];
				}
				heigth_ii += extra;
				__syncthreads();
			}
			if (y == 0){
				unsigned int loc = blockIdx.y*pitch_width + tidx;
				dev_cos_value[loc] = sha_cos_value[0][x];
				dev_sin_value[loc] = sha_sin_value[0][x];
			}
		}
	}
}

__global__ void cudaNoiseSoSSum(float *dev_cos_value, float *dev_sin_value,
	unsigned int pitch_width, unsigned int width, unsigned int heigth, float sum_amp){
	unsigned int loc;
	float reg_cos_value, reg_sin_value;

	for (unsigned int tidx = blockIdx.x*blockDim.x + threadIdx.x;
		tidx < width; tidx += gridDim.x*blockDim.x){
		reg_cos_value = 0;
		reg_sin_value = 0;
		for (unsigned int heigth_ii = 0; heigth_ii < heigth; heigth_ii++){
			loc = heigth_ii*pitch_width + tidx;
			reg_cos_value += dev_cos_value[loc];
			reg_sin_value += dev_sin_value[loc];
		}

		dev_cos_value[tidx] = sum_amp * reg_cos_value;
		dev_sin_value[tidx] = sum_amp * reg_sin_value;
	}
}
	.file	"tmpxft_002a3b4c_00000000-6_dev_noise.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2224:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2224:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z50__device_stub__Z20cudaNoiseGeneWithSoSPfS_jjyfffffPfS_jjyfffff
	.type	_Z50__device_stub__Z20cudaNoiseGeneWithSoSPfS_jjyfffffPfS_jjyfffff, @function
_Z50__device_stub__Z20cudaNoiseGeneWithSoSPfS_jjyfffffPfS_jjyfffff:
.LFB2246:
	.cfi_startproc
	endbr64
	subq	$232, %rsp
	.cfi_def_cfa_offset 240
	movq	%rdi, 56(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	100(%rsp), %rsi
	movl	%edx, 44(%rsp)
	leaq	72(%rsp), %rdx
	movl	%ecx, 40(%rsp)
	leaq	80(%rsp), %rcx
	movq	%r8, 32(%rsp)
	movss	%xmm0, 28(%rsp)
	movss	%xmm1, 24(%rsp)
	movss	%xmm2, 20(%rsp)
	movss	%xmm3, 16(%rsp)
	movss	%xmm4, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 216(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 96(%rsp)
	movq	%rax, 136(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	44(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 208(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	80(%rsp)
	.cfi_def_cfa_offset 248
	leaq	_Z20cudaNoiseGeneWithSoSPfS_jjyfffff(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 256
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 248
	popq	%rdx
	.cfi_def_cfa_offset 240
.L2:
	movq	216(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$232, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2246:
	.size	_Z50__device_stub__Z20cudaNoiseGeneWithSoSPfS_jjyfffffPfS_jjyfffff, .-_Z50__device_stub__Z20cudaNoiseGeneWithSoSPfS_jjyfffffPfS_jjyfffff
	.globl	_Z20cudaNoiseGeneWithSoSPfS_jjyfffff
	.type	_Z20cudaNoiseGeneWithSoSPfS_jjyfffff, @function
_Z20cudaNoiseGeneWithSoSPfS_jjyfffff:
.LFB2247:
	.cfi_startproc
	endbr64
	jmp	_Z50__device_stub__Z20cudaNoiseGeneWithSoSPfS_jjyfffffPfS_jjyfffff
	.cfi_endproc
.LFE2247:
	.size	_Z20cudaNoiseGeneWithSoSPfS_jjyfffff, .-_Z20cudaNoiseGeneWithSoSPfS_jjyfffff
	.globl	_Z44__device_stub__Z18cudaNoiseOmegaCulcPfS_jfffPfS_jfff
	.type	_Z44__device_stub__Z18cudaNoiseOmegaCulcPfS_jfffPfS_jfff, @function
_Z44__device_stub__Z18cudaNoiseOmegaCulcPfS_jfffPfS_jfff:
.LFB2248:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	40(%rsp), %rdx
	movss	%xmm0, 8(%rsp)
	movss	%xmm1, 4(%rsp)
	movss	%xmm2, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%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	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 136(%rsp)
	movq	%rsp, %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	48(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z18cudaNoiseOmegaCulcPfS_jfff(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 192
	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 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L8:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2248:
	.size	_Z44__device_stub__Z18cudaNoiseOmegaCulcPfS_jfffPfS_jfff, .-_Z44__device_stub__Z18cudaNoiseOmegaCulcPfS_jfffPfS_jfff
	.globl	_Z18cudaNoiseOmegaCulcPfS_jfff
	.type	_Z18cudaNoiseOmegaCulcPfS_jfff, @function
_Z18cudaNoiseOmegaCulcPfS_jfff:
.LFB2249:
	.cfi_startproc
	endbr64
	jmp	_Z44__device_stub__Z18cudaNoiseOmegaCulcPfS_jfffPfS_jfff
	.cfi_endproc
.LFE2249:
	.size	_Z18cudaNoiseOmegaCulcPfS_jfff, .-_Z18cudaNoiseOmegaCulcPfS_jfff
	.globl	_Z48__device_stub__Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_PfS_jjjfS_S_S_
	.type	_Z48__device_stub__Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_PfS_jjjfS_S_S_, @function
_Z48__device_stub__Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_PfS_jjjfS_S_S_:
.LFB2250:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	224(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	80(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	92(%rsp), %rsi
	movq	%rax, 16(%rsp)
	movq	232(%rsp), %rax
	movl	%edx, 44(%rsp)
	leaq	64(%rsp), %rdx
	movl	%ecx, 40(%rsp)
	leaq	72(%rsp), %rcx
	movl	%r8d, 36(%rsp)
	movq	%r9, 24(%rsp)
	movq	%rax, 8(%rsp)
	movss	%xmm0, 32(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 88(%rsp)
	movq	%rax, 128(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	44(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 80(%rsp)
	movq	%rax, 92(%rsp)
	movl	$1, 100(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	72(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_(%rip), %rdi
	pushq	72(%rsp)
	.cfi_def_cfa_offset 240
	movq	108(%rsp), %rcx
	movl	116(%rsp), %r8d
	movq	96(%rsp), %rsi
	movl	104(%rsp), %edx
	leaq	144(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L13:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2250:
	.size	_Z48__device_stub__Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_PfS_jjjfS_S_S_, .-_Z48__device_stub__Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_PfS_jjjfS_S_S_
	.globl	_Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_
	.type	_Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_, @function
_Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_:
.LFB2251:
	.cfi_startproc
	endbr64
	jmp	_Z48__device_stub__Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_PfS_jjjfS_S_S_
	.cfi_endproc
.LFE2251:
	.size	_Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_, .-_Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_
	.globl	_Z41__device_stub__Z15cudaNoiseSoSSumPfS_jjjfPfS_jjjf
	.type	_Z41__device_stub__Z15cudaNoiseSoSSumPfS_jjjfPfS_jjjf, @function
_Z41__device_stub__Z15cudaNoiseSoSSumPfS_jjjfPfS_jjjf:
.LFB2252:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movss	%xmm0, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%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	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 136(%rsp)
	movq	%rsp, %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	pushq	48(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z15cudaNoiseSoSSumPfS_jjjf(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 192
	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 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L18:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2252:
	.size	_Z41__device_stub__Z15cudaNoiseSoSSumPfS_jjjfPfS_jjjf, .-_Z41__device_stub__Z15cudaNoiseSoSSumPfS_jjjfPfS_jjjf
	.globl	_Z15cudaNoiseSoSSumPfS_jjjf
	.type	_Z15cudaNoiseSoSSumPfS_jjjf, @function
_Z15cudaNoiseSoSSumPfS_jjjf:
.LFB2253:
	.cfi_startproc
	endbr64
	jmp	_Z41__device_stub__Z15cudaNoiseSoSSumPfS_jjjfPfS_jjjf
	.cfi_endproc
.LFE2253:
	.size	_Z15cudaNoiseSoSSumPfS_jjjf, .-_Z15cudaNoiseSoSSumPfS_jjjf
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z15cudaNoiseSoSSumPfS_jjjf"
.LC1:
	.string	"_Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_"
.LC2:
	.string	"_Z18cudaNoiseOmegaCulcPfS_jfff"
.LC3:
	.string	"_Z20cudaNoiseGeneWithSoSPfS_jjyfffff"
.LC4:
	.string	"precalc_xorwow_matrix"
.LC5:
	.string	"precalc_xorwow_offset_matrix"
.LC6:
	.string	"mrg32k3aM1"
.LC7:
	.string	"mrg32k3aM2"
.LC8:
	.string	"mrg32k3aM1SubSeq"
.LC9:
	.string	"mrg32k3aM2SubSeq"
.LC10:
	.string	"mrg32k3aM1Seq"
.LC11:
	.string	"mrg32k3aM2Seq"
.LC12:
	.string	"__cr_lgamma_table"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2255:
	.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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z15cudaNoiseSoSSumPfS_jjjf(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z16cudaNoiseSoSCulcPfS_jjjfS_S_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z18cudaNoiseOmegaCulcPfS_jfff(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z20cudaNoiseGeneWithSoSPfS_jjyfffff(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC4(%rip), %rdx
	movl	$102400, %r9d
	leaq	_ZL21precalc_xorwow_matrix(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	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	$0
	.cfi_def_cfa_offset 32
	leaq	.LC5(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$102400, %r9d
	leaq	_ZL28precalc_xorwow_offset_matrix(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC6(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL10mrg32k3aM1(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC7(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL10mrg32k3aM2(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC8(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$2016, %r9d
	leaq	_ZL16mrg32k3aM1SubSeq(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rax
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC9(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$2016, %r9d
	leaq	_ZL16mrg32k3aM2SubSeq(%rip), %rsi
	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	$0
	.cfi_def_cfa_offset 32
	leaq	.LC10(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL13mrg32k3aM1Seq(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC11(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL13mrg32k3aM2Seq(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC12(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$72, %r9d
	movq	%rdx, %rcx
	leaq	_ZL17__cr_lgamma_table(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.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
.LFE2255:
	.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	_ZL17__cr_lgamma_table
	.comm	_ZL17__cr_lgamma_table,72,32
	.local	_ZL13mrg32k3aM2Seq
	.comm	_ZL13mrg32k3aM2Seq,2304,32
	.local	_ZL13mrg32k3aM1Seq
	.comm	_ZL13mrg32k3aM1Seq,2304,32
	.local	_ZL16mrg32k3aM2SubSeq
	.comm	_ZL16mrg32k3aM2SubSeq,2016,32
	.local	_ZL16mrg32k3aM1SubSeq
	.comm	_ZL16mrg32k3aM1SubSeq,2016,32
	.local	_ZL10mrg32k3aM2
	.comm	_ZL10mrg32k3aM2,2304,32
	.local	_ZL10mrg32k3aM1
	.comm	_ZL10mrg32k3aM1,2304,32
	.local	_ZL28precalc_xorwow_offset_matrix
	.comm	_ZL28precalc_xorwow_offset_matrix,102400,32
	.local	_ZL21precalc_xorwow_matrix
	.comm	_ZL21precalc_xorwow_matrix,102400,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:

﻿#pragma once
#include "hip/hip_runtime.h"

#define BLOCK_DIM_X 32
#define BLOCK_DIM_Y 8
#define GRID_DIM_X 65535

__global__ void cudaNoiseGeneWithSoS(float *dev_cos_value, float *dev_sin_value, unsigned int length, unsigned int path_num,
	unsigned long long uniform_seed, float omega_amp, float delta_alpha, float delta_omega, float delta_t, float sum_amp);

__global__ void cudaNoiseOmegaCulc(float *dev_omega_n_I, float *dev_omega_n_Q, unsigned int length,
	float omega_amp, float delta_alpha, float delta_omega);

__global__ void cudaNoiseSoSCulc(float *dev_cos_value, float *dev_sin_value,
	unsigned int pitch_width, unsigned int width, unsigned int heigth, float delta_t,
	float *dev_omega_n_I, float *dev_omega_n_Q, float *dev_phi_n);

__global__ void cudaNoiseSoSSum(float *dev_cos_value, float *dev_sin_value,
	unsigned int pitch_width, unsigned int width, unsigned int heigth, float sum_amp);
	.text
	.file	"dev_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
