diff --git a/cuFFT/lto_callback_window_1d/CMakeLists.txt b/cuFFT/lto_callback_window_1d/CMakeLists.txt index 787a0095..5716da3c 100644 --- a/cuFFT/lto_callback_window_1d/CMakeLists.txt +++ b/cuFFT/lto_callback_window_1d/CMakeLists.txt @@ -41,59 +41,104 @@ 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) +add_executable(r2c_c2r_padding_lto_callback_example) +add_executable(r2c_c2r_padding_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_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 + ${PROJECT_SOURCE_DIR}/src/common.cpp + ${PROJECT_SOURCE_DIR}/src/r2c_c2r_windowing_reference.cu + ${CMAKE_BINARY_DIR}/r2c_c2r_windowing_lto_callback_device_fatbin.h) + +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_windowing_reference.cu) -target_sources(r2c_c2r_lto_callback_example - PRIVATE ${PROJECT_SOURCE_DIR}/src/r2c_c2r_lto_callback_example.cpp +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 - ${CMAKE_BINARY_DIR}/r2c_c2r_lto_callback_device_fatbin.h) + ${PROJECT_SOURCE_DIR}/src/r2c_c2r_windowing_reference.cu) -target_sources(r2c_c2r_lto_nvrtc_callback_example - PRIVATE ${PROJECT_SOURCE_DIR}/src/r2c_c2r_lto_nvrtc_callback_example.cpp +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_reference.cu) + ${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_legacy_callback_example - PRIVATE ${PROJECT_SOURCE_DIR}/src/r2c_c2r_legacy_callback_example.cu +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_reference.cu) + ${PROJECT_SOURCE_DIR}/src/r2c_c2r_padding_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) + +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_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} +) + +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_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) +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 a7e4b639..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: -* `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. +### 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_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 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,16 +77,24 @@ 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 +``` + +### Padding examples ``` -./bin/r2c_c2r_lto_callback_example -./bin/r2c_c2r_lto_nvrtc_callback_example -./bin/r2c_c2r_callback_example +./bin/r2c_c2r_padding_lto_callback_example +./bin/r2c_c2r_padding_legacy_callback_example ``` -Sample of output +Sample of output (windowing) ``` -$ ./bin/r2c_c2r_lto_callback_example +$ ./bin/r2c_c2r_windowing_lto_callback_example Transforming signal cufftExecR2C Transforming signal cufftExecC2R Transforming reference cufftExecR2C @@ -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..437f1a6a 100644 --- a/cuFFT/lto_callback_window_1d/src/callback_params.h +++ b/cuFFT/lto_callback_window_1d/src/callback_params.h @@ -25,12 +25,23 @@ 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; 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/common.cpp b/cuFFT/lto_callback_window_1d/src/common.cpp index ae6ffcdc..a220dcaf 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,5 @@ void init_input_signals(unsigned batches, unsigned signal_size, float* signals) } } + #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..49266457 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,23 @@ 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); #endif // _COMMON__H_ \ No newline at end of file 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..322444e8 --- /dev/null +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_legacy_callback_example.cu @@ -0,0 +1,185 @@ +/* + * 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" + +// 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; + cufftDestroy(forward_plan); + cufftDestroy(inverse_plan); + cudaFree(device_signals); + cudaFree(device_complex); + cudaFree(device_params); + 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..30cb2d47 --- /dev/null +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_callback_example.cpp @@ -0,0 +1,159 @@ +/* + * 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" + +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; + cufftDestroy(forward_plan); + cufftDestroy(inverse_plan); + cudaFree(device_signals); + cudaFree(device_complex); + cudaFree(device_params); + 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..7a23d4ba --- /dev/null +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_load_callback_device.cu @@ -0,0 +1,42 @@ +/* + * 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 +#include "callback_params.h" + +// 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..756b8906 --- /dev/null +++ b/cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_store_callback_device.cu @@ -0,0 +1,42 @@ +/* + * 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 +#include "callback_params.h" + +// 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_ 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 97% 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..9b05c064 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 @@ -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_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 95% 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..523e12c9 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"); @@ -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_lto_nvrtc_callback_example.cpp b/cuFFT/lto_callback_window_1d/src/r2c_c2r_windowing_lto_nvrtc_callback_example.cpp similarity index 96% 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..ff1f97e0 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; @@ -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; 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