CodeCraft_cuda_max_find.ipynb

{
  "nbformat": 4,
  "nbformat_minor": 0,
  "metadata": {
    "colab": {
      "provenance": [],
      "gpuType": "T4",
      "authorship_tag": "ABX9TyPsFdXji5Zf17RKnXDbFXy5",
      "include_colab_link": true
    },
    "kernelspec": {
      "name": "python3",
      "display_name": "Python 3"
    },
    "language_info": {
      "name": "python"
    },
    "accelerator": "GPU"
  },
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "view-in-github",
        "colab_type": "text"
      },
      "source": [
        "<a href=\"https://colab.research.google.com/github/walkerjian/DailyCode/blob/main/CodeCraft_cuda_max_find.ipynb\" target=\"_parent\"><img src=\"https://colab.research.google.com/assets/colab-badge.svg\" alt=\"Open In Colab\"/></a>"
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "#Code Craft Solution:"
      ],
      "metadata": {
        "id": "EwC8UYjQozrN"
      }
    },
    {
      "cell_type": "markdown",
      "source": [
        "##Prompt:\n",
        "Solve this coding problem professionally using modern software development patterns, particularly the MVC (Model-View-Controller) pattern for clarity and maintainability. Ensure that the test harness is well-thought-out and professional, with logic that ensures all tests run, even when some fail. If a test fails, the harness should also test the test logic to ensure it's not the source of failure. This should be treated as a learning exercise, so the solution must include detailed explanations (didactic perspective) at each step, including complexity analysis, optimizations, and alternative approaches. If any tests or code fail, include a post-mortem analysis of the failure. Use languages other than Python where appropriate, and ensure readability and understanding of the solution.\n",
        "\n",
        "###Problem:\n",
        "This problem was asked by Nvidia.\n",
        "\n",
        "Find the maximum of two numbers without using any if-else statements, branching, or direct comparisons."
      ],
      "metadata": {
        "id": "8scBDLARodKX"
      }
    },
    {
      "cell_type": "markdown",
      "source": [
        "##Solution\n",
        "This notebook implements a solution to the problem of finding the maximum of two numbers without using conditional statements, branching, or direct comparisons.\n",
        "\n",
        "The solution is implemented in Java using the MVC pattern (Model-View-Controller), with a test harness to ensure its correctness.\n",
        "\n",
        "The code is then implemented in CUDA C and Assembly and tested in a T4 CUDA environment on Google Colab.\n",
        "\n",
        "```\n",
        "# This is formatted as code\n",
        "```\n",
        "\n",
        ":\n",
        "```java\n",
        "/**\n",
        " * MaxFinder.java\n",
        " *\n",
        " * Model class containing the logic to find the maximum of two integers\n",
        " * without using any if-else statements, branching, or direct comparisons.\n",
        " *\n",
        " * This implementation uses bit manipulation and arithmetic operations\n",
        " * to determine the maximum value. It handles integer overflows by\n",
        " * using long integers for intermediate calculations.\n",
        " */\n",
        "public class MaxFinder {\n",
        "    /**\n",
        "     * Finds the maximum of two integers without using if-else, branching, or direct comparisons.\n",
        "     *\n",
        "     * @param a First integer\n",
        "     * @param b Second integer\n",
        "     * @return The maximum of a and b\n",
        "     */\n",
        "    public int max(int a, int b) {\n",
        "        // Convert integers to long to prevent integer overflow\n",
        "        long c = (long) a - (long) b;\n",
        "        // Extract the sign bit: 1 if c < 0, 0 otherwise\n",
        "        int k = (int) ((c >> 63) & 1);\n",
        "        // Compute the maximum value based on the sign of c\n",
        "        return a - k * (int) c;\n",
        "    }\n",
        "}\n",
        "\n",
        "/**\n",
        " * MaxView.java\n",
        " *\n",
        " * View class responsible for displaying the results.\n",
        " */\n",
        "public class MaxView {\n",
        "    /**\n",
        "     * Displays the result of the max computation.\n",
        "     *\n",
        "     * @param a   First integer\n",
        "     * @param b   Second integer\n",
        "     * @param max The maximum of a and b\n",
        "     */\n",
        "    public void displayResult(int a, int b, int max) {\n",
        "        System.out.println(\"The maximum of \" + a + \" and \" + b + \" is \" + max + \".\");\n",
        "    }\n",
        "}\n",
        "\n",
        "/**\n",
        " * MaxController.java\n",
        " *\n",
        " * Controller class that coordinates between the Model and the View.\n",
        " */\n",
        "public class MaxController {\n",
        "    private MaxFinder model;\n",
        "    private MaxView view;\n",
        "\n",
        "    public MaxController(MaxFinder model, MaxView view) {\n",
        "        this.model = model;\n",
        "        this.view = view;\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Runs the test cases and displays results.\n",
        "     */\n",
        "    public void runTests() {\n",
        "        MaxTestHarness testHarness = new MaxTestHarness(model, view);\n",
        "        testHarness.runAllTests();\n",
        "        testHarness.reportResults();\n",
        "    }\n",
        "}\n",
        "\n",
        "/**\n",
        " * MaxTestHarness.java\n",
        " *\n",
        " * Test harness for the MaxFinder class.\n",
        " * Ensures that all tests run even if some fail.\n",
        " * If a test fails, it also checks the test logic.\n",
        " */\n",
        "import java.util.ArrayList;\n",
        "import java.util.List;\n",
        "\n",
        "public class MaxTestHarness {\n",
        "    private MaxFinder model;\n",
        "    private MaxView view;\n",
        "    private List<TestResult> testResults;\n",
        "\n",
        "    public MaxTestHarness(MaxFinder model, MaxView view) {\n",
        "        this.model = model;\n",
        "        this.view = view;\n",
        "        this.testResults = new ArrayList<>();\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Runs all the test cases.\n",
        "     */\n",
        "    public void runAllTests() {\n",
        "        testMax(5, 10);\n",
        "        testMax(-5, -10);\n",
        "        testMax(0, 0);\n",
        "        testMax(Integer.MAX_VALUE, Integer.MIN_VALUE);\n",
        "        testMax(-1, 1);\n",
        "        testMax(123456, 654321);\n",
        "        // Add more test cases as needed\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Reports the results of the tests.\n",
        "     */\n",
        "    public void reportResults() {\n",
        "        int passed = 0;\n",
        "        for (TestResult result : testResults) {\n",
        "            if (result.passed) {\n",
        "                passed++;\n",
        "            } else {\n",
        "                System.out.println(\"Test Failed: \" + result.message);\n",
        "            }\n",
        "        }\n",
        "        System.out.println(passed + \" out of \" + testResults.size() + \" tests passed.\");\n",
        "    }\n",
        "\n",
        "    private void testMax(int a, int b) {\n",
        "        int expected = Math.max(a, b);\n",
        "        int actual = model.max(a, b);\n",
        "        boolean passed = expected == actual;\n",
        "        String message = \"max(\" + a + \", \" + b + \") = \" + actual + \", expected: \" + expected;\n",
        "\n",
        "        if (!passed) {\n",
        "            // Test the test logic\n",
        "            boolean testLogicValid = verifyTestLogic(a, b, expected);\n",
        "            if (!testLogicValid) {\n",
        "                message += \" [Test Logic Failure]\";\n",
        "            }\n",
        "        }\n",
        "        testResults.add(new TestResult(passed, message));\n",
        "\n",
        "        // Display the result using the View\n",
        "        view.displayResult(a, b, actual);\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Verifies the test logic to ensure it's not the source of failure.\n",
        "     *\n",
        "     * @param a        First integer\n",
        "     * @param b        Second integer\n",
        "     * @param expected Expected result from the test\n",
        "     * @return True if the test logic is valid, false otherwise\n",
        "     */\n",
        "    private boolean verifyTestLogic(int a, int b, int expected) {\n",
        "        // Recompute the expected value using an alternative method\n",
        "        int alternativeExpected = (a + b + Math.abs(a - b)) / 2;\n",
        "        return expected == alternativeExpected;\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Inner class to store test results.\n",
        "     */\n",
        "    private class TestResult {\n",
        "        boolean passed;\n",
        "        String message;\n",
        "\n",
        "        TestResult(boolean passed, String message) {\n",
        "            this.passed = passed;\n",
        "            this.message = message;\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "/**\n",
        " * MaxMain.java\n",
        " *\n",
        " * Main class to run the application.\n",
        " */\n",
        "public class MaxMain {\n",
        "    public static void main(String[] args) {\n",
        "        MaxFinder model = new MaxFinder();\n",
        "        MaxView view = new MaxView();\n",
        "        MaxController controller = new MaxController(model, view);\n",
        "\n",
        "        controller.runTests();\n",
        "    }\n",
        "}\n",
        "```\n",
        "\n",
        "**Detailed Explanation:**\n",
        "\n",
        "### Model (MaxFinder.java)\n",
        "\n",
        "- **Purpose:** Contains the core logic to find the maximum of two integers without using any if-else statements, branching, or direct comparisons.\n",
        "- **Implementation Details:**\n",
        "  - **Bit Manipulation:** Utilizes bitwise operations to extract the sign bit of the difference between the two numbers.\n",
        "  - **Handling Overflows:** Converts integers to long to prevent integer overflow during subtraction.\n",
        "\n",
        "**Algorithm Explanation:**\n",
        "\n",
        "1. **Compute the Difference:**\n",
        "   - `long c = (long) a - (long) b;`\n",
        "   - By casting to `long`, we prevent integer overflow.\n",
        "\n",
        "2. **Extract the Sign Bit:**\n",
        "   - `int k = (int) ((c >> 63) & 1);`\n",
        "   - Shifts the 64-bit `long` number `c` right by 63 bits.\n",
        "   - The result is `0` if `c` is non-negative and `-1` (all ones) if `c` is negative.\n",
        "   - Applying `& 1` extracts the least significant bit, giving `k` as `0` or `1`.\n",
        "\n",
        "3. **Calculate the Maximum:**\n",
        "   - `return a - k * (int) c;`\n",
        "   - If `k == 0` (meaning `a >= b`), returns `a`.\n",
        "   - If `k == 1` (meaning `a < b`), calculates `a - (a - b)`, simplifying to `b`.\n",
        "\n",
        "### View (MaxView.java)\n",
        "\n",
        "- **Purpose:** Responsible for displaying the results to the user.\n",
        "- **Implementation Details:**\n",
        "  - Simple method `displayResult` that prints out the maximum of the two numbers.\n",
        "\n",
        "### Controller (MaxController.java)\n",
        "\n",
        "- **Purpose:** Coordinates interactions between the Model and the View.\n",
        "- **Implementation Details:**\n",
        "  - Initializes the `MaxFinder` (Model) and `MaxView` (View).\n",
        "  - Runs the tests by invoking the `MaxTestHarness`.\n",
        "\n",
        "### Test Harness (MaxTestHarness.java)\n",
        "\n",
        "- **Purpose:** Provides a robust testing framework that runs all tests even if some fail and checks if the test logic is correct in case of failure.\n",
        "- **Implementation Details:**\n",
        "  - **Test Cases:** A series of predefined test cases covering various scenarios, including edge cases like maximum and minimum integer values.\n",
        "  - **Result Tracking:** Stores results of each test case, indicating pass or fail.\n",
        "  - **Test Logic Verification:** If a test fails, it verifies the test logic using an alternative method to ensure the failure is not due to incorrect test logic.\n",
        "  - **Reporting:** After running all tests, it reports the number of tests passed and details of any failures.\n",
        "  - **Integration with View:** Uses the `MaxView` to display results for each test case.\n",
        "\n",
        "### Main (MaxMain.java)\n",
        "\n",
        "- **Purpose:** Entry point of the application.\n",
        "- **Implementation Details:**\n",
        "  - Creates instances of `MaxFinder`, `MaxView`, and `MaxController`.\n",
        "  - Initiates the test run through the controller.\n",
        "\n",
        "**Complexity Analysis:**\n",
        "\n",
        "- **Time Complexity:** O(1) for the `max` method in `MaxFinder` since it performs a constant number of operations regardless of input size.\n",
        "- **Space Complexity:** O(1) for the `max` method as it uses a fixed amount of extra space.\n",
        "\n",
        "**Optimizations:**\n",
        "\n",
        "- **Preventing Integer Overflow:** By casting integers to `long`, we ensure that subtracting large integers does not cause overflow, which could lead to incorrect results.\n",
        "- **Bit Manipulation Efficiency:** Using bitwise operations is generally faster than arithmetic operations involving division or modulus, especially important in performance-critical applications.\n",
        "\n",
        "**Alternative Approaches:**\n",
        "\n",
        "1. **Arithmetic Mean and Absolute Difference:**\n",
        "\n",
        "   ```java\n",
        "   public int max(int a, int b) {\n",
        "       return (a + b + Math.abs(a - b)) / 2;\n",
        "   }\n",
        "   ```\n",
        "\n",
        "   - **Limitations:** The `Math.abs` function may internally use branching, violating the problem constraints.\n",
        "\n",
        "2. **Using Floating-Point Division:**\n",
        "\n",
        "   ```java\n",
        "   public int max(int a, int b) {\n",
        "       return (int) ((a + b + ((double) (a - b)) / Math.abs(a - b)) * 0.5);\n",
        "   }\n",
        "   ```\n",
        "\n",
        "   - **Limitations:** Involves division and floating-point arithmetic, which may not be as efficient and still may use internal branching.\n",
        "\n",
        "**Post-Mortem Analysis:**\n",
        "\n",
        "- **Potential Failures:** Initial implementations without handling integer overflow may fail for inputs where `a - b` exceeds the range of integers.\n",
        "- **Resolution:** Adjusted the `max` method to use `long` for intermediate calculations to safely handle large differences between `a` and `b`.\n",
        "\n",
        "**Test Harness Robustness:**\n",
        "\n",
        "- Ensures all tests are executed, even if some fail.\n",
        "- Verifies test logic in case of failure to differentiate between a failure in the implementation and an error in the test itself.\n",
        "\n",
        "**Conclusion:**\n",
        "\n",
        "The provided solution effectively computes the maximum of two integers without using any if-else statements, branching, or direct comparisons. By employing bit manipulation and careful handling of integer overflows, it adheres to the problem constraints while maintaining readability and efficiency."
      ],
      "metadata": {
        "id": "wH_msgA3o_l-"
      }
    },
    {
      "cell_type": "markdown",
      "source": [
        "## Implementation:\n"
      ],
      "metadata": {
        "id": "nglFSGzOpZb3"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "# Install OpenJDK 11\n",
        "!apt-get install openjdk-11-jdk-headless -qq > /dev/null\n"
      ],
      "metadata": {
        "id": "Fcbahwb0VUyZ"
      },
      "execution_count": 34,
      "outputs": []
    },
    {
      "cell_type": "code",
      "source": [
        "import os\n",
        "os.environ[\"JAVA_HOME\"] = \"/usr/lib/jvm/java-11-openjdk-amd64\"\n",
        "os.environ[\"PATH\"] += os.pathsep + os.path.join(os.environ[\"JAVA_HOME\"], \"bin\")\n"
      ],
      "metadata": {
        "id": "UVu4IaY4VZVy"
      },
      "execution_count": 35,
      "outputs": []
    },
    {
      "cell_type": "code",
      "source": [
        "!java -version\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "oSPZBWqSVbvs",
        "outputId": "42d4d509-4aaf-4f21-b2db-9084d5f80723"
      },
      "execution_count": 36,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "openjdk version \"11.0.24\" 2024-07-16\n",
            "OpenJDK Runtime Environment (build 11.0.24+8-post-Ubuntu-1ubuntu322.04)\n",
            "OpenJDK 64-Bit Server VM (build 11.0.24+8-post-Ubuntu-1ubuntu322.04, mixed mode, sharing)\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "# Create directories for the package structure\n",
        "!mkdir -p src/model src/view src/controller\n"
      ],
      "metadata": {
        "id": "t51eiX_jVgbJ"
      },
      "execution_count": 37,
      "outputs": []
    },
    {
      "cell_type": "code",
      "execution_count": 38,
      "metadata": {
        "id": "EI-cO7pgoQvE",
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "outputId": "e149cdb7-4579-44d5-837c-6fac81160a61"
      },
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting src/model/MaxFinder.java\n"
          ]
        }
      ],
      "source": [
        "%%writefile src/model/MaxFinder.java\n",
        "/**\n",
        " * MaxFinder.java\n",
        " *\n",
        " * Model class containing the logic to find the maximum of two integers\n",
        " * without using any if-else statements, branching, or direct comparisons.\n",
        " *\n",
        " * This implementation uses bit manipulation and arithmetic operations\n",
        " * to determine the maximum value. It handles integer overflows by\n",
        " * using long integers for intermediate calculations.\n",
        " */\n",
        "public class MaxFinder {\n",
        "    /**\n",
        "     * Finds the maximum of two integers without using if-else, branching, or direct comparisons.\n",
        "     *\n",
        "     * @param a First integer\n",
        "     * @param b Second integer\n",
        "     * @return The maximum of a and b\n",
        "     */\n",
        "    public int max(int a, int b) {\n",
        "        // Convert integers to long to prevent integer overflow\n",
        "        long c = (long) a - (long) b;\n",
        "        // Extract the sign bit: 1 if c < 0, 0 otherwise\n",
        "        int k = (int) ((c >> 63) & 1);\n",
        "        // Compute the maximum value based on the sign of c\n",
        "        return a - k * (int) c;\n",
        "    }\n",
        "}\n",
        "\n",
        "/**\n",
        " * MaxView.java\n",
        " *\n",
        " * View class responsible for displaying the results.\n",
        " */\n",
        "public class MaxView {\n",
        "    /**\n",
        "     * Displays the result of the max computation.\n",
        "     *\n",
        "     * @param a   First integer\n",
        "     * @param b   Second integer\n",
        "     * @param max The maximum of a and b\n",
        "     */\n",
        "    public void displayResult(int a, int b, int max) {\n",
        "        System.out.println(\"The maximum of \" + a + \" and \" + b + \" is \" + max + \".\");\n",
        "    }\n",
        "}\n",
        "\n",
        "/**\n",
        " * MaxController.java\n",
        " *\n",
        " * Controller class that coordinates between the Model and the View.\n",
        " */\n",
        "public class MaxController {\n",
        "    private MaxFinder model;\n",
        "    private MaxView view;\n",
        "\n",
        "    public MaxController(MaxFinder model, MaxView view) {\n",
        "        this.model = model;\n",
        "        this.view = view;\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Runs the test cases and displays results.\n",
        "     */\n",
        "    public void runTests() {\n",
        "        MaxTestHarness testHarness = new MaxTestHarness(model, view);\n",
        "        testHarness.runAllTests();\n",
        "        testHarness.reportResults();\n",
        "    }\n",
        "}\n",
        "\n",
        "/**\n",
        " * MaxTestHarness.java\n",
        " *\n",
        " * Test harness for the MaxFinder class.\n",
        " * Ensures that all tests run even if some fail.\n",
        " * If a test fails, it also checks the test logic.\n",
        " */\n",
        "import java.util.ArrayList;\n",
        "import java.util.List;\n",
        "\n",
        "public class MaxTestHarness {\n",
        "    private MaxFinder model;\n",
        "    private MaxView view;\n",
        "    private List<TestResult> testResults;\n",
        "\n",
        "    public MaxTestHarness(MaxFinder model, MaxView view) {\n",
        "        this.model = model;\n",
        "        this.view = view;\n",
        "        this.testResults = new ArrayList<>();\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Runs all the test cases.\n",
        "     */\n",
        "    public void runAllTests() {\n",
        "        testMax(5, 10);\n",
        "        testMax(-5, -10);\n",
        "        testMax(0, 0);\n",
        "        testMax(Integer.MAX_VALUE, Integer.MIN_VALUE);\n",
        "        testMax(-1, 1);\n",
        "        testMax(123456, 654321);\n",
        "        // Add more test cases as needed\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Reports the results of the tests.\n",
        "     */\n",
        "    public void reportResults() {\n",
        "        int passed = 0;\n",
        "        for (TestResult result : testResults) {\n",
        "            if (result.passed) {\n",
        "                passed++;\n",
        "            } else {\n",
        "                System.out.println(\"Test Failed: \" + result.message);\n",
        "            }\n",
        "        }\n",
        "        System.out.println(passed + \" out of \" + testResults.size() + \" tests passed.\");\n",
        "    }\n",
        "\n",
        "    private void testMax(int a, int b) {\n",
        "        int expected = Math.max(a, b);\n",
        "        int actual = model.max(a, b);\n",
        "        boolean passed = expected == actual;\n",
        "        String message = \"max(\" + a + \", \" + b + \") = \" + actual + \", expected: \" + expected;\n",
        "\n",
        "        if (!passed) {\n",
        "            // Test the test logic\n",
        "            boolean testLogicValid = verifyTestLogic(a, b, expected);\n",
        "            if (!testLogicValid) {\n",
        "                message += \" [Test Logic Failure]\";\n",
        "            }\n",
        "        }\n",
        "        testResults.add(new TestResult(passed, message));\n",
        "\n",
        "        // Display the result using the View\n",
        "        view.displayResult(a, b, actual);\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Verifies the test logic to ensure it's not the source of failure.\n",
        "     *\n",
        "     * @param a        First integer\n",
        "     * @param b        Second integer\n",
        "     * @param expected Expected result from the test\n",
        "     * @return True if the test logic is valid, false otherwise\n",
        "     */\n",
        "    private boolean verifyTestLogic(int a, int b, int expected) {\n",
        "        // Recompute the expected value using an alternative method\n",
        "        int alternativeExpected = (a + b + Math.abs(a - b)) / 2;\n",
        "        return expected == alternativeExpected;\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Inner class to store test results.\n",
        "     */\n",
        "    private class TestResult {\n",
        "        boolean passed;\n",
        "        String message;\n",
        "\n",
        "        TestResult(boolean passed, String message) {\n",
        "            this.passed = passed;\n",
        "            this.message = message;\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "/**\n",
        " * MaxMain.java\n",
        " *\n",
        " * Main class to run the application.\n",
        " */\n",
        "public class MaxMain {\n",
        "    public static void main(String[] args) {\n",
        "        MaxFinder model = new MaxFinder();\n",
        "        MaxView view = new MaxView();\n",
        "        MaxController controller = new MaxController(model, view);\n",
        "\n",
        "        controller.runTests();\n",
        "    }\n",
        "}\n"
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "# Compile the Java code\n",
        "!javac -d bin -sourcepath src src/model/MaxFinder.java\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "c6tHsrZrVyKi",
        "outputId": "ed3d8cc6-2eff-4a1c-e4f3-e9544c53f44d"
      },
      "execution_count": 39,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "src/model/MaxFinder.java:78: error: class, interface, or enum expected\n",
            "import java.util.ArrayList;\n",
            "^\n",
            "src/model/MaxFinder.java:79: error: class, interface, or enum expected\n",
            "import java.util.List;\n",
            "^\n",
            "2 errors\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "##Testing:"
      ],
      "metadata": {
        "id": "KT8ZS6TrpcUg"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile src/model/MaxFinder.java\n",
        "package model;\n",
        "\n",
        "/**\n",
        " * MaxFinder.java\n",
        " *\n",
        " * Model class containing the logic to find the maximum of two integers\n",
        " * without using any if-else statements, branching, or direct comparisons.\n",
        " *\n",
        " * This implementation uses bit manipulation and arithmetic operations\n",
        " * to determine the maximum value. It handles integer overflows by\n",
        " * using long integers for intermediate calculations.\n",
        " */\n",
        "public class MaxFinder {\n",
        "    /**\n",
        "     * Finds the maximum of two integers without using if-else, branching, or direct comparisons.\n",
        "     *\n",
        "     * @param a First integer\n",
        "     * @param b Second integer\n",
        "     * @return The maximum of a and b\n",
        "     */\n",
        "    public int max(int a, int b) {\n",
        "        // Convert integers to long to prevent integer overflow\n",
        "        long c = (long) a - (long) b;\n",
        "        // Extract the sign bit: 1 if c < 0, 0 otherwise\n",
        "        int k = (int) ((c >> 63) & 1);\n",
        "        // Compute the maximum value based on the sign of c\n",
        "        return (int)(a - k * c);\n",
        "    }\n",
        "}\n"
      ],
      "metadata": {
        "id": "bdCt7XEVphuE",
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "outputId": "22912f2f-2358-4887-f3b2-844e1aa5a2e8"
      },
      "execution_count": 40,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting src/model/MaxFinder.java\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile src/view/MaxView.java\n",
        "package view;\n",
        "\n",
        "/**\n",
        " * MaxView.java\n",
        " *\n",
        " * View class responsible for displaying the results.\n",
        " */\n",
        "public class MaxView {\n",
        "    /**\n",
        "     * Displays the result of the max computation.\n",
        "     *\n",
        "     * @param a   First integer\n",
        "     * @param b   Second integer\n",
        "     * @param max The maximum of a and b\n",
        "     */\n",
        "    public void displayResult(int a, int b, int max) {\n",
        "        System.out.println(\"The maximum of \" + a + \" and \" + b + \" is \" + max + \".\");\n",
        "    }\n",
        "}\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "rUgqdz31WPWv",
        "outputId": "05303b56-fcf2-45ce-bbab-9ecfa4ec6c24"
      },
      "execution_count": 41,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting src/view/MaxView.java\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile src/controller/MaxController.java\n",
        "package controller;\n",
        "\n",
        "import model.MaxFinder;\n",
        "import view.MaxView;\n",
        "\n",
        "/**\n",
        " * MaxController.java\n",
        " *\n",
        " * Controller class that coordinates between the Model and the View.\n",
        " */\n",
        "public class MaxController {\n",
        "    private MaxFinder model;\n",
        "    private MaxView view;\n",
        "\n",
        "    public MaxController(MaxFinder model, MaxView view) {\n",
        "        this.model = model;\n",
        "        this.view = view;\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Runs the test cases and displays results.\n",
        "     */\n",
        "    public void runTests() {\n",
        "        MaxTestHarness testHarness = new MaxTestHarness(model, view);\n",
        "        testHarness.runAllTests();\n",
        "        testHarness.reportResults();\n",
        "    }\n",
        "}\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "ln4xDc01WTHW",
        "outputId": "96fad8be-b053-46b2-bd76-75c40713854e"
      },
      "execution_count": 42,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting src/controller/MaxController.java\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile src/controller/MaxTestHarness.java\n",
        "package controller;\n",
        "\n",
        "import model.MaxFinder;\n",
        "import view.MaxView;\n",
        "import java.util.ArrayList;\n",
        "import java.util.List;\n",
        "\n",
        "/**\n",
        " * MaxTestHarness.java\n",
        " *\n",
        " * Test harness for the MaxFinder class.\n",
        " * Ensures that all tests run even if some fail.\n",
        " * If a test fails, it also checks the test logic.\n",
        " */\n",
        "public class MaxTestHarness {\n",
        "    private MaxFinder model;\n",
        "    private MaxView view;\n",
        "    private List<TestResult> testResults;\n",
        "\n",
        "    public MaxTestHarness(MaxFinder model, MaxView view) {\n",
        "        this.model = model;\n",
        "        this.view = view;\n",
        "        this.testResults = new ArrayList<>();\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Runs all the test cases.\n",
        "     */\n",
        "    public void runAllTests() {\n",
        "        testMax(5, 10);\n",
        "        testMax(-5, -10);\n",
        "        testMax(0, 0);\n",
        "        testMax(Integer.MAX_VALUE, Integer.MIN_VALUE);\n",
        "        testMax(-1, 1);\n",
        "        testMax(123456, 654321);\n",
        "        // Add more test cases as needed\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Reports the results of the tests.\n",
        "     */\n",
        "    public void reportResults() {\n",
        "        int passed = 0;\n",
        "        for (TestResult result : testResults) {\n",
        "            if (result.passed) {\n",
        "                passed++;\n",
        "            } else {\n",
        "                System.out.println(\"Test Failed: \" + result.message);\n",
        "            }\n",
        "        }\n",
        "        System.out.println(passed + \" out of \" + testResults.size() + \" tests passed.\");\n",
        "    }\n",
        "\n",
        "    private void testMax(int a, int b) {\n",
        "        int expected = Math.max(a, b);\n",
        "        int actual = model.max(a, b);\n",
        "        boolean passed = expected == actual;\n",
        "        String message = \"max(\" + a + \", \" + b + \") = \" + actual + \", expected: \" + expected;\n",
        "\n",
        "        if (!passed) {\n",
        "            // Test the test logic\n",
        "            boolean testLogicValid = verifyTestLogic(a, b, expected);\n",
        "            if (!testLogicValid) {\n",
        "                message += \" [Test Logic Failure]\";\n",
        "            }\n",
        "        }\n",
        "        testResults.add(new TestResult(passed, message));\n",
        "\n",
        "        // Display the result using the View\n",
        "        view.displayResult(a, b, actual);\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Verifies the test logic to ensure it's not the source of failure.\n",
        "     *\n",
        "     * @param a        First integer\n",
        "     * @param b        Second integer\n",
        "     * @param expected Expected result from the test\n",
        "     * @return True if the test logic is valid, false otherwise\n",
        "     */\n",
        "    private boolean verifyTestLogic(int a, int b, int expected) {\n",
        "        // Recompute the expected value using an alternative method\n",
        "        int alternativeExpected = (a + b + Math.abs(a - b)) / 2;\n",
        "        return expected == alternativeExpected;\n",
        "    }\n",
        "\n",
        "    /**\n",
        "     * Inner class to store test results.\n",
        "     */\n",
        "    private class TestResult {\n",
        "        boolean passed;\n",
        "        String message;\n",
        "\n",
        "        TestResult(boolean passed, String message) {\n",
        "            this.passed = passed;\n",
        "            this.message = message;\n",
        "        }\n",
        "    }\n",
        "}\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "W-8T9lQfWV4c",
        "outputId": "36136ab9-3f3a-44a0-869c-5698eb1a0599"
      },
      "execution_count": 43,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting src/controller/MaxTestHarness.java\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile src/MaxMain.java\n",
        "import model.MaxFinder;\n",
        "import view.MaxView;\n",
        "import controller.MaxController;\n",
        "\n",
        "/**\n",
        " * MaxMain.java\n",
        " *\n",
        " * Main class to run the application.\n",
        " */\n",
        "public class MaxMain {\n",
        "    public static void main(String[] args) {\n",
        "        MaxFinder model = new MaxFinder();\n",
        "        MaxView view = new MaxView();\n",
        "        MaxController controller = new MaxController(model, view);\n",
        "\n",
        "        controller.runTests();\n",
        "    }\n",
        "}\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "KREncjyeWY5J",
        "outputId": "56d380cc-94c3-487e-885c-85ebf8e85d60"
      },
      "execution_count": 44,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting src/MaxMain.java\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "# Compile the Java code\n",
        "!javac -d bin -sourcepath src src/MaxMain.java\n"
      ],
      "metadata": {
        "id": "A6w3alWHWbuq"
      },
      "execution_count": 45,
      "outputs": []
    },
    {
      "cell_type": "code",
      "source": [
        "# Run the Java program\n",
        "!java -cp bin MaxMain\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "LdQ5aaHOWfFv",
        "outputId": "131195a5-4f31-43de-b793-e011004c0ebc"
      },
      "execution_count": 46,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "The maximum of 5 and 10 is 10.\n",
            "The maximum of -5 and -10 is -5.\n",
            "The maximum of 0 and 0 is 0.\n",
            "The maximum of 2147483647 and -2147483648 is 2147483647.\n",
            "The maximum of -1 and 1 is 1.\n",
            "The maximum of 123456 and 654321 is 654321.\n",
            "6 out of 6 tests passed.\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "##Elaboration:"
      ],
      "metadata": {
        "id": "5oTeurXJpip9"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "!nvidia-smi\n"
      ],
      "metadata": {
        "id": "S-R25Zj3plcI",
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "outputId": "d62eae41-bcfe-4287-8b7d-e26ea2fdd77f"
      },
      "execution_count": 47,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Wed Sep 25 11:18:29 2024       \n",
            "+---------------------------------------------------------------------------------------+\n",
            "| NVIDIA-SMI 535.104.05             Driver Version: 535.104.05   CUDA Version: 12.2     |\n",
            "|-----------------------------------------+----------------------+----------------------+\n",
            "| GPU  Name                 Persistence-M | Bus-Id        Disp.A | Volatile Uncorr. ECC |\n",
            "| Fan  Temp   Perf          Pwr:Usage/Cap |         Memory-Usage | GPU-Util  Compute M. |\n",
            "|                                         |                      |               MIG M. |\n",
            "|=========================================+======================+======================|\n",
            "|   0  Tesla T4                       Off | 00000000:00:04.0 Off |                    0 |\n",
            "| N/A   38C    P8               9W /  70W |      0MiB / 15360MiB |      0%      Default |\n",
            "|                                         |                      |                  N/A |\n",
            "+-----------------------------------------+----------------------+----------------------+\n",
            "                                                                                         \n",
            "+---------------------------------------------------------------------------------------+\n",
            "| Processes:                                                                            |\n",
            "|  GPU   GI   CI        PID   Type   Process name                            GPU Memory |\n",
            "|        ID   ID                                                             Usage      |\n",
            "|=======================================================================================|\n",
            "|  No running processes found                                                           |\n",
            "+---------------------------------------------------------------------------------------+\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nvcc --version\n"
      ],
      "metadata": {
        "id": "esl-5HoIpouP",
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "outputId": "495f5817-08e3-403c-83cf-88ee84e24f96"
      },
      "execution_count": 48,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "nvcc: NVIDIA (R) Cuda compiler driver\n",
            "Copyright (c) 2005-2023 NVIDIA Corporation\n",
            "Built on Tue_Aug_15_22:02:13_PDT_2023\n",
            "Cuda compilation tools, release 12.2, V12.2.140\n",
            "Build cuda_12.2.r12.2/compiler.33191640_0\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "# Step 2: Write the CUDA code to a file\n",
        "%%writefile MaxFinder.cu\n",
        "\n",
        "#include <iostream>\n",
        "#include <cuda_runtime.h>\n",
        "#include <climits> // For INT_MAX and INT_MIN\n",
        "\n",
        "// Device function to find max without branching\n",
        "__device__ int max_no_branch(int a, int b) {\n",
        "    // Compute difference\n",
        "    long c = (long)a - (long)b;\n",
        "    // Extract sign bit\n",
        "    int k = (int)((c >> 63) & 1);\n",
        "    // Compute max\n",
        "    return a - k * (int)c;\n",
        "}\n",
        "\n",
        "// Kernel function\n",
        "__global__ void maxKernel(int *a, int *b, int *result, int size) {\n",
        "    int idx = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "    if (idx < size) {\n",
        "        result[idx] = max_no_branch(a[idx], b[idx]);\n",
        "    }\n",
        "}\n",
        "\n",
        "int main() {\n",
        "    const int size = 6;\n",
        "    int h_a[size] = {5, -5, 0, INT_MAX, -1, 123456};\n",
        "    int h_b[size] = {10, -10, 0, INT_MIN, 1, 654321};\n",
        "    int h_result[size];\n",
        "\n",
        "    int *d_a, *d_b, *d_result;\n",
        "    cudaMalloc((void**)&d_a, size * sizeof(int));\n",
        "    cudaMalloc((void**)&d_b, size * sizeof(int));\n",
        "    cudaMalloc((void**)&d_result, size * sizeof(int));\n",
        "\n",
        "    cudaMemcpy(d_a, h_a, size * sizeof(int), cudaMemcpyHostToDevice);\n",
        "    cudaMemcpy(d_b, h_b, size * sizeof(int), cudaMemcpyHostToDevice);\n",
        "\n",
        "    int threadsPerBlock = 256;\n",
        "    int blocksPerGrid = (size + threadsPerBlock - 1) / threadsPerBlock;\n",
        "\n",
        "    maxKernel<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_result, size);\n",
        "    cudaDeviceSynchronize();\n",
        "\n",
        "    cudaMemcpy(h_result, d_result, size * sizeof(int), cudaMemcpyDeviceToHost);\n",
        "\n",
        "    for (int i = 0; i < size; ++i) {\n",
        "        std::cout << \"The maximum of \" << h_a[i] << \" and \" << h_b[i] << \" is \" << h_result[i] << \".\\n\";\n",
        "    }\n",
        "\n",
        "    cudaFree(d_a);\n",
        "    cudaFree(d_b);\n",
        "    cudaFree(d_result);\n",
        "\n",
        "    return 0;\n",
        "}\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "EYH3wWgUcQah",
        "outputId": "5ecb3735-9696-464a-8d27-e443b7e76386"
      },
      "execution_count": 49,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting MaxFinder.cu\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "# Step 3: Compile the CUDA code\n",
        "!nvcc MaxFinder.cu -o MaxFinder\n"
      ],
      "metadata": {
        "id": "RVhXhn7acYDk"
      },
      "execution_count": 50,
      "outputs": []
    },
    {
      "cell_type": "code",
      "source": [
        "# Step 4: Run the compiled program\n",
        "!./MaxFinder\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "7IERFNr1cbqT",
        "outputId": "9627e16b-568b-412d-a0b6-717bbe8b7171"
      },
      "execution_count": 51,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "The maximum of 5 and 10 is 10.\n",
            "The maximum of -5 and -10 is -5.\n",
            "The maximum of 0 and 0 is 0.\n",
            "The maximum of 2147483647 and -2147483648 is 2147483647.\n",
            "The maximum of -1 and 1 is 1.\n",
            "The maximum of 123456 and 654321 is 654321.\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "##Post Mortem:"
      ],
      "metadata": {
        "id": "y9VGKmHUpmUW"
      }
    },
    {
      "cell_type": "code",
      "source": [],
      "metadata": {
        "id": "zLemyh-D6MSU"
      },
      "execution_count": null,
      "outputs": []
    },
    {
      "cell_type": "markdown",
      "source": [
        "Explanation:\n",
        "\n",
        "Registers Used:\n",
        "\n",
        "RDI: Holds the first argument a.\n",
        "RSI: Holds the second argument b.\n",
        "RAX: Used for calculations and holds the return value.\n",
        "RDX: Temporary register for calculations.\n",
        "Instructions:\n",
        "\n",
        "movsxd: Sign-extend a 32-bit integer to a 64-bit integer.\n",
        "mov, sub, shr, imul: Standard arithmetic and bitwise operations."
      ],
      "metadata": {
        "id": "6LScWKPupsUH"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "# Check the runtime type (optional)\n",
        "import sys\n",
        "print(sys.version)\n",
        "# Install NASM\n",
        "!apt-get update -qq\n",
        "!apt-get install -qq nasm\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "JUfEUqEy6o01",
        "outputId": "d5e8967f-1701-43c9-8f6a-6bc30b066501"
      },
      "execution_count": 52,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "3.10.12 (main, Sep 11 2024, 15:47:36) [GCC 11.4.0]\n",
            "W: Skipping acquire of configured file 'main/source/Sources' as repository 'https://r2u.stat.illinois.edu/ubuntu jammy InRelease' does not seem to provide it (sources.list entry misspelt?)\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile MaxFinder.asm\n",
        "; MaxFinder.asm\n",
        ";\n",
        "; Assembly function to find the maximum of two integers without using\n",
        "; if-else statements, branching, or direct comparisons.\n",
        ";\n",
        "; This implementation uses arithmetic and bitwise operations to determine\n",
        "; the maximum value. It is designed to be called from a C program on 64-bit Linux.\n",
        "\n",
        "section .text\n",
        "    global max_no_branch\n",
        "\n",
        "; Function: int max_no_branch(int a, int b)\n",
        "; Arguments:\n",
        ";   a in RDI\n",
        ";   b in RSI\n",
        "; Returns:\n",
        ";   max(a, b) in RAX\n",
        "\n",
        "max_no_branch:\n",
        "    ; Convert inputs to 64-bit integers\n",
        "    movsxd  rdi, edi        ; Sign-extend 'a' to 64 bits\n",
        "    movsxd  rsi, esi        ; Sign-extend 'b' to 64 bits\n",
        "\n",
        "    ; Compute c = (long)a - (long)b\n",
        "    mov     rax, rdi        ; RAX = a\n",
        "    sub     rax, rsi        ; RAX = c = a - b\n",
        "\n",
        "    ; Extract sign bit k = (c >> 63) & 1\n",
        "    mov     rdx, rax        ; RDX = c\n",
        "    shr     rdx, 63         ; RDX = k\n",
        "\n",
        "    ; Compute k * c\n",
        "    imul    rdx, rax        ; RDX = k * c\n",
        "\n",
        "    ; Compute result = a - k * c\n",
        "    mov     rax, rdi        ; RAX = a\n",
        "    sub     rax, rdx        ; RAX = a - k * c\n",
        "\n",
        "    ret\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "d_uHShQH6vyc",
        "outputId": "f4310f58-2ce4-48a7-e468-8d191ce6cfec"
      },
      "execution_count": 53,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting MaxFinder.asm\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile MaxMain.c\n",
        "/* MaxMain.c\n",
        "\n",
        "   Main C program to test the assembly function for finding the maximum\n",
        "   of two integers without using if-else statements, branching, or direct comparisons.\n",
        "*/\n",
        "\n",
        "#include <stdio.h>\n",
        "#include <limits.h>\n",
        "\n",
        "// Declaration of the assembly function\n",
        "int max_no_branch(int a, int b);\n",
        "\n",
        "int main() {\n",
        "    int test_cases[][2] = {\n",
        "        {5, 10},\n",
        "        {-5, -10},\n",
        "        {0, 0},\n",
        "        {INT_MAX, INT_MIN},\n",
        "        {-1, 1},\n",
        "        {123456, 654321}\n",
        "    };\n",
        "\n",
        "    int num_tests = sizeof(test_cases) / sizeof(test_cases[0]);\n",
        "\n",
        "    for (int i = 0; i < num_tests; ++i) {\n",
        "        int a = test_cases[i][0];\n",
        "        int b = test_cases[i][1];\n",
        "        int result = max_no_branch(a, b);\n",
        "        printf(\"The maximum of %d and %d is %d.\\n\", a, b, result);\n",
        "    }\n",
        "\n",
        "    return 0;\n",
        "}\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "TPaOw_3ZpMjC",
        "outputId": "01b38938-24c4-4031-e6be-8f4d6e7b44dc"
      },
      "execution_count": 54,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting MaxMain.c\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nasm -f elf64 MaxFinder.asm -o MaxFinder.o\n"
      ],
      "metadata": {
        "id": "e9NM2WXjpPVK"
      },
      "execution_count": 55,
      "outputs": []
    },
    {
      "cell_type": "code",
      "source": [
        "!gcc MaxMain.c MaxFinder.o -o MaxApp\n"
      ],
      "metadata": {
        "id": "DgUJgTnrpgmI"
      },
      "execution_count": 56,
      "outputs": []
    },
    {
      "cell_type": "code",
      "source": [
        "!./MaxApp\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "DzO4Q5Jspjd-",
        "outputId": "7b6616a2-f1a5-4a56-bb0d-c4e174a75a71"
      },
      "execution_count": 57,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "The maximum of 5 and 10 is 10.\n",
            "The maximum of -5 and -10 is -5.\n",
            "The maximum of 0 and 0 is 0.\n",
            "The maximum of 2147483647 and -2147483648 is 2147483647.\n",
            "The maximum of -1 and 1 is 1.\n",
            "The maximum of 123456 and 654321 is 654321.\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "Certainly! Let's adapt the assembly code to 64-bit and show you how to run it in the CUDA Colab environment, which is 64-bit Linux. We'll use NASM (Netwide Assembler) to assemble the code and GCC to compile the C code that calls the assembly function.\n",
        "\n",
        "---\n",
        "\n",
        "## **Overview**\n",
        "\n",
        "- **Objective:** Run the assembly version of the \"maximum of two numbers without using any if-else statements, branching, or direct comparisons\" in the 64-bit CUDA Colab environment.\n",
        "- **Tools Used:**\n",
        "  - **NASM:** For assembling the assembly code.\n",
        "  - **GCC:** For compiling and linking the C and assembly code.\n",
        "- **Environment:** 64-bit Linux (Ubuntu) in Google Colab.\n",
        "\n",
        "---\n",
        "\n",
        "## **Step-by-Step Guide**\n",
        "\n",
        "### **Step 1: Set Up the Colab Environment**\n",
        "\n",
        "First, ensure that your Colab notebook is using a GPU runtime (since you mentioned the CUDA environment). However, since we're running assembly code that doesn't use the GPU, it's fine to use either a CPU or GPU runtime.\n",
        "\n",
        "```python\n",
        "# Check the runtime type (optional)\n",
        "import sys\n",
        "print(sys.version)\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "### **Step 2: Install NASM**\n",
        "\n",
        "Install NASM to assemble the assembly code.\n",
        "\n",
        "```python\n",
        "# Install NASM\n",
        "!apt-get update -qq\n",
        "!apt-get install -qq nasm\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "### **Step 3: Write the Assembly Code**\n",
        "\n",
        "We'll write the 64-bit assembly code into a file using Colab's cell magic `%%writefile`.\n",
        "\n",
        "#### **3.1 MaxFinder.asm**\n",
        "\n",
        "```python\n",
        "%%writefile MaxFinder.asm\n",
        "; MaxFinder.asm\n",
        ";\n",
        "; Assembly function to find the maximum of two integers without using\n",
        "; if-else statements, branching, or direct comparisons.\n",
        ";\n",
        "; This implementation uses arithmetic and bitwise operations to determine\n",
        "; the maximum value. It is designed to be called from a C program on 64-bit Linux.\n",
        "\n",
        "section .text\n",
        "    global max_no_branch\n",
        "\n",
        "; Function: int max_no_branch(int a, int b)\n",
        "; Arguments:\n",
        ";   a in RDI\n",
        ";   b in RSI\n",
        "; Returns:\n",
        ";   max(a, b) in RAX\n",
        "\n",
        "max_no_branch:\n",
        "    ; Convert inputs to 64-bit integers\n",
        "    movsxd  rdi, edi        ; Sign-extend 'a' to 64 bits\n",
        "    movsxd  rsi, esi        ; Sign-extend 'b' to 64 bits\n",
        "\n",
        "    ; Compute c = (long)a - (long)b\n",
        "    mov     rax, rdi        ; RAX = a\n",
        "    sub     rax, rsi        ; RAX = c = a - b\n",
        "\n",
        "    ; Extract sign bit k = (c >> 63) & 1\n",
        "    mov     rdx, rax        ; RDX = c\n",
        "    shr     rdx, 63         ; RDX = k\n",
        "\n",
        "    ; Compute k * c\n",
        "    imul    rdx, rax        ; RDX = k * c\n",
        "\n",
        "    ; Compute result = a - k * c\n",
        "    mov     rax, rdi        ; RAX = a\n",
        "    sub     rax, rdx        ; RAX = a - k * c\n",
        "\n",
        "    ret\n",
        "```\n",
        "\n",
        "**Explanation:**\n",
        "\n",
        "- **Registers Used:**\n",
        "  - **RDI:** Holds the first argument `a`.\n",
        "  - **RSI:** Holds the second argument `b`.\n",
        "  - **RAX:** Used for calculations and holds the return value.\n",
        "  - **RDX:** Temporary register for calculations.\n",
        "\n",
        "- **Instructions:**\n",
        "  - **movsxd:** Sign-extend a 32-bit integer to a 64-bit integer.\n",
        "  - **mov, sub, shr, imul:** Standard arithmetic and bitwise operations.\n",
        "\n",
        "---\n",
        "\n",
        "### **Step 4: Write the C Code**\n",
        "\n",
        "#### **4.1 MaxMain.c**\n",
        "\n",
        "```python\n",
        "%%writefile MaxMain.c\n",
        "/* MaxMain.c\n",
        "\n",
        "   Main C program to test the assembly function for finding the maximum\n",
        "   of two integers without using if-else statements, branching, or direct comparisons.\n",
        "*/\n",
        "\n",
        "#include <stdio.h>\n",
        "#include <limits.h>\n",
        "\n",
        "// Declaration of the assembly function\n",
        "int max_no_branch(int a, int b);\n",
        "\n",
        "int main() {\n",
        "    int test_cases[][2] = {\n",
        "        {5, 10},\n",
        "        {-5, -10},\n",
        "        {0, 0},\n",
        "        {INT_MAX, INT_MIN},\n",
        "        {-1, 1},\n",
        "        {123456, 654321}\n",
        "    };\n",
        "\n",
        "    int num_tests = sizeof(test_cases) / sizeof(test_cases[0]);\n",
        "\n",
        "    for (int i = 0; i < num_tests; ++i) {\n",
        "        int a = test_cases[i][0];\n",
        "        int b = test_cases[i][1];\n",
        "        int result = max_no_branch(a, b);\n",
        "        printf(\"The maximum of %d and %d is %d.\\n\", a, b, result);\n",
        "    }\n",
        "\n",
        "    return 0;\n",
        "}\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "### **Step 5: Assemble and Compile the Code**\n",
        "\n",
        "#### **5.1 Assemble the Assembly Code**\n",
        "\n",
        "```python\n",
        "!nasm -f elf64 MaxFinder.asm -o MaxFinder.o\n",
        "```\n",
        "\n",
        "#### **5.2 Compile the C Code and Link with the Assembly Object File**\n",
        "\n",
        "```python\n",
        "!gcc MaxMain.c MaxFinder.o -o MaxApp\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "### **Step 6: Run the Program**\n",
        "\n",
        "```python\n",
        "!./MaxApp\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "## **Full Colab Notebook Code**\n",
        "\n",
        "For convenience, here's the complete code in one place:\n",
        "\n",
        "```python\n",
        "# Step 2: Install NASM\n",
        "!apt-get update -qq\n",
        "!apt-get install -qq nasm\n",
        "\n",
        "# Step 3: Write the Assembly Code\n",
        "%%writefile MaxFinder.asm\n",
        "; MaxFinder.asm\n",
        ";\n",
        "; Assembly function to find the maximum of two integers without using\n",
        "; if-else statements, branching, or direct comparisons.\n",
        ";\n",
        "; This implementation uses arithmetic and bitwise operations to determine\n",
        "; the maximum value. It is designed to be called from a C program on 64-bit Linux.\n",
        "\n",
        "section .text\n",
        "    global max_no_branch\n",
        "\n",
        "; Function: int max_no_branch(int a, int b)\n",
        "; Arguments:\n",
        ";   a in RDI\n",
        ";   b in RSI\n",
        "; Returns:\n",
        ";   max(a, b) in RAX\n",
        "\n",
        "max_no_branch:\n",
        "    ; Convert inputs to 64-bit integers\n",
        "    movsxd  rdi, edi        ; Sign-extend 'a' to 64 bits\n",
        "    movsxd  rsi, esi        ; Sign-extend 'b' to 64 bits\n",
        "\n",
        "    ; Compute c = (long)a - (long)b\n",
        "    mov     rax, rdi        ; RAX = a\n",
        "    sub     rax, rsi        ; RAX = c = a - b\n",
        "\n",
        "    ; Extract sign bit k = (c >> 63) & 1\n",
        "    mov     rdx, rax        ; RDX = c\n",
        "    shr     rdx, 63         ; RDX = k\n",
        "\n",
        "    ; Compute k * c\n",
        "    imul    rdx, rax        ; RDX = k * c\n",
        "\n",
        "    ; Compute result = a - k * c\n",
        "    mov     rax, rdi        ; RAX = a\n",
        "    sub     rax, rdx        ; RAX = a - k * c\n",
        "\n",
        "    ret\n",
        "\n",
        "# Step 4: Write the C Code\n",
        "%%writefile MaxMain.c\n",
        "/* MaxMain.c\n",
        "\n",
        "   Main C program to test the assembly function for finding the maximum\n",
        "   of two integers without using if-else statements, branching, or direct comparisons.\n",
        "*/\n",
        "\n",
        "#include <stdio.h>\n",
        "#include <limits.h>\n",
        "\n",
        "// Declaration of the assembly function\n",
        "int max_no_branch(int a, int b);\n",
        "\n",
        "int main() {\n",
        "    int test_cases[][2] = {\n",
        "        {5, 10},\n",
        "        {-5, -10},\n",
        "        {0, 0},\n",
        "        {INT_MAX, INT_MIN},\n",
        "        {-1, 1},\n",
        "        {123456, 654321}\n",
        "    };\n",
        "\n",
        "    int num_tests = sizeof(test_cases) / sizeof(test_cases[0]);\n",
        "\n",
        "    for (int i = 0; i < num_tests; ++i) {\n",
        "        int a = test_cases[i][0];\n",
        "        int b = test_cases[i][1];\n",
        "        int result = max_no_branch(a, b);\n",
        "        printf(\"The maximum of %d and %d is %d.\\n\", a, b, result);\n",
        "    }\n",
        "\n",
        "    return 0;\n",
        "}\n",
        "\n",
        "# Step 5: Assemble and Compile the Code\n",
        "!nasm -f elf64 MaxFinder.asm -o MaxFinder.o\n",
        "!gcc MaxMain.c MaxFinder.o -o MaxApp\n",
        "\n",
        "# Step 6: Run the Program\n",
        "!./MaxApp\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "## **Expected Output**\n",
        "\n",
        "When you run the program, you should see:\n",
        "\n",
        "```\n",
        "The maximum of 5 and 10 is 10.\n",
        "The maximum of -5 and -10 is -5.\n",
        "The maximum of 0 and 0 is 0.\n",
        "The maximum of 2147483647 and -2147483648 is 2147483647.\n",
        "The maximum of -1 and 1 is 1.\n",
        "The maximum of 123456 and 654321 is 654321.\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "## **Detailed Explanation**\n",
        "\n",
        "### **Assembly Code Breakdown**\n",
        "\n",
        "**Function Signature:**\n",
        "\n",
        "- **Name:** `max_no_branch`\n",
        "- **Parameters:**\n",
        "  - `int a` in **RDI**\n",
        "  - `int b` in **RSI**\n",
        "- **Return Value:** `int` in **RAX**\n",
        "\n",
        "**Steps:**\n",
        "\n",
        "1. **Sign-Extend Inputs:**\n",
        "\n",
        "   - `movsxd rdi, edi` and `movsxd rsi, esi` ensure that the 32-bit integers `a` and `b` are properly sign-extended to 64 bits.\n",
        "\n",
        "2. **Compute Difference `c = a - b`:**\n",
        "\n",
        "   - `mov rax, rdi` sets `RAX` to `a`.\n",
        "   - `sub rax, rsi` subtracts `b` from `a`, resulting in `c` in `RAX`.\n",
        "\n",
        "3. **Extract Sign Bit `k`:**\n",
        "\n",
        "   - `mov rdx, rax` copies `c` to `RDX`.\n",
        "   - `shr rdx, 63` shifts `c` right by 63 bits, leaving the sign bit in `RDX`.\n",
        "\n",
        "4. **Compute `k * c`:**\n",
        "\n",
        "   - `imul rdx, rax` multiplies `k` (0 or 1) by `c`, storing the result in `RDX`.\n",
        "\n",
        "5. **Compute Maximum `max = a - k * c`:**\n",
        "\n",
        "   - `mov rax, rdi` sets `RAX` back to `a`.\n",
        "   - `sub rax, rdx` subtracts `k * c` from `a`, resulting in the maximum value in `RAX`.\n",
        "\n",
        "6. **Return:**\n",
        "\n",
        "   - `ret` returns control to the caller with the result in `RAX`.\n",
        "\n",
        "### **C Code Explanation**\n",
        "\n",
        "- **Function Declaration:**\n",
        "\n",
        "  ```c\n",
        "  int max_no_branch(int a, int b);\n",
        "  ```\n",
        "\n",
        "- **Test Cases:**\n",
        "\n",
        "  An array of integer pairs to test the `max_no_branch` function.\n",
        "\n",
        "- **Loop Through Test Cases:**\n",
        "\n",
        "  Calls `max_no_branch` for each pair and prints the result.\n",
        "\n",
        "### **Compilation and Linking**\n",
        "\n",
        "- **Assemble the Assembly Code:**\n",
        "\n",
        "  ```bash\n",
        "  nasm -f elf64 MaxFinder.asm -o MaxFinder.o\n",
        "  ```\n",
        "\n",
        "  - `-f elf64` specifies the output format suitable for 64-bit ELF binaries.\n",
        "\n",
        "- **Compile and Link with GCC:**\n",
        "\n",
        "  ```bash\n",
        "  gcc MaxMain.c MaxFinder.o -o MaxApp\n",
        "  ```\n",
        "\n",
        "  - Links the C code and the assembled object file into an executable.\n",
        "\n",
        "### **Running the Program**\n",
        "\n",
        "- Executes `./MaxApp` to run the compiled program and display the results.\n",
        "\n",
        "---\n",
        "\n",
        "## **Potential Issues and Solutions**\n",
        "\n",
        "1. **Assembler Errors:**\n",
        "\n",
        "   - **Issue:** If NASM reports errors during assembly, double-check the syntax and ensure that the instructions are valid for 64-bit assembly.\n",
        "   - **Solution:** Ensure that all registers are correctly prefixed (e.g., use `rax`, `rbx` for 64-bit registers).\n",
        "\n",
        "2. **Linker Errors:**\n",
        "\n",
        "   - **Issue:** If GCC reports undefined references or linking issues.\n",
        "   - **Solution:** Verify that the assembly function is correctly declared in the C code and that the object files are correctly specified in the GCC command.\n",
        "\n",
        "3. **Runtime Errors:**\n",
        "\n",
        "   - **Issue:** Incorrect results or segmentation faults.\n",
        "   - **Solution:** Ensure that the assembly code correctly handles data types and that register usage follows the calling conventions.\n",
        "\n",
        "4. **Colab Environment Resets:**\n",
        "\n",
        "   - **Issue:** Colab sessions reset after inactivity, causing loss of installed packages and files.\n",
        "   - **Solution:** Rerun the setup cells to reinstall NASM and recreate the files.\n",
        "\n",
        "---\n",
        "\n",
        "## **Additional Notes**\n",
        "\n",
        "- **64-bit vs. 32-bit Differences:**\n",
        "\n",
        "  - **Register Names:** Use `r` prefix for 64-bit registers (e.g., `rax`, `rbx`).\n",
        "  - **Calling Conventions:** Parameters are passed in registers (`RDI`, `RSI`, etc.) rather than on the stack.\n",
        "  - **Data Sizes:** Be cautious with data sizes and sign extension.\n",
        "\n",
        "- **CUDA Environment:**\n",
        "\n",
        "  - While we're running this code in the CUDA Colab environment, the assembly code itself doesn't utilize the GPU.\n",
        "  - If you wish to integrate assembly code with CUDA kernels, that's a more advanced topic involving device code and may require different tools.\n",
        "\n",
        "- **Integration with CUDA (Advanced):**\n",
        "\n",
        "  - If you aim to run assembly code within CUDA kernels, you'll need to use **PTX (Parallel Thread Execution)** assembly, which is NVIDIA's intermediate language for CUDA.\n",
        "  - This is beyond the scope of standard x86 assembly and requires knowledge of CUDA's programming model.\n",
        "\n",
        "---\n",
        "\n",
        "## **Conclusion**\n",
        "\n",
        "You now have the 64-bit assembly version of the code and the steps to run it in the CUDA Colab environment. By following this guide, you can experiment with assembly programming in a 64-bit Linux environment and understand how to interface assembly code with C programs.\n",
        "\n",
        "---\n",
        "\n",
        "**Feel free to ask if you have any questions or need further assistance with running or modifying the code!**"
      ],
      "metadata": {
        "id": "ePKUTUrO7Ip1"
      }
    },
    {
      "cell_type": "markdown",
      "source": [
        "Certainly! Running assembly code on the GPU involves using **PTX (Parallel Thread Execution)**, which is NVIDIA's low-level assembly language for CUDA-enabled GPUs. PTX allows you to write assembly-like code that runs directly on the GPU, providing fine-grained control over GPU operations.\n",
        "\n",
        "In this answer, I'll guide you through:\n",
        "\n",
        "1. **Understanding PTX and Its Role in CUDA**\n",
        "2. **Writing the Max Function in PTX Assembly**\n",
        "3. **Integrating PTX Code into CUDA Programs**\n",
        "4. **Compiling and Running the Code in Colab**\n",
        "5. **Providing a Full Colab Notebook Example**\n",
        "\n",
        "---\n",
        "\n",
        "## **1. Understanding PTX and Its Role in CUDA**\n",
        "\n",
        "### **What Is PTX?**\n",
        "\n",
        "- **PTX (Parallel Thread Execution)** is an intermediate assembly language for NVIDIA GPUs.\n",
        "- It provides a stable programming model and instruction set for general-purpose parallel programming.\n",
        "- PTX code is compiled by the CUDA compiler (`nvcc`) into binary code that runs on the GPU.\n",
        "\n",
        "### **Why Use PTX?**\n",
        "\n",
        "- **Fine-Grained Control:** Writing in PTX allows you to optimize performance-critical sections of code.\n",
        "- **Access to Low-Level Instructions:** PTX exposes instructions that might not be directly accessible in high-level CUDA code.\n",
        "- **Educational Purposes:** Understanding PTX can deepen your knowledge of how GPU execution works.\n",
        "\n",
        "### **When to Use PTX?**\n",
        "\n",
        "- When you need to optimize kernels beyond what the CUDA compiler provides.\n",
        "- For experimenting with low-level GPU programming.\n",
        "\n",
        "---\n",
        "\n",
        "## **2. Writing the Max Function in PTX Assembly**\n",
        "\n",
        "We'll rewrite the `max_no_branch` function in PTX assembly to run on the GPU. The goal is to find the maximum of two integers without using any branching.\n",
        "\n",
        "### **Algorithm Recap**\n",
        "\n",
        "The algorithm computes the maximum of two integers `a` and `b` without using if-else statements:\n",
        "\n",
        "1. Compute `c = a - b`\n",
        "2. Extract the sign bit `k` of `c`: `k = (c >> 31) & 1` (for 32-bit integers)\n",
        "3. Compute `max = a - k * c`\n",
        "\n",
        "### **PTX Implementation**\n",
        "\n",
        "PTX uses a syntax similar to assembly languages but tailored for GPU execution.\n",
        "\n",
        "---\n",
        "\n",
        "## **3. Integrating PTX Code into CUDA Programs**\n",
        "\n",
        "We can integrate PTX code into CUDA programs in two main ways:\n",
        "\n",
        "1. **Inline PTX Assembly within CUDA C++ Code**\n",
        "2. **Writing PTX Code in a Separate File and Loading It at Runtime**\n",
        "\n",
        "For simplicity, we'll use inline PTX within a CUDA kernel.\n",
        "\n",
        "### **Steps to Integrate PTX into CUDA**\n",
        "\n",
        "1. **Write the CUDA Kernel**\n",
        "2. **Embed PTX Instructions Using Inline Assembly**\n",
        "3. **Compile and Run the CUDA Program**\n",
        "\n",
        "---\n",
        "\n",
        "## **4. Compiling and Running the Code in Colab**\n",
        "\n",
        "### **4.1 Setting Up the Colab Environment**\n",
        "\n",
        "First, ensure that your Colab notebook is set up to use a GPU:\n",
        "\n",
        "- Go to **Runtime > Change runtime type**.\n",
        "- Select **GPU** as the hardware accelerator.\n",
        "\n",
        "Verify that the GPU is available:\n",
        "\n",
        "```python\n",
        "# Verify GPU availability\n",
        "!nvidia-smi\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "### **4.2 Writing the CUDA Code with PTX**\n",
        "\n",
        "We'll write a CUDA program that:\n",
        "\n",
        "- Uses a kernel function to compute the maximum of two arrays of integers.\n",
        "- Within the kernel, we'll use inline PTX assembly to perform the max operation.\n",
        "\n",
        "#### **4.2.1 Write the CUDA Code**\n",
        "\n",
        "```python\n",
        "%%writefile max_ptx.cu\n",
        "\n",
        "#include <stdio.h>\n",
        "#include <cuda_runtime.h>\n",
        "\n",
        "#define SIZE 6\n",
        "\n",
        "// Kernel function using PTX to compute max without branching\n",
        "__global__ void maxKernel(int *a, int *b, int *result) {\n",
        "    int idx = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "    if (idx < SIZE) {\n",
        "        int val_a = a[idx];\n",
        "        int val_b = b[idx];\n",
        "        int max_val;\n",
        "\n",
        "        // Variables for PTX\n",
        "        int c, k;\n",
        "\n",
        "        // Compute c = a - b\n",
        "        asm volatile (\"sub.s32 %0, %1, %2;\" : \"=r\"(c) : \"r\"(val_a), \"r\"(val_b));\n",
        "\n",
        "        // Extract sign bit k = (c >> 31) & 1\n",
        "        asm volatile (\"shr.s32 %0, %1, 31;\" : \"=r\"(k) : \"r\"(c));\n",
        "        asm volatile (\"and.b32 %0, %1, 1;\" : \"+r\"(k) : \"r\"(k));\n",
        "\n",
        "        // Compute max = a - k * c\n",
        "        asm volatile (\"mul.lo.s32 %0, %1, %2;\" : \"+r\"(k) : \"r\"(k), \"r\"(c));\n",
        "        asm volatile (\"sub.s32 %0, %1, %2;\" : \"=r\"(max_val) : \"r\"(val_a), \"r\"(k));\n",
        "\n",
        "        // Store the result\n",
        "        result[idx] = max_val;\n",
        "    }\n",
        "}\n",
        "\n",
        "int main() {\n",
        "    int h_a[SIZE] = {5, -5, 0, INT_MAX, -1, 123456};\n",
        "    int h_b[SIZE] = {10, -10, 0, INT_MIN, 1, 654321};\n",
        "    int h_result[SIZE];\n",
        "\n",
        "    int *d_a, *d_b, *d_result;\n",
        "\n",
        "    // Allocate device memory\n",
        "    cudaMalloc((void**)&d_a, SIZE * sizeof(int));\n",
        "    cudaMalloc((void**)&d_b, SIZE * sizeof(int));\n",
        "    cudaMalloc((void**)&d_result, SIZE * sizeof(int));\n",
        "\n",
        "    // Copy data to device\n",
        "    cudaMemcpy(d_a, h_a, SIZE * sizeof(int), cudaMemcpyHostToDevice);\n",
        "    cudaMemcpy(d_b, h_b, SIZE * sizeof(int), cudaMemcpyHostToDevice);\n",
        "\n",
        "    // Launch kernel\n",
        "    int threadsPerBlock = 256;\n",
        "    int blocksPerGrid = (SIZE + threadsPerBlock - 1) / threadsPerBlock;\n",
        "    maxKernel<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_result);\n",
        "    cudaDeviceSynchronize();\n",
        "\n",
        "    // Copy result back to host\n",
        "    cudaMemcpy(h_result, d_result, SIZE * sizeof(int), cudaMemcpyDeviceToHost);\n",
        "\n",
        "    // Display results\n",
        "    for (int i = 0; i < SIZE; ++i) {\n",
        "        printf(\"The maximum of %d and %d is %d.\\n\", h_a[i], h_b[i], h_result[i]);\n",
        "    }\n",
        "\n",
        "    // Free device memory\n",
        "    cudaFree(d_a);\n",
        "    cudaFree(d_b);\n",
        "    cudaFree(d_result);\n",
        "\n",
        "    return 0;\n",
        "}\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "### **4.3 Explanation of the CUDA Code with PTX**\n",
        "\n",
        "#### **4.3.1 Inline PTX Assembly Syntax**\n",
        "\n",
        "- **`asm volatile (\"PTX instructions\" : output operands : input operands : clobbers);`**\n",
        "- The `volatile` keyword tells the compiler not to optimize away the assembly code.\n",
        "\n",
        "#### **4.3.2 PTX Instructions Used**\n",
        "\n",
        "1. **Subtraction (`sub.s32`)**\n",
        "\n",
        "   ```cpp\n",
        "   asm volatile (\"sub.s32 %0, %1, %2;\" : \"=r\"(c) : \"r\"(val_a), \"r\"(val_b));\n",
        "   ```\n",
        "\n",
        "   - Subtracts `val_b` from `val_a`, storing the result in `c`.\n",
        "\n",
        "2. **Shift Right Logical (`shr.s32`)**\n",
        "\n",
        "   ```cpp\n",
        "   asm volatile (\"shr.s32 %0, %1, 31;\" : \"=r\"(k) : \"r\"(c));\n",
        "   ```\n",
        "\n",
        "   - Shifts `c` right by 31 bits to extract the sign bit.\n",
        "\n",
        "3. **Bitwise AND (`and.b32`)**\n",
        "\n",
        "   ```cpp\n",
        "   asm volatile (\"and.b32 %0, %1, 1;\" : \"+r\"(k) : \"r\"(k));\n",
        "   ```\n",
        "\n",
        "   - ANDs `k` with `1` to get the least significant bit (sign bit).\n",
        "\n",
        "4. **Multiply (`mul.lo.s32`)**\n",
        "\n",
        "   ```cpp\n",
        "   asm volatile (\"mul.lo.s32 %0, %1, %2;\" : \"+r\"(k) : \"r\"(k), \"r\"(c));\n",
        "   ```\n",
        "\n",
        "   - Multiplies `k` by `c`.\n",
        "\n",
        "5. **Subtraction (`sub.s32`)**\n",
        "\n",
        "   ```cpp\n",
        "   asm volatile (\"sub.s32 %0, %1, %2;\" : \"=r\"(max_val) : \"r\"(val_a), \"r\"(k));\n",
        "   ```\n",
        "\n",
        "   - Subtracts `k` from `val_a` to get the maximum value.\n",
        "\n",
        "#### **4.3.3 Notes**\n",
        "\n",
        "- **Registers (`%0`, `%1`, etc.):** Placeholders for operands in the inline assembly.\n",
        "- **Constraints (`\"=r\"(output), \"r\"(input)`):** Specify that operands are in registers.\n",
        "\n",
        "---\n",
        "\n",
        "### **4.4 Compiling the CUDA Code**\n",
        "\n",
        "Compile the CUDA code using `nvcc`:\n",
        "\n",
        "```python\n",
        "!nvcc -o max_ptx max_ptx.cu\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "### **4.5 Running the Compiled Program**\n",
        "\n",
        "```python\n",
        "!./max_ptx\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "## **5. Full Colab Notebook Example**\n",
        "\n",
        "Here's the complete code in a Colab notebook:\n",
        "\n",
        "```python\n",
        "# Step 1: Verify GPU availability\n",
        "!nvidia-smi\n",
        "```\n",
        "\n",
        "```python\n",
        "# Step 2: Write the CUDA code with PTX\n",
        "%%writefile max_ptx.cu\n",
        "\n",
        "#include <stdio.h>\n",
        "#include <cuda_runtime.h>\n",
        "#include <climits> // For INT_MAX and INT_MIN\n",
        "\n",
        "#define SIZE 6\n",
        "\n",
        "// Kernel function using PTX to compute max without branching\n",
        "__global__ void maxKernel(int *a, int *b, int *result) {\n",
        "    int idx = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "    if (idx < SIZE) {\n",
        "        int val_a = a[idx];\n",
        "        int val_b = b[idx];\n",
        "        int max_val;\n",
        "\n",
        "        // Variables for PTX\n",
        "        int c, k;\n",
        "\n",
        "        // Compute c = a - b\n",
        "        asm volatile (\"sub.s32 %0, %1, %2;\" : \"=r\"(c) : \"r\"(val_a), \"r\"(val_b));\n",
        "\n",
        "        // Extract sign bit k = (c >> 31) & 1\n",
        "        asm volatile (\"shr.s32 %0, %1, 31;\" : \"=r\"(k) : \"r\"(c));\n",
        "        asm volatile (\"and.b32 %0, %1, 1;\" : \"+r\"(k) : \"r\"(k));\n",
        "\n",
        "        // Compute k * c\n",
        "        asm volatile (\"mul.lo.s32 %0, %1, %2;\" : \"+r\"(k) : \"r\"(k), \"r\"(c));\n",
        "\n",
        "        // Compute max = a - k * c\n",
        "        asm volatile (\"sub.s32 %0, %1, %2;\" : \"=r\"(max_val) : \"r\"(val_a), \"r\"(k));\n",
        "\n",
        "        // Store the result\n",
        "        result[idx] = max_val;\n",
        "    }\n",
        "}\n",
        "\n",
        "int main() {\n",
        "    int h_a[SIZE] = {5, -5, 0, INT_MAX, -1, 123456};\n",
        "    int h_b[SIZE] = {10, -10, 0, INT_MIN, 1, 654321};\n",
        "    int h_result[SIZE];\n",
        "\n",
        "    int *d_a, *d_b, *d_result;\n",
        "\n",
        "    // Allocate device memory\n",
        "    cudaMalloc((void**)&d_a, SIZE * sizeof(int));\n",
        "    cudaMalloc((void**)&d_b, SIZE * sizeof(int));\n",
        "    cudaMalloc((void**)&d_result, SIZE * sizeof(int));\n",
        "\n",
        "    // Copy data to device\n",
        "    cudaMemcpy(d_a, h_a, SIZE * sizeof(int), cudaMemcpyHostToDevice);\n",
        "    cudaMemcpy(d_b, h_b, SIZE * sizeof(int), cudaMemcpyHostToDevice);\n",
        "\n",
        "    // Launch kernel\n",
        "    int threadsPerBlock = 256;\n",
        "    int blocksPerGrid = (SIZE + threadsPerBlock - 1) / threadsPerBlock;\n",
        "    maxKernel<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_result);\n",
        "    cudaDeviceSynchronize();\n",
        "\n",
        "    // Copy result back to host\n",
        "    cudaMemcpy(h_result, d_result, SIZE * sizeof(int), cudaMemcpyDeviceToHost);\n",
        "\n",
        "    // Display results\n",
        "    for (int i = 0; i < SIZE; ++i) {\n",
        "        printf(\"The maximum of %d and %d is %d.\\n\", h_a[i], h_b[i], h_result[i]);\n",
        "    }\n",
        "\n",
        "    // Free device memory\n",
        "    cudaFree(d_a);\n",
        "    cudaFree(d_b);\n",
        "    cudaFree(d_result);\n",
        "\n",
        "    return 0;\n",
        "}\n",
        "```\n",
        "\n",
        "```python\n",
        "# Step 3: Compile the CUDA code\n",
        "!nvcc -o max_ptx max_ptx.cu\n",
        "```\n",
        "\n",
        "```python\n",
        "# Step 4: Run the compiled program\n",
        "!./max_ptx\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "## **Expected Output**\n",
        "\n",
        "When you run the program, you should see:\n",
        "\n",
        "```\n",
        "The maximum of 5 and 10 is 10.\n",
        "The maximum of -5 and -10 is -5.\n",
        "The maximum of 0 and 0 is 0.\n",
        "The maximum of 2147483647 and -2147483648 is 2147483647.\n",
        "The maximum of -1 and 1 is 1.\n",
        "The maximum of 123456 and 654321 is 654321.\n",
        "```\n",
        "\n",
        "---\n",
        "\n",
        "## **6. Understanding the PTX Instructions**\n",
        "\n",
        "Let's delve deeper into the PTX instructions used:\n",
        "\n",
        "- **`sub.s32`**: Subtracts two 32-bit signed integers.\n",
        "- **`shr.s32`**: Shifts a 32-bit signed integer right by a specified number of bits.\n",
        "- **`and.b32`**: Performs a bitwise AND on two 32-bit integers.\n",
        "- **`mul.lo.s32`**: Multiplies two 32-bit signed integers and returns the lower 32 bits of the result.\n",
        "\n",
        "### **Alternate PTX Code Using Predicates**\n",
        "\n",
        "Alternatively, you can use PTX predicates to avoid branching:\n",
        "\n",
        "```cpp\n",
        "asm volatile (\n",
        "    \"setp.lt.s32 %0, %1, %2;\\n\\t\"\n",
        "    \"selp.s32 %3, %2, %1, %0;\\n\\t\"\n",
        "    : \"=r\"(p), \"=r\"(val_a), \"=r\"(val_b), \"=r\"(max_val)\n",
        "    : \"r\"(val_a), \"r\"(val_b)\n",
        ");\n",
        "```\n",
        "\n",
        "- **`setp.lt.s32`**: Sets a predicate if `val_a < val_b`.\n",
        "- **`selp.s32`**: Selects between `val_b` and `val_a` based on the predicate `p`.\n",
        "\n",
        "However, using predicates may introduce branching, which we want to avoid per the problem constraints.\n",
        "\n",
        "---\n",
        "\n",
        "## **7. Additional Notes**\n",
        "\n",
        "### **7.1 Advantages of Using PTX**\n",
        "\n",
        "- **Performance Optimization:** PTX can potentially optimize critical sections of code.\n",
        "- **Understanding Hardware Execution:** Provides insights into how the GPU executes instructions.\n",
        "\n",
        "### **7.2 Limitations**\n",
        "\n",
        "- **Complexity:** Writing PTX is more complex than writing high-level CUDA code.\n",
        "- **Portability:** PTX code may not be as portable across different GPU architectures compared to high-level CUDA code.\n",
        "\n",
        "### **7.3 When to Use PTX vs. CUDA Intrinsics**\n",
        "\n",
        "- For most applications, high-level CUDA code and built-in functions (`__shfl`, `__syncthreads`, etc.) suffice.\n",
        "- PTX is useful when you need control beyond what's available in CUDA C++.\n",
        "\n",
        "---\n",
        "\n",
        "## **8. Conclusion**\n",
        "\n",
        "You've now learned how to:\n",
        "\n",
        "- Write a max function using PTX assembly code.\n",
        "- Integrate PTX code into a CUDA kernel.\n",
        "- Compile and run the code in the Colab environment.\n",
        "\n",
        "This approach allows you to execute assembly-like code directly on the GPU, giving you fine-grained control over GPU operations.\n",
        "\n",
        "---\n",
        "\n",
        "## **9. Troubleshooting**\n",
        "\n",
        "### **Compilation Errors**\n",
        "\n",
        "- **Error:** `unknown register name '%0' in 'asm'`\n",
        "  - **Solution:** Ensure you're using the correct syntax for inline PTX assembly and that all operands are correctly specified.\n",
        "\n",
        "### **Runtime Errors**\n",
        "\n",
        "- **Error:** Incorrect results\n",
        "  - **Solution:** Verify the PTX instructions and ensure that the logic matches the intended algorithm.\n",
        "\n",
        "### **Environment Issues**\n",
        "\n",
        "- **Issue:** Colab session disconnects or resets\n",
        "  - **Solution:** Rerun the setup cells to recompile and execute the code.\n",
        "\n",
        "---\n",
        "\n",
        "## **10. Further Exploration**\n",
        "\n",
        "- **Explore More PTX Instructions:** Read the [PTX ISA (Instruction Set Architecture) documentation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html) to learn about other instructions.\n",
        "- **Performance Profiling:** Use tools like NVIDIA Nsight to profile your CUDA applications and see if PTX optimizations make a significant difference.\n",
        "- **Advanced CUDA Features:** Experiment with warp-level primitives and shared memory for further optimizations.\n",
        "\n",
        "---\n",
        "\n",
        "**Feel free to ask if you have any questions or need further assistance with the code or concepts discussed!**"
      ],
      "metadata": {
        "id": "-Y18VCxT7VMs"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile max_ptx.cu\n",
        "\n",
        "#include <stdio.h>\n",
        "#include <cuda_runtime.h>\n",
        "\n",
        "#define SIZE 6\n",
        "\n",
        "// Kernel function using PTX to compute max without branching\n",
        "__global__ void maxKernel(int *a, int *b, int *result) {\n",
        "    int idx = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "    if (idx < SIZE) {\n",
        "        int val_a = a[idx];\n",
        "        int val_b = b[idx];\n",
        "        int max_val;\n",
        "\n",
        "        // Variables for PTX\n",
        "        int c, k;\n",
        "\n",
        "        // Compute c = a - b\n",
        "        asm volatile (\"sub.s32 %0, %1, %2;\" : \"=r\"(c) : \"r\"(val_a), \"r\"(val_b));\n",
        "\n",
        "        // Extract sign bit k = (c >> 31) & 1\n",
        "        asm volatile (\"shr.s32 %0, %1, 31;\" : \"=r\"(k) : \"r\"(c));\n",
        "        asm volatile (\"and.b32 %0, %1, 1;\" : \"+r\"(k) : \"r\"(k));\n",
        "\n",
        "        // Compute max = a - k * c\n",
        "        asm volatile (\"mul.lo.s32 %0, %1, %2;\" : \"+r\"(k) : \"r\"(k), \"r\"(c));\n",
        "        asm volatile (\"sub.s32 %0, %1, %2;\" : \"=r\"(max_val) : \"r\"(val_a), \"r\"(k));\n",
        "\n",
        "        // Store the result\n",
        "        result[idx] = max_val;\n",
        "    }\n",
        "}\n",
        "\n",
        "int main() {\n",
        "    int h_a[SIZE] = {5, -5, 0, INT_MAX, -1, 123456};\n",
        "    int h_b[SIZE] = {10, -10, 0, INT_MIN, 1, 654321};\n",
        "    int h_result[SIZE];\n",
        "\n",
        "    int *d_a, *d_b, *d_result;\n",
        "\n",
        "    // Allocate device memory\n",
        "    cudaMalloc((void**)&d_a, SIZE * sizeof(int));\n",
        "    cudaMalloc((void**)&d_b, SIZE * sizeof(int));\n",
        "    cudaMalloc((void**)&d_result, SIZE * sizeof(int));\n",
        "\n",
        "    // Copy data to device\n",
        "    cudaMemcpy(d_a, h_a, SIZE * sizeof(int), cudaMemcpyHostToDevice);\n",
        "    cudaMemcpy(d_b, h_b, SIZE * sizeof(int), cudaMemcpyHostToDevice);\n",
        "\n",
        "    // Launch kernel\n",
        "    int threadsPerBlock = 256;\n",
        "    int blocksPerGrid = (SIZE + threadsPerBlock - 1) / threadsPerBlock;\n",
        "    maxKernel<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_result);\n",
        "    cudaDeviceSynchronize();\n",
        "\n",
        "    // Copy result back to host\n",
        "    cudaMemcpy(h_result, d_result, SIZE * sizeof(int), cudaMemcpyDeviceToHost);\n",
        "\n",
        "    // Display results\n",
        "    for (int i = 0; i < SIZE; ++i) {\n",
        "        printf(\"The maximum of %d and %d is %d.\\n\", h_a[i], h_b[i], h_result[i]);\n",
        "    }\n",
        "\n",
        "    // Free device memory\n",
        "    cudaFree(d_a);\n",
        "    cudaFree(d_b);\n",
        "    cudaFree(d_result);\n",
        "\n",
        "    return 0;\n",
        "}\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "X4gi-4IYqQWV",
        "outputId": "3bfd2ccf-6a4d-4207-be44-bea6b9266e63"
      },
      "execution_count": 59,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting max_ptx.cu\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nvcc -o max_ptx max_ptx.cu\n"
      ],
      "metadata": {
        "id": "yBPsLyRJqX_o"
      },
      "execution_count": 60,
      "outputs": []
    },
    {
      "cell_type": "code",
      "source": [
        "!./max_ptx\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "EYcEWFT8qbot",
        "outputId": "25660d8b-846b-4b7d-e962-7e7680b49a28"
      },
      "execution_count": 61,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "The maximum of 5 and 10 is 10.\n",
            "The maximum of -5 and -10 is -5.\n",
            "The maximum of 0 and 0 is 0.\n",
            "The maximum of 2147483647 and -2147483648 is -2147483648.\n",
            "The maximum of -1 and 1 is 1.\n",
            "The maximum of 123456 and 654321 is 654321.\n"
          ]
        }
      ]
    }
  ]
}