/***********************************************************************
!                        Version 1:  04/16 SCR                         *
!                                                                      *
!   CUDA/C optimization of:                                            *
!                                                                      *
!   SNSWP3D  - This routine calculates angular fluxes for a single     *
!              direction for an upstream corner-balance spatial        *
!              discretization in 3D.                                   *
!                                                                      *
!   Input:                                                             *
!                                                                      *
!   Output:                                                            *
!                                                                      *
!   Questions:                                                         *
!              1. Should the hyperplane pass over zones or corners?    *
!                 A: Corners                                           *
!              2.                                                      *
!                                                                      *
!**********************************************************************/

/*
  TODO:
       explicit setting of #hyperplanes

       
*/

#include<stdio.h>
//#include<omp.h>
//#include"nvToolsExt.h"

#define GROUPSET 16
#define NUMFACES 3

#define NUMSM 90

#define fouralpha 1.82
#define fouralpha4 5.82

//#include "snutil.cu"
//#include "snsweep.cu"

#define Connect(a,b,c) Connect[ a + 3 * ( b + mC * c ) ]


extern "C" {

  void callCudaHostMalloc ();
//  void CC_sweep(
__global__ void GPU_sweep(
          int  size_maxCorner,
          int  size_maxcf,
          int  nAngle,
          int  nzones,
          int  ncornr,
          int  Groups,
          int  nbelem,
          int* Angle,
       double* soa_omega,
          int* nextZ,
          int* next,
          int* soa_nCorner,
          int* soa_nCFaces,
          int* soa_c0,
       double* soa_STotal,
       double* soa_STime,
       double* soa_SigtInv,
       double* soa_Volume,
       double* soa_Sigt,
       double* soa_A_fp,
       double* soa_A_ez,
          int* soa_Connect,
       double* psic,
       double* psib );

  /*
    Simple C file to handle calling CUDA.  Easily accessible from Fortran.  
    This removes any need to compile any other part of the program with nvcc.
    Just this file needs nvcc.  Everything else just uses regular intel compilers.
    Simple.
   */

  void callcudasub (
                     int *numzones, 
		     int *numgroups,
		     int *ncornr,
		     int *numAngles,
                     int *AngleOrder,
		     int *maxcorners, 
		     int *maxfaces, 
		     int *octant,  //=binRecv
		     int *NangBin, 
                     int *nbelem,
		     double *omega,
                     int    *nCorner,
                     int    *nCFaces,
                     int    *c0,
		     double *A_fp, 
		     double *A_ez, 
		     int    *Connect,
		     double *STotal,
		     double *STime,
		     double *Volume,
                     double *psic,
                     double *psib,
		     int *next, 
                     int *nextZ,
		     double *Sigt,
		     double *SigtInv
		     ) 
  {
    //dump data to check
    static int dump_cnt=0;
    int zone,ic;
    static    int* d_AngleOrder;
    static double* d_omega;
    static    int* d_nextZ;
    static    int* d_next;
    static    int* d_nCorner;
    static    int* d_nCFaces;
    static    int* d_c0;
    static double* d_STotal;
    static double* d_STime;
    static double* d_SigtInv;
    static double* d_Volume;
    static double* d_Sigt;
    static double* d_A_fp;
    static double* d_A_ez;
    static    int* d_Connect;
    static double* d_psic;
    static double* d_psib;

    int nZ = *numzones;
    int nA = *numAngles;
    int mC = *maxcorners;
    int mF = *maxfaces;
    int nG = *numgroups;
    int nC = *ncornr;
    int nBe = *nbelem;


    if ( dump_cnt < 5 )
    {
//      for(zone=0;zone<nZ;zone++)
//      {
//        for(ic=0;ic<nCorner[zone]; ic++)
//        {
//          printf(" zone,corner,connect3 = %d,%d,%d \n",zone,ic,Connect(2,ic,zone) );
//        }
//      }

      printf("max faces=%d\n",mF);

      if ( dump_cnt == 0 )
      {
        cudaMalloc(&d_AngleOrder,sizeof(int)*8*nA);
        cudaMalloc(&d_omega,sizeof(double)*3*nA);
        cudaMalloc(&d_nextZ,sizeof(int)*nZ*nA);
        cudaMalloc(&d_next,sizeof(int)*(nC+1)*nA);
        cudaMalloc(&d_nCorner,sizeof(int)*nZ);
        cudaMalloc(&d_nCFaces,sizeof(int)*nZ);
        cudaMalloc(&d_c0,sizeof(int)*nZ);
        cudaMalloc(&d_STotal,sizeof(double)*nZ*nG*mC);
        cudaMalloc(&d_STime,sizeof(double)*nZ*nA*nG*mC);
        cudaMalloc(&d_SigtInv,sizeof(double)*nZ*nG);
        cudaMalloc(&d_Volume,sizeof(double)*nZ*mC);
        cudaMalloc(&d_Sigt,sizeof(double)*nZ*nG);
        cudaMalloc(&d_A_fp,sizeof(double)*3*nZ*mC*mF);
        cudaMalloc(&d_A_ez,sizeof(double)*3*nZ*mC*mF);
        cudaMalloc(&d_Connect,sizeof(int)*3*nZ*mC*mF);
        cudaMalloc(&d_psic,sizeof(double)*nG*nC*nA);
        cudaMalloc(&d_psib,sizeof(double)*nG*nBe*nA);
      }



      cudaMemcpy(d_AngleOrder,AngleOrder,sizeof(int)*8*nA,cudaMemcpyHostToDevice);
      cudaMemcpy(d_omega,omega,sizeof(double)*3*nA,cudaMemcpyHostToDevice);
      cudaMemcpy(d_nextZ,nextZ,sizeof(int)*nZ*nA,cudaMemcpyHostToDevice);
      cudaMemcpy(d_next,next,sizeof(int)*(nC+1)*nA,cudaMemcpyHostToDevice);
      cudaMemcpy(d_nCorner,nCorner,sizeof(int)*nZ,cudaMemcpyHostToDevice);
      cudaMemcpy(d_nCFaces,nCFaces,sizeof(int)*nZ,cudaMemcpyHostToDevice);
      cudaMemcpy(d_c0,c0,sizeof(int)*nZ,cudaMemcpyHostToDevice);
      cudaMemcpy(d_STotal,STotal,sizeof(double)*nZ*nG*mC,cudaMemcpyHostToDevice);
      cudaMemcpy(d_STime,STime,sizeof(double)*nZ*nA*nG*mC,cudaMemcpyHostToDevice);
      cudaMemcpy(d_SigtInv,SigtInv,sizeof(double)*nZ*nG,cudaMemcpyHostToDevice);
      cudaMemcpy(d_Volume,Volume,sizeof(double)*nZ*mC,cudaMemcpyHostToDevice);
      cudaMemcpy(d_Sigt,Sigt,sizeof(double)*nZ*nG,cudaMemcpyHostToDevice);
      cudaMemcpy(d_A_fp,A_fp,sizeof(double)*3*nZ*mC*mF,cudaMemcpyHostToDevice);
      cudaMemcpy(d_A_ez,A_ez,sizeof(double)*3*nZ*mC*mF,cudaMemcpyHostToDevice);
      cudaMemcpy(d_Connect,Connect,sizeof(int)*3*nZ*mC*mF,cudaMemcpyHostToDevice);
      cudaMemcpy(d_psib,psib,sizeof(double)*nG*nBe*nA,cudaMemcpyHostToDevice);
      cudaMemcpy(d_psic,psic,sizeof(double)*nG*nC*nA,cudaMemcpyHostToDevice);

//
//      
      GPU_sweep<<<nA,32>>>(
                       mC,                 
                       mF,       //                 
                       nA,                  
                       nZ,                   
                       nC,                          
                       nG,                  
                       nBe,                         
                       d_AngleOrder,                
                       d_omega,                     
                       d_nextZ,                     
                       d_next,                      
                       d_nCorner,                   
                       d_nCFaces,                   
                       d_c0,                        
                       d_STotal,                    
                       d_STime,                     
                       d_SigtInv,                   
                       d_Volume,                    
                       d_Sigt,                      
                       d_A_fp,                      
                       d_A_ez,                      
                       d_Connect,                   
                       d_psic,                  
                       d_psib        
                          );

      cudaMemcpy(psic,d_psic,sizeof(double)*nG*nC*nA,cudaMemcpyDeviceToHost);



//      CC_sweep(
//                  *maxcorners,      
//                  mF,       //*maxfaces,      
//                  *numAngles,      
//                  *numzones,      
//                  nC,      
//                  *numgroups,      
//                  nBe,      
//                  AngleOrder,      
//                  omega,      
//                  nextZ,      
//                  next,      
//                  nCorner,      
//                  nCFaces,      
//                  c0,      
//                  STotal,      
//                  STime,      
//                  SigtInv,      
//                  Volume,      
//                  Sigt,      
//                  A_fp,      
//                  A_ez,      
//                  Connect,      
//                  psic,      
//                  psib               
//             );
                  

      dump_cnt++;
    }
  } //callcuasub

}  // extern C


	.file	"tmpxft_0020969f_00000000-6_callcudasub.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2030:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2030:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0, @function
_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0:
.LFB2170:
	.cfi_startproc
	jmp	cudaMalloc@PLT
	.cfi_endproc
.LFE2170:
	.size	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0, .-_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"max faces=%d\n"
	.text
	.globl	callcudasub
	.type	callcudasub, @function
callcudasub:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$168, %rsp
	.cfi_def_cfa_offset 224
	cmpl	$4, _ZZ11callcudasubE8dump_cnt(%rip)
	movq	%r8, 112(%rsp)
	jg	.L3
	movl	(%r9), %ebx
	movq	%rdx, %rax
	movq	224(%rsp), %rdx
	movl	(%rax), %r12d
	movq	248(%rsp), %rax
	movl	(%rcx), %ebp
	movl	(%rdi), %r13d
	movl	%ebx, 44(%rsp)
	movl	$2, %edi
	movl	(%rdx), %ebx
	movl	(%rsi), %r14d
	leaq	.LC0(%rip), %rsi
	movl	(%rax), %r15d
	xorl	%eax, %eax
	movl	%ebx, %edx
	call	__printf_chk@PLT
	movslq	%ebp, %rcx
	movslq	44(%rsp), %rdi
	imulq	$24, %rcx, %rax
	movq	%rcx, %r9
	movq	%rcx, %rdx
	salq	$5, %rdx
	movq	%rax, 48(%rsp)
	movslq	%r13d, %rax
	imulq	%rax, %r9
	movq	%rax, %r10
	leaq	0(,%r9,4), %rsi
	movq	%rsi, 56(%rsp)
	leal	1(%r12), %esi
	movslq	%esi, %rsi
	imulq	%rcx, %rsi
	salq	$2, %rsi
	movq	%rsi, 64(%rsp)
	leaq	0(,%rax,4), %rsi
	movq	%rsi, 8(%rsp)
	movslq	%r14d, %rsi
	imulq	%rdi, %rax
	imulq	%rsi, %r10
	movq	%r10, %r8
	imulq	%rdi, %r8
	leaq	0(,%r8,8), %r11
	movq	%rsi, %r8
	imulq	%rdi, %r8
	leaq	0(,%rax,8), %rdi
	movq	%r11, 72(%rsp)
	movq	%rdi, 88(%rsp)
	movslq	%ebx, %rdi
	imulq	%rdi, %rax
	imulq	%r9, %r8
	imulq	$24, %rax, %rdi
	imulq	$12, %rax, %rax
	leaq	0(,%r8,8), %r11
	movq	%r11, 80(%rsp)
	leaq	0(,%r10,8), %r11
	movq	%r11, 16(%rsp)
	movq	%rax, 96(%rsp)
	movslq	%r12d, %rax
	imulq	%rsi, %rax
	movq	%rdi, 24(%rsp)
	imulq	%rcx, %rax
	salq	$3, %rax
	movq	%rax, 32(%rsp)
	movslq	%r15d, %rax
	imulq	%rsi, %rax
	imulq	%rcx, %rax
	salq	$3, %rax
	cmpl	$0, _ZZ11callcudasubE8dump_cnt(%rip)
	movq	%rax, 104(%rsp)
	jne	.L5
	movq	%rdx, %rsi
	leaq	_ZZ11callcudasubE12d_AngleOrder(%rip), %rdi
	movq	%rdx, 120(%rsp)
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	48(%rsp), %rsi
	leaq	_ZZ11callcudasubE7d_omega(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	56(%rsp), %rsi
	leaq	_ZZ11callcudasubE7d_nextZ(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	64(%rsp), %rsi
	leaq	_ZZ11callcudasubE6d_next(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	8(%rsp), %rsi
	leaq	_ZZ11callcudasubE9d_nCorner(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	8(%rsp), %rsi
	leaq	_ZZ11callcudasubE9d_nCFaces(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	8(%rsp), %rsi
	leaq	_ZZ11callcudasubE4d_c0(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	72(%rsp), %rsi
	leaq	_ZZ11callcudasubE8d_STotal(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	80(%rsp), %rsi
	leaq	_ZZ11callcudasubE7d_STime(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	16(%rsp), %rsi
	leaq	_ZZ11callcudasubE9d_SigtInv(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	88(%rsp), %rsi
	leaq	_ZZ11callcudasubE8d_Volume(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	16(%rsp), %rsi
	leaq	_ZZ11callcudasubE6d_Sigt(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	24(%rsp), %rsi
	leaq	_ZZ11callcudasubE6d_A_fp(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	24(%rsp), %rsi
	leaq	_ZZ11callcudasubE6d_A_ez(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	96(%rsp), %rsi
	leaq	_ZZ11callcudasubE9d_Connect(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	32(%rsp), %rsi
	leaq	_ZZ11callcudasubE6d_psic(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	104(%rsp), %rsi
	leaq	_ZZ11callcudasubE6d_psib(%rip), %rdi
	call	_Z10cudaMallocIiE9cudaErrorPPT_m.isra.0
	movq	120(%rsp), %rdx
.L5:
	movq	112(%rsp), %rsi
	movq	_ZZ11callcudasubE12d_AngleOrder(%rip), %rdi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	48(%rsp), %rdx
	movl	$1, %ecx
	movq	256(%rsp), %rsi
	movq	_ZZ11callcudasubE7d_omega(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	56(%rsp), %rdx
	movl	$1, %ecx
	movq	360(%rsp), %rsi
	movq	_ZZ11callcudasubE7d_nextZ(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	64(%rsp), %rdx
	movl	$1, %ecx
	movq	352(%rsp), %rsi
	movq	_ZZ11callcudasubE6d_next(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdx
	movl	$1, %ecx
	movq	264(%rsp), %rsi
	movq	_ZZ11callcudasubE9d_nCorner(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdx
	movl	$1, %ecx
	movq	272(%rsp), %rsi
	movq	_ZZ11callcudasubE9d_nCFaces(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdx
	movl	$1, %ecx
	movq	280(%rsp), %rsi
	movq	_ZZ11callcudasubE4d_c0(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	72(%rsp), %rdx
	movl	$1, %ecx
	movq	312(%rsp), %rsi
	movq	_ZZ11callcudasubE8d_STotal(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	80(%rsp), %rdx
	movl	$1, %ecx
	movq	320(%rsp), %rsi
	movq	_ZZ11callcudasubE7d_STime(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdx
	movl	$1, %ecx
	movq	376(%rsp), %rsi
	movq	_ZZ11callcudasubE9d_SigtInv(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	88(%rsp), %rdx
	movl	$1, %ecx
	movq	328(%rsp), %rsi
	movq	_ZZ11callcudasubE8d_Volume(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdx
	movl	$1, %ecx
	movq	368(%rsp), %rsi
	movq	_ZZ11callcudasubE6d_Sigt(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	24(%rsp), %rdx
	movl	$1, %ecx
	movq	288(%rsp), %rsi
	movq	_ZZ11callcudasubE6d_A_fp(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	24(%rsp), %rdx
	movl	$1, %ecx
	movq	296(%rsp), %rsi
	movq	_ZZ11callcudasubE6d_A_ez(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	96(%rsp), %rdx
	movl	$1, %ecx
	movq	304(%rsp), %rsi
	movq	_ZZ11callcudasubE9d_Connect(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	104(%rsp), %rdx
	movl	$1, %ecx
	movq	344(%rsp), %rsi
	movq	_ZZ11callcudasubE6d_psib(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	32(%rsp), %rdx
	movl	$1, %ecx
	movq	336(%rsp), %rsi
	movq	_ZZ11callcudasubE6d_psic(%rip), %rdi
	call	cudaMemcpy@PLT
	movl	$134217729, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	%ebp, 136(%rsp)
	salq	$5, %rdx
	movl	$1, %ecx
	movabsq	$4294967297, %rax
	movq	%rax, 140(%rsp)
	movq	136(%rsp), %rdi
	movl	144(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L6
	pushq	_ZZ11callcudasubE6d_psib(%rip)
	.cfi_def_cfa_offset 232
	movl	%r14d, %r9d
	movl	%r12d, %r8d
	movl	%r13d, %ecx
	pushq	_ZZ11callcudasubE6d_psic(%rip)
	.cfi_def_cfa_offset 240
	movl	%ebp, %edx
	movl	%ebx, %esi
	pushq	_ZZ11callcudasubE9d_Connect(%rip)
	.cfi_def_cfa_offset 248
	pushq	_ZZ11callcudasubE6d_A_ez(%rip)
	.cfi_def_cfa_offset 256
	pushq	_ZZ11callcudasubE6d_A_fp(%rip)
	.cfi_def_cfa_offset 264
	pushq	_ZZ11callcudasubE6d_Sigt(%rip)
	.cfi_def_cfa_offset 272
	pushq	_ZZ11callcudasubE8d_Volume(%rip)
	.cfi_def_cfa_offset 280
	pushq	_ZZ11callcudasubE9d_SigtInv(%rip)
	.cfi_def_cfa_offset 288
	pushq	_ZZ11callcudasubE7d_STime(%rip)
	.cfi_def_cfa_offset 296
	pushq	_ZZ11callcudasubE8d_STotal(%rip)
	.cfi_def_cfa_offset 304
	pushq	_ZZ11callcudasubE4d_c0(%rip)
	.cfi_def_cfa_offset 312
	pushq	_ZZ11callcudasubE9d_nCFaces(%rip)
	.cfi_def_cfa_offset 320
	pushq	_ZZ11callcudasubE9d_nCorner(%rip)
	.cfi_def_cfa_offset 328
	pushq	_ZZ11callcudasubE6d_next(%rip)
	.cfi_def_cfa_offset 336
	pushq	_ZZ11callcudasubE7d_nextZ(%rip)
	.cfi_def_cfa_offset 344
	pushq	_ZZ11callcudasubE7d_omega(%rip)
	.cfi_def_cfa_offset 352
	pushq	_ZZ11callcudasubE12d_AngleOrder(%rip)
	.cfi_def_cfa_offset 360
	pushq	%r15
	.cfi_def_cfa_offset 368
	movl	188(%rsp), %edi
	call	GPU_sweep@PLT
	addq	$144, %rsp
	.cfi_def_cfa_offset 224
.L6:
	movq	32(%rsp), %rdx
	movq	_ZZ11callcudasubE6d_psic(%rip), %rsi
	movl	$2, %ecx
	movq	336(%rsp), %rdi
	call	cudaMemcpy@PLT
	incl	_ZZ11callcudasubE8dump_cnt(%rip)
.L3:
	addq	$168, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	callcudasub, .-callcudasub
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2053:
	.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
.LFE2053:
	.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	_ZZ11callcudasubE6d_psib
	.comm	_ZZ11callcudasubE6d_psib,8,8
	.local	_ZZ11callcudasubE6d_psic
	.comm	_ZZ11callcudasubE6d_psic,8,8
	.local	_ZZ11callcudasubE9d_Connect
	.comm	_ZZ11callcudasubE9d_Connect,8,8
	.local	_ZZ11callcudasubE6d_A_ez
	.comm	_ZZ11callcudasubE6d_A_ez,8,8
	.local	_ZZ11callcudasubE6d_A_fp
	.comm	_ZZ11callcudasubE6d_A_fp,8,8
	.local	_ZZ11callcudasubE6d_Sigt
	.comm	_ZZ11callcudasubE6d_Sigt,8,8
	.local	_ZZ11callcudasubE8d_Volume
	.comm	_ZZ11callcudasubE8d_Volume,8,8
	.local	_ZZ11callcudasubE9d_SigtInv
	.comm	_ZZ11callcudasubE9d_SigtInv,8,8
	.local	_ZZ11callcudasubE7d_STime
	.comm	_ZZ11callcudasubE7d_STime,8,8
	.local	_ZZ11callcudasubE8d_STotal
	.comm	_ZZ11callcudasubE8d_STotal,8,8
	.local	_ZZ11callcudasubE4d_c0
	.comm	_ZZ11callcudasubE4d_c0,8,8
	.local	_ZZ11callcudasubE9d_nCFaces
	.comm	_ZZ11callcudasubE9d_nCFaces,8,8
	.local	_ZZ11callcudasubE9d_nCorner
	.comm	_ZZ11callcudasubE9d_nCorner,8,8
	.local	_ZZ11callcudasubE6d_next
	.comm	_ZZ11callcudasubE6d_next,8,8
	.local	_ZZ11callcudasubE7d_nextZ
	.comm	_ZZ11callcudasubE7d_nextZ,8,8
	.local	_ZZ11callcudasubE7d_omega
	.comm	_ZZ11callcudasubE7d_omega,8,8
	.local	_ZZ11callcudasubE12d_AngleOrder
	.comm	_ZZ11callcudasubE12d_AngleOrder,8,8
	.local	_ZZ11callcudasubE8dump_cnt
	.comm	_ZZ11callcudasubE8dump_cnt,4,4
	.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:

/***********************************************************************
!                        Version 1:  04/16 SCR                         *
!                                                                      *
!   CUDA/C optimization of:                                            *
!                                                                      *
!   SNSWP3D  - This routine calculates angular fluxes for a single     *
!              direction for an upstream corner-balance spatial        *
!              discretization in 3D.                                   *
!                                                                      *
!   Input:                                                             *
!                                                                      *
!   Output:                                                            *
!                                                                      *
!   Questions:                                                         *
!              1. Should the hyperplane pass over zones or corners?    *
!                 A: Corners                                           *
!              2.                                                      *
!                                                                      *
!**********************************************************************/

/*
  TODO:
       explicit setting of #hyperplanes

       
*/


#include <hip/hip_runtime.h>
#include<stdio.h>
//#include<omp.h>
//#include"nvToolsExt.h"

#define GROUPSET 16
#define NUMFACES 3

#define NUMSM 90

#define fouralpha 1.82
#define fouralpha4 5.82

//#include "snutil.cu"
//#include "snsweep.cu"

#define Connect(a,b,c) Connect[ a + 3 * ( b + mC * c ) ]


extern "C" {

  void callCudaHostMalloc ();
//  void CC_sweep(
__global__ void GPU_sweep(
          int  size_maxCorner,
          int  size_maxcf,
          int  nAngle,
          int  nzones,
          int  ncornr,
          int  Groups,
          int  nbelem,
          int* Angle,
       double* soa_omega,
          int* nextZ,
          int* next,
          int* soa_nCorner,
          int* soa_nCFaces,
          int* soa_c0,
       double* soa_STotal,
       double* soa_STime,
       double* soa_SigtInv,
       double* soa_Volume,
       double* soa_Sigt,
       double* soa_A_fp,
       double* soa_A_ez,
          int* soa_Connect,
       double* psic,
       double* psib );

  /*
    Simple C file to handle calling CUDA.  Easily accessible from Fortran.  
    This removes any need to compile any other part of the program with nvcc.
    Just this file needs nvcc.  Everything else just uses regular intel compilers.
    Simple.
   */

  void callcudasub (
                     int *numzones, 
		     int *numgroups,
		     int *ncornr,
		     int *numAngles,
                     int *AngleOrder,
		     int *maxcorners, 
		     int *maxfaces, 
		     int *octant,  //=binRecv
		     int *NangBin, 
                     int *nbelem,
		     double *omega,
                     int    *nCorner,
                     int    *nCFaces,
                     int    *c0,
		     double *A_fp, 
		     double *A_ez, 
		     int    *Connect,
		     double *STotal,
		     double *STime,
		     double *Volume,
                     double *psic,
                     double *psib,
		     int *next, 
                     int *nextZ,
		     double *Sigt,
		     double *SigtInv
		     ) 
  {
    //dump data to check
    static int dump_cnt=0;
    int zone,ic;
    static    int* d_AngleOrder;
    static double* d_omega;
    static    int* d_nextZ;
    static    int* d_next;
    static    int* d_nCorner;
    static    int* d_nCFaces;
    static    int* d_c0;
    static double* d_STotal;
    static double* d_STime;
    static double* d_SigtInv;
    static double* d_Volume;
    static double* d_Sigt;
    static double* d_A_fp;
    static double* d_A_ez;
    static    int* d_Connect;
    static double* d_psic;
    static double* d_psib;

    int nZ = *numzones;
    int nA = *numAngles;
    int mC = *maxcorners;
    int mF = *maxfaces;
    int nG = *numgroups;
    int nC = *ncornr;
    int nBe = *nbelem;


    if ( dump_cnt < 5 )
    {
//      for(zone=0;zone<nZ;zone++)
//      {
//        for(ic=0;ic<nCorner[zone]; ic++)
//        {
//          printf(" zone,corner,connect3 = %d,%d,%d \n",zone,ic,Connect(2,ic,zone) );
//        }
//      }

      printf("max faces=%d\n",mF);

      if ( dump_cnt == 0 )
      {
        hipMalloc(&d_AngleOrder,sizeof(int)*8*nA);
        hipMalloc(&d_omega,sizeof(double)*3*nA);
        hipMalloc(&d_nextZ,sizeof(int)*nZ*nA);
        hipMalloc(&d_next,sizeof(int)*(nC+1)*nA);
        hipMalloc(&d_nCorner,sizeof(int)*nZ);
        hipMalloc(&d_nCFaces,sizeof(int)*nZ);
        hipMalloc(&d_c0,sizeof(int)*nZ);
        hipMalloc(&d_STotal,sizeof(double)*nZ*nG*mC);
        hipMalloc(&d_STime,sizeof(double)*nZ*nA*nG*mC);
        hipMalloc(&d_SigtInv,sizeof(double)*nZ*nG);
        hipMalloc(&d_Volume,sizeof(double)*nZ*mC);
        hipMalloc(&d_Sigt,sizeof(double)*nZ*nG);
        hipMalloc(&d_A_fp,sizeof(double)*3*nZ*mC*mF);
        hipMalloc(&d_A_ez,sizeof(double)*3*nZ*mC*mF);
        hipMalloc(&d_Connect,sizeof(int)*3*nZ*mC*mF);
        hipMalloc(&d_psic,sizeof(double)*nG*nC*nA);
        hipMalloc(&d_psib,sizeof(double)*nG*nBe*nA);
      }



      hipMemcpy(d_AngleOrder,AngleOrder,sizeof(int)*8*nA,hipMemcpyHostToDevice);
      hipMemcpy(d_omega,omega,sizeof(double)*3*nA,hipMemcpyHostToDevice);
      hipMemcpy(d_nextZ,nextZ,sizeof(int)*nZ*nA,hipMemcpyHostToDevice);
      hipMemcpy(d_next,next,sizeof(int)*(nC+1)*nA,hipMemcpyHostToDevice);
      hipMemcpy(d_nCorner,nCorner,sizeof(int)*nZ,hipMemcpyHostToDevice);
      hipMemcpy(d_nCFaces,nCFaces,sizeof(int)*nZ,hipMemcpyHostToDevice);
      hipMemcpy(d_c0,c0,sizeof(int)*nZ,hipMemcpyHostToDevice);
      hipMemcpy(d_STotal,STotal,sizeof(double)*nZ*nG*mC,hipMemcpyHostToDevice);
      hipMemcpy(d_STime,STime,sizeof(double)*nZ*nA*nG*mC,hipMemcpyHostToDevice);
      hipMemcpy(d_SigtInv,SigtInv,sizeof(double)*nZ*nG,hipMemcpyHostToDevice);
      hipMemcpy(d_Volume,Volume,sizeof(double)*nZ*mC,hipMemcpyHostToDevice);
      hipMemcpy(d_Sigt,Sigt,sizeof(double)*nZ*nG,hipMemcpyHostToDevice);
      hipMemcpy(d_A_fp,A_fp,sizeof(double)*3*nZ*mC*mF,hipMemcpyHostToDevice);
      hipMemcpy(d_A_ez,A_ez,sizeof(double)*3*nZ*mC*mF,hipMemcpyHostToDevice);
      hipMemcpy(d_Connect,Connect,sizeof(int)*3*nZ*mC*mF,hipMemcpyHostToDevice);
      hipMemcpy(d_psib,psib,sizeof(double)*nG*nBe*nA,hipMemcpyHostToDevice);
      hipMemcpy(d_psic,psic,sizeof(double)*nG*nC*nA,hipMemcpyHostToDevice);

//
//      
      GPU_sweep<<<nA,32>>>(
                       mC,                 
                       mF,       //                 
                       nA,                  
                       nZ,                   
                       nC,                          
                       nG,                  
                       nBe,                         
                       d_AngleOrder,                
                       d_omega,                     
                       d_nextZ,                     
                       d_next,                      
                       d_nCorner,                   
                       d_nCFaces,                   
                       d_c0,                        
                       d_STotal,                    
                       d_STime,                     
                       d_SigtInv,                   
                       d_Volume,                    
                       d_Sigt,                      
                       d_A_fp,                      
                       d_A_ez,                      
                       d_Connect,                   
                       d_psic,                  
                       d_psib        
                          );

      hipMemcpy(psic,d_psic,sizeof(double)*nG*nC*nA,hipMemcpyDeviceToHost);



//      CC_sweep(
//                  *maxcorners,      
//                  mF,       //*maxfaces,      
//                  *numAngles,      
//                  *numzones,      
//                  nC,      
//                  *numgroups,      
//                  nBe,      
//                  AngleOrder,      
//                  omega,      
//                  nextZ,      
//                  next,      
//                  nCorner,      
//                  nCFaces,      
//                  c0,      
//                  STotal,      
//                  STime,      
//                  SigtInv,      
//                  Volume,      
//                  Sigt,      
//                  A_fp,      
//                  A_ez,      
//                  Connect,      
//                  psic,      
//                  psib               
//             );
                  

      dump_cnt++;
    }
  } //callcuasub

}  // extern C


	.text
	.file	"callcudasub.hip"
	.globl	callcudasub                     # -- Begin function callcudasub
	.type	callcudasub,@function
callcudasub:                            # @callcudasub
	.cfi_startproc
# %bb.0:
	cmpl	$4, _ZZ11callcudasubE8dump_cnt(%rip)
	jg	.LBB0_7
# %bb.1:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$136, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r8, 120(%rsp)                  # 8-byte Spill
	movq	216(%rsp), %rax
	movq	192(%rsp), %r8
	movslq	(%rdi), %r12
	movslq	(%rcx), %r15
	movslq	(%r9), %rcx
	movq	%rcx, 8(%rsp)                   # 8-byte Spill
	movslq	(%r8), %rcx
	movslq	(%rsi), %rbp
	movslq	(%rdx), %rbx
	movslq	(%rax), %rax
	movq	%rax, 24(%rsp)                  # 8-byte Spill
	movl	$.L.str, %edi
	movq	%rcx, 16(%rsp)                  # 8-byte Spill
	movl	%ecx, %esi
	xorl	%eax, %eax
	callq	printf
	movq	%r15, %rdx
	shlq	$5, %rdx
	cmpl	$0, _ZZ11callcudasubE8dump_cnt(%rip)
	movq	%r12, 40(%rsp)                  # 8-byte Spill
	movq	%r15, 128(%rsp)                 # 8-byte Spill
	movq	%rbp, 112(%rsp)                 # 8-byte Spill
	movq	%rbx, 32(%rsp)                  # 8-byte Spill
	je	.LBB0_3
# %bb.2:                                # %._crit_edge
	leaq	(,%r15,8), %rax
	leaq	(%rax,%rax,2), %rax
	movq	%rax, 80(%rsp)                  # 8-byte Spill
	leaq	(,%r12,4), %rax
	movq	%rax, 88(%rsp)                  # 8-byte Spill
	imulq	%r15, %rax
	movq	%rax, 72(%rsp)                  # 8-byte Spill
	movq	%rbx, %rdi
	incq	%rbx
	imulq	%r15, %rbx
	shlq	$2, %rbx
	leaq	(,%r12,8), %r13
	movq	%r13, %r14
	imulq	%rbp, %r14
	movq	%r14, %rax
	movq	8(%rsp), %rcx                   # 8-byte Reload
	imulq	%rcx, %rax
	movq	%rax, 64(%rsp)                  # 8-byte Spill
	imulq	%rcx, %r13
	imulq	%rbp, %r15
	movq	%r13, %rax
	imulq	%r15, %rax
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	movq	%rcx, %rax
	imulq	%r12, %rax
	imulq	16(%rsp), %rax                  # 8-byte Folded Reload
	leaq	(,%rax,8), %rcx
	leaq	(%rcx,%rcx,2), %rcx
	movq	%rcx, 56(%rsp)                  # 8-byte Spill
	shlq	$2, %rax
	leaq	(%rax,%rax,2), %rbp
	shlq	$3, %r15
	movq	%r15, %r12
	imulq	24(%rsp), %r12                  # 8-byte Folded Reload
	imulq	%rdi, %r15
	jmp	.LBB0_4
.LBB0_3:
	movl	$_ZZ11callcudasubE12d_AngleOrder, %edi
	movq	%rdx, %rsi
	movq	%rdx, 104(%rsp)                 # 8-byte Spill
	callq	hipMalloc
	leaq	(,%r15,8), %rax
	leaq	(%rax,%rax,2), %rsi
	movl	$_ZZ11callcudasubE7d_omega, %edi
	movq	%rsi, 80(%rsp)                  # 8-byte Spill
	callq	hipMalloc
	leaq	(,%r12,4), %r14
	movq	%r14, %rsi
	imulq	%r15, %rsi
	movl	$_ZZ11callcudasubE7d_nextZ, %edi
	movq	%rsi, 72(%rsp)                  # 8-byte Spill
	callq	hipMalloc
	leaq	1(%rbx), %rsi
	imulq	%r15, %rsi
	shlq	$2, %rsi
	movq	%rsi, 96(%rsp)                  # 8-byte Spill
	movl	$_ZZ11callcudasubE6d_next, %edi
	callq	hipMalloc
	movl	$_ZZ11callcudasubE9d_nCorner, %edi
	movq	%r14, %rsi
	callq	hipMalloc
	movl	$_ZZ11callcudasubE9d_nCFaces, %edi
	movq	%r14, %rsi
	callq	hipMalloc
	movl	$_ZZ11callcudasubE4d_c0, %edi
	movq	%r14, 88(%rsp)                  # 8-byte Spill
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	(,%r12,8), %r13
	movq	%r13, %r14
	imulq	%rbp, %r14
	movq	%r14, %rsi
	movq	8(%rsp), %rbx                   # 8-byte Reload
	imulq	%rbx, %rsi
	movl	$_ZZ11callcudasubE8d_STotal, %edi
	movq	%rsi, 64(%rsp)                  # 8-byte Spill
	callq	hipMalloc
	imulq	%rbx, %r13
	movq	%r15, %r12
	imulq	%rbp, %r12
	movq	%r13, %rsi
	imulq	%r12, %rsi
	movl	$_ZZ11callcudasubE7d_STime, %edi
	movq	%rsi, 48(%rsp)                  # 8-byte Spill
	callq	hipMalloc
	movl	$_ZZ11callcudasubE9d_SigtInv, %edi
	movq	%r14, %rsi
	callq	hipMalloc
	movl	$_ZZ11callcudasubE8d_Volume, %edi
	movq	%r13, %rsi
	callq	hipMalloc
	movl	$_ZZ11callcudasubE6d_Sigt, %edi
	movq	%r14, %rsi
	callq	hipMalloc
	movq	%rbx, %r15
	movq	96(%rsp), %rbx                  # 8-byte Reload
	imulq	40(%rsp), %r15                  # 8-byte Folded Reload
	imulq	16(%rsp), %r15                  # 8-byte Folded Reload
	leaq	(,%r15,8), %rax
	leaq	(%rax,%rax,2), %rbp
	movl	$_ZZ11callcudasubE6d_A_fp, %edi
	movq	%rbp, %rsi
	callq	hipMalloc
	movl	$_ZZ11callcudasubE6d_A_ez, %edi
	movq	%rbp, 56(%rsp)                  # 8-byte Spill
	movq	%rbp, %rsi
	callq	hipMalloc
	shlq	$2, %r15
	leaq	(%r15,%r15,2), %rsi
	movl	$_ZZ11callcudasubE9d_Connect, %edi
	movq	%rsi, %rbp
	callq	hipMalloc
	shlq	$3, %r12
	movq	%r12, %r15
	imulq	32(%rsp), %r15                  # 8-byte Folded Reload
	movl	$_ZZ11callcudasubE6d_psic, %edi
	movq	%r15, %rsi
	callq	hipMalloc
	imulq	24(%rsp), %r12                  # 8-byte Folded Reload
	movl	$_ZZ11callcudasubE6d_psib, %edi
	movq	%r12, %rsi
	callq	hipMalloc
	movq	104(%rsp), %rdx                 # 8-byte Reload
.LBB0_4:
	movq	_ZZ11callcudasubE12d_AngleOrder(%rip), %rdi
	movq	120(%rsp), %rsi                 # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE7d_omega(%rip), %rdi
	movq	224(%rsp), %rsi
	movq	80(%rsp), %rdx                  # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE7d_nextZ(%rip), %rdi
	movq	328(%rsp), %rsi
	movq	72(%rsp), %rdx                  # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE6d_next(%rip), %rdi
	movq	320(%rsp), %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE9d_nCorner(%rip), %rdi
	movq	232(%rsp), %rsi
	movq	88(%rsp), %rbx                  # 8-byte Reload
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE9d_nCFaces(%rip), %rdi
	movq	240(%rsp), %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE4d_c0(%rip), %rdi
	movq	248(%rsp), %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE8d_STotal(%rip), %rdi
	movq	280(%rsp), %rsi
	movq	64(%rsp), %rdx                  # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE7d_STime(%rip), %rdi
	movq	288(%rsp), %rsi
	movq	48(%rsp), %rdx                  # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE9d_SigtInv(%rip), %rdi
	movq	344(%rsp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE8d_Volume(%rip), %rdi
	movq	296(%rsp), %rsi
	movq	%r13, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE6d_Sigt(%rip), %rdi
	movq	336(%rsp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE6d_A_fp(%rip), %rdi
	movq	256(%rsp), %rsi
	movq	56(%rsp), %rbx                  # 8-byte Reload
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE6d_A_ez(%rip), %rdi
	movq	264(%rsp), %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE9d_Connect(%rip), %rdi
	movq	272(%rsp), %rsi
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE6d_psib(%rip), %rdi
	movq	312(%rsp), %rsi
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	_ZZ11callcudasubE6d_psic(%rip), %rdi
	movq	304(%rsp), %rbx
	movq	%rbx, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	128(%rsp), %r14                 # 8-byte Reload
	movl	%r14d, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$32, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_6
# %bb.5:
	movq	8(%rsp), %rdi                   # 8-byte Reload
                                        # kill: def $edi killed $edi killed $rdi
	movq	16(%rsp), %rsi                  # 8-byte Reload
                                        # kill: def $esi killed $esi killed $rsi
	movl	%r14d, %edx
	movq	40(%rsp), %rcx                  # 8-byte Reload
                                        # kill: def $ecx killed $ecx killed $rcx
	movq	32(%rsp), %r8                   # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	movq	112(%rsp), %r9                  # 8-byte Reload
                                        # kill: def $r9d killed $r9d killed $r9
	pushq	_ZZ11callcudasubE6d_psib(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE6d_psic(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE9d_Connect(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE6d_A_ez(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE6d_A_fp(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE6d_Sigt(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE8d_Volume(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE9d_SigtInv(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE7d_STime(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE8d_STotal(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE4d_c0(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE9d_nCFaces(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE9d_nCorner(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE6d_next(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE7d_nextZ(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE7d_omega(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	_ZZ11callcudasubE12d_AngleOrder(%rip)
	.cfi_adjust_cfa_offset 8
	pushq	160(%rsp)                       # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	callq	__device_stub__GPU_sweep
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
.LBB0_6:
	movq	_ZZ11callcudasubE6d_psic(%rip), %rsi
	movq	%rbx, %rdi
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	incl	_ZZ11callcudasubE8dump_cnt(%rip)
	addq	$136, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r12
	.cfi_restore %r13
	.cfi_restore %r14
	.cfi_restore %r15
	.cfi_restore %rbp
.LBB0_7:
	retq
.Lfunc_end0:
	.size	callcudasub, .Lfunc_end0-callcudasub
	.cfi_endproc
                                        # -- End function
	.type	_ZZ11callcudasubE8dump_cnt,@object # @_ZZ11callcudasubE8dump_cnt
	.local	_ZZ11callcudasubE8dump_cnt
	.comm	_ZZ11callcudasubE8dump_cnt,4,4
	.type	_ZZ11callcudasubE12d_AngleOrder,@object # @_ZZ11callcudasubE12d_AngleOrder
	.local	_ZZ11callcudasubE12d_AngleOrder
	.comm	_ZZ11callcudasubE12d_AngleOrder,8,8
	.type	_ZZ11callcudasubE7d_omega,@object # @_ZZ11callcudasubE7d_omega
	.local	_ZZ11callcudasubE7d_omega
	.comm	_ZZ11callcudasubE7d_omega,8,8
	.type	_ZZ11callcudasubE7d_nextZ,@object # @_ZZ11callcudasubE7d_nextZ
	.local	_ZZ11callcudasubE7d_nextZ
	.comm	_ZZ11callcudasubE7d_nextZ,8,8
	.type	_ZZ11callcudasubE6d_next,@object # @_ZZ11callcudasubE6d_next
	.local	_ZZ11callcudasubE6d_next
	.comm	_ZZ11callcudasubE6d_next,8,8
	.type	_ZZ11callcudasubE9d_nCorner,@object # @_ZZ11callcudasubE9d_nCorner
	.local	_ZZ11callcudasubE9d_nCorner
	.comm	_ZZ11callcudasubE9d_nCorner,8,8
	.type	_ZZ11callcudasubE9d_nCFaces,@object # @_ZZ11callcudasubE9d_nCFaces
	.local	_ZZ11callcudasubE9d_nCFaces
	.comm	_ZZ11callcudasubE9d_nCFaces,8,8
	.type	_ZZ11callcudasubE4d_c0,@object  # @_ZZ11callcudasubE4d_c0
	.local	_ZZ11callcudasubE4d_c0
	.comm	_ZZ11callcudasubE4d_c0,8,8
	.type	_ZZ11callcudasubE8d_STotal,@object # @_ZZ11callcudasubE8d_STotal
	.local	_ZZ11callcudasubE8d_STotal
	.comm	_ZZ11callcudasubE8d_STotal,8,8
	.type	_ZZ11callcudasubE7d_STime,@object # @_ZZ11callcudasubE7d_STime
	.local	_ZZ11callcudasubE7d_STime
	.comm	_ZZ11callcudasubE7d_STime,8,8
	.type	_ZZ11callcudasubE9d_SigtInv,@object # @_ZZ11callcudasubE9d_SigtInv
	.local	_ZZ11callcudasubE9d_SigtInv
	.comm	_ZZ11callcudasubE9d_SigtInv,8,8
	.type	_ZZ11callcudasubE8d_Volume,@object # @_ZZ11callcudasubE8d_Volume
	.local	_ZZ11callcudasubE8d_Volume
	.comm	_ZZ11callcudasubE8d_Volume,8,8
	.type	_ZZ11callcudasubE6d_Sigt,@object # @_ZZ11callcudasubE6d_Sigt
	.local	_ZZ11callcudasubE6d_Sigt
	.comm	_ZZ11callcudasubE6d_Sigt,8,8
	.type	_ZZ11callcudasubE6d_A_fp,@object # @_ZZ11callcudasubE6d_A_fp
	.local	_ZZ11callcudasubE6d_A_fp
	.comm	_ZZ11callcudasubE6d_A_fp,8,8
	.type	_ZZ11callcudasubE6d_A_ez,@object # @_ZZ11callcudasubE6d_A_ez
	.local	_ZZ11callcudasubE6d_A_ez
	.comm	_ZZ11callcudasubE6d_A_ez,8,8
	.type	_ZZ11callcudasubE9d_Connect,@object # @_ZZ11callcudasubE9d_Connect
	.local	_ZZ11callcudasubE9d_Connect
	.comm	_ZZ11callcudasubE9d_Connect,8,8
	.type	_ZZ11callcudasubE6d_psic,@object # @_ZZ11callcudasubE6d_psic
	.local	_ZZ11callcudasubE6d_psic
	.comm	_ZZ11callcudasubE6d_psic,8,8
	.type	_ZZ11callcudasubE6d_psib,@object # @_ZZ11callcudasubE6d_psib
	.local	_ZZ11callcudasubE6d_psib
	.comm	_ZZ11callcudasubE6d_psib,8,8
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"max faces=%d\n"
	.size	.L.str, 14

	.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 _ZZ11callcudasubE12d_AngleOrder
	.addrsig_sym _ZZ11callcudasubE7d_omega
	.addrsig_sym _ZZ11callcudasubE7d_nextZ
	.addrsig_sym _ZZ11callcudasubE6d_next
	.addrsig_sym _ZZ11callcudasubE9d_nCorner
	.addrsig_sym _ZZ11callcudasubE9d_nCFaces
	.addrsig_sym _ZZ11callcudasubE4d_c0
	.addrsig_sym _ZZ11callcudasubE8d_STotal
	.addrsig_sym _ZZ11callcudasubE7d_STime
	.addrsig_sym _ZZ11callcudasubE9d_SigtInv
	.addrsig_sym _ZZ11callcudasubE8d_Volume
	.addrsig_sym _ZZ11callcudasubE6d_Sigt
	.addrsig_sym _ZZ11callcudasubE6d_A_fp
	.addrsig_sym _ZZ11callcudasubE6d_A_ez
	.addrsig_sym _ZZ11callcudasubE9d_Connect
	.addrsig_sym _ZZ11callcudasubE6d_psic
	.addrsig_sym _ZZ11callcudasubE6d_psib
	.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
