// Compile with:
// nvcc --std=c++11 fft_stream.cu -o fft_stream -lcufft

#include <iostream>

#include <cuda.h>
#include <cuda_runtime.h>

#include <cufft.h>

#include <chrono>
#include <iomanip>

using namespace std;
using namespace std::chrono;

// Print file name, line number, and error code when a CUDA error occurs.
#define check_cuda_errors(val)  __check_cuda_errors__ ( (val), #val, __FILE__, __LINE__ )

template <typename T>
inline void __check_cuda_errors__(T code, const char *func, const char *file, int line) {
    if (code) {
    std::cout << "CUDA error at "
          << file << ":" << line << std::endl
          << "error code: " << (unsigned int) code
          << " type: \""  << cudaGetErrorString(cudaGetLastError()) << "\"" << std::endl
          << "func: \"" << func << "\""
          << std::endl;
    cudaDeviceReset();
    exit(EXIT_FAILURE);
    }
}

int main(int argc, char *argv[]) {

    // Number of FFTs to compute.
    const int NUM_DATAs[] = {64*1024, 64*1024, 32*1024, 32*1024, 16*1024, 16*1024, 8*1024, 4*1024, 4*1024};

    // Length of each FFT.
    const int Ns[] = {128, 256, 512, 1024, 4096, 8192, 16636, 32768, 65536};

    // Number of GPU streams across which to distribute the FFTs.
    const int NUM_STREAMSs[] = { 32, 32, 32, 32, 32, 32, 32, 32, 32 };

    for (auto nn = 0; nn < 9; ++nn)
    {

       auto NUM_DATA = NUM_DATAs[nn];
       auto N = Ns[nn];
       auto NUM_STREAMS = NUM_STREAMSs[nn];

       // Allocate and initialize host input data.
       float2 **h_in = new float2 *[NUM_STREAMS];
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
           h_in[ii] = new float2[N];
           for (int jj = 0; jj < N; ++jj) {
               h_in[ii][jj].x = (float) 1.f;
               h_in[ii][jj].y = (float) 0.f;
           }
       }

       // Allocate and initialize host output data.
       float2 **h_out = new float2 *[NUM_STREAMS];
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
       h_out[ii] = new float2[N];
       for (int jj = 0; jj < N; ++jj) {
               h_out[ii][jj].x = 0.f;
               h_out[ii][jj].y = 0.f;
           }
       }

       // Pin host input and output memory for cudaMemcpyAsync.
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
           check_cuda_errors(cudaHostRegister(h_in[ii], N*sizeof(float2), cudaHostRegisterPortable));
           check_cuda_errors(cudaHostRegister(h_out[ii], N*sizeof(float2), cudaHostRegisterPortable));
       }

       // Allocate pointers to device input and output arrays.
       float2 **d_in = new float2 *[NUM_STREAMS];
       float2 **d_out = new float2 *[NUM_STREAMS];

       // Allocate intput and output arrays on device.
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
           check_cuda_errors(cudaMalloc((void**)&d_in[ii], N*sizeof(float2)));
           check_cuda_errors(cudaMalloc((void**)&d_out[ii], N*sizeof(float2)));
       }

       // Create CUDA streams.
       cudaStream_t streams[NUM_STREAMS];
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
           check_cuda_errors(cudaStreamCreate(&streams[ii]));
       }

       // Creates cuFFT plans and sets them in streams
       cufftHandle* plans = (cufftHandle*) malloc(sizeof(cufftHandle)*NUM_STREAMS);
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
           cufftPlan1d(&plans[ii], N, CUFFT_C2C, 1);
           cufftSetStream(plans[ii], streams[ii]);
       }

       steady_clock::time_point before = steady_clock::now();

       // Fill streams with async memcopies and FFTs.
       for (int ii = 0; ii < NUM_DATA; ii++) {
           int jj = ii % NUM_STREAMS;
           check_cuda_errors(cudaMemcpyAsync(d_in[jj], h_in[jj], N*sizeof(float2), cudaMemcpyHostToDevice, streams[jj]));
           cufftExecC2C(plans[jj], (cufftComplex*)d_in[jj], (cufftComplex*)d_out[jj], CUFFT_FORWARD);
           check_cuda_errors(cudaMemcpyAsync(h_out[jj], d_out[jj], N*sizeof(float2), cudaMemcpyDeviceToHost, streams[jj]));
       }

       // Wait for calculations to complete.
       for(int ii = 0; ii < NUM_STREAMS; ii++) {
           check_cuda_errors(cudaStreamSynchronize(streams[ii]));
       }

       steady_clock::time_point after = steady_clock::now();

       // Free memory and streams.
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
           check_cuda_errors(cudaHostUnregister(h_in[ii]));
           check_cuda_errors(cudaHostUnregister(h_out[ii]));
           check_cuda_errors(cudaFree(d_in[ii]));
           check_cuda_errors(cudaFree(d_out[ii]));
           delete[] h_in[ii];
           delete[] h_out[ii];
           check_cuda_errors(cudaStreamDestroy(streams[ii]));
       }

       delete plans;

       cudaDeviceReset();

       auto totalTime = duration<double>(after - before).count();
       auto timePer = totalTime / NUM_DATA;
       auto sampPerSec = N / timePer;

       cout << "===================================================" << endl;
       cout << "NUM_DATA:       " << NUM_DATA << endl;
       cout << "NUM_STREAMS:    " << NUM_STREAMS << endl;
       cout << "N:              " << N << endl;
       cout << "Total Time:     " << totalTime << endl;
       cout << "Time Per FFT:   " << timePer << endl;
       cout << "Samps Per Sec:  " << sampPerSec << endl;
    }

    return 0;
}

	.file	"tmpxft_00384c8a_00000000-6_fft_stream.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4044:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4044:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4067:
	.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
.LFE4067:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata._Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i.str1.1,"aMS",@progbits,1
.LC3:
	.string	"CUDA error at "
.LC4:
	.string	":"
.LC5:
	.string	"error code: "
.LC6:
	.string	" type: \""
.LC7:
	.string	"\""
.LC8:
	.string	"func: \""
	.section	.text._Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i,"axG",@progbits,_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i,comdat
	.weak	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	.type	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i, @function
_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i:
.LFB4409:
	.cfi_startproc
	endbr64
	testl	%edi, %edi
	je	.L4
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movq	%rdx, %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movl	%ecx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rsi, %rbp
	leaq	.LC3(%rip), %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	%edi, %ebx
	leaq	_ZSt4cout(%rip), %rdi
	pushq	%rax
	.cfi_def_cfa_offset 48
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%r13, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	.LC4(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%r12d, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC5(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	leaq	.LC6(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rbx
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	movq	%rbx, %rdi
	leaq	.LC7(%rip), %rbx
	movq	%rax, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rbx, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC8(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rbp, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rbx, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	call	cudaDeviceReset@PLT
	movl	$1, %edi
	call	exit@PLT
.L4:
	.cfi_def_cfa_offset 8
	.cfi_restore 3
	.cfi_restore 6
	.cfi_restore 12
	.cfi_restore 13
	ret
	.cfi_endproc
.LFE4409:
	.size	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i, .-_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC11:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/savagesmc/cuda_stuff/master/example/fft_stream.cu"
.LC12:
	.string	"cudaHostRegister(h_in[ii], N*sizeof(float2), cudaHostRegisterPortable)"
.LC13:
	.string	"cudaHostRegister(h_out[ii], N*sizeof(float2), cudaHostRegisterPortable)"
.LC14:
	.string	"cudaMalloc((void**)&d_in[ii], N*sizeof(float2))"
.LC15:
	.string	"cudaMalloc((void**)&d_out[ii], N*sizeof(float2))"
.LC16:
	.string	"cudaStreamCreate(&streams[ii])"
.LC17:
	.string	"cudaMemcpyAsync(d_in[jj], h_in[jj], N*sizeof(float2), cudaMemcpyHostToDevice, streams[jj])"
.LC18:
	.string	"cudaMemcpyAsync(h_out[jj], d_out[jj], N*sizeof(float2), cudaMemcpyDeviceToHost, streams[jj])"
.LC19:
	.string	"cudaStreamSynchronize(streams[ii])"
.LC20:
	.string	"cudaHostUnregister(h_in[ii])"
.LC21:
	.string	"cudaHostUnregister(h_out[ii])"
.LC22:
	.string	"cudaFree(d_in[ii])"
.LC23:
	.string	"cudaFree(d_out[ii])"
.LC24:
	.string	"cudaStreamDestroy(streams[ii])"
.LC26:
	.string	"==================================================="
.LC27:
	.string	"NUM_DATA:       "
.LC28:
	.string	"NUM_STREAMS:    "
.LC29:
	.string	"N:              "
.LC30:
	.string	"Total Time:     "
.LC31:
	.string	"Time Per FFT:   "
.LC32:
	.string	"Samps Per Sec:  "
	.section	.rodata
	.align 32
.LC0:
	.long	65536
	.long	65536
	.long	32768
	.long	32768
	.long	16384
	.long	16384
	.long	8192
	.long	4096
	.long	4096
	.align 32
.LC1:
	.long	128
	.long	256
	.long	512
	.long	1024
	.long	4096
	.long	8192
	.long	16636
	.long	32768
	.long	65536
	.align 32
.LC2:
	.long	32
	.long	32
	.long	32
	.long	32
	.long	32
	.long	32
	.long	32
	.long	32
	.long	32
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB4038:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	leaq	.LC0(%rip), %rsi
	movl	$9, %ecx
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	pushq	%r14
	leaq	-164(%rbp), %rdi
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$232, %rsp
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	.cfi_offset 13, -40
	.cfi_offset 12, -48
	.cfi_offset 3, -56
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	xorl	%eax, %eax
	rep movsl
	leaq	-128(%rbp), %rdi
	leaq	.LC1(%rip), %rsi
	movl	$9, %ecx
	rep movsl
	leaq	-92(%rbp), %rdi
	leaq	.LC2(%rip), %rsi
	movl	$9, %ecx
	rep movsl
	movq	%rax, -200(%rbp)
.L41:
	movq	-200(%rbp), %rax
	movq	%rsp, -240(%rbp)
	movl	-164(%rbp,%rax), %eax
	movl	%eax, -192(%rbp)
	movq	-200(%rbp), %rax
	movl	-128(%rbp,%rax), %eax
	movl	%eax, -188(%rbp)
	movq	-200(%rbp), %rax
	movslq	-92(%rbp,%rax), %r15
	movabsq	$1152921504606846975, %rax
	movq	%r15, %rbx
	cmpq	%r15, %rax
	jb	.L11
	leaq	0(,%r15,8), %rax
	xorl	%r12d, %r12d
	movq	%rax, %rdi
	movq	%rax, -208(%rbp)
	call	_Znam@PLT
	movslq	-188(%rbp), %r14
	movq	%rax, %r13
	leaq	0(,%r14,8), %rax
	movq	%rax, -184(%rbp)
	jmp	.L12
.L11:
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	jne	.L57
	call	__cxa_throw_bad_array_new_length@PLT
.L59:
	incq	%r12
.L12:
	cmpl	%r12d, %ebx
	jle	.L58
	movabsq	$1152921504606846975, %rax
	cmpq	%r14, %rax
	jb	.L11
	movq	-184(%rbp), %rdi
	call	_Znam@PLT
	movss	.LC9(%rip), %xmm0
	xorl	%edx, %edx
	movq	%rax, 0(%r13,%r12,8)
.L14:
	cmpl	%edx, -188(%rbp)
	jle	.L59
	movl	$0x00000000, 4(%rax,%rdx,8)
	movss	%xmm0, (%rax,%rdx,8)
	incq	%rdx
	jmp	.L14
.L58:
	movq	-208(%rbp), %rdi
	xorl	%r12d, %r12d
	call	_Znam@PLT
	movq	%rax, -216(%rbp)
.L17:
	cmpl	%r12d, %ebx
	jle	.L60
	movabsq	$1152921504606846975, %rax
	cmpq	%r14, %rax
	jb	.L11
	movq	-184(%rbp), %rdi
	call	_Znam@PLT
	movq	-216(%rbp), %rcx
	xorl	%edx, %edx
	movq	%rax, (%rcx,%r12,8)
.L18:
	cmpl	%edx, -188(%rbp)
	jle	.L61
	movl	$0x00000000, (%rax,%rdx,8)
	movl	$0x00000000, 4(%rax,%rdx,8)
	incq	%rdx
	jmp	.L18
.L61:
	incq	%r12
	jmp	.L17
.L60:
	xorl	%r12d, %r12d
	leaq	.LC11(%rip), %r14
.L21:
	cmpl	%r12d, %ebx
	jle	.L62
	movq	0(%r13,%r12,8), %rdi
	movq	-184(%rbp), %rsi
	movl	$1, %edx
	call	cudaHostRegister@PLT
	movl	$74, %ecx
	movq	%r14, %rdx
	leaq	.LC12(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	movq	-216(%rbp), %rax
	movl	$1, %edx
	movq	-184(%rbp), %rsi
	movq	(%rax,%r12,8), %rdi
	incq	%r12
	call	cudaHostRegister@PLT
	movl	$75, %ecx
	movq	%r14, %rdx
	leaq	.LC13(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	jmp	.L21
.L62:
	movq	-208(%rbp), %rdi
	xorl	%r12d, %r12d
	call	_Znam@PLT
	movq	-208(%rbp), %rdi
	movq	%rax, -224(%rbp)
	call	_Znam@PLT
	movq	-224(%rbp), %r9
	movq	%rax, -232(%rbp)
	movq	%rax, %r14
.L23:
	cmpl	%ebx, %r12d
	jge	.L63
	movq	-184(%rbp), %rsi
	movq	%r9, %rdi
	movq	%r9, -248(%rbp)
	incl	%r12d
	call	cudaMalloc@PLT
	movl	$84, %ecx
	leaq	.LC11(%rip), %rdx
	leaq	.LC14(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	movq	-184(%rbp), %rsi
	movq	%r14, %rdi
	addq	$8, %r14
	call	cudaMalloc@PLT
	movl	$85, %ecx
	leaq	.LC11(%rip), %rdx
	leaq	.LC15(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	movq	-248(%rbp), %r9
	addq	$8, %r9
	jmp	.L23
.L63:
	movq	-208(%rbp), %r12
	movq	%rsp, %rdx
	addq	$15, %r12
	movq	%r12, %rax
	andq	$-4096, %r12
	andq	$-16, %rax
	subq	%r12, %rdx
.L25:
	cmpq	%rdx, %rsp
	je	.L26
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L25
.L26:
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L27
	orq	$0, -8(%rsp,%rax)
.L27:
	movq	%rsp, %r14
	xorl	%r12d, %r12d
	movq	%r14, %r9
.L28:
	cmpl	%ebx, %r12d
	jge	.L64
	movq	%r9, %rdi
	movq	%r9, -208(%rbp)
	incl	%r12d
	call	cudaStreamCreate@PLT
	movl	$91, %ecx
	leaq	.LC11(%rip), %rdx
	leaq	.LC16(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	movq	-208(%rbp), %r9
	addq	$8, %r9
	jmp	.L28
.L64:
	leaq	0(,%r15,4), %rdi
	xorl	%r15d, %r15d
	call	malloc@PLT
	movq	%rax, -208(%rbp)
	movq	%rax, %r12
.L30:
	cmpl	%r15d, %ebx
	jle	.L65
	movl	-188(%rbp), %esi
	movq	%r12, %rdi
	movl	$1, %ecx
	movl	$41, %edx
	addq	$4, %r12
	call	cufftPlan1d@PLT
	movq	(%r14,%r15,8), %rsi
	movl	-4(%r12), %edi
	incq	%r15
	call	cufftSetStream@PLT
	jmp	.L30
.L65:
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	xorl	%r15d, %r15d
	movq	%rax, -248(%rbp)
.L32:
	movl	-192(%rbp), %eax
	cmpl	%eax, %r15d
	jge	.L66
	movl	%r15d, %eax
	movl	$1, %ecx
	incl	%r15d
	cltd
	idivl	%ebx
	movq	-224(%rbp), %rax
	movslq	%edx, %r12
	movq	-184(%rbp), %rdx
	leaq	0(,%r12,8), %r9
	movq	(%r14,%r12,8), %r8
	movq	0(%r13,%r12,8), %rsi
	leaq	(%rax,%r9), %r10
	movq	%r9, -256(%rbp)
	movq	(%r10), %rdi
	movq	%r10, -264(%rbp)
	call	cudaMemcpyAsync@PLT
	movl	$106, %ecx
	leaq	.LC11(%rip), %rdx
	leaq	.LC17(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	movq	-232(%rbp), %rax
	movq	-256(%rbp), %r9
	orl	$-1, %ecx
	movq	-264(%rbp), %r10
	addq	%rax, %r9
	movq	-208(%rbp), %rax
	movq	(%r10), %rsi
	movq	(%r9), %rdx
	movq	%r9, -256(%rbp)
	movl	(%rax,%r12,4), %edi
	call	cufftExecC2C@PLT
	movq	-256(%rbp), %r9
	movq	(%r14,%r12,8), %r8
	movl	$2, %ecx
	movq	-216(%rbp), %rax
	movq	-184(%rbp), %rdx
	movq	(%r9), %rsi
	movq	(%rax,%r12,8), %rdi
	call	cudaMemcpyAsync@PLT
	movl	$108, %ecx
	leaq	.LC11(%rip), %rdx
	leaq	.LC18(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	jmp	.L32
.L66:
	xorl	%r15d, %r15d
	leaq	.LC11(%rip), %r12
.L34:
	cmpl	%r15d, %ebx
	jle	.L67
	movq	(%r14,%r15,8), %rdi
	incq	%r15
	call	cudaStreamSynchronize@PLT
	movl	$113, %ecx
	movq	%r12, %rdx
	leaq	.LC19(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	jmp	.L34
.L67:
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	xorl	%r12d, %r12d
	movq	%rax, -184(%rbp)
.L36:
	cmpl	%r12d, %ebx
	jle	.L68
	movq	0(%r13,%r12,8), %r8
	movq	%r8, %rdi
	movq	%r8, -256(%rbp)
	call	cudaHostUnregister@PLT
	movl	$120, %ecx
	leaq	.LC11(%rip), %rdx
	leaq	.LC20(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	movq	-216(%rbp), %rax
	movq	(%rax,%r12,8), %r15
	movq	%r15, %rdi
	call	cudaHostUnregister@PLT
	movl	$121, %ecx
	leaq	.LC11(%rip), %rdx
	leaq	.LC21(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	movq	-224(%rbp), %rax
	movq	(%rax,%r12,8), %rdi
	call	cudaFree@PLT
	movl	$122, %ecx
	leaq	.LC11(%rip), %rdx
	leaq	.LC22(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	movq	-232(%rbp), %rax
	movq	(%rax,%r12,8), %rdi
	call	cudaFree@PLT
	movl	$123, %ecx
	leaq	.LC11(%rip), %rdx
	leaq	.LC23(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	movq	-256(%rbp), %r8
	testq	%r8, %r8
	je	.L37
	movq	%r8, %rdi
	call	_ZdaPv@PLT
.L37:
	testq	%r15, %r15
	je	.L38
	movq	%r15, %rdi
	call	_ZdaPv@PLT
.L38:
	movq	(%r14,%r12,8), %rdi
	incq	%r12
	call	cudaStreamDestroy@PLT
	movl	$126, %ecx
	leaq	.LC11(%rip), %rdx
	leaq	.LC24(%rip), %rsi
	movl	%eax, %edi
	call	_Z21__check_cuda_errors__I9cudaErrorEvT_PKcS3_i
	jmp	.L36
.L68:
	cmpq	$0, -208(%rbp)
	je	.L40
	movq	-208(%rbp), %rdi
	movl	$4, %esi
	call	_ZdlPvm@PLT
.L40:
	call	cudaDeviceReset@PLT
	movq	-248(%rbp), %rcx
	cvtsi2sdl	-192(%rbp), %xmm1
	leaq	_ZSt4cout(%rip), %r12
	movq	-184(%rbp), %rax
	leaq	.LC26(%rip), %rsi
	movq	%r12, %rdi
	subq	%rcx, %rax
	cvtsi2sdq	%rax, %xmm0
	divsd	.LC25(%rip), %xmm0
	movaps	%xmm0, %xmm2
	movsd	%xmm0, -208(%rbp)
	divsd	%xmm1, %xmm2
	cvtsi2sdl	-188(%rbp), %xmm1
	divsd	%xmm2, %xmm1
	movsd	%xmm2, -184(%rbp)
	movq	%xmm1, %r13
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC27(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	-192(%rbp), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC28(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC29(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	-188(%rbp), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC30(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	-208(%rbp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC31(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	-184(%rbp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC32(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%r13, %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	-240(%rbp), %rsp
	addq	$4, -200(%rbp)
	cmpq	$36, -200(%rbp)
	jne	.L41
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L42
.L57:
	call	__stack_chk_fail@PLT
.L42:
	leaq	-40(%rbp), %rsp
	xorl	%eax, %eax
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4038:
	.size	main, .-main
	.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
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC9:
	.long	1065353216
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC25:
	.long	0
	.long	1104006501
	.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:

// Compile with:
// nvcc --std=c++11 fft_stream.cu -o fft_stream -lcufft

#include <iostream>

#include <hip/hip_runtime.h>


#include <hipfft/hipfft.h>

#include <chrono>
#include <iomanip>

using namespace std;
using namespace std::chrono;

// Print file name, line number, and error code when a CUDA error occurs.
#define check_cuda_errors(val)  __check_cuda_errors__ ( (val), #val, __FILE__, __LINE__ )

template <typename T>
inline void __check_cuda_errors__(T code, const char *func, const char *file, int line) {
    if (code) {
    std::cout << "CUDA error at "
          << file << ":" << line << std::endl
          << "error code: " << (unsigned int) code
          << " type: \""  << hipGetErrorString(hipGetLastError()) << "\"" << std::endl
          << "func: \"" << func << "\""
          << std::endl;
    hipDeviceReset();
    exit(EXIT_FAILURE);
    }
}

int main(int argc, char *argv[]) {

    // Number of FFTs to compute.
    const int NUM_DATAs[] = {64*1024, 64*1024, 32*1024, 32*1024, 16*1024, 16*1024, 8*1024, 4*1024, 4*1024};

    // Length of each FFT.
    const int Ns[] = {128, 256, 512, 1024, 4096, 8192, 16636, 32768, 65536};

    // Number of GPU streams across which to distribute the FFTs.
    const int NUM_STREAMSs[] = { 32, 32, 32, 32, 32, 32, 32, 32, 32 };

    for (auto nn = 0; nn < 9; ++nn)
    {

       auto NUM_DATA = NUM_DATAs[nn];
       auto N = Ns[nn];
       auto NUM_STREAMS = NUM_STREAMSs[nn];

       // Allocate and initialize host input data.
       float2 **h_in = new float2 *[NUM_STREAMS];
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
           h_in[ii] = new float2[N];
           for (int jj = 0; jj < N; ++jj) {
               h_in[ii][jj].x = (float) 1.f;
               h_in[ii][jj].y = (float) 0.f;
           }
       }

       // Allocate and initialize host output data.
       float2 **h_out = new float2 *[NUM_STREAMS];
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
       h_out[ii] = new float2[N];
       for (int jj = 0; jj < N; ++jj) {
               h_out[ii][jj].x = 0.f;
               h_out[ii][jj].y = 0.f;
           }
       }

       // Pin host input and output memory for cudaMemcpyAsync.
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
           check_cuda_errors(hipHostRegister(h_in[ii], N*sizeof(float2), hipHostRegisterPortable));
           check_cuda_errors(hipHostRegister(h_out[ii], N*sizeof(float2), hipHostRegisterPortable));
       }

       // Allocate pointers to device input and output arrays.
       float2 **d_in = new float2 *[NUM_STREAMS];
       float2 **d_out = new float2 *[NUM_STREAMS];

       // Allocate intput and output arrays on device.
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
           check_cuda_errors(hipMalloc((void**)&d_in[ii], N*sizeof(float2)));
           check_cuda_errors(hipMalloc((void**)&d_out[ii], N*sizeof(float2)));
       }

       // Create CUDA streams.
       hipStream_t streams[NUM_STREAMS];
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
           check_cuda_errors(hipStreamCreate(&streams[ii]));
       }

       // Creates cuFFT plans and sets them in streams
       hipfftHandle* plans = (hipfftHandle*) malloc(sizeof(hipfftHandle)*NUM_STREAMS);
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
           hipfftPlan1d(&plans[ii], N, HIPFFT_C2C, 1);
           hipfftSetStream(plans[ii], streams[ii]);
       }

       steady_clock::time_point before = steady_clock::now();

       // Fill streams with async memcopies and FFTs.
       for (int ii = 0; ii < NUM_DATA; ii++) {
           int jj = ii % NUM_STREAMS;
           check_cuda_errors(hipMemcpyAsync(d_in[jj], h_in[jj], N*sizeof(float2), hipMemcpyHostToDevice, streams[jj]));
           hipfftExecC2C(plans[jj], (hipfftComplex*)d_in[jj], (hipfftComplex*)d_out[jj], HIPFFT_FORWARD);
           check_cuda_errors(hipMemcpyAsync(h_out[jj], d_out[jj], N*sizeof(float2), hipMemcpyDeviceToHost, streams[jj]));
       }

       // Wait for calculations to complete.
       for(int ii = 0; ii < NUM_STREAMS; ii++) {
           check_cuda_errors(hipStreamSynchronize(streams[ii]));
       }

       steady_clock::time_point after = steady_clock::now();

       // Free memory and streams.
       for (int ii = 0; ii < NUM_STREAMS; ii++) {
           check_cuda_errors(hipHostUnregister(h_in[ii]));
           check_cuda_errors(hipHostUnregister(h_out[ii]));
           check_cuda_errors(hipFree(d_in[ii]));
           check_cuda_errors(hipFree(d_out[ii]));
           delete[] h_in[ii];
           delete[] h_out[ii];
           check_cuda_errors(hipStreamDestroy(streams[ii]));
       }

       delete plans;

       hipDeviceReset();

       auto totalTime = duration<double>(after - before).count();
       auto timePer = totalTime / NUM_DATA;
       auto sampPerSec = N / timePer;

       cout << "===================================================" << endl;
       cout << "NUM_DATA:       " << NUM_DATA << endl;
       cout << "NUM_STREAMS:    " << NUM_STREAMS << endl;
       cout << "N:              " << N << endl;
       cout << "Total Time:     " << totalTime << endl;
       cout << "Time Per FFT:   " << timePer << endl;
       cout << "Samps Per Sec:  " << sampPerSec << endl;
    }

    return 0;
}

	.text
	.file	"fft_stream.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI0_0:
	.long	65536                           # 0x10000
	.long	65536                           # 0x10000
	.long	32768                           # 0x8000
	.long	32768                           # 0x8000
.LCPI0_1:
	.long	16384                           # 0x4000
	.long	16384                           # 0x4000
	.long	8192                            # 0x2000
	.long	4096                            # 0x1000
.LCPI0_2:
	.long	128                             # 0x80
	.long	256                             # 0x100
	.long	512                             # 0x200
	.long	1024                            # 0x400
.LCPI0_3:
	.long	4096                            # 0x1000
	.long	8192                            # 0x2000
	.long	16636                           # 0x40fc
	.long	32768                           # 0x8000
.LCPI0_4:
	.long	32                              # 0x20
	.long	32                              # 0x20
	.long	32                              # 0x20
	.long	32                              # 0x20
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI0_5:
	.quad	0x41cdcd6500000000              # double 1.0E+9
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset %rbp, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register %rbp
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$248, %rsp
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	movaps	.LCPI0_0(%rip), %xmm0           # xmm0 = [65536,65536,32768,32768]
	movaps	%xmm0, -288(%rbp)
	movaps	.LCPI0_1(%rip), %xmm0           # xmm0 = [16384,16384,8192,4096]
	movaps	%xmm0, -272(%rbp)
	movl	$4096, -256(%rbp)               # imm = 0x1000
	movaps	.LCPI0_2(%rip), %xmm0           # xmm0 = [128,256,512,1024]
	movaps	%xmm0, -240(%rbp)
	movaps	.LCPI0_3(%rip), %xmm0           # xmm0 = [4096,8192,16636,32768]
	movaps	%xmm0, -224(%rbp)
	movl	$65536, -208(%rbp)              # imm = 0x10000
	movaps	.LCPI0_4(%rip), %xmm0           # xmm0 = [32,32,32,32]
	movaps	%xmm0, -192(%rbp)
	movaps	%xmm0, -176(%rbp)
	movl	$32, -160(%rbp)
	xorl	%eax, %eax
	movq	$-1, %r14
.LBB0_1:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_8 Depth 2
                                        #       Child Loop BB0_10 Depth 3
                                        #     Child Loop BB0_17 Depth 2
                                        #     Child Loop BB0_14 Depth 2
                                        #     Child Loop BB0_23 Depth 2
                                        #     Child Loop BB0_34 Depth 2
                                        #     Child Loop BB0_36 Depth 2
                                        #     Child Loop BB0_27 Depth 2
                                        #     Child Loop BB0_30 Depth 2
                                        #     Child Loop BB0_43 Depth 2
	movl	-288(%rbp,%rax,4), %ecx
	movl	%ecx, -44(%rbp)                 # 4-byte Spill
	movl	-240(%rbp,%rax,4), %r12d
	movq	%rax, -152(%rbp)                # 8-byte Spill
	movslq	-192(%rbp,%rax,4), %rbx
	leaq	(,%rbx,8), %r15
	testq	%rbx, %rbx
	movq	%r15, -56(%rbp)                 # 8-byte Spill
	cmovsq	%r14, %r15
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, -88(%rbp)                 # 8-byte Spill
	movslq	%r12d, %rax
	movl	%ebx, %r13d
	testq	%rbx, %rbx
	movq	%r13, -64(%rbp)                 # 8-byte Spill
	movq	%rax, -96(%rbp)                 # 8-byte Spill
	jle	.LBB0_2
# %bb.7:                                # %.lr.ph150
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	%r15, -72(%rbp)                 # 8-byte Spill
	leaq	(,%rax,8), %r15
	testl	%eax, %eax
	movq	%r15, %r13
	cmovsq	%r14, %r13
	xorl	%ebx, %ebx
.LBB0_8:                                #   Parent Loop BB0_1 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB0_10 Depth 3
	movq	%r13, %rdi
	callq	_Znam
	movq	-88(%rbp), %rcx                 # 8-byte Reload
	movq	%rax, (%rcx,%rbx,8)
	testl	%r12d, %r12d
	jle	.LBB0_11
# %bb.9:                                # %.lr.ph.preheader
                                        #   in Loop: Header=BB0_8 Depth=2
	xorl	%ecx, %ecx
.LBB0_10:                               # %.lr.ph
                                        #   Parent Loop BB0_1 Depth=1
                                        #     Parent Loop BB0_8 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movq	$1065353216, (%rax,%rcx,8)      # imm = 0x3F800000
	incq	%rcx
	cmpq	%rcx, %r12
	jne	.LBB0_10
.LBB0_11:                               # %._crit_edge
                                        #   in Loop: Header=BB0_8 Depth=2
	incq	%rbx
	cmpq	-64(%rbp), %rbx                 # 8-byte Folded Reload
	jne	.LBB0_8
# %bb.5:                                # %._crit_edge151
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	-72(%rbp), %rdi                 # 8-byte Reload
	callq	_Znam
	movq	%rax, -80(%rbp)                 # 8-byte Spill
	movq	-64(%rbp), %r13                 # 8-byte Reload
	testl	%r13d, %r13d
	jle	.LBB0_6
# %bb.16:                               # %.lr.ph158
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	-96(%rbp), %rax                 # 8-byte Reload
	testl	%eax, %eax
	movq	%r15, %r13
	cmovsq	%r14, %r13
	movl	%eax, %ebx
	shlq	$3, %rbx
	xorl	%r14d, %r14d
.LBB0_17:                               #   Parent Loop BB0_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%r13, %rdi
	callq	_Znam
	movq	-80(%rbp), %rcx                 # 8-byte Reload
	movq	%rax, (%rcx,%r14,8)
	cmpl	$0, -96(%rbp)                   # 4-byte Folded Reload
	jle	.LBB0_19
# %bb.18:                               # %.lr.ph154.preheader
                                        #   in Loop: Header=BB0_17 Depth=2
	movq	%rax, %rdi
	xorl	%esi, %esi
	movq	%rbx, %rdx
	callq	memset@PLT
.LBB0_19:                               # %._crit_edge155
                                        #   in Loop: Header=BB0_17 Depth=2
	incq	%r14
	cmpq	%r14, -64(%rbp)                 # 8-byte Folded Reload
	jne	.LBB0_17
# %bb.12:                               # %.preheader146
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	-64(%rbp), %r13                 # 8-byte Reload
	testl	%r13d, %r13d
	jle	.LBB0_6
# %bb.13:                               # %.lr.ph160
                                        #   in Loop: Header=BB0_1 Depth=1
	xorl	%ebx, %ebx
	movq	-88(%rbp), %r14                 # 8-byte Reload
	movq	-80(%rbp), %r13                 # 8-byte Reload
.LBB0_14:                               #   Parent Loop BB0_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	(%r14,%rbx,8), %rdi
	movq	%r15, %rsi
	movl	$1, %edx
	callq	hipHostRegister
	movl	$.L.str, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$74, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	movq	(%r13,%rbx,8), %rdi
	movq	%r15, %rsi
	movl	$1, %edx
	callq	hipHostRegister
	movl	$.L.str.2, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$75, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	incq	%rbx
	cmpq	%rbx, -64(%rbp)                 # 8-byte Folded Reload
	jne	.LBB0_14
# %bb.15:                               #   in Loop: Header=BB0_1 Depth=1
	movq	-72(%rbp), %r15                 # 8-byte Reload
	movq	-64(%rbp), %r13                 # 8-byte Reload
	jmp	.LBB0_3
.LBB0_6:                                #   in Loop: Header=BB0_1 Depth=1
	movq	-72(%rbp), %r15                 # 8-byte Reload
	jmp	.LBB0_3
.LBB0_2:                                # %._crit_edge151.thread
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, -80(%rbp)                 # 8-byte Spill
.LBB0_3:                                # %._crit_edge161
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, %rbx
	movq	%r15, %rdi
	callq	_Znam
	testl	%r13d, %r13d
	movq	%rbx, -128(%rbp)                # 8-byte Spill
	movq	%rax, -136(%rbp)                # 8-byte Spill
	movq	%r12, -112(%rbp)                # 8-byte Spill
	jle	.LBB0_4
# %bb.22:                               # %.lr.ph164
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	%rax, %rcx
	movq	-96(%rbp), %rax                 # 8-byte Reload
	leaq	(,%rax,8), %r15
	movq	%r13, %r14
	movq	%rbx, %r13
	movq	%rcx, %rbx
.LBB0_23:                               #   Parent Loop BB0_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%r13, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.3, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$84, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	movq	%rbx, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.4, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$85, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	addq	$8, %rbx
	addq	$8, %r13
	decq	%r14
	jne	.LBB0_23
# %bb.20:                               # %._crit_edge165
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	%rsp, -120(%rbp)                # 8-byte Spill
	movq	%rsp, %r14
	movq	-64(%rbp), %r13                 # 8-byte Reload
	leaq	15(,%r13,8), %rax
	andq	$-16, %rax
	subq	%rax, %r14
	movq	%r14, %rsp
	testl	%r13d, %r13d
	jle	.LBB0_21
# %bb.33:                               # %.lr.ph168.preheader
                                        #   in Loop: Header=BB0_1 Depth=1
	leaq	(,%r13,8), %rbx
	xorl	%r15d, %r15d
.LBB0_34:                               # %.lr.ph168
                                        #   Parent Loop BB0_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	leaq	(%r14,%r15), %rdi
	callq	hipStreamCreate
	movl	$.L.str.5, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$91, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	addq	$8, %r15
	cmpq	%r15, %rbx
	jne	.LBB0_34
# %bb.24:                               # %._crit_edge169
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	-56(%rbp), %rdi                 # 8-byte Reload
	callq	malloc
	movq	%rax, -72(%rbp)                 # 8-byte Spill
	testl	%r13d, %r13d
	movq	%r14, -104(%rbp)                # 8-byte Spill
	jle	.LBB0_25
# %bb.35:                               # %.lr.ph172.preheader
                                        #   in Loop: Header=BB0_1 Depth=1
	xorl	%r15d, %r15d
	movq	-72(%rbp), %r13                 # 8-byte Reload
.LBB0_36:                               # %.lr.ph172
                                        #   Parent Loop BB0_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	leaq	(%r15,%r13), %rdi
	movl	%r12d, %esi
	movl	$41, %edx
	movl	$1, %ecx
	callq	hipfftPlan1d
	movq	(%r13,%r15), %rdi
	movq	(%r14,%r15), %rsi
	callq	hipfftSetStream
	addq	$8, %r15
	cmpq	%r15, %rbx
	jne	.LBB0_36
	jmp	.LBB0_25
.LBB0_4:                                # %._crit_edge165.thread
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	%rsp, -120(%rbp)                # 8-byte Spill
	movq	%rsp, %r14
	leaq	15(,%r13,8), %rax
	andq	$-16, %rax
	subq	%rax, %r14
	movq	%r14, %rsp
.LBB0_21:                               #   in Loop: Header=BB0_1 Depth=1
	movq	-56(%rbp), %rdi                 # 8-byte Reload
	movq	%r14, -104(%rbp)                # 8-byte Spill
	callq	malloc
	movq	%rax, -72(%rbp)                 # 8-byte Spill
.LBB0_25:                               #   in Loop: Header=BB0_1 Depth=1
	movl	-44(%rbp), %ebx                 # 4-byte Reload
	callq	_ZNSt6chrono3_V212steady_clock3nowEv
	movq	%rax, -144(%rbp)                # 8-byte Spill
	testl	%ebx, %ebx
	movq	-64(%rbp), %r15                 # 8-byte Reload
	movq	-88(%rbp), %rcx                 # 8-byte Reload
	movq	-128(%rbp), %r12                # 8-byte Reload
	movq	-104(%rbp), %r13                # 8-byte Reload
	jle	.LBB0_28
# %bb.26:                               # %.lr.ph176
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	-96(%rbp), %rax                 # 8-byte Reload
	shlq	$3, %rax
	movq	%rax, -56(%rbp)                 # 8-byte Spill
	xorl	%r14d, %r14d
.LBB0_27:                               #   Parent Loop BB0_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	%r14d, %eax
	cltd
	idivl	%r15d
	movl	%edx, %ebx
	movq	(%r12,%rbx,8), %rdi
	movq	(%rcx,%rbx,8), %rsi
	movq	(%r13,%rbx,8), %r8
	movq	-56(%rbp), %rdx                 # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpyAsync
	movl	$.L.str.6, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$106, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	movq	-72(%rbp), %rax                 # 8-byte Reload
	movq	(%rax,%rbx,8), %rdi
	movq	(%r12,%rbx,8), %rsi
	movq	-136(%rbp), %r15                # 8-byte Reload
	movq	(%r15,%rbx,8), %rdx
	movl	$-1, %ecx
	callq	hipfftExecC2C
	movq	-80(%rbp), %rax                 # 8-byte Reload
	movq	(%rax,%rbx,8), %rdi
	movq	(%r15,%rbx,8), %rsi
	movq	(%r13,%rbx,8), %r8
	movl	-44(%rbp), %ebx                 # 4-byte Reload
	movq	-56(%rbp), %rdx                 # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpyAsync
	movl	$.L.str.7, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$108, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	movq	-88(%rbp), %rcx                 # 8-byte Reload
	movq	-64(%rbp), %r15                 # 8-byte Reload
	incl	%r14d
	cmpl	%r14d, %ebx
	jne	.LBB0_27
.LBB0_28:                               # %.preheader
                                        #   in Loop: Header=BB0_1 Depth=1
	testl	%r15d, %r15d
	jle	.LBB0_37
# %bb.29:                               # %.lr.ph178.preheader
                                        #   in Loop: Header=BB0_1 Depth=1
	xorl	%ebx, %ebx
	movq	-104(%rbp), %r14                # 8-byte Reload
.LBB0_30:                               # %.lr.ph178
                                        #   Parent Loop BB0_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	(%r14,%rbx,8), %rdi
	callq	hipStreamSynchronize
	movl	$.L.str.8, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$113, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	incq	%rbx
	cmpq	%rbx, %r15
	jne	.LBB0_30
# %bb.31:                               # %._crit_edge179
                                        #   in Loop: Header=BB0_1 Depth=1
	callq	_ZNSt6chrono3_V212steady_clock3nowEv
	movq	%rax, -56(%rbp)                 # 8-byte Spill
	testl	%r15d, %r15d
	movq	-80(%rbp), %r13                 # 8-byte Reload
	movq	-136(%rbp), %r14                # 8-byte Reload
	movq	-112(%rbp), %r12                # 8-byte Reload
	jle	.LBB0_32
# %bb.42:                               # %.lr.ph182.preheader
                                        #   in Loop: Header=BB0_1 Depth=1
	xorl	%ebx, %ebx
.LBB0_43:                               # %.lr.ph182
                                        #   Parent Loop BB0_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	-88(%rbp), %rax                 # 8-byte Reload
	movq	(%rax,%rbx,8), %r15
	movq	%r15, %rdi
	callq	hipHostUnregister
	movl	$.L.str.9, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$120, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	movq	(%r13,%rbx,8), %r13
	movq	%r13, %rdi
	callq	hipHostUnregister
	movl	$.L.str.10, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$121, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	movq	-128(%rbp), %rax                # 8-byte Reload
	movq	(%rax,%rbx,8), %rdi
	callq	hipFree
	movl	$.L.str.11, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$122, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	movq	(%r14,%rbx,8), %rdi
	callq	hipFree
	movl	$.L.str.12, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$123, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	testq	%r15, %r15
	je	.LBB0_45
# %bb.44:                               #   in Loop: Header=BB0_43 Depth=2
	movq	%r15, %rdi
	callq	_ZdaPv
.LBB0_45:                               #   in Loop: Header=BB0_43 Depth=2
	testq	%r13, %r13
	je	.LBB0_47
# %bb.46:                               #   in Loop: Header=BB0_43 Depth=2
	movq	%r13, %rdi
	callq	_ZdaPv
.LBB0_47:                               #   in Loop: Header=BB0_43 Depth=2
	movq	-104(%rbp), %rax                # 8-byte Reload
	movq	(%rax,%rbx,8), %rdi
	callq	hipStreamDestroy
	movl	$.L.str.13, %esi
	movl	$.L.str.1, %edx
	movl	%eax, %edi
	movl	$126, %ecx
	callq	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	incq	%rbx
	movq	-64(%rbp), %r15                 # 8-byte Reload
	cmpq	%rbx, %r15
	movq	-80(%rbp), %r13                 # 8-byte Reload
	jne	.LBB0_43
.LBB0_32:                               #   in Loop: Header=BB0_1 Depth=1
	movl	-44(%rbp), %ebx                 # 4-byte Reload
	movq	-56(%rbp), %r13                 # 8-byte Reload
	jmp	.LBB0_38
.LBB0_37:                               # %._crit_edge179.thread
                                        #   in Loop: Header=BB0_1 Depth=1
	callq	_ZNSt6chrono3_V212steady_clock3nowEv
	movq	%rax, %r13
	movq	-112(%rbp), %r12                # 8-byte Reload
.LBB0_38:                               # %._crit_edge183
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	-72(%rbp), %rdi                 # 8-byte Reload
	testq	%rdi, %rdi
	movq	$-1, %r14
	je	.LBB0_40
# %bb.39:                               #   in Loop: Header=BB0_1 Depth=1
	callq	_ZdlPv
.LBB0_40:                               #   in Loop: Header=BB0_1 Depth=1
	callq	hipDeviceReset
	subq	-144(%rbp), %r13                # 8-byte Folded Reload
	cvtsi2sd	%r13, %xmm1
	divsd	.LCPI0_5(%rip), %xmm1
	movsd	%xmm1, -56(%rbp)                # 8-byte Spill
	cvtsi2sd	%ebx, %xmm0
	divsd	%xmm0, %xmm1
	movsd	%xmm1, -80(%rbp)                # 8-byte Spill
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r12d, %xmm0
	divsd	%xmm1, %xmm0
	movsd	%xmm0, -88(%rbp)                # 8-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.14, %esi
	movl	$51, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	movl	$_ZSt4cout, %eax
	addq	%rax, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.15, %esi
	movl	$16, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%ebx, %esi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.16, %esi
	movl	$16, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%r15d, %esi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.17, %esi
	movl	$16, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movq	-96(%rbp), %rsi                 # 8-byte Reload
                                        # kill: def $esi killed $esi killed $rsi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.18, %esi
	movl	$16, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	-56(%rbp), %xmm0                # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.19, %esi
	movl	$16, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	-80(%rbp), %xmm0                # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.20, %esi
	movl	$16, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	-88(%rbp), %xmm0                # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	-120(%rbp), %rsp                # 8-byte Reload
	movq	-152(%rbp), %rax                # 8-byte Reload
	incq	%rax
	cmpq	$9, %rax
	jne	.LBB0_1
# %bb.41:
	xorl	%eax, %eax
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa %rsp, 8
	retq
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.section	.text._Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i,"axG",@progbits,_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i,comdat
	.weak	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i # -- Begin function _Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	.type	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i,@function
_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i: # @_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	.cfi_startproc
# %bb.0:
	testl	%edi, %edi
	jne	.LBB1_2
# %bb.1:
	retq
.LBB1_2:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%ecx, %r14d
	movq	%rdx, %r15
	movq	%rsi, %rbx
	movl	%edi, %ebp
	movl	$_ZSt4cout, %edi
	movl	$.L.str.21, %esi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movq	%rax, %rdi
	movq	%r15, %rsi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movl	$.L.str.22, %esi
	movq	%rax, %rdi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movq	%rax, %rdi
	movl	%r14d, %esi
	callq	_ZNSolsEi
	movq	%rax, %rdi
	callq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_
	movl	$.L.str.23, %esi
	movq	%rax, %rdi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movq	%rax, %rdi
	movl	%ebp, %esi
	callq	_ZNSolsEj
	movl	$.L.str.24, %esi
	movq	%rax, %rdi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movq	%rax, %r14
	callq	hipGetLastError
	movl	%eax, %edi
	callq	hipGetErrorString
	movq	%r14, %rdi
	movq	%rax, %rsi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movl	$.L.str.25, %esi
	movq	%rax, %rdi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movq	%rax, %rdi
	callq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_
	movl	$.L.str.26, %esi
	movq	%rax, %rdi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movq	%rax, %rdi
	movq	%rbx, %rsi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movl	$.L.str.25, %esi
	movq	%rax, %rdi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movq	%rax, %rdi
	callq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_
	callq	hipDeviceReset
	movl	$1, %edi
	callq	exit
.Lfunc_end1:
	.size	_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i, .Lfunc_end1-_Z21__check_cuda_errors__I10hipError_tEvT_PKcS3_i
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"hipHostRegister(h_in[ii], N*sizeof(float2), hipHostRegisterPortable)"
	.size	.L.str, 69

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"/home/ubuntu/Datasets/stackv2/train-structured-repos-hip-sm89/savagesmc/cuda_stuff/master/example/fft_stream.hip"
	.size	.L.str.1, 113

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"hipHostRegister(h_out[ii], N*sizeof(float2), hipHostRegisterPortable)"
	.size	.L.str.2, 70

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"hipMalloc((void**)&d_in[ii], N*sizeof(float2))"
	.size	.L.str.3, 47

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"hipMalloc((void**)&d_out[ii], N*sizeof(float2))"
	.size	.L.str.4, 48

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"hipStreamCreate(&streams[ii])"
	.size	.L.str.5, 30

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"hipMemcpyAsync(d_in[jj], h_in[jj], N*sizeof(float2), hipMemcpyHostToDevice, streams[jj])"
	.size	.L.str.6, 89

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"hipMemcpyAsync(h_out[jj], d_out[jj], N*sizeof(float2), hipMemcpyDeviceToHost, streams[jj])"
	.size	.L.str.7, 91

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"hipStreamSynchronize(streams[ii])"
	.size	.L.str.8, 34

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"hipHostUnregister(h_in[ii])"
	.size	.L.str.9, 28

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"hipHostUnregister(h_out[ii])"
	.size	.L.str.10, 29

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"hipFree(d_in[ii])"
	.size	.L.str.11, 18

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"hipFree(d_out[ii])"
	.size	.L.str.12, 19

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"hipStreamDestroy(streams[ii])"
	.size	.L.str.13, 30

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"==================================================="
	.size	.L.str.14, 52

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	"NUM_DATA:       "
	.size	.L.str.15, 17

	.type	.L.str.16,@object               # @.str.16
.L.str.16:
	.asciz	"NUM_STREAMS:    "
	.size	.L.str.16, 17

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"N:              "
	.size	.L.str.17, 17

	.type	.L.str.18,@object               # @.str.18
.L.str.18:
	.asciz	"Total Time:     "
	.size	.L.str.18, 17

	.type	.L.str.19,@object               # @.str.19
.L.str.19:
	.asciz	"Time Per FFT:   "
	.size	.L.str.19, 17

	.type	.L.str.20,@object               # @.str.20
.L.str.20:
	.asciz	"Samps Per Sec:  "
	.size	.L.str.20, 17

	.type	.L.str.21,@object               # @.str.21
.L.str.21:
	.asciz	"CUDA error at "
	.size	.L.str.21, 15

	.type	.L.str.22,@object               # @.str.22
.L.str.22:
	.asciz	":"
	.size	.L.str.22, 2

	.type	.L.str.23,@object               # @.str.23
.L.str.23:
	.asciz	"error code: "
	.size	.L.str.23, 13

	.type	.L.str.24,@object               # @.str.24
.L.str.24:
	.asciz	" type: \""
	.size	.L.str.24, 9

	.type	.L.str.25,@object               # @.str.25
.L.str.25:
	.asciz	"\""
	.size	.L.str.25, 2

	.type	.L.str.26,@object               # @.str.26
.L.str.26:
	.asciz	"func: \""
	.size	.L.str.26, 8

	.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 _ZSt4cout
	.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
