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 →
1.3 KiB48 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
18void launch_precision_rocFFT_double(void* inputC, void* output_rocFFT, int device_id, uint64_t* dims)
19{
20	hipSetDevice(device_id);
21	hipfftHandle planZ2Z;
22	hipfftDoubleComplex* dataC;
23	hipMalloc((void**)&dataC, sizeof(hipfftDoubleComplex) * dims[0] * dims[1] * dims[2]);
24	hipMemcpy(dataC, inputC, sizeof(hipfftDoubleComplex) * dims[0] * dims[1] * dims[2], hipMemcpyHostToDevice);
25	if (hipGetLastError() != hipSuccess) {
26		fprintf(stderr, "ROCM error: Failed to allocate\n");
27		return;
28	}
29	switch (dims[4]) {
30	case 1:
31		hipfftPlan1d(&planZ2Z, dims[0], HIPFFT_Z2Z, 1);
32		break;
33	case 2:
34		hipfftPlan2d(&planZ2Z, dims[1], dims[0], HIPFFT_Z2Z);
35		break;
36	case 3:
37		hipfftPlan3d(&planZ2Z, dims[2], dims[1], dims[0], HIPFFT_Z2Z);
38		break;
39	}
40	for (int i = 0; i < 1; i++) {
41		hipfftExecZ2Z(planZ2Z, dataC, dataC, -1);
42	}
43	hipDeviceSynchronize();
44	hipMemcpy(output_rocFFT, dataC, sizeof(hipfftDoubleComplex) * dims[0] * dims[1] * dims[2], hipMemcpyDeviceToHost);
45	hipDeviceSynchronize();
46	hipfftDestroy(planZ2Z);
47	hipFree(dataC);
48}