#include "kernel.cuh"
#define TX 32
#define TY 32
#define RAD 1

#define divUp(a,b) ((a) + (b) - 1) / (b)

__device__
unsigned char clip(int n){ return n > 255 ? 255 : (n < 0 ? 0 : n);}

__device__
int idxClip(int idx,int idxMax){
    return idx > idxMax ? idxMax : (idx < 0 ? 0 : idx);
}

__device__
int flatten(int col,int row,int width,int height){
    return idxClip(col,width) + idxClip(row,height)*width;
}

__global__
void resetKernel(float *d_temp,int w,int h,BC bc){
    const int col = blockIdx.x*blockDim.x + threadIdx.x;
    const int row = blockIdx.y*blockDim.y + threadIdx.y;
    if((col >= w) || (row >= h)) return;
    d_temp[row*w + col] = bc.t_a;
}

__global__
void tempKernel(uchar4* d_out,float *d_temp,int w,int h,BC bc){
    extern __shared__ float s_in[];
    const int col = threadIdx.x + blockDim.x*blockIdx.x;
    const int row = threadIdx.y + blockDim.y*blockIdx.y;

    if((col >= w) || (row >= h)) return;
    const int idx = flatten(col,row,w,h);
    const int s_w = blockDim.x + 2*RAD;
    const int s_h = blockDim.y + 2*RAD;

    const int s_col = threadIdx.x + RAD;
    const int s_row = threadIdx.y + RAD;
    const int s_idx = flatten(s_col,s_row,s_w,s_h);

    d_out[idx].x = 0;
    d_out[idx].y = 0;
    d_out[idx].z = 0;
    d_out[idx].w = 255;

    s_in[s_idx] = d_temp[idx];

    if(threadIdx.x < RAD){
        s_in[flatten(s_col - RAD,s_row,s_w,s_h)] = d_temp[flatten(col - RAD,row,w,h)];
        s_in[flatten(s_col + blockDim.x,s_row,s_w,s_h)] = d_temp[flatten(col + blockDim.x,row,w,h)];
    }

    if(threadIdx.y < RAD){
        s_in[flatten(s_col,s_row - RAD,s_w,s_h)] = d_temp[flatten(col,row - RAD,w,h)];
        s_in[flatten(s_col,s_row + blockDim.y,s_w,s_h)] = d_temp[flatten(col,row + blockDim.y,w,h)];
    }

    float dSq = (col - bc.x)*(col - bc.x) + (row - bc.y)*(row - bc.y);
    if(dSq < bc.rad*bc.rad){
        d_temp[idx] = bc.t_s;
        return;
    }

    if((col == 0) || (col == w - 1) || (row == 0) || (col + row < bc.chamfer) || (col - row > w - bc.chamfer)){
        d_temp[idx] = bc.t_a;
        return;
    }

    if(row == h-1){
        d_temp[idx] = bc.t_g;
        return;
    }

    __syncthreads();

    float temp = 0.25f*(s_in[flatten(s_col - 1,s_row,s_w,s_h)] +
                        s_in[flatten(s_col + 1,s_row,s_w,s_h)] +
                        s_in[flatten(s_col,s_row - 1,s_w,s_h)] +
                        s_in[flatten(s_col,s_row + 1,s_w,s_h)]);

    d_temp[idx] = temp;
    const unsigned char intensity = clip((int)temp);
    d_out[idx].x = intensity;
    d_out[idx].z = 255 - intensity;
}

void kernelLauncher(uchar4 *d_out,float *d_temp,int w,int h,BC bc){
    const dim3 blockSize(TX,TY);
    const dim3 gridSize(divUp(w,TX),divUp(h,TY));
    const size_t smSz = (TX + 2*RAD)*(TY + 2*RAD)*sizeof(float);
    tempKernel<<<gridSize,blockSize,smSz>>>(d_out,d_temp,w,h,bc);
}

void resetTemperature(float *d_temp,int w,int h, BC bc){
    const dim3 blockSize(TX,TY);
    const dim3 gridSize(divUp(w,TX),divUp(h,TY));
    resetKernel<<<gridSize,blockSize>>>(d_temp,w,h,bc);
}

	.file	"tmpxft_001fdc14_00000000-6_kernel.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2015:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2015:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z4clipi
	.type	_Z4clipi, @function
_Z4clipi:
.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	_Z4clipi, .-_Z4clipi
	.globl	_Z7idxClipii
	.type	_Z7idxClipii, @function
_Z7idxClipii:
.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	_Z7idxClipii, .-_Z7idxClipii
	.globl	_Z7flatteniiii
	.type	_Z7flatteniiii, @function
_Z7flatteniiii:
.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	_Z7flatteniiii, .-_Z7flatteniiii
	.globl	_Z36__device_stub__Z11resetKernelPfii2BCPfiiR2BC
	.type	_Z36__device_stub__Z11resetKernelPfii2BCPfiiR2BC, @function
_Z36__device_stub__Z11resetKernelPfii2BCPfiiR2BC:
.LFB2037:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 8(%rsp)
	leaq	40(%rsp), %rdi
	movl	%esi, 4(%rsp)
	leaq	52(%rsp), %rsi
	movl	%edx, (%rsp)
	leaq	24(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movq	%rcx, 112(%rsp)
	leaq	32(%rsp), %rcx
	movq	%rax, 88(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 96(%rsp)
	movq	%rsp, %rax
	movq	%rax, 104(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movl	$1, 48(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	32(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z11resetKernelPfii2BC(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 160
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L8:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2037:
	.size	_Z36__device_stub__Z11resetKernelPfii2BCPfiiR2BC, .-_Z36__device_stub__Z11resetKernelPfii2BCPfiiR2BC
	.globl	_Z11resetKernelPfii2BC
	.type	_Z11resetKernelPfii2BC, @function
_Z11resetKernelPfii2BC:
.LFB2038:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	16(%rsp), %rcx
	call	_Z36__device_stub__Z11resetKernelPfii2BCPfiiR2BC
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2038:
	.size	_Z11resetKernelPfii2BC, .-_Z11resetKernelPfii2BC
	.globl	_Z16resetTemperaturePfii2BC
	.type	_Z16resetTemperaturePfii2BC, @function
_Z16resetTemperaturePfii2BC:
.LFB2012:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	$32, %ecx
	movq	%rdi, %r12
	xorl	%r9d, %r9d
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movl	%edx, %ebp
	xorl	%r8d, %r8d
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movl	%esi, %ebx
	subq	$64, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	leal	31(%rsi), %eax
	movl	$1, %esi
	cltd
	idivl	%ecx
	movl	%eax, 16(%rsp)
	leal	31(%rbp), %eax
	cltd
	idivl	%ecx
	movabsq	$137438953504, %rdx
	movl	$1, %ecx
	movl	%eax, 20(%rsp)
	movq	16(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L15
	movups	96(%rsp), %xmm0
	movups	108(%rsp), %xmm1
	movl	%ebp, %edx
	movl	%ebx, %esi
	leaq	28(%rsp), %rcx
	movq	%r12, %rdi
	movups	%xmm0, 28(%rsp)
	movups	%xmm1, 40(%rsp)
	call	_Z36__device_stub__Z11resetKernelPfii2BCPfiiR2BC
.L15:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L17
	call	__stack_chk_fail@PLT
.L17:
	addq	$64, %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
.LFE2012:
	.size	_Z16resetTemperaturePfii2BC, .-_Z16resetTemperaturePfii2BC
	.globl	_Z43__device_stub__Z10tempKernelP6uchar4Pfii2BCP6uchar4PfiiR2BC
	.type	_Z43__device_stub__Z10tempKernelP6uchar4Pfii2BCP6uchar4PfiiR2BC, @function
_Z43__device_stub__Z10tempKernelP6uchar4Pfii2BCP6uchar4PfiiR2BC:
.LFB2039:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movq	%r8, 128(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movl	$1, 56(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z10tempKernelP6uchar4Pfii2BC(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L19:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2039:
	.size	_Z43__device_stub__Z10tempKernelP6uchar4Pfii2BCP6uchar4PfiiR2BC, .-_Z43__device_stub__Z10tempKernelP6uchar4Pfii2BCP6uchar4PfiiR2BC
	.globl	_Z10tempKernelP6uchar4Pfii2BC
	.type	_Z10tempKernelP6uchar4Pfii2BC, @function
_Z10tempKernelP6uchar4Pfii2BC:
.LFB2040:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	16(%rsp), %r8
	call	_Z43__device_stub__Z10tempKernelP6uchar4Pfii2BCP6uchar4PfiiR2BC
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2040:
	.size	_Z10tempKernelP6uchar4Pfii2BC, .-_Z10tempKernelP6uchar4Pfii2BC
	.globl	_Z14kernelLauncherP6uchar4Pfii2BC
	.type	_Z14kernelLauncherP6uchar4Pfii2BC, @function
_Z14kernelLauncherP6uchar4Pfii2BC:
.LFB2011:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	xorl	%r9d, %r9d
	movq	%rsi, %r13
	movl	$4624, %r8d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rdi, %r12
	movl	$1, %esi
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movl	%ecx, %ebp
	movl	$32, %ecx
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	%edx, %ebx
	subq	$72, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	leal	31(%rdx), %eax
	cltd
	idivl	%ecx
	movl	%eax, 16(%rsp)
	leal	31(%rbp), %eax
	cltd
	idivl	%ecx
	movabsq	$137438953504, %rdx
	movl	$1, %ecx
	movl	%eax, 20(%rsp)
	movq	16(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L25
	movups	112(%rsp), %xmm0
	movups	124(%rsp), %xmm1
	movl	%ebp, %ecx
	movl	%ebx, %edx
	leaq	28(%rsp), %r8
	movq	%r13, %rsi
	movq	%r12, %rdi
	movups	%xmm0, 28(%rsp)
	movups	%xmm1, 40(%rsp)
	call	_Z43__device_stub__Z10tempKernelP6uchar4Pfii2BCP6uchar4PfiiR2BC
.L25:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L27
	call	__stack_chk_fail@PLT
.L27:
	addq	$72, %rsp
	.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
.LFE2011:
	.size	_Z14kernelLauncherP6uchar4Pfii2BC, .-_Z14kernelLauncherP6uchar4Pfii2BC
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z10tempKernelP6uchar4Pfii2BC"
.LC1:
	.string	"_Z11resetKernelPfii2BC"
	.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
	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	_Z10tempKernelP6uchar4Pfii2BC(%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	.LC1(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z11resetKernelPfii2BC(%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
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.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
	.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:

//
// Created by shuliang on 19-7-18.
//

#ifndef HARMONICCOORDINATEDEMO_KERNEL_CUH
#define HARMONICCOORDINATEDEMO_KERNEL_CUH

struct uchar4;
typedef struct{
    int x,y;
    float rad;
    int chamfer;
    float t_s,t_a,t_g;
}BC;

void kernelLauncher(uchar4 *d_out,float *d_temp,int w,int h,BC bc);

void resetTemperature(float *d_temp,int w,int h,BC bc);

#endif //HARMONICCOORDINATEDEMO_KERNEL_CUH

	.text
	.file	"kernel.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
