#include "simple_particle.cuh"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include "math_functions.h"
#include <stdio.h>

__constant__ simpleParticleSystem d_sps[1];

__global__ void generateParticles();

__global__ void renderParticles(uchar4* devPtr, int img_width, int img_height);

__global__ void updateParticles(float passed_time);

__device__ float2 get_normal_vector(float rand_num);

__device__ float get_energy(float2 p1, float2 p2, float dist_bound_powerd);

__device__ uchar4 get_color_from_energy(float energy);

__device__ float2 get_acceleration(int index);

__device__ void update_particle_velocity(int index, float2 acc, float passed_time);

__device__ int update_particle_position(int index, float passed_time); //return whether the particle is dead

void init_particles_cuda(simpleParticleSystem &sps) {
	int max_num_particles = sps.MAX_PARTICLE_SIZE;
	int one_batch_num_particles = sps.ONE_BATCH_PARTICLE_SIZE;

	cudaMalloc((void**)&sps.energy, sizeof(*sps.energy)*max_num_particles);
	cudaMalloc((void**)&sps.position, sizeof(*sps.position)*max_num_particles);
	cudaMalloc((void**)&sps.velocity, sizeof(*sps.velocity)*max_num_particles);
	cudaMalloc((void**)&sps.remain_time, sizeof(*sps.remain_time)*max_num_particles);
	cudaMalloc((void**)&sps.rand_data, sizeof(*sps.rand_data)*one_batch_num_particles*3);

	cudaError_t err = cudaGetLastError();
	if (err != cudaSuccess)
		printf("Memory Allocation Error: %s\n", cudaGetErrorString(err));
}

void destroy_particles_cuda(simpleParticleSystem &sps) {
	cudaError_t er;

	er = cudaFree(sps.energy);
	er = cudaFree(sps.position);
	er = cudaFree(sps.velocity);
	er = cudaFree(sps.remain_time);
	er = cudaFree(sps.rand_data);

	cudaError_t err = cudaGetLastError();
	if (err != cudaSuccess)
		printf("Memory Free Error: %s\n", cudaGetErrorString(err));
}

void copy_to_device_sps(simpleParticleSystem &sps) {
	cudaError_t err = cudaMemcpyToSymbol(d_sps, &sps, sizeof(simpleParticleSystem));

	if (err != cudaSuccess)
		printf("Constant Memory Copy Error: %s\n", cudaGetErrorString(err));
}

void generate_particles(int generate_size) {
	generateParticles << < 1, generate_size >> > ();
	//generateParticlesLine <<< 1, sps.ONE_BATCH_PARTICLE_SIZE >>> (
	//	sps.position, sps.velocity_orientation, sps.velocity, sps.remain_time, sps.rand_data, sps.ONE_BATCH_PARTICLE_SIZE,
	//	sps.MAX_PARTICLE_SIZE, sps.generator_line[0], sps.generator_line[1], sps.MAX_VELOCITY, sps.MIN_VELOCITY, sps.LIFE_TIME
	//);
}

void updata_particles(int generate_size, float passed_time) {
	updateParticles << < 1, generate_size >> > (passed_time);
}

void render_particles(uchar4* devPtr, int img_width, int img_height) {
	int thread_dim = 16;
	int grid_dim_x = (img_width + thread_dim - 1) / thread_dim;
	int grid_dim_y = (img_height + thread_dim - 1) / thread_dim;
	dim3 grids(grid_dim_x, grid_dim_y);
	dim3 threads(thread_dim, thread_dim);
	renderParticles << <grids, threads >> > (devPtr, img_width, img_height);
}

__global__ void generateParticles()
{
	float2 *position = (*d_sps).position;
	float2 *velocity = (*d_sps).velocity;
	float *remain_time = (*d_sps).remain_time;
	float *rand = (*d_sps).rand_data;
	int generate_size = (*d_sps).ONE_BATCH_PARTICLE_SIZE;
	int max_size = (*d_sps).MAX_PARTICLE_SIZE;

	unsigned int index = blockIdx.x*blockDim.x + threadIdx.x;

	__shared__ unsigned int generate_start_index;

	//get the particle generate block pos
	if (index == 0) {
		max_size -= generate_size;
		generate_start_index = 0;
		while (generate_start_index <= max_size) {
			if (remain_time[generate_start_index] == 0)
				break;
			generate_start_index += generate_size;
		}
	}

	__syncthreads();

	if (generate_start_index > max_size)
		return;

	int pid = generate_start_index + index; 
	float x;
	float y;
	float2 velocity_orientation;
	float n_velocity;

	//generate rand position and velocity
	switch ((*d_sps).TYPE)
	{
	case LineGenerator:
		x = rand[index] * ((*d_sps).generator_line[0].x - (*d_sps).generator_line[1].x) + (*d_sps).generator_line[1].x;
		y = rand[index] * ((*d_sps).generator_line[0].y - (*d_sps).generator_line[1].y) + (*d_sps).generator_line[1].y;
		position[pid] = make_float2(x, y);

		rand += generate_size;
		pid = generate_start_index + index;
		velocity_orientation = get_normal_vector(rand[index]);

		rand += generate_size;
		n_velocity = rand[index] * ((*d_sps).MAX_VELOCITY - (*d_sps).MIN_VELOCITY) + (*d_sps).MIN_VELOCITY;
		velocity[pid].x = n_velocity * velocity_orientation.x;
		velocity[pid].y = n_velocity * velocity_orientation.y;
		break;
	case CircleGenerator:
		float rand_pos = rand[index];
		float2 vec = get_normal_vector(rand_pos);
		x = vec.x * (*d_sps).generator_radius.x + (*d_sps).generator_center.x;
		y = vec.y * (*d_sps).generator_radius.y + (*d_sps).generator_center.y;
		position[pid] = make_float2(x, y);

		rand += generate_size;
		pid = generate_start_index + index;
		float rand_orient = rand[index];
		rand_orient = rand_pos + (rand_orient / 2 - rand_orient / 4);
		velocity_orientation = get_normal_vector(rand_orient);

		rand += generate_size;
		n_velocity = rand[index] * ((*d_sps).MAX_VELOCITY - (*d_sps).MIN_VELOCITY) + (*d_sps).MIN_VELOCITY;
		velocity[pid].x = n_velocity * velocity_orientation.x;
		velocity[pid].y = n_velocity * velocity_orientation.y;
		break;
	default:
		break;
	}

	//generate remain time
	remain_time[pid] = (*d_sps).LIFE_TIME;
}

__global__ void renderParticles(uchar4* devPtr, int img_width, int img_height) {
	int x = threadIdx.x + blockIdx.x * blockDim.x;
	int y = threadIdx.y + blockIdx.y * blockDim.y;

	if (x >= img_width || y >= img_height)
		return;

	
	if (!(x >= (*d_sps).LIFE_BOUND[0] && x <= (*d_sps).LIFE_BOUND[2]
		&& y <= (*d_sps).LIFE_BOUND[1] && y >= (*d_sps).LIFE_BOUND[3]))
		return;

	int generate_size = (*d_sps).ONE_BATCH_PARTICLE_SIZE;
	int max_size = (*d_sps).MAX_PARTICLE_SIZE;
	float energy = 0;
	float dist_bound_powerd = (*d_sps).ENERGY_SCOPE * (*d_sps).ENERGY_SCOPE;
	float2 pos = make_float2(x, y); 
	for (int start_index = 0; start_index < max_size - generate_size; start_index += generate_size)
	{
		if ((*d_sps).remain_time[start_index] != 0) {
			//here we do not render the first particle of the batch
			for (int index = start_index + 1; index < start_index + generate_size; ++index) {
				if ((*d_sps).remain_time[index] != 0) {
					energy += get_energy((*d_sps).position[index], pos, dist_bound_powerd);
					if (energy >= 1) {
						energy = 1;
						break;
					}
				}
			}
			if (energy >= 1) {
				break;
			}
		}
	}


	int offset = x + y * img_width;
	devPtr[offset] = get_color_from_energy(energy);
}


__global__ void updateParticles(float passed_time) {
	unsigned int index = blockIdx.x*blockDim.x + threadIdx.x;
	unsigned int strip = gridDim.x * blockDim.x;
	unsigned int start_index = blockIdx.x*blockDim.x;

	__shared__ int living_particle_num;

	while (index < (*d_sps).MAX_PARTICLE_SIZE) {
		living_particle_num = 0;
		//__syncthreads();

		if ((*d_sps).remain_time[start_index] != 0) {
			if (index != start_index) {
				float2 acc = get_acceleration(index);
				update_particle_velocity(index, acc, passed_time);
				int is_living = update_particle_position(index, passed_time);
				if (is_living) {
					living_particle_num += 1;
				}
			}

			__syncthreads();

			if (index == start_index) {
				if(living_particle_num == 0)
					(*d_sps).remain_time[index] = 0;
				else
					(*d_sps).remain_time[index] = 1.0;
			}
		}

		index += strip;
		start_index += strip;
	}
}


__device__ float2 get_normal_vector(float rand_num) {
	float x, y;
	sincosf(rand_num*2*PI, &y, &x);

	return make_float2(x, y);
}

__device__ float get_energy(float2 p1, float2 p2, float dist_bound_powerd) {
	float dx = p1.x - p2.x;
	float dy = p1.y - p2.y;
	float dist_powered = dx*dx + dy*dy;

	if (dist_powered > dist_bound_powerd)
		return 0;
	if (dist_powered == 0)
		return 0.1;
	return 0.1 * sqrtf(dist_bound_powerd - dist_powered)/sqrtf(dist_bound_powerd);
}

__device__ uchar4 get_color_from_energy(float energy) {
	if (energy == 0)
		return make_uchar4(0, 0, 0, 0);

	unsigned char r = 90 * energy + 160;
	unsigned char g = 180 * energy;
	unsigned char b = 60 * energy;
	unsigned char w = 255 * energy;

	return make_uchar4(r, g, b, w);
}

__device__ float2 get_acceleration(int index) {
	return make_float2(0.0, 40.0);
}

__device__ void update_particle_velocity(int index, float2 acc, float passed_time) {
	(*d_sps).velocity[index].x += acc.x * passed_time;
	(*d_sps).velocity[index].y += acc.y * passed_time;
}

__device__ int update_particle_position(int index, float passed_time) {
	(*d_sps).remain_time[index] -= passed_time;
	if ((*d_sps).remain_time[index] <= 0) {
		(*d_sps).remain_time[index] = 0;
		return 0;
	}
	
	float2 *pos = &(*d_sps).position[index];
	(*pos).x += (*d_sps).velocity[index].x * passed_time;
	(*pos).y += (*d_sps).velocity[index].y * passed_time;
	float x = (*pos).x;
	float y = (*pos).y;

	if (x > (*d_sps).LIFE_BOUND[0] && x < (*d_sps).LIFE_BOUND[2]
		&& y < (*d_sps).LIFE_BOUND[1] && y > (*d_sps).LIFE_BOUND[3]) {

		//if (x < (*d_sps).BOUND_BOX[0] + (*d_sps).ENERGY_SCOPE)
		//	(*d_sps).BOUND_BOX[0] = x - (*d_sps).ENERGY_SCOPE;
		//else if (x >(*d_sps).BOUND_BOX[2] - (*d_sps).ENERGY_SCOPE)
		//	(*d_sps).BOUND_BOX[2] = x + (*d_sps).ENERGY_SCOPE;
		//if (y < (*d_sps).BOUND_BOX[3] + (*d_sps).ENERGY_SCOPE)
		//	(*d_sps).BOUND_BOX[3] = y - (*d_sps).ENERGY_SCOPE;
		//else if (y >(*d_sps).BOUND_BOX[1] - (*d_sps).ENERGY_SCOPE)
		//	(*d_sps).BOUND_BOX[1] = y + (*d_sps).ENERGY_SCOPE;

		return 1;
	}

	(*d_sps).remain_time[index] = 0;
	return 0;
}
	.file	"tmpxft_002a631b_00000000-6_simple_particle.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2041:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2041:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Memory Allocation Error: %s\n"
	.text
	.globl	_Z19init_particles_cudaR20simpleParticleSystem
	.type	_Z19init_particles_cudaR20simpleParticleSystem, @function
_Z19init_particles_cudaR20simpleParticleSystem:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	addq	$16, %rdi
	pushq	%rsi
	.cfi_def_cfa_offset 48
	movslq	24(%rdi), %r12
	movl	28(%rdi), %ebp
	leaq	0(,%r12,4), %r13
	salq	$3, %r12
	movq	%r13, %rsi
	call	cudaMalloc@PLT
	movq	%r12, %rsi
	movq	%rbx, %rdi
	call	cudaMalloc@PLT
	leaq	8(%rbx), %rdi
	movq	%r12, %rsi
	call	cudaMalloc@PLT
	leaq	24(%rbx), %rdi
	movq	%r13, %rsi
	call	cudaMalloc@PLT
	movslq	%ebp, %rsi
	leaq	32(%rbx), %rdi
	imulq	$12, %rsi, %rsi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	je	.L2
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	popq	%rcx
	.cfi_remember_state
	.cfi_def_cfa_offset 40
	movl	$2, %edi
	popq	%rbx
	.cfi_def_cfa_offset 32
	movq	%rax, %rdx
	popq	%rbp
	.cfi_def_cfa_offset 24
	leaq	.LC0(%rip), %rsi
	popq	%r12
	.cfi_def_cfa_offset 16
	xorl	%eax, %eax
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	__printf_chk@PLT
.L2:
	.cfi_restore_state
	popq	%rax
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z19init_particles_cudaR20simpleParticleSystem, .-_Z19init_particles_cudaR20simpleParticleSystem
	.section	.rodata.str1.1
.LC1:
	.string	"Memory Free Error: %s\n"
	.text
	.globl	_Z22destroy_particles_cudaR20simpleParticleSystem
	.type	_Z22destroy_particles_cudaR20simpleParticleSystem, @function
_Z22destroy_particles_cudaR20simpleParticleSystem:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	movq	16(%rdi), %rdi
	call	cudaFree@PLT
	movq	(%rbx), %rdi
	call	cudaFree@PLT
	movq	8(%rbx), %rdi
	call	cudaFree@PLT
	movq	24(%rbx), %rdi
	call	cudaFree@PLT
	movq	32(%rbx), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	je	.L6
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	popq	%rbx
	.cfi_remember_state
	.cfi_def_cfa_offset 8
	movq	%rax, %rdx
	xorl	%eax, %eax
	jmp	__printf_chk@PLT
.L6:
	.cfi_restore_state
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z22destroy_particles_cudaR20simpleParticleSystem, .-_Z22destroy_particles_cudaR20simpleParticleSystem
	.section	.rodata.str1.1
.LC2:
	.string	"Constant Memory Copy Error: %s\n"
	.text
	.globl	_Z18copy_to_device_spsR20simpleParticleSystem
	.type	_Z18copy_to_device_spsR20simpleParticleSystem, @function
_Z18copy_to_device_spsR20simpleParticleSystem:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%rsi
	.cfi_def_cfa_offset 16
	xorl	%ecx, %ecx
	movq	%rdi, %rsi
	movl	$1, %r8d
	movl	$136, %edx
	leaq	_ZL5d_sps(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	testl	%eax, %eax
	je	.L9
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	popq	%rcx
	.cfi_remember_state
	.cfi_def_cfa_offset 8
	movq	%rax, %rdx
	xorl	%eax, %eax
	jmp	__printf_chk@PLT
.L9:
	.cfi_restore_state
	popq	%rax
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z18copy_to_device_spsR20simpleParticleSystem, .-_Z18copy_to_device_spsR20simpleParticleSystem
	.globl	_Z17get_normal_vectorf
	.type	_Z17get_normal_vectorf, @function
_Z17get_normal_vectorf:
.LFB2033:
	.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
.LFE2033:
	.size	_Z17get_normal_vectorf, .-_Z17get_normal_vectorf
	.globl	_Z10get_energy6float2S_f
	.type	_Z10get_energy6float2S_f, @function
_Z10get_energy6float2S_f:
.LFB2034:
	.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
.LFE2034:
	.size	_Z10get_energy6float2S_f, .-_Z10get_energy6float2S_f
	.globl	_Z21get_color_from_energyf
	.type	_Z21get_color_from_energyf, @function
_Z21get_color_from_energyf:
.LFB2035:
	.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
.LFE2035:
	.size	_Z21get_color_from_energyf, .-_Z21get_color_from_energyf
	.globl	_Z16get_accelerationi
	.type	_Z16get_accelerationi, @function
_Z16get_accelerationi:
.LFB2036:
	.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
.LFE2036:
	.size	_Z16get_accelerationi, .-_Z16get_accelerationi
	.globl	_Z24update_particle_velocityi6float2f
	.type	_Z24update_particle_velocityi6float2f, @function
_Z24update_particle_velocityi6float2f:
.LFB2037:
	.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
.LFE2037:
	.size	_Z24update_particle_velocityi6float2f, .-_Z24update_particle_velocityi6float2f
	.globl	_Z24update_particle_positionif
	.type	_Z24update_particle_positionif, @function
_Z24update_particle_positionif:
.LFB2038:
	.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
.LFE2038:
	.size	_Z24update_particle_positionif, .-_Z24update_particle_positionif
	.globl	_Z36__device_stub__Z17generateParticlesvv
	.type	_Z36__device_stub__Z17generateParticlesvv, @function
_Z36__device_stub__Z17generateParticlesvv:
.LFB2063:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	movabsq	$4294967297, %rax
	leaq	8(%rsp), %rcx
	movq	%rsp, %rdx
	leaq	28(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, 24(%rsp)
	movl	$1, 36(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 28(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L24
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z17generateParticlesv(%rip), %rdi
	pushq	8(%rsp)
	.cfi_def_cfa_offset 112
	movq	44(%rsp), %rcx
	movl	52(%rsp), %r8d
	movq	32(%rsp), %rsi
	movl	40(%rsp), %edx
	leaq	80(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L24:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L26
	call	__stack_chk_fail@PLT
.L26:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2063:
	.size	_Z36__device_stub__Z17generateParticlesvv, .-_Z36__device_stub__Z17generateParticlesvv
	.globl	_Z17generateParticlesv
	.type	_Z17generateParticlesv, @function
_Z17generateParticlesv:
.LFB2064:
	.cfi_startproc
	endbr64
	jmp	_Z36__device_stub__Z17generateParticlesvv
	.cfi_endproc
.LFE2064:
	.size	_Z17generateParticlesv, .-_Z17generateParticlesv
	.globl	_Z18generate_particlesi
	.type	_Z18generate_particlesi, @function
_Z18generate_particlesi:
.LFB2030:
	.cfi_startproc
	endbr64
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %esi
	movl	%edi, 20(%rsp)
	movabsq	$4294967297, %rdi
	movq	%rdi, 24(%rsp)
	movl	28(%rsp), %ecx
	movq	20(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L29
	addq	$40, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 8
	jmp	_Z36__device_stub__Z17generateParticlesvv
.L29:
	.cfi_restore_state
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2030:
	.size	_Z18generate_particlesi, .-_Z18generate_particlesi
	.globl	_Z43__device_stub__Z15renderParticlesP6uchar4iiP6uchar4ii
	.type	_Z43__device_stub__Z15renderParticlesP6uchar4iiP6uchar4ii, @function
_Z43__device_stub__Z15renderParticlesP6uchar4iiP6uchar4ii:
.LFB2065:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	32(%rsp), %rdi
	movl	%esi, 4(%rsp)
	leaq	44(%rsp), %rsi
	movl	%edx, (%rsp)
	leaq	16(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L32
	pushq	24(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z15renderParticlesP6uchar4ii(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 144
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L32:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L34
	call	__stack_chk_fail@PLT
.L34:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2065:
	.size	_Z43__device_stub__Z15renderParticlesP6uchar4iiP6uchar4ii, .-_Z43__device_stub__Z15renderParticlesP6uchar4iiP6uchar4ii
	.globl	_Z15renderParticlesP6uchar4ii
	.type	_Z15renderParticlesP6uchar4ii, @function
_Z15renderParticlesP6uchar4ii:
.LFB2066:
	.cfi_startproc
	endbr64
	jmp	_Z43__device_stub__Z15renderParticlesP6uchar4iiP6uchar4ii
	.cfi_endproc
.LFE2066:
	.size	_Z15renderParticlesP6uchar4ii, .-_Z15renderParticlesP6uchar4ii
	.globl	_Z16render_particlesP6uchar4ii
	.type	_Z16render_particlesP6uchar4ii, @function
_Z16render_particlesP6uchar4ii:
.LFB2032:
	.cfi_startproc
	endbr64
	leal	15(%rsi), %eax
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	$16, %ecx
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movl	%edx, %ebp
	cltd
	xorl	%r9d, %r9d
	idivl	%ecx
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movl	%esi, %ebx
	xorl	%r8d, %r8d
	movl	$1, %esi
	movl	%eax, 8(%rsp)
	leal	15(%rbp), %eax
	cltd
	idivl	%ecx
	movabsq	$68719476752, %rdx
	movl	$1, %ecx
	movl	%eax, 12(%rsp)
	movq	8(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L37
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 32
	movl	%ebp, %edx
	movl	%ebx, %esi
	movq	%r12, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	_Z43__device_stub__Z15renderParticlesP6uchar4iiP6uchar4ii
.L37:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2032:
	.size	_Z16render_particlesP6uchar4ii, .-_Z16render_particlesP6uchar4ii
	.globl	_Z34__device_stub__Z15updateParticlesff
	.type	_Z34__device_stub__Z15updateParticlesff, @function
_Z34__device_stub__Z15updateParticlesff:
.LFB2067:
	.cfi_startproc
	endbr64
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movss	%xmm0, 12(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rax
	leaq	44(%rsp), %rsi
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	leaq	32(%rsp), %rdi
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L40
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z15updateParticlesf(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L40:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L42
	call	__stack_chk_fail@PLT
.L42:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2067:
	.size	_Z34__device_stub__Z15updateParticlesff, .-_Z34__device_stub__Z15updateParticlesff
	.globl	_Z15updateParticlesf
	.type	_Z15updateParticlesf, @function
_Z15updateParticlesf:
.LFB2068:
	.cfi_startproc
	endbr64
	jmp	_Z34__device_stub__Z15updateParticlesff
	.cfi_endproc
.LFE2068:
	.size	_Z15updateParticlesf, .-_Z15updateParticlesf
	.globl	_Z16updata_particlesif
	.type	_Z16updata_particlesif, @function
_Z16updata_particlesif:
.LFB2031:
	.cfi_startproc
	endbr64
	subq	$56, %rsp
	.cfi_def_cfa_offset 64
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %esi
	movl	%edi, 36(%rsp)
	movabsq	$4294967297, %rdi
	movq	%rdi, 40(%rsp)
	movl	44(%rsp), %ecx
	movq	36(%rsp), %rdx
	movss	%xmm0, 12(%rsp)
	call	__cudaPushCallConfiguration@PLT
	movss	12(%rsp), %xmm0
	testl	%eax, %eax
	jne	.L45
	addq	$56, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 8
	jmp	_Z34__device_stub__Z15updateParticlesff
.L45:
	.cfi_restore_state
	addq	$56, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2031:
	.size	_Z16updata_particlesif, .-_Z16updata_particlesif
	.section	.rodata.str1.1
.LC3:
	.string	"_Z15updateParticlesf"
.LC4:
	.string	"_Z15renderParticlesP6uchar4ii"
.LC5:
	.string	"_Z17generateParticlesv"
.LC6:
	.string	"d_sps"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2070:
	.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	.LC3(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z15updateParticlesf(%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
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC4(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z15renderParticlesP6uchar4ii(%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
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC5(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z17generateParticlesv(%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
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC6(%rip), %rdx
	movl	$136, %r9d
	leaq	_ZL5d_sps(%rip), %rsi
	pushq	$1
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.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
.LFE2070:
	.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	_ZL5d_sps
	.comm	_ZL5d_sps,136,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"

#define PARTICLE_LIFE_MS 3000
#define PI 3.1415926

enum ParticleGeneratorType
{
	LineGenerator,
	CircleGenerator
};

typedef struct
{
	////physical attribute////
	float2 *position;
	float2 *velocity;
	float *energy;

	////particle life////
	float *remain_time;

	////rand data used for generating particles////
	float *rand_data;

	////common attributes////
	int MAX_PARTICLE_SIZE;
	int ONE_BATCH_PARTICLE_SIZE;
	float ENERGY_SCOPE;
	int LIFE_BOUND[4]; //left, top, right, bottom
	int BOUND_BOX[4];
	float MAX_VELOCITY; //pixels per second
	float MIN_VELOCITY;
	float LIFE_TIME; //second

	/////particle generator type related////
	enum ParticleGeneratorType TYPE;
	//LineGenerator
	float2 generator_line[2];
	//CircleGenerator
	float2 generator_center;
	float2 generator_radius;

} simpleParticleSystem;

void init_particles_cuda(simpleParticleSystem &sps);
void destroy_particles_cuda(simpleParticleSystem &sps);
void generate_particles(int generate_size);
void copy_to_device_sps(simpleParticleSystem &sps);
void render_particles(uchar4* devPtr, int img_width, int img_height);
void updata_particles(int generate_size, float passed_time);
	.text
	.file	"simple_particle.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
