glitchFocus.cpp

#include "glitchFocus.h"
#include "glitchFocus_launcher.cuh"
#include "glitchFocus_OpenCL.h"
#include "globals.h"
#include "utils.h"

bool GlitchFocusProcessor::firstTime = true;

void (GlitchFocusProcessor::* GlitchFocusProcessor::processFunc)() const = nullptr;
void (GlitchFocusProcessor::* GlitchFocusProcessor::processFuncRGBA)() const = nullptr;

cl_kernel GlitchFocusProcessor::m_openclKernel = nullptr;
cl_kernel GlitchFocusProcessor::m_openclKernelRGBA = nullptr;

GlitchFocusProcessor::GlitchFocusProcessor(int size, int rows, int cols, float focusX, float focusY, float shiftX, float shiftY, float intensity)
    : m_rows(rows), m_cols(cols), m_focusX(focusX), m_focusY(focusY), m_shiftX(shiftX), m_shiftY(shiftY), m_intensity(intensity) {
    imgSize = size;

    if (GlitchFocusProcessor::firstTime) {
        GlitchFocusProcessor::init();
        if (BaseProcessor::firstTime) {
            BaseProcessor::init();
        }
    }

    // Allocate buffers for CUDA or OpenCL
    if (isCudaAvailable()) {
        cudaStreamCreate(&this->m_cudaStream);
        allocateCUDA();
        
        // Allocate copy buffer for CUDA
        cudaMalloc(&d_img_copy, size);

        blockSizeX = 16;
        blockSizeY = 16;
        gridSizeX = (cols + blockSizeX - 1) / blockSizeX;
        gridSizeY = (rows + blockSizeY - 1) / blockSizeY;
    }
    else {
        globalWorkSize[0] = cols;
        globalWorkSize[1] = rows;
        allocateOpenCL();
        
        // Allocate copy buffer for OpenCL
        m_imgBuf_copy = clCreateBuffer(globalContextOpenCL, CL_MEM_READ_WRITE, size, nullptr, nullptr);
    }
}

GlitchFocusProcessor::~GlitchFocusProcessor() {
    if (isCudaAvailable()) {
        cudaFree(d_img);
        cudaFree(d_img_copy);
        cudaStreamDestroy(m_cudaStream);
    }
    else {
        clReleaseMemObject(m_imgBuf);
        clReleaseMemObject(m_imgBuf_copy);
    }
}

void GlitchFocusProcessor::process() const {
    (this->*processFunc)();
}

void GlitchFocusProcessor::processRGBA() const {
    (this->*processFuncRGBA)();
}

void GlitchFocusProcessor::processCUDA() const {
    // Copy original image to copy buffer
    cudaMemcpyAsync(d_img_copy, d_img, imgSize, cudaMemcpyDeviceToDevice, m_cudaStream);
    
    dim3 grid(gridSizeX, gridSizeY);
    dim3 block(blockSizeX, blockSizeY);
    glitchFocus_CUDA(gridSizeX, gridSizeY, blockSizeX, blockSizeY, m_cudaStream, d_img, d_img_copy, m_rows, m_cols, m_focusX, m_focusY, m_shiftX, m_shiftY, m_intensity);
}

void GlitchFocusProcessor::processRGBA_CUDA() const {
    // Copy original image to copy buffer
    cudaMemcpyAsync(d_img_copy, d_img, imgSize, cudaMemcpyDeviceToDevice, m_cudaStream);
    
    dim3 grid(gridSizeX, gridSizeY);
    dim3 block(blockSizeX, blockSizeY);
    glitchFocusRGBA_CUDA(gridSizeX, gridSizeY, blockSizeX, blockSizeY, m_cudaStream, d_img, d_img_copy, m_rows, m_cols, m_focusX, m_focusY, m_shiftX, m_shiftY, m_intensity);
}

void GlitchFocusProcessor::processOpenCL() const {
    // Copy original image to copy buffer
    clEnqueueCopyBuffer(globalQueueOpenCL, m_imgBuf, m_imgBuf_copy, 0, 0, imgSize, 0, nullptr, nullptr);
    
    clSetKernelArg(m_openclKernel, 0, sizeof(cl_mem), &m_imgBuf);
    clSetKernelArg(m_openclKernel, 1, sizeof(cl_mem), &m_imgBuf_copy);
    clSetKernelArg(m_openclKernel, 2, sizeof(int), &m_rows);
    clSetKernelArg(m_openclKernel, 3, sizeof(int), &m_cols);
    clSetKernelArg(m_openclKernel, 4, sizeof(float), &m_focusX);
    clSetKernelArg(m_openclKernel, 5, sizeof(float), &m_focusY);
    clSetKernelArg(m_openclKernel, 6, sizeof(int), &m_shiftX);
    clSetKernelArg(m_openclKernel, 7, sizeof(int), &m_shiftY);
    clSetKernelArg(m_openclKernel, 8, sizeof(float), &m_intensity);

    cl_event kernelEvent;
    clEnqueueNDRangeKernel(globalQueueOpenCL,
        m_openclKernel,
        2,
        nullptr,
        globalWorkSize,
        nullptr,
        0, nullptr, &kernelEvent);

    // ⏳ Wait for kernel to complete
    clWaitForEvents(1, &kernelEvent);

    // Release the kernel event
    clReleaseEvent(kernelEvent);
}

void GlitchFocusProcessor::processRGBA_OpenCL() const {
    // Copy original image to copy buffer
    clEnqueueCopyBuffer(globalQueueOpenCL, m_imgBuf, m_imgBuf_copy, 0, 0, imgSize, 0, nullptr, nullptr);
    
    clSetKernelArg(m_openclKernelRGBA, 0, sizeof(cl_mem), &m_imgBuf);
    clSetKernelArg(m_openclKernelRGBA, 1, sizeof(cl_mem), &m_imgBuf_copy);
    clSetKernelArg(m_openclKernelRGBA, 2, sizeof(int), &m_rows);
    clSetKernelArg(m_openclKernelRGBA, 3, sizeof(int), &m_cols);
    clSetKernelArg(m_openclKernelRGBA, 4, sizeof(float), &m_focusX);
    clSetKernelArg(m_openclKernelRGBA, 5, sizeof(float), &m_focusY);
    clSetKernelArg(m_openclKernelRGBA, 6, sizeof(int), &m_shiftX);
    clSetKernelArg(m_openclKernelRGBA, 7, sizeof(int), &m_shiftY);
    clSetKernelArg(m_openclKernelRGBA, 8, sizeof(float), &m_intensity);

    cl_event kernelEvent;
    clEnqueueNDRangeKernel(globalQueueOpenCL,
        m_openclKernelRGBA,
        2,
        nullptr,
        globalWorkSize,
        nullptr,
        0, nullptr, &kernelEvent);

    // ⏳ Wait for kernel to complete
    clWaitForEvents(1, &kernelEvent);

    // Release the kernel event
    clReleaseEvent(kernelEvent);
}

void GlitchFocusProcessor::init() {
    if (GlitchFocusProcessor::firstTime) {
        if (isCudaAvailable()) {
            GlitchFocusProcessor::processFunc = &GlitchFocusProcessor::processCUDA;
            GlitchFocusProcessor::processFuncRGBA = &GlitchFocusProcessor::processRGBA_CUDA;
        }
        else {
            GlitchFocusProcessor::processFunc = &GlitchFocusProcessor::processOpenCL;
            GlitchFocusProcessor::processFuncRGBA = &GlitchFocusProcessor::processRGBA_OpenCL;

            GlitchFocusProcessor::m_openclKernel = openclUtils::createKernelFromSource(globalContextOpenCL, globalDeviceOpenCL, glitchFocusOpenCLKernelSource, "glitchFocus_kernel");
            GlitchFocusProcessor::m_openclKernelRGBA = openclUtils::createKernelFromSource(globalContextOpenCL, globalDeviceOpenCL, glitchFocusRGBAOpenCLKernelSource, "glitchFocusRGBA_kernel");
        }
        GlitchFocusProcessor::firstTime = false;
        if (BaseProcessor::firstTime) {
            BaseProcessor::init();
        }
    }
}