echo_effect.cu

#include <iostream>
#include <vector>
#include <cuda_runtime.h> // Required for CUDA functions
#include <cmath>          // For fmod, if needed, or just integer modulo

// Macro for checking CUDA errors
#define CUDA_CHECK(call)                                                          \
    do {                                                                          \
        cudaError_t err = call;                                                   \
        if (err != cudaSuccess) {                                                 \
            fprintf(stderr, "CUDA Error at %s:%d - %s\n", __FILE__, __LINE__,     \
                    cudaGetErrorString(err));                                     \
            exit(EXIT_FAILURE);                                                   \
        }                                                                         \
    } while (0)

// CUDA Kernel for applying echo effect to an audio block
// Each thread processes one audio sample.
__global__ void echoKernel(float* d_circularBuffer,   // Global circular buffer on device
                           float* d_inputBlock,       // Current input audio block
                           float* d_outputBlock,      // Output audio block
                           int bufferSize,            // Total size of the circular buffer
                           int delaySamples,          // Number of samples for delay
                           float decayCoefficient,    // Decay factor (0.0 to 1.0)
                           int blockStartWriteIdx,    // Starting index in circular buffer for current block
                           int blockSize)             // Size of the current audio block
{
    // Calculate global thread ID
    int tid = blockIdx.x * blockDim.x + threadIdx.x;

    // Process only if thread ID is within the current block size
    if (tid < blockSize) {
        // Get the current input sample
        float input_sample = d_inputBlock[tid];

        // Calculate the absolute index in the circular buffer for the current sample
        // This is where the current input_sample (or mixed output) will be stored.
        int current_buffer_idx = (blockStartWriteIdx + tid) % bufferSize;

        // Calculate the read pointer for the delayed sample.
        // We need to subtract delaySamples from the current position.
        // Adding bufferSize before modulo ensures a positive result for negative differences.
        int read_buffer_idx = (current_buffer_idx - delaySamples + bufferSize) % bufferSize;

        // Read the delayed sample from the circular buffer
        float delayed_sample = d_circularBuffer[read_buffer_idx];

        // Calculate the echo component (delayed and attenuated)
        float echo_component = delayed_sample * decayCoefficient;

        // Calculate the final output sample by mixing input and echo
        float output_sample = input_sample + echo_component;

        // Store the *output_sample* back into the circular buffer at the current position.
        // This creates the feedback loop for subsequent echoes.
        d_circularBuffer[current_buffer_idx] = output_sample;

        // Store the output sample in the output block to be transferred back to host
        d_outputBlock[tid] = output_sample;
    }
}

int main() {
    // --- Configuration Parameters ---
    const int SAMPLE_RATE = 44100;                 // Samples per second
    const float MAX_DELAY_TIME_SECONDS = 2.0f;     // Maximum echo delay
    const int CIRCULAR_BUFFER_SIZE = (int)(MAX_DELAY_TIME_SECONDS * SAMPLE_RATE);

    const float DESIRED_DELAY_SECONDS = 0.5f;      // Desired echo delay
    const int DELAY_SAMPLES = (int)(DESIRED_DELAY_SECONDS * SAMPLE_RATE);

    const float DECAY_COEFFICIENT = 0.6f;          // Echo decay factor (0.0 to 1.0)

    const int BLOCK_SIZE = 1024;                   // Number of samples processed per kernel launch
    const int NUM_BLOCKS_TO_SIMULATE = 50;         // Number of audio blocks to simulate

    // --- Host Variables ---
    std::vector<float> h_inputBlock(BLOCK_SIZE);
    std::vector<float> h_outputBlock(BLOCK_SIZE);

    // --- Device Pointers ---
    float* d_circularBuffer;
    float* d_inputBlock;
    float* d_outputBlock;

    // --- Allocate Device Memory ---
    CUDA_CHECK(cudaMalloc((void**)&d_circularBuffer, CIRCULAR_BUFFER_SIZE * sizeof(float)));
    CUDA_CHECK(cudaMalloc((void**)&d_inputBlock, BLOCK_SIZE * sizeof(float)));
    CUDA_CHECK(cudaMalloc((void**)&d_outputBlock, BLOCK_SIZE * sizeof(float)));

    // --- Initialize Circular Buffer on Device to Zeros ---
    CUDA_CHECK(cudaMemset(d_circularBuffer, 0, CIRCULAR_BUFFER_SIZE * sizeof(float)));

    // --- Host-side Pointer for Circular Buffer (tracks start of next write) ---
    int host_circular_buffer_write_idx = 0;

    // --- CUDA Kernel Launch Configuration ---
    // Calculate grid and block dimensions for the kernel
    // Each block will have BLOCK_SIZE threads, and we need enough blocks to cover BLOCK_SIZE samples.
    int threadsPerBlock = 256; // Common practice, can be tuned
    int blocksPerGrid = (BLOCK_SIZE + threadsPerBlock - 1) / threadsPerBlock;

    std::cout << "Simulating Real-Time Echo Effect with CUDA..." << std::endl;
    std::cout << "Circular Buffer Size: " << CIRCULAR_BUFFER_SIZE << " samples" << std::endl;
    std::cout << "Delay: " << DESIRED_DELAY_SECONDS << " seconds (" << DELAY_SAMPLES << " samples)" << std::endl;
    std::cout << "Decay Coefficient: " << DECAY_COEFFICIENT << std::endl;
    std::cout << "Block Size: " << BLOCK_SIZE << " samples" << std::endl;
    std::cout << "Simulating " << NUM_BLOCKS_TO_SIMULATE << " blocks of audio." << std::endl;
    std::cout << "---------------------------------------------------" << std::endl;

    // --- Simulate Audio Stream Processing ---
    for (int i = 0; i < NUM_BLOCKS_TO_SIMULATE; ++i) {
        // --- Simulate Input Audio (e.g., a sine wave or random noise) ---
        // In a real application, you would read this from an audio input device.
        for (int j = 0; j < BLOCK_SIZE; ++j) {
            // Simple sine wave burst to demonstrate sound
            h_inputBlock[j] = 0.5f * sinf(2.0f * M_PI * 440.0f * (float)(i * BLOCK_SIZE + j) / SAMPLE_RATE);
            // Add some noise for more dynamic input
            // h_inputBlock[j] += ((float)rand() / RAND_MAX - 0.5f) * 0.1f;
        }

        // --- Copy Input Block from Host to Device ---
        CUDA_CHECK(cudaMemcpy(d_inputBlock, h_inputBlock.data(), BLOCK_SIZE * sizeof(float), cudaMemcpyHostToDevice));

        // --- Launch CUDA Kernel ---
        echoKernel<<<blocksPerGrid, threadsPerBlock>>>(
            d_circularBuffer,
            d_inputBlock,
            d_outputBlock,
            CIRCULAR_BUFFER_SIZE,
            DELAY_SAMPLES,
            DECAY_COEFFICIENT,
            host_circular_buffer_write_idx, // Pass the starting write index for this block
            BLOCK_SIZE
        );
        CUDA_CHECK(cudaGetLastError()); // Check for errors during kernel launch

        // --- Synchronize and Copy Output Block from Device to Host ---
        CUDA_CHECK(cudaDeviceSynchronize()); // Wait for kernel to complete
        CUDA_CHECK(cudaMemcpy(h_outputBlock.data(), d_outputBlock, BLOCK_SIZE * sizeof(float), cudaMemcpyDeviceToHost));

        // --- Update Host-side Circular Buffer Write Index ---
        host_circular_buffer_write_idx = (host_circular_buffer_write_idx + BLOCK_SIZE) % CIRCULAR_BUFFER_SIZE;

        // --- (Optional) Print/Process Output Block ---
        // In a real application, you would send h_outputBlock to an audio output device.
        std::cout << "Processed Block " << i << ". First few samples of output: ";
        for (int j = 0; j < std::min(5, BLOCK_SIZE); ++j) {
            std::cout << h_outputBlock[j] << " ";
        }
        std::cout << "..." << std::endl;
    }

    std::cout << "---------------------------------------------------" << std::endl;
    std::cout << "Simulation complete." << std::endl;

    // --- Free Device Memory ---
    CUDA_CHECK(cudaFree(d_circularBuffer));
    CUDA_CHECK(cudaFree(d_inputBlock));
    CUDA_CHECK(cudaFree(d_outputBlock));

    return 0;
}