Build an Expression-to-PTX Compiler

Introduction

30 Second Summary

A short formula can hide a surprising amount of decision-making before a computer can run it. Parentheses and operator order must become a precise plan before hardware can perform the arithmetic.

In this project, you will build MiniPTX as a C++20 command-line compiler that transforms arithmetic expressions into PTX. The result exposes the full path from source text to an Ada GPU binary.

What You'll Build

A MiniPTX demo takes x * (2 + 3) from source text to an optimized tree before the NVIDIA CUDA Toolkit assembles generated.ptx into an sm_89 GPU binary.

By the end of this project, you'll have:

• A precedence-aware AST that displays Binary(*) above a nested Binary(+).
• A constant folding pass that turns (2 + 3) into Number(5) before GPU code generation.
• A GPU binary trail you can inspect from generated.ptx to an sm_89 CUBIN with reported register usage.
• Secret Mission: A SASS disassembly that traces your PTX arithmetic into Ada machine instructions.

Are there any prerequisites?

This project assumes intermediate C++ experience. You also need administrator access to a Windows laptop with a GeForce RTX 4050 Laptop GPU.

Before We Start

This step is your commitment to the compiler you are about to build. MiniPTX compiles expressions containing x, numeric constants, addition, multiplication, and parentheses into PTX. This demonstrates a complete compiler pipeline beyond writing a CUDA kernel directly.

Set Up the Windows Compiler Toolchain

MiniPTX needs a C++20 host compiler before it can turn source text into a tree. It also needs NVIDIA's assembler path before emitted PTX can become a GPU binary.

This step removes that toolchain uncertainty before you write any code. Visual Studio Community 2026 supplies the host compiler. The NVIDIA CUDA Toolkit 13.4.2 supplies the GPU compilation tools.

In this step, get ready to:

• Verify the Visual Studio C++ compiler from an x64 Native Tools Command Prompt.
• Confirm that the NVIDIA driver detects your GPU.
• Verify the CUDA tools before opening an empty MiniPTX workspace in Visual Studio Code.

Verify the C++20 host compiler

The x64 Native Tools Command Prompt loads the environment that lets Windows find the Microsoft C++ compiler. Searching for that prompt also reveals whether the required Visual Studio installation is available.

• Press the Windows key to open Windows search.
• Type x64 Native Tools Command Prompt into the search field.
• Choose the tab that matches what you see.

✔️ I see the expected developer prompt

• Select the x64 Native Tools Command Prompt result.

You now have a developer shell with the Visual Studio compiler environment loaded.

ⓧ I only see an older Visual Studio prompt

The older prompt belongs to a different Visual Studio installation. Install the required release so CUDA and the host compiler use the supported toolchain.

• Press the Windows key.
• Type Visual Studio Installer into the search field.
• Press Enter to open the installer.

Windows may request administrator approval. This approval only authorizes the local compiler installation.

• Use the official Visual Studio downloads page if Visual Studio Community 2026 is not listed.
• Choose Visual Studio Community 2026 in the installer.
• Select the Desktop development with C++ workload.
• Complete the installation with that workload selected.

The installation may take several minutes while Windows adds the compiler components. A quiet progress screen during that time is normal.

• Return to Windows search after the installation finishes.
• Type x64 Native Tools Command Prompt into the search field.
• Select the new developer prompt result.

ⓧ I do not see a developer prompt

A missing developer prompt means the required Visual Studio installation is unavailable. The Community edition is free for eligible individual use under its license terms.

Windows may request administrator approval during installation. This approval allows the installer to add the compiler locally.

• Open the official Visual Studio downloads page.
• Download Visual Studio Community 2026.
• Run the downloaded installer.
• Select the Desktop development with C++ workload.
• Complete the installation with that workload selected.

The installation may take several minutes while Windows adds the C++ toolchain. Let it finish before opening the developer prompt.

• Return to Windows search after the installation finishes.
• Type x64 Native Tools Command Prompt into the search field.
• Select the developer prompt result.

• Check that the developer shell can find the host compiler by running this command:

cl

What does this check?

The cl command starts Microsoft's C++ compiler. Its banner proves that the developer prompt loaded the compiler path.

You should see the Microsoft C/C++ compiler banner followed by usage information. Your host compiler is ready to build C++20 code.

Compiler command still unavailable?

• Confirm that the prompt title includes x64 Native Tools.
• Return to Visual Studio Installer to confirm that Desktop development with C++ is selected.
• Close the existing prompt after changing the installation.
• Return to a fresh x64 Native Tools Command Prompt.

Ask for help with the exact state you see: Help me diagnose why cl is unavailable in my x64 Native Tools Command Prompt after installing the Desktop development with C++ workload.

That is the host side confirmed. MiniPTX now has a compiler capable of building its front end.

Verify the NVIDIA driver and CUDA tools

The NVIDIA Graphics Driver exposes your physical GPU to Windows. The CUDA Toolkit provides nvcc for compiling PTX into a GPU binary. It also provides cuobjdump for inspecting that binary.

• Check whether the driver detects your GPU by running this command in the developer prompt:

nvidia-smi

What does this check?

The nvidia-smi command asks the installed driver for GPU status information. A listed GPU proves that Windows can communicate with the NVIDIA hardware.

✔️ I see my NVIDIA GPU

You should see GeForce RTX 4050 Laptop GPU in the output. This confirms that the driver detects the hardware MiniPTX targets.

ⓧ The GPU is not detected

Driver installation changes a system component. Save any unrelated work first because Windows may request a restart.

• Open the official GeForce driver page.
• Select GeForce RTX 4050 Laptop GPU as the product.
• Select your installed Windows version.
• Download the current NVIDIA Graphics Driver.
• Run the downloaded driver installer.
• Restart Windows if the installer requests it.
• Return to an x64 Native Tools Command Prompt after installation.
• Check GPU detection again by running:

nvidia-smi

What does this check?

This second driver query confirms whether the installation connected Windows to your NVIDIA GPU. You should now see GeForce RTX 4050 Laptop GPU in the output.

• Check the installed CUDA compiler version by running:

nvcc -V

What does this check?

The nvcc -V command prints CUDA compiler version information. The output identifies which Toolkit installation the developer shell can access.

✔️ I see the Toolkit version I need

Confirm that the output reports CUDA Toolkit 13.4.2. The compiler side of the Toolkit is ready.

• Check that the CUDA binary inspection tool is also available by running:

cuobjdump -V

What does this check?

The cuobjdump -V command prints the binary utility's version information. This proves that the disassembly tool is available for the later Secret Mission.

You should see CUDA binary utility version information without a command lookup failure.

ⓧ I see an older Toolkit version

The older Toolkit does not match this project's pinned environment. Install CUDA Toolkit 13.4.2 so the later commands use the expected compiler release.

The Toolkit installer can take several minutes. Windows may also request administrator approval before it changes the local toolchain.

• Open the CUDA Toolkit 13.4.2 download page.
• Select Windows as the operating system.
• Select x86_64 as the architecture.
• Choose the Network Installer or Full Installer.
• Download the selected installer.
• Run the downloaded installer.
• Approve the administrator request if Windows shows one.
• Complete the Toolkit installation.

Close the old developer prompt after installation. A fresh prompt loads the new Toolkit paths.

• Return to an x64 Native Tools Command Prompt.
• Verify the CUDA compiler and binary utility by running:

nvcc -V
cuobjdump -V

What do these checks prove?

• The first command confirms that the fresh prompt finds CUDA Toolkit 13.4.2.
• The second command confirms that the Toolkit's binary inspection utility is available.

You should see version information from both commands.

ⓧ The CUDA command is unavailable

The Toolkit is missing from the current environment. NVIDIA distributes the Windows driver separately from the Toolkit, so a working GPU driver does not provide nvcc.

The Toolkit installer can take several minutes. Windows may request administrator approval before installation begins.

• Open the CUDA Toolkit 13.4.2 download page.
• Select Windows as the operating system.
• Select x86_64 as the architecture.
• Choose the Network Installer or Full Installer.
• Download the selected installer.
• Run the downloaded installer.
• Approve the administrator request if Windows shows one.
• Complete the Toolkit installation.

Close the old developer prompt after installation. A fresh prompt loads the newly installed CUDA tools.

• Return to an x64 Native Tools Command Prompt.
• Verify the CUDA compiler and binary utility by running:

nvcc -V
cuobjdump -V

What do these checks prove?

• The first command confirms that the developer shell can find CUDA Toolkit 13.4.2.
• The second command confirms that the CUDA binary inspection utility is installed.

You should see version information from both tools without a command lookup failure.

CUDA tools still unavailable?

• Close every developer prompt that was open during the Toolkit installation.
• Return to a fresh x64 Native Tools Command Prompt.
• Confirm that CUDA Toolkit 13.4.2 completed installation before testing again.

Ask for help with your exact output: Help me diagnose why nvcc or cuobjdump is unavailable after installing CUDA Toolkit 13.4.2 on Windows.

The GPU side is now connected. Your driver can see the hardware. Your CUDA tools can compile and inspect GPU binaries.

Create the MiniPTX workspace and run the final check

A VS Code workspace keeps the compiler source under one named folder. The empty src directory reserves the location for the C++ file you create in the next step.

• Press the Windows key to open Windows search.
• Type Visual Studio Code into the search field.
• Press Enter to open VS Code.

VS Code opens without a project folder. You can now create the workspace directly from its folder dialog.

• Click File in the top menu.
• Select Open Folder.
• Select Desktop as the location.
• Click New Folder.
• Enter MiniPTX as the folder name.
• Click Select Folder to open the workspace.

VS Code may show a Workspace Trust prompt because this folder is newly opened. You created the empty folder yourself, so you can approve trust for this workspace.

• Approve the workspace trust prompt if it appears.

The Explorer now shows MiniPTX as the workspace folder. That folder is the local home for the compiler you build next.

• Right-click MiniPTX in the Explorer sidebar.
• Select New Folder from the context menu.
• Type src as the folder name.
• Press Enter to create the folder.

You should see an empty src folder nested beneath MiniPTX in the Explorer.

Workspace checkpoint

• MiniPTX: Created local project workspace opened in Visual Studio Code.
• MiniPTX/src: Created empty source directory for the C++ compiler source file.

Before you run the final check, predict whether all four tools will answer from the same developer shell.

• Switch back to the x64 Native Tools Command Prompt from earlier.
• Verify the complete host and NVIDIA toolchain by running these commands:

cl
nvcc -V
nvidia-smi
cuobjdump -V

What does the final check prove?

• The cl output proves that the C++ host compiler is available.
• The nvcc -V output proves that CUDA Toolkit 13.4.2 is available.
• The nvidia-smi output proves that the driver detects the GeForce RTX 4050 Laptop GPU.
• The cuobjdump -V output proves that the CUDA binary inspection tool is available.

You should see compiler information, CUDA version information, GPU status information, and binary utility version information. None of the four commands should fail because Windows cannot find the command.

One final check still fails?

• Match the failing command to the setup tab that introduced its tool.
• Return to a fresh x64 Native Tools Command Prompt after any installation change.
• Run only the failing command again before repeating the full check.

Ask for a focused diagnosis: Help me identify which MiniPTX toolchain component is missing from my final Windows verification.

Your toolchain is ready. The host compiler can build MiniPTX. NVIDIA's tools can turn its future PTX output into an Ada GPU binary.

Your Windows compiler workspace is ready. Next up, you will build the lexer, parser, and visible abstract syntax tree that give MiniPTX its front end.

Parse an Expression into an AST

Your Windows toolchain is ready. MiniPTX now needs a front end that can turn an expression into a structure the compiler can understand.

Raw source text does not expose precedence. A lexer and a precedence-aware parser solve that problem by producing an abstract syntax tree that you can inspect in the terminal.

In this step, get ready to:

• Convert expression text into tokens with lexical analysis.
• Represent numbers, variables, and binary operations in a recursive AST.
• Parse operators according to precedence with a recursive-descent parser.

Build the lexer

A lexer reads source text from left to right. It groups characters into tokens such as numbers, identifiers, operators, and parentheses.

• Select the src folder in the VS Code Explorer sidebar.
• Click the New File icon in the Explorer toolbar.
• Enter main.cpp as the file name.
• Add the required standard-library headers and token definitions by pasting this code into src\main.cpp:

#include <cctype>
#include <iostream>
#include <memory>
#include <stdexcept>
#include <string>
#include <string_view>
#include <type_traits>
#include <utility>
#include <variant>

enum class TokenKind {
    End,
    Identifier,
    Number,
    Plus,
    Star,
    LeftParen,
    RightParen
};

struct Token {
    TokenKind kind;
    std::string text;
    float number = 0.0f;
};

What do these definitions represent?

• Token kinds give each supported piece of syntax a distinct category.
• Each Token stores its category and original text.
• Number tokens also carry a parsed floating-point value for the AST.

• Save src\main.cpp.
• Check that the token definitions compile without linking by running this command in the x64 Native Tools Command Prompt:

cl /std:c++20 /EHsc /W4 /permissive- /c src\main.cpp

The compiler should finish without reporting a C++ error. This proves the token model is valid before the lexer grows around it.

Seeing a compiler error?

Confirm that the terminal is still inside the MiniPTX folder. Check that src\main.cpp is saved.

Compare the braces around TokenKind and Token with the code above.

Help me fix the token-definition compiler error.

The lexer keeps its current position inside a std::string_view. Each call to next() advances that position and returns one token.

• Add the initial Lexer class below the Token structure by pasting this code:

class Lexer {
public:
    explicit Lexer(std::string_view source) : source_(source) {}

    Token next() {
        skip_whitespace();
        if (position_ >= source_.size()) return {TokenKind::End, ""};

        const char current = source_[position_];
        switch (current) {
        case '+': ++position_; return {TokenKind::Plus, "+"};
        case '*': ++position_; return {TokenKind::Star, "*"};
        case '(': ++position_; return {TokenKind::LeftParen, "("};
        case ')': ++position_; return {TokenKind::RightParen, ")"};
        default: break;
        }

        throw std::runtime_error("Unexpected character: " + std::string(1, current));
    }

private:
    void skip_whitespace() {
        while (position_ < source_.size() && std::isspace(static_cast<unsigned char>(source_[position_]))) ++position_;
    }

    std::string_view source_;
    std::size_t position_ = 0;
};

How does the lexer advance?

• The constructor stores a view of the source expression without copying it.
• The switch handles every supported single-character token.
• The position moves forward whenever the lexer consumes a character.
• Unexpected characters stop compilation with a specific source-level error.

• Save src\main.cpp.
• Verify the initial lexer by running the compile-only check again:

cl /std:c++20 /EHsc /W4 /permissive- /c src\main.cpp

The command should complete without a C++ error. The lexer can now recognize operators and parentheses.

Does the lexer fail to compile?

Check that the Lexer class appears after the complete Token structure. Confirm that every switch case ends with the same punctuation shown above.

Help me diagnose the Lexer compiler error.

Identifiers and numbers contain several characters, so they need scanning loops. Each loop stops as soon as the next character belongs to a different token.

• Find the line that throws the unexpected-character error inside Lexer::next().
• Insert the identifier and number scanners directly above that line by pasting this code:

        if (std::isalpha(static_cast<unsigned char>(current)) || current == '_') {
            const std::size_t start = position_++;
            while (position_ < source_.size()) {
                const char next_character = source_[position_];
                if (!std::isalnum(static_cast<unsigned char>(next_character)) && next_character != '_') break;
                ++position_;
            }
            return {TokenKind::Identifier, std::string(source_.substr(start, position_ - start))};
        }

        if (std::isdigit(static_cast<unsigned char>(current)) || current == '.') {
            const std::size_t start = position_;
            bool seen_decimal_point = false;
            while (position_ < source_.size()) {
                const char next_character = source_[position_];
                if (next_character == '.') {
                    if (seen_decimal_point) break;
                    seen_decimal_point = true;
                    ++position_;
                    continue;
                }
                if (!std::isdigit(static_cast<unsigned char>(next_character))) break;
                ++position_;
            }
            const std::string text(source_.substr(start, position_ - start));
            return {TokenKind::Number, text, std::stof(text)};
        }

How do the scanning loops work?

• The identifier branch accepts letters and underscores at the start.
• Later identifier characters may also contain digits.
• The number branch tracks whether it has already consumed a decimal point.
• The final substring becomes both the token text and its floating-point value.

• Save src\main.cpp.
• Confirm that the completed lexer compiles by running:

cl /std:c++20 /EHsc /W4 /permissive- /c src\main.cpp

The compiler should finish without a C++ error. MiniPTX can now tokenize every expression form required by this project.

Is the scanner block causing errors?

Confirm that both scanner branches sit inside Lexer::next(). They must appear after the switch and before the unexpected-character error.

Help me place the scanner branches correctly.

Model the recursive AST

An AST stores the meaning of the expression as nested nodes. std::variant holds one node shape at a time, while std::unique_ptr gives each binary node ownership of its children.

• Add the AST types and construction helpers below the Lexer class by pasting this code:

struct Expr;
using ExprPtr = std::unique_ptr<Expr>;
struct Number { float value; };
struct Variable {};
struct Binary { char operation; ExprPtr left; ExprPtr right; };

struct Expr {
    using Node = std::variant<Number, Variable, Binary>;
    explicit Expr(Number number) : node(number) {}
    explicit Expr(Variable variable) : node(variable) {}
    explicit Expr(Binary binary) : node(std::move(binary)) {}
    Node node;
};

ExprPtr make_number(float value) { return std::make_unique<Expr>(Number{value}); }
ExprPtr make_variable() { return std::make_unique<Expr>(Variable{}); }
ExprPtr make_binary(char operation, ExprPtr left, ExprPtr right) {
    return std::make_unique<Expr>(Binary{operation, std::move(left), std::move(right)});
}

How can the AST contain itself?

• The forward declaration allows ExprPtr to refer to Expr before its full definition appears.
• A Number stores one floating-point constant.
• A Variable represents the supported input variable x.
• A Binary owns its left and right subexpressions.

• Save src\main.cpp.
• Verify the recursive types by running the compile-only check:

cl /std:c++20 /EHsc /W4 /permissive- /c src\main.cpp

The compiler should accept the recursive ownership model without an error. The front end now has a tree structure ready for parsed tokens.

Seeing an incomplete-type error?

Check that struct Expr; appears before the ExprPtr alias. Keep the full Expr definition after Number, Variable, and Binary.

Help me fix the recursive AST types.

Parse and print the expression

Recursive descent assigns one function to each precedence level. Primary expressions sit at the bottom, multiplication binds above addition, and the resulting call structure determines the tree shape.

• Add the initial Parser class below the AST helper functions by pasting this code:

class Parser {
public:
    explicit Parser(std::string_view source) : lexer_(source), current_(lexer_.next()) {}

private:
    ExprPtr parse_primary() {
        if (current_.kind == TokenKind::Number) {
            const float value = current_.number;
            advance();
            return make_number(value);
        }
        if (current_.kind == TokenKind::Identifier) {
            if (current_.text != "x") throw std::runtime_error("Only the variable 'x' is supported, received: " + current_.text);
            advance();
            return make_variable();
        }
        throw std::runtime_error("Expected a number, 'x', or '('");
    }

    void advance() { current_ = lexer_.next(); }
    Lexer lexer_;
    Token current_;
};

What does primary parsing handle?

• The parser starts with the first token from the lexer.
• A number token becomes a Number node.
• The identifier x becomes a Variable node.
• The advance() helper keeps token movement in one place.

• Save src\main.cpp.
• Confirm that primary parsing compiles by running:

cl /std:c++20 /EHsc /W4 /permissive- /c src\main.cpp

The compiler should complete without a C++ error. Number and variable tokens can now become AST nodes.

Parentheses restart parsing at the addition level. This allows the enclosed expression to build its own subtree before the outer operation continues.

• Find the final throw statement inside parse_primary().
• Insert the parenthesized-expression branch directly above that statement by pasting this code:

        if (current_.kind == TokenKind::LeftParen) {
            advance();
            ExprPtr expression = parse_addition();
            if (current_.kind != TokenKind::RightParen) throw std::runtime_error("Expected ')'");
            advance();
            return expression;
        }

Why call addition from parentheses?

An opening parenthesis begins a complete nested expression. Calling parse_addition() lets that nested expression use every supported precedence level.

Multiplication and addition each use a loop so repeated operators build from left to right. Their order inside the call chain gives multiplication higher precedence.

• Insert the addition and multiplication methods immediately above parse_primary() by pasting this code:

    ExprPtr parse_addition() {
        ExprPtr left = parse_multiplication();
        while (current_.kind == TokenKind::Plus) {
            advance();
            left = make_binary('+', std::move(left), parse_multiplication());
        }
        return left;
    }

    ExprPtr parse_multiplication() {
        ExprPtr left = parse_primary();
        while (current_.kind == TokenKind::Star) {
            advance();
            left = make_binary('*', std::move(left), parse_primary());
        }
        return left;
    }

How does precedence emerge?

• Addition asks multiplication to produce each operand.
• Multiplication asks primary parsing to produce each operand.
• A multiplication subtree therefore finishes before addition can attach it to a larger tree.

• Save src\main.cpp.
• Check the precedence methods by running:

cl /std:c++20 /EHsc /W4 /permissive- /c src\main.cpp

The compiler should complete without a C++ error. The parser now contains every precedence level required by MiniPTX.

Does parse_addition appear undefined?

Confirm that parse_addition() and parse_multiplication() appear before parse_primary() inside the private section of Parser.

Help me fix the parser method order.

The public parser entry point starts at the lowest-precedence operation. It also checks that no unconsumed token remains after the expression.

• Insert the public parse() method directly below the Parser constructor by pasting this code:

    ExprPtr parse() {
        ExprPtr expression = parse_addition();
        if (current_.kind != TokenKind::End) throw std::runtime_error("Unexpected token: " + current_.text);
        return expression;
    }

What does the entry point guarantee?

The parser returns one complete expression tree. The end-token check prevents valid prefixes from hiding unsupported trailing input.

• Save src\main.cpp.
• Confirm that the complete parser compiles by running:

cl /std:c++20 /EHsc /W4 /permissive- /c src\main.cpp

The compiler should finish without a C++ error. MiniPTX can now turn a complete token stream into an AST.

A tree printer makes precedence observable. Indentation shows which nodes belong beneath each binary operation.

• Add print_ast() below the Parser class by pasting this code:

void print_ast(const Expr& expression, int indentation = 0) {
    const std::string padding(static_cast<std::size_t>(indentation), ' ');
    std::visit([&](const auto& node) {
        using NodeType = std::decay_t<decltype(node)>;
        if constexpr (std::is_same_v<NodeType, Number>) {
            std::cout << padding << "Number(" << node.value << ")\n";
        } else if constexpr (std::is_same_v<NodeType, Variable>) {
            std::cout << padding << "Variable(x)\n";
        } else {
            std::cout << padding << "Binary(" << node.operation << ")\n";
            print_ast(*node.left, indentation + 2);
            print_ast(*node.right, indentation + 2);
        }
    }, expression.node);
}

How does the printer visit each node?

• The visitor selects behavior from the active std::variant alternative.
• Leaf nodes print their stored value or variable name.
• Binary nodes print their operation before recursively printing both children.
• Each recursive call adds two spaces of indentation.

• Add the program entry point below print_ast() by pasting this code:

int main(int argument_count, char* arguments[]) {
    try {
        std::string source = "x * (2 + 3)";
        for (int index = 1; index < argument_count; ++index) {
            const std::string argument = arguments[index];
            source = argument;
        }
        Parser parser(source);
        ExprPtr expression = parser.parse();
        std::cout << "Source: " << source << "\n\nRaw AST:\n";
        print_ast(*expression);
        return 0;
    } catch (const std::exception& error) {
        std::cerr << "MiniPTX error: " << error.what() << '\n';
        return 1;
    }
}

What does the program entry point do?

• The default source provides a useful expression when no argument is supplied.
• A command-line argument replaces that default source.
• The parser builds the expression tree before the printer displays it.
• The exception handler turns lexer and parser failures into readable terminal output.

• Save src\main.cpp.
• Build mini_ptx.exe by running this command in the x64 Native Tools Command Prompt:

cl /std:c++20 /EHsc /W4 /permissive- /Fe:mini_ptx.exe src\main.cpp

What does the build command select?

• The language option enables C++20 features used by the AST.
• The exception option enables standard C++ exception unwinding.
• Warning level four exposes suspicious code during compilation.
• Conformance mode keeps the source aligned with standard C++ behavior.

The build should finish successfully and create mini_ptx.exe in the MiniPTX folder. Your compiler front end is ready for its first visible test.

Before you run MiniPTX, which operation do you expect at the root of the tree?

• Run the precedence test with this command:

.\mini_ptx.exe "x * (2 + 3)"

What should the tree show?

You should see Binary(*) at the root. Its right child should be a nested Binary(+) containing Number(2) and Number(3).

That shape proves the parser respected both parentheses and operator precedence. MiniPTX has turned source text into a visible compiler data structure.

Does the AST look different?

If the executable is missing, return to the build output and fix the first reported C++ error. If the tree root is Binary(+), check that parse_addition() calls parse_multiplication() for each operand.

If parentheses produce an error, confirm that their branch calls parse_addition() before checking for RightParen.

Help me debug the MiniPTX AST output.

Compare your finished source with the complete front end for this step.

✔️ Awesome, I've got everything!

Your saved src\main.cpp now builds into a precedence-aware expression parser.

ⓧ I'd like to double check the full code

#include <cctype>
#include <iostream>
#include <memory>
#include <stdexcept>
#include <string>
#include <string_view>
#include <type_traits>
#include <utility>
#include <variant>

enum class TokenKind {
    End,
    Identifier,
    Number,
    Plus,
    Star,
    LeftParen,
    RightParen
};

struct Token {
    TokenKind kind;
    std::string text;
    float number = 0.0f;
};

class Lexer {
public:
    explicit Lexer(std::string_view source) : source_(source) {}

    Token next() {
        skip_whitespace();
        if (position_ >= source_.size()) return {TokenKind::End, ""};

        const char current = source_[position_];
        switch (current) {
        case '+': ++position_; return {TokenKind::Plus, "+"};
        case '*': ++position_; return {TokenKind::Star, "*"};
        case '(': ++position_; return {TokenKind::LeftParen, "("};
        case ')': ++position_; return {TokenKind::RightParen, ")"};
        default: break;
        }

        if (std::isalpha(static_cast<unsigned char>(current)) || current == '_') {
            const std::size_t start = position_++;
            while (position_ < source_.size()) {
                const char next_character = source_[position_];
                if (!std::isalnum(static_cast<unsigned char>(next_character)) && next_character != '_') break;
                ++position_;
            }
            return {TokenKind::Identifier, std::string(source_.substr(start, position_ - start))};
        }

        if (std::isdigit(static_cast<unsigned char>(current)) || current == '.') {
            const std::size_t start = position_;
            bool seen_decimal_point = false;
            while (position_ < source_.size()) {
                const char next_character = source_[position_];
                if (next_character == '.') {
                    if (seen_decimal_point) break;
                    seen_decimal_point = true;
                    ++position_;
                    continue;
                }
                if (!std::isdigit(static_cast<unsigned char>(next_character))) break;
                ++position_;
            }
            const std::string text(source_.substr(start, position_ - start));
            return {TokenKind::Number, text, std::stof(text)};
        }

        throw std::runtime_error("Unexpected character: " + std::string(1, current));
    }

private:
    void skip_whitespace() {
        while (position_ < source_.size() && std::isspace(static_cast<unsigned char>(source_[position_]))) ++position_;
    }

    std::string_view source_;
    std::size_t position_ = 0;
};

struct Expr;
using ExprPtr = std::unique_ptr<Expr>;
struct Number { float value; };
struct Variable {};
struct Binary { char operation; ExprPtr left; ExprPtr right; };

struct Expr {
    using Node = std::variant<Number, Variable, Binary>;
    explicit Expr(Number number) : node(number) {}
    explicit Expr(Variable variable) : node(variable) {}
    explicit Expr(Binary binary) : node(std::move(binary)) {}
    Node node;
};

ExprPtr make_number(float value) { return std::make_unique<Expr>(Number{value}); }
ExprPtr make_variable() { return std::make_unique<Expr>(Variable{}); }
ExprPtr make_binary(char operation, ExprPtr left, ExprPtr right) {
    return std::make_unique<Expr>(Binary{operation, std::move(left), std::move(right)});
}

class Parser {
public:
    explicit Parser(std::string_view source) : lexer_(source), current_(lexer_.next()) {}

    ExprPtr parse() {
        ExprPtr expression = parse_addition();
        if (current_.kind != TokenKind::End) throw std::runtime_error("Unexpected token: " + current_.text);
        return expression;
    }

private:
    ExprPtr parse_addition() {
        ExprPtr left = parse_multiplication();
        while (current_.kind == TokenKind::Plus) {
            advance();
            left = make_binary('+', std::move(left), parse_multiplication());
        }
        return left;
    }

    ExprPtr parse_multiplication() {
        ExprPtr left = parse_primary();
        while (current_.kind == TokenKind::Star) {
            advance();
            left = make_binary('*', std::move(left), parse_primary());
        }
        return left;
    }

    ExprPtr parse_primary() {
        if (current_.kind == TokenKind::Number) {
            const float value = current_.number;
            advance();
            return make_number(value);
        }
        if (current_.kind == TokenKind::Identifier) {
            if (current_.text != "x") throw std::runtime_error("Only the variable 'x' is supported, received: " + current_.text);
            advance();
            return make_variable();
        }
        if (current_.kind == TokenKind::LeftParen) {
            advance();
            ExprPtr expression = parse_addition();
            if (current_.kind != TokenKind::RightParen) throw std::runtime_error("Expected ')'");
            advance();
            return expression;
        }
        throw std::runtime_error("Expected a number, 'x', or '('");
    }

    void advance() { current_ = lexer_.next(); }
    Lexer lexer_;
    Token current_;
};

void print_ast(const Expr& expression, int indentation = 0) {
    const std::string padding(static_cast<std::size_t>(indentation), ' ');
    std::visit([&](const auto& node) {
        using NodeType = std::decay_t<decltype(node)>;
        if constexpr (std::is_same_v<NodeType, Number>) {
            std::cout << padding << "Number(" << node.value << ")\n";
        } else if constexpr (std::is_same_v<NodeType, Variable>) {
            std::cout << padding << "Variable(x)\n";
        } else {
            std::cout << padding << "Binary(" << node.operation << ")\n";
            print_ast(*node.left, indentation + 2);
            print_ast(*node.right, indentation + 2);
        }
    }, expression.node);
}

int main(int argument_count, char* arguments[]) {
    try {
        std::string source = "x * (2 + 3)";
        for (int index = 1; index < argument_count; ++index) {
            const std::string argument = arguments[index];
            source = argument;
        }
        Parser parser(source);
        ExprPtr expression = parser.parse();
        std::cout << "Source: " << source << "\n\nRaw AST:\n";
        print_ast(*expression);
        return 0;
    } catch (const std::exception& error) {
        std::cerr << "MiniPTX error: " << error.what() << '\n';
        return 1;
    }
}

Your compiler front end now turns expression text into a visible precedence-aware AST. Next, you will make MiniPTX optimize that tree by folding constant subexpressions.

Fold Constant Expressions

Your parser now turns source text into a precedence-aware abstract syntax tree that you can inspect. The tree still carries 2 + 3 as work that could happen at runtime.

Constant folding evaluates arithmetic whose inputs are already known. This step replaces the numeric addition with Number(5) before later compiler stages use the tree.

In this step, get ready to:

• Walk binary expressions from their children toward their parent.
• Replace numeric addition or multiplication subtrees with number nodes.
• Compare the raw tree with the optimized tree in the terminal.

Optimize child nodes first

A parent can become a constant only after both children have been optimized. This post-order walk lets an inner expression collapse before its result reaches the parent.

• In the Explorer sidebar, select src\main.cpp.
• Find the closing brace of print_ast().
• Add fold_constants() immediately below it by pasting this code:

ExprPtr fold_constants(ExprPtr expression) {
    auto* binary = std::get_if<Binary>(&expression->node);
    if (binary == nullptr) return expression;
    binary->left = fold_constants(std::move(binary->left));
    binary->right = fold_constants(std::move(binary->right));
    const auto* left_number = std::get_if<Number>(&binary->left->node);
    const auto* right_number = std::get_if<Number>(&binary->right->node);
    if (left_number != nullptr && right_number != nullptr) {
        if (binary->operation == '+') return make_number(left_number->value + right_number->value);
        if (binary->operation == '*') return make_number(left_number->value * right_number->value);
    }
    return expression;
}

What does this code do?

• The first check leaves number and variable nodes unchanged because they have no children to optimize.
• The recursive calls optimize both child pointers before the function examines their parent.
• The two std::get_if<Number>() calls test whether both optimized children hold numeric values.
• A numeric + or * creates a replacement number node through make_number().

• Save src\main.cpp.
• Rebuild mini_ptx.exe from the x64 Native Tools Command Prompt by running this command:

cl /std:c++20 /EHsc /W4 /permissive- /Fe:mini_ptx.exe src\main.cpp

What does this build prove?

The compiler checks the recursive ownership moves and the std::variant type queries. A successful build proves the optimization pass fits the existing C++20 AST model.

You should return to the command prompt without a compiler error. Your rebuilt executable now contains the constant-folding pass.

Seeing a compiler error?

Check that fold_constants() sits after the complete print_ast() function. A misplaced brace can accidentally nest one function inside the other.

Compare every std::move call with the snippet. The tree uses unique ownership, so copying an ExprPtr causes a compilation failure.

Help me debug my constant-folding function.

Print both tree versions

The optimizer consumes the current ExprPtr because the tree owns its children through std::unique_ptr. It returns either the original tree or a replacement number node.

• In src\main.cpp, find the line ExprPtr expression = parser.parse(); inside main().
• Replace the output statements immediately below that line with this code:

std::cout << "Source: " << source << "\n\nRaw AST:\n";
print_ast(*expression);
expression = fold_constants(std::move(expression));
std::cout << "\nOptimized AST:\n";
print_ast(*expression);

Why print both trees?

The first call to print_ast() captures the parser's original structure. The second call reveals exactly what changed after constant folding.

Keeping both views makes the optimization observable. You can now prove that parsing and optimization are separate compiler stages.

• Save src\main.cpp.

✔️ Awesome, I've got everything!

Your source now defines fold_constants() and prints the tree before and after optimization. Keep src\main.cpp saved before rebuilding.

ⓧ I'd like to double check the full code

#include <cctype>
#include <iostream>
#include <memory>
#include <stdexcept>
#include <string>
#include <string_view>
#include <type_traits>
#include <utility>
#include <variant>

enum class TokenKind {
    End,
    Identifier,
    Number,
    Plus,
    Star,
    LeftParen,
    RightParen
};

struct Token {
    TokenKind kind;
    std::string text;
    float number = 0.0f;
};

class Lexer {
public:
    explicit Lexer(std::string_view source) : source_(source) {}

    Token next() {
        skip_whitespace();
        if (position_ >= source_.size()) return {TokenKind::End, ""};

        const char current = source_[position_];
        switch (current) {
        case '+': ++position_; return {TokenKind::Plus, "+"};
        case '*': ++position_; return {TokenKind::Star, "*"};
        case '(': ++position_; return {TokenKind::LeftParen, "("};
        case ')': ++position_; return {TokenKind::RightParen, ")"};
        default: break;
        }

        if (std::isalpha(static_cast<unsigned char>(current)) || current == '_') {
            const std::size_t start = position_++;
            while (position_ < source_.size()) {
                const char next_character = source_[position_];
                if (!std::isalnum(static_cast<unsigned char>(next_character)) && next_character != '_') break;
                ++position_;
            }
            return {TokenKind::Identifier, std::string(source_.substr(start, position_ - start))};
        }

        if (std::isdigit(static_cast<unsigned char>(current)) || current == '.') {
            const std::size_t start = position_;
            bool seen_decimal_point = false;
            while (position_ < source_.size()) {
                const char next_character = source_[position_];
                if (next_character == '.') {
                    if (seen_decimal_point) break;
                    seen_decimal_point = true;
                    ++position_;
                    continue;
                }
                if (!std::isdigit(static_cast<unsigned char>(next_character))) break;
                ++position_;
            }
            const std::string text(source_.substr(start, position_ - start));
            return {TokenKind::Number, text, std::stof(text)};
        }

        throw std::runtime_error("Unexpected character: " + std::string(1, current));
    }

private:
    void skip_whitespace() {
        while (position_ < source_.size() && std::isspace(static_cast<unsigned char>(source_[position_]))) ++position_;
    }

    std::string_view source_;
    std::size_t position_ = 0;
};

struct Expr;
using ExprPtr = std::unique_ptr<Expr>;
struct Number { float value; };
struct Variable {};
struct Binary { char operation; ExprPtr left; ExprPtr right; };

struct Expr {
    using Node = std::variant<Number, Variable, Binary>;
    explicit Expr(Number number) : node(number) {}
    explicit Expr(Variable variable) : node(variable) {}
    explicit Expr(Binary binary) : node(std::move(binary)) {}
    Node node;
};

ExprPtr make_number(float value) { return std::make_unique<Expr>(Number{value}); }
ExprPtr make_variable() { return std::make_unique<Expr>(Variable{}); }
ExprPtr make_binary(char operation, ExprPtr left, ExprPtr right) {
    return std::make_unique<Expr>(Binary{operation, std::move(left), std::move(right)});
}

class Parser {
public:
    explicit Parser(std::string_view source) : lexer_(source), current_(lexer_.next()) {}

    ExprPtr parse() {
        ExprPtr expression = parse_addition();
        if (current_.kind != TokenKind::End) throw std::runtime_error("Unexpected token: " + current_.text);
        return expression;
    }

private:
    ExprPtr parse_addition() {
        ExprPtr left = parse_multiplication();
        while (current_.kind == TokenKind::Plus) {
            advance();
            left = make_binary('+', std::move(left), parse_multiplication());
        }
        return left;
    }

    ExprPtr parse_multiplication() {
        ExprPtr left = parse_primary();
        while (current_.kind == TokenKind::Star) {
            advance();
            left = make_binary('*', std::move(left), parse_primary());
        }
        return left;
    }

    ExprPtr parse_primary() {
        if (current_.kind == TokenKind::Number) {
            const float value = current_.number;
            advance();
            return make_number(value);
        }
        if (current_.kind == TokenKind::Identifier) {
            if (current_.text != "x") throw std::runtime_error("Only the variable 'x' is supported, received: " + current_.text);
            advance();
            return make_variable();
        }
        if (current_.kind == TokenKind::LeftParen) {
            advance();
            ExprPtr expression = parse_addition();
            if (current_.kind != TokenKind::RightParen) throw std::runtime_error("Expected ')'");
            advance();
            return expression;
        }
        throw std::runtime_error("Expected a number, 'x', or '('");
    }

    void advance() { current_ = lexer_.next(); }
    Lexer lexer_;
    Token current_;
};

void print_ast(const Expr& expression, int indentation = 0) {
    const std::string padding(static_cast<std::size_t>(indentation), ' ');
    std::visit([&](const auto& node) {
        using NodeType = std::decay_t<decltype(node)>;
        if constexpr (std::is_same_v<NodeType, Number>) {
            std::cout << padding << "Number(" << node.value << ")\n";
        } else if constexpr (std::is_same_v<NodeType, Variable>) {
            std::cout << padding << "Variable(x)\n";
        } else {
            std::cout << padding << "Binary(" << node.operation << ")\n";
            print_ast(*node.left, indentation + 2);
            print_ast(*node.right, indentation + 2);
        }
    }, expression.node);
}

ExprPtr fold_constants(ExprPtr expression) {
    auto* binary = std::get_if<Binary>(&expression->node);
    if (binary == nullptr) return expression;
    binary->left = fold_constants(std::move(binary->left));
    binary->right = fold_constants(std::move(binary->right));
    const auto* left_number = std::get_if<Number>(&binary->left->node);
    const auto* right_number = std::get_if<Number>(&binary->right->node);
    if (left_number != nullptr && right_number != nullptr) {
        if (binary->operation == '+') return make_number(left_number->value + right_number->value);
        if (binary->operation == '*') return make_number(left_number->value * right_number->value);
    }
    return expression;
}

int main(int argument_count, char* arguments[]) {
    try {
        std::string source = "x * (2 + 3)";
        bool body_only = false;
        for (int index = 1; index < argument_count; ++index) {
            const std::string argument = arguments[index];
            if (argument == "--body-only") body_only = true;
            else source = argument;
        }
        Parser parser(source);
        ExprPtr expression = parser.parse();
        std::cout << "Source: " << source << "\n\nRaw AST:\n";
        print_ast(*expression);
        expression = fold_constants(std::move(expression));
        std::cout << "\nOptimized AST:\n";
        print_ast(*expression);
        return 0;
    } catch (const std::exception& error) {
        std::cerr << "MiniPTX error: " << error.what() << '\n';
        return 1;
    }
}

How the pieces now fit together

The lexer and parser still build the original tree. The new optimization pass rewrites that tree before the second printout exposes the result.

Rebuild and compare the trees

The two labels now give you a before-and-after view of the same expression. The final run checks that the inner addition disappears only from the optimized tree.

• Rebuild the updated executable by running this command:

cl /std:c++20 /EHsc /W4 /permissive- /Fe:mini_ptx.exe src\main.cpp

What does this command do?

The command compiles the updated C++20 source into mini_ptx.exe. It also checks the optimization code with warning level 4 and conformance mode enabled.

You should return to the command prompt without a compiler error. The executable is ready for the comparison run.

Before you run MiniPTX, picture the optimized tree and predict what replaces the inner addition.

• Run the rebuilt compiler with the original expression by entering this command:

.\mini_ptx.exe "x * (2 + 3)"

What should you see?

The Raw AST section contains a Binary(*) root with a nested Binary(+) node. This is the parser's untouched result.

The Optimized AST section keeps the Binary(*) root. Its numeric child is now Number(5).

That transformation is working. MiniPTX now removes arithmetic whose result is known before runtime.

Still seeing Binary(+) in both trees?

Confirm that the assignment to expression calls fold_constants() between the two print_ast() calls.

Rebuild mini_ptx.exe after saving src\main.cpp. Running an older executable leaves both trees unchanged.

Help me find why my optimized AST still contains Binary(+).

Your front end now parses expressions and performs a visible compiler optimization. Next, you will translate the optimized tree into PTX instructions.

Emit Body-Only PTX

Your optimizer now turns (2 + 3) into 5 before runtime. The next stage must translate that optimized tree into instructions a GPU toolchain can process.

You will emit a PTX kernel body that looks complete at first glance. You will then send it to NVIDIA's compiler to discover which larger contract the body does not satisfy.

In this step, get ready to:

• Translate optimized AST nodes into PTX arithmetic instructions.
• Write a body-only generated.ptx file containing a transform entry function.
• Send the body-only file through nvcc to expose the intended compilation failure.

Translate the optimized AST into PTX

The emitter walks the optimized tree recursively. Each numeric result receives a virtual floating-point register such as %f2.

The kernel body also calculates one thread's memory address. It loads that thread's input value before evaluating the expression.

• Switch back to src\main.cpp in the MiniPTX workspace from earlier.
• Add the PTX emitter code after fold_constants(ExprPtr expression) by matching the cumulative file in the comparison tab below.
• Update main by matching the same cumulative file.

✔️ Awesome, I've got everything!

Great. Confirm that your changes are saved in src\main.cpp.

ⓧ I'd like to double check the full code

#include <bit>
#include <cctype>
#include <cstdint>
#include <fstream>
#include <iomanip>
#include <iostream>
#include <memory>
#include <sstream>
#include <stdexcept>
#include <string>
#include <string_view>
#include <type_traits>
#include <utility>
#include <variant>
#include <vector>

enum class TokenKind {
    End,
    Identifier,
    Number,
    Plus,
    Star,
    LeftParen,
    RightParen
};

struct Token {
    TokenKind kind;
    std::string text;
    float number = 0.0f;
};

class Lexer {
public:
    explicit Lexer(std::string_view source) : source_(source) {}

    Token next() {
        skip_whitespace();
        if (position_ >= source_.size()) return {TokenKind::End, ""};

        const char current = source_[position_];
        switch (current) {
        case '+': ++position_; return {TokenKind::Plus, "+"};
        case '*': ++position_; return {TokenKind::Star, "*"};
        case '(': ++position_; return {TokenKind::LeftParen, "("};
        case ')': ++position_; return {TokenKind::RightParen, ")"};
        default: break;
        }

        if (std::isalpha(static_cast<unsigned char>(current)) || current == '_') {
            const std::size_t start = position_++;
            while (position_ < source_.size()) {
                const char next_character = source_[position_];
                if (!std::isalnum(static_cast<unsigned char>(next_character)) && next_character != '_') break;
                ++position_;
            }
            return {TokenKind::Identifier, std::string(source_.substr(start, position_ - start))};
        }

        if (std::isdigit(static_cast<unsigned char>(current)) || current == '.') {
            const std::size_t start = position_;
            bool seen_decimal_point = false;
            while (position_ < source_.size()) {
                const char next_character = source_[position_];
                if (next_character == '.') {
                    if (seen_decimal_point) break;
                    seen_decimal_point = true;
                    ++position_;
                    continue;
                }
                if (!std::isdigit(static_cast<unsigned char>(next_character))) break;
                ++position_;
            }
            const std::string text(source_.substr(start, position_ - start));
            return {TokenKind::Number, text, std::stof(text)};
        }

        throw std::runtime_error("Unexpected character: " + std::string(1, current));
    }

private:
    void skip_whitespace() {
        while (position_ < source_.size() && std::isspace(static_cast<unsigned char>(source_[position_]))) ++position_;
    }

    std::string_view source_;
    std::size_t position_ = 0;
};

struct Expr;
using ExprPtr = std::unique_ptr<Expr>;
struct Number { float value; };
struct Variable {};
struct Binary { char operation; ExprPtr left; ExprPtr right; };

struct Expr {
    using Node = std::variant<Number, Variable, Binary>;
    explicit Expr(Number number) : node(number) {}
    explicit Expr(Variable variable) : node(variable) {}
    explicit Expr(Binary binary) : node(std::move(binary)) {}
    Node node;
};

ExprPtr make_number(float value) { return std::make_unique<Expr>(Number{value}); }
ExprPtr make_variable() { return std::make_unique<Expr>(Variable{}); }
ExprPtr make_binary(char operation, ExprPtr left, ExprPtr right) {
    return std::make_unique<Expr>(Binary{operation, std::move(left), std::move(right)});
}

class Parser {
public:
    explicit Parser(std::string_view source) : lexer_(source), current_(lexer_.next()) {}

    ExprPtr parse() {
        ExprPtr expression = parse_addition();
        if (current_.kind != TokenKind::End) throw std::runtime_error("Unexpected token: " + current_.text);
        return expression;
    }

private:
    ExprPtr parse_addition() {
        ExprPtr left = parse_multiplication();
        while (current_.kind == TokenKind::Plus) {
            advance();
            left = make_binary('+', std::move(left), parse_multiplication());
        }
        return left;
    }

    ExprPtr parse_multiplication() {
        ExprPtr left = parse_primary();
        while (current_.kind == TokenKind::Star) {
            advance();
            left = make_binary('*', std::move(left), parse_primary());
        }
        return left;
    }

    ExprPtr parse_primary() {
        if (current_.kind == TokenKind::Number) {
            const float value = current_.number;
            advance();
            return make_number(value);
        }
        if (current_.kind == TokenKind::Identifier) {
            if (current_.text != "x") throw std::runtime_error("Only the variable 'x' is supported, received: " + current_.text);
            advance();
            return make_variable();
        }
        if (current_.kind == TokenKind::LeftParen) {
            advance();
            ExprPtr expression = parse_addition();
            if (current_.kind != TokenKind::RightParen) throw std::runtime_error("Expected ')'");
            advance();
            return expression;
        }
        throw std::runtime_error("Expected a number, 'x', or '('");
    }

    void advance() { current_ = lexer_.next(); }
    Lexer lexer_;
    Token current_;
};

void print_ast(const Expr& expression, int indentation = 0) {
    const std::string padding(static_cast<std::size_t>(indentation), ' ');
    std::visit([&](const auto& node) {
        using NodeType = std::decay_t<decltype(node)>;
        if constexpr (std::is_same_v<NodeType, Number>) {
            std::cout << padding << "Number(" << node.value << ")\n";
        } else if constexpr (std::is_same_v<NodeType, Variable>) {
            std::cout << padding << "Variable(x)\n";
        } else {
            std::cout << padding << "Binary(" << node.operation << ")\n";
            print_ast(*node.left, indentation + 2);
            print_ast(*node.right, indentation + 2);
        }
    }, expression.node);
}

ExprPtr fold_constants(ExprPtr expression) {
    auto* binary = std::get_if<Binary>(&expression->node);
    if (binary == nullptr) return expression;
    binary->left = fold_constants(std::move(binary->left));
    binary->right = fold_constants(std::move(binary->right));
    const auto* left_number = std::get_if<Number>(&binary->left->node);
    const auto* right_number = std::get_if<Number>(&binary->right->node);
    if (left_number != nullptr && right_number != nullptr) {
        if (binary->operation == '+') return make_number(left_number->value + right_number->value);
        if (binary->operation == '*') return make_number(left_number->value * right_number->value);
    }
    return expression;
}

class PtxEmitter {
public:
    std::string emit(const Expr& expression, bool body_only) {
        instructions_.clear();
        next_float_register_ = 2;
        const std::string result_register = emit_expression(expression);
        std::ostringstream body;
        body << ".visible .entry transform(\n"
             << "    .param .u64 transform_param_0,\n"
             << "    .param .u64 transform_param_1\n"
             << ") {\n"
             << "    .reg .f32 %f<" << next_float_register_ << ">;\n"
             << "    .reg .b32 %r<5>;\n"
             << "    .reg .b64 %rd<8>;\n\n"
             << "    ld.param.u64 %rd1, [transform_param_0];\n"
             << "    ld.param.u64 %rd2, [transform_param_1];\n"
             << "    cvta.to.global.u64 %rd3, %rd2;\n"
             << "    cvta.to.global.u64 %rd4, %rd1;\n"
             << "    mov.u32 %r1, %ctaid.x;\n"
             << "    mov.u32 %r2, %ntid.x;\n"
             << "    mov.u32 %r3, %tid.x;\n"
             << "    mad.lo.s32 %r4, %r2, %r1, %r3;\n"
             << "    mul.wide.u32 %rd5, %r4, 4;\n"
             << "    add.s64 %rd6, %rd4, %rd5;\n"
             << "    ld.global.f32 %f1, [%rd6];\n";
        for (const std::string& instruction : instructions_) body << "    " << instruction << '\n';
        body << "    add.s64 %rd7, %rd3, %rd5;\n"
             << "    st.global.f32 [%rd7], " << result_register << ";\n"
             << "    ret;\n"
             << "}\n";
        if (body_only) return body.str();
        return std::string(".version 7.0\n.target sm_50\n.address_size 64\n\n") + body.str();
    }

private:
    std::string emit_expression(const Expr& expression) {
        if (const auto* number = std::get_if<Number>(&expression.node)) {
            const std::string destination = allocate_float_register();
            instructions_.push_back("mov.f32 " + destination + ", " + ptx_immediate(number->value) + ";");
            return destination;
        }
        if (std::holds_alternative<Variable>(expression.node)) return "%f1";
        const auto& binary = std::get<Binary>(expression.node);
        const std::string left = emit_expression(*binary.left);
        const std::string right = emit_expression(*binary.right);
        const std::string destination = allocate_float_register();
        const std::string operation = binary.operation == '+' ? "add.f32" : "mul.f32";
        instructions_.push_back(operation + " " + destination + ", " + left + ", " + right + ";");
        return destination;
    }

    std::string allocate_float_register() { return "%f" + std::to_string(next_float_register_++); }
    static std::string ptx_immediate(float value) {
        const std::uint32_t bits = std::bit_cast<std::uint32_t>(value);
        std::ostringstream output;
        output << "0f" << std::hex << std::setw(8) << std::setfill('0') << bits;
        return output.str();
    }

    int next_float_register_ = 2;
    std::vector<std::string> instructions_;
};

int main(int argument_count, char* arguments[]) {
    try {
        std::string source = "x * (2 + 3)";
        bool body_only = false;
        for (int index = 1; index < argument_count; ++index) {
            const std::string argument = arguments[index];
            if (argument == "--body-only") body_only = true;
            else source = argument;
        }
        Parser parser(source);
        ExprPtr expression = parser.parse();
        std::cout << "Source: " << source << "\n\nRaw AST:\n";
        print_ast(*expression);
        expression = fold_constants(std::move(expression));
        std::cout << "\nOptimized AST:\n";
        print_ast(*expression);
        PtxEmitter emitter;
        const std::string ptx = emitter.emit(*expression, body_only);
        std::ofstream output("generated.ptx", std::ios::binary);
        if (!output) throw std::runtime_error("Could not open generated.ptx for writing");
        output << ptx;
        std::cout << "\nWrote generated.ptx (" << (body_only ? "body only" : "complete module") << ").\n";
        return 0;
    } catch (const std::exception& error) {
        std::cerr << "MiniPTX error: " << error.what() << '\n';
        return 1;
    }
}

What does this code do?

• The emit_expression method translates each AST node into a result register.
• A Number node emits mov.f32 with an exact hexadecimal floating-point value.
• A Variable node reuses %f1 because the kernel loads the input element there.
• A Binary node recursively emits its children before adding add.f32 or mul.f32.
• The emit method wraps the arithmetic in a transform entry function with parameter loads and per-thread address calculations.

Generate the body-only kernel

The --body-only option selects the deliberate test path. MiniPTX still parses and optimizes the expression before it writes the kernel body.

• Save src\main.cpp.
• Rebuild mini_ptx.exe from the x64 Native Tools Command Prompt by running:

cl /std:c++20 /EHsc /W4 /permissive- /Fe:mini_ptx.exe src\main.cpp

What does this command do?

The command rebuilds MiniPTX with C++20 enabled. It also enables standard exception unwinding and warning level 4.

The conformance option tightens Microsoft-specific language behavior. The output option writes the rebuilt program to mini_ptx.exe.

The host compiler should finish successfully. Your rebuilt executable can now turn the optimized AST into PTX text.

Does the rebuild fail?

Check the first compiler diagnostic against the matching line in the full-file comparison. A missing brace near PtxEmitter can cause later lines to look incorrect too.

Confirm that the new standard library headers appear at the top of src\main.cpp.

Help me diagnose the first C++ compiler error in my MiniPTX PtxEmitter code.

Before you run MiniPTX, which optimized node do you expect the emitter to translate for the parenthesized addition?

• Generate the body-only PTX file by running:

.\mini_ptx.exe "x * (2 + 3)" --body-only

What should I see?

MiniPTX prints the raw AST first. It then prints the optimized AST with Number(5) beneath Binary(*).

The final status reports that MiniPTX wrote generated.ptx as body-only output. That confirms the emitter received the optimized tree.

• Select generated.ptx from the MiniPTX file list in VS Code.
• Confirm that the file starts with .visible .entry transform(.
• Find the mov.f32 instruction that materializes the folded constant.
• Find the mul.f32 instruction that multiplies %f1 by that constant.
• Confirm that the file ends with ret inside the entry function.

You can now inspect a complete thread-level calculation from parameter loads through the output store. The AST transformation has become GPU-oriented instruction text.

Does generated.ptx look different?

Confirm that --body-only appears after the quoted expression in the command. The option controls which output path MiniPTX selects.

If generated.ptx is missing, check whether the terminal reports that the file could not be opened for writing.

Help me compare my generated PTX with the expected body-only transform entry.

Test the body with NVIDIA's compiler

A kernel body describes registers and instructions. The NVIDIA compilation phase validates whether the surrounding PTX input provides everything required to build a CUBIN.

Before you run the compilation command, do you think a plausible kernel body gives the toolchain enough information to create the binary?

• Send generated.ptx through the Ada compilation path by running:

nvcc --cubin generated.ptx --gpu-architecture=sm_89 --resource-usage --output-file generated.cubin

What just happened?

• The --cubin phase asks nvcc to assemble PTX into a device-only binary.
• The --gpu-architecture=sm_89 option selects the Ada architecture used by the GeForce RTX 4050 Laptop GPU.
• The --resource-usage option requests register and GPU memory usage if assembly succeeds.
• The body-only input fails PTX compilation before a usable generated.cubin can be produced.

That failure is the result you wanted. Your emitter produced convincing kernel logic while the assembler exposed the missing module-level contract.

• Return to the MiniPTX file list in VS Code.
• Confirm that no usable generated.cubin exists after the failed compilation.

Did you miss the intended PTX failure?

Confirm that you ran MiniPTX with --body-only immediately before the NVIDIA compilation command. An older generated.ptx file can produce a different result.

If the NVIDIA command itself is unavailable, return to the x64 Native Tools Command Prompt where you verified the CUDA Toolkit earlier.

Help me reproduce the intended body-only PTX compilation failure.

You have reached the boundary between valid-looking instructions and a valid PTX module. Next up, you will satisfy that module contract and turn MiniPTX output into a real Ada binary.

Build a Complete PTX Module

Your body-only PTX reached NVIDIA's assembler in the previous step. The failure proved that a valid kernel body still needs a module contract.

You will add that contract without changing the kernel body. The finished module becomes a real CUBIN for your Ada GPU.

In this step, get ready to:

• Add the required module directives to MiniPTX's complete-module output.
• Regenerate PTX with folded arithmetic visible in the kernel.
• Assemble the PTX into an sm_89 GPU binary with a resource report.

Add the PTX module directives

A complete PTX module begins with version information followed by its target. MiniPTX can prepend those directives while preserving the body-only mode that exposed the failure.

• Return to src\main.cpp in Visual Studio Code.
• Find this line near the end of PtxEmitter::emit():

if (body_only) return body.str();

What does this branch preserve?

The condition returns only the entry body when body_only is true. This keeps the designed failure available for comparison.

• Add this complete-module return statement directly below that line:

return std::string(".version 7.0\n.target sm_50\n.address_size 64\n\n") + body.str();

What completes the module?

• The .version 7.0 directive declares the PTX language version used by the module.
• The .target sm_50 directive sets the module's baseline target.
• The .address_size 64 directive establishes 64-bit addressing for the pointer parameters.
• The existing body.str() value places the unchanged transform entry after the header.

• Save src\main.cpp.
• Rebuild mini_ptx.exe by running this command in the x64 Native Tools Command Prompt:

cl /std:c++20 /EHsc /W4 /permissive- /Fe:mini_ptx.exe src\main.cpp

What does this command check?

The command rebuilds MiniPTX as a C++20 program with conformance checks plus warning level 4. A successful build confirms that the new return path compiles.

Good progress. Your compiler can now choose between the deliberately incomplete body and a complete PTX module.

Does the rebuild fail?

Check that the new return statement sits after the body_only condition inside PtxEmitter::emit(). Confirm that each newline inside the C++ string uses \n.

Compare the quotes plus parentheses against the reference below. Help me diagnose the C++ compiler output from this PtxEmitter return statement.

✔️ Awesome, I've got everything!

Your saved src\main.cpp now emits the module header unless body_only is true.

ⓧ I'd like to double check the full code

#include <bit>
#include <cctype>
#include <cstdint>
#include <fstream>
#include <iomanip>
#include <iostream>
#include <memory>
#include <sstream>
#include <stdexcept>
#include <string>
#include <string_view>
#include <type_traits>
#include <utility>
#include <variant>
#include <vector>

enum class TokenKind {
    End,
    Identifier,
    Number,
    Plus,
    Star,
    LeftParen,
    RightParen
};

struct Token {
    TokenKind kind;
    std::string text;
    float number = 0.0f;
};

class Lexer {
public:
    explicit Lexer(std::string_view source) : source_(source) {}

    Token next() {
        skip_whitespace();
        if (position_ >= source_.size()) return {TokenKind::End, ""};

        const char current = source_[position_];
        switch (current) {
        case '+': ++position_; return {TokenKind::Plus, "+"};
        case '*': ++position_; return {TokenKind::Star, "*"};
        case '(': ++position_; return {TokenKind::LeftParen, "("};
        case ')': ++position_; return {TokenKind::RightParen, ")"};
        default: break;
        }

        if (std::isalpha(static_cast<unsigned char>(current)) || current == '_') {
            const std::size_t start = position_++;
            while (position_ < source_.size()) {
                const char next_character = source_[position_];
                if (!std::isalnum(static_cast<unsigned char>(next_character)) && next_character != '_') break;
                ++position_;
            }
            return {TokenKind::Identifier, std::string(source_.substr(start, position_ - start))};
        }

        if (std::isdigit(static_cast<unsigned char>(current)) || current == '.') {
            const std::size_t start = position_;
            bool seen_decimal_point = false;
            while (position_ < source_.size()) {
                const char next_character = source_[position_];
                if (next_character == '.') {
                    if (seen_decimal_point) break;
                    seen_decimal_point = true;
                    ++position_;
                    continue;
                }
                if (!std::isdigit(static_cast<unsigned char>(next_character))) break;
                ++position_;
            }
            const std::string text(source_.substr(start, position_ - start));
            return {TokenKind::Number, text, std::stof(text)};
        }

        throw std::runtime_error("Unexpected character: " + std::string(1, current));
    }

private:
    void skip_whitespace() {
        while (position_ < source_.size() && std::isspace(static_cast<unsigned char>(source_[position_]))) ++position_;
    }

    std::string_view source_;
    std::size_t position_ = 0;
};

struct Expr;
using ExprPtr = std::unique_ptr<Expr>;
struct Number { float value; };
struct Variable {};
struct Binary { char operation; ExprPtr left; ExprPtr right; };

struct Expr {
    using Node = std::variant<Number, Variable, Binary>;
    explicit Expr(Number number) : node(number) {}
    explicit Expr(Variable variable) : node(variable) {}
    explicit Expr(Binary binary) : node(std::move(binary)) {}
    Node node;
};

ExprPtr make_number(float value) { return std::make_unique<Expr>(Number{value}); }
ExprPtr make_variable() { return std::make_unique<Expr>(Variable{}); }
ExprPtr make_binary(char operation, ExprPtr left, ExprPtr right) {
    return std::make_unique<Expr>(Binary{operation, std::move(left), std::move(right)});
}

class Parser {
public:
    explicit Parser(std::string_view source) : lexer_(source), current_(lexer_.next()) {}

    ExprPtr parse() {
        ExprPtr expression = parse_addition();
        if (current_.kind != TokenKind::End) throw std::runtime_error("Unexpected token: " + current_.text);
        return expression;
    }

private:
    ExprPtr parse_addition() {
        ExprPtr left = parse_multiplication();
        while (current_.kind == TokenKind::Plus) {
            advance();
            left = make_binary('+', std::move(left), parse_multiplication());
        }
        return left;
    }

    ExprPtr parse_multiplication() {
        ExprPtr left = parse_primary();
        while (current_.kind == TokenKind::Star) {
            advance();
            left = make_binary('*', std::move(left), parse_primary());
        }
        return left;
    }

    ExprPtr parse_primary() {
        if (current_.kind == TokenKind::Number) {
            const float value = current_.number;
            advance();
            return make_number(value);
        }
        if (current_.kind == TokenKind::Identifier) {
            if (current_.text != "x") throw std::runtime_error("Only the variable 'x' is supported, received: " + current_.text);
            advance();
            return make_variable();
        }
        if (current_.kind == TokenKind::LeftParen) {
            advance();
            ExprPtr expression = parse_addition();
            if (current_.kind != TokenKind::RightParen) throw std::runtime_error("Expected ')'");
            advance();
            return expression;
        }
        throw std::runtime_error("Expected a number, 'x', or '('");
    }

    void advance() { current_ = lexer_.next(); }
    Lexer lexer_;
    Token current_;
};

void print_ast(const Expr& expression, int indentation = 0) {
    const std::string padding(static_cast<std::size_t>(indentation), ' ');
    std::visit([&](const auto& node) {
        using NodeType = std::decay_t<decltype(node)>;
        if constexpr (std::is_same_v<NodeType, Number>) {
            std::cout << padding << "Number(" << node.value << ")\n";
        } else if constexpr (std::is_same_v<NodeType, Variable>) {
            std::cout << padding << "Variable(x)\n";
        } else {
            std::cout << padding << "Binary(" << node.operation << ")\n";
            print_ast(*node.left, indentation + 2);
            print_ast(*node.right, indentation + 2);
        }
    }, expression.node);
}

ExprPtr fold_constants(ExprPtr expression) {
    auto* binary = std::get_if<Binary>(&expression->node);
    if (binary == nullptr) return expression;
    binary->left = fold_constants(std::move(binary->left));
    binary->right = fold_constants(std::move(binary->right));
    const auto* left_number = std::get_if<Number>(&binary->left->node);
    const auto* right_number = std::get_if<Number>(&binary->right->node);
    if (left_number != nullptr && right_number != nullptr) {
        if (binary->operation == '+') return make_number(left_number->value + right_number->value);
        if (binary->operation == '*') return make_number(left_number->value * right_number->value);
    }
    return expression;
}

class PtxEmitter {
public:
    std::string emit(const Expr& expression, bool body_only) {
        instructions_.clear();
        next_float_register_ = 2;
        const std::string result_register = emit_expression(expression);
        std::ostringstream body;
        body << ".visible .entry transform(\n"
             << "    .param .u64 transform_param_0,\n"
             << "    .param .u64 transform_param_1\n"
             << ") {\n"
             << "    .reg .f32 %f<" << next_float_register_ << ">;\n"
             << "    .reg .b32 %r<5>;\n"
             << "    .reg .b64 %rd<8>;\n\n"
             << "    ld.param.u64 %rd1, [transform_param_0];\n"
             << "    ld.param.u64 %rd2, [transform_param_1];\n"
             << "    cvta.to.global.u64 %rd3, %rd2;\n"
             << "    cvta.to.global.u64 %rd4, %rd1;\n"
             << "    mov.u32 %r1, %ctaid.x;\n"
             << "    mov.u32 %r2, %ntid.x;\n"
             << "    mov.u32 %r3, %tid.x;\n"
             << "    mad.lo.s32 %r4, %r2, %r1, %r3;\n"
             << "    mul.wide.u32 %rd5, %r4, 4;\n"
             << "    add.s64 %rd6, %rd4, %rd5;\n"
             << "    ld.global.f32 %f1, [%rd6];\n";
        for (const std::string& instruction : instructions_) body << "    " << instruction << '\n';
        body << "    add.s64 %rd7, %rd3, %rd5;\n"
             << "    st.global.f32 [%rd7], " << result_register << ";\n"
             << "    ret;\n"
             << "}\n";
        if (body_only) return body.str();
        return std::string(".version 7.0\n.target sm_50\n.address_size 64\n\n") + body.str();
    }

private:
    std::string emit_expression(const Expr& expression) {
        if (const auto* number = std::get_if<Number>(&expression.node)) {
            const std::string destination = allocate_float_register();
            instructions_.push_back("mov.f32 " + destination + ", " + ptx_immediate(number->value) + ";");
            return destination;
        }
        if (std::holds_alternative<Variable>(expression.node)) return "%f1";
        const auto& binary = std::get<Binary>(expression.node);
        const std::string left = emit_expression(*binary.left);
        const std::string right = emit_expression(*binary.right);
        const std::string destination = allocate_float_register();
        const std::string operation = binary.operation == '+' ? "add.f32" : "mul.f32";
        instructions_.push_back(operation + " " + destination + ", " + left + ", " + right + ";");
        return destination;
    }

    std::string allocate_float_register() { return "%f" + std::to_string(next_float_register_++); }
    static std::string ptx_immediate(float value) {
        const std::uint32_t bits = std::bit_cast<std::uint32_t>(value);
        std::ostringstream output;
        output << "0f" << std::hex << std::setw(8) << std::setfill('0') << bits;
        return output.str();
    }

    int next_float_register_ = 2;
    std::vector<std::string> instructions_;
};

int main(int argument_count, char* arguments[]) {
    try {
        std::string source = "x * (2 + 3)";
        bool body_only = false;
        for (int index = 1; index < argument_count; ++index) {
            const std::string argument = arguments[index];
            if (argument == "--body-only") body_only = true;
            else source = argument;
        }
        Parser parser(source);
        ExprPtr expression = parser.parse();
        std::cout << "Source: " << source << "\n\nRaw AST:\n";
        print_ast(*expression);
        expression = fold_constants(std::move(expression));
        std::cout << "\nOptimized AST:\n";
        print_ast(*expression);
        PtxEmitter emitter;
        const std::string ptx = emitter.emit(*expression, body_only);
        std::ofstream output("generated.ptx", std::ios::binary);
        if (!output) throw std::runtime_error("Could not open generated.ptx for writing");
        output << ptx;
        std::cout << "\nWrote generated.ptx (" << (body_only ? "body only" : "complete module") << ").\n";
        return 0;
    } catch (const std::exception& error) {
        std::cerr << "MiniPTX error: " << error.what() << '\n';
        return 1;
    }
}

What does the full file confirm?

The reference combines the lexer, parser, recursive AST, constant folder, PTX emitter, module modes, plus file output in one source file. Compare it with your saved src\main.cpp before continuing.

Regenerate and inspect the PTX

The optimized AST already replaces 2 + 3 with 5. Regenerating without the body-only option lets you verify that both the module header and folded arithmetic reach the emitted file.

• Regenerate generated.ptx in complete-module mode by running:

.\mini_ptx.exe "x * (2 + 3)"

What should the compiler report?

MiniPTX prints the raw AST followed by the optimized AST. It finishes by reporting that it wrote generated.ptx as a complete module.

• Return to the existing generated.ptx file in Visual Studio Code.
• Confirm that its first three directives are .version 7.0, .target sm_50, plus .address_size 64.
• Locate the mov.f32 instruction that loads the folded constant.
• Locate the mul.f32 instruction that multiplies the input by that constant.
• Confirm that the optimized expression has no add.f32 instruction for 2 + 3.

You now have visible evidence of two compiler stages in one file. The header satisfies the module contract while the shortened instruction sequence proves that constant folding happened before emission.

Still seeing body-only PTX?

Confirm that you ran MiniPTX without --body-only. Check that Visual Studio Code is showing the regenerated generated.ptx inside the MiniPTX folder.

If the header is still missing, rebuild mini_ptx.exe from the saved source. Help me find why MiniPTX still emits body-only PTX.

Compile the PTX into a CUBIN

The CUDA toolchain can now assemble your complete virtual instruction module for the Ada architecture. The resource report also exposes the register plus memory requirements chosen by NVIDIA's backend.

Before you run this, do you expect the previous module-format failure or a successful assembly?

• Compile generated.ptx for sm_89 by running:

nvcc --cubin generated.ptx --gpu-architecture=sm_89 --resource-usage --output-file generated.cubin

What does this command produce?

• The --cubin phase assembles the PTX input into a device-only binary.
• The --gpu-architecture=sm_89 option targets the Ada architecture used by your GeForce RTX 4050 Laptop GPU.
• The --resource-usage option displays register plus GPU memory usage.
• The --output-file generated.cubin option names the binary artifact.

The command now completes without the earlier PTX compilation failure. You will see GPU resource information in the terminal.

• Return to the Visual Studio Code Explorer sidebar.
• Refresh the MiniPTX folder from the Explorer toolbar.
• Confirm that generated.cubin is listed beside generated.ptx.

That completes your source-to-binary compiler pipeline. MiniPTX now turns an expression into optimized PTX that NVIDIA's backend accepts for sm_89.

Does the CUBIN build still fail?

Confirm that generated.ptx begins with the three module directives. Regenerate it without the --body-only option if the file still begins with the entry function.

Confirm that you are still using the x64 Native Tools Command Prompt where nvcc worked earlier. Help me diagnose why nvcc cannot assemble my complete generated.ptx file.

Secret mission

Trace Your PTX into Ada SASS

Your compiler has produced a CUBIN, but the GPU executes architecture-specific machine instructions. In this mission, you will disassemble generated.cubin to trace the transform data path through Ada SASS. You will explain how NVIDIA's backend reshapes PTX during lowering.

Clean Up Your Resources

Clean Up Your Resources

Choose whether to keep your MiniPTX workspace ready, close it for later, or remove it. This project runs locally, so its files create no ongoing service charges.

Resources you used:

• Your local MiniPTX workspace containing src\main.cpp, mini_ptx.exe, generated.ptx, and generated.cubin.

Keep everything running

No action is needed. Choose this option if you want to keep experimenting with expressions or extending the compiler.

• Keep the MiniPTX folder in its current location.
• Keep src\main.cpp as your editable compiler source.
• Keep generated.ptx as the readable output from your compiler.
• Keep generated.cubin as the assembled GPU binary you inspected.
• Keep the workspace available in Visual Studio Code for your next compiler experiment.

Pause - I'll come back to this later

Close your editing session while preserving every project file. Your source and generated artifacts remain ready for your return.

• Save any open changes in src\main.cpp.
• Close Visual Studio Code.
• Close the x64 Native Tools Command Prompt window.
• Leave the MiniPTX folder in place.

Delete - I don't want to use this again

This option permanently removes the local compiler workspace. Use it when you no longer need the source or generated artifacts.

Deletion Removes the Whole Workspace

Removing a complete compiler project can feel drastic. The commands below target only the local MiniPTX folder from this project.

• Close Visual Studio Code so Windows releases the project files.
• Switch back to the x64 Native Tools Command Prompt from earlier.
• Move to the parent folder and remove MiniPTX by running these commands:

cd ..
rmdir /s /q MiniPTX

What Do These Commands Do?

• The first line moves the prompt from MiniPTX to its parent folder.
• The second line removes the MiniPTX folder with every file inside it.

• Check that the workspace is gone by running this command:

dir MiniPTX

What Should You See?

Command Prompt checks for a folder named MiniPTX. You should see a message indicating that the path is unavailable, which confirms the local workspace is gone.

Folder Still Present?

• Confirm that Visual Studio Code is closed.
• Close any File Explorer window showing the MiniPTX folder.
• Retry the deletion commands from the parent folder.

Help me safely remove my MiniPTX folder when Windows says it is still in use.

Nice Work!

Nice Work!

You did it! MiniPTX now carries an expression through a complete compiler pipeline into an sm_89 GPU binary.

You've learned how to:

• Build a C++20 lexer plus a precedence-aware recursive-descent parser. Your compiler turns x * (2 + 3) into a printable AST backed by std::variant plus std::unique_ptr.
• Implement constant folding to replace compile-time arithmetic with a single value. Emit the optimized expression as a register-based PTX kernel.
• Reproduce a genuine PTX module-format failure from --body-only output. Fix the module with its required directives. Compile the result into an sm_89 CUBIN with a resource-usage report.
• Secret Mission: Disassemble generated.cubin into Ada SASS. Identify the machine instructions behind the thread index calculation. Locate the global load, floating-point arithmetic, global store, plus function exit. Explain how NVIDIA's backend lowers PTX operations into architecture-specific instructions.

Ready to quiz yourself?