Home / C/C++ / Miscellaneous

VkFFT - Vulkan Fast Fourier Transform library

VkFFT is an efficient GPU-accelerated multidimensional Fast Fourier Transform library for Vulkan projects. VkFFT aims to provide community with an open-source alternative to Nvidia's cuFFT library, while achieving better performance.
Information
Category: C/C++ / Miscellaneous
Watchers: 2
Star: 19
Fork: 1
Last update: Aug 4, 2020
Resource links
https://github.com/DTolm/VkFFT

Build Status

VkFFT - Vulkan Fast Fourier Transform library

VkFFT is an efficient GPU-accelerated multidimensional Fast Fourier Transform library for Vulkan projects. VkFFT aims to provide community with an open-source alternative to Nvidia's cuFFT library, while achieving better performance.

Currently supported features:

  • 1D/2D/3D systems
  • Forward and inverse directions of FFT
  • Maximum dimension size is 4096, 32-bit float
  • Radix-2/4/8 FFT, only power of two systems
  • All transformations are performed in-place with no performance loss
  • Complex to complex (C2C), real to complex (R2C) and complex to real (C2R) transformations. R2C and C2R are optimized to run up to 2x times faster than C2C (2D and 3D case only)
  • 1x1, 2x2, 3x3 convolutions with symmetric or nonsymmetric kernel
  • Header-only (+precompiled shaders) library with Vulkan interface, which allows to append VkFFT directly to user's command buffer

Future release plan

  • Almost ready:
    • Zero padding support
    • Double and half-precision arithmetics
    • 8192 and 16384 dimension sizes
  • Planned
    • Publication based on implemented optimizations
    • Mobile and integrated GPU support
  • Ambitious
    • Multiple GPU job splitting

Installation

Include the vkFFT.h file and specify path to the shaders folder in CMake or from C interface. Sample CMakeLists.txt file configures project based on Vulkan_FFT.cpp file, which contains two examples on how to use VkFFT to perform FFT, iFFT and convolution calculations.

How to use VkFFT

VkFFT.h is a library which can append FFT, iFFT or convolution calculation to the user defined command buffer. It operates on storage buffers allocated by user and doesn't require any additional memory by itself. All computations are fully based on Vulkan compute shaders with no CPU usage except for FFT planning. VkFFT creates and optimizes memory layout by itself and performs FFT with the best chosen parameters. For an example application, see Vulkan_FFT.cpp file, which has comments explaining the VkFFT configuration process.
Picture below shows how data is restructured during the R2C transform depending on the system dimensions. This layout has minimal transfers between on-chip memory and graphics card (one read and one write per FFT axis + transposition if axis dimension is ≥ 256). If convolution is performed, it is embedded into the last FFT axis, which reduces memory transfers even further. alt text

Benchmark results in comparison to cuFFT

To measure how Vulkan FFT implementation works in comparison to cuFFT, we will perform a number of 2D and 3D tests. The test will consist of performing R2C FFT and inverse C2R FFT consecutively multiple times to calculate average time required. cuFFT uses out-of-place configuration while VkFFT uses in-place. The results are obtained on Nvidia 1660 Ti graphics card with no other GPU load. Launching example 0 from Vulkan_FFT.cpp performs VkFFT benchmark, benchmark_cuFFT.cu file contains similar benchmark script for cuFFT library. alt text alt text

Contact information

Initial version of VkFFT is developed by Tolmachev Dmitrii.
Email: [email protected]

0 Open More issues
Closed

Accuracy issue with double-precision (iFFT on Bluestein transforms with norm=1 ; some specific sizes.,..)

Hi @DTolm, I am implementing a systematic accuracy test for pyvkfft, following fftw's methodology, starting from a random complex array betwen -0.5 and +0.5 (for both real and imaginary parts), and comparing the result from the one obtained using long double accuracy in fftw (see https://github.com/vincefn/pyvkfft/blob/master/examples/accuracy.ipynb).

I am implementing the same approach in the systematic unit tests, and I found a discrepancy for Bluestein double precision transforms with norm=1. Here is an example output from my tests:

              shape    precision   ndim     norm                n2                ninf         tol       ninf<tol?
               (30,)    complex64 ndim=1 norm=    0  iFFT: n2=1.495756e-07 ninf=1.334992e-07 < 1.590849e-06 True
               (30,)    complex64 ndim=1 norm=    1   FFT: n2=1.744849e-07 ninf=2.263902e-07 < 1.590849e-06 True
               (30,)    complex64 ndim=1 norm=    1  iFFT: n2=1.448747e-07 ninf=2.058743e-07 < 1.590849e-06 True
               (30,)    complex64 ndim=1 norm=ortho   FFT: n2=1.744849e-07 ninf=2.263902e-07 < 1.590849e-06 True
               (30,)    complex64 ndim=1 norm=ortho  iFFT: n2=1.495756e-07 ninf=1.334992e-07 < 1.590849e-06 True
               (30,)   complex128 ndim=1 norm=    0   FFT: n2=2.707221e-16 ninf=4.774906e-16 < 1.443136e-15 True
               (30,)   complex128 ndim=1 norm=    0  iFFT: n2=2.696311e-16 ninf=3.399735e-16 < 1.443136e-15 True
               (30,)   complex128 ndim=1 norm=    1   FFT: n2=2.707221e-16 ninf=4.774906e-16 < 1.443136e-15 True
               (30,)   complex128 ndim=1 norm=    1  iFFT: n2=2.302078e-16 ninf=3.183271e-16 < 1.443136e-15 True
               (30,)   complex128 ndim=1 norm=ortho   FFT: n2=2.707221e-16 ninf=4.774906e-16 < 1.443136e-15 True
               (30,)   complex128 ndim=1 norm=ortho  iFFT: n2=2.696311e-16 ninf=3.399735e-16 < 1.443136e-15 True
            (30, 30)    complex64 ndim=1 norm=    0   FFT: n2=1.861849e-07 ninf=1.957236e-07 < 2.181697e-06 True
            (30, 30)    complex64 ndim=1 norm=    0  iFFT: n2=1.947846e-07 ninf=1.946365e-07 < 2.181697e-06 True
            (30, 30)    complex64 ndim=1 norm=    1   FFT: n2=1.861849e-07 ninf=1.957236e-07 < 2.181697e-06 True
            (30, 30)    complex64 ndim=1 norm=    1  iFFT: n2=1.662961e-07 ninf=1.756136e-07 < 2.181697e-06 True
            (30, 30)    complex64 ndim=1 norm=ortho   FFT: n2=1.861849e-07 ninf=1.957236e-07 < 2.181697e-06 True
            (30, 30)    complex64 ndim=1 norm=ortho  iFFT: n2=1.947846e-07 ninf=1.946365e-07 < 2.181697e-06 True
[...]
               (34,)    complex64 ndim=1 norm=    0   FFT: n2=3.737796e-07 ninf=3.612878e-07 < 1.612592e-06 True
               (34,)    complex64 ndim=1 norm=    0  iFFT: n2=3.329764e-07 ninf=3.350309e-07 < 1.612592e-06 True
               (34,)    complex64 ndim=1 norm=    1   FFT: n2=3.737796e-07 ninf=3.612878e-07 < 1.612592e-06 True
               (34,)    complex64 ndim=1 norm=    1  iFFT: n2=3.321449e-07 ninf=3.143022e-07 < 1.612592e-06 True
               (34,)    complex64 ndim=1 norm=ortho   FFT: n2=3.737796e-07 ninf=3.612878e-07 < 1.612592e-06 True
               (34,)    complex64 ndim=1 norm=ortho  iFFT: n2=3.329764e-07 ninf=3.350309e-07 < 1.612592e-06 True
               (34,)   complex128 ndim=1 norm=    0   FFT: n2=4.929609e-16 ninf=4.179421e-16 < 1.459444e-15 True
               (34,)   complex128 ndim=1 norm=    0  iFFT: n2=5.060519e-16 ninf=4.168635e-16 < 1.459444e-15 True
               (34,)   complex128 ndim=1 norm=    1   FFT: n2=4.929609e-16 ninf=4.179421e-16 < 1.459444e-15 True
               (34,)   complex128 ndim=1 norm=    1  iFFT: n2=1.797103e-14 ninf=1.813050e-14 < 1.459444e-15 False
               (34,)   complex128 ndim=1 norm=ortho   FFT: n2=4.929609e-16 ninf=4.179421e-16 < 1.459444e-15 True
               (34,)   complex128 ndim=1 norm=ortho  iFFT: n2=5.060519e-16 ninf=4.168635e-16 < 1.459444e-15 True
            (34, 34)    complex64 ndim=1 norm=    0   FFT: n2=3.373936e-07 ninf=3.704312e-07 < 2.225183e-06 True
            (34, 34)    complex64 ndim=1 norm=    0  iFFT: n2=3.231635e-07 ninf=3.330752e-07 < 2.225183e-06 True
            (34, 34)    complex64 ndim=1 norm=    1   FFT: n2=3.373936e-07 ninf=3.704312e-07 < 2.225183e-06 True
            (34, 34)    complex64 ndim=1 norm=    1  iFFT: n2=3.303084e-07 ninf=3.815846e-07 < 2.225183e-06 True
            (34, 34)    complex64 ndim=1 norm=ortho   FFT: n2=3.373936e-07 ninf=3.704312e-07 < 2.225183e-06 True
            (34, 34)    complex64 ndim=1 norm=ortho  iFFT: n2=3.231635e-07 ninf=3.330752e-07 < 2.225183e-06 True
            (34, 34)   complex128 ndim=1 norm=    0   FFT: n2=5.693998e-16 ninf=8.358774e-16 < 1.918887e-15 True
            (34, 34)   complex128 ndim=1 norm=    0  iFFT: n2=5.495357e-16 ninf=7.050855e-16 < 1.918887e-15 True
            (34, 34)   complex128 ndim=1 norm=    1   FFT: n2=5.693998e-16 ninf=8.358774e-16 < 1.918887e-15 True
            (34, 34)   complex128 ndim=1 norm=    1  iFFT: n2=1.800159e-14 ninf=1.810448e-14 < 1.918887e-15 False
            (34, 34)   complex128 ndim=1 norm=ortho   FFT: n2=5.693998e-16 ninf=8.358774e-16 < 1.918887e-15 True
            (34, 34)   complex128 ndim=1 norm=ortho  iFFT: n2=5.495357e-16 ninf=7.050855e-16 < 1.918887e-15 True
[...]
              (808,)    complex64 ndim=1 norm=    0   FFT: n2=6.011864e-07 ninf=7.164535e-07 < 2.162965e-06 True
              (808,)    complex64 ndim=1 norm=    0  iFFT: n2=6.040269e-07 ninf=7.427027e-07 < 2.162965e-06 True
              (808,)    complex64 ndim=1 norm=    1   FFT: n2=6.011864e-07 ninf=7.164535e-07 < 2.162965e-06 True
              (808,)    complex64 ndim=1 norm=    1  iFFT: n2=5.930688e-07 ninf=6.351446e-07 < 2.162965e-06 True
              (808,)    complex64 ndim=1 norm=ortho   FFT: n2=6.011864e-07 ninf=7.164535e-07 < 2.162965e-06 True
              (808,)    complex64 ndim=1 norm=ortho  iFFT: n2=6.040269e-07 ninf=7.427027e-07 < 2.162965e-06 True
              (808,)   complex128 ndim=1 norm=    0   FFT: n2=7.388737e-16 ninf=8.785366e-16 < 1.872223e-15 True
              (808,)   complex128 ndim=1 norm=    0  iFFT: n2=6.731794e-16 ninf=8.549905e-16 < 1.872223e-15 True
              (808,)   complex128 ndim=1 norm=    1   FFT: n2=7.388737e-16 ninf=8.785366e-16 < 1.872223e-15 True
              (808,)   complex128 ndim=1 norm=    1  iFFT: n2=4.587314e-12 ninf=4.587143e-12 < 1.872223e-15 False
              (808,)   complex128 ndim=1 norm=ortho   FFT: n2=7.388737e-16 ninf=8.785366e-16 < 1.872223e-15 True
              (808,)   complex128 ndim=1 norm=ortho  iFFT: n2=6.731794e-16 ninf=8.549905e-16 < 1.872223e-15 True
          (808, 808)    complex64 ndim=1 norm=    0   FFT: n2=6.119184e-07 ninf=7.602659e-07 < 3.325929e-06 True
          (808, 808)    complex64 ndim=1 norm=    0  iFFT: n2=6.122331e-07 ninf=6.903897e-07 < 3.325929e-06 True
          (808, 808)    complex64 ndim=1 norm=    1   FFT: n2=6.119184e-07 ninf=7.602659e-07 < 3.325929e-06 True
          (808, 808)    complex64 ndim=1 norm=    1  iFFT: n2=5.972787e-07 ninf=7.561319e-07 < 3.325929e-06 True
          (808, 808)    complex64 ndim=1 norm=ortho   FFT: n2=6.119184e-07 ninf=7.602659e-07 < 3.325929e-06 True
          (808, 808)    complex64 ndim=1 norm=ortho  iFFT: n2=6.122331e-07 ninf=6.903897e-07 < 3.325929e-06 True
          (808, 808)   complex128 ndim=1 norm=    0   FFT: n2=7.482011e-16 ninf=1.366856e-15 < 2.744447e-15 True
          (808, 808)   complex128 ndim=1 norm=    0  iFFT: n2=6.868058e-16 ninf=1.017457e-15 < 2.744447e-15 True
          (808, 808)   complex128 ndim=1 norm=    1   FFT: n2=7.482011e-16 ninf=1.366856e-15 < 2.744447e-15 True
          (808, 808)   complex128 ndim=1 norm=    1  iFFT: n2=4.587295e-12 ninf=4.587149e-12 < 2.744447e-15 False
          (808, 808)   complex128 ndim=1 norm=ortho   FFT: n2=7.482011e-16 ninf=1.366856e-15 < 2.744447e-15 True
          (808, 808)   complex128 ndim=1 norm=ortho  iFFT: n2=6.868058e-16 ninf=1.017457e-15 < 2.744447e-15 True
[...]

As you can see in the above only the Bluestein inverse fft with double precision (complex128) and norm=1 fail the accuracy test.

Note: the expected tolerance is currently written as 1e-6 + 4e-7 * np.log10(d0.size) for single precision and 1e-15 + 3e-16 * np.log10(d0.size) for double, this should have a fair margin. norm=ortho corresponds to norm=0 in VkFFT followed by a separate normalisation in pyvkfft.

I have only tested C2C transforms for this yet.

opened Dec 30, 2021 by vincefn 19
Closed

MetalFFT

I saw your comment in Apple Developer Forums and gave some insight. If you’re interested in carrying on my work, we can discuss here which is more convenient than Apple Developer Forums.

opened Feb 3, 2022 by philipturner 18
Closed

Titan V FP32 & FP64 results..

Hi, just tested your awesome library with a Titan V (Volta) which has only 2x slower FP64 vs FP32.. using NVIDIA 455.26.01 and CUDA 10.1 (can update CUDA SDK to 11.1 if requested: faster CUFFT?) Titan V: VKFFT single: titanv.txt CUFFT single: titanvcufft.txt VKFFT double: titanvfp64.txt CUFFT double: titanvcufftdouble.txt

roughly speaking your library is well optimized for double precision as it slowdowns 2x and similar score to CUFFT..

some perf issues in vkFFT in double precision vs single compared to CUFFT like for example in 4kx4kx8 case:

VkFFT System: 4096x4096x8 Buffer: 2048 MB avg_time_per_step: 75.931 ms cuFFT System: 4096x4096x8 Buffer: 2048 MB avg_time_per_step: 58.844 ms

note similar "big" double precision FFTs are not affected:

VkFFT System: 2048x256x256 Buffer: 2048 MB avg_time_per_step: 48.205 ms std_error: 0.329 batch: 1 cuFFT System: 2048x256x256 Buffer: 2048 MB avg_time_per_step: 45.925 ms std_error: 0.215 batch: 1

as you can see from fp32 case:

cuFFT System: 4096x4096x8 Buffer: 1024 MB avg_time_per_step: 29.235 ms std_error: 0.179 batch: 4 VkFFT System: 4096x4096x8 Buffer: 1024 MB avg_time_per_step: 29.923 ms std_error: 0.613 batch: 4

4kx4kx8 was performing similar to CUFFT in fp32 case : 29ms .. and CUFFT scales well from 29ms to 58ms..

as said 2kx256x256 scales well to fp64: goes from 23-24ms in fp32 to 46-48 ms in fp64 case

VkFFT System: 2048x256x256 Buffer: 1024 MB avg_time_per_step: 24.031 ms std_error: 0.662 batch: 4 benchmark: 43633 cuFFT System: 2048x256x256 Buffer: 1024 MB avg_time_per_step: 23.424 ms std_error: 0.085 batch: 4 benchmark: 44765

Plots: [EDIT: in next post]

opened Oct 11, 2020 by oscarbg 8
Closed

Cleaned up CMake file, added .gitignore and added shader compilation directly into the CMake file.

Cleaned up CMake file, added .gitignore and added shader compilation directly into the CMake file. There's now an option (default OFF) to compile the shaders via CMake, also option to manually specify the glslangvalidator path, default just uses glslangvalidator, like the bat file.

opened Aug 12, 2020 by Cazadorro 8
Closed

Questions on a potential SYCL port

I have been working quite a bit on FFTs in the past, but only started recently to develop for GPUs, using SYCL (since that's closest to the C++ world I'm used to, and promises to be somewhat more portable than CUDA and friends).

One thing that SYCL is sorely lacking is an FFT implementation, and I'm wondering if a port of VkFFT to SYCL is possible, at least in theory. As far as I understand, VkFFT currently generates GPU code on the fly for any requested FFT, whereas in SYCL all the functions would already have to be present during compilation, so that's a fairly major change.

Does the performance of VkFFT depend on the fact that you can hardwire lots of constants when you produce code for a requested FFT, or do you think the algorithms can perform well even without custom code generation on the fly?

In any case I am really looking forward to your planned publication! I'm sure there are many interesting optimization details which are only hard to spot when looking at the source code.

opened Mar 9, 2022 by mreineck 7
Closed

2D and 3D R2C/C2R get wrong results with HIP backend

Hello, I'm using vkfft with the hip backend. I extracted a simple test case from the source file Vulkan_FFT.cpp。On AMD MI50 GPU and ROCm 3.9/4.0，the 1D R2C/C2R results are right, but those of 2D and 3D are partially wrong. I'm not sure if it is caused by vkfft itself or I used it incorrectly. Here is the code:

#include<cstdio>
#include<cstdlib>

#include <hip/hip_runtime.h>
#include <hip/hiprtc.h>
#include <hip/hip_complex.h>
#include "vkFFT.h"

#include <chrono>

#define DIM 3
#define SIZE 52

template < typename T>
void initialData(T *ip, const int size, const int max_value)
{
    int i;

    for(i = 0; i < size; i++)
    {
        ip[i] = (T)(rand() % max_value);
    }
}

void checkResult(float *backRef, float *originalRef, const int N)
{
    double epsilon = 1.0E-2;
    bool match = 1;
    const float overN = 1.0 / pow(SIZE, DIM);

    for (int i = 0; i < N; i++)
    {
        if (fabs(backRef[i] * overN - originalRef[i]) > epsilon)
        {
            match = 0;
            printf("Arrays do not match!\n");
            printf("back result %5.2f original result %5.2f at %d\n", backRef[i]* overN, originalRef[i], i);
            break;
        }
    }

    if (match) printf("Arrays match.\n\n");
}

typedef struct {
    hipDevice_t device;
    hipCtx_t context;
    uint32_t device_id;
}VkGPU;

int main()
{
    hipError_t res = hipSuccess;
    VkGPU vkGPU = {};
    hipInit(0);
    if (res != hipSuccess) return VKFFT_ERROR_FAILED_TO_INITIALIZE;
    hipSetDevice(vkGPU.device_id);
    if (res != hipSuccess) return VKFFT_ERROR_FAILED_TO_SET_DEVICE_ID;
    hipDeviceGet(&vkGPU.device, vkGPU.device_id);
    if (res != hipSuccess) return VKFFT_ERROR_FAILED_TO_GET_DEVICE;
    hipCtxCreate(&vkGPU.context, 0, vkGPU.device);
    if (res != hipSuccess) return VKFFT_ERROR_FAILED_TO_CREATE_CONTEXT;
    VkFFTConfiguration configuration = {};
    VkFFTApplication app = {};
    configuration.FFTdim = 3;
    configuration.size[0] = SIZE;
    configuration.size[1] = SIZE;
    configuration.size[2] = SIZE;
    configuration.device = &(vkGPU.device);
    configuration.performR2C = true;
    uint64_t buffersize = sizeof(hipFloatComplex) * pow(SIZE, DIM);

    float *h_realGrid = (float *)malloc(sizeof(float) * pow(SIZE, DIM));
    float *h_realGrid_original = (float *)malloc(sizeof(float) * pow(SIZE, DIM));
    hipFloatComplex *h_complexGrid = (hipFloatComplex *)malloc(sizeof(hipFloatComplex) * pow(SIZE, DIM));

    hipFloatComplex *d_complexGrid;
    //hipMalloc((void **)&d_realGrid, sizeof(float) * pow(SIZE, DIM));
    hipMalloc((void **)&d_complexGrid, sizeof(hipFloatComplex) * pow(SIZE, DIM));
 
    configuration.buffer = (void**)&d_complexGrid;
    configuration.bufferSize = &buffersize;

    initialData(h_realGrid_original, pow(SIZE, DIM), 10);
    //for(int i = 0; i < pow(SIZE, DIM); ++i)
    //{
    //    h_complexGrid[i].x = h_realGrid_original[i];
    //    h_complexGrid[i].y = 0.0f;
    //}
    memset(h_complexGrid, 0, sizeof(hipFloatComplex) * pow(SIZE, DIM));
    memcpy(h_complexGrid, h_realGrid_original, sizeof(float) * pow(SIZE, DIM));

    hipMemcpy(d_complexGrid, h_complexGrid, sizeof(hipFloatComplex) * pow(SIZE, DIM), hipMemcpyHostToDevice);
 
    std::chrono::steady_clock::time_point timeSubmit = std::chrono::steady_clock::now(); 
    VkFFTResult resFFT;
    resFFT = initializeVkFFT(&app, configuration);
    if (resFFT != VKFFT_SUCCESS) 
    {
        printf("VkFFT error code %d\n", resFFT);
        return resFFT;
    }
    VkFFTLaunchParams launchParams = {};
  
 
    resFFT = VkFFTAppend(&app, -1, &launchParams); 
    if (resFFT != VKFFT_SUCCESS) 
    {
        printf("vkfft error code %d\n", resFFT);
        return resFFT;
    }
    resFFT = VkFFTAppend(&app, 1, &launchParams);
    if (resFFT != VKFFT_SUCCESS) 
    {
        printf("vkfft error code %d\n", resFFT);
        return resFFT;
    }
    res = hipDeviceSynchronize();
    if (res != hipSuccess) 
    {
        printf("hip error code %d\n", res);
        return VKFFT_ERROR_FAILED_TO_SYNCHRONIZE;
    }
    
    std::chrono::steady_clock::time_point timeEnd = std::chrono::steady_clock::now();
    float totTime = std::chrono::duration_cast<std::chrono::microseconds>(timeEnd - timeSubmit).count() * 0.001; 
    printf("time: %.3f ms\n", totTime);
    
    hipMemcpy(h_complexGrid, d_complexGrid, sizeof(hipFloatComplex) * pow(SIZE, DIM), hipMemcpyDeviceToHost);
    //for(int i = 0; i < pow(SIZE, DIM); ++i)
    //{
    //    h_realGrid[i] = h_complexGrid[i].x;
    //}
    memcpy(h_realGrid, h_complexGrid, sizeof(float) * pow(SIZE, DIM));
 
    printf("h_realGrid: ");
    for(int i = 0; i < 30; ++i)
    {
         printf("%.6f  ", h_realGrid[i] / pow(SIZE, DIM));
    }
    printf("\nh_realGrid_original: ");
    for(int i = 0; i < 30; ++i)
    {
         printf("%.6f  ", h_realGrid_original[i]);
    }

    printf("\n");
    checkResult(h_realGrid, h_realGrid_original, pow(SIZE, DIM));

    deleteVkFFT(&app);
    free(h_realGrid);
    free(h_realGrid_original);
    free(h_complexGrid);

    //hipFree(d_realGrid);
    hipFree(d_complexGrid);
    
    return 0;
}

I compile it using the command:

hipcc -o planMany-vkfft -DVKFFT_BACKEND=2  planMany-vkfft.cpp

and my results:

time: 11023.094 ms
h_realGrid: 2.999996  5.999995  6.999996  4.999996  2.999997  4.999995  5.999996  1.999996  8.999998  0.999997  1.999998  6.999994  -0.000003  8.999997  2.999996  5.999996  -0.000004  5.999995  1.999997  5.999997  0.999996  7.999999  6.999996  8.999996  1.999997  -0.000003  1.999996  2.999999  6.999996  4.999996
h_realGrid_original: 3.000000  6.000000  7.000000  5.000000  3.000000  5.000000  6.000000  2.000000  9.000000  1.000000  2.000000  7.000000  0.000000  9.000000  3.000000  6.000000  0.000000  6.000000  2.000000  6.000000  1.000000  8.000000  7.000000  9.000000  2.000000  0.000000  2.000000  3.000000  7.000000  5.000000
back result -0.46 original result  1.00 at 52
back result  0.08 original result  7.00 at 106
back result  0.00 original result  2.00 at 107
back result -0.23 original result  9.00 at 160
back result -0.00 original result  8.00 at 161
back result  0.35 original result  7.00 at 214
back result  0.00 original result  6.00 at 215
back result -0.31 original result  2.00 at 268
back result  0.00 original result  4.00 at 269
back result  0.56 original result  3.00 at 322
back result -0.00 original result  2.00 at 323
back result -0.02 original result  0.00 at 376
back result -0.00 original result  5.00 at 377
......

Could you please help me with this problem? Thank you!

opened May 7, 2021 by ipe-zhangyz 7
Open

vkFFT changes locale

Hi,

as vkFFT uses sprintf for generating the shader code, it is dependent on the locale that is set globally for the program. To circumvent printing e.g. commas for floating point numbers given certain locales, vkFFT currently changes the locale "temporarily" using setlocale. But it looks like the implementation for this is bugged. Starting from this locale (result of setlocale(LC_ALL, 0);) before using vkFFT: LC_COLLATE=C;LC_CTYPE=German_Germany.1252;LC_MONETARY=C;LC_NUMERIC=C;LC_TIME=C it is set to C as soon as a vkFFT plan is initialized through initializeVkFFT and stays like that. Sometimes it switches back to German_Germany.1252 after initializing another vkFFT plan, but it still is not the locale it was before and kind of random if that happens or not.

I think the problem is storing a pointer to the result of setlocale as oldLocale - which probably is a bad idea, as this is not a memory location under vkFFTs control, so it likely is changed whenever setlocale is called again. Probably storing a copy of the string it returns would already fix some problems. But generally, I'd like to get rid of setting the locale in vkFFT, as this is a general problem in applications, that might rely on the globally set locale in multiple threads.

With the common sprintf implementation this seems to be not possible (though there are platform dependent versions that accept a locale as parameter, like _sprintf_s_l on windows). But for vkFFT it should be simple to at least optionally replace the sprintf calls with e.g. fmt::sprintf or stbsp_sprintf, both of which are locale agnostic.

If you are interested, I prototyped this for fmt::sprintf, which works fine so far. Maybe you could include something into the library to drop in a sprintf replacement e.g. via some define?

Thanks!

The prototypical modifications:

// some cpp using vkfft:

#include <fmt/printf.h>
#include <utility>

template<typename... Ts>
int vkfftSprintfReplacement(char * str, char const *fmt, Ts&&... args) {
    std::string formatted = fmt::sprintf(fmt, std::forward<Ts>(args)...);
    std::copy(formatted.begin(), formatted.end(), str);
    auto result = formatted.length();
    str[result] = '\0';
    return (int)result;
}

#define VKFFT_REPLACE_SPRINTF vkfftSprintfReplacement

#include <vkFFT.h>

// code ...

// in vkFFT.h:

// vkFFT.h includes go here

#ifdef VKFFT_REPLACE_SPRINTF
# define sprintf VKFFT_REPLACE_SPRINTF
#endif

// vkFFT.h body goes here, the setlocale calls are guarded with #ifndef VKFFT_REPLACE_SPRINTF

// vkFFT.h at the very bottom

#ifdef VKFFT_REPLACE_SPRINTF
# undef sprintf
#endif
opened May 19, 2022 by pborsutzki 1
Open

Fail to compile

hi, it seems increasing kernel sizes any number above 8192 floats (w 2 batches) causes a parsing error (VKFFT_ERROR_FAILED_SHADER_PARSE) for batched convolutions. i've reproduced this in many setting but perhaps the simplest is by changing the kernel size in sample #52 (code below). Below is also the exact error msg. I tried checking any publicized limits but didnt find any limit that applies (hardware: AMD Radeon Pro 5300M 4 GB, macos). Any help would be much appreciated, thank you.

  • error msg

#version 450

layout (local_size_x = 128, local_size_y = 1, local_size_z = 1) in; const float loc_PI = 3.1415926535897932384626433832795f; const float loc_SQRT1_2 = 0.70710678118654752440084436210485f; layout(std430, binding = 0) buffer DataIn{ vec2 inputs[16386]; };

layout(std430, binding = 1) buffer DataOut{ vec2 outputs[16386]; };

layout(std430, binding = 2) buffer Kernel_FFT{ vec2 kernel_obj[16386]; };

void main() { uint id_x = gl_GlobalInvocationID.x % 4096; uint id_y = (gl_GlobalInvocationID.x / 4096) % 1; uint id_z = (gl_GlobalInvocationID.x / 4096) / 1; if (gl_GlobalInvocationID.x < 4096){ uint inoutID = (id_x + id_y8193 +id_z8193) + (null) * 8193; uint inoutID2; uint inoutID3; vec2 t0 = inputs[inoutID]; vec2 tf; if (id_x == 0) { inoutID2 = (8192 + id_y8193 +id_z8193) + (null) * 8193; inoutID3 = (4096 + id_y8193 +id_z8193) + (null) * 8193; tf = inputs[inoutID3]; } else { inoutID2 = ((8192-id_x) + id_y8193 +id_z8193) + (null) * 8193; } vec2 t1 = inputs[inoutID2]; vec2 t2; vec2 t3; if (id_x == 0) { t2.x = (t0.x+t1.x); t2.y = (t0.x-t1.x); tf.y = -tf.y; tf.x = tf.x * 2; tf.y = tf.y * 2; outputs[inoutID] = t2; outputs[inoutID3] = tf; } else { t2.x = t0.x + t1.x; t2.y = t0.y + t1.y; t3.x = t0.x - t1.x; t3.y = t0.y - t1.y; float angle = (loc_PIid_x)/8192; tf.x = cos(angle); tf.y = sin(angle); t0.x = tf.xt2.y+tf.yt3.x; t0.y = -tf.yt2.y+tf.x*t3.x; t1.x = t2.x+t0.x; t1.y = -t3.y+t0.y; t0.x = t2.x-t0.x; t0.y = t3.y+t0.y; outputs[inoutID] = t0; outputs[inoutID2] = t1; } } }

ERROR: 0:23: 'null' : undeclared identifier ERROR: 0:23: '=' : cannot convert from ' temp float' to ' temp highp uint' ERROR: 0:23: '' : compilation terminated ERROR: 3 compilation errors. No code generated.

VkFFT shader type: 0

  • error reproduction by changing sample 52:

VkFFTResult resFFT = VKFFT_SUCCESS; VkResult res = VK_SUCCESS; //Configuration + FFT application. VkFFTConfiguration configuration = {}; VkFFTConfiguration convolution_configuration = {}; VkFFTApplication app_convolution = {}; VkFFTApplication app_kernel = {}; //Convolution sample code //Setting up FFT configuration. FFT is performed in-place with no performance loss.

std::cout << "size:" << 2*real_kernel.size() << std::endl;

configuration.FFTdim = 1; //FFT dimension, 1D, 2D or 3D (default 1).
configuration.size[0] = 8192*2; //Multidimensional FFT dimensions sizes (default 1). For best performance (and stability), order dimensions in descendant size order as: x>y>z.
configuration.size[1] = 1;
configuration.size[2] = 1;

configuration.kernelConvolution = true; //specify if this plan is used to create kernel for convolution
configuration.performR2C = true; //Perform R2C/C2R transform. Can be combined with all other options. Reduces memory requirements by a factor of 2. Requires special input data alignment: for x*y*z system pad x*y plane to (x+2)*y with last 2*y elements reserved, total array dimensions are (x*y+2y)*z. Memory layout after R2C and before C2R can be found on github.
configuration.coordinateFeatures = 1; //Specify dimensionality of the input feature vector (default 1). Each component is stored not as a vector, but as a separate system and padded on it's own according to other options (i.e. for x*y system of 3-vector, first x*y elements correspond to the first dimension, then goes x*y for the second, etc).
//coordinateFeatures number is an important constant for convolution. If we perform 1x1 convolution, it is equal to number of features, but matrixConvolution should be equal to 1. For matrix convolution, it must be equal to matrixConvolution parameter. If we perform 2x2 convolution, it is equal to 3 for symmetric kernel (stored as xx, xy, yy) and 4 for nonsymmetric (stored as xx, xy, yx, yy). Similarly, 6 (stored as xx, xy, xz, yy, yz, zz) and 9 (stored as xx, xy, xz, yx, yy, yz, zx, zy, zz) for 3x3 convolutions.
configuration.normalize = 1;//normalize iFFT

configuration.numberBatches = 2;
//After this, configuration file contains pointers to Vulkan objects needed to work with the GPU: VkDevice* device - created device, [uint64_t *bufferSize, VkBuffer *buffer, VkDeviceMemory* bufferDeviceMemory] - allocated GPU memory FFT is performed on. [uint64_t *kernelSize, VkBuffer *kernel, VkDeviceMemory* kernelDeviceMemory] - allocated GPU memory, where kernel for convolution is stored.
configuration.device = &vkGPU->device;

#if(VKFFT_BACKEND==0) configuration.queue = &vkGPU->queue; //to allocate memory for LUT, we have to pass a queue, vkGPU->fence, commandPool and physicalDevice pointers configuration.fence = &vkGPU->fence; configuration.commandPool = &vkGPU->commandPool; configuration.physicalDevice = &vkGPU->physicalDevice; configuration.isCompilerInitialized = 1;//compiler can be initialized before VkFFT plan creation. if not, VkFFT will create and destroy one after initialization #elif(VKFFT_BACKEND==3) configuration.platform = &vkGPU->platform; configuration.context = &vkGPU->context; #endif //In this example, we perform a convolution for a real vectorfield (3vector) with a symmetric kernel (6 values). We use configuration to initialize convolution kernel first from real data, then we create convolution_configuration for convolution. The buffer object from configuration is passed to convolution_configuration as kernel object. //1. Kernel forward FFT. uint64_t kernelSize = ((uint64_t)configuration.numberBatches) * configuration.coordinateFeatures * sizeof(float) * 2 * (configuration.size[0] / 2 + 1) * configuration.size[1] * configuration.size[2];;

#if(VKFFT_BACKEND==0) VkBuffer kernel = {}; VkDeviceMemory kernelDeviceMemory = {}; resFFT = allocateBuffer(vkGPU, &kernel, &kernelDeviceMemory, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT, VK_MEMORY_HEAP_DEVICE_LOCAL_BIT, kernelSize); if (resFFT != VKFFT_SUCCESS) return resFFT; configuration.buffer = &kernel; #elif(VKFFT_BACKEND==1) cuFloatComplex* kernel = 0; res = cudaMalloc((void**)&kernel, kernelSize); if (res != cudaSuccess) return VKFFT_ERROR_FAILED_TO_ALLOCATE; configuration.buffer = (void**)&kernel; #elif(VKFFT_BACKEND==2) hipFloatComplex* kernel = 0; res = hipMalloc((void**)&kernel, kernelSize); if (res != hipSuccess) return VKFFT_ERROR_FAILED_TO_ALLOCATE; configuration.buffer = (void**)&kernel; #elif(VKFFT_BACKEND==3) cl_mem kernel = 0; kernel = clCreateBuffer(vkGPU->context, CL_MEM_READ_WRITE, kernelSize, 0, &res); if (res != CL_SUCCESS) return VKFFT_ERROR_FAILED_TO_ALLOCATE; configuration.buffer = &kernel; #endif

configuration.bufferSize = &kernelSize;

printf("Total memory needed for kernel: %" PRIu64 " MB\n", kernelSize / 1024 / 1024);

//Fill kernel on CPU.
float* kernel_input = (float*)malloc(kernelSize);
if (!kernel_input) return VKFFT_ERROR_MALLOC_FAILED;
for (uint64_t f = 0; f < configuration.numberBatches; f++) {
    for (uint64_t v = 0; v < configuration.coordinateFeatures; v++) {
        for (uint64_t k = 0; k < configuration.size[2]; k++) {
            for (uint64_t j = 0; j < configuration.size[1]; j++) {

                //Below is the test identity kernel for 1x1 nonsymmetric FFT, multiplied by (f * configuration.coordinateFeatures + v + 1);
                for (uint64_t i = 0; i < configuration.size[0] / 2 + 1; i++) {

                    kernel_input[2 * i + j * (configuration.size[0] + 2) + k * (configuration.size[0] + 2) * configuration.size[1] + v * (configuration.size[0] + 2) * configuration.size[1] * configuration.size[2] + f * configuration.coordinateFeatures * (configuration.size[0] + 2) * configuration.size[1] * configuration.size[2]] = (float)(f * configuration.coordinateFeatures + v + 1.0);
                    kernel_input[2 * i + 1 + j * (configuration.size[0] + 2) + k * (configuration.size[0] + 2) * configuration.size[1] + v * (configuration.size[0] + 2) * configuration.size[1] * configuration.size[2] + f * configuration.coordinateFeatures * (configuration.size[0] + 2) * configuration.size[1] * configuration.size[2]] = 0;

                }
            }
        }
    }
}
//Sample buffer transfer tool. Uses staging buffer of the same size as destination buffer, which can be reduced if transfer is done sequentially in small buffers.

#if(VKFFT_BACKEND==0) resFFT = transferDataFromCPU(vkGPU, kernel_input, &kernel, kernelSize); if (resFFT != VKFFT_SUCCESS) return resFFT; #elif(VKFFT_BACKEND==1) res = cudaMemcpy(kernel, kernel_input, kernelSize, cudaMemcpyHostToDevice); if (res != cudaSuccess) return VKFFT_ERROR_FAILED_TO_COPY; #elif(VKFFT_BACKEND==2) res = hipMemcpy(kernel, kernel_input, kernelSize, hipMemcpyHostToDevice); if (res != hipSuccess) return VKFFT_ERROR_FAILED_TO_COPY; #elif(VKFFT_BACKEND==3) res = clEnqueueWriteBuffer(vkGPU->commandQueue, kernel, CL_TRUE, 0, kernelSize, kernel_input, 0, NULL, NULL); if (res != CL_SUCCESS) return VKFFT_ERROR_FAILED_TO_COPY; #endif //Initialize application responsible for the kernel. This function loads shaders, creates pipeline and configures FFT based on configuration file. No buffer allocations inside VkFFT library. resFFT = initializeVkFFT(&app_kernel, configuration); if (resFFT != VKFFT_SUCCESS) return resFFT; //Sample forward FFT command buffer allocation + execution performed on kernel. Second number determines how many times perform application in one submit. FFT can also be appended to user defined command buffers.

//Uncomment the line below if you want to perform kernel FFT. In this sample we use predefined identitiy kernel.
//performVulkanFFT(vkGPU, &app_kernel, -1, 1);

//The kernel has been trasnformed.


//2. Buffer convolution with transformed kernel.
//Copy configuration, as it mostly remains unchanged. Change specific parts.
convolution_configuration = configuration;
convolution_configuration.kernelConvolution = false;
convolution_configuration.performConvolution = true;
convolution_configuration.symmetricKernel = false;//Specify if convolution kernel is symmetric. In this case we only pass upper triangle part of it in the form of: (xx, xy, yy) for 2d and (xx, xy, xz, yy, yz, zz) for 3d.

#if(VKFFT_BACKEND==0) convolution_configuration.kernel = &kernel; #elif(VKFFT_BACKEND==1) convolution_configuration.kernel = (void**)&kernel; #elif(VKFFT_BACKEND==2) convolution_configuration.kernel = (void**)&kernel; #elif(VKFFT_BACKEND==3) convolution_configuration.kernel = &kernel; #endif

convolution_configuration.kernelSize = &kernelSize;
convolution_configuration.numberBatches = 1;//one batch - numberKernels convolutions
convolution_configuration.numberKernels = configuration.numberBatches;// number of convolutions on a single input
//Allocate separate buffer for the input data.
uint64_t inputBufferSize = ((uint64_t)convolution_configuration.coordinateFeatures) * sizeof(float) * 2 * (convolution_configuration.size[0] / 2 + 1) * convolution_configuration.size[1] * convolution_configuration.size[2];;
uint64_t bufferSize = convolution_configuration.numberKernels * convolution_configuration.coordinateFeatures * sizeof(float) * 2 * (convolution_configuration.size[0] / 2 + 1) * convolution_configuration.size[1] * convolution_configuration.size[2];;
convolution_configuration.isInputFormatted = true; //if input is a different buffer, it doesn't have to be zeropadded/R2C padded

#if(VKFFT_BACKEND==0) VkBuffer inputBuffer = {}; VkBuffer buffer = {}; VkDeviceMemory inputBufferDeviceMemory = {}; VkDeviceMemory bufferDeviceMemory = {}; resFFT = allocateBuffer(vkGPU, &inputBuffer, &inputBufferDeviceMemory, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT, VK_MEMORY_HEAP_DEVICE_LOCAL_BIT, inputBufferSize); if (resFFT != VKFFT_SUCCESS) return resFFT; resFFT = allocateBuffer(vkGPU, &buffer, &bufferDeviceMemory, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT, VK_MEMORY_HEAP_DEVICE_LOCAL_BIT, bufferSize); if (resFFT != VKFFT_SUCCESS) return resFFT; convolution_configuration.inputBuffer = &inputBuffer; convolution_configuration.buffer = &buffer; #elif(VKFFT_BACKEND==1) cuFloatComplex* inputBuffer = 0; cuFloatComplex* buffer = 0; res = cudaMalloc((void**)&inputBuffer, inputBufferSize); if (res != cudaSuccess) return VKFFT_ERROR_FAILED_TO_ALLOCATE; res = cudaMalloc((void**)&buffer, bufferSize); if (res != cudaSuccess) return VKFFT_ERROR_FAILED_TO_ALLOCATE; convolution_configuration.inputBuffer = (void**)&inputBuffer; convolution_configuration.buffer = (void**)&buffer; #elif(VKFFT_BACKEND==2) hipFloatComplex* inputBuffer = 0; hipFloatComplex* buffer = 0; res = hipMalloc((void**)&inputBuffer, inputBufferSize); if (res != hipSuccess) return VKFFT_ERROR_FAILED_TO_ALLOCATE; res = hipMalloc((void**)&buffer, bufferSize); if (res != hipSuccess) return VKFFT_ERROR_FAILED_TO_ALLOCATE; convolution_configuration.inputBuffer = (void**)&inputBuffer; convolution_configuration.buffer = (void**)&buffer; #elif(VKFFT_BACKEND==3) cl_mem inputBuffer = 0; cl_mem buffer = 0; inputBuffer = clCreateBuffer(vkGPU->context, CL_MEM_READ_WRITE, inputBufferSize, 0, &res); if (res != CL_SUCCESS) return VKFFT_ERROR_FAILED_TO_ALLOCATE; buffer = clCreateBuffer(vkGPU->context, CL_MEM_READ_WRITE, bufferSize, 0, &res); if (res != CL_SUCCESS) return VKFFT_ERROR_FAILED_TO_ALLOCATE; convolution_configuration.inputBuffer = &inputBuffer; convolution_configuration.buffer = &buffer; #endif

convolution_configuration.inputBufferSize = &inputBufferSize;
convolution_configuration.bufferSize = &bufferSize;


 printf("Total memory needed for buffer: %" PRIu64 " MB\n", bufferSize / 1024 / 1024);
//Fill data on CPU. It is best to perform all operations on GPU after initial upload.
float* buffer_input = (float*)malloc(inputBufferSize);
if (!buffer_input) return VKFFT_ERROR_MALLOC_FAILED;
for (uint64_t v = 0; v < convolution_configuration.coordinateFeatures; v++) {
    for (uint64_t k = 0; k < convolution_configuration.size[2]; k++) {
        for (uint64_t j = 0; j < convolution_configuration.size[1]; j++) {
            for (uint64_t i = 0; i < convolution_configuration.size[0]; i++) {
                buffer_input[i + j * (convolution_configuration.size[0] + 2) + k * (convolution_configuration.size[0] + 2) * convolution_configuration.size[1] + v * (convolution_configuration.size[0] + 2) * convolution_configuration.size[1] * convolution_configuration.size[2]] = 1;
            }
        }
    }
}
//Transfer data to GPU using staging buffer.

#if(VKFFT_BACKEND==0) resFFT = transferDataFromCPU(vkGPU, buffer_input, &inputBuffer, inputBufferSize); if (resFFT != VKFFT_SUCCESS) return resFFT; #elif(VKFFT_BACKEND==1) res = cudaMemcpy(inputBuffer, buffer_input, inputBufferSize, cudaMemcpyHostToDevice); if (res != cudaSuccess) return VKFFT_ERROR_FAILED_TO_COPY; #elif(VKFFT_BACKEND==2) res = hipMemcpy(inputBuffer, buffer_input, inputBufferSize, hipMemcpyHostToDevice); if (res != hipSuccess) return VKFFT_ERROR_FAILED_TO_COPY; #elif(VKFFT_BACKEND==3) res = clEnqueueWriteBuffer(vkGPU->commandQueue, inputBuffer, CL_TRUE, 0, inputBufferSize, buffer_input, 0, NULL, NULL); if (res != CL_SUCCESS) return VKFFT_ERROR_FAILED_TO_COPY; #endif

//Initialize application responsible for the convolution.
resFFT = initializeVkFFT(&app_convolution, convolution_configuration);
if (resFFT != VKFFT_SUCCESS) return resFFT;
//Sample forward FFT command buffer allocation + execution performed on kernel. FFT can also be appended to user defined command buffers.
VkFFTLaunchParams launchParams = {};
resFFT = performVulkanFFT(vkGPU, &app_convolution, &launchParams, -1, 1);
if (resFFT != VKFFT_SUCCESS) return resFFT;
//The kernel has been trasnformed.

float* buffer_output = (float*)malloc(bufferSize);
if (!buffer_output) return VKFFT_ERROR_MALLOC_FAILED;
//Transfer data from GPU using staging buffer.

#if(VKFFT_BACKEND==0) resFFT = transferDataToCPU(vkGPU, buffer_output, &buffer, bufferSize); if (resFFT != VKFFT_SUCCESS) return resFFT; #elif(VKFFT_BACKEND==1) res = cudaMemcpy(buffer_output, buffer, bufferSize, cudaMemcpyDeviceToHost); if (res != cudaSuccess) return VKFFT_ERROR_FAILED_TO_COPY; #elif(VKFFT_BACKEND==2) res = hipMemcpy(buffer_output, buffer, bufferSize, hipMemcpyDeviceToHost); if (res != hipSuccess) return VKFFT_ERROR_FAILED_TO_COPY; #elif(VKFFT_BACKEND==3) res = clEnqueueReadBuffer(vkGPU->commandQueue, buffer, CL_TRUE, 0, bufferSize, buffer_output, 0, NULL, NULL); if (res != CL_SUCCESS) return VKFFT_ERROR_FAILED_TO_COPY; #endif

//Print data, if needed.
for (uint64_t f = 0; f < convolution_configuration.numberKernels; f++) {
    printf("\nKernel id: %" PRIu64 "\n\n", f);
    for (uint64_t v = 0; v < convolution_configuration.coordinateFeatures; v++) {
        printf("\ncoordinate: %" PRIu64 "\n\n", v);
        for (uint64_t k = 0; k < convolution_configuration.size[2]; k++) {
            for (uint64_t j = 0; j < convolution_configuration.size[1]; j++) {
                for (uint64_t i = 0; i < convolution_configuration.size[0]; i++) {

                    printf("%.6f ", buffer_output[i + j * (convolution_configuration.size[0] + 2) + k * (convolution_configuration.size[0] + 2) * convolution_configuration.size[1] + v * (convolution_configuration.size[0] + 2) * convolution_configuration.size[1] * convolution_configuration.size[2] + f * convolution_configuration.coordinateFeatures * (convolution_configuration.size[0] + 2) * convolution_configuration.size[1] * convolution_configuration.size[2]]);
                }
                std::cout << "\n";
            }
        }
    }
}
free(kernel_input);
free(buffer_input);
free(buffer_output);

#if(VKFFT_BACKEND==0) vkDestroyBuffer(vkGPU->device, inputBuffer, NULL); vkFreeMemory(vkGPU->device, inputBufferDeviceMemory, NULL); vkDestroyBuffer(vkGPU->device, buffer, NULL); vkFreeMemory(vkGPU->device, bufferDeviceMemory, NULL); vkDestroyBuffer(vkGPU->device, kernel, NULL); vkFreeMemory(vkGPU->device, kernelDeviceMemory, NULL); #elif(VKFFT_BACKEND==1) cudaFree(inputBuffer); cudaFree(buffer); cudaFree(kernel); #elif(VKFFT_BACKEND==2) hipFree(inputBuffer); hipFree(buffer); hipFree(kernel); #elif(VKFFT_BACKEND==3) clReleaseMemObject(inputBuffer); clReleaseMemObject(buffer); clReleaseMemObject(kernel); #endif deleteVkFFT(&app_kernel); deleteVkFFT(&app_convolution); return resFFT;

opened May 18, 2022 by ribeur 2
Open

Batch Mode Data Ordering?

Hello,

I'm very interested in adding VkFFT to my project here: https://github.com/spinicist/riesling as an addition alongside FFTW.

However I'm confused about the data orderings supported for batched FFTs in VkFFT. Currently my batch dimension is the innermost/fastest varying dimension for my input array, and the FFT dimensions are the outermost. You can see how I set up an FFTW plan here: https://github.com/spinicist/riesling/blob/main/src/fft/cpu.hpp#L58 - note I specify N FFTs that are spaced 1 element apart in memory.

Does VkFFT support the same data ordering? I have looked at the docs but confess I am confused as to how batch mode works - they way I read them, it seems the batches have to be the outermost dimension. I am hoping I misunderstood.

Thanks in advance.

opened May 7, 2022 by spinicist 4
Open

Add OpenCL event management

Here we create a chain of events that starts with a marker event that is started when VkFFTAppend is called and each opencl kernel invocation depends on the previous event. If the user gives an address for an event, the last event is returned to the user.

cc @zachjweiner, @yves-surrel, @inducer

Fixes #39

opened May 1, 2022 by isuruf 1
Open

OpenGL compute shaders

It'd be nice to have an OpenGL version, enabling comparison vs Vulkan drivers for compute

opened Mar 28, 2022 by LifeIsStrange 1
Open

2D C2C Convolution

Hi there,

Apologies in advance if this is not the proper platform for my question.

I'm trying to have vkfft perform a simple 2D C2C convolution and leverage the no-reorder and zero-padding optimizations. I'm confused by the documentation and all the examples that show some form of matrix x vector convolution.

In my application, the kernel is not padded (the circular convolution is desired), but the input data is. I haven't enabled zeropadding optimization yet, but instead pad the input data beforehand.

The kernel generation and actual convolution are done in 2 separate steps, by the same function: I re-use the same plan for kernel generation and convolution. This might be my first mistake?

The function sets up the dimension, sizes, FFT precision, C2C, the same way in both cases.

I have the following, pre-allocated, device buffers:

  • VkBuffer devBuffer; // FFT buffer
  • VkBuffer kerBuffer; // Kernel buffer

1. Kernel generation

configuration.kernelConvolution = true; configuration.coordinateFeatures = 1; configuration.buffer = &devBuffer;

performVulkanFFT(vkGPU, &app, &launchParams, -1, 1);

copyBuffer(devBuffer, kerBuffer);

2. Convolution

configuration.kernelConvolution = false; configuration.performConvolution = true; configuration.symmetricKernel = false; configuration.coordinateFeatures = 1; configuration.conjugateConvolution = 0; configuration.matrixConvolution = 0; configuration.crossPowerSpectrumNormalization = 0; configuration.numberKernels = 1;

configuration.buffer = &devBuffer; configuration.kernel = &kerBuffer;

resFFT = performVulkanFFT(vkGPU, &app, &launchParams, -1, 1);

Questions

Anything wrong with sharing the same plan for kernel generation and convolution? Could I enable zeropadding optimization without rebuilding the plan? (The input data would be half the size (in both dimensions) as the allocated buffer and the 0-padding would fill up the rest) Shouldn't I run performVulkanFFTiFFT, instead? Does VulkanFFT perform the inverse transfrom automatically when .performConvolution = true?

On a different note: any plan to support zero-padding around the data, instead of only before/after?

opened Mar 10, 2022 by Clemasteredge 16

Contents

    fotisk07 An implementation of the Fourier Transform using Python
    zongyi-li Use Fourier transform to learn operators in differential equations.
    dfm JAX bindings to the Flatiron Institute Non-uniform Fast Fourier Transform (FINUFFT) library
    JodiTheTigger A simple, C99, header only, 0-Clause BSD Licensed, fast fourier transform (FFT).
    gasgiant Fast Fourier Transform ocean waves simulation for Unity.
    nihui waifu2x converter ncnn version, runs fast on intel / amd / nvidia GPU with vulkan
    chalharu chfft - Fastest Fourier Transform library implemented with pure Rust
    simonboots A Fourier Series visualisation written in Swift/SwiftUI
    aliemen Here is some Python code that allows you to read in SVG files and approximate their paths using a Fourier series.
    GPUOpen-LibrariesAndSDKs V-EZ is designed to be a C based light-weight layer around Vulkan
    hexops mach-glfw Vulkan example
    EQMG A high speed C++17 Vulkan game engine
    jgbit VUDA is a header-only library based on Vulkan that provides a CUDA Runtime API interface for writing GPU-accelerated applications.
    philipturner 1D, 2D, and 3D variations of Fast Fourier Transforms for a Metal S4TF backend
    nvpro-pipeline Experimental High Level Framework for Vulkan
    gabdube Standalone python program that draws a triangle using vulkan on Windows and Linux
    fynnfluegge OpenGL/Vulkan Java 3D Engine
    gabdube A 3D rendering demo powered by Python and Vulkan
    gwihlidal vk-sync - Simplification of core Vulkan synchronization mechanisms such as pipeline barriers and events
    sebbbi Rust & Vulkan test projects
    nvpro-samples Vulkan video samples
    Share this repo

    Related Repos


    Miscellaneous
    420
    For recording and retrieving metadata associated with ML developer and data scientist work...
    google ML Metadata ML Metadata (MLMD) is a library for recording and retrieving metadata associated with ML developer and data scientist workflows. NOTE: ML
     
    Miscellaneous
    100
    Hotcaml: an interpreter with watching and reloading
    let-def Hotcaml: an OCaml interpreter with watching and reloading Hotcaml is an OCaml interpreter that starts from a source file and loads its dependencies. W
     
    Miscellaneous
    50
    A React Native library for accessing an ArrayBuffer of a Blob instance.
    mrousavy react-native-blob-jsi-helper A React Native library for directly accessing an ArrayBuffer of a Blob instance. Note: This is a workaround. A long-term
     
    Miscellaneous
    14
    This is a OOSDK port of flat's ps4_remote_pkg_installer, all credit goes to them and their...
    Backporter ps4_remote_pkg_installer-OOSDK This is a OOSDK port of flat's ps4_remote_pkg_installer, all credit goes to them and their amazing work! NOTES Some fun
     
    Miscellaneous
    731
    General purpose GPU compute framework for cross vendor graphics cards (AMD, Qualcomm, NVID...
    KomputeProject General purpose GPU compute framework for cross vendor graphics cards (AMD, Qualcomm, NVIDIA & friends). Blazing fast, mobile-enabled, asynchronous and optimized for advanced GPU data processing usecases based on Vulkan compute. Backed by the Linux Foundation.
     
    Miscellaneous
    31
    An open source machine learning library for performing regression tasks using RVM techniqu...
    siavashserver Introduction neonrvm is an open source machine learning library for performing regression tasks using RVM technique. It is written in C programming la
     
    Miscellaneous
    108
    Pure C ONNX runtime with zero dependancies for embedded devices
    alrevuelta 🤖 cONNXr C ONNX Runtime A onnx runtime written in pure C99 with zero dependencies focused on embedded devices. Run inference on your machine learning
     
    Miscellaneous
    125
    Multi-armed bandit simulation library
    jkomiyama BanditLib: a simple multi-armed bandit library .-. .-') ('-. .-') _ _ .-') _ .-') _ .-. .-') \ ( OO )
     
    Miscellaneous
    466
    oneAPI Data Analytics Library (oneDAL)
    oneapi-src oneAPI Data Analytics Library Installation   |   Documentation   |   Support   |   Examples   |   Samples   |   How to Contribute    oneAPI Data Analy
     
    Miscellaneous
    1.4k
    Library for factorization machines
    srendle libFM Library for factorization machines web: http://www.libfm.org/ forum: https://groups.google.com/forum/#!forum/libfm Factorization machines (FM) a
     
    Miscellaneous
    617
    ThunderGBM: Fast GBDTs and Random Forests on GPUs
    Xtra-Computing Documentations | Installation | Parameters | Python (scikit-learn) interface What's new? ThunderGBM won 2019 Best Paper Award from IEEE Transactions o
     
    Miscellaneous
    34
    Low dependency（C++11 STL only）, good portability, header-only, deep neural networks for em...
    mosdeo LKYDeepNN LKYDeepNN 可訓練的深度類神經網路 (Deep Neural Network) 函式庫。 輕量，核心部份只依賴 C++11 標準函式庫，低相依性、好移植，方便在嵌入式系統上使用。 Class diagram 附有訓練視覺化 demo 程式 訓練視覺化程式以 OpenCV
     
    Miscellaneous
    115
    Colibri core is an NLP tool as well as a C++ and Python library for working with basic lin...
    proycon Colibri Core is software to quickly and efficiently count and extract patterns from large corpus data, to extract various statistics on the extracted patterns, and to compute relations between the extracted patterns.
     
    Miscellaneous
    67
    Frog is an integration of memory-based natural language processing (NLP) modules developed...
    LanguageMachines Frog is an integration of memory-based natural language processing (NLP) modules developed for Dutch. All NLP modules are based on Timbl, the Tilburg memory-based learning software package.
     
    Miscellaneous
    14
    FoLiA library for C++
    LanguageMachines Libfolia: FoLiA Library for C++ Libfolia (c) CLS/ILK 2010 - 2022 Centre for Language Studies, Radboud University Nijmegen Induction of Linguistic Know