diff --git a/GPU_puzzlers.ipynb b/GPU_puzzlers.ipynb
index 034ee2c..097fa08 100644
--- a/GPU_puzzlers.ipynb
+++ b/GPU_puzzlers.ipynb
@@ -1,6546 +1,1227 @@
{
- "cells": [
- {
- "cell_type": "markdown",
- "id": "23454034",
- "metadata": {
- "lines_to_next_cell": 2,
- "id": "23454034"
- },
- "source": [
- "# GPU Puzzles\n",
- "- by [Sasha Rush](http://rush-nlp.com) - [srush_nlp](https://twitter.com/srush_nlp)"
- ]
- },
- {
- "cell_type": "markdown",
- "id": "09ffa5eb",
- "metadata": {
- "id": "09ffa5eb"
- },
- "source": [
- "![](https://github.com/srush/GPU-Puzzles/raw/main/cuda.png)"
- ]
- },
- {
- "cell_type": "markdown",
- "id": "c9064f5c",
- "metadata": {
- "id": "c9064f5c"
- },
- "source": [
- "GPU architectures are critical to machine learning, and seem to be\n",
- "becoming even more important every day. However, you can be an expert\n",
- "in machine learning without ever touching GPU code. It is hard to gain\n",
- "intuition working through abstractions. "
- ]
- },
- {
- "cell_type": "markdown",
- "id": "d4af993a",
- "metadata": {
- "id": "d4af993a"
- },
- "source": [
- "This notebook is an attempt to teach beginner GPU programming in a\n",
- "completely interactive fashion. Instead of providing text with\n",
- "concepts, it throws you right into coding and building GPU\n",
- "kernels. The exercises use NUMBA which directly maps Python\n",
- "code to CUDA kernels. It looks like Python but is basically\n",
- "identical to writing low-level CUDA code. \n",
- "In a few hours, I think you can go from basics to\n",
- "understanding the real algorithms that power 99% of deep learning\n",
- "today. If you do want to read the manual, it is here:"
- ]
- },
- {
- "cell_type": "markdown",
- "id": "cb974c6b",
- "metadata": {
- "id": "cb974c6b"
- },
- "source": [
- "[NUMBA CUDA Guide](https://numba.readthedocs.io/en/stable/cuda/index.html)"
- ]
- },
- {
- "cell_type": "markdown",
- "id": "783f5a51",
- "metadata": {
- "id": "783f5a51"
- },
- "source": [
- "I recommend doing these in Colab, as it is easy to get started. Be\n",
- "sure to make your own copy, turn on GPU mode in the settings (`Runtime / Change runtime type`, then set `Hardware accelerator` to `GPU`), and\n",
- "then get to coding."
- ]
- },
- {
- "cell_type": "markdown",
- "id": "743b3805",
- "metadata": {
- "id": "743b3805"
- },
- "source": [
- "[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/srush/GPU-Puzzles/blob/main/GPU_puzzlers.ipynb)"
- ]
- },
- {
- "cell_type": "markdown",
- "id": "044bd802",
- "metadata": {
- "id": "044bd802"
- },
- "source": [
- "(If you are into this style of puzzle, also check out my [Tensor\n",
- "Puzzles](https://github.com/srush/Tensor-Puzzles) for PyTorch.)"
- ]
- },
- {
- "cell_type": "code",
- "execution_count": 5,
- "id": "4c078331",
- "metadata": {
- "execution": {
- "iopub.execute_input": "2022-08-07T21:24:28.090989Z",
- "iopub.status.busy": "2022-08-07T21:24:28.090509Z",
- "iopub.status.idle": "2022-08-07T21:24:40.218920Z",
- "shell.execute_reply": "2022-08-07T21:24:40.217864Z"
- },
- "lines_to_next_cell": 2,
- "id": "4c078331",
- "outputId": "f3711c44-23da-464b-90fe-b615880d1a14",
- "colab": {
- "base_uri": "https://localhost:8080/"
- }
- },
- "outputs": [
- {
- "output_type": "stream",
- "name": "stdout",
- "text": [
- " Installing build dependencies ... \u001b[?25l\u001b[?25hdone\n",
- " Getting requirements to build wheel ... \u001b[?25l\u001b[?25hdone\n",
- " Preparing metadata (pyproject.toml) ... \u001b[?25l\u001b[?25hdone\n",
- " Building wheel for chalk-diagrams (pyproject.toml) ... \u001b[?25l\u001b[?25hdone\n"
- ]
- }
- ],
- "source": [
- "!pip install -qqq git+https://github.com/chalk-diagrams/planar git+https://github.com/danoneata/chalk@srush-patch-1\n",
- "!wget -q https://github.com/srush/GPU-Puzzles/raw/main/robot.png https://github.com/srush/GPU-Puzzles/raw/main/lib.py"
- ]
- },
- {
- "cell_type": "code",
- "execution_count": null,
- "id": "53cba43c",
- "metadata": {
- "execution": {
- "iopub.execute_input": "2022-08-07T21:24:40.223406Z",
- "iopub.status.busy": "2022-08-07T21:24:40.222945Z",
- "iopub.status.idle": "2022-08-07T21:24:40.697137Z",
- "shell.execute_reply": "2022-08-07T21:24:40.696287Z"
- },
- "id": "53cba43c"
- },
- "outputs": [],
- "source": [
- "import numba\n",
- "import numpy as np\n",
- "import warnings\n",
- "from lib import CudaProblem, Coord"
- ]
- },
- {
- "cell_type": "code",
- "execution_count": null,
- "id": "76a82629",
- "metadata": {
- "execution": {
- "iopub.execute_input": "2022-08-07T21:24:40.701357Z",
- "iopub.status.busy": "2022-08-07T21:24:40.700827Z",
- "iopub.status.idle": "2022-08-07T21:24:40.705366Z",
- "shell.execute_reply": "2022-08-07T21:24:40.704574Z"
- },
- "lines_to_next_cell": 2,
- "id": "76a82629"
- },
- "outputs": [],
- "source": [
- "warnings.filterwarnings(\n",
- " action=\"ignore\", category=numba.NumbaPerformanceWarning, module=\"numba\"\n",
- ")"
- ]
- },
- {
- "cell_type": "markdown",
- "id": "d0ad1b7a",
- "metadata": {
- "lines_to_next_cell": 2,
- "id": "d0ad1b7a"
- },
- "source": [
- "## Puzzle 1: Map\n",
- "\n",
- "Implement a \"kernel\" (GPU function) that adds 10 to each position of vector `a`\n",
- "and stores it in vector `out`. You have 1 thread per position."
- ]
- },
- {
- "cell_type": "markdown",
- "id": "2bcbce0d",
- "metadata": {
- "lines_to_next_cell": 2,
- "id": "2bcbce0d"
- },
- "source": [
- "**Warning** This code looks like Python but it is really CUDA! You cannot use\n",
- "standard python tools like list comprehensions or ask for Numpy properties\n",
- "like shape or size (if you need the size, it is given as an argument).\n",
- "The puzzles only require doing simple operations, basically\n",
- "+, *, simple array indexing, for loops, and if statements.\n",
- "You are allowed to use local variables. \n",
- "If you get an\n",
- "error it is probably because you did something fancy :). "
- ]
- },
- {
- "cell_type": "markdown",
- "id": "cb5e1d85",
- "metadata": {
- "id": "cb5e1d85"
- },
- "source": [
- "*Tip: Think of the function `call` as being run 1 time for each thread.\n",
- "The only difference is that `cuda.threadIdx.x` changes each time.*"
- ]
- },
- {
- "cell_type": "code",
- "execution_count": null,
- "id": "ab0899de",
- "metadata": {
- "execution": {
- "iopub.execute_input": "2022-08-07T21:24:40.710192Z",
- "iopub.status.busy": "2022-08-07T21:24:40.709745Z",
- "iopub.status.idle": "2022-08-07T21:24:40.779921Z",
- "shell.execute_reply": "2022-08-07T21:24:40.779104Z"
- },
- "id": "ab0899de",
- "outputId": "a98a5828-bf1d-402f-dd61-68cfe9b69d9d"
- },
- "outputs": [
- {
- "name": "stdout",
- "output_type": "stream",
- "text": [
- "# Map\n",
- " \n",
- " Score (Max Per Thread):\n",
- " | Global Reads | Global Writes | Shared Reads | Shared Writes |\n",
- " | 0 | 0 | 0 | 0 | \n",
- "\n"
- ]
- },
- {
- "data": {
- "image/svg+xml": [
- "\n",
- ""
- ],
- "text/plain": [
- "Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Rectangle(width=11.616000000000001, height=12.936255200000002, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -5.643127600000001), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Primitive(shape=Spacer(width=0, height=1.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.75), diagram=Primitive(shape=Text(text='Map', font_size=1), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 1.2501), diagram=Primitive(shape=Spacer(width=0, height=1), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 2.2500999999999998), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -1.6500000000000001,\n",
- " 0.0, 1.0, -0.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 1.1,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 2.2,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 3.3000000000000003,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0)))))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 3.2500999999999998), diagram=Primitive(shape=Spacer(width=0, height=1), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 7.380166), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -0.0), diagram=Compose(envelope=, diagram1=Empty(), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=ApplyTransform(transform=Affine(1.0, 0.0, -3.5,\n",
- " 0.0, 1.0, -0.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Primitive(shape=Spacer(width=1, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 3.5,\n",
- " 0.0, 1.0, 0.0), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, 0.5000499999999999), diagram=Compose(envelope=, diagram1=ApplyTransform(transform=Affine(1.0, 0.0, -2.5,\n",
- " 0.0, 1.0, -0.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, 0.25), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.05), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -1.5), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=ApplyName(dname=('a', 0), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyName(dname=('a', 1), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 1.0)))), diagram2=ApplyName(dname=('a', 2), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 2.0)))), diagram2=ApplyName(dname=('a', 3), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 3.0)))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -2.25), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.0), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Text(text='a', font_size=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 1.5,\n",
- " 0.0, 1.0, 0.0), diagram=Primitive(shape=Spacer(width=2.0, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 2.5,\n",
- " 0.0, 1.0, 0.0), diagram=Empty())), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 3.5,\n",
- " 0.0, 1.0, 0.0), diagram=Primitive(shape=Spacer(width=2.0, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 5.0,\n",
- " 0.0, 1.0, 0.0), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, 0.25), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.05), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -1.5), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=ApplyName(dname=('out', 0), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyName(dname=('out', 1), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 1.0)))), diagram2=ApplyName(dname=('out', 2), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 2.0)))), diagram2=ApplyName(dname=('out', 3), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 3.0)))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -2.25), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.0), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Text(text='out', font_size=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))))))), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty()))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -3.0000999999999998), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.25), diagram=ApplyStyle(style=Style(line_width=(, 0), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Text(text='Block 0 0', font_size=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.5001), diagram=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))))))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 7.0,\n",
- " 0.0, 1.0, 0.0), diagram=Primitive(shape=Spacer(width=1, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))))), diagram2=Empty()), diagram2=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.0, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Rectangle(width=9.6, height=6.6001199999999995, radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))))), diagram2=Empty())))))), diagram2=Empty())))"
- ]
- },
- "execution_count": 4,
- "metadata": {},
- "output_type": "execute_result"
- }
- ],
- "source": [
- "def map_spec(a):\n",
- " return a + 10\n",
- "\n",
- "\n",
- "def map_test(cuda):\n",
- " def call(out, a) -> None:\n",
- " local_i = cuda.threadIdx.x\n",
- " # FILL ME IN (roughly 1 lines)\n",
- "\n",
- " return call\n",
- "\n",
- "\n",
- "SIZE = 4\n",
- "out = np.zeros((SIZE,))\n",
- "a = np.arange(SIZE)\n",
- "problem = CudaProblem(\n",
- " \"Map\", map_test, [a], out, threadsperblock=Coord(SIZE, 1), spec=map_spec\n",
- ")\n",
- "problem.show()"
- ]
- },
- {
- "cell_type": "code",
- "execution_count": null,
- "id": "0a5520a6",
- "metadata": {
- "execution": {
- "iopub.execute_input": "2022-08-07T21:24:40.783843Z",
- "iopub.status.busy": "2022-08-07T21:24:40.783394Z",
- "iopub.status.idle": "2022-08-07T21:24:41.318965Z",
- "shell.execute_reply": "2022-08-07T21:24:41.318117Z"
- },
- "lines_to_end_of_cell_marker": 0,
- "lines_to_next_cell": 1,
- "id": "0a5520a6",
- "outputId": "d3f1e273-d5e0-4e4b-ccaf-fd28840f899f"
- },
- "outputs": [
- {
- "name": "stdout",
- "output_type": "stream",
- "text": [
- "Failed Tests.\n",
- "Yours: [0. 0. 0. 0.]\n",
- "Spec : [10 11 12 13]\n"
- ]
- }
- ],
- "source": [
- "problem.check()"
- ]
- },
- {
- "cell_type": "markdown",
- "id": "59aecc28",
- "metadata": {
- "id": "59aecc28"
- },
- "source": [
- "## Puzzle 2 - Zip\n",
- "\n",
- "Implement a kernel that adds together each position of `a` and `b` and stores it in `out`.\n",
- "You have 1 thread per position."
- ]
- },
- {
- "cell_type": "code",
- "execution_count": null,
- "id": "619b4208",
- "metadata": {
- "execution": {
- "iopub.execute_input": "2022-08-07T21:24:41.322974Z",
- "iopub.status.busy": "2022-08-07T21:24:41.322540Z",
- "iopub.status.idle": "2022-08-07T21:24:41.408046Z",
- "shell.execute_reply": "2022-08-07T21:24:41.407300Z"
- },
- "lines_to_next_cell": 0,
- "id": "619b4208",
- "outputId": "9c06a417-9d69-4787-ca33-3c222496053e"
- },
- "outputs": [
- {
- "name": "stdout",
- "output_type": "stream",
- "text": [
- "# Zip\n",
- " \n",
- " Score (Max Per Thread):\n",
- " | Global Reads | Global Writes | Shared Reads | Shared Writes |\n",
- " | 0 | 0 | 0 | 0 | \n",
- "\n"
- ]
- },
- {
- "data": {
- "image/svg+xml": [
- "\n",
- ""
- ],
- "text/plain": [
- "Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Rectangle(width=11.616000000000001, height=22.3742552, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -10.3621276), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Primitive(shape=Spacer(width=0, height=1.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.75), diagram=Primitive(shape=Text(text='Zip', font_size=1), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 1.2501), diagram=Primitive(shape=Spacer(width=0, height=1), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 2.2500999999999998), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -1.6500000000000001,\n",
- " 0.0, 1.0, -0.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 1.1,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 2.2,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 3.3000000000000003,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0)))))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 3.2500999999999998), diagram=Primitive(shape=Spacer(width=0, height=1), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 11.670166), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -0.0), diagram=Compose(envelope=, diagram1=Empty(), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=ApplyTransform(transform=Affine(1.0, 0.0, -3.5,\n",
- " 0.0, 1.0, -0.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Primitive(shape=Spacer(width=1, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 3.5,\n",
- " 0.0, 1.0, 0.0), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, 0.5000499999999999), diagram=Compose(envelope=, diagram1=ApplyTransform(transform=Affine(1.0, 0.0, -2.5,\n",
- " 0.0, 1.0, -0.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -3.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.05), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -1.5), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=ApplyName(dname=('a', 0), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyName(dname=('a', 1), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 1.0)))), diagram2=ApplyName(dname=('a', 2), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 2.0)))), diagram2=ApplyName(dname=('a', 3), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 3.0)))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -2.25), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.0), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Text(text='a', font_size=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 3.0), diagram=Primitive(shape=Spacer(width=2.0, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0, 1.0, 0.0,\n",
- " -1.0, 0.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 6.5), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.05), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -1.5), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=ApplyName(dname=('b', 0), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyName(dname=('b', 1), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 1.0)))), diagram2=ApplyName(dname=('b', 2), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 2.0)))), diagram2=ApplyName(dname=('b', 3), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 3.0)))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -2.25), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.0), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Text(text='b', font_size=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 1.5,\n",
- " 0.0, 1.0, 0.0), diagram=Primitive(shape=Spacer(width=2.0, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 2.5,\n",
- " 0.0, 1.0, 0.0), diagram=Empty())), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 3.5,\n",
- " 0.0, 1.0, 0.0), diagram=Primitive(shape=Spacer(width=2.0, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 5.0,\n",
- " 0.0, 1.0, 0.0), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, 0.25), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.05), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -1.5), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=ApplyName(dname=('out', 0), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyName(dname=('out', 1), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 1.0)))), diagram2=ApplyName(dname=('out', 2), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 2.0)))), diagram2=ApplyName(dname=('out', 3), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 3.0)))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -2.25), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.0), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Text(text='out', font_size=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))))))), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty()), diagram2=Empty()))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -6.2501), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.25), diagram=ApplyStyle(style=Style(line_width=(, 0), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Text(text='Block 0 0', font_size=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.5001), diagram=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))))))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 7.0,\n",
- " 0.0, 1.0, 0.0), diagram=Primitive(shape=Spacer(width=1, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))))), diagram2=Empty()), diagram2=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.0, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Rectangle(width=9.6, height=14.40012, radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))))), diagram2=Empty())))))), diagram2=Empty())))"
- ]
- },
- "execution_count": 6,
- "metadata": {},
- "output_type": "execute_result"
- }
- ],
- "source": [
- "def zip_spec(a, b):\n",
- " return a + b\n",
- "\n",
- "\n",
- "def zip_test(cuda):\n",
- " def call(out, a, b) -> None:\n",
- " local_i = cuda.threadIdx.x\n",
- " # FILL ME IN (roughly 1 lines)\n",
- "\n",
- " return call\n",
- "\n",
- "\n",
- "SIZE = 4\n",
- "out = np.zeros((SIZE,))\n",
- "a = np.arange(SIZE)\n",
- "b = np.arange(SIZE)\n",
- "problem = CudaProblem(\n",
- " \"Zip\", zip_test, [a, b], out, threadsperblock=Coord(SIZE, 1), spec=zip_spec\n",
- ")\n",
- "problem.show()"
- ]
- },
- {
- "cell_type": "code",
- "execution_count": null,
- "id": "93ddda39",
- "metadata": {
- "id": "93ddda39"
- },
- "outputs": [],
- "source": []
- },
- {
- "cell_type": "code",
- "execution_count": null,
- "id": "f5037564",
- "metadata": {
- "execution": {
- "iopub.execute_input": "2022-08-07T21:24:41.411997Z",
- "iopub.status.busy": "2022-08-07T21:24:41.411569Z",
- "iopub.status.idle": "2022-08-07T21:24:41.450428Z",
- "shell.execute_reply": "2022-08-07T21:24:41.449609Z"
- },
- "lines_to_end_of_cell_marker": 0,
- "lines_to_next_cell": 1,
- "id": "f5037564",
- "outputId": "22798fcf-0f3c-47a0-98c5-dbe671f25f57"
- },
- "outputs": [
- {
- "name": "stdout",
- "output_type": "stream",
- "text": [
- "Failed Tests.\n",
- "Yours: [0. 0. 0. 0.]\n",
- "Spec : [0 2 4 6]\n"
- ]
- }
- ],
- "source": [
- "problem.check()"
- ]
- },
- {
- "cell_type": "markdown",
- "id": "bb90ab5b",
- "metadata": {
- "id": "bb90ab5b"
- },
- "source": [
- "## Puzzle 3 - Guards\n",
- "\n",
- "Implement a kernel that adds 10 to each position of `a` and stores it in `out`.\n",
- "You have more threads than positions."
- ]
- },
- {
- "cell_type": "code",
- "execution_count": null,
- "id": "f5c83704",
- "metadata": {
- "execution": {
- "iopub.execute_input": "2022-08-07T21:24:41.454648Z",
- "iopub.status.busy": "2022-08-07T21:24:41.454179Z",
- "iopub.status.idle": "2022-08-07T21:24:41.527083Z",
- "shell.execute_reply": "2022-08-07T21:24:41.526268Z"
- },
- "id": "f5c83704",
- "outputId": "2ca047fc-ab05-43f6-e13c-bf0c5504a54a"
- },
- "outputs": [
- {
- "name": "stdout",
- "output_type": "stream",
- "text": [
- "# Guard\n",
- " \n",
- " Score (Max Per Thread):\n",
- " | Global Reads | Global Writes | Shared Reads | Shared Writes |\n",
- " | 0 | 0 | 0 | 0 | \n",
- "\n"
- ]
- },
- {
- "data": {
- "image/svg+xml": [
- "\n",
- ""
- ],
- "text/plain": [
- "Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Rectangle(width=11.616000000000001, height=12.936255200000002, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -5.643127600000001), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Primitive(shape=Spacer(width=0, height=1.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.75), diagram=Primitive(shape=Text(text='Guard', font_size=1), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 1.2501), diagram=Primitive(shape=Spacer(width=0, height=1), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 2.2500999999999998), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -3.8500000000000005,\n",
- " 0.0, 1.0, -0.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 1.1,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 2.2,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 3.3000000000000003,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 4.4,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 5.5,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 6.6000000000000005,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 7.700000000000001,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.7, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Circle(radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Image(local_path='robot.png', url_path='https://raw.githubusercontent.com/minitorch/diagrams/main/robot.png'), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(0.0008333333333333334, 0.0, 0.0,\n",
- " 0.0, 0.0008333333333333334, 0.0)))))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 3.2500999999999998), diagram=Primitive(shape=Spacer(width=0, height=1), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 7.380166), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -0.0), diagram=Compose(envelope=, diagram1=Empty(), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=ApplyTransform(transform=Affine(1.0, 0.0, -3.5,\n",
- " 0.0, 1.0, -0.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Primitive(shape=Spacer(width=1, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 3.5,\n",
- " 0.0, 1.0, 0.0), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, 0.5000499999999999), diagram=Compose(envelope=, diagram1=ApplyTransform(transform=Affine(1.0, 0.0, -2.5,\n",
- " 0.0, 1.0, -0.0), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, 0.25), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.05), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -1.5), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=ApplyName(dname=('a', 0), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyName(dname=('a', 1), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 1.0)))), diagram2=ApplyName(dname=('a', 2), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 2.0)))), diagram2=ApplyName(dname=('a', 3), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 3.0)))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -2.25), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.0), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Text(text='a', font_size=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 1.5,\n",
- " 0.0, 1.0, 0.0), diagram=Primitive(shape=Spacer(width=2.0, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 2.5,\n",
- " 0.0, 1.0, 0.0), diagram=Empty())), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 3.5,\n",
- " 0.0, 1.0, 0.0), diagram=Primitive(shape=Spacer(width=2.0, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 5.0,\n",
- " 0.0, 1.0, 0.0), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, 0.25), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.05), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -1.5), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=ApplyName(dname=('out', 0), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=ApplyName(dname=('out', 1), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 1.0)))), diagram2=ApplyName(dname=('out', 2), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 2.0)))), diagram2=ApplyName(dname=('out', 3), diagram=Primitive(shape=Rectangle(width=1, height=1, radius=None), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 3.0)))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -2.25), diagram=Compose(envelope=, diagram1=ApplyStyle(style=Style(line_width=(, 0.0), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Text(text='out', font_size=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))), diagram2=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))))))), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty())), diagram2=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Empty(), diagram2=Empty()), diagram2=Empty()))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, -0.0,\n",
- " 0.0, 1.0, -3.0000999999999998), diagram=Compose(envelope=, diagram1=Compose(envelope=, diagram1=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.25), diagram=ApplyStyle(style=Style(line_width=(, 0), line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Text(text='Block 0 0', font_size=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0)))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.5001), diagram=Primitive(shape=Spacer(width=0, height=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))))))))), diagram2=ApplyTransform(transform=Affine(1.0, 0.0, 7.0,\n",
- " 0.0, 1.0, 0.0), diagram=Primitive(shape=Spacer(width=1, height=0), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))))), diagram2=Empty()), diagram2=ApplyStyle(style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=0.0, dashing=None, output_size=None), diagram=ApplyStyle(style=Style(line_width=None, line_color=, fill_color=None, fill_opacity=None, dashing=None, output_size=None), diagram=Primitive(shape=Rectangle(width=9.6, height=6.6001199999999995, radius=0.5), style=Style(line_width=None, line_color=None, fill_color=None, fill_opacity=None, dashing=None, output_size=None), transform=Affine(1.0, 0.0, 0.0,\n",
- " 0.0, 1.0, 0.0))))), diagram2=Empty())))))), diagram2=Empty())))"
- ]
- },
- "execution_count": 8,
- "metadata": {},
- "output_type": "execute_result"
- }
- ],
- "source": [
- "def map_guard_test(cuda):\n",
- " def call(out, a, size) -> None:\n",
- " local_i = cuda.threadIdx.x\n",
- " # FILL ME IN (roughly 2 lines)\n",
- "\n",
- " return call\n",
- "\n",
- "\n",
- "SIZE = 4\n",
- "out = np.zeros((SIZE,))\n",
- "a = np.arange(SIZE)\n",
- "problem = CudaProblem(\n",
- " \"Guard\",\n",
- " map_guard_test,\n",
- " [a],\n",
- " out,\n",
- " [SIZE],\n",
- " threadsperblock=Coord(8, 1),\n",
- " spec=map_spec,\n",
- ")\n",
- "problem.show()"
- ]
- },
- {
- "cell_type": "code",
- "execution_count": null,
- "id": "2e34629f",
- "metadata": {
- "execution": {
- "iopub.execute_input": "2022-08-07T21:24:41.530990Z",
- "iopub.status.busy": "2022-08-07T21:24:41.530541Z",
- "iopub.status.idle": "2022-08-07T21:24:41.561067Z",
- "shell.execute_reply": "2022-08-07T21:24:41.560144Z"
- },
- "lines_to_end_of_cell_marker": 0,
- "lines_to_next_cell": 1,
- "id": "2e34629f",
- "outputId": "4624d5c3-436e-4a3b-d532-c67538ed40e0"
- },
- "outputs": [
- {
- "name": "stdout",
- "output_type": "stream",
- "text": [
- "Failed Tests.\n",
- "Yours: [0. 0. 0. 0.]\n",
- "Spec : [10 11 12 13]\n"
- ]
- }
- ],
- "source": [
- "problem.check()"
- ]
- },
- {
- "cell_type": "markdown",
- "id": "ec33ed08",
- "metadata": {
- "id": "ec33ed08"
- },
- "source": [
- "## Puzzle 4 - Map 2D\n",
- "\n",
- "Implement a kernel that adds 10 to each position of `a` and stores it in `out`.\n",
- "Input `a` is 2D and square. You have more threads than positions."
- ]
- },
- {
- "cell_type": "code",
- "execution_count": null,
- "id": "75d5eff5",
- "metadata": {
- "execution": {
- "iopub.execute_input": "2022-08-07T21:24:41.565190Z",
- "iopub.status.busy": "2022-08-07T21:24:41.564731Z",
- "iopub.status.idle": "2022-08-07T21:24:41.639943Z",
- "shell.execute_reply": "2022-08-07T21:24:41.639094Z"
- },
- "id": "75d5eff5",
- "outputId": "358b8e1c-5601-4a90-d1f0-2bfbbff26b9b"
- },
- "outputs": [
- {
- "name": "stdout",
- "output_type": "stream",
- "text": [
- "# Map 2D\n",
- " \n",
- " Score (Max Per Thread):\n",
- " | Global Reads | Global Writes | Shared Reads | Shared Writes |\n",
- " | 0 | 0 | 0 | 0 | \n",
- "\n"
- ]
- },
- {
- "data": {
- "image/svg+xml": [
- "\n",
- "