GitHub - rubber-duck-debug/gpu-lite: A lightweight CUDA header-only runtime wrapper with dynamic symbol resolution.

#include <gpulite/gpulite.hpp>

int main() {
    // Check if CUDA is available
    if (!CUDA_DRIVER_INSTANCE.loaded() || !NVRTC_INSTANCE.loaded()) {
        throw std::runtime_error("CUDA runtime not available");
    }

    // Your CUDA kernel code as a string
    const char* kernel_source = R"(
        extern "C" __global__ void vector_add(float* a, float* b, float* c, int n) {
            int idx = blockIdx.x * blockDim.x + threadIdx.x;
            if (idx < n) {
                c[idx] = a[idx] + b[idx];
            }
        }
    )";

    // Create kernel factory and cache the kernel
    auto& factory = KernelFactory::instance();
    auto* kernel = factory.create(
        "vector_add",           // kernel name
        kernel_source,          // kernel source code
        "vector_add.cu",        // virtual source filename
        {"-std=c++17"}          // compilation options
    );

    // Allocate device memory and launch kernel
    float *d_a, *d_b, *d_c;
    int n = 1024;

    // Use the cleaner GPULITE_CUDART_CALL macro (hides the global instance)
    GPULITE_CUDART_CALL(cudaMalloc(reinterpret_cast<void**>(&d_a), n * sizeof(float)));
    GPULITE_CUDART_CALL(cudaMalloc(reinterpret_cast<void**>(&d_b), n * sizeof(float)));
    GPULITE_CUDART_CALL(cudaMalloc(reinterpret_cast<void**>(&d_c), n * sizeof(float)));

    // Prepare kernel arguments
    std::vector<void*> args = {&d_a, &d_b, &d_c, &n};

    // Launch kernel
    kernel->launch(
        dim3((n + 127) / 128),  // grid size
        dim3(128),              // block size
        0,                      // shared memory size
        nullptr,                // stream
        args,                   // kernel arguments
        true                    // synchronize after launch
    );

    // Clean up
    GPULITE_CUDART_CALL(cudaFree(d_a));
    GPULITE_CUDART_CALL(cudaFree(d_b));
    GPULITE_CUDART_CALL(cudaFree(d_c));

    return 0;
}

# Save the above code as main.cpp, then compile:
g++ -std=c++17 main.cpp -ldl -o my_gpu_app

# Run the application:
./my_gpu_app

# Save the above code as main.cpp, then compile:
cl /std:c++17 main.cpp /Fe:my_gpu_app.exe

# Run the application:
my_gpu_app.exe

// Load kernel from a .cu file
auto* kernel = KernelFactory::instance().createFromSource(
    "my_kernel",
    "/path/to/kernel.cu",
    "kernel.cu",
    {"-std=c++17"}
);

// For a templated kernel like: template<typename T> __global__ void process(T* data)
std::string kernel_name = getKernelName<float>("process");
auto* kernel = factory.create(kernel_name, source, "template_kernel.cu", {});

cmake_minimum_required(VERSION 3.14)
project(MyProject LANGUAGES CXX)

include(FetchContent)

FetchContent_Declare(
    gpulite
    GIT_REPOSITORY https://github.com/rubber-duck-debug/gpu-lite.git
    GIT_TAG v1.0.0
)

FetchContent_MakeAvailable(gpulite)

add_executable(my_app main.cpp)
target_link_libraries(my_app PRIVATE gpulite)

cmake_minimum_required(VERSION 3.12)
project(MyProject LANGUAGES CXX)

add_subdirectory(external/gpulite)

add_executable(my_app main.cpp)
target_link_libraries(my_app PRIVATE gpulite)

cmake_minimum_required(VERSION 3.12)
project(MyProject LANGUAGES CXX)

set(CMAKE_CXX_STANDARD 17)

add_executable(my_app main.cpp)
target_include_directories(my_app PRIVATE path/to/gpulite)

if(NOT WIN32)
    target_link_libraries(my_app PRIVATE ${CMAKE_DL_LIBS})
endif()

std::vector<std::string> options = {
    "-std=c++17",                   // C++ standard
    "--use_fast_math",              // Fast math operations
    "-DBLOCK_SIZE=256",             // Preprocessor definitions
    "--maxrregcount=32"             // Maximum register count
};

auto* kernel = factory.create("optimized_kernel", source, "kernel.cu", options);

// Set maximum dynamic shared memory
kernel->setFuncAttribute(CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, 8192);

// Get kernel register usage
int reg_count = kernel->getFuncAttribute(CU_FUNC_ATTRIBUTE_NUM_REGS);
std::cout << "Kernel uses " << reg_count << " registers per thread" << std::endl;

// Create CUDA stream
cudaStream_t stream;
GPULITE_CUDART_CALL(cudaStreamCreate(&stream));

// Launch kernel asynchronously
kernel->launch(
    grid, block, shared_mem_size,
    reinterpret_cast<void*>(stream),
    args,
    false  // don't synchronize
);

// Do other work...

// Synchronize when needed
GPULITE_CUDART_CALL(cudaStreamSynchronize(stream));
GPULITE_CUDART_CALL(cudaStreamDestroy(stream));

try {
    auto* kernel = factory.create("my_kernel", source, "kernel.cu", {});
    kernel->launch(grid, block, 0, nullptr, args);
} catch (const std::runtime_error& e) {
    std::cerr << "CUDA Error: " << e.what() << std::endl;
    // Error message includes file and line information
}

# Run from the gpulite root directory
./scripts/build_examples_linux.sh

# Run from the gpulite root directory
scripts\build_examples_windows.bat

# Create build directory
mkdir build && cd build

# Generate Makefiles
cmake ..

# Build the project
make -j$(nproc)

# Create build directory
mkdir build
cd build

# Generate Visual Studio project files
cmake .. -G "Visual Studio 16 2019" -A x64

# Build the project
cmake --build . --config Release