From 978a55fca774723389ebd6afd02eae7d234340fa Mon Sep 17 00:00:00 2001 From: Thibault Cimic Date: Fri, 3 Jul 2026 15:19:10 +0200 Subject: [PATCH 1/6] Rename windowing files to include windowing prefix - Rename all windowing example files to include 'windowing' in their names: - r2c_c2r_legacy_callback_example.cu -> r2c_c2r_windowing_legacy_callback_example.cu - r2c_c2r_lto_callback_example.cpp -> r2c_c2r_windowing_lto_callback_example.cpp - r2c_c2r_lto_nvrtc_callback_example.cpp -> r2c_c2r_windowing_lto_nvrtc_callback_example.cpp - r2c_c2r_lto_callback_device.cu -> r2c_c2r_windowing_lto_callback_device.cu - r2c_c2r_reference.* -> r2c_c2r_windowing_reference.* - Update CMakeLists.txt, README.md, and includes to use new names - Rename executables to include windowing prefix Signed-off-by: Thibault Cimic --- cuFFT/lto_callback_window_1d/CMakeLists.txt | 57 ++++++++++--------- cuFFT/lto_callback_window_1d/README.md | 18 +++--- ..._c2r_windowing_legacy_callback_example.cu} | 2 +- ... r2c_c2r_windowing_lto_callback_device.cu} | 0 ...2c_c2r_windowing_lto_callback_example.cpp} | 4 +- ..._windowing_lto_nvrtc_callback_example.cpp} | 4 +- ...ence.cu => r2c_c2r_windowing_reference.cu} | 2 +- ...erence.h => r2c_c2r_windowing_reference.h} | 0 8 files changed, 44 insertions(+), 43 deletions(-) rename cuFFT/lto_callback_window_1d/src/{r2c_c2r_legacy_callback_example.cu => r2c_c2r_windowing_legacy_callback_example.cu} (99%) rename cuFFT/lto_callback_window_1d/src/{r2c_c2r_lto_callback_device.cu => r2c_c2r_windowing_lto_callback_device.cu} (100%) rename cuFFT/lto_callback_window_1d/src/{r2c_c2r_lto_callback_example.cpp => r2c_c2r_windowing_lto_callback_example.cpp} (98%) rename cuFFT/lto_callback_window_1d/src/{r2c_c2r_lto_nvrtc_callback_example.cpp => r2c_c2r_windowing_lto_nvrtc_callback_example.cpp} (98%) rename cuFFT/lto_callback_window_1d/src/{r2c_c2r_reference.cu => r2c_c2r_windowing_reference.cu} (98%) rename cuFFT/lto_callback_window_1d/src/{r2c_c2r_reference.h => r2c_c2r_windowing_reference.h} (100%) diff --git a/cuFFT/lto_callback_window_1d/CMakeLists.txt b/cuFFT/lto_callback_window_1d/CMakeLists.txt index 787a0095..ccfd098e 100644 --- a/cuFFT/lto_callback_window_1d/CMakeLists.txt +++ b/cuFFT/lto_callback_window_1d/CMakeLists.txt @@ -41,59 +41,60 @@ set(GENCODE_FLAGS_LTO --generate-code arch=compute_${CUDA_LTO_ARCHITECTURE},code set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin) -add_executable(r2c_c2r_lto_callback_example) -add_executable(r2c_c2r_lto_nvrtc_callback_example) -add_executable(r2c_c2r_legacy_callback_example) +add_executable(r2c_c2r_windowing_lto_callback_example) +add_executable(r2c_c2r_windowing_lto_nvrtc_callback_example) +add_executable(r2c_c2r_windowing_legacy_callback_example) -target_include_directories(r2c_c2r_lto_callback_example +target_include_directories(r2c_c2r_windowing_lto_callback_example PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} ${CMAKE_SOURCE_DIR}/src ${CMAKE_BINARY_DIR}) -target_include_directories(r2c_c2r_lto_nvrtc_callback_example +target_include_directories(r2c_c2r_windowing_lto_nvrtc_callback_example PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} ${CMAKE_SOURCE_DIR}/src ${CMAKE_BINARY_DIR}) -target_include_directories(r2c_c2r_legacy_callback_example +target_include_directories(r2c_c2r_windowing_legacy_callback_example PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} ${CMAKE_SOURCE_DIR}/src ${CMAKE_BINARY_DIR}) -target_sources(r2c_c2r_lto_callback_example - PRIVATE ${PROJECT_SOURCE_DIR}/src/r2c_c2r_lto_callback_example.cpp +target_sources(r2c_c2r_windowing_lto_callback_example + PRIVATE ${PROJECT_SOURCE_DIR}/src/r2c_c2r_windowing_lto_callback_example.cpp ${PROJECT_SOURCE_DIR}/src/common.cpp - ${PROJECT_SOURCE_DIR}/src/r2c_c2r_reference.cu - ${CMAKE_BINARY_DIR}/r2c_c2r_lto_callback_device_fatbin.h) + ${PROJECT_SOURCE_DIR}/src/r2c_c2r_windowing_reference.cu + ${CMAKE_BINARY_DIR}/r2c_c2r_windowing_lto_callback_device_fatbin.h) -target_sources(r2c_c2r_lto_nvrtc_callback_example - PRIVATE ${PROJECT_SOURCE_DIR}/src/r2c_c2r_lto_nvrtc_callback_example.cpp +target_sources(r2c_c2r_windowing_lto_nvrtc_callback_example + PRIVATE ${PROJECT_SOURCE_DIR}/src/r2c_c2r_windowing_lto_nvrtc_callback_example.cpp ${PROJECT_SOURCE_DIR}/src/common.cpp - ${PROJECT_SOURCE_DIR}/src/r2c_c2r_reference.cu) + ${PROJECT_SOURCE_DIR}/src/r2c_c2r_windowing_reference.cu) -target_sources(r2c_c2r_legacy_callback_example - PRIVATE ${PROJECT_SOURCE_DIR}/src/r2c_c2r_legacy_callback_example.cu +target_sources(r2c_c2r_windowing_legacy_callback_example + PRIVATE ${PROJECT_SOURCE_DIR}/src/r2c_c2r_windowing_legacy_callback_example.cu ${PROJECT_SOURCE_DIR}/src/common.cpp - ${PROJECT_SOURCE_DIR}/src/r2c_c2r_reference.cu) + ${PROJECT_SOURCE_DIR}/src/r2c_c2r_windowing_reference.cu) if(CB_USE_CONSTANT_MEMORY) - target_compile_definitions(r2c_c2r_legacy_callback_example PRIVATE CB_USE_CONSTANT_MEMORY) - target_compile_definitions(r2c_c2r_lto_callback_example PRIVATE CB_USE_CONSTANT_MEMORY) - target_compile_definitions(r2c_c2r_lto_nvrtc_callback_example PRIVATE CB_USE_CONSTANT_MEMORY) + target_compile_definitions(r2c_c2r_windowing_legacy_callback_example PRIVATE CB_USE_CONSTANT_MEMORY) + target_compile_definitions(r2c_c2r_windowing_lto_callback_example PRIVATE CB_USE_CONSTANT_MEMORY) + target_compile_definitions(r2c_c2r_windowing_lto_nvrtc_callback_example PRIVATE CB_USE_CONSTANT_MEMORY) endif() -target_compile_definitions(r2c_c2r_lto_nvrtc_callback_example PRIVATE CUDA_ARCH=${CUDA_LTO_ARCHITECTURE} CUDA_PATH=${CUDAToolkit_BIN_DIR}/.. -DSOURCE_PATH=${CMAKE_SOURCE_DIR}/src) +target_compile_definitions(r2c_c2r_windowing_lto_nvrtc_callback_example PRIVATE CUDA_ARCH=${CUDA_LTO_ARCHITECTURE} CUDA_PATH=${CUDAToolkit_BIN_DIR}/.. -DSOURCE_PATH=${CMAKE_SOURCE_DIR}/src) -set_target_properties(r2c_c2r_legacy_callback_example PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON +set_target_properties(r2c_c2r_windowing_legacy_callback_example PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON CUDA_SEPARABLE_COMPILATION ON) message(STATUS "CMAKE_BINARY_DIR: ${CMAKE_BINARY_DIR}") add_custom_command( - OUTPUT ${CMAKE_BINARY_DIR}/r2c_c2r_lto_callback_device.fatbin ${CMAKE_BINARY_DIR}/r2c_c2r_lto_callback_device_fatbin.h - DEPENDS ${CMAKE_SOURCE_DIR}/src/r2c_c2r_lto_callback_device.cu - COMMAND ${CMAKE_CUDA_COMPILER} --std=c++11 ${GENCODE_FLAGS_LTO} -dc -fatbin ${CMAKE_SOURCE_DIR}/src/r2c_c2r_lto_callback_device.cu -o ${CMAKE_BINARY_DIR}/r2c_c2r_lto_callback_device.fatbin - COMMAND ${CUDAToolkit_BIN_DIR}/bin2c --name window_callback --type longlong ${CMAKE_BINARY_DIR}/r2c_c2r_lto_callback_device.fatbin > ${CMAKE_BINARY_DIR}/r2c_c2r_lto_callback_device_fatbin.h + OUTPUT ${CMAKE_BINARY_DIR}/r2c_c2r_windowing_lto_callback_device.fatbin ${CMAKE_BINARY_DIR}/r2c_c2r_windowing_lto_callback_device_fatbin.h + DEPENDS ${CMAKE_SOURCE_DIR}/src/r2c_c2r_windowing_lto_callback_device.cu + COMMAND ${CMAKE_CUDA_COMPILER} --std=c++11 ${GENCODE_FLAGS_LTO} -dc -fatbin ${CMAKE_SOURCE_DIR}/src/r2c_c2r_windowing_lto_callback_device.cu -o ${CMAKE_BINARY_DIR}/r2c_c2r_windowing_lto_callback_device.fatbin + COMMAND ${CUDAToolkit_BIN_DIR}/bin2c --name window_callback --type longlong ${CMAKE_BINARY_DIR}/r2c_c2r_windowing_lto_callback_device.fatbin > ${CMAKE_BINARY_DIR}/r2c_c2r_windowing_lto_callback_device_fatbin.h WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} ) -target_link_libraries(r2c_c2r_lto_callback_example PRIVATE CUDA::cufft CUDA::cudart CUDA::nvJitLink) -target_link_libraries(r2c_c2r_lto_nvrtc_callback_example PRIVATE CUDA::cufft CUDA::cudart CUDA::nvJitLink CUDA::nvrtc) -target_link_libraries(r2c_c2r_legacy_callback_example PRIVATE CUDA::cufft_static CUDA::cudart CUDA::culibos) \ No newline at end of file + +target_link_libraries(r2c_c2r_windowing_lto_callback_example PRIVATE CUDA::cufft CUDA::cudart CUDA::nvJitLink) +target_link_libraries(r2c_c2r_windowing_lto_nvrtc_callback_example PRIVATE CUDA::cufft CUDA::cudart CUDA::nvJitLink CUDA::nvrtc) +target_link_libraries(r2c_c2r_windowing_legacy_callback_example PRIVATE CUDA::cufft_static CUDA::cudart CUDA::culibos) \ No newline at end of file diff --git a/cuFFT/lto_callback_window_1d/README.md b/cuFFT/lto_callback_window_1d/README.md index a7e4b639..b340f52f 100644 --- a/cuFFT/lto_callback_window_1d/README.md +++ b/cuFFT/lto_callback_window_1d/README.md @@ -5,13 +5,13 @@ In this example, we apply a low-pass filter to a batch of signals in the frequency domain. Specifically, the sample code creates a forward (R2C, Real-To-Complex) plan and an inverse (C2R, Complex-To-Real) plan. The low-pass filter is implemented via a load callback in the inverse plan, where values under a certain threshold (specified by the user as a window dimension) are loaded normally, and every other value is zeroed. Three versions of the code are provided: -* `r2c_c2r_lto_callback_example.cpp` contains the sample code using a load callback with LTO to compute the window function. The LTO callback is compiled offline using nvcc. -* `r2c_c2r_lto_nvrtc_callback_example.cpp` contains the sample code using a load callback with LTO to compute the window function. The LTO callback is compiled at runtime using NVRTC. -* `r2c_c2r_legacy_callback_example.cu` contains the sample code using a 'legacy' (non-LTO) load callback to compute the window function. The callback does not use LTO and requires separate device linking against the cuFFT static library. +* `r2c_c2r_windowing_lto_callback_example.cpp` contains the sample code using a load callback with LTO to compute the window function. The LTO callback is compiled offline using nvcc. +* `r2c_c2r_windowing_lto_nvrtc_callback_example.cpp` contains the sample code using a load callback with LTO to compute the window function. The LTO callback is compiled at runtime using NVRTC. +* `r2c_c2r_windowing_legacy_callback_example.cu` contains the sample code using a 'legacy' (non-LTO) load callback to compute the window function. The callback does not use LTO and requires separate device linking against the cuFFT static library. Other source files included: -* `r2c_c2r_lto_callback_device.cu` contains the callback device function used in the LTO and LTO + NVRTC examples. -* `r2c_c2r_reference.cu` contains the code used as reference for the samples. The reference computes the window function using a separate kernel, rather than callbacks. +* `r2c_c2r_windowing_lto_callback_device.cu` contains the callback device function used in the LTO and LTO + NVRTC examples. +* `r2c_c2r_windowing_reference.cu` contains the code used as reference for the samples. The reference computes the window function using a separate kernel, rather than callbacks. * `nvrtc_helper.h` contains the required code to do runtime compilation of the LTO callback using NVRTC. * `common.cpp` and `common.h` include some helper functions, like methods to perform the initialization of the signal in the time domain.. @@ -63,15 +63,15 @@ make ## Running the examples ``` -./bin/r2c_c2r_lto_callback_example -./bin/r2c_c2r_lto_nvrtc_callback_example -./bin/r2c_c2r_callback_example +./bin/r2c_c2r_windowing_lto_callback_example +./bin/r2c_c2r_windowing_lto_nvrtc_callback_example +./bin/r2c_c2r_windowing_legacy_callback_example ``` Sample of output ``` -$ ./bin/r2c_c2r_lto_callback_example +$ ./bin/r2c_c2r_windowing_lto_callback_example Transforming signal cufftExecR2C Transforming signal cufftExecC2R Transforming reference cufftExecR2C diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_legacy_callback_example.cu b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_legacy_callback_example.cu similarity index 99% rename from cuFFT/lto_callback_window_1d/src/r2c_c2r_legacy_callback_example.cu rename to cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_legacy_callback_example.cu index 7d838d02..abe0205f 100644 --- a/cuFFT/lto_callback_window_1d/src/r2c_c2r_legacy_callback_example.cu +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_legacy_callback_example.cu @@ -26,7 +26,7 @@ #include #include #include "common.h" -#include "r2c_c2r_reference.h" +#include "r2c_c2r_windowing_reference.h" #include "callback_params.h" // This is the store callback routine. It filters high frequencies diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_lto_callback_device.cu b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_callback_device.cu similarity index 100% rename from cuFFT/lto_callback_window_1d/src/r2c_c2r_lto_callback_device.cu rename to cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_callback_device.cu diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_lto_callback_example.cpp b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_callback_example.cpp similarity index 98% rename from cuFFT/lto_callback_window_1d/src/r2c_c2r_lto_callback_example.cpp rename to cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_callback_example.cpp index 8b35b94c..8745d1b5 100644 --- a/cuFFT/lto_callback_window_1d/src/r2c_c2r_lto_callback_example.cpp +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_callback_example.cpp @@ -25,12 +25,12 @@ #include #include -#include "r2c_c2r_reference.h" +#include "r2c_c2r_windowing_reference.h" #include "common.h" #include "callback_params.h" // NOTE: Header containing the compiled LTO callback device function in a C array, generated with bin2c -#include "r2c_c2r_lto_callback_device_fatbin.h" +#include "r2c_c2r_windowing_lto_callback_device_fatbin.h" static_assert(window_size < (signal_size/2 + 1), "The window size must be smaller than the signal size in complex space"); diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_lto_nvrtc_callback_example.cpp b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_nvrtc_callback_example.cpp similarity index 98% rename from cuFFT/lto_callback_window_1d/src/r2c_c2r_lto_nvrtc_callback_example.cpp rename to cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_nvrtc_callback_example.cpp index 4ec57e6b..0069d4d0 100644 --- a/cuFFT/lto_callback_window_1d/src/r2c_c2r_lto_nvrtc_callback_example.cpp +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_nvrtc_callback_example.cpp @@ -25,7 +25,7 @@ #include #include -#include "r2c_c2r_reference.h" +#include "r2c_c2r_windowing_reference.h" #include "common.h" #include "nvrtc_helper.h" #include "callback_params.h" @@ -48,7 +48,7 @@ int test_r2c_window_c2r() { // NOTE: Use NVRTC to compile the callback function to LTO std::vector callback_buffer; - compile_file_to_lto(callback_buffer, CALLBACK_CODE_PATH("r2c_c2r_lto_callback_device.cu")); + compile_file_to_lto(callback_buffer, CALLBACK_CODE_PATH("r2c_c2r_windowing_lto_callback_device.cu")); // Create a CUFFT plan for the forward transform, and a cuFFT plan for the inverse transform with load callback cufftHandle forward_plan, inverse_plan_cb; diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_reference.cu b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_reference.cu similarity index 98% rename from cuFFT/lto_callback_window_1d/src/r2c_c2r_reference.cu rename to cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_reference.cu index b040bfd5..b029f54f 100644 --- a/cuFFT/lto_callback_window_1d/src/r2c_c2r_reference.cu +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_reference.cu @@ -26,7 +26,7 @@ #include #include #include -#include "r2c_c2r_reference.h" +#include "r2c_c2r_windowing_reference.h" __global__ void windowing(unsigned nbatches, unsigned complex_signal_size, unsigned window_size, float2* buffer) { diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_reference.h b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_reference.h similarity index 100% rename from cuFFT/lto_callback_window_1d/src/r2c_c2r_reference.h rename to cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_reference.h From 7ce081e20760e4087e4a379e5c3ee56279db0454 Mon Sep 17 00:00:00 2001 From: Thibault Cimic Date: Fri, 3 Jul 2026 15:19:10 +0200 Subject: [PATCH 2/6] Rename compute_error to compute_error_windowing for windowing examples - Add compute_error_windowing for complex buffer layout (windowing) - Add compute_error_padding for real buffer layout (padding) - Update windowing examples to use compute_error_windowing - Update common.h with both error computation functions Signed-off-by: Thibault Cimic --- cuFFT/lto_callback_window_1d/src/common.cpp | 31 +++++++++++++++++++ cuFFT/lto_callback_window_1d/src/common.h | 20 +++++++++++- ...c_c2r_windowing_legacy_callback_example.cu | 2 +- ...r2c_c2r_windowing_lto_callback_example.cpp | 2 +- ...r_windowing_lto_nvrtc_callback_example.cpp | 2 +- 5 files changed, 53 insertions(+), 4 deletions(-) diff --git a/cuFFT/lto_callback_window_1d/src/common.cpp b/cuFFT/lto_callback_window_1d/src/common.cpp index ae6ffcdc..ddfa10ac 100644 --- a/cuFFT/lto_callback_window_1d/src/common.cpp +++ b/cuFFT/lto_callback_window_1d/src/common.cpp @@ -30,6 +30,7 @@ constexpr float sampling_dt = 1e-3; // Initialize the input signal as a composite of sine waves // with random amplitudes and frequencies +// Uses complex buffer layout (stride = 2 * complex_signal_size) void init_input_signals(unsigned batches, unsigned signal_size, float* signals) { std::mt19937 e2(0); @@ -60,4 +61,34 @@ void init_input_signals(unsigned batches, unsigned signal_size, float* signals) } } +// Initialize the input signal for padding examples +// Uses real buffer layout (stride = signal_size) +void init_input_signals_padding(unsigned batches, unsigned signal_size, float* signals) { + std::mt19937 e2(0); + + std::uniform_real_distribution<> A_dist(0., signal_max_A); + std::uniform_real_distribution<> f_dist(0., signal_max_f); + + for(unsigned batch = 0; batch < batches; ++batch) { + std::vector wave_amplitudes; + std::vector wave_frequencies; + + // Generate the amplitudes and frequencies of the waves + for(unsigned i = 0; i < waves; ++i) { + wave_amplitudes.push_back(A_dist(e2)); + wave_frequencies.push_back(f_dist(e2)); + } + + // Compose the signal + float time = 0.; + for(unsigned s = 0; s < signal_size; ++s) { + for(unsigned i = 0; i < waves; ++i) { + unsigned idx = batch * signal_size + s; + signals[idx] += wave_amplitudes[i] * sin(2. * PI * wave_frequencies[i] * time) ; + } + time += sampling_dt; + } + } +} + #endif // _COMMON__CPP_ \ No newline at end of file diff --git a/cuFFT/lto_callback_window_1d/src/common.h b/cuFFT/lto_callback_window_1d/src/common.h index 6bcbaccc..b1fc6d1f 100644 --- a/cuFFT/lto_callback_window_1d/src/common.h +++ b/cuFFT/lto_callback_window_1d/src/common.h @@ -49,7 +49,7 @@ inline int checkErrors(cufftResult error, int line_number) { #define CHECK_ERROR(error) checkErrors(error, __LINE__) template -double compute_error(T* ref, T* out, unsigned batches, unsigned signal_size){ +double compute_error_windowing(T* ref, T* out, unsigned batches, unsigned signal_size){ double squared_diff = 0; double squared_norm = 0; const unsigned batch_offset = 2 * (signal_size / 2 + 1); @@ -63,6 +63,24 @@ double compute_error(T* ref, T* out, unsigned batches, unsigned signal_size){ return std::sqrt(squared_diff / squared_norm); } +// Compute error for real buffers (padding examples) +template +double compute_error_padding(T* ref, T* out, unsigned batches, unsigned signal_size){ + double squared_diff = 0; + double squared_norm = 0; + const unsigned batch_offset = signal_size; + + for (unsigned b = 0; b < batches; b++) { + for (unsigned i = 0; i < signal_size; i++) { + unsigned ref_idx = b * batch_offset + i; + squared_diff += std::norm(ref[ref_idx] - out[ref_idx]); + squared_norm += std::norm(ref[ref_idx]); + } + } + return std::sqrt(squared_diff / squared_norm); +} + // Initialize input signals void init_input_signals(unsigned batches, unsigned signal_size, float* signals); +void init_input_signals_padding(unsigned batches, unsigned signal_size, float* signals); #endif // _COMMON__H_ \ No newline at end of file diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_legacy_callback_example.cu b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_legacy_callback_example.cu index abe0205f..9b05c064 100644 --- a/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_legacy_callback_example.cu +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_legacy_callback_example.cu @@ -125,7 +125,7 @@ int test_r2c_window_c2r() { return ERROR_VALUE; } - double l2_error = compute_error(&reference[0][0], &output_signals[0][0], batches, signal_size); + double l2_error = compute_error_windowing(&reference[0][0], &output_signals[0][0], batches, signal_size); printf("L2 error: %e\n", l2_error); return (l2_error < threshold) ? PASS_VALUE : ERROR_VALUE; diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_callback_example.cpp b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_callback_example.cpp index 8745d1b5..523e12c9 100644 --- a/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_callback_example.cpp +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_callback_example.cpp @@ -110,7 +110,7 @@ int test_r2c_window_c2r() { return ERROR_VALUE; }; - double l2_error = compute_error(&reference[0][0], &output_signals[0][0], batches, signal_size); + double l2_error = compute_error_windowing(&reference[0][0], &output_signals[0][0], batches, signal_size); printf("L2 error: %e\n", l2_error); return (l2_error < threshold) ? PASS_VALUE : ERROR_VALUE; diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_nvrtc_callback_example.cpp b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_nvrtc_callback_example.cpp index 0069d4d0..ff1f97e0 100644 --- a/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_nvrtc_callback_example.cpp +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_nvrtc_callback_example.cpp @@ -109,7 +109,7 @@ int test_r2c_window_c2r() { return ERROR_VALUE; } - double l2_error = compute_error(reference[0], output_signals[0], batches, signal_size); + double l2_error = compute_error_windowing(reference[0], output_signals[0], batches, signal_size); printf("L2 error: %e\n", l2_error); return (l2_error < threshold) ? PASS_VALUE : ERROR_VALUE; From cf0604fc160fa83002a54b7b7f075c79f7d7f636 Mon Sep 17 00:00:00 2001 From: Thibault Cimic Date: Fri, 3 Jul 2026 15:19:10 +0200 Subject: [PATCH 3/6] Add padding callback examples demonstrating load and store callbacks New files: - r2c_c2r_padding_legacy_callback_example.cu: Legacy API with load/store callbacks - r2c_c2r_padding_lto_callback_example.cpp: LTO API with load/store callbacks - r2c_c2r_padding_lto_load_callback_device.cu: LTO load callback device code - r2c_c2r_padding_lto_store_callback_device.cu: LTO store callback device code - r2c_c2r_padding_reference.cu/h: Reference implementation without callbacks Modified files: - CMakeLists.txt: Add padding targets and fatbin generation - README.md: Document padding examples - callback_params.h: Add padding parameters - common.h: Add compute_error_padding function Signed-off-by: Thibault Cimic --- cuFFT/lto_callback_window_1d/CMakeLists.txt | 46 ++++- cuFFT/lto_callback_window_1d/README.md | 46 ++++- .../src/callback_params.h | 5 + ...r2c_c2r_padding_legacy_callback_example.cu | 186 ++++++++++++++++++ .../r2c_c2r_padding_lto_callback_example.cpp | 160 +++++++++++++++ ...2c_c2r_padding_lto_load_callback_device.cu | 47 +++++ ...c_c2r_padding_lto_store_callback_device.cu | 47 +++++ .../src/r2c_c2r_padding_reference.cu | 97 +++++++++ .../src/r2c_c2r_padding_reference.h | 26 +++ 9 files changed, 653 insertions(+), 7 deletions(-) create mode 100644 cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_legacy_callback_example.cu create mode 100644 cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_callback_example.cpp create mode 100644 cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_load_callback_device.cu create mode 100644 cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_store_callback_device.cu create mode 100644 cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_reference.cu create mode 100644 cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_reference.h diff --git a/cuFFT/lto_callback_window_1d/CMakeLists.txt b/cuFFT/lto_callback_window_1d/CMakeLists.txt index ccfd098e..5716da3c 100644 --- a/cuFFT/lto_callback_window_1d/CMakeLists.txt +++ b/cuFFT/lto_callback_window_1d/CMakeLists.txt @@ -44,6 +44,8 @@ set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin) add_executable(r2c_c2r_windowing_lto_callback_example) add_executable(r2c_c2r_windowing_lto_nvrtc_callback_example) add_executable(r2c_c2r_windowing_legacy_callback_example) +add_executable(r2c_c2r_padding_lto_callback_example) +add_executable(r2c_c2r_padding_legacy_callback_example) target_include_directories(r2c_c2r_windowing_lto_callback_example PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} @@ -57,6 +59,14 @@ target_include_directories(r2c_c2r_windowing_legacy_callback_example PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} ${CMAKE_SOURCE_DIR}/src ${CMAKE_BINARY_DIR}) +target_include_directories(r2c_c2r_padding_lto_callback_example + PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} + ${CMAKE_SOURCE_DIR}/src + ${CMAKE_BINARY_DIR}) +target_include_directories(r2c_c2r_padding_legacy_callback_example + PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} + ${CMAKE_SOURCE_DIR}/src + ${CMAKE_BINARY_DIR}) target_sources(r2c_c2r_windowing_lto_callback_example PRIVATE ${PROJECT_SOURCE_DIR}/src/r2c_c2r_windowing_lto_callback_example.cpp @@ -74,6 +84,18 @@ target_sources(r2c_c2r_windowing_legacy_callback_example ${PROJECT_SOURCE_DIR}/src/common.cpp ${PROJECT_SOURCE_DIR}/src/r2c_c2r_windowing_reference.cu) +target_sources(r2c_c2r_padding_lto_callback_example + PRIVATE ${PROJECT_SOURCE_DIR}/src/r2c_c2r_padding_lto_callback_example.cpp + ${PROJECT_SOURCE_DIR}/src/common.cpp + ${PROJECT_SOURCE_DIR}/src/r2c_c2r_padding_reference.cu + ${CMAKE_BINARY_DIR}/r2c_c2r_padding_load_callback_fatbin.h + ${CMAKE_BINARY_DIR}/r2c_c2r_padding_store_callback_fatbin.h) + +target_sources(r2c_c2r_padding_legacy_callback_example + PRIVATE ${PROJECT_SOURCE_DIR}/src/r2c_c2r_padding_legacy_callback_example.cu + ${PROJECT_SOURCE_DIR}/src/common.cpp + ${PROJECT_SOURCE_DIR}/src/r2c_c2r_padding_reference.cu) + if(CB_USE_CONSTANT_MEMORY) target_compile_definitions(r2c_c2r_windowing_legacy_callback_example PRIVATE CB_USE_CONSTANT_MEMORY) target_compile_definitions(r2c_c2r_windowing_lto_callback_example PRIVATE CB_USE_CONSTANT_MEMORY) @@ -85,6 +107,9 @@ target_compile_definitions(r2c_c2r_windowing_lto_nvrtc_callback_example PRIVATE set_target_properties(r2c_c2r_windowing_legacy_callback_example PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON CUDA_SEPARABLE_COMPILATION ON) +set_target_properties(r2c_c2r_padding_legacy_callback_example PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON + CUDA_SEPARABLE_COMPILATION ON) + message(STATUS "CMAKE_BINARY_DIR: ${CMAKE_BINARY_DIR}") add_custom_command( OUTPUT ${CMAKE_BINARY_DIR}/r2c_c2r_windowing_lto_callback_device.fatbin ${CMAKE_BINARY_DIR}/r2c_c2r_windowing_lto_callback_device_fatbin.h @@ -94,7 +119,26 @@ add_custom_command( WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} ) +add_custom_command( + OUTPUT ${CMAKE_BINARY_DIR}/r2c_c2r_padding_load_callback.fatbin ${CMAKE_BINARY_DIR}/r2c_c2r_padding_load_callback_fatbin.h + DEPENDS ${CMAKE_SOURCE_DIR}/src/r2c_c2r_padding_lto_load_callback_device.cu + COMMAND ${CMAKE_CUDA_COMPILER} --std=c++11 ${GENCODE_FLAGS_LTO} -dc -fatbin ${CMAKE_SOURCE_DIR}/src/r2c_c2r_padding_lto_load_callback_device.cu -o ${CMAKE_BINARY_DIR}/r2c_c2r_padding_load_callback.fatbin + COMMAND ${CUDAToolkit_BIN_DIR}/bin2c --name padding_load_callback --type longlong ${CMAKE_BINARY_DIR}/r2c_c2r_padding_load_callback.fatbin > ${CMAKE_BINARY_DIR}/r2c_c2r_padding_load_callback_fatbin.h + WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} + COMMENT "Compiling padding LTO load callback device code" +) + +add_custom_command( + OUTPUT ${CMAKE_BINARY_DIR}/r2c_c2r_padding_store_callback.fatbin ${CMAKE_BINARY_DIR}/r2c_c2r_padding_store_callback_fatbin.h + DEPENDS ${CMAKE_SOURCE_DIR}/src/r2c_c2r_padding_lto_store_callback_device.cu + COMMAND ${CMAKE_CUDA_COMPILER} --std=c++11 ${GENCODE_FLAGS_LTO} -dc -fatbin ${CMAKE_SOURCE_DIR}/src/r2c_c2r_padding_lto_store_callback_device.cu -o ${CMAKE_BINARY_DIR}/r2c_c2r_padding_store_callback.fatbin + COMMAND ${CUDAToolkit_BIN_DIR}/bin2c --name padding_store_callback --type longlong ${CMAKE_BINARY_DIR}/r2c_c2r_padding_store_callback.fatbin > ${CMAKE_BINARY_DIR}/r2c_c2r_padding_store_callback_fatbin.h + WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} + COMMENT "Compiling padding LTO store callback device code" +) target_link_libraries(r2c_c2r_windowing_lto_callback_example PRIVATE CUDA::cufft CUDA::cudart CUDA::nvJitLink) target_link_libraries(r2c_c2r_windowing_lto_nvrtc_callback_example PRIVATE CUDA::cufft CUDA::cudart CUDA::nvJitLink CUDA::nvrtc) -target_link_libraries(r2c_c2r_windowing_legacy_callback_example PRIVATE CUDA::cufft_static CUDA::cudart CUDA::culibos) \ No newline at end of file +target_link_libraries(r2c_c2r_windowing_legacy_callback_example PRIVATE CUDA::cufft_static CUDA::cudart CUDA::culibos) +target_link_libraries(r2c_c2r_padding_lto_callback_example PRIVATE CUDA::cufft CUDA::cudart CUDA::nvJitLink) +target_link_libraries(r2c_c2r_padding_legacy_callback_example PRIVATE CUDA::cufft_static CUDA::cudart CUDA::culibos) diff --git a/cuFFT/lto_callback_window_1d/README.md b/cuFFT/lto_callback_window_1d/README.md index b340f52f..2faae93e 100644 --- a/cuFFT/lto_callback_window_1d/README.md +++ b/cuFFT/lto_callback_window_1d/README.md @@ -2,18 +2,33 @@ ## Description -In this example, we apply a low-pass filter to a batch of signals in the frequency domain. Specifically, the sample code creates a forward (R2C, Real-To-Complex) plan and an inverse (C2R, Complex-To-Real) plan. The low-pass filter is implemented via a load callback in the inverse plan, where values under a certain threshold (specified by the user as a window dimension) are loaded normally, and every other value is zeroed. +This directory contains examples demonstrating the use of cuFFT callbacks with both legacy and LTO (Link-Time Optimization) approaches. Two different use cases are provided: windowing and zero-padding. -Three versions of the code are provided: +### Windowing Examples + +In the windowing examples, we apply a low-pass filter to a batch of signals in the frequency domain. Specifically, the sample code creates a forward (R2C, Real-To-Complex) plan and an inverse (C2R, Complex-To-Real) plan. The low-pass filter is implemented via a **load callback** in the inverse plan, where values under a certain threshold (specified by the user as a window dimension) are loaded normally, and every other value is zeroed. + +Three versions of the windowing code are provided: * `r2c_c2r_windowing_lto_callback_example.cpp` contains the sample code using a load callback with LTO to compute the window function. The LTO callback is compiled offline using nvcc. * `r2c_c2r_windowing_lto_nvrtc_callback_example.cpp` contains the sample code using a load callback with LTO to compute the window function. The LTO callback is compiled at runtime using NVRTC. * `r2c_c2r_windowing_legacy_callback_example.cu` contains the sample code using a 'legacy' (non-LTO) load callback to compute the window function. The callback does not use LTO and requires separate device linking against the cuFFT static library. +### Padding Examples + +In the padding examples, we demonstrate the use of **both load and store callbacks** to implement zero-padding. The load callback is applied on the forward (R2C) plan to pad the input with zeros, and the store callback is applied on the inverse (C2R) plan to truncate the output to the original size and normalize. + +Two versions of the padding code are provided: +* `r2c_c2r_padding_lto_callback_example.cpp` contains the sample code using LTO callbacks to implement zero-padding. The LTO callback is compiled offline using nvcc. +* `r2c_c2r_padding_legacy_callback_example.cu` contains the sample code using legacy callbacks to implement zero-padding. The callback does not use LTO and requires separate device linking against the cuFFT static library. + Other source files included: -* `r2c_c2r_windowing_lto_callback_device.cu` contains the callback device function used in the LTO and LTO + NVRTC examples. -* `r2c_c2r_windowing_reference.cu` contains the code used as reference for the samples. The reference computes the window function using a separate kernel, rather than callbacks. +* `r2c_c2r_windowing_lto_callback_device.cu` contains the callback device function used in the windowing LTO and LTO + NVRTC examples. +* `r2c_c2r_padding_lto_load_callback_device.cu` contains the load callback device function used in the padding LTO example. +* `r2c_c2r_padding_lto_store_callback_device.cu` contains the store callback device function used in the padding LTO example. +* `r2c_c2r_windowing_reference.cu` contains the code used as reference for the windowing samples. The reference computes the window function using a separate kernel, rather than callbacks. +* `r2c_c2r_padding_reference.cu` contains the code used as reference for the padding samples. The reference computes the zero-padding using separate kernels, rather than callbacks. * `nvrtc_helper.h` contains the required code to do runtime compilation of the LTO callback using NVRTC. -* `common.cpp` and `common.h` include some helper functions, like methods to perform the initialization of the signal in the time domain.. +* `common.cpp` and `common.h` include some helper functions, like methods to perform the initialization of the signal in the time domain. ## Supported SM Architectures @@ -62,13 +77,21 @@ make ``` ## Running the examples + +### Windowing examples ``` ./bin/r2c_c2r_windowing_lto_callback_example ./bin/r2c_c2r_windowing_lto_nvrtc_callback_example ./bin/r2c_c2r_windowing_legacy_callback_example ``` -Sample of output +### Padding examples +``` +./bin/r2c_c2r_padding_lto_callback_example +./bin/r2c_c2r_padding_legacy_callback_example +``` + +Sample of output (windowing) ``` $ ./bin/r2c_c2r_windowing_lto_callback_example @@ -79,6 +102,17 @@ Transforming reference cufftExecC2R L2 error: 0.000000e+00 ``` +Sample of output (padding) + +``` +$ ./bin/r2c_c2r_padding_lto_callback_example +Transforming signal cufftExecR2C with padding load callback +Transforming signal cufftExecC2R with padding store callback +Transforming reference cufftExecR2C +Transforming reference cufftExecC2R +L2 error: 0.000000e+00 +``` + ## Troubleshooting ### I am getting an error with the message: "error: ‘cufftXtSetJITCallback’ was not declared in this scope". Please, make sure you are using cuFFT from CUDA Toolkit 12.6 Update 2 or newer. LTO callbacks are not available on older versions of cuFFT. diff --git a/cuFFT/lto_callback_window_1d/src/callback_params.h b/cuFFT/lto_callback_window_1d/src/callback_params.h index f07ac06b..9f203a2f 100644 --- a/cuFFT/lto_callback_window_1d/src/callback_params.h +++ b/cuFFT/lto_callback_window_1d/src/callback_params.h @@ -31,6 +31,11 @@ constexpr unsigned signal_size = 328; constexpr unsigned window_size = 32; constexpr unsigned complex_signal_size = signal_size / 2 + 1; +// Padding example parameters - use single batch for simplicity +constexpr unsigned padding_batches = 1; +constexpr unsigned padded_signal_size = 2 * signal_size; +constexpr unsigned padding_complex_signal_size = padded_signal_size / 2 + 1; + // Precision threshold constexpr float threshold = 1e-6; diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_legacy_callback_example.cu b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_legacy_callback_example.cu new file mode 100644 index 00000000..7d2ce82e --- /dev/null +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_legacy_callback_example.cu @@ -0,0 +1,186 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +/* + * Example showing the use of legacy callbacks with CUFFT to perform + * zero-padding on a single 1D signal. + * + */ + +#include +#include +#include "common.h" +#include "r2c_c2r_padding_reference.h" +#include "callback_params.h" + +// Callback parameters structure for padding +struct PaddingCallbackParams { + unsigned int signal_size; + unsigned int padded_signal_size; +}; + +// Load callback for R2C: returns 0 for indices beyond signal_size +__device__ cufftReal padding_load_callback_r2c(void* input, + size_t index, + void* info, + void* sharedMem) { + PaddingCallbackParams* params = (PaddingCallbackParams*)info; + cufftReal* in = (cufftReal*)input; + + // If index is beyond the original data, return zero + if (index >= params->signal_size) { + return 0.0f; + } + + return in[index]; +} + +// Store callback for C2R: only stores first signal_size elements, zeros the rest +__device__ void padding_store_callback_c2r(void* output, + size_t index, + cufftReal in, + void* info, + void* sharedMem) { + PaddingCallbackParams* params = (PaddingCallbackParams*)info; + cufftReal* out = (cufftReal*)output; + + // Only store if within original size, zero and normalize otherwise + if (index < params->signal_size) { + out[index] = in / static_cast(params->padded_signal_size); + } else { + out[index] = 0.0f; + } +} + +// Device pointers to callback functions (required for legacy callbacks) +__device__ cufftCallbackLoadR d_padding_load_callback_ptr = padding_load_callback_r2c; +__device__ cufftCallbackStoreR d_padding_store_callback_ptr = padding_store_callback_c2r; + +static_assert(padded_signal_size > signal_size, "The padded size must be larger than the signal size"); + +int test_r2c_padding_c2r() { + + // Allocate and initialize host input (real data) - only signal_size elements + float *input_signals = new float[signal_size]; + float *output_signals = new float[signal_size]; + float *reference = new float[signal_size]; + + // Initialize with simple linear ramp for debugging + for (unsigned i = 0; i < signal_size; i++) { + input_signals[i] = static_cast(i); + } + + const size_t padded_real_size_bytes = padded_signal_size * sizeof(float); + const size_t complex_size_bytes = padding_complex_signal_size * 2 * sizeof(float); + + // Allocate device memory for padded real data + float *device_signals; + CHECK_ERROR(cudaMalloc((void **)&device_signals, padded_real_size_bytes)); + CHECK_ERROR(cudaMemcpy(device_signals, input_signals, signal_size * sizeof(float), cudaMemcpyHostToDevice)); + + // Allocate device memory for complex data + cufftComplex *device_complex; + CHECK_ERROR(cudaMalloc((void **)&device_complex, complex_size_bytes)); + + // Create callback parameters on host + PaddingCallbackParams h_params; + h_params.signal_size = signal_size; + h_params.padded_signal_size = padded_signal_size; + + // Allocate and copy callback parameters to device + PaddingCallbackParams *device_params; + CHECK_ERROR(cudaMalloc((void **)&device_params, sizeof(PaddingCallbackParams))); + CHECK_ERROR(cudaMemcpy(device_params, &h_params, sizeof(PaddingCallbackParams), cudaMemcpyHostToDevice)); + + // Get device pointers to callback functions + cufftCallbackLoadR h_load_callback_ptr; + cufftCallbackStoreR h_store_callback_ptr; + CHECK_ERROR(cudaMemcpyFromSymbol(&h_load_callback_ptr, d_padding_load_callback_ptr, sizeof(h_load_callback_ptr))); + CHECK_ERROR(cudaMemcpyFromSymbol(&h_store_callback_ptr, d_padding_store_callback_ptr, sizeof(h_store_callback_ptr))); + + // Create cuFFT plans + cufftHandle forward_plan, inverse_plan; + size_t work_size; + + CHECK_ERROR(cufftCreate(&forward_plan)); + CHECK_ERROR(cufftCreate(&inverse_plan)); + + // Create forward plan (R2C) with load callback + CHECK_ERROR(cufftMakePlan1d(forward_plan, padded_signal_size, CUFFT_R2C, 1, &work_size)); + CHECK_ERROR(cufftXtSetCallback(forward_plan, (void **)&h_load_callback_ptr, + CUFFT_CB_LD_REAL, (void **)&device_params)); + + // Create inverse plan (C2R) with store callback + CHECK_ERROR(cufftMakePlan1d(inverse_plan, padded_signal_size, CUFFT_C2R, 1, &work_size)); + CHECK_ERROR(cufftXtSetCallback(inverse_plan, (void **)&h_store_callback_ptr, + CUFFT_CB_ST_REAL, (void **)&device_params)); + + // Execute forward FFT (R2C) with load callback + printf("Transforming signal cufftExecR2C with padding load callback\n"); + CHECK_ERROR(cufftExecR2C(forward_plan, (cufftReal *)device_signals, (cufftComplex *)device_complex)); + + // Execute inverse FFT (C2R) with store callback + printf("Transforming signal cufftExecC2R with padding store callback\n"); + CHECK_ERROR(cufftExecC2R(inverse_plan, (cufftComplex *)device_complex, (cufftReal *)device_signals)); + + // Copy result back to host (only first signal_size real values) + CHECK_ERROR(cudaMemcpy(output_signals, device_signals, signal_size * sizeof(float), cudaMemcpyDeviceToHost)); + + // Compute reference + if (reference_r2c_padding_c2r(1, signal_size, padded_signal_size, input_signals, reference) != PASS_VALUE) { + printf("Failed to compute the reference"); + delete[] input_signals; + delete[] output_signals; + delete[] reference; + return ERROR_VALUE; + } + + double l2_error = compute_error_padding(reference, output_signals, 1, signal_size); + printf("L2 error: %e\n", l2_error); + + // Cleanup host memory + delete[] input_signals; + delete[] output_signals; + delete[] reference; + + // Destroy CUFFT context + CHECK_ERROR(cufftDestroy(forward_plan)); + CHECK_ERROR(cufftDestroy(inverse_plan)); + + // Cleanup memory + CHECK_ERROR(cudaFree(device_signals)); + CHECK_ERROR(cudaFree(device_complex)); + CHECK_ERROR(cudaFree(device_params)); + + return (l2_error < threshold) ? PASS_VALUE : ERROR_VALUE; +} + +//////////////////////////////////////////////////////////////////////////////// +// Program main +//////////////////////////////////////////////////////////////////////////////// +int main(int argc, char **argv) { + struct cudaDeviceProp properties; + int device; + CHECK_ERROR(cudaGetDevice(&device)); + CHECK_ERROR(cudaGetDeviceProperties(&properties, device)); + if (!(properties.major >= 5)) { + printf("cuFFT with callbacks requires CUDA architecture SM5.0 or higher\n"); + return ERROR_VALUE; + } + + return test_r2c_padding_c2r(); +} diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_callback_example.cpp b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_callback_example.cpp new file mode 100644 index 00000000..2e945390 --- /dev/null +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_callback_example.cpp @@ -0,0 +1,160 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +/* + * Example showing the use of LTO callbacks with CUFFT to perform + * zero-padding on a single 1D signal. + * + */ + +#include +#include +#include "r2c_c2r_padding_reference.h" +#include "common.h" +#include "callback_params.h" + +// NOTE: Headers containing the compiled LTO callback device functions in C arrays, generated with bin2c +#include "r2c_c2r_padding_load_callback_fatbin.h" +#include "r2c_c2r_padding_store_callback_fatbin.h" + +// Callback parameters structure for padding +struct PaddingCallbackParams { + unsigned int signal_size; + unsigned int padded_signal_size; +}; + +static_assert(padded_signal_size > signal_size, "The padded size must be larger than the signal size"); + +int test_r2c_padding_c2r() { + + // Allocate and initialize host input (real data) - only signal_size elements + float *input_signals = new float[signal_size]; + float *output_signals = new float[signal_size]; + float *reference = new float[signal_size]; + + // Initialize with simple linear ramp for debugging + for (unsigned i = 0; i < signal_size; i++) { + input_signals[i] = static_cast(i); + } + + const size_t padded_real_size_bytes = padded_signal_size * sizeof(float); + const size_t complex_size_bytes = padding_complex_signal_size * 2 * sizeof(float); + + // Allocate device memory for padded real data + float *device_signals; + CHECK_ERROR(cudaMalloc((void **)&device_signals, padded_real_size_bytes)); + CHECK_ERROR(cudaMemcpy(device_signals, input_signals, signal_size * sizeof(float), cudaMemcpyHostToDevice)); + + // Allocate device memory for complex data + cufftComplex *device_complex; + CHECK_ERROR(cudaMalloc((void **)&device_complex, complex_size_bytes)); + + // Create callback parameters on host + PaddingCallbackParams h_params; + h_params.signal_size = signal_size; + h_params.padded_signal_size = padded_signal_size; + + // Allocate and copy callback parameters to device + PaddingCallbackParams *device_params; + CHECK_ERROR(cudaMalloc((void **)&device_params, sizeof(PaddingCallbackParams))); + CHECK_ERROR(cudaMemcpy(device_params, &h_params, sizeof(PaddingCallbackParams), cudaMemcpyHostToDevice)); + + // Create cuFFT plans + cufftHandle forward_plan, inverse_plan; + size_t work_size; + + CHECK_ERROR(cufftCreate(&forward_plan)); + CHECK_ERROR(cufftCreate(&inverse_plan)); + + // NOTE: LTO callbacks must be set before plan creation and cannot be unset (yet) + // Set load callback for forward plan (R2C) + size_t lto_load_callback_fatbin_size = sizeof(padding_load_callback); + CHECK_ERROR(cufftXtSetJITCallback(forward_plan, + "padding_load_callback", + (void*)padding_load_callback, + lto_load_callback_fatbin_size, + CUFFT_CB_LD_REAL, + (void **)&device_params)); + + // Set store callback for inverse plan (C2R) + size_t lto_store_callback_fatbin_size = sizeof(padding_store_callback); + CHECK_ERROR(cufftXtSetJITCallback(inverse_plan, + "padding_store_callback", + (void*)padding_store_callback, + lto_store_callback_fatbin_size, + CUFFT_CB_ST_REAL, + (void **)&device_params)); + + // Create plans AFTER setting callbacks + CHECK_ERROR(cufftMakePlan1d(forward_plan, padded_signal_size, CUFFT_R2C, 1, &work_size)); + CHECK_ERROR(cufftMakePlan1d(inverse_plan, padded_signal_size, CUFFT_C2R, 1, &work_size)); + + // Execute forward FFT (R2C) with load callback + printf("Transforming signal cufftExecR2C with padding load callback\n"); + CHECK_ERROR(cufftExecR2C(forward_plan, (cufftReal *)device_signals, (cufftComplex *)device_complex)); + + // Execute inverse FFT (C2R) with store callback + printf("Transforming signal cufftExecC2R with padding store callback\n"); + CHECK_ERROR(cufftExecC2R(inverse_plan, (cufftComplex *)device_complex, (cufftReal *)device_signals)); + + // Copy result back to host (only first signal_size real values) + CHECK_ERROR(cudaMemcpy(output_signals, device_signals, signal_size * sizeof(float), cudaMemcpyDeviceToHost)); + + // Compute reference + if (reference_r2c_padding_c2r(1, signal_size, padded_signal_size, input_signals, reference) != PASS_VALUE) { + printf("Failed to compute the reference"); + delete[] input_signals; + delete[] output_signals; + delete[] reference; + return ERROR_VALUE; + } + + double l2_error = compute_error_padding(reference, output_signals, 1, signal_size); + printf("L2 error: %e\n", l2_error); + + // Cleanup host memory + delete[] input_signals; + delete[] output_signals; + delete[] reference; + + // Destroy CUFFT context + CHECK_ERROR(cufftDestroy(forward_plan)); + CHECK_ERROR(cufftDestroy(inverse_plan)); + + // Cleanup memory + CHECK_ERROR(cudaFree(device_signals)); + CHECK_ERROR(cudaFree(device_complex)); + CHECK_ERROR(cudaFree(device_params)); + + return (l2_error < threshold) ? PASS_VALUE : ERROR_VALUE; +} + +//////////////////////////////////////////////////////////////////////////////// +// Program main +//////////////////////////////////////////////////////////////////////////////// +int main(int argc, char **argv) { + struct cudaDeviceProp properties; + int device; + CHECK_ERROR(cudaGetDevice(&device)); + CHECK_ERROR(cudaGetDeviceProperties(&properties, device)); + if (!(properties.major >= 5)) { + printf("cuFFT with LTO requires CUDA architecture SM5.0 or higher\n"); + return ERROR_VALUE; + } + + return test_r2c_padding_c2r(); +} diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_load_callback_device.cu b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_load_callback_device.cu new file mode 100644 index 00000000..5e147b77 --- /dev/null +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_load_callback_device.cu @@ -0,0 +1,47 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +/* + * Device callback for LTO + nvcc offline compilation example + * Load callback for zero-padding + * + */ + +#include + +// Callback parameters structure for padding +struct PaddingCallbackParams { + unsigned int signal_size; + unsigned int padded_signal_size; +}; + +// Load callback for R2C: returns 0 for indices beyond signal_size +__device__ cufftReal padding_load_callback(void* dataIn, + unsigned long long offset, + void* callerInfo, + void* sharedPointer) { + PaddingCallbackParams* params = static_cast(callerInfo); + cufftReal* in = static_cast(dataIn); + + // If offset is beyond the original data, return zero + if (offset >= params->signal_size) { + return 0.0f; + } + + // Return the real input data (non-interleaved format) + return in[offset]; +} diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_store_callback_device.cu b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_store_callback_device.cu new file mode 100644 index 00000000..e0202f9a --- /dev/null +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_store_callback_device.cu @@ -0,0 +1,47 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +/* + * Device callback for LTO + nvcc offline compilation example + * Store callback for zero-padding + * + */ + +#include + +// Callback parameters structure for padding +struct PaddingCallbackParams { + unsigned int signal_size; + unsigned int padded_signal_size; +}; + +// Store callback for C2R: only stores first signal_size elements, zeros the rest +__device__ void padding_store_callback(void* dataOut, + unsigned long long offset, + cufftReal element, + void* callerInfo, + void* sharedPointer) { + PaddingCallbackParams* params = static_cast(callerInfo); + cufftReal* out = static_cast(dataOut); + + // Only store if within original size, zero and normalize otherwise + if (offset < params->signal_size) { + out[offset] = element / params->padded_signal_size; + } else { + out[offset] = 0.0f; + } +} diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_reference.cu b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_reference.cu new file mode 100644 index 00000000..47e35270 --- /dev/null +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_reference.cu @@ -0,0 +1,97 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +/* + * Reference for the padding example, + * performs zero-padding manually without callbacks. + * + */ + +#include +#include +#include +#include "r2c_c2r_padding_reference.h" + +int reference_r2c_padding_c2r(unsigned batches, unsigned signal_size, unsigned padded_signal_size, float* input_signals, float* output_signals) { + const size_t padded_real_size_bytes = batches * padded_signal_size * sizeof(float); + const size_t complex_size_bytes = batches * (padded_signal_size / 2 + 1) * 2 * sizeof(float); + + // Allocate device memory for padded real input + float *device_padded_input; + CHECK_ERROR(cudaMalloc((void **)&device_padded_input, padded_real_size_bytes)); + + // Allocate device memory for complex data + float *device_complex; + CHECK_ERROR(cudaMalloc((void **)&device_complex, complex_size_bytes)); + + // Allocate device memory for padded real output + float *device_padded_output; + CHECK_ERROR(cudaMalloc((void **)&device_padded_output, padded_real_size_bytes)); + + // Zero-fill the padded input buffer + CHECK_ERROR(cudaMemset(device_padded_input, 0, padded_real_size_bytes)); + + // Copy input to device (only first signal_size elements per batch) + for (unsigned b = 0; b < batches; b++) { + CHECK_ERROR(cudaMemcpy(device_padded_input + b * padded_signal_size, + input_signals + b * signal_size, + signal_size * sizeof(float), + cudaMemcpyHostToDevice)); + } + + // Create cuFFT plans + cufftHandle forward_plan, inverse_plan; + size_t work_size; + + CHECK_ERROR(cufftCreate(&forward_plan)); + CHECK_ERROR(cufftCreate(&inverse_plan)); + + CHECK_ERROR(cufftMakePlan1d(forward_plan, padded_signal_size, CUFFT_R2C, batches, &work_size)); + CHECK_ERROR(cufftMakePlan1d(inverse_plan, padded_signal_size, CUFFT_C2R, batches, &work_size)); + + // Transform signal forward (R2C) on padded data + printf("Transforming reference cufftExecR2C\n"); + CHECK_ERROR(cufftExecR2C(forward_plan, (cufftReal *)device_padded_input, (cufftComplex *)device_complex)); + + // Transform signal inverse (C2R) + printf("Transforming reference cufftExecC2R\n"); + CHECK_ERROR(cufftExecC2R(inverse_plan, (cufftComplex *)device_complex, (cufftReal *)device_padded_output)); + + // Copy padded output to host for normalization + float *host_padded_output = new float[batches * padded_signal_size]; + CHECK_ERROR(cudaMemcpy(host_padded_output, device_padded_output, batches * padded_signal_size * sizeof(float), cudaMemcpyDeviceToHost)); + + // Normalize and copy only first signal_size elements to output + for (unsigned b = 0; b < batches; b++) { + for (unsigned i = 0; i < signal_size; i++) { + output_signals[b * signal_size + i] = host_padded_output[b * padded_signal_size + i] / static_cast(padded_signal_size); + } + } + + delete[] host_padded_output; + + // Destroy CUFFT context + CHECK_ERROR(cufftDestroy(forward_plan)); + CHECK_ERROR(cufftDestroy(inverse_plan)); + + // Cleanup memory + CHECK_ERROR(cudaFree(device_padded_input)); + CHECK_ERROR(cudaFree(device_complex)); + CHECK_ERROR(cudaFree(device_padded_output)); + + return PASS_VALUE; +} diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_reference.h b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_reference.h new file mode 100644 index 00000000..18a2b74d --- /dev/null +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_reference.h @@ -0,0 +1,26 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + + +#ifndef R2C_C2R_PADDING_REFERENCE__H_ +#define R2C_C2R_PADDING_REFERENCE__H_ + +#include "common.h" + +int reference_r2c_padding_c2r(unsigned batches, unsigned signal_size, unsigned padded_signal_size, float* input_signals, float* output_signals); + +#endif // R2C_C2R_PADDING_REFERENCE__H_ From 06b0e12234199efb8173116f21835504ac57636e Mon Sep 17 00:00:00 2001 From: Thibault Cimic Date: Fri, 3 Jul 2026 15:19:10 +0200 Subject: [PATCH 4/6] fix(padding): add missing resource cleanup in error paths In the error path when reference_r2c_padding_c2r() fails, the code was only freeing host memory but not CUDA resources. This adds the missing: - cufftDestroy() for both forward and inverse plans - cudaFree() for device_signals, device_complex, and device_params Files: - cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_legacy_callback_example.cu - cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_callback_example.cpp Signed-off-by: Thibault Cimic --- .../src/r2c_c2r_padding_legacy_callback_example.cu | 11 +++++------ .../src/r2c_c2r_padding_lto_callback_example.cpp | 11 +++++------ 2 files changed, 10 insertions(+), 12 deletions(-) diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_legacy_callback_example.cu b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_legacy_callback_example.cu index 7d2ce82e..322444e8 100644 --- a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_legacy_callback_example.cu +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_legacy_callback_example.cu @@ -27,12 +27,6 @@ #include "r2c_c2r_padding_reference.h" #include "callback_params.h" -// Callback parameters structure for padding -struct PaddingCallbackParams { - unsigned int signal_size; - unsigned int padded_signal_size; -}; - // Load callback for R2C: returns 0 for indices beyond signal_size __device__ cufftReal padding_load_callback_r2c(void* input, size_t index, @@ -146,6 +140,11 @@ int test_r2c_padding_c2r() { delete[] input_signals; delete[] output_signals; delete[] reference; + cufftDestroy(forward_plan); + cufftDestroy(inverse_plan); + cudaFree(device_signals); + cudaFree(device_complex); + cudaFree(device_params); return ERROR_VALUE; } diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_callback_example.cpp b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_callback_example.cpp index 2e945390..30cb2d47 100644 --- a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_callback_example.cpp +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_callback_example.cpp @@ -31,12 +31,6 @@ #include "r2c_c2r_padding_load_callback_fatbin.h" #include "r2c_c2r_padding_store_callback_fatbin.h" -// Callback parameters structure for padding -struct PaddingCallbackParams { - unsigned int signal_size; - unsigned int padded_signal_size; -}; - static_assert(padded_signal_size > signal_size, "The padded size must be larger than the signal size"); int test_r2c_padding_c2r() { @@ -120,6 +114,11 @@ int test_r2c_padding_c2r() { delete[] input_signals; delete[] output_signals; delete[] reference; + cufftDestroy(forward_plan); + cufftDestroy(inverse_plan); + cudaFree(device_signals); + cudaFree(device_complex); + cudaFree(device_params); return ERROR_VALUE; } From fcabf332efd78e2c88c7b6a91fef5efe2fc0ed80 Mon Sep 17 00:00:00 2001 From: Thibault Cimic Date: Fri, 3 Jul 2026 15:19:10 +0200 Subject: [PATCH 5/6] refactor(padding): consolidate PaddingCallbackParams into shared header The PaddingCallbackParams struct was duplicated in 4 files. This consolidates it into callback_params.h and updates all files to include the header. Files: - cuFFT/lto_callback_window_1d/src/callback_params.h (added struct) - cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_legacy_callback_example.cu - cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_callback_example.cpp - cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_load_callback_device.cu - cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_store_callback_device.cu Signed-off-by: Thibault Cimic --- cuFFT/lto_callback_window_1d/src/callback_params.h | 6 ++++++ .../src/r2c_c2r_padding_lto_load_callback_device.cu | 7 +------ .../src/r2c_c2r_padding_lto_store_callback_device.cu | 7 +------ 3 files changed, 8 insertions(+), 12 deletions(-) diff --git a/cuFFT/lto_callback_window_1d/src/callback_params.h b/cuFFT/lto_callback_window_1d/src/callback_params.h index 9f203a2f..437f1a6a 100644 --- a/cuFFT/lto_callback_window_1d/src/callback_params.h +++ b/cuFFT/lto_callback_window_1d/src/callback_params.h @@ -25,6 +25,12 @@ struct cb_params { unsigned signal_size; }; +// Callback parameters structure for padding +struct PaddingCallbackParams { + unsigned int signal_size; + unsigned int padded_signal_size; +}; + // Problem input parameters constexpr unsigned batches = 830; constexpr unsigned signal_size = 328; diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_load_callback_device.cu b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_load_callback_device.cu index 5e147b77..7a23d4ba 100644 --- a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_load_callback_device.cu +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_load_callback_device.cu @@ -22,12 +22,7 @@ */ #include - -// Callback parameters structure for padding -struct PaddingCallbackParams { - unsigned int signal_size; - unsigned int padded_signal_size; -}; +#include "callback_params.h" // Load callback for R2C: returns 0 for indices beyond signal_size __device__ cufftReal padding_load_callback(void* dataIn, diff --git a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_store_callback_device.cu b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_store_callback_device.cu index e0202f9a..756b8906 100644 --- a/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_store_callback_device.cu +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_store_callback_device.cu @@ -22,12 +22,7 @@ */ #include - -// Callback parameters structure for padding -struct PaddingCallbackParams { - unsigned int signal_size; - unsigned int padded_signal_size; -}; +#include "callback_params.h" // Store callback for C2R: only stores first signal_size elements, zeros the rest __device__ void padding_store_callback(void* dataOut, From 522914bb1406a5e835a0c3d155bcf322bcaed746 Mon Sep 17 00:00:00 2001 From: Thibault Cimic Date: Fri, 3 Jul 2026 15:19:10 +0200 Subject: [PATCH 6/6] refactor(common): remove unused init_input_signals_padding The init_input_signals_padding function was never called in any example and had a hidden bug (used += without zeroing the buffer first). This removes the unused function and its declaration. Files: - cuFFT/lto_callback_window_1d/src/common.cpp - cuFFT/lto_callback_window_1d/src/common.h Signed-off-by: Thibault Cimic --- cuFFT/lto_callback_window_1d/src/common.cpp | 29 --------------------- cuFFT/lto_callback_window_1d/src/common.h | 1 - 2 files changed, 30 deletions(-) diff --git a/cuFFT/lto_callback_window_1d/src/common.cpp b/cuFFT/lto_callback_window_1d/src/common.cpp index ddfa10ac..a220dcaf 100644 --- a/cuFFT/lto_callback_window_1d/src/common.cpp +++ b/cuFFT/lto_callback_window_1d/src/common.cpp @@ -61,34 +61,5 @@ void init_input_signals(unsigned batches, unsigned signal_size, float* signals) } } -// Initialize the input signal for padding examples -// Uses real buffer layout (stride = signal_size) -void init_input_signals_padding(unsigned batches, unsigned signal_size, float* signals) { - std::mt19937 e2(0); - - std::uniform_real_distribution<> A_dist(0., signal_max_A); - std::uniform_real_distribution<> f_dist(0., signal_max_f); - - for(unsigned batch = 0; batch < batches; ++batch) { - std::vector wave_amplitudes; - std::vector wave_frequencies; - - // Generate the amplitudes and frequencies of the waves - for(unsigned i = 0; i < waves; ++i) { - wave_amplitudes.push_back(A_dist(e2)); - wave_frequencies.push_back(f_dist(e2)); - } - - // Compose the signal - float time = 0.; - for(unsigned s = 0; s < signal_size; ++s) { - for(unsigned i = 0; i < waves; ++i) { - unsigned idx = batch * signal_size + s; - signals[idx] += wave_amplitudes[i] * sin(2. * PI * wave_frequencies[i] * time) ; - } - time += sampling_dt; - } - } -} #endif // _COMMON__CPP_ \ No newline at end of file diff --git a/cuFFT/lto_callback_window_1d/src/common.h b/cuFFT/lto_callback_window_1d/src/common.h index b1fc6d1f..49266457 100644 --- a/cuFFT/lto_callback_window_1d/src/common.h +++ b/cuFFT/lto_callback_window_1d/src/common.h @@ -82,5 +82,4 @@ double compute_error_padding(T* ref, T* out, unsigned batches, unsigned signal_s // Initialize input signals void init_input_signals(unsigned batches, unsigned signal_size, float* signals); -void init_input_signals_padding(unsigned batches, unsigned signal_size, float* signals); #endif // _COMMON__H_ \ No newline at end of file