#include "my_device_function.cuh"



__global__ void sigmoid(float *a,float *z,long n)
{
  long tid = blockIdx.x*blockDim.x + threadIdx.x;
  while(tid < n)
  {
      a[tid] = 1/(1+expf(-z[tid]));
      
      tid+= blockDim.x * gridDim.x;
  }
}

__global__ void sigmoid_inv(float *a,float *z,long n)
{
  long tid = blockIdx.x*blockDim.x + threadIdx.x;
  while(tid < n)
  {
      a[tid] = (1/(1+expf(-z[tid])))*(1 - 1/(1+expf(-z[tid])));
      tid+= blockDim.x * gridDim.x;
  }
}

__global__ void relu(float *a,float *z,long n)
{
    long tid = blockIdx.x*blockDim.x + threadIdx.x;
    while(tid < n)
    {
        if(z[tid] > 0 ) a[tid] = z[tid];
        else a[tid] = 0.0;      

        tid+= blockDim.x * gridDim.x;
    }
}

__global__ void relu_inv(float *a,float *z,long n)
{
    long tid = blockIdx.x*blockDim.x + threadIdx.x;
    while(tid < n)
    { 
        if(z[tid] > 0 ) a[tid] = 1.0;
        else a[tid] = 0.0;      
        tid+= blockDim.x * gridDim.x;
    }
}


__global__ void new_activation(float *a,float *z,long n)
{
 long tid = blockIdx.x*blockDim.x + threadIdx.x;
    while(tid < n)
    { 
//        if(z[tid] > 0 ) a[tid] = 1.0;
//        else a[tid] = 0.0;      
        tid+= blockDim.x * gridDim.x;
    }

}
__global__ void new_activation_inv(float *a,float *z,long n)
{
 long tid = blockIdx.x*blockDim.x + threadIdx.x;
    while(tid < n)
    { 
//        if(z[tid] > 0 ) a[tid] = 1.0;
//        else a[tid] = 0.0;      
        tid+= blockDim.x * gridDim.x;
    }

}











__global__ void deliver_front_to_rear(float *front,float *rear,long n)
{
  long tid = blockIdx.x*blockDim.x + threadIdx.x;
 
  while(tid < n)
  {
      rear[tid] = front[tid];  
      tid+= blockDim.x * gridDim.x;
  }
}


__global__ void add_bias(float *z,float *b,long column,long n)
{
    long tid = blockIdx.x*blockDim.x + threadIdx.x;
    while(tid < n)
    {
        z[tid] += b[tid % column];  

        tid+= blockDim.x * gridDim.x;
    }
}
__global__ void last_delta_before_transpose(float *temp, float *y,float *T,long batch_size,long n)
{
  long tid = blockIdx.x*blockDim.x + threadIdx.x;
  while(tid < n)
  {
      temp[tid] = (y[tid]-T[tid])/(2*batch_size);   
      tid+= blockDim.x * gridDim.x;
  }
}
__global__ void transpose(float *after, float *before,long before_columns,long before_rows)
{
  long tid = blockIdx.x*blockDim.x + threadIdx.x;
  long x,y;
  
  while(tid < before_columns*before_rows)
  {
      y = tid % before_columns;
      x = tid / before_columns;
      after[IDX2C(x,y,before_rows)] = before[IDX2C(y,x,before_columns)];
      tid+= blockDim.x * gridDim.x;
  }
}
__global__ void basic_multi(float *a,float *b,float *c, long n)
{
  long tid = blockIdx.x*blockDim.x + threadIdx.x;
  while(tid < n)
  {
      c[tid] = a[tid]*b[tid]; 
      tid+= blockDim.x * gridDim.x;
  }
}


__global__ void loss_cross_entropy(float *target,float *y, float * result,long last_neural,long batch_size)
{
  long tid = blockIdx.x*blockDim.x + threadIdx.x;
 
  while(tid < last_neural*batch_size)
  {
      result[tid] = -0.5*(target[tid]*logf(y[tid] + 0.000000001) + (1.0 - target[tid])*logf(1-y[tid] + 0.000000001));      
      tid+= blockDim.x * gridDim.x;
  }
}


__global__ void matching(float *target,float *y, float * result,long last_neural,long batch_size)
{
    long tid = blockIdx.x*blockDim.x + threadIdx.x;

    int target_inx;
    int y_inx;

    while(tid < batch_size)
    {
        float max = 0.0;
        for(int i = 0 ; i < last_neural ; i++)
        {
            if(target[IDX2C(i,tid,last_neural)] > 0.9)
            {
                target_inx = i;
            }
            if(y[IDX2C(i,tid,last_neural)] > max)
            {
                max = y[IDX2C(i,tid,last_neural)];
                y_inx = i;
            }
        }
        if(target_inx == y_inx) result[tid] = 1.0;
        else result[tid] = 0.0;

        tid+= blockDim.x * gridDim.x;
    }
}



__global__ void weight_update(float *w,float *delta_w, float alpha,float ramda,long n)
{
  long tid = blockIdx.x*blockDim.x + threadIdx.x;
 
  while(tid < n)
  {
      w[tid] = w[tid] - alpha*(delta_w[tid] + ramda*w[tid]);    
      tid+= blockDim.x * gridDim.x;
  }
}
__global__ void bias_update(float *b,float *delta_b, float alpha,long n)
{
  long tid = blockIdx.x*blockDim.x + threadIdx.x;
 
  while(tid < n)
  {
      b[tid] = b[tid] - alpha*delta_b[tid];    
      tid+= blockDim.x * gridDim.x;
  }
}








	.file	"tmpxft_002c897e_00000000-6_my_device_function.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2010:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2010:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0, @function
_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0:
.LFB2179:
	.cfi_startproc
	jmp	cudaLaunchKernel@PLT
	.cfi_endproc
.LFE2179:
	.size	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0, .-_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	.globl	_Z29__device_stub__Z7sigmoidPfS_lPfS_l
	.type	_Z29__device_stub__Z7sigmoidPfS_lPfS_l, @function
_Z29__device_stub__Z7sigmoidPfS_lPfS_l:
.LFB2032:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L3
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z7sigmoidPfS_l(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L3:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L5
	call	__stack_chk_fail@PLT
.L5:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2032:
	.size	_Z29__device_stub__Z7sigmoidPfS_lPfS_l, .-_Z29__device_stub__Z7sigmoidPfS_lPfS_l
	.globl	_Z7sigmoidPfS_l
	.type	_Z7sigmoidPfS_l, @function
_Z7sigmoidPfS_l:
.LFB2033:
	.cfi_startproc
	endbr64
	jmp	_Z29__device_stub__Z7sigmoidPfS_lPfS_l
	.cfi_endproc
.LFE2033:
	.size	_Z7sigmoidPfS_l, .-_Z7sigmoidPfS_l
	.globl	_Z34__device_stub__Z11sigmoid_invPfS_lPfS_l
	.type	_Z34__device_stub__Z11sigmoid_invPfS_lPfS_l, @function
_Z34__device_stub__Z11sigmoid_invPfS_lPfS_l:
.LFB2034:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z11sigmoid_invPfS_l(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L9:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L11
	call	__stack_chk_fail@PLT
.L11:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2034:
	.size	_Z34__device_stub__Z11sigmoid_invPfS_lPfS_l, .-_Z34__device_stub__Z11sigmoid_invPfS_lPfS_l
	.globl	_Z11sigmoid_invPfS_l
	.type	_Z11sigmoid_invPfS_l, @function
_Z11sigmoid_invPfS_l:
.LFB2035:
	.cfi_startproc
	endbr64
	jmp	_Z34__device_stub__Z11sigmoid_invPfS_lPfS_l
	.cfi_endproc
.LFE2035:
	.size	_Z11sigmoid_invPfS_l, .-_Z11sigmoid_invPfS_l
	.globl	_Z26__device_stub__Z4reluPfS_lPfS_l
	.type	_Z26__device_stub__Z4reluPfS_lPfS_l, @function
_Z26__device_stub__Z4reluPfS_lPfS_l:
.LFB2036:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L14
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z4reluPfS_l(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L14:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L16
	call	__stack_chk_fail@PLT
.L16:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2036:
	.size	_Z26__device_stub__Z4reluPfS_lPfS_l, .-_Z26__device_stub__Z4reluPfS_lPfS_l
	.globl	_Z4reluPfS_l
	.type	_Z4reluPfS_l, @function
_Z4reluPfS_l:
.LFB2037:
	.cfi_startproc
	endbr64
	jmp	_Z26__device_stub__Z4reluPfS_lPfS_l
	.cfi_endproc
.LFE2037:
	.size	_Z4reluPfS_l, .-_Z4reluPfS_l
	.globl	_Z30__device_stub__Z8relu_invPfS_lPfS_l
	.type	_Z30__device_stub__Z8relu_invPfS_lPfS_l, @function
_Z30__device_stub__Z8relu_invPfS_lPfS_l:
.LFB2038:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z8relu_invPfS_l(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L19:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2038:
	.size	_Z30__device_stub__Z8relu_invPfS_lPfS_l, .-_Z30__device_stub__Z8relu_invPfS_lPfS_l
	.globl	_Z8relu_invPfS_l
	.type	_Z8relu_invPfS_l, @function
_Z8relu_invPfS_l:
.LFB2039:
	.cfi_startproc
	endbr64
	jmp	_Z30__device_stub__Z8relu_invPfS_lPfS_l
	.cfi_endproc
.LFE2039:
	.size	_Z8relu_invPfS_l, .-_Z8relu_invPfS_l
	.globl	_Z37__device_stub__Z14new_activationPfS_lPfS_l
	.type	_Z37__device_stub__Z14new_activationPfS_lPfS_l, @function
_Z37__device_stub__Z14new_activationPfS_lPfS_l:
.LFB2040:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L24
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z14new_activationPfS_l(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L24:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L26
	call	__stack_chk_fail@PLT
.L26:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2040:
	.size	_Z37__device_stub__Z14new_activationPfS_lPfS_l, .-_Z37__device_stub__Z14new_activationPfS_lPfS_l
	.globl	_Z14new_activationPfS_l
	.type	_Z14new_activationPfS_l, @function
_Z14new_activationPfS_l:
.LFB2041:
	.cfi_startproc
	endbr64
	jmp	_Z37__device_stub__Z14new_activationPfS_lPfS_l
	.cfi_endproc
.LFE2041:
	.size	_Z14new_activationPfS_l, .-_Z14new_activationPfS_l
	.globl	_Z41__device_stub__Z18new_activation_invPfS_lPfS_l
	.type	_Z41__device_stub__Z18new_activation_invPfS_lPfS_l, @function
_Z41__device_stub__Z18new_activation_invPfS_lPfS_l:
.LFB2042:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L29
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z18new_activation_invPfS_l(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L29:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L31
	call	__stack_chk_fail@PLT
.L31:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2042:
	.size	_Z41__device_stub__Z18new_activation_invPfS_lPfS_l, .-_Z41__device_stub__Z18new_activation_invPfS_lPfS_l
	.globl	_Z18new_activation_invPfS_l
	.type	_Z18new_activation_invPfS_l, @function
_Z18new_activation_invPfS_l:
.LFB2043:
	.cfi_startproc
	endbr64
	jmp	_Z41__device_stub__Z18new_activation_invPfS_lPfS_l
	.cfi_endproc
.LFE2043:
	.size	_Z18new_activation_invPfS_l, .-_Z18new_activation_invPfS_l
	.globl	_Z44__device_stub__Z21deliver_front_to_rearPfS_lPfS_l
	.type	_Z44__device_stub__Z21deliver_front_to_rearPfS_lPfS_l, @function
_Z44__device_stub__Z21deliver_front_to_rearPfS_lPfS_l:
.LFB2044:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L34
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z21deliver_front_to_rearPfS_l(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L34:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L36
	call	__stack_chk_fail@PLT
.L36:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2044:
	.size	_Z44__device_stub__Z21deliver_front_to_rearPfS_lPfS_l, .-_Z44__device_stub__Z21deliver_front_to_rearPfS_lPfS_l
	.globl	_Z21deliver_front_to_rearPfS_l
	.type	_Z21deliver_front_to_rearPfS_l, @function
_Z21deliver_front_to_rearPfS_l:
.LFB2045:
	.cfi_startproc
	endbr64
	jmp	_Z44__device_stub__Z21deliver_front_to_rearPfS_lPfS_l
	.cfi_endproc
.LFE2045:
	.size	_Z21deliver_front_to_rearPfS_l, .-_Z21deliver_front_to_rearPfS_l
	.globl	_Z31__device_stub__Z8add_biasPfS_llPfS_ll
	.type	_Z31__device_stub__Z8add_biasPfS_llPfS_ll, @function
_Z31__device_stub__Z8add_biasPfS_llPfS_ll:
.LFB2046:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L39
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z8add_biasPfS_ll(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L39:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L41
	call	__stack_chk_fail@PLT
.L41:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2046:
	.size	_Z31__device_stub__Z8add_biasPfS_llPfS_ll, .-_Z31__device_stub__Z8add_biasPfS_llPfS_ll
	.globl	_Z8add_biasPfS_ll
	.type	_Z8add_biasPfS_ll, @function
_Z8add_biasPfS_ll:
.LFB2047:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z8add_biasPfS_llPfS_ll
	.cfi_endproc
.LFE2047:
	.size	_Z8add_biasPfS_ll, .-_Z8add_biasPfS_ll
	.globl	_Z53__device_stub__Z27last_delta_before_transposePfS_S_llPfS_S_ll
	.type	_Z53__device_stub__Z27last_delta_before_transposePfS_S_llPfS_S_ll, @function
_Z53__device_stub__Z27last_delta_before_transposePfS_S_llPfS_S_ll:
.LFB2048:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 96(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L44
	pushq	64(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z27last_delta_before_transposePfS_S_ll(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 176
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	112(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L44:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L46
	call	__stack_chk_fail@PLT
.L46:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2048:
	.size	_Z53__device_stub__Z27last_delta_before_transposePfS_S_llPfS_S_ll, .-_Z53__device_stub__Z27last_delta_before_transposePfS_S_llPfS_S_ll
	.globl	_Z27last_delta_before_transposePfS_S_ll
	.type	_Z27last_delta_before_transposePfS_S_ll, @function
_Z27last_delta_before_transposePfS_S_ll:
.LFB2049:
	.cfi_startproc
	endbr64
	jmp	_Z53__device_stub__Z27last_delta_before_transposePfS_S_llPfS_S_ll
	.cfi_endproc
.LFE2049:
	.size	_Z27last_delta_before_transposePfS_S_ll, .-_Z27last_delta_before_transposePfS_S_ll
	.globl	_Z32__device_stub__Z9transposePfS_llPfS_ll
	.type	_Z32__device_stub__Z9transposePfS_llPfS_ll, @function
_Z32__device_stub__Z9transposePfS_llPfS_ll:
.LFB2050:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L49
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z9transposePfS_ll(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L49:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L51
	call	__stack_chk_fail@PLT
.L51:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2050:
	.size	_Z32__device_stub__Z9transposePfS_llPfS_ll, .-_Z32__device_stub__Z9transposePfS_llPfS_ll
	.globl	_Z9transposePfS_ll
	.type	_Z9transposePfS_ll, @function
_Z9transposePfS_ll:
.LFB2051:
	.cfi_startproc
	endbr64
	jmp	_Z32__device_stub__Z9transposePfS_llPfS_ll
	.cfi_endproc
.LFE2051:
	.size	_Z9transposePfS_ll, .-_Z9transposePfS_ll
	.globl	_Z36__device_stub__Z11basic_multiPfS_S_lPfS_S_l
	.type	_Z36__device_stub__Z11basic_multiPfS_S_lPfS_S_l, @function
_Z36__device_stub__Z11basic_multiPfS_S_lPfS_S_l:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L54
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z11basic_multiPfS_S_l(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L54:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L56
	call	__stack_chk_fail@PLT
.L56:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z36__device_stub__Z11basic_multiPfS_S_lPfS_S_l, .-_Z36__device_stub__Z11basic_multiPfS_S_lPfS_S_l
	.globl	_Z11basic_multiPfS_S_l
	.type	_Z11basic_multiPfS_S_l, @function
_Z11basic_multiPfS_S_l:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z36__device_stub__Z11basic_multiPfS_S_lPfS_S_l
	.cfi_endproc
.LFE2053:
	.size	_Z11basic_multiPfS_S_l, .-_Z11basic_multiPfS_S_l
	.globl	_Z44__device_stub__Z18loss_cross_entropyPfS_S_llPfS_S_ll
	.type	_Z44__device_stub__Z18loss_cross_entropyPfS_S_llPfS_S_ll, @function
_Z44__device_stub__Z18loss_cross_entropyPfS_S_llPfS_S_ll:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 96(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L59
	pushq	64(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z18loss_cross_entropyPfS_S_ll(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 176
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	112(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L59:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L61
	call	__stack_chk_fail@PLT
.L61:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z44__device_stub__Z18loss_cross_entropyPfS_S_llPfS_S_ll, .-_Z44__device_stub__Z18loss_cross_entropyPfS_S_llPfS_S_ll
	.globl	_Z18loss_cross_entropyPfS_S_ll
	.type	_Z18loss_cross_entropyPfS_S_ll, @function
_Z18loss_cross_entropyPfS_S_ll:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z44__device_stub__Z18loss_cross_entropyPfS_S_llPfS_S_ll
	.cfi_endproc
.LFE2055:
	.size	_Z18loss_cross_entropyPfS_S_ll, .-_Z18loss_cross_entropyPfS_S_ll
	.globl	_Z33__device_stub__Z8matchingPfS_S_llPfS_S_ll
	.type	_Z33__device_stub__Z8matchingPfS_S_llPfS_S_ll, @function
_Z33__device_stub__Z8matchingPfS_S_llPfS_S_ll:
.LFB2056:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 96(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L64
	pushq	64(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z8matchingPfS_S_ll(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 176
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	112(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L64:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L66
	call	__stack_chk_fail@PLT
.L66:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2056:
	.size	_Z33__device_stub__Z8matchingPfS_S_llPfS_S_ll, .-_Z33__device_stub__Z8matchingPfS_S_llPfS_S_ll
	.globl	_Z8matchingPfS_S_ll
	.type	_Z8matchingPfS_S_ll, @function
_Z8matchingPfS_S_ll:
.LFB2057:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z8matchingPfS_S_llPfS_S_ll
	.cfi_endproc
.LFE2057:
	.size	_Z8matchingPfS_S_ll, .-_Z8matchingPfS_S_ll
	.globl	_Z38__device_stub__Z13weight_updatePfS_fflPfS_ffl
	.type	_Z38__device_stub__Z13weight_updatePfS_fflPfS_ffl, @function
_Z38__device_stub__Z13weight_updatePfS_fflPfS_ffl:
.LFB2058:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, (%rsp)
	leaq	40(%rsp), %rdx
	movss	%xmm0, 12(%rsp)
	movss	%xmm1, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 96(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 104(%rsp)
	movq	%rsp, %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L69
	pushq	48(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z13weight_updatePfS_ffl(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 160
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L69:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L71
	call	__stack_chk_fail@PLT
.L71:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2058:
	.size	_Z38__device_stub__Z13weight_updatePfS_fflPfS_ffl, .-_Z38__device_stub__Z13weight_updatePfS_fflPfS_ffl
	.globl	_Z13weight_updatePfS_ffl
	.type	_Z13weight_updatePfS_ffl, @function
_Z13weight_updatePfS_ffl:
.LFB2059:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z13weight_updatePfS_fflPfS_ffl
	.cfi_endproc
.LFE2059:
	.size	_Z13weight_updatePfS_ffl, .-_Z13weight_updatePfS_ffl
	.globl	_Z35__device_stub__Z11bias_updatePfS_flPfS_fl
	.type	_Z35__device_stub__Z11bias_updatePfS_flPfS_fl, @function
_Z35__device_stub__Z11bias_updatePfS_flPfS_fl:
.LFB2060:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, (%rsp)
	leaq	32(%rsp), %rdx
	movss	%xmm0, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L74
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z11bias_updatePfS_fl(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L74:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L76
	call	__stack_chk_fail@PLT
.L76:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2060:
	.size	_Z35__device_stub__Z11bias_updatePfS_flPfS_fl, .-_Z35__device_stub__Z11bias_updatePfS_flPfS_fl
	.globl	_Z11bias_updatePfS_fl
	.type	_Z11bias_updatePfS_fl, @function
_Z11bias_updatePfS_fl:
.LFB2061:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z11bias_updatePfS_flPfS_fl
	.cfi_endproc
.LFE2061:
	.size	_Z11bias_updatePfS_fl, .-_Z11bias_updatePfS_fl
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z11bias_updatePfS_fl"
.LC1:
	.string	"_Z13weight_updatePfS_ffl"
.LC2:
	.string	"_Z8matchingPfS_S_ll"
.LC3:
	.string	"_Z18loss_cross_entropyPfS_S_ll"
.LC4:
	.string	"_Z11basic_multiPfS_S_l"
.LC5:
	.string	"_Z9transposePfS_ll"
.LC6:
	.string	"_Z27last_delta_before_transposePfS_S_ll"
.LC7:
	.string	"_Z8add_biasPfS_ll"
.LC8:
	.string	"_Z21deliver_front_to_rearPfS_l"
.LC9:
	.string	"_Z18new_activation_invPfS_l"
.LC10:
	.string	"_Z14new_activationPfS_l"
.LC11:
	.string	"_Z8relu_invPfS_l"
.LC12:
	.string	"_Z4reluPfS_l"
.LC13:
	.string	"_Z11sigmoid_invPfS_l"
.LC14:
	.string	"_Z7sigmoidPfS_l"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2063:
	.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	_Z11bias_updatePfS_fl(%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	_Z13weight_updatePfS_ffl(%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	_Z8matchingPfS_S_ll(%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	_Z18loss_cross_entropyPfS_S_ll(%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	.LC4(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z11basic_multiPfS_S_l(%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	.LC5(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z9transposePfS_ll(%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	.LC6(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z27last_delta_before_transposePfS_S_ll(%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	.LC7(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z8add_biasPfS_ll(%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
	leaq	.LC8(%rip), %rdx
	pushq	$0
	.cfi_def_cfa_offset 24
	movq	%rdx, %rcx
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	_Z21deliver_front_to_rearPfS_l(%rip), %rsi
	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	.LC9(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z18new_activation_invPfS_l(%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	.LC10(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z14new_activationPfS_l(%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	.LC11(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z8relu_invPfS_l(%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	.LC12(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z4reluPfS_l(%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	.LC13(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z11sigmoid_invPfS_l(%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	.LC14(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z7sigmoidPfS_l(%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
.LFE2063:
	.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:

#ifndef __MY_DEVICE_FUNCTION_CU__
#define __MY_DEVICE_FUNCTION_CU__

#define IDX2C(i,j,Id)       (((j)*(Id))+(i)) // j -> row, i -> column



__global__ void sigmoid(float *a,float *z,long n);
__global__ void sigmoid_inv(float *a,float *z,long n);

__global__ void relu(float *a,float *z,long n);
__global__ void relu_inv(float *a,float *z,long n);

__global__ void new_activation(float *a,float *z,long n);
__global__ void new_activation_inv(float *a,float *z,long n);




// front -> rear 변수에 전달 
__global__ void deliver_front_to_rear(float *front,float *rear,long n);
// z = z + b
__global__ void add_bias(float *z,float *b,long column,long n);
// temp = (y-T)*(2*batch_size) 
__global__ void last_delta_before_transpose(float *temp, float *y,float *T,long batch_size,long n);
// before -> after transpose
__global__ void transpose(float *after, float *before,long before_columns,long before_rows);
// c = a .* b
__global__ void basic_multi(float *a,float *b,float *c, long n);
//temp = -0.5*(T*log(y) + (1-T)*log(1-y))
__global__ void loss_cross_entropy(float *target,float *y, float * result,long last_neural,long batch_size);
// result = matching(target, y) -> 같으면 1 아니면 0
__global__ void matching(float *target,float *y, float * result,long last_neural,long batch_size);
// w = w - alpha*(delta_w + ramda*w)
__global__ void weight_update(float *w,float *delta_w, float alpha,float ramda,long n);
// b = b - alpha * delta_b
__global__ void bias_update(float *b,float *delta_b, float alpha,long n);


#endif

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