Add padding callback examples demonstrating load and store callbacks

[ThibaultGH]

Add padding callback examples demonstrating load and store callbacks

This PR adds new examples to demonstrate the use of both load and store callbacks with cuFFT, complementing the existing windowing examples that only show load callbacks.

New Examples

• 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_reference.cu: Reference implementation without callbacks

Key Differences from Windowing Examples

• Load callback on forward FFT (R2C): Implements zero-padding by returning 0 for indices beyond the original signal size
• Store callback on inverse FFT (C2R): Truncates output to original size and normalizes
• Demonstrates complete round-trip with both callback types

Implementation Notes

• Followed the existing code structure and naming conventions from the windowing examples
• Happy to make adjustments if the structure or implementation doesn't match NVIDIA's preferences
• Currently only includes legacy and LTO (nvcc) variants
• Note: LTO+NVRTC variant not included in this PR, but can be added later if needed

Testing

• All examples build cleanly with no warnings
• All examples run successfully with L2 error ≈ 0 (within threshold)
• Existing windowing examples remain unaffected

[greptile-apps]

Greptile Summary

This PR adds zero-padding examples to the existing lto_callback_window_1d cuFFT sample directory, complementing the windowing examples with a demonstration of both load and store callbacks used together. The existing source files are also renamed with a windowing_ prefix to disambiguate them from the new padding_ files.

• New files added: r2c_c2r_padding_legacy_callback_example.cu, r2c_c2r_padding_lto_callback_example.cpp, two LTO device callback files, and a reference implementation — all following the structure of the existing windowing examples.
• callback_params.h is updated to define the shared PaddingCallbackParams struct and padding-specific constants; all four padding files include this header rather than duplicating the struct.
• CMakeLists.txt adds new build targets and two add_custom_command blocks (for the load and store fatbins), and all existing targets are renamed accordingly.

Confidence Score: 5/5

Safe to merge. The new padding examples are logically correct and follow the established windowing example patterns without introducing regressions.

The load and store callbacks implement zero-padding correctly: the load callback returns 0 for indices beyond the original signal, the store callback normalizes and truncates the C2R output, and the reference implementation performs the same operations explicitly. The CMakeLists.txt fatbin symbol names match the bin2c --name arguments and the LTO example string literals. The PaddingCallbackParams struct is now shared via callback_params.h. Existing windowing examples are renamed cleanly with no functional changes.

No files require special attention. The reference implementation and both callback examples are consistent with each other.

Important Files Changed

Filename	Overview
cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_legacy_callback_example.cu	New legacy-callback padding example. Load callback correctly zero-pads R2C input; store callback correctly truncates and normalizes C2R output.
cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_callback_example.cpp	New LTO-callback padding example. Correctly calls cufftXtSetJITCallback before cufftMakePlan1d, fatbin symbol names match bin2c --name arguments in CMakeLists.txt.
cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_reference.cu	Reference implementation correctly zero-fills, transforms, and normalizes. Normalization divisor and output truncation match what the callbacks do.
cuFFT/lto_callback_window_1d/src/callback_params.h	Adds PaddingCallbackParams struct and padding constants in a single shared header; all four padding translation units include this header.
cuFFT/lto_callback_window_1d/src/common.h	Renames compute_error to compute_error_windowing and adds compute_error_padding using signal_size stride; both templates use std::norm correctly.
cuFFT/lto_callback_window_1d/CMakeLists.txt	Adds two new add_custom_command blocks for load/store fatbins, new executable targets, and renames all windowing targets correctly.
cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_load_callback_device.cu	Device-side load callback; correctly returns 0 for offsets >= signal_size without reading uninitialized memory.
cuFFT/lto_callback_window_1d/src/r2c_c2r_padding_lto_store_callback_device.cu	Device-side store callback; division by padded_signal_size (unsigned int) is correctly promoted to float before dividing cufftReal.

Sequence Diagram

%%{init: {'theme': 'neutral'}}%%
sequenceDiagram
    participant Host
    participant LoadCB as Load Callback (R2C)
    participant cuFFT_R2C as cuFFT R2C Plan
    participant cuFFT_C2R as cuFFT C2R Plan
    participant StoreCB as Store Callback (C2R)

    Host->>cuFFT_R2C: cufftMakePlan1d(padded_signal_size)
    Host->>cuFFT_R2C: cufftXtSetJITCallback(CUFFT_CB_LD_REAL)
    Host->>cuFFT_C2R: cufftMakePlan1d(padded_signal_size)
    Host->>cuFFT_C2R: cufftXtSetJITCallback(CUFFT_CB_ST_REAL)

    Host->>cuFFT_R2C: cufftExecR2C(device_signals to device_complex)
    loop For each index 0..padded_signal_size-1
        cuFFT_R2C->>LoadCB: read element at offset
        alt "offset < signal_size"
            LoadCB-->>cuFFT_R2C: return in[offset]
        else "offset >= signal_size"
            LoadCB-->>cuFFT_R2C: return 0.0f (virtual zero-pad)
        end
    end

    Host->>cuFFT_C2R: cufftExecC2R(device_complex to device_signals)
    loop For each index 0..padded_signal_size-1
        cuFFT_C2R->>StoreCB: write element at offset
        alt "offset < signal_size"
            StoreCB-->>cuFFT_C2R: "out[offset] = element / padded_signal_size"
        else "offset >= signal_size"
            StoreCB-->>cuFFT_C2R: "out[offset] = 0.0f (truncate)"
        end
    end

    Host->>Host: cudaMemcpy(output_signals, device_signals, signal_size)
    Host->>Host: compute_error_padding(reference, output_signals)

Loading

%%{init: {'theme': 'base', 'themeVariables': {"darkMode": true, "background": "#0d1117", "primaryColor": "#21262d", "primaryTextColor": "#e6edf3", "primaryBorderColor": "#8b949e", "lineColor": "#8b949e", "textColor": "#e6edf3", "edgeLabelBackground": "#161b22", "actorBkg": "#21262d", "actorBorder": "#8b949e", "actorTextColor": "#e6edf3", "actorLineColor": "#8b949e", "signalColor": "#8b949e", "signalTextColor": "#e6edf3", "noteBkgColor": "#373320", "noteBorderColor": "#d4a72c", "noteTextColor": "#f0e6c0", "labelBoxBkgColor": "#21262d", "labelBoxBorderColor": "#8b949e", "labelTextColor": "#e6edf3", "loopTextColor": "#e6edf3", "activationBkgColor": "#30363d", "activationBorderColor": "#8b949e"}}}%%
sequenceDiagram
    participant Host
    participant LoadCB as Load Callback (R2C)
    participant cuFFT_R2C as cuFFT R2C Plan
    participant cuFFT_C2R as cuFFT C2R Plan
    participant StoreCB as Store Callback (C2R)

    Host->>cuFFT_R2C: cufftMakePlan1d(padded_signal_size)
    Host->>cuFFT_R2C: cufftXtSetJITCallback(CUFFT_CB_LD_REAL)
    Host->>cuFFT_C2R: cufftMakePlan1d(padded_signal_size)
    Host->>cuFFT_C2R: cufftXtSetJITCallback(CUFFT_CB_ST_REAL)

    Host->>cuFFT_R2C: cufftExecR2C(device_signals to device_complex)
    loop For each index 0..padded_signal_size-1
        cuFFT_R2C->>LoadCB: read element at offset
        alt "offset < signal_size"
            LoadCB-->>cuFFT_R2C: return in[offset]
        else "offset >= signal_size"
            LoadCB-->>cuFFT_R2C: return 0.0f (virtual zero-pad)
        end
    end

    Host->>cuFFT_C2R: cufftExecC2R(device_complex to device_signals)
    loop For each index 0..padded_signal_size-1
        cuFFT_C2R->>StoreCB: write element at offset
        alt "offset < signal_size"
            StoreCB-->>cuFFT_C2R: "out[offset] = element / padded_signal_size"
        else "offset >= signal_size"
            StoreCB-->>cuFFT_C2R: "out[offset] = 0.0f (truncate)"
        end
    end

    Host->>Host: cudaMemcpy(output_signals, device_signals, signal_size)
    Host->>Host: compute_error_padding(reference, output_signals)

Loading

_{Reviews (3): Last reviewed commit: "refactor(common): remove unused init_inp..." | Re-trigger Greptile}