//
// Created by rafa on 2/5/21.
//

#include "counter.cuh"

__global__
void
doPairCount(long long np, const double *array, const long long *particlesInGrid, const long long *offset,
            double boxsize, int nside, double minsep, double maxsep, const int nbins, const double *rbins2,
            unsigned long long *paircounts, const int *key, const int *iwrap) {
    long long index = blockIdx.x * blockDim.x + threadIdx.x;
    long long stride = blockDim.x * gridDim.x;

    double agrid = boxsize / nside;
    int index_max = (int) (maxsep / agrid) + 1;

    long long i, j;
    int i1, j1, k1;
    double x1, y1, z1;
    double x11, y11, z11;
    double dx, dy, dz;
    double r2;
    int iwrapx, iwrapy, iwrapz;
    int i2, j2, k2;
    int s;
    double factor = nside / boxsize;

//    __shared__  unsigned long long private_histogram[20];
//    if (threadIdx.x < 20) {
//        private_histogram[threadIdx.x] = 0;
//    }
//    __syncthreads();

    for (i = index; i < np; i += stride) {
        x1 = array[3 * i + 0];
        y1 = array[3 * i + 1];
        z1 = array[3 * i + 2];

        i1 = (int) (x1 * factor);
        j1 = (int) (y1 * factor);
        k1 = (int) (z1 * factor);

        i1 %= nside;
        j1 %= nside;
        k1 %= nside;

        for (int idx = -index_max; idx <= index_max; idx++) {
            i2 = i1 + idx;
            iwrapx = iwrap[i2 + index_max];
            i2 = key[i2 + index_max];

            for (int idy = -index_max; idy <= index_max; idy++) {
                j2 = j1 + idy;
                iwrapy = iwrap[j2 + index_max];
                j2 = key[j2 + index_max];

                for (int idz = -index_max; idz <= index_max; idz++) {
                    k2 = k1 + idz;
                    iwrapz = iwrap[k2 + index_max];
                    k2 = key[k2 + index_max];

                    s = nside * nside * k2 + nside * j2 + i2;

                    x11 = x1 - boxsize * iwrapx;
                    y11 = y1 - boxsize * iwrapy;
                    z11 = z1 - boxsize * iwrapz;

                    for (j = 0; j < particlesInGrid[s]; j += 1) {
                        dx = x11 - array[3 * offset[s] + 3 * j + 0];
                        dy = y11 - array[3 * offset[s] + 3 * j + 1];
                        dz = z11 - array[3 * offset[s] + 3 * j + 2];

                        r2 = dx * dx + dy * dy + dz * dz;

                        if (r2 > rbins2[nbins] || r2 < rbins2[0]) continue;

                        for (int k = nbins - 1; k >= 0; k--) {
                            if (r2 >= rbins2[k]) {
                                atomicAdd(&(paircounts[k]), 1);
//                                atomicAdd(&(private_histogram[k]), 1);
//                                paircounts[k] += 1;
                                break;
                            }
                        }
                    } //loop over neighbor grid2 particles
                } //loop over z neighbors
            } //loop over y neighbors
        } //loop over x neighbors
    } //loop over catalog1 particles

//    __syncthreads();
//
//    if (threadIdx.x < 20) {
//        atomicAdd(&(paircounts[threadIdx.x]), private_histogram[threadIdx.x]);
//    }

//    return paircounts;
}
	.file	"tmpxft_00369e99_00000000-6_counter.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
	.globl	_Z55__device_stub__Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_xPKdPKxS2_diddiS0_PyPKiS5_
	.type	_Z55__device_stub__Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_xPKdPKxS2_diddiS0_PyPKiS5_, @function
_Z55__device_stub__Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_xPKdPKxS2_diddiS0_PyPKiS5_:
.LFB2032:
	.cfi_startproc
	endbr64
	subq	$280, %rsp
	.cfi_def_cfa_offset 288
	movq	288(%rsp), %rax
	movq	%rdi, 88(%rsp)
	leaq	112(%rsp), %rdi
	movq	%rsi, 80(%rsp)
	leaq	124(%rsp), %rsi
	movq	%rax, 24(%rsp)
	movq	296(%rsp), %rax
	movq	%rdx, 72(%rsp)
	leaq	96(%rsp), %rdx
	movq	%rax, 16(%rsp)
	movq	304(%rsp), %rax
	movq	%rcx, 64(%rsp)
	leaq	104(%rsp), %rcx
	movq	%rax, 8(%rsp)
	movq	312(%rsp), %rax
	movl	%r8d, 52(%rsp)
	movl	%r9d, 48(%rsp)
	movq	%rax, (%rsp)
	movsd	%xmm0, 56(%rsp)
	movsd	%xmm1, 40(%rsp)
	movsd	%xmm2, 32(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 264(%rsp)
	xorl	%eax, %eax
	leaq	88(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	52(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 248(%rsp)
	movq	%rsp, %rax
	movq	%rax, 256(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 112(%rsp)
	movl	$1, 120(%rsp)
	movq	%rax, 124(%rsp)
	movl	$1, 132(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	104(%rsp)
	.cfi_def_cfa_offset 296
	leaq	_Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_(%rip), %rdi
	pushq	104(%rsp)
	.cfi_def_cfa_offset 304
	movq	140(%rsp), %rcx
	movl	148(%rsp), %r8d
	movq	128(%rsp), %rsi
	movl	136(%rsp), %edx
	leaq	176(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 296
	popq	%rdx
	.cfi_def_cfa_offset 288
.L2:
	movq	264(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$280, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2032:
	.size	_Z55__device_stub__Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_xPKdPKxS2_diddiS0_PyPKiS5_, .-_Z55__device_stub__Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_xPKdPKxS2_diddiS0_PyPKiS5_
	.globl	_Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_
	.type	_Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_, @function
_Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_:
.LFB2033:
	.cfi_startproc
	endbr64
	jmp	_Z55__device_stub__Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_xPKdPKxS2_diddiS0_PyPKiS5_
	.cfi_endproc
.LFE2033:
	.size	_Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_, .-_Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2035:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 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
	leaq	_Z11doPairCountxPKdPKxS2_diddiS0_PyPKiS5_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	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	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2035:
	.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 rafa on 2/5/21.
//

#ifndef CUPAC_COUNTER_CUH
#define CUPAC_COUNTER_CUH

__global__
void doPairCount(long long np, const double *array, const long long *particlesInGrid, const long long *offset,
                 double boxsize, int nside, double minsep, double maxsep, int nbins, const double *rbins2,
                 unsigned long long *paircounts, const int *key, const int *iwrap);

#endif //CUPAC_COUNTER_CUH

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