BraneTechnologies Documentation

Compiling for NVIDIA GPUs (CUDA & OpenCL) #

For NVIDIA GPUs, the Brane SDK offers comprehensive support for both CUDA (using NVIDIA’s nvcc compiler) and OpenCL for vendor-agnostic GPU programming. CUDA provides direct access to NVIDIA’s hardware features, while OpenCL enables cross-platform compatibility.

saxpy/
├── build.gradle
└── src
    └── cuda
        └── cpp
            └── main.cpp

#include <stdio.h>

__global__
void saxpy(int n, float a, float *x, float *y)
{
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if (i < n) y[i] = a*x[i] + y[i];
}

int main(void)
{
  int N = 1<<20;
  float *x, *y, *d_x, *d_y;
  x = (float*)malloc(N*sizeof(float));
  y = (float*)malloc(N*sizeof(float));

  cudaMalloc(&d_x, N*sizeof(float)); 
  cudaMalloc(&d_y, N*sizeof(float));

  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);

  // Perform SAXPY on 1M elements
  saxpy<<<(N+255)/256, 256>>>(N, 2.0f, d_x, d_y);

  cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);

  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = max(maxError, abs(y[i]-4.0f));
  printf("Max error: %f\n", maxError);

  cudaFree(d_x);
  cudaFree(d_y);
  free(x);
  free(y);
}

plugins {
    id 'com.brane.accelerator.gpu.cuda.cpp-application' // Plugin for HIP-based C++ applications
}

application {
    // Define the target machines for this application
    targetMachines = [machines.linux.x86_64, machines.windows.x86_64]
}

Example: Vector Addition with OpenCL

This example demonstrates how to implement and compile a vector addition kernel for execution on a Kalray MPPA processor.

Project Structure

/vectorAddition/
├── build.gradle
├── README.md
└── src
    └── main
        ├── cpp
        │   └── main.cpp
        ├── headers
        └── opencl
            └── vector_add.cl

Example OpenCL Kernel code:

// Define a kernel function for asynchronous vector addition.
__kernel void vector_add(__global int *a, __global int *b, __global int *c) {
	int i = get_global_id(0); // Get the global ID to identify the work-item.
	int j = get_local_id(0); // Get the local ID within the work-group.
	int nb_pe = get_local_size(0); // Get the number of processing elements (work-items) in the local work-group.

    // Define local memory to store a segment of the input vectors and the result vector.
	__local int la[16];
	__local int lb[16];
	__local int lc[16];

    // Asynchronously copy a segment of the 'a' vector from global to local memory.
    // 'eva' will hold the event associated with this operation.
	event_t eva = async_work_group_copy(la, &a[i], nb_pe, 0);

    // Asynchronously copy a segment of the 'b' vector from global to local memory.
    // 'evb' will hold the event associated with this operation.	
	event_t evb = async_work_group_copy(lb, &b[i], nb_pe, 0);

    // Wait for the asynchronous copy of the 'a' (and then 'b') vector to complete.
	wait_group_events(1, &eva);
	wait_group_events(1, &evb);

    // Perform the addition operation locally on the PE (processing element).
	lc[j] = la[j] + lb[j];

    // Synchronize all work-items in the work-group to ensure addition is completed.
	barrier(CLK_LOCAL_MEM_FENCE);

    // Asynchronously copy the result vector 'lc' from local to global memory.
    // 'evc' will hold the event associated with this operation.
	event_t evc = async_work_group_copy(&c[i], lc, nb_pe, 0);

    // Wait for the asynchronous copy of the 'c' vector to complete.
	wait_group_events(1, &evc);
}

This kernel showcases Kalray-specific optimizations, particularly the use of asynchronous memory operations and local memory to improve performance on the MPPA architecture.

Example main c++ code:

/*
 * Copyright (C) 2025 Brane Technologies LLC - All Rights Reserved
 * Licensed under the Apache License, Version 2.0 (the "License");
 */

#include <iostream>
#include <vector>
#include <fstream>
#include <sstream>
#include "opencl.hpp" // Include OpenCL-Wrapper

const int ARRAY_SIZE = 1024;

// Function to read OpenCL kernel file
std::string read_file(const std::string& filename) {
    std::ifstream file(filename);
    if (!file.is_open()) {
        throw std::runtime_error("Failed to open kernel file: " + filename);
    }
    std::ostringstream content;
    content << file.rdbuf();
    return content.str();
}

int main() {
    try {
        // List available OpenCL devices
        std::vector<Device_Info> devices = get_devices();
        for (const auto& device : devices) {
            print_device_info(device);
        }

        // Select the best device (highest FLOPS)
        Device_Info selected_device = select_device_with_most_flops(devices);
        std::cout << "\nSelected Device:\n";
        print_device_info(selected_device);

        // Load and compile OpenCL kernel
        std::string kernel_code = read_file("vector_add.cl");
        Device device(selected_device, kernel_code); // Create OpenCL device

        // Allocate memory on both Host & Device
        std::vector<int> a(ARRAY_SIZE), b(ARRAY_SIZE), result(ARRAY_SIZE, 0);
        for (int i = 0; i < ARRAY_SIZE; i++) {
            a[i] = i;
            b[i] = i * 2;
        }

        Memory<int> bufferA(device, ARRAY_SIZE, 1, a.data());
        Memory<int> bufferB(device, ARRAY_SIZE, 1, b.data());
        Memory<int> bufferC(device, ARRAY_SIZE, 1);  // Output buffer

        // Determine optimal workgroup size
        size_t max_workgroup_size;
        selected_device.cl_device.getInfo(CL_DEVICE_MAX_WORK_GROUP_SIZE, &max_workgroup_size);
        size_t workgroup_size = std::min(max_workgroup_size, (size_t)64);

        // Create and run the kernel
        Kernel vector_add_kernel(device, ARRAY_SIZE, workgroup_size, "vector_add", bufferA, bufferB, bufferC);

        bufferA.write_to_device();
        bufferB.write_to_device();

        #define NB_ITER (10)
        for (int i = 0; i < NB_ITER; ++i) {
            vector_add_kernel.run();

            bufferC.read_from_device();
            int* dst = bufferC();  // Access results

            // Verify results
            bool success = true;
            for (int j = 0; j < ARRAY_SIZE; ++j) {
                int expected = a[j] + b[j];
                if (dst[j] != expected) {
                    std::cerr << "Mismatch at index " << j << ": " << expected << " != " << dst[j] << std::endl;
                    success = false;
                }
            }

            if (!success) {
                std::cerr << "Verification failed in iteration " << i << "!" << std::endl;
                return EXIT_FAILURE;
            }

            std::cout << "Iteration " << i << " passed." << std::endl;

            // Increment input buffers for next iteration
            for (int j = 0; j < ARRAY_SIZE; ++j) {
                a[j]++;
                b[j]++;
            }
            bufferA.write_to_device();
            bufferB.write_to_device();
        }

        std::cout << "Success: All iterations completed correctly!" << std::endl;
    } catch (const std::exception &e) {
        std::cerr << "Error: " << e.what() << std::endl;
        return EXIT_FAILURE;
    }

    return EXIT_SUCCESS;
}

The host code demonstrates how to discover OpenCL devices, select the most powerful one, set up memory buffers, and coordinate kernel execution with appropriate error checking.

Example build.gradle configuration:

plugins {
    id 'com.brane.coolidge.cpp'
}

application {
    targetMachines = [machines.linux.x86_64,
    coolidgeMachines.clusterOS.architecture("MPPA_v2")
    ]
}

Configuration Explanation:

• Dual-Target Setup: This configuration enables cross-compilation for both standard x86_64 CPUs and Kalray MPPA accelerators, allowing the application to be executed on the Kalray Coolidge.
• Kalray-specific Plugin: The com.brane.coolidge.cpp plugin loads the Kalray Coolidge toolchain and manages the complex compilation process for MPPA processors.

Kalray specific features #

The Brane SDK’s Kalray plugin provides several specialized features for MPPA development.

For example, it automates:

• compileOclKernels: Automatically compiles OpenCL kernels for the Kalray Coolidge.
• emulate: Runs an emulation of the OpenCL code execution on an x86_64 host and/or Kalray Coolidge CPU. Useful for rapid testing without hardware execution.
• simulate: Simulates OpenCL execution on x86_64 or Kalray Coolidge. The simulator provides higher accuracy than emulation but is significantly slower.
• jtagRun / linuxRun: Executes the OpenCL kernel on the Kalray Coolidge processor via JTAG (for low-level debugging) or directly on a Linux-based host/server.

Example: Bitcoin Mining Core #

This example demonstrates an FPGA-based Bitcoin mining implementation using Verilog.

Project Structure

/BitcoinMiner/
├── build.gradle
├── README.md
├── scripts
│   └── mine
│       ├── config.example.tcl
│       ├── json_rpc.tcl
│       ├── jtag_comm.tcl
│       ├── mine.bat
│       ├── mine.tcl
│       └── utils.tcl
└── src
    ├── IP
    │   ├── vio_0
    │   │   ├── doc
    │   │   │   └── vio_v3_0_changelog.txt
    │   │   ├── hdl
    │   │   │   ├── ltlib_v1_0_vl_rfs.v
    │   │   │   ├── verilog
    │   │   │   │   ├── ltlib_v1_0_2_lib_fn.vh
    │   │   │   │   ├── ltlib_v1_0_2_ver.vh
    │   │   │   │   ├── xsdbs_v1_0_4_i2x.vh
    │   │   │   │   └── xsdbs_v1_0_4_in.vh
    │   │   │   ├── vio_v3_0_26_vio_include.v
    │   │   │   ├── vio_v3_0_syn_rfs.v
    │   │   │   └── xsdbs_v1_0_vl_rfs.v
    │   │   ├── sim
    │   │   │   └── vio_0.v
    │   │   ├── synth
    │   │   │   └── vio_0.v
    │   │   ├── vio_0.bmj
    │   │   ├── vio_0_bmstub.v
    │   │   ├── vio_0.dcp
    │   │   ├── vio_0_ooc.xdc
    │   │   ├── vio_0_sim_netlist.v
    │   │   ├── vio_0_sim_netlist.vhdl
    │   │   ├── vio_0_stub.v
    │   │   ├── vio_0_stub.vhdl
    │   │   ├── vio_0.veo
    │   │   ├── vio_0.vho
    │   │   ├── vio_0.xdc
    │   │   └── vio_0.xml
    │   └── vio_0.xci
    ├── main
    │   └── verilog
    │       ├── fpgaminer_top.v
    │       ├── main_pll.v
    │       ├── PLLs
    │       │   ├── altera_pll.v
    │       │   └── xilinx_pll.v
    │       ├── Probes
    │       │   ├── altera_virtual_wire.v
    │       │   └── xilinx_virtual_wire.v
    │       ├── sha-256-functions.v
    │       ├── sha256_transform.v
    │       ├── uart_comm.v
    │       ├── uart_rx.v
    │       ├── uart_tx_fifo.v
    │       ├── uart_tx.v
    │       └── virtual_wire.v
    └── test
        ├── resources
        │   ├── reverse_hex.tcl
        │   ├── test_data.txt
        │   └── test_hasher.py
        └── verilog
            └── test_fpgaminer_top.v

Example build.gradle configuration:

plugins {
    id 'com.brane.hdl'
    id 'com.brane.hdl.vivado'
}

HardwarePlugin {
    top = "fpgaminer_top"
    sim = "Test_Blinking"
    revision = "rev1"
}

sourceSets {
    main {
        verilog {
            srcDirs += [] // 
        }
    }
    test {
        verilog {
            srcDirs += [] //
        }
    }
}

Configuration Explanation:

• HDL Plugin: The com.brane.hdl plugin provides general support for HDL-based designs.
• Vivado Integration: The com.brane.hdl.vivado plugin integrates the Xilinx Vivado toolchain for FPGA synthesis, implementation, and bitstream generation.
• Hardware Configuration:
  • top: Defines the top-level module of the FPGA design (fpgaminer_top)
  • sim: Specifies the testbench module for simulation
  • revision: Helps manage different hardware revisions
• Source Sets: Define the locations of Verilog source files for both implementation and testing.

Features #

The Brane SDK’s HDL plugins transform FPGA development by automating the complex, time-consuming tasks typically required when working with hardware description languages. This automation creates a streamlined development experience that allows hardware engineers to focus on their designs rather than struggling with toolchain configuration and management.

The SDK handles these essential FPGA development tasks automatically:

• Project Generation: Creates complete Vivado or Quartus projects from your source files.
• Synthesis and implementation: Runs HDL compilation and synthesis operations for Xilinx or Intel/Altera platforms.
• Simulation: Supports HDL simulation for functional verification before hardware deployment

The Brane SDK exposes several specialized Gradle tasks for FPGA development workflows:

• generateXML: Generates a generic project specification in XML format, which can be used as an intermediate representation or for documentation
• createHDLConfiguration: Creates a configuration file tailored for either Xilinx Vivado or Intel Quartus, setting up appropriate parameters for your specific FPGA target
• detectX_main: Automatically detects and catalogs main HDL source files in your project structure
• detectX_test: Automatically identifies test bench files and verification components
• checkProjectFile: Verifies that Xilinx Vivado or Intel Quartus software is installed correctly and accessible.
• generateProject: Creates a complete project in either Vivado or Quartus format.
• compileProject: Launches a full compilation and synthesis process using the target vendor’s tools
• verify: Executes simulation-based verification of your FPGA design

Compiling a Multi-Target Project

A single Brane SDK project can target multiple architectures simultaneously through a unified build configuration. The example below demonstrates a comprehensive multi-target setup:

Project Structure

multi-target/
├── accelerators
│   ├── coolidge
│   │   ├── basics
│   │   │   ├── openCL
│   │   │   │   └── vectorAddition
│   │   │   │       ├── build.gradle
│   │   │   │       ├── README.md
│   │   │   │       └── src
│   │   │   │           └── main
│   │   │   │               ├── cpp
│   │   │   │               │   └── main.cpp
│   │   │   │               ├── headers
│   │   │   │               └── opencl
│   │   │   │                   └── vector_add.cl
│   │   │   └── MPPABenchmarkPapi
│   │   │       ├── build.gradle
│   │   │       ├── README.md
│   │   │       └── src
│   │   │           ├── main
│   │   │           │   ├── c
│   │   │           │   └── headers
│   │   │           ├── Makefile
│   │   │           └── mppa
│   │   │               └── c
│   │   │                   └── main.c
│   ├── fpga
│   │   └── blinkingLeds
│   │       ├── build.gradle
│   │       └── src
│   │           ├── main
│   │           │   └── verilog
│   │           │       └── Blinking.v
│   │           └── test
│   │               └── verilog
│   │                   ├── Test_Blinking.v
│   │                   └── vector_test.txt
│   └── gpu
│       └── matrix-multiply
│           ├── build.gradle
│           └── src
│               └── hip
│                   └── cpp
│                       └── main.cpp
└── x86_64_CPU
    ├── fft
        ├── build.gradle
        ├── LICENSE.md
        └── src
            └── main
                ├── c
                │   ├── fft.c
                │   └── main.c
                └── headers
                    └── fft.h

Example Root build.gradle configuration:

This configuration defines multiple subprojects targeting different architectures:

rootProject.name = 'multi-target'

include 'multi-target:x86_64_CPU:fft'
include 'multi-target:accelerators:gpu:matrix-multiply'
include 'multi-target:accelerators:coolidge:basics:openCL:vectorAddition'
include 'multi-target:accelerators:coolidge:benchmark:MPPABenchmarkPapi'
include 'multi-target:accelerators:fpga:blinkingLeds'

Configuration Breakdown:

• x86_64_CPU:fft: Fast Fourier Transform application compiled with GCC for Intel or AMD CPUs
• gpu:matrix-multiply: Matrix multiplication using HIP for AMD GPUs
• coolidge:openCL:vectorAddition: Vector addition using OpenCL targeting Kalray Coolidge processors
• coolidge:benchmark:MPPABenchmarkPapi: Performance benchmarking application for Kalray MPPA
• fpga:blinkingLeds: Verilog project for FPGA implementation using Xilinx toolchain

The Brane SDK automates the compilation process by detecting the appropriate toolchains (GCC, NVCC, HIP Clang, kvx-gcc, Vivado/Quartus) for each target architecture.

Features #

The Brane SDK’s HDL plugins automate numerous tasks in the FPGA development workflow.

The Brane SDK’s build system intelligently manages the compilation process for multi-target projects:

1. Toolchain Detection: Automatically identifies the appropriate compilers and tools for each target
2. Parallel Execution: Compiles different targets concurrently to maximize efficiency
3. Resource Management: Optimizes system resource usage during compilation
4. Dependency Resolution: Handles inter-module dependencies across different architectures

When running (build) assemble, the system compiles all specified targets in parallel, typically completing CPU, GPU, and accelerator builds first, while FPGA synthesis (which generally takes longer) continues in the background.

Advanced Build Options and Customization

The Brane SDK provides extensive customization options for specialized build requirements. The build system can be tailored to specific project needs through:

• Custom Compiler Flags: Fine-tune optimization levels and target-specific options
• Conditional Compilation: Apply different build settings based on target architecture
• Library Integration: Specify and manage external dependencies
• Build Hooks: Insert custom build steps at various points in the process

For detailed information on advanced build configurations, refer to the SDK’s Gradle plugin documentation.