//general parts
#include <stdio.h>
#include <vector>
#include <memory>
#include <string.h>
#include <chrono>
#include <thread>
#include <iostream>
#ifndef __STDC_FORMAT_MACROS
#define __STDC_FORMAT_MACROS
#endif
#include <inttypes.h>

//CUDA parts
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cufft.h>

#define GROUP 1


void sample_6_benchmark_cuFFT_single_r2c(bool file_output, FILE* output, int device_id)
{
	if (file_output)
		fprintf(output, "6 - cuFFT FFT + iFFT R2C/C2R multidimensional benchmark in single precision\n");
	printf("6 - cuFFT FFT + iFFT R2C/C2R multidimensional benchmark in single precision\n");
	cudaSetDevice(device_id);
	const int num_benchmark_samples = 24;
	const int num_runs = 3;
	//printf("First %" PRIu64 " runs are a warmup\n", num_runs);
	uint64_t benchmark_dimensions[num_benchmark_samples][4] = { {1024, 1024, 1, 2}, {64, 64, 1, 2}, {256, 256, 1, 2}, {1024, 256, 1, 2}, {512, 512, 1, 2}, {1024, 1024, 1, 2},  {4096, 256, 1, 2}, {2048, 1024, 1, 2},{4096, 2048, 1, 2}, {4096, 4096, 1, 2}, {720, 480, 1, 2},{1280, 720, 1, 2},{1920, 1080, 1, 2}, {2560, 1440, 1, 2},{3840, 2160, 1, 2},
																{32, 32, 32, 3}, {64, 64, 64, 3}, {256, 256, 32, 3},  {1024, 256, 32, 3},  {256, 256, 256, 3}, {2048, 1024, 8, 3},  {512, 512, 128, 3}, {2048, 256, 256, 3}, {4096, 512, 8, 3} };
	
	double benchmark_result[2] = { 0,0 };//averaged result = sum(system_size/iteration_time)/num_benchmark_samples
	cufftReal* inputC = (cufftReal*)malloc((uint64_t)sizeof(cufftReal)*pow(2, 27));
	for (uint64_t i = 0; i < pow(2, 27); i++) {
		inputC[i] = 2 * ((float)rand()) / RAND_MAX - 1.0;
	}
	for (int n = 0; n < num_benchmark_samples; n++) {
		double run_time[num_runs][2];
		for (int r = 0; r < num_runs; r++) {
			cufftHandle planR2C;
			cufftHandle planC2R;
			cufftReal* dataR;
			cufftComplex* dataC;

			uint64_t dims[3] = { benchmark_dimensions[n][0] , benchmark_dimensions[n][1] ,benchmark_dimensions[n][2] };

			cudaMalloc((void**)&dataR, sizeof(cufftComplex) * (dims[0] / 2 + 1) * dims[1] * dims[2]);
			cudaMalloc((void**)&dataC, sizeof(cufftComplex) * (dims[0] / 2 + 1) * dims[1] * dims[2]);

			cudaMemcpy(dataR, inputC, sizeof(cufftReal) * dims[0] * dims[1] * dims[2], cudaMemcpyHostToDevice);
			if (cudaGetLastError() != cudaSuccess) {
				fprintf(stderr, "Cuda error: Failed to allocate\n");
				return;
			}
			switch (benchmark_dimensions[n][3]) {
			case 1:
				cufftPlan1d(&planR2C, dims[0], CUFFT_R2C, 1);
				cufftPlan1d(&planC2R, dims[0], CUFFT_C2R, 1);
				break;
			case 2:
				cufftPlan2d(&planR2C, dims[1], dims[0], CUFFT_R2C);
				cufftPlan2d(&planC2R, dims[1], dims[0], CUFFT_C2R);
				break;
			case 3:
				cufftPlan3d(&planR2C, dims[2], dims[1], dims[0], CUFFT_R2C);
				cufftPlan3d(&planC2R, dims[2], dims[1], dims[0], CUFFT_C2R);
				break;
			}

			float totTime = 0;
			uint64_t cuBufferSize = sizeof(float) * 2 * (dims[0]/2+1) * dims[1] * dims[2];
			uint64_t num_iter = ((4096 * 1024.0 * 1024.0) / cuBufferSize > 1000) ? 1000 : (4096 * 1024.0 * 1024.0) / cuBufferSize;
			if (num_iter == 0) num_iter = 1;
			std::chrono::steady_clock::time_point timeSubmit = std::chrono::steady_clock::now();
			for (int i = 0; i < num_iter; i++) {

				cufftExecR2C(planR2C, dataR, dataC);
				cufftExecC2R(planC2R, dataC, dataR);
			}
			cudaDeviceSynchronize();
			std::chrono::steady_clock::time_point timeEnd = std::chrono::steady_clock::now();
			totTime = (std::chrono::duration_cast<std::chrono::microseconds>(timeEnd - timeSubmit).count() * 0.001) / num_iter;
			run_time[r][0] = totTime;
			if (n > 0) {
				if (r == num_runs - 1) {
					double std_error = 0;
					double avg_time = 0;
					for (uint64_t t = 0; t < num_runs; t++) {
						avg_time += run_time[t][0];
					}
					avg_time /= num_runs;
					for (uint64_t t = 0; t < num_runs; t++) {
						std_error += (run_time[t][0] - avg_time) * (run_time[t][0] - avg_time);
					}
					std_error = sqrt(std_error / num_runs);
					if (file_output)
						fprintf(output, "cuFFT System: %" PRIu64 "x%" PRIu64 "x%" PRIu64 " Buffer: %" PRIu64 " MB avg_time_per_step: %0.3f ms std_error: %0.3f num_iter: %" PRIu64 " benchmark: %" PRIu64 "\n", benchmark_dimensions[n][0], benchmark_dimensions[n][1], benchmark_dimensions[n][2], cuBufferSize / 1024 / 1024, avg_time, std_error, num_iter, (uint64_t)(((double)cuBufferSize / 1024) / avg_time));

					printf("cuFFT System: %" PRIu64 "x%" PRIu64 "x%" PRIu64 " Buffer: %" PRIu64 " MB avg_time_per_step: %0.3f ms std_error: %0.3f num_iter: %" PRIu64 " benchmark: %" PRIu64 "\n", benchmark_dimensions[n][0], benchmark_dimensions[n][1], benchmark_dimensions[n][2], cuBufferSize / 1024 / 1024, avg_time, std_error, num_iter, (uint64_t)(((double)cuBufferSize / 1024) / avg_time));
					benchmark_result[0] += ((double)cuBufferSize / 1024) / avg_time;
				}

			}
			cufftDestroy(planR2C);
			cufftDestroy(planC2R);
			cudaFree(dataC);
			cudaFree(dataR);
			cudaDeviceSynchronize();
			//cufftComplex* output_cuFFT = (cufftComplex*)(malloc(sizeof(cufftComplex) * dims[0] * dims[1] * dims[2]));
			//cudaMemcpy(output_cuFFT, dataC, sizeof(cufftComplex) * dims[0] * dims[1] * dims[2], cudaMemcpyDeviceToHost);
			//cudaDeviceSynchronize();
			

		}
	}
	free(inputC);
	benchmark_result[0] /= (num_benchmark_samples - 1);
	if (file_output)
		fprintf(output, "Benchmark score cuFFT: %" PRIu64 "\n", (uint64_t)(benchmark_result[0]));
	printf("Benchmark score cuFFT: %" PRIu64 "\n", (uint64_t)(benchmark_result[0]));

}

	.file	"tmpxft_00364622_00000000-6_sample_6_benchmark_cuFFT_single_r2c.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4751:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4751:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC2:
	.string	"6 - cuFFT FFT + iFFT R2C/C2R multidimensional benchmark in single precision\n"
.LC6:
	.string	"Cuda error: Failed to allocate\n"
.LC12:
	.string	"cuFFT System: %lux%lux%lu Buffer: %lu MB avg_time_per_step: %0.3f ms std_error: %0.3f num_iter: %lu benchmark: %lu\n"
.LC15:
	.string	"Benchmark score cuFFT: %lu\n"
	.section	.rodata
	.align 32
.LC0:
	.quad	1024
	.quad	1024
	.quad	1
	.quad	2
	.quad	64
	.quad	64
	.quad	1
	.quad	2
	.quad	256
	.quad	256
	.quad	1
	.quad	2
	.quad	1024
	.quad	256
	.quad	1
	.quad	2
	.quad	512
	.quad	512
	.quad	1
	.quad	2
	.quad	1024
	.quad	1024
	.quad	1
	.quad	2
	.quad	4096
	.quad	256
	.quad	1
	.quad	2
	.quad	2048
	.quad	1024
	.quad	1
	.quad	2
	.quad	4096
	.quad	2048
	.quad	1
	.quad	2
	.quad	4096
	.quad	4096
	.quad	1
	.quad	2
	.quad	720
	.quad	480
	.quad	1
	.quad	2
	.quad	1280
	.quad	720
	.quad	1
	.quad	2
	.quad	1920
	.quad	1080
	.quad	1
	.quad	2
	.quad	2560
	.quad	1440
	.quad	1
	.quad	2
	.quad	3840
	.quad	2160
	.quad	1
	.quad	2
	.quad	32
	.quad	32
	.quad	32
	.quad	3
	.quad	64
	.quad	64
	.quad	64
	.quad	3
	.quad	256
	.quad	256
	.quad	32
	.quad	3
	.quad	1024
	.quad	256
	.quad	32
	.quad	3
	.quad	256
	.quad	256
	.quad	256
	.quad	3
	.quad	2048
	.quad	1024
	.quad	8
	.quad	3
	.quad	512
	.quad	512
	.quad	128
	.quad	3
	.quad	2048
	.quad	256
	.quad	256
	.quad	3
	.quad	4096
	.quad	512
	.quad	8
	.quad	3
	.text
	.globl	_Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi
	.type	_Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi, @function
_Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi:
.LFB4745:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	pushq	%r14
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	movl	%edi, %r14d
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	.cfi_offset 13, -40
	.cfi_offset 12, -48
	.cfi_offset 3, -56
	movl	%edx, %ebx
	subq	$968, %rsp
	movq	%rsi, -984(%rbp)
	movb	%dil, -973(%rbp)
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	testb	%dil, %dil
	je	.L3
	movq	-984(%rbp), %rdi
	leaq	.LC2(%rip), %rdx
	movl	$2, %esi
	call	__fprintf_chk@PLT
.L3:
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	%ebx, %edi
	xorl	%ebx, %ebx
	call	cudaSetDevice@PLT
	leaq	-824(%rbp), %rdi
	movl	$192, %ecx
	leaq	.LC0(%rip), %rsi
	rep movsl
	movl	$536870912, %edi
	call	malloc@PLT
	movq	%rax, -968(%rbp)
.L4:
	cvtsi2sdq	%rbx, %xmm0
	movsd	.LC5(%rip), %xmm6
	comisd	%xmm0, %xmm6
	jbe	.L48
	call	rand@PLT
	cvtsi2ssl	%eax, %xmm0
	movq	-968(%rbp), %rax
	addss	%xmm0, %xmm0
	mulss	.LC3(%rip), %xmm0
	subss	.LC4(%rip), %xmm0
	movss	%xmm0, (%rax,%rbx,4)
	incq	%rbx
	jmp	.L4
.L48:
	leaq	-824(%rbp), %rax
	xorl	%edi, %edi
	movq	$0x000000000, -960(%rbp)
	movq	%rax, -936(%rbp)
	movl	%edi, -972(%rbp)
.L6:
	xorl	%eax, %eax
	movq	%rax, -952(%rbp)
.L24:
	movq	-936(%rbp), %rax
	leaq	-896(%rbp), %rdi
	movq	(%rax), %rax
	movq	%rax, -920(%rbp)
	movq	-936(%rbp), %rax
	movq	-920(%rbp), %rbx
	movq	8(%rax), %r12
	movq	16(%rax), %rax
	shrq	%rbx
	movq	%rax, -928(%rbp)
	imulq	%r12, %rax
	incq	%rbx
	imulq	%rax, %rbx
	salq	$3, %rbx
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-888(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movq	-920(%rbp), %rdx
	movq	-928(%rbp), %rax
	movl	$1, %ecx
	movq	-968(%rbp), %rsi
	movq	-896(%rbp), %rdi
	imulq	%r12, %rdx
	imulq	%rax, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	je	.L7
	movq	stderr(%rip), %rdi
	xorl	%eax, %eax
	movl	$2, %esi
	leaq	.LC6(%rip), %rdx
	call	__fprintf_chk@PLT
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L32
	jmp	.L45
.L7:
	movq	-936(%rbp), %rax
	movq	24(%rax), %rax
	cmpq	$2, %rax
	je	.L9
	cmpq	$3, %rax
	je	.L10
	decq	%rax
	jne	.L11
	movl	-920(%rbp), %r13d
	movl	$1, %ecx
	movl	$42, %edx
	leaq	-904(%rbp), %rdi
	movl	%r13d, %esi
	call	cufftPlan1d@PLT
	leaq	-900(%rbp), %rdi
	movl	$1, %ecx
	movl	%r13d, %esi
	movl	$44, %edx
	call	cufftPlan1d@PLT
	jmp	.L11
.L9:
	movl	-920(%rbp), %r15d
	leaq	-904(%rbp), %rdi
	movl	%r12d, %esi
	movl	$42, %ecx
	movl	%r15d, %edx
	call	cufftPlan2d@PLT
	movl	$44, %ecx
	movl	%r15d, %edx
	movl	%r12d, %esi
	leaq	-900(%rbp), %rdi
	call	cufftPlan2d@PLT
	jmp	.L11
.L10:
	movl	-920(%rbp), %edx
	movl	-928(%rbp), %r13d
	leaq	-904(%rbp), %rdi
	movl	$42, %r8d
	movl	%edx, %ecx
	movl	%edx, -944(%rbp)
	movl	%r13d, %esi
	movl	%r12d, %edx
	call	cufftPlan3d@PLT
	movl	-944(%rbp), %ecx
	movl	%r12d, %edx
	movl	%r13d, %esi
	leaq	-900(%rbp), %rdi
	movl	$44, %r8d
	call	cufftPlan3d@PLT
.L11:
	testq	%rbx, %rbx
	js	.L12
	cvtsi2sdq	%rbx, %xmm7
	movsd	%xmm7, -944(%rbp)
	jmp	.L13
.L12:
	movq	%rbx, %rax
	shrq	%rax
	cvtsi2sdq	%rax, %xmm0
	addsd	%xmm0, %xmm0
	movsd	%xmm0, -944(%rbp)
.L13:
	movsd	.LC7(%rip), %xmm0
	divsd	-944(%rbp), %xmm0
	comisd	.LC8(%rip), %xmm0
	movl	$1000, %r15d
	ja	.L14
	movsd	.LC9(%rip), %xmm1
	comisd	%xmm1, %xmm0
	jnb	.L15
	cvttsd2siq	%xmm0, %r15
	jmp	.L16
.L15:
	subsd	%xmm1, %xmm0
	cvttsd2siq	%xmm0, %r15
	btcq	$63, %r15
.L16:
	testq	%r15, %r15
	movl	$1, %eax
	cmove	%rax, %r15
.L14:
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	xorl	%r13d, %r13d
	movq	%rax, -992(%rbp)
.L17:
	movq	-888(%rbp), %rdx
	movq	-896(%rbp), %rsi
	incq	%r13
	movl	-904(%rbp), %edi
	call	cufftExecR2C@PLT
	movq	-896(%rbp), %rdx
	movq	-888(%rbp), %rsi
	movl	-900(%rbp), %edi
	call	cufftExecC2R@PLT
	cmpq	%r13, %r15
	jne	.L17
	call	cudaDeviceSynchronize@PLT
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	movq	-992(%rbp), %rsi
	cvtsi2sdq	%r13, %xmm1
	movq	-952(%rbp), %rcx
	subq	%rsi, %rax
	movl	$1000, %esi
	salq	$4, %rcx
	cqto
	idivq	%rsi
	cmpl	$2, -952(%rbp)
	cvtsi2sdq	%rax, %xmm0
	mulsd	.LC10(%rip), %xmm0
	divsd	%xmm1, %xmm0
	cvtsd2ss	%xmm0, %xmm0
	cvtss2sd	%xmm0, %xmm0
	movsd	%xmm0, -880(%rbp,%rcx)
	jne	.L18
	cmpl	$0, -972(%rbp)
	je	.L18
	movsd	-880(%rbp), %xmm3
	xorps	%xmm5, %xmm5
	movsd	-864(%rbp), %xmm0
	shrq	$20, %rbx
	movsd	-848(%rbp), %xmm1
	movsd	.LC11(%rip), %xmm4
	movaps	%xmm3, %xmm2
	addsd	%xmm5, %xmm2
	addsd	%xmm0, %xmm2
	addsd	%xmm1, %xmm2
	divsd	%xmm4, %xmm2
	subsd	%xmm2, %xmm3
	subsd	%xmm2, %xmm0
	movsd	%xmm2, -992(%rbp)
	subsd	%xmm2, %xmm1
	mulsd	%xmm3, %xmm3
	mulsd	%xmm0, %xmm0
	mulsd	%xmm1, %xmm1
	addsd	%xmm5, %xmm3
	addsd	%xmm3, %xmm0
	addsd	%xmm1, %xmm0
	divsd	%xmm4, %xmm0
	call	sqrt@PLT
	cmpb	$0, -973(%rbp)
	movsd	.LC9(%rip), %xmm4
	movsd	-992(%rbp), %xmm2
	movaps	%xmm0, %xmm1
	je	.L19
	movsd	.LC13(%rip), %xmm0
	pushq	%rsi
	mulsd	-944(%rbp), %xmm0
	divsd	%xmm2, %xmm0
	comisd	%xmm4, %xmm0
	jnb	.L20
	cvttsd2siq	%xmm0, %rax
	jmp	.L21
.L20:
	subsd	%xmm4, %xmm0
	cvttsd2siq	%xmm0, %rax
	btcq	$63, %rax
.L21:
	pushq	%rax
	movaps	%xmm2, %xmm0
	movq	%r12, %r8
	leaq	.LC12(%rip), %rdx
	pushq	%r13
	movl	$2, %esi
	movb	$2, %al
	pushq	%rbx
	movq	-928(%rbp), %r9
	movq	-920(%rbp), %rcx
	movq	-984(%rbp), %rdi
	movsd	%xmm1, -1000(%rbp)
	movsd	%xmm2, -992(%rbp)
	call	__fprintf_chk@PLT
	movsd	.LC9(%rip), %xmm4
	addq	$32, %rsp
	movsd	-1000(%rbp), %xmm1
	movsd	-992(%rbp), %xmm2
.L19:
	movsd	.LC13(%rip), %xmm3
	mulsd	-944(%rbp), %xmm3
	divsd	%xmm2, %xmm3
	comisd	%xmm4, %xmm3
	jnb	.L22
	cvttsd2siq	%xmm3, %rax
	jmp	.L23
.L22:
	movaps	%xmm3, %xmm0
	subsd	%xmm4, %xmm0
	cvttsd2siq	%xmm0, %rax
	btcq	$63, %rax
.L23:
	movq	%r12, %rcx
	movaps	%xmm2, %xmm0
	movq	%rbx, %r9
	movl	$2, %edi
	movsd	%xmm3, -944(%rbp)
	leaq	.LC12(%rip), %rsi
	pushq	%rax
	movb	$2, %al
	pushq	%r13
	movq	-920(%rbp), %rdx
	movq	-928(%rbp), %r8
	call	__printf_chk@PLT
	movsd	-944(%rbp), %xmm3
	addsd	-960(%rbp), %xmm3
	movsd	%xmm3, -960(%rbp)
	popq	%rdx
	popq	%rcx
.L18:
	movl	-904(%rbp), %edi
	call	cufftDestroy@PLT
	movl	-900(%rbp), %edi
	call	cufftDestroy@PLT
	movq	-888(%rbp), %rdi
	call	cudaFree@PLT
	movq	-896(%rbp), %rdi
	call	cudaFree@PLT
	call	cudaDeviceSynchronize@PLT
	incq	-952(%rbp)
	cmpq	$3, -952(%rbp)
	jne	.L24
	incl	-972(%rbp)
	addq	$32, -936(%rbp)
	cmpl	$24, -972(%rbp)
	jne	.L6
	movq	-968(%rbp), %rdi
	leaq	.LC15(%rip), %rbx
	call	free@PLT
	testb	%r14b, %r14b
	movsd	-960(%rbp), %xmm0
	movsd	.LC9(%rip), %xmm1
	divsd	.LC14(%rip), %xmm0
	je	.L26
	comisd	%xmm1, %xmm0
	jnb	.L27
	cvttsd2siq	%xmm0, %rcx
	jmp	.L28
.L27:
	movaps	%xmm0, %xmm2
	subsd	%xmm1, %xmm2
	cvttsd2siq	%xmm2, %rcx
	btcq	$63, %rcx
.L28:
	movq	-984(%rbp), %rdi
	movq	%rbx, %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	movsd	%xmm0, -920(%rbp)
	call	__fprintf_chk@PLT
	movsd	.LC9(%rip), %xmm1
	movsd	-920(%rbp), %xmm0
.L26:
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	jne	.L45
	comisd	%xmm1, %xmm0
	jnb	.L30
	cvttsd2siq	%xmm0, %rdx
	jmp	.L31
.L30:
	subsd	%xmm1, %xmm0
	cvttsd2siq	%xmm0, %rdx
	btcq	$63, %rdx
.L31:
	leaq	-40(%rbp), %rsp
	movq	%rbx, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_remember_state
	.cfi_def_cfa 7, 8
	jmp	__printf_chk@PLT
.L45:
	.cfi_restore_state
	call	__stack_chk_fail@PLT
.L32:
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4745:
	.size	_Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi, .-_Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4774:
	.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
.LFE4774:
	.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
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC3:
	.long	805306368
	.align 4
.LC4:
	.long	1065353216
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC5:
	.long	0
	.long	1101004800
	.align 8
.LC7:
	.long	0
	.long	1106247680
	.align 8
.LC8:
	.long	0
	.long	1083129856
	.align 8
.LC9:
	.long	0
	.long	1138753536
	.align 8
.LC10:
	.long	-755914244
	.long	1062232653
	.align 8
.LC11:
	.long	0
	.long	1074266112
	.align 8
.LC13:
	.long	0
	.long	1062207488
	.align 8
.LC14:
	.long	0
	.long	1077346304
	.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:

//general parts
#include <stdio.h>
#include <vector>
#include <memory>
#include <string.h>
#include <chrono>
#include <thread>
#include <iostream>
#ifndef __STDC_FORMAT_MACROS
#define __STDC_FORMAT_MACROS
#endif
#include <inttypes.h>

//CUDA parts
#include "hip/hip_runtime.h"

#include <hipfft/hipfft.h>

#define GROUP 1


void sample_6_benchmark_cuFFT_single_r2c(bool file_output, FILE* output, int device_id)
{
	if (file_output)
		fprintf(output, "6 - cuFFT FFT + iFFT R2C/C2R multidimensional benchmark in single precision\n");
	printf("6 - cuFFT FFT + iFFT R2C/C2R multidimensional benchmark in single precision\n");
	hipSetDevice(device_id);
	const int num_benchmark_samples = 24;
	const int num_runs = 3;
	//printf("First %" PRIu64 " runs are a warmup\n", num_runs);
	uint64_t benchmark_dimensions[num_benchmark_samples][4] = { {1024, 1024, 1, 2}, {64, 64, 1, 2}, {256, 256, 1, 2}, {1024, 256, 1, 2}, {512, 512, 1, 2}, {1024, 1024, 1, 2},  {4096, 256, 1, 2}, {2048, 1024, 1, 2},{4096, 2048, 1, 2}, {4096, 4096, 1, 2}, {720, 480, 1, 2},{1280, 720, 1, 2},{1920, 1080, 1, 2}, {2560, 1440, 1, 2},{3840, 2160, 1, 2},
																{32, 32, 32, 3}, {64, 64, 64, 3}, {256, 256, 32, 3},  {1024, 256, 32, 3},  {256, 256, 256, 3}, {2048, 1024, 8, 3},  {512, 512, 128, 3}, {2048, 256, 256, 3}, {4096, 512, 8, 3} };
	
	double benchmark_result[2] = { 0,0 };//averaged result = sum(system_size/iteration_time)/num_benchmark_samples
	hipfftReal* inputC = (hipfftReal*)malloc((uint64_t)sizeof(hipfftReal)*pow(2, 27));
	for (uint64_t i = 0; i < pow(2, 27); i++) {
		inputC[i] = 2 * ((float)rand()) / RAND_MAX - 1.0;
	}
	for (int n = 0; n < num_benchmark_samples; n++) {
		double run_time[num_runs][2];
		for (int r = 0; r < num_runs; r++) {
			hipfftHandle planR2C;
			hipfftHandle planC2R;
			hipfftReal* dataR;
			hipfftComplex* dataC;

			uint64_t dims[3] = { benchmark_dimensions[n][0] , benchmark_dimensions[n][1] ,benchmark_dimensions[n][2] };

			hipMalloc((void**)&dataR, sizeof(hipfftComplex) * (dims[0] / 2 + 1) * dims[1] * dims[2]);
			hipMalloc((void**)&dataC, sizeof(hipfftComplex) * (dims[0] / 2 + 1) * dims[1] * dims[2]);

			hipMemcpy(dataR, inputC, sizeof(hipfftReal) * dims[0] * dims[1] * dims[2], hipMemcpyHostToDevice);
			if (hipGetLastError() != hipSuccess) {
				fprintf(stderr, "Cuda error: Failed to allocate\n");
				return;
			}
			switch (benchmark_dimensions[n][3]) {
			case 1:
				hipfftPlan1d(&planR2C, dims[0], HIPFFT_R2C, 1);
				hipfftPlan1d(&planC2R, dims[0], HIPFFT_C2R, 1);
				break;
			case 2:
				hipfftPlan2d(&planR2C, dims[1], dims[0], HIPFFT_R2C);
				hipfftPlan2d(&planC2R, dims[1], dims[0], HIPFFT_C2R);
				break;
			case 3:
				hipfftPlan3d(&planR2C, dims[2], dims[1], dims[0], HIPFFT_R2C);
				hipfftPlan3d(&planC2R, dims[2], dims[1], dims[0], HIPFFT_C2R);
				break;
			}

			float totTime = 0;
			uint64_t cuBufferSize = sizeof(float) * 2 * (dims[0]/2+1) * dims[1] * dims[2];
			uint64_t num_iter = ((4096 * 1024.0 * 1024.0) / cuBufferSize > 1000) ? 1000 : (4096 * 1024.0 * 1024.0) / cuBufferSize;
			if (num_iter == 0) num_iter = 1;
			std::chrono::steady_clock::time_point timeSubmit = std::chrono::steady_clock::now();
			for (int i = 0; i < num_iter; i++) {

				hipfftExecR2C(planR2C, dataR, dataC);
				hipfftExecC2R(planC2R, dataC, dataR);
			}
			hipDeviceSynchronize();
			std::chrono::steady_clock::time_point timeEnd = std::chrono::steady_clock::now();
			totTime = (std::chrono::duration_cast<std::chrono::microseconds>(timeEnd - timeSubmit).count() * 0.001) / num_iter;
			run_time[r][0] = totTime;
			if (n > 0) {
				if (r == num_runs - 1) {
					double std_error = 0;
					double avg_time = 0;
					for (uint64_t t = 0; t < num_runs; t++) {
						avg_time += run_time[t][0];
					}
					avg_time /= num_runs;
					for (uint64_t t = 0; t < num_runs; t++) {
						std_error += (run_time[t][0] - avg_time) * (run_time[t][0] - avg_time);
					}
					std_error = sqrt(std_error / num_runs);
					if (file_output)
						fprintf(output, "cuFFT System: %" PRIu64 "x%" PRIu64 "x%" PRIu64 " Buffer: %" PRIu64 " MB avg_time_per_step: %0.3f ms std_error: %0.3f num_iter: %" PRIu64 " benchmark: %" PRIu64 "\n", benchmark_dimensions[n][0], benchmark_dimensions[n][1], benchmark_dimensions[n][2], cuBufferSize / 1024 / 1024, avg_time, std_error, num_iter, (uint64_t)(((double)cuBufferSize / 1024) / avg_time));

					printf("cuFFT System: %" PRIu64 "x%" PRIu64 "x%" PRIu64 " Buffer: %" PRIu64 " MB avg_time_per_step: %0.3f ms std_error: %0.3f num_iter: %" PRIu64 " benchmark: %" PRIu64 "\n", benchmark_dimensions[n][0], benchmark_dimensions[n][1], benchmark_dimensions[n][2], cuBufferSize / 1024 / 1024, avg_time, std_error, num_iter, (uint64_t)(((double)cuBufferSize / 1024) / avg_time));
					benchmark_result[0] += ((double)cuBufferSize / 1024) / avg_time;
				}

			}
			hipfftDestroy(planR2C);
			hipfftDestroy(planC2R);
			hipFree(dataC);
			hipFree(dataR);
			hipDeviceSynchronize();
			//cufftComplex* output_cuFFT = (cufftComplex*)(malloc(sizeof(cufftComplex) * dims[0] * dims[1] * dims[2]));
			//cudaMemcpy(output_cuFFT, dataC, sizeof(cufftComplex) * dims[0] * dims[1] * dims[2], cudaMemcpyDeviceToHost);
			//cudaDeviceSynchronize();
			

		}
	}
	free(inputC);
	benchmark_result[0] /= (num_benchmark_samples - 1);
	if (file_output)
		fprintf(output, "Benchmark score cuFFT: %" PRIu64 "\n", (uint64_t)(benchmark_result[0]));
	printf("Benchmark score cuFFT: %" PRIu64 "\n", (uint64_t)(benchmark_result[0]));

}

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

                                        # End of file scope inline assembly
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi
.LCPI0_0:
	.long	0x30000000                      # float 4.65661287E-10
.LCPI0_1:
	.long	0xbf800000                      # float -1
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0
.LCPI0_2:
	.long	1127219200                      # 0x43300000
	.long	1160773632                      # 0x45300000
	.long	0                               # 0x0
	.long	0                               # 0x0
.LCPI0_3:
	.quad	0x4330000000000000              # double 4503599627370496
	.quad	0x4530000000000000              # double 1.9342813113834067E+25
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI0_4:
	.quad	0x41f0000000000000              # double 4294967296
.LCPI0_5:
	.quad	0x408f400000000000              # double 1000
.LCPI0_6:
	.quad	0x43e0000000000000              # double 9.2233720368547758E+18
.LCPI0_7:
	.quad	0x3f50000000000000              # double 9.765625E-4
.LCPI0_8:
	.quad	0x3f50624dd2f1a9fc              # double 0.001
.LCPI0_9:
	.quad	0x4008000000000000              # double 3
.LCPI0_10:
	.quad	0x4037000000000000              # double 23
.LCPI0_11:
	.quad	0x0000000000000000              # double 0
	.text
	.globl	_Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi
	.type	_Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi,@function
_Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi: # @_Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi
	.cfi_startproc
# %bb.0:
	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	$248, %rsp
	.cfi_def_cfa_offset 304
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edx, %ebx
	movq	%rsi, 48(%rsp)                  # 8-byte Spill
	movl	%edi, 12(%rsp)                  # 4-byte Spill
	testl	%edi, %edi
	je	.LBB0_2
# %bb.1:
	movl	$.L.str, %edi
	movl	$76, %esi
	movl	$1, %edx
	movq	48(%rsp), %rcx                  # 8-byte Reload
	callq	fwrite@PLT
.LBB0_2:
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	%ebx, %edi
	callq	hipSetDevice
	movl	$536870912, %edi                # imm = 0x20000000
	callq	malloc
	movq	%rax, %r14
	xorl	%ebx, %ebx
.LBB0_3:                                # =>This Inner Loop Header: Depth=1
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	addss	%xmm0, %xmm0
	mulss	.LCPI0_0(%rip), %xmm0
	addss	.LCPI0_1(%rip), %xmm0
	movss	%xmm0, (%r14,%rbx,4)
	incq	%rbx
	cmpq	$134217728, %rbx                # imm = 0x8000000
	jne	.LBB0_3
# %bb.4:                                # %.split.preheader
	xorl	%eax, %eax
	xorps	%xmm0, %xmm0
	movaps	%xmm0, 112(%rsp)                # 16-byte Spill
	movq	%r14, 96(%rsp)                  # 8-byte Spill
.LBB0_5:                                # %.split
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_6 Depth 2
                                        #       Child Loop BB0_15 Depth 3
                                        #       Child Loop BB0_20 Depth 3
                                        #       Child Loop BB0_22 Depth 3
	movq	%rax, 104(%rsp)                 # 8-byte Spill
	shlq	$5, %rax
	leaq	.L__const._Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi.benchmark_dimensions+24(%rax), %rax
	movq	-24(%rax), %r13
	movq	-16(%rax), %rcx
	movq	%rax, 128(%rsp)                 # 8-byte Spill
	movq	-8(%rax), %rdx
	leaq	(,%r13,4), %rax
	movq	%rax, %r15
	andq	$-8, %r15
	addq	$8, %r15
	movq	%rdx, 56(%rsp)                  # 8-byte Spill
	movq	%rcx, 16(%rsp)                  # 8-byte Spill
	imulq	%rcx, %rdx
	imulq	%rdx, %r15
	imulq	%rax, %rdx
	movq	%rdx, 136(%rsp)                 # 8-byte Spill
	movq	%r15, %xmm0
	movdqa	.LCPI0_2(%rip), %xmm2           # xmm2 = [1127219200,1160773632,0,0]
	punpckldq	%xmm2, %xmm0            # xmm0 = xmm0[0],xmm2[0],xmm0[1],xmm2[1]
	movapd	.LCPI0_3(%rip), %xmm3           # xmm3 = [4.503599627370496E+15,1.9342813113834067E+25]
	subpd	%xmm3, %xmm0
	movapd	%xmm0, %xmm4
	unpckhpd	%xmm0, %xmm4                    # xmm4 = xmm4[1],xmm0[1]
	addsd	%xmm0, %xmm4
	movsd	.LCPI0_4(%rip), %xmm0           # xmm0 = mem[0],zero
	divsd	%xmm4, %xmm0
	movsd	.LCPI0_5(%rip), %xmm1           # xmm1 = mem[0],zero
	minsd	%xmm0, %xmm1
	cvttsd2si	%xmm1, %rax
	movq	%rax, %rcx
	subsd	.LCPI0_6(%rip), %xmm1
	cvttsd2si	%xmm1, %rbx
	sarq	$63, %rcx
	andq	%rcx, %rbx
	orq	%rax, %rbx
	cmpq	$1, %rbx
	adcq	$0, %rbx
	movq	%rbx, %xmm0
	punpckldq	%xmm2, %xmm0            # xmm0 = xmm0[0],xmm2[0],xmm0[1],xmm2[1]
	subpd	%xmm3, %xmm0
	movapd	%xmm0, %xmm1
	unpckhpd	%xmm0, %xmm1                    # xmm1 = xmm1[1],xmm0[1]
	addsd	%xmm0, %xmm1
	movapd	%xmm1, 176(%rsp)                # 16-byte Spill
	movq	%r15, %rax
	shrq	$20, %rax
	movq	%rax, 80(%rsp)                  # 8-byte Spill
	mulsd	.LCPI0_7(%rip), %xmm4
	movapd	%xmm4, 144(%rsp)                # 16-byte Spill
	xorl	%ebp, %ebp
.LBB0_6:                                #   Parent Loop BB0_5 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB0_15 Depth 3
                                        #       Child Loop BB0_20 Depth 3
                                        #       Child Loop BB0_22 Depth 3
	leaq	24(%rsp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	leaq	64(%rsp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movq	24(%rsp), %rdi
	movq	%r14, %rsi
	movq	136(%rsp), %rdx                 # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB0_7
# %bb.8:                                #   in Loop: Header=BB0_6 Depth=2
	movq	128(%rsp), %rax                 # 8-byte Reload
	movq	(%rax), %rax
	cmpq	$3, %rax
	je	.LBB0_13
# %bb.9:                                #   in Loop: Header=BB0_6 Depth=2
	cmpq	$2, %rax
	je	.LBB0_12
# %bb.10:                               #   in Loop: Header=BB0_6 Depth=2
	cmpq	$1, %rax
	jne	.LBB0_14
# %bb.11:                               #   in Loop: Header=BB0_6 Depth=2
	leaq	40(%rsp), %rdi
	movl	%r13d, %esi
	movl	$42, %edx
	movl	$1, %ecx
	callq	hipfftPlan1d
	leaq	32(%rsp), %rdi
	movl	%r13d, %esi
	movl	$44, %edx
	movl	$1, %ecx
	callq	hipfftPlan1d
	jmp	.LBB0_14
.LBB0_12:                               #   in Loop: Header=BB0_6 Depth=2
	leaq	40(%rsp), %rdi
	movq	16(%rsp), %r14                  # 8-byte Reload
	movl	%r14d, %esi
	movl	%r13d, %edx
	movl	$42, %ecx
	callq	hipfftPlan2d
	leaq	32(%rsp), %rdi
	movl	%r14d, %esi
	movl	%r13d, %edx
	movl	$44, %ecx
	callq	hipfftPlan2d
	jmp	.LBB0_14
.LBB0_13:                               #   in Loop: Header=BB0_6 Depth=2
	leaq	40(%rsp), %rdi
	movq	56(%rsp), %r12                  # 8-byte Reload
	movl	%r12d, %esi
	movq	16(%rsp), %r14                  # 8-byte Reload
	movl	%r14d, %edx
	movl	%r13d, %ecx
	movl	$42, %r8d
	callq	hipfftPlan3d
	leaq	32(%rsp), %rdi
	movl	%r12d, %esi
	movl	%r14d, %edx
	movl	%r13d, %ecx
	movl	$44, %r8d
	callq	hipfftPlan3d
.LBB0_14:                               #   in Loop: Header=BB0_6 Depth=2
	callq	_ZNSt6chrono3_V212steady_clock3nowEv
	movq	%rax, %r12
	movq	%rbx, %r14
.LBB0_15:                               #   Parent Loop BB0_5 Depth=1
                                        #     Parent Loop BB0_6 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movq	40(%rsp), %rdi
	movq	24(%rsp), %rsi
	movq	64(%rsp), %rdx
	callq	hipfftExecR2C
	movq	32(%rsp), %rdi
	movq	64(%rsp), %rsi
	movq	24(%rsp), %rdx
	callq	hipfftExecC2R
	decq	%r14
	jne	.LBB0_15
# %bb.16:                               #   in Loop: Header=BB0_6 Depth=2
	callq	hipDeviceSynchronize
	callq	_ZNSt6chrono3_V212steady_clock3nowEv
	subq	%r12, %rax
	movabsq	$2361183241434822607, %rcx      # imm = 0x20C49BA5E353F7CF
	imulq	%rcx
	movq	%rdx, %rax
	shrq	$63, %rax
	sarq	$7, %rdx
	addq	%rax, %rdx
	movq	%rbp, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rdx, %xmm0
	mulsd	.LCPI0_8(%rip), %xmm0
	divsd	176(%rsp), %xmm0                # 16-byte Folded Reload
	cvtsd2ss	%xmm0, %xmm0
	shlq	$4, %rax
	cvtss2sd	%xmm0, %xmm0
	movsd	%xmm0, 192(%rsp,%rax)
	cmpq	$0, 104(%rsp)                   # 8-byte Folded Reload
	je	.LBB0_17
# %bb.18:                               #   in Loop: Header=BB0_6 Depth=2
	cmpq	$2, %rbp
	movq	96(%rsp), %r14                  # 8-byte Reload
	jne	.LBB0_29
# %bb.19:                               # %.preheader.preheader
                                        #   in Loop: Header=BB0_6 Depth=2
	pxor	%xmm2, %xmm2
	xorl	%eax, %eax
.LBB0_20:                               # %.preheader
                                        #   Parent Loop BB0_5 Depth=1
                                        #     Parent Loop BB0_6 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	addsd	192(%rsp,%rax), %xmm2
	addq	$16, %rax
	cmpq	$48, %rax
	jne	.LBB0_20
# %bb.21:                               #   in Loop: Header=BB0_6 Depth=2
	divsd	.LCPI0_9(%rip), %xmm2
	xorps	%xmm0, %xmm0
	xorl	%eax, %eax
.LBB0_22:                               #   Parent Loop BB0_5 Depth=1
                                        #     Parent Loop BB0_6 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movsd	192(%rsp,%rax), %xmm1           # xmm1 = mem[0],zero
	subsd	%xmm2, %xmm1
	mulsd	%xmm1, %xmm1
	addsd	%xmm1, %xmm0
	addq	$16, %rax
	cmpq	$48, %rax
	jne	.LBB0_22
# %bb.23:                               #   in Loop: Header=BB0_6 Depth=2
	divsd	.LCPI0_9(%rip), %xmm0
	ucomisd	.LCPI0_11(%rip), %xmm0
	movsd	%xmm2, 88(%rsp)                 # 8-byte Spill
	jb	.LBB0_25
# %bb.24:                               #   in Loop: Header=BB0_6 Depth=2
	xorps	%xmm1, %xmm1
	sqrtsd	%xmm0, %xmm1
	jmp	.LBB0_26
.LBB0_17:                               #   in Loop: Header=BB0_6 Depth=2
	movq	96(%rsp), %r14                  # 8-byte Reload
	jmp	.LBB0_29
.LBB0_25:                               # %call.sqrt
                                        #   in Loop: Header=BB0_6 Depth=2
	callq	sqrt
	movsd	88(%rsp), %xmm2                 # 8-byte Reload
                                        # xmm2 = mem[0],zero
	movapd	%xmm0, %xmm1
.LBB0_26:                               # %.split160
                                        #   in Loop: Header=BB0_6 Depth=2
	movapd	144(%rsp), %xmm0                # 16-byte Reload
	divsd	%xmm2, %xmm0
	cvttsd2si	%xmm0, %rax
	movq	%rax, %rcx
	movapd	%xmm0, 160(%rsp)                # 16-byte Spill
	subsd	.LCPI0_6(%rip), %xmm0
	cvttsd2si	%xmm0, %r12
	sarq	$63, %rcx
	andq	%rcx, %r12
	orq	%rax, %r12
	cmpb	$0, 12(%rsp)                    # 1-byte Folded Reload
	je	.LBB0_28
# %bb.27:                               #   in Loop: Header=BB0_6 Depth=2
	movl	$.L.str.2, %esi
	movq	48(%rsp), %rdi                  # 8-byte Reload
	movq	%r13, %rdx
	movq	16(%rsp), %rcx                  # 8-byte Reload
	movq	56(%rsp), %r8                   # 8-byte Reload
	movq	80(%rsp), %r9                   # 8-byte Reload
	movapd	%xmm2, %xmm0
	movsd	%xmm1, 72(%rsp)                 # 8-byte Spill
	movsd	72(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	movb	$2, %al
	pushq	%r12
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	callq	fprintf
	movsd	88(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	movsd	104(%rsp), %xmm2                # 8-byte Reload
                                        # xmm2 = mem[0],zero
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
.LBB0_28:                               # %._crit_edge
                                        #   in Loop: Header=BB0_6 Depth=2
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	$.L.str.2, %edi
	movq	%r13, %rsi
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movq	64(%rsp), %rcx                  # 8-byte Reload
	movq	88(%rsp), %r8                   # 8-byte Reload
	movapd	%xmm2, %xmm0
	movq	%rbx, %r9
	movb	$2, %al
	pushq	%r12
	.cfi_adjust_cfa_offset 8
	callq	printf
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movapd	112(%rsp), %xmm0                # 16-byte Reload
	addsd	160(%rsp), %xmm0                # 16-byte Folded Reload
	movapd	%xmm0, 112(%rsp)                # 16-byte Spill
.LBB0_29:                               #   in Loop: Header=BB0_6 Depth=2
	movq	40(%rsp), %rdi
	callq	hipfftDestroy
	movq	32(%rsp), %rdi
	callq	hipfftDestroy
	movq	64(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	callq	hipDeviceSynchronize
	incq	%rbp
	cmpq	$3, %rbp
	jne	.LBB0_6
# %bb.30:                               # %.thread110
                                        #   in Loop: Header=BB0_5 Depth=1
	movq	104(%rsp), %rax                 # 8-byte Reload
	incq	%rax
	cmpq	$24, %rax
	jne	.LBB0_5
# %bb.31:                               # %.thread115
	movq	%r14, %rdi
	callq	free
	movapd	112(%rsp), %xmm0                # 16-byte Reload
	divsd	.LCPI0_10(%rip), %xmm0
	cvttsd2si	%xmm0, %rax
	movq	%rax, %rcx
	subsd	.LCPI0_6(%rip), %xmm0
	cvttsd2si	%xmm0, %rbx
	sarq	$63, %rcx
	andq	%rcx, %rbx
	orq	%rax, %rbx
	cmpb	$0, 12(%rsp)                    # 1-byte Folded Reload
	je	.LBB0_33
# %bb.32:
	movl	$.L.str.3, %esi
	movq	48(%rsp), %rdi                  # 8-byte Reload
	movq	%rbx, %rdx
	xorl	%eax, %eax
	callq	fprintf
.LBB0_33:                               # %.thread115._crit_edge
	movl	$.L.str.3, %edi
	movq	%rbx, %rsi
	xorl	%eax, %eax
	callq	printf
	jmp	.LBB0_34
.LBB0_7:                                # %.thread113
	movq	stderr(%rip), %rcx
	movl	$.L.str.1, %edi
	movl	$31, %esi
	movl	$1, %edx
	callq	fwrite@PLT
.LBB0_34:
	addq	$248, %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
	retq
.Lfunc_end0:
	.size	_Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi, .Lfunc_end0-_Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"6 - cuFFT FFT + iFFT R2C/C2R multidimensional benchmark in single precision\n"
	.size	.L.str, 77

	.type	.L__const._Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi.benchmark_dimensions,@object # @__const._Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi.benchmark_dimensions
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
.L__const._Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi.benchmark_dimensions:
	.quad	1024                            # 0x400
	.quad	1024                            # 0x400
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	64                              # 0x40
	.quad	64                              # 0x40
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	256                             # 0x100
	.quad	256                             # 0x100
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	1024                            # 0x400
	.quad	256                             # 0x100
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	512                             # 0x200
	.quad	512                             # 0x200
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	1024                            # 0x400
	.quad	1024                            # 0x400
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	4096                            # 0x1000
	.quad	256                             # 0x100
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	2048                            # 0x800
	.quad	1024                            # 0x400
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	4096                            # 0x1000
	.quad	2048                            # 0x800
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	4096                            # 0x1000
	.quad	4096                            # 0x1000
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	720                             # 0x2d0
	.quad	480                             # 0x1e0
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	1280                            # 0x500
	.quad	720                             # 0x2d0
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	1920                            # 0x780
	.quad	1080                            # 0x438
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	2560                            # 0xa00
	.quad	1440                            # 0x5a0
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	3840                            # 0xf00
	.quad	2160                            # 0x870
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	32                              # 0x20
	.quad	32                              # 0x20
	.quad	32                              # 0x20
	.quad	3                               # 0x3
	.quad	64                              # 0x40
	.quad	64                              # 0x40
	.quad	64                              # 0x40
	.quad	3                               # 0x3
	.quad	256                             # 0x100
	.quad	256                             # 0x100
	.quad	32                              # 0x20
	.quad	3                               # 0x3
	.quad	1024                            # 0x400
	.quad	256                             # 0x100
	.quad	32                              # 0x20
	.quad	3                               # 0x3
	.quad	256                             # 0x100
	.quad	256                             # 0x100
	.quad	256                             # 0x100
	.quad	3                               # 0x3
	.quad	2048                            # 0x800
	.quad	1024                            # 0x400
	.quad	8                               # 0x8
	.quad	3                               # 0x3
	.quad	512                             # 0x200
	.quad	512                             # 0x200
	.quad	128                             # 0x80
	.quad	3                               # 0x3
	.quad	2048                            # 0x800
	.quad	256                             # 0x100
	.quad	256                             # 0x100
	.quad	3                               # 0x3
	.quad	4096                            # 0x1000
	.quad	512                             # 0x200
	.quad	8                               # 0x8
	.quad	3                               # 0x3
	.size	.L__const._Z35sample_6_benchmark_cuFFT_single_r2cbP8_IO_FILEi.benchmark_dimensions, 768

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"Cuda error: Failed to allocate\n"
	.size	.L.str.1, 32

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"cuFFT System: %lux%lux%lu Buffer: %lu MB avg_time_per_step: %0.3f ms std_error: %0.3f num_iter: %lu benchmark: %lu\n"
	.size	.L.str.2, 116

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Benchmark score cuFFT: %lu\n"
	.size	.L.str.3, 28

	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"6 - cuFFT FFT + iFFT R2C/C2R multidimensional benchmark in single precision"
	.size	.Lstr, 76

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
