Skip to content

Folders and files

NameName
Last commit message
Last commit date

Latest commit

 

History

52 Commits
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 

Repository files navigation

math-gpu-dsl

Optimized Math-DSL compiler to IR generation followed by NVRTC JIT-compiler to PTX and execution by Cuda Driver API.

Building & Requirements

Requirements

  1. C++23 compatible compiler:
    • GCC 15+
    • Clang 19+
    • MSVC 19.40+ (VS 2022 v17.10 and newer)
  2. CMake 3.30+
  3. CUDA Toolkit 13.3+

Build

  1. Clone the repository and navigate to the project directory:
git clone https://github.com/DanSynko/math-gpu-dsl.git
cd math-gpu-dsl
  1. Generate & build:
cmake -B build
cmake --build build
  1. Run the executable*:
./build/math-gpu-dsl

*The project also supports flags, see "CLI arguments" section.

CLI arguments

The application supports the following command-line options:

Flags

  • -h, --help
  • -f, --fast-math

Example

Happy Path
Welcome to all-math-gpu-calculator! This project was created for performing calculations from basic arithmetic to advanced mathematics.)
Enter:
 1) 'exit' to close the program.
 2) 'help' to show the expression printing guide and supported operators and symbols.

The entered mathematical expression will be displayed in a clean mathematical format.

Enter any mathematical expression or command:
2 * 3 * -5 ^ 3
┌──────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┐
                                                       1. TOKENS
 [index: 0;            type: Number;                 value: 2 ]
 [index: 1;            type: MultiplicationSign;     value: * ]
 [index: 2;            type: Number;                 value: 3 ]
 [index: 3;            type: MultiplicationSign;     value: * ]
 [index: 4;            type: NegationSign;           value: - ]
 [index: 5;            type: Number;                 value: 5 ]
 [index: 6;            type: PowerSign;              value: ^ ]
 [index: 7;            type: Number;                 value: 3 ]
 [index: 8;            type: EndOfFile;              value:   ]

                                          2. ABSTRACT SYNTAX TREE VALIDATION.
 Math expression is valid!
                                                   3. IR GENERATION.

v0 = ldc f64 2
v1 = ldc f64 3
v2 = mul f64 v0, v1
v3 = ldc f64 5
v4 = neg f64 v3
v5 = ldc f64 3
v6 = pow f64 v4, v5
v7 = mul f64 v2, v6
RET v7

                                                  4. IR OPTIMIZATION.

v7 = ldc f64 -750
RET v7
                                                  5 EVALUATE IN GPU.

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-37862127
// Cuda compilation tools, release 13.3, V13.3.33
// Based on NVVM 7.0.1
//

.version 9.3
.target sm_75
.address_size 64

        // .globl       custom_kernel

.visible .entry custom_kernel(
        .param .u64 custom_kernel_param_0
)
{
        .reg .b64       %rd<4>;


        ld.param.u64    %rd1, [custom_kernel_param_0];
        cvta.to.global.u64      %rd2, %rd1;
        mov.u64         %rd3, -4573563751269138432;
        st.global.u64   [%rd2], %rd3;
        ret;

}


                                                    6. RESULT:
                                                       -750

Enter any mathematical expression or command:
Error Handling
Welcome to all-math-gpu-calculator! This project was created for performing calculations from basic arithmetic to advanced mathematics.)
Enter:
 1) 'exit' to close the program.
 2) 'help' to show the expression printing guide and supported operators and symbols.

The entered mathematical expression will be displayed in a clean mathematical format.

Enter any mathematical expression or command:
2 ++ 2
┌──────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┐
                                                       1. TOKENS
 [index: 0;            type: Number;                 value: 2 ]
 [index: 1;            type: Plus;                   value: + ]
 [index: 2;            type: Plus;                   value: + ]
 [index: 3;            type: Number;                 value: 2 ]
 [index: 4;            type: EndOfFile;              value:   ]

                                          2. ABSTRACT SYNTAX TREE VALIDATION.
 Math expression is invalid. Reason:
 [index: 2;            type: Plus;                   value: + ]
Invalid operator was detected: A binary-operator have less than two operands or a postfix-operator is before the operand.

Enter any mathematical expression or command:

Feature Support & Roadmap

Implemented math

  1. General math: plus, minus, multiplication, division, power, percent.

Goals

  • General Math: variables, equalities, approximate equalities, functions, square roots, logarithms, moduli, inequalities, not equal sign.
  • Trigonometry: degrees, π, sine, cosine, tangent, cotangent, arcsine, arccosine, arctangent, atan2, arccotangent, secant.
  • Calculus: limits, right arrow("approaches", →), Euler's number, natural logarithm, derivative(and ′, ″), Partial derivative, integral, double- and triple integral, contour integral, summation sign, product sign, infinity, Nabla, Gradient, detla(δ).
  • Complex numbers: complex variable "z", imaginary unit "i", complex variable "w", imaginary unit "j", real part, imaginary part, complex conjugate(z-bar), argumant.

WIP

Variables and equal implement.

Project implementation

Program Pipeline

  1. Lexer
  2. Pratt Parser that builds Abstract Syntax Tree.
  3. IR generator that builds Three-address code(TAC) + Static Single Assignment(SSA) intermediate Representation.
  4. IR optimizer that do Constant Folding and Dead Code Elimination.
  5. Codegeneration to CUDA kernel that emitted to NVRTC.
  6. PTX from NVRTC emitted to CUDA Driver API for evaluation.

Data-Oriented Design (SoA AST)

The Abstract Syntax Tree is implemented using a Data-Oriented Design (DOD) approach as an N-ary Structure of Arrays (SoA) backed by flat std::vector buffers, replacing the traditional node-based OOP structure.

  • Traditional OOP Drawbacks: Standard pointer-based trees require dynamic memory allocation for every node. Pointer chasing leads to frequent cache misses, forcing expensive DRAM accesses (100+ clock cycles per node). Because nodes are scattered in memory, spatial locality is lost and the hardware prefetcher cannot operate efficiently.
  • DOD / SoA Advantages: Storing AST data in contiguous arrays ensures maximum cache line utilization and high spatial locality. Data is fetched into fast L1/L2/L3 CPU caches in just a few clock cycles, while predictable memory access patterns allow the hardware prefetcher to run at full efficiency.

Errors-as-values

All compilation stages (except IR generation) rely on std::expected for zero-overhead error handling instead of costly exception stack unwinding.

Compile-time bitwise lookup table

The program uses a compile-time bitwise lookup table instead of direct symbol comparison.

Monadic Error Handling

The execution pipeline in main.cpp utilizes C++23 monadic operations (and_then(), transform_error()) on std::expected to maintain a clean "Happy Path" control flow, fully decoupling error handling from the core compilation logic.

Help section

Help messages are generated at compile-time using constexpr static arrays