hey! this is a static git repo viewer. commit diffs, file history, older commits and other revisions won't load without javascript.

you are still able to view source files on branch tips, but the experience is degraded.

yum-mirror/VkFFT

git clone https://git.yummers.dev/yum-mirror/VkFFT

Tolmachev DmitriiMetal support in VkFFTba7001c2022-10-06

1//general parts
2#include <stdio.h>
3#include <vector>
4#include <memory>
5#include <string.h>
6#include <chrono>
7#include <thread>
8#include <iostream>
9#ifndef __STDC_FORMAT_MACROS
10#define __STDC_FORMAT_MACROS
11#endif
12#include <inttypes.h>
13
14//ROCM parts
15#include "hip/hip_runtime.h"
16#include <hipfft.h>
17
18#include "user_benchmark_rocFFT.h"
19
20void user_benchmark_rocFFT(bool file_output, FILE* output, rocFFTUserSystemParameters* userParams, int device_id)
21{
22	hipSetDevice(device_id);
23	const int num_runs = 7;
24	double benchmark_result[2] = { 0,0 };//averaged result = sum(system_size/iteration_time)/num_benchmark_samples
25	uint64_t storageComplexSize=8;
26	switch (userParams->P) {
27	case 0:
28		storageComplexSize = (2 * sizeof(float));
29		break;
30	case 1:
31		storageComplexSize = (2 * sizeof(double));
32		break;
33	case 2:
34		storageComplexSize = (2 * 2);
35		break;
36	default:
37		storageComplexSize = (2 * sizeof(float));
38		break;
39	}
40    uint64_t bufferSize = 0;
41    if (userParams->R2C) {
42        bufferSize = (uint64_t)(storageComplexSize / 2) * (userParams->X + 2) * userParams->Y * userParams->Z * userParams->B;
43    }
44    else {
45		bufferSize = (uint64_t)storageComplexSize * userParams->X * userParams->Y * userParams->Z * userParams->B;
46    }
47    
48    float* buffer_input = (float*)malloc(bufferSize);
49    for (uint64_t i = 0; i < bufferSize/sizeof(float); i++) {
50        buffer_input[i] = (float)(2 * ((float)rand()) / RAND_MAX - 1.0);
51    }
52	for (int n = 0; n < 2; n++) {
53		double run_time[num_runs][2];
54		for (int r = 0; r < num_runs; r++) {
55			hipfftHandle plan;
56			hipfftHandle plan2;
57			void* dataC;
58
59			int dims[3];
60			int FFTdim = 1;
61			if (userParams->Y > 1) FFTdim++;
62			if (userParams->Z > 1) FFTdim++;
63			switch (FFTdim) {
64			case 1:
65				dims[0] = userParams->X;
66				dims[1] = 1;
67				dims[2] = 1;
68				break;
69			case 2:
70				dims[0] = userParams->Y;
71				dims[1] = userParams->X;
72				dims[2] = 1;
73				break;
74			case 3:
75				dims[0] = userParams->Z;
76				dims[1] = userParams->Y;
77				dims[2] = userParams->X;
78				break;
79			}
80		
81			hipMalloc((void**)&dataC, bufferSize);
82			
83			if (hipGetLastError() != hipSuccess) {
84				fprintf(stderr, "ROCM error: Failed to allocate\n");
85				return;
86			}
87			hipMemcpy(dataC, buffer_input, bufferSize, hipMemcpyHostToDevice);
88			
89			//forward + inverse
90			int iembed[2][3];
91			int istride[2] = { 1, 1 };
92			int idist[2] = {bufferSize / userParams->B / storageComplexSize, bufferSize / userParams->B / storageComplexSize};
93			if (userParams->R2C) idist[0] *= 2;
94			int oembed[2][3];
95			int ostride[2] = { 1, 1 };
96			int odist[2] = { bufferSize / userParams->B / storageComplexSize, bufferSize / userParams->B / storageComplexSize };
97			if (userParams->R2C) odist[1] *= 2;
98			switch (FFTdim) {
99			case 1:
100				iembed[0][0] = (userParams->R2C) ? dims[0] + 2 : dims[0];
101				oembed[0][0] = (userParams->R2C) ? (dims[0] + 2) / 2 : dims[0];
102
103				iembed[1][0] = (userParams->R2C) ? (dims[0] + 2) / 2 : dims[0];
104				oembed[1][0] = (userParams->R2C) ? dims[0] + 2 : dims[0];
105				break;
106			case 2:
107				iembed[0][0] = dims[0];
108				iembed[0][1] = (userParams->R2C) ? dims[1] + 2 : dims[1];
109				oembed[0][0] = dims[0];
110				oembed[0][1] = (userParams->R2C) ? (dims[1] + 2) / 2 : dims[1];
111
112				iembed[1][0] = dims[0];
113				iembed[1][1] = (userParams->R2C) ? (dims[1] + 2) / 2 : dims[1];
114				oembed[1][0] = dims[0];
115				oembed[1][1] = (userParams->R2C) ? dims[1] + 2 : dims[1];
116				break;
117			case 3:
118				iembed[0][0] = idist[0];
119				iembed[0][1] = dims[1];
120				iembed[0][2] = (userParams->R2C) ? dims[2] + 2 : dims[2];
121				oembed[0][0] = odist[0];
122				oembed[0][1] = dims[1];
123				oembed[0][2] = (userParams->R2C) ? (dims[2] + 2)/2 : dims[2];
124
125				iembed[1][0] = idist[0];
126				iembed[1][1] = dims[1];
127				iembed[1][2] = (userParams->R2C) ? (dims[2] + 2)/2 : dims[2];
128				oembed[1][0] = odist[0];
129				oembed[1][1] = dims[1];
130				oembed[1][2] = (userParams->R2C) ? dims[2] + 2 : dims[2];
131				break;
132			}
133			switch (userParams->P) {
134			case 0:
135				if (userParams->R2C){
136					hipfftPlanMany(&plan, FFTdim, dims, iembed[0], istride[0], idist[0], oembed[0], ostride[0], odist[0], HIPFFT_R2C, userParams->B);
137					hipfftPlanMany(&plan2, FFTdim, dims, iembed[1], istride[1], idist[1], oembed[1], ostride[1], odist[1], HIPFFT_C2R, userParams->B);
138				}else{
139					hipfftPlanMany(&plan, FFTdim, dims, iembed[0], istride[0], idist[0], oembed[0], ostride[0], odist[0], HIPFFT_C2C, userParams->B);
140				}
141				break;
142			case 1:
143				if (userParams->R2C){
144					hipfftPlanMany(&plan, FFTdim, dims,iembed[0], istride[0], idist[0], oembed[0], ostride[0], odist[0],  HIPFFT_D2Z, userParams->B);
145					hipfftPlanMany(&plan2, FFTdim, dims, iembed[1], istride[1], idist[1], oembed[1], ostride[1], odist[1], HIPFFT_Z2D, userParams->B);
146				}else
147					hipfftPlanMany(&plan, FFTdim, dims, iembed[0], istride[0], idist[0], oembed[0], ostride[0], odist[0],  HIPFFT_Z2Z, userParams->B);
148				break;
149			}
150
151			float totTime = 0;
152			std::chrono::steady_clock::time_point timeSubmit = std::chrono::steady_clock::now();
153			for (int i = 0; i < userParams->N; i++) {
154				switch (userParams->P) {
155				case 0:
156					if (userParams->R2C){
157						hipfftExecR2C(plan, (hipfftReal*) dataC, (hipfftComplex*) dataC);
158						hipfftExecC2R(plan2, (hipfftComplex*) dataC, (hipfftReal*) dataC);
159					}else{
160						hipfftExecC2C(plan, (hipfftComplex*) dataC, (hipfftComplex*) dataC, -1);
161						hipfftExecC2C(plan, (hipfftComplex*) dataC, (hipfftComplex*) dataC, 1);
162					}
163					break;
164				case 1:
165					if (userParams->R2C){
166						hipfftExecD2Z(plan, (hipfftDoubleReal*) dataC, (hipfftDoubleComplex*) dataC);
167						hipfftExecZ2D(plan2, (hipfftDoubleComplex*) dataC, (hipfftDoubleReal*) dataC);
168					}else{
169						hipfftExecZ2Z(plan, (hipfftDoubleComplex*) dataC, (hipfftDoubleComplex*) dataC, -1);
170						hipfftExecZ2Z(plan, (hipfftDoubleComplex*) dataC, (hipfftDoubleComplex*) dataC, 1);
171					}
172					break;
173				}
174			}
175			hipDeviceSynchronize();
176			std::chrono::steady_clock::time_point timeEnd = std::chrono::steady_clock::now();
177			totTime = (std::chrono::duration_cast<std::chrono::microseconds>(timeEnd - timeSubmit).count() * 0.001) / userParams->N;
178				
179			run_time[r][0] = totTime;
180			if (n > 0) {
181				if (r == num_runs - 1) {
182					double std_error = 0;
183					double avg_time = 0;
184					for (uint64_t t = 2; t < num_runs; t++) {
185						avg_time += run_time[t][0];
186					}
187					avg_time /= (num_runs-2);
188					for (uint64_t t = 2; t < num_runs; t++) {
189						std_error += (run_time[t][0] - avg_time) * (run_time[t][0] - avg_time);
190					}
191					std_error = sqrt(std_error / (num_runs-2));
192					if (file_output)
193						fprintf(output, "rocFFT System: %" PRIu64 "x%" PRIu64 "x%" PRIu64 " Batch: %" PRIu64 " Buffer: %" PRIu64 " MB avg_time_per_step: %0.3f ms std_error: %0.3f num_iter: %" PRIu64 " benchmark: %" PRIu64 " scaled bandwidth: %0.1f\n",  userParams->X, userParams->Y, userParams->Z, userParams->B, bufferSize / 1024 / 1024, avg_time, std_error, userParams->N, (uint64_t)(((double)bufferSize / 1024) / avg_time), ((double)bufferSize / 1024.0 / 1024.0 / 1.024 * 4 * FFTdim / avg_time));
194
195					printf("rocFFT System: %" PRIu64 "x%" PRIu64 "x%" PRIu64 " Batch: %" PRIu64 " Buffer: %" PRIu64 " MB avg_time_per_step: %0.3f ms std_error: %0.3f num_iter: %" PRIu64 " benchmark: %" PRIu64 " scaled bandwidth: %0.1f\n", userParams->X, userParams->Y, userParams->Z, userParams->B, bufferSize / 1024 / 1024, avg_time, std_error, userParams->N, (uint64_t)(((double)bufferSize / 1024) / avg_time), ((double)bufferSize / 1024.0 / 1024.0 / 1.024 * 4 * FFTdim / avg_time));
196					benchmark_result[0] += ((double)bufferSize / 1024) / avg_time;
197				}
198
199			}
200			hipfftDestroy(plan);
201			if (userParams->R2C)
202				hipfftDestroy(plan2);
203			hipFree(dataC);
204			hipDeviceSynchronize();
205			//hipfftComplex* output_rocFFT = (hipfftComplex*)(malloc(sizeof(hipfftComplex) * dims[0] * dims[1] * dims[2]));
206			//hipMemcpy(output_rocFFT, dataC, sizeof(hipfftComplex) * dims[0] * dims[1] * dims[2], hipMemcpyDeviceToHost);
207			//hipDeviceSynchronize();
208		}
209	}
210	free(buffer_input);
211}