/*
number of mathematical operations (only floating point)
      operation  flo/o total
      +-* : 27      1    27
      /   :  2      4     8
      pow :  1     13    13
      sum                48
*/
#define M2L_KERNEL_CORE \
    for(i=0;i<3;i++) nc[i]=0;\
    nb=je-1;\
    k=0;\
    i=1;\
    while(nb!=0){\
      j=2-k;\
      nc[j]=nc[j]+nb%2*i;\
      nb=nb/2;\
      j=k+1;\
      k=j%3;\
      if(k==0) i=i*2;\
    }\
    nd=nc[0];\
    nc[0]=nc[1];\
    nc[1]=nc[2];\
    nc[2]=nd;\
    xijc=(nc[0]-3)*rb;\
    yijc=(nc[1]-3)*rb;\
    zijc=(nc[2]-3)*rb;\
    rh=sqrt(xijc*xijc+yijc*yijc+zijc*zijc)+eps;\
    jbase=(je-1)*mpdnm;\
    n=ng[tx];\
    m=mg[tx];\
    nms=n*(n+1)/2+m;\
    for(i=0;i<2;i++) vecd[i]=0;\
    for(k=-n;k<0;k++){\
      nks=n*(n+1)/2-k;\
      nmk=jbase+(4*n*n*n+6*n*n+5*n)/3+m*(2*n+1)+k;\
      dnmre=lvec[2*nmk+0];\
      dnmim=lvec[2*nmk+1];\
      vecd[0]+=dnmre*vecj[2*nks+0];\
      vecd[0]+=dnmim*vecj[2*nks+1];\
      vecd[1]-=dnmre*vecj[2*nks+1];\
      vecd[1]+=dnmim*vecj[2*nks+0];\
    }\
    for(k=0;k<=n;k++){\
      nks=n*(n+1)/2+k;\
      nmk=jbase+(4*n*n*n+6*n*n+5*n)/3+m*(2*n+1)+k;\
      dnmre=lvec[2*nmk+0];\
      dnmim=lvec[2*nmk+1];\
      vecd[0]+=dnmre*vecj[2*nks+0];\
      vecd[0]-=dnmim*vecj[2*nks+1];\
      vecd[1]+=dnmre*vecj[2*nks+1];\
      vecd[1]+=dnmim*vecj[2*nks+0];\
    }\
    __syncthreads();\
    for(i=0;i<2;i++) vecj[2*nms+i]=vecd[i];\
    __syncthreads();\
    j=ng[tx];\
    k=mg[tx];\
    jks=j*(j+1)/2+k;\
    for(i=0;i<2;i++) vecd[i]=0;\
    fnmm=1.0;\
    for(i=0;i<j-k;i++) fnmm=fnmm*(i+1);\
    fnpm=1.0;\
    for(i=0;i<j+k;i++) fnpm=fnpm*(i+1);\
    ajk=pow(-1.0,j)*rsqrtf(fnmm*fnpm);\
    for(n=abs(k);n<mp;n++){\
      nks=n*(n+1)/2+k;\
      jnk=(j+n)*(j+n)+j+n;\
      fnmm=1.0;\
      for(i=0;i<n-k;i++) fnmm=fnmm*(i+1);\
      fnpm=1.0;\
      for(i=0;i<n+k;i++) fnpm=fnpm*(i+1);\
      ank=pow(-1.0,n)*rsqrtf(fnmm*fnpm);\
      fnpm=1.0;\
      for(i=0;i<j+n;i++) fnpm=fnpm*(i+1);\
      ajn=pow(-1.0,j+n)/fnpm;\
      sr=pow(-1.0,j+k)*ank*ajk/ajn;\
      cnmre=sr*ynmre[jnk]/pow(rh,j+n+1);\
      cnmim=sr*ynmim[jnk]/pow(rh,j+n+1);\
      vecd[0]+=vecj[2*nks+0]*cnmre;\
      vecd[0]-=vecj[2*nks+1]*cnmim;\
      vecd[1]+=vecj[2*nks+0]*cnmim;\
      vecd[1]+=vecj[2*nks+1]*cnmre;\
    }\
    __syncthreads();\
    for(i=0;i<2;i++) vecj[2*jks+i]=vecd[i];\
    __syncthreads();\
    jbase=(je+nrbm-1)*mpdnm;\
    n=ng[tx];\
    m=mg[tx];\
    nms=n*(n+1)/2+m;\
    for(i=0;i<2;i++) vecd[i]=0;\
    for(k=-n;k<0;k++){\
      nks=n*(n+1)/2-k;\
      nmk=jbase+(4*n*n*n+6*n*n+5*n)/3+m*(2*n+1)+k;\
      dnmre=lvec[2*nmk+0];\
      dnmim=lvec[2*nmk+1];\
      vecd[0]+=dnmre*vecj[2*nks+0];\
      vecd[0]+=dnmim*vecj[2*nks+1];\
      vecd[1]-=dnmre*vecj[2*nks+1];\
      vecd[1]+=dnmim*vecj[2*nks+0];\
    }\
    for(k=0;k<=n;k++){\
      nks=n*(n+1)/2+k;\
      nmk=jbase+(4*n*n*n+6*n*n+5*n)/3+m*(2*n+1)+k;\
      dnmre=lvec[2*nmk+0];\
      dnmim=lvec[2*nmk+1];\
      vecd[0]+=dnmre*vecj[2*nks+0];\
      vecd[0]-=dnmim*vecj[2*nks+1];\
      vecd[1]+=dnmre*vecj[2*nks+1];\
      vecd[1]+=dnmim*vecj[2*nks+0];\
    }

	.file	"tmpxft_0038ba37_00000000-6_m2l.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
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2033:
	.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:


#include <hip/hip_runtime.h>
/*
number of mathematical operations (only floating point)
      operation  flo/o total
      +-* : 27      1    27
      /   :  2      4     8
      pow :  1     13    13
      sum                48
*/
#define M2L_KERNEL_CORE \
    for(i=0;i<3;i++) nc[i]=0;\
    nb=je-1;\
    k=0;\
    i=1;\
    while(nb!=0){\
      j=2-k;\
      nc[j]=nc[j]+nb%2*i;\
      nb=nb/2;\
      j=k+1;\
      k=j%3;\
      if(k==0) i=i*2;\
    }\
    nd=nc[0];\
    nc[0]=nc[1];\
    nc[1]=nc[2];\
    nc[2]=nd;\
    xijc=(nc[0]-3)*rb;\
    yijc=(nc[1]-3)*rb;\
    zijc=(nc[2]-3)*rb;\
    rh=sqrt(xijc*xijc+yijc*yijc+zijc*zijc)+eps;\
    jbase=(je-1)*mpdnm;\
    n=ng[tx];\
    m=mg[tx];\
    nms=n*(n+1)/2+m;\
    for(i=0;i<2;i++) vecd[i]=0;\
    for(k=-n;k<0;k++){\
      nks=n*(n+1)/2-k;\
      nmk=jbase+(4*n*n*n+6*n*n+5*n)/3+m*(2*n+1)+k;\
      dnmre=lvec[2*nmk+0];\
      dnmim=lvec[2*nmk+1];\
      vecd[0]+=dnmre*vecj[2*nks+0];\
      vecd[0]+=dnmim*vecj[2*nks+1];\
      vecd[1]-=dnmre*vecj[2*nks+1];\
      vecd[1]+=dnmim*vecj[2*nks+0];\
    }\
    for(k=0;k<=n;k++){\
      nks=n*(n+1)/2+k;\
      nmk=jbase+(4*n*n*n+6*n*n+5*n)/3+m*(2*n+1)+k;\
      dnmre=lvec[2*nmk+0];\
      dnmim=lvec[2*nmk+1];\
      vecd[0]+=dnmre*vecj[2*nks+0];\
      vecd[0]-=dnmim*vecj[2*nks+1];\
      vecd[1]+=dnmre*vecj[2*nks+1];\
      vecd[1]+=dnmim*vecj[2*nks+0];\
    }\
    __syncthreads();\
    for(i=0;i<2;i++) vecj[2*nms+i]=vecd[i];\
    __syncthreads();\
    j=ng[tx];\
    k=mg[tx];\
    jks=j*(j+1)/2+k;\
    for(i=0;i<2;i++) vecd[i]=0;\
    fnmm=1.0;\
    for(i=0;i<j-k;i++) fnmm=fnmm*(i+1);\
    fnpm=1.0;\
    for(i=0;i<j+k;i++) fnpm=fnpm*(i+1);\
    ajk=pow(-1.0,j)*rsqrtf(fnmm*fnpm);\
    for(n=abs(k);n<mp;n++){\
      nks=n*(n+1)/2+k;\
      jnk=(j+n)*(j+n)+j+n;\
      fnmm=1.0;\
      for(i=0;i<n-k;i++) fnmm=fnmm*(i+1);\
      fnpm=1.0;\
      for(i=0;i<n+k;i++) fnpm=fnpm*(i+1);\
      ank=pow(-1.0,n)*rsqrtf(fnmm*fnpm);\
      fnpm=1.0;\
      for(i=0;i<j+n;i++) fnpm=fnpm*(i+1);\
      ajn=pow(-1.0,j+n)/fnpm;\
      sr=pow(-1.0,j+k)*ank*ajk/ajn;\
      cnmre=sr*ynmre[jnk]/pow(rh,j+n+1);\
      cnmim=sr*ynmim[jnk]/pow(rh,j+n+1);\
      vecd[0]+=vecj[2*nks+0]*cnmre;\
      vecd[0]-=vecj[2*nks+1]*cnmim;\
      vecd[1]+=vecj[2*nks+0]*cnmim;\
      vecd[1]+=vecj[2*nks+1]*cnmre;\
    }\
    __syncthreads();\
    for(i=0;i<2;i++) vecj[2*jks+i]=vecd[i];\
    __syncthreads();\
    jbase=(je+nrbm-1)*mpdnm;\
    n=ng[tx];\
    m=mg[tx];\
    nms=n*(n+1)/2+m;\
    for(i=0;i<2;i++) vecd[i]=0;\
    for(k=-n;k<0;k++){\
      nks=n*(n+1)/2-k;\
      nmk=jbase+(4*n*n*n+6*n*n+5*n)/3+m*(2*n+1)+k;\
      dnmre=lvec[2*nmk+0];\
      dnmim=lvec[2*nmk+1];\
      vecd[0]+=dnmre*vecj[2*nks+0];\
      vecd[0]+=dnmim*vecj[2*nks+1];\
      vecd[1]-=dnmre*vecj[2*nks+1];\
      vecd[1]+=dnmim*vecj[2*nks+0];\
    }\
    for(k=0;k<=n;k++){\
      nks=n*(n+1)/2+k;\
      nmk=jbase+(4*n*n*n+6*n*n+5*n)/3+m*(2*n+1)+k;\
      dnmre=lvec[2*nmk+0];\
      dnmim=lvec[2*nmk+1];\
      vecd[0]+=dnmre*vecj[2*nks+0];\
      vecd[0]-=dnmim*vecj[2*nks+1];\
      vecd[1]+=dnmre*vecj[2*nks+1];\
      vecd[1]+=dnmim*vecj[2*nks+0];\
    }

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