yum-mirror/VkFFT

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

Dmitrii TolmachevNew register assignment logic -Implemented radix codelets up to 47. -Implemented composite radix codelets for arbitrary composite stage sizes. -Implemented new register assignment logic, aimed at optimizing shared memory transfers, register usage and warp utilization. -Performance improvements for all system sizes - please report regressions if they happen (especially for vendors other than Nvidia and AMD). -All double pointers passed to VkFFT now make local copy of their contents (https://github.com/DTolm/VkFFT/issues/184, https://github.com/DTolm/VkFFT/issues/185) -Fixed locale setting for code generator (https://github.com/vincefn/pyvkfft/issues/38)9a96811

develop
branchesdevelopmasterall refs →
7.8 KiB211 linesraw
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/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}