notebook.ipynb

{
  "nbformat": 4,
  "nbformat_minor": 0,
  "metadata": {
    "colab": {
      "name": "leandro3513mateus3489.ipynb",
      "provenance": [],
      "collapsed_sections": []
    },
    "kernelspec": {
      "name": "python3",
      "display_name": "Python 3"
    },
    "accelerator": "GPU"
  },
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "tpqcDD8Pd6hp",
        "colab_type": "text"
      },
      "source": [
        "**CCF​353 - ORGANIZAÇÃO DE COMPUTADORES II**\n",
        "\n",
        "**Bacharelado em Ciência da Computação**\n",
        "\n",
        "**Universidade Federal de Viçosa - Campus Florestal**\n",
        "\n",
        "**Prof. José Augusto​ M. Nacif - jnacif@ufv.br**\n",
        "\n",
        "# Atividades Práticas em GPU usando CUDA\n",
        "\n",
        "Este conteúdo foi baseado no [*colab*](https://colab.research.google.com/drive/1caBu4aCskJMyojYbU55a0aexxGZY0L_M) do Professor Ricardo Ferreira, da UFV - Campus Viçosa, e adaptado com exemplos do livro Professional CUDA C Programming. Ao longo dos exemplos, serão deixadas referências aos capítulos do livro que abordam cada assunto, para uma consulta mais detalhada.\n"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "MF40R8miyWKl",
        "colab_type": "text"
      },
      "source": [
        "## Introdução // Usando o Google Colab\n",
        "\n",
        "Esta plataforma é um serviço de nuvem gratuito hospedado pelo Google baseado no Jupyter Notebook, possuindo suporte ao Python 2.7 e 3.6. Além disso, você tem acesso grátis a uma GPU, que será nosso alvo de estudo.\n",
        "\n",
        "Depois de criar uma cópia deste colab para seu drive, você poderá editar as células de texto, editar as células de código e executá-las. Há suporte para código em Python, comandos bash, e qualquer outra linguagem, desde haja os plugins adequados, usando os comandos adequados.\n",
        "\n",
        "Para criar uma nova célula de texto ou código, basta aproximar o mouse do fim ou início de uma célula já existente, e você poderá escolher que tipo de célula criar.\n",
        "\n",
        "Para executar as células de código, basta clicar no botão 'play' no canto superior esquerdo dessas células. As células de texto têm suporte à Markdown e HTML."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "Hcx-5WlMTw_9",
        "colab_type": "text"
      },
      "source": [
        "## Introdução // Preparando o ambiente\n",
        "Primeiramente, certifique-se que este colab está configurado com uma GPU: acesse *Runtime > Change runtime type > Hardware accelerator* e selecione GPU.\n",
        "\n",
        "Usaremos um plugin desenvolvido por Andrei Nechaev para executar códigos do CUDA. Para isso, você deve executar os comandos a seguir, para baixar o plugin pelo git e depois ativá-lo com o comando `load_ext`.\n",
        "\n",
        "Observação: você deve executar estes comandos sempre que começar a usar este colab."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "ayHgoYeF91O7",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 218
        },
        "outputId": "f6db3efc-9a5a-4634-f9e9-25b7cc19c0ec"
      },
      "source": [
        "!pip install git+git://github.com/andreinechaev/nvcc4jupyter.git\n",
        "%load_ext nvcc_plugin"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Collecting git+git://github.com/andreinechaev/nvcc4jupyter.git\n",
            "  Cloning git://github.com/andreinechaev/nvcc4jupyter.git to /tmp/pip-req-build-uq53d1rl\n",
            "  Running command git clone -q git://github.com/andreinechaev/nvcc4jupyter.git /tmp/pip-req-build-uq53d1rl\n",
            "Building wheels for collected packages: NVCCPlugin\n",
            "  Building wheel for NVCCPlugin (setup.py) ... \u001b[?25l\u001b[?25hdone\n",
            "  Created wheel for NVCCPlugin: filename=NVCCPlugin-0.0.2-cp36-none-any.whl size=4307 sha256=35816bef59572300769d0d7496930fd47b09e57cc54101f6baaa2ad257428cd3\n",
            "  Stored in directory: /tmp/pip-ephem-wheel-cache-g9ttoh6n/wheels/10/c2/05/ca241da37bff77d60d31a9174f988109c61ba989e4d4650516\n",
            "Successfully built NVCCPlugin\n",
            "Installing collected packages: NVCCPlugin\n",
            "Successfully installed NVCCPlugin-0.0.2\n",
            "created output directory at /content/src\n",
            "Out bin /content/result.out\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "k0F90UTq-Uut",
        "colab_type": "text"
      },
      "source": [
        "Pronto! Agora você pode adicionar células de código CUDA, basta digitar `%%cu` na primeira linha do código."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "fj2exCIRiVJR",
        "colab_type": "text"
      },
      "source": [
        "## Introdução // Programação com CUDA\n",
        "\n",
        "_Referência: Capítulo 1 - Heterogeneous Parallel Computing with CUDA_\n",
        "\n",
        "É importante entender alguns termos e conceitos da programação com CUDA para seguir adequadamente os exemplos e tarefas.\n",
        "\n",
        "A ideia do CUDA é permitir ao programador trabalhar com uma arquitetura computacional heterogênea, que contém partes de código a ser executado na CPU e partes de código a ser executado na GPU.\n",
        "\n",
        "As GPU's não são substitutas das CPU's, muito pelo contrário, o melhor é tirar proveito das especialidades de cada uma, usando uma GPU em conjunto com uma CPU. Por esse motivo, a grosso modo, damos o nome de _host_ (hospedeiro) à CPU, e _device_ (dispositivo) à GPU, pois ela será um _dispositivo_ acelerando seu _hospedeiro_.\n",
        "\n",
        "![Diagrama de Arquitetura Heterogênea](https://nanxiao.me/wp-content/uploads/2016/12/Capture.jpg)\n",
        "\n",
        "Um programa do CUDA então é divido em duas partes principais, o _host code_ - executado na CPU - e o _device code_ - executado na GPU. O _device code_ é escrito com palavras reservadas específicas para demarcar funções com paralelismo de dados, e a estas é dado o nome de _kernels_.\n",
        "\n",
        "Uma métrica importante é a _capability_ de uma GPU, que mede seu poder computacional. Para simplificar, a NVIDIA criou uma numeração de _capability_ para as versões seus dispositivos: 1.0, 2.0, 3.0, 3.5, etc. A Tesla K40, por exemplo, possui _capability_ 3.5, outras GPU's da arquitetur Turing tem _capability_ 7.5. Todos exemplos mostrados serão compatíveis com _capability_ 2.0 ou superior.\n",
        "\n",
        "Outro conceito a se entender é o de blocos e threads. Cada bloco contém várias threads e cada thread executa um mesmo trecho de código. No exemplo a seguir será exibido o funcionamento das threads."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "19gMZsEpqdak",
        "colab_type": "text"
      },
      "source": [
        "## Exemplo 1 // Hello World!\n",
        "\n",
        "_Referência: Capítulo 1 - Heterogeneous Parallel Computing with CUDA_\n",
        "\n",
        "Vamos começar com um exemplo simples do livro base. Segue abaixo o código do \"Hello World\" modificado para imprimir o número do thread e o número do bloco de onde está sendo executado.\n",
        "\n",
        "Observação: a macro `CHECK` serve para verificar erros nas chamadas de funções do CUDA. Essa macro será utilizada apenas aqui, mas você pode copiá-la e usá-la para verificar erros em qualquer chamada à API do CUDA. Neste exemplo, a chamada `cudaDeviceReset()` é verificada."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "xoTswU9LsHU0",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 454
        },
        "outputId": "ea826f64-86c8-4ad3-c83f-11942d8bf460"
      },
      "source": [
        "%%cu\n",
        "#include <stdio.h>\n",
        "#include <stdlib.h>\n",
        "\n",
        "#define CHECK(call) { \\\n",
        "  const cudaError_t error = call; \\\n",
        "  if (error != cudaSuccess) { \\\n",
        "    printf(\"Error: %s:%d, \", __FILE__, __LINE__); \\\n",
        "    printf(\"code:%d, reason: %s\\n\", error, cudaGetErrorString(error)); \\\n",
        "    exit(1); \\\n",
        "  } \\\n",
        "}\n",
        "\n",
        "__global__ void helloFromGPU() {\n",
        "  printf(\"Hello World from GPU! %d %d\\n\", blockIdx.x, threadIdx.x);\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"Hello World from CPU!\\n\");\n",
        "  helloFromGPU<<<4, 6>>>();\n",
        "  CHECK(cudaDeviceReset());\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Hello World from CPU!\n",
            "Hello World from GPU! 3 0\n",
            "Hello World from GPU! 3 1\n",
            "Hello World from GPU! 3 2\n",
            "Hello World from GPU! 3 3\n",
            "Hello World from GPU! 3 4\n",
            "Hello World from GPU! 3 5\n",
            "Hello World from GPU! 0 0\n",
            "Hello World from GPU! 0 1\n",
            "Hello World from GPU! 0 2\n",
            "Hello World from GPU! 0 3\n",
            "Hello World from GPU! 0 4\n",
            "Hello World from GPU! 0 5\n",
            "Hello World from GPU! 1 0\n",
            "Hello World from GPU! 1 1\n",
            "Hello World from GPU! 1 2\n",
            "Hello World from GPU! 1 3\n",
            "Hello World from GPU! 1 4\n",
            "Hello World from GPU! 1 5\n",
            "Hello World from GPU! 2 0\n",
            "Hello World from GPU! 2 1\n",
            "Hello World from GPU! 2 2\n",
            "Hello World from GPU! 2 3\n",
            "Hello World from GPU! 2 4\n",
            "Hello World from GPU! 2 5\n",
            "\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "5xsjAETltDSD",
        "colab_type": "text"
      },
      "source": [
        "Podemos ver que este é um código em C com uma diferença nas chamadas de funções executadas pela GPU.\n",
        "\n",
        "No exemplo acima, a função é chamada com a sintaxe `<<<4, 6>>>`, ou seja, com 4 blocos, sendo que cada bloco tem 6 threads, o que fará com que sejam executadas 4 * 6 = 24 cópias da função. Cada cópia terá um ID único, composto pelo número do bloco (0, 1, 2 ou 3) e número do thread (0, 1, 2, 3, 4 ou 5).\n",
        "\n",
        "O especificador `__global__` na definição da função diz ao compilador que a função seráá chamada pela CPU e executada na GPU.\n",
        "\n",
        "A execução das várias instâncias da função `helloFromGPU` são em paralelo com a CPU. No exemplo, a CPU irá aguardar pois chama a função `cudaDeviceReset`. Porém qualquer código que for inserido entre a chamada `cudaDeviceReset` e `helloFromGPU` será executado em paralelo.\n",
        "\n",
        "A função `cudaDeviceReset()` vai explicitamente destruir e limpar todos os recursos associados ao dispositivo usado no processo corrente."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "d9Tkuc4NUyPX",
        "colab_type": "text"
      },
      "source": [
        "Se, ao invés de executar diretamente o código, você quiser salvá-lo, use o comando `%%writefile caminho/para/o/arquivo`. Por padrão, o plugin cria a pasta `/content/src/` para colocar os códigos. Você pode então substituir o comando `%%cu` por `%%writefile /content/src/hello.cu`, por exemplo.\n",
        "\n",
        "Isso é útil se você quiser especificar opções de compilação ou extrair informações sobre a execução do seu programa - o que fará parte dos exercícios. Por exemplo, após salvar o exemplo anterior com o comando `writefile` mostrado, você poderia compilar da seguinte forma:"
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "colab_type": "code",
        "id": "-esyL1wdklmn",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 34
        },
        "outputId": "1f2f901a-ab44-42c9-8322-7506ccb93c02"
      },
      "source": [
        "%%writefile /content/src/hello.cu\n",
        "#include <stdio.h>\n",
        "#include <stdlib.h>\n",
        "\n",
        "#define CHECK(call) { \\\n",
        "  const cudaError_t error = call; \\\n",
        "  if (error != cudaSuccess) { \\\n",
        "    printf(\"Error: %s:%d, \", __FILE__, __LINE__); \\\n",
        "    printf(\"code:%d, reason: %s\\n\", error, cudaGetErrorString(error)); \\\n",
        "    exit(1); \\\n",
        "  } \\\n",
        "}\n",
        "\n",
        "__global__ void helloFromGPU() {\n",
        "  printf(\"Hello World from GPU! %d %d\\n\", blockIdx.x, threadIdx.x);\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"Hello World from CPU!\\n\");\n",
        "  helloFromGPU<<<4, 6>>>();\n",
        "  CHECK(cudaDeviceReset());\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Overwriting /content/src/hello.cu\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "YPEFfKGjWzM1",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 454
        },
        "outputId": "8cf88e52-23ff-4a9f-8498-4c3de4f08d6c"
      },
      "source": [
        "%cd /content/src\n",
        "!nvcc -O2 -arch=sm_35 -o hello hello.cu\n",
        "!./hello"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "Hello World from CPU!\n",
            "Hello World from GPU! 3 0\n",
            "Hello World from GPU! 3 1\n",
            "Hello World from GPU! 3 2\n",
            "Hello World from GPU! 3 3\n",
            "Hello World from GPU! 3 4\n",
            "Hello World from GPU! 3 5\n",
            "Hello World from GPU! 0 0\n",
            "Hello World from GPU! 0 1\n",
            "Hello World from GPU! 0 2\n",
            "Hello World from GPU! 0 3\n",
            "Hello World from GPU! 0 4\n",
            "Hello World from GPU! 0 5\n",
            "Hello World from GPU! 1 0\n",
            "Hello World from GPU! 1 1\n",
            "Hello World from GPU! 1 2\n",
            "Hello World from GPU! 1 3\n",
            "Hello World from GPU! 1 4\n",
            "Hello World from GPU! 1 5\n",
            "Hello World from GPU! 2 0\n",
            "Hello World from GPU! 2 1\n",
            "Hello World from GPU! 2 2\n",
            "Hello World from GPU! 2 3\n",
            "Hello World from GPU! 2 4\n",
            "Hello World from GPU! 2 5\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "3Wy4qJnIVAmJ",
        "colab_type": "text"
      },
      "source": [
        "Com a opção `-arch`, você pode selecionar uma arquitetura. Usando uma arquitetura com _capability_ maior ou igual à 2.0, você pode usar o `printf` dentro do código da GPU - que chamammos de kernel. No comando de exemplo, foi usada a capability 3.5. Porém, como a GPU pode executar muitos threads - na ordem dos bilhões -, para exemplos maiores, é melhor não usá-lo, exceto se for condicionalmente para verificar apenas alguns threads. Por exemplo:"
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "hCxsmB9VV2P3",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 67
        },
        "outputId": "0778d731-8ddf-456f-bf06-55e921036c40"
      },
      "source": [
        "%%cu\n",
        "#include <stdio.h>\n",
        "#include <stdlib.h>\n",
        "\n",
        "__global__ void helloFromGPU() {\n",
        "  if (blockIdx.x == 2048 && threadIdx.x == 512)\n",
        "    printf(\"Hello World from GPU! Block X: %d, Thread X: %d\\n\",\n",
        "           blockIdx.x, threadIdx.x);\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"Hello World from CPU!\\n\");\n",
        "  helloFromGPU<<<64 * 1024 * 1024, 1024>>>();\n",
        "  cudaDeviceReset();\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Hello World from CPU!\n",
            "Hello World from GPU! Block X: 2048, Thread X: 512\n",
            "\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "8wQ7bhjqabq7",
        "colab_type": "text"
      },
      "source": [
        "O trecho acima irá executar com 64 M (2^26) blocos com 1024 threads em cada bloco, ou seja, serão 64 G (2^36) execuções da função `helloFromGPU`. Ao longo de todo o texto, os prefixos K, M, G, etc. sempre vão se referir aos seus valores em potência de 2. Ou seja, K = 2^10, M = 2^20, G = 2^30, etc."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "ZCQoUXTea-67",
        "colab_type": "text"
      },
      "source": [
        "## Tarefa 1 // Hello World!\n",
        "\n",
        "1. Modifique o exemplo para executar 1 bloco com 1024 threads. Imprima e verifique a ordem de execução dos threads. É sequencial? É determinística? É agrupada?\n"
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "lBZ0-m1FbNLl",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 1000
        },
        "outputId": "a802c5ca-22fb-4334-930d-3a5c0a333382"
      },
      "source": [
        "%%cu\n",
        "#include <stdio.h>\n",
        "#include <stdlib.h>\n",
        "\n",
        "#define CHECK(call) { \\\n",
        "  const cudaError_t error = call; \\\n",
        "  if (error != cudaSuccess) { \\\n",
        "    printf(\"Error: %s:%d, \", __FILE__, __LINE__); \\\n",
        "    printf(\"code:%d, reason: %s\\n\", error, cudaGetErrorString(error)); \\\n",
        "    exit(1); \\\n",
        "  } \\\n",
        "}\n",
        "\n",
        "__global__ void helloFromGPU() {\n",
        "  printf(\"Hello World from GPU! %d %d\\n\", blockIdx.x, threadIdx.x);\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"Hello World from CPU!\\n\");\n",
        "  helloFromGPU<<<1, 1024>>>();\n",
        "  CHECK(cudaDeviceReset());\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Hello World from CPU!\n",
            "Hello World from GPU! 0 192\n",
            "Hello World from GPU! 0 193\n",
            "Hello World from GPU! 0 194\n",
            "Hello World from GPU! 0 195\n",
            "Hello World from GPU! 0 196\n",
            "Hello World from GPU! 0 197\n",
            "Hello World from GPU! 0 198\n",
            "Hello World from GPU! 0 199\n",
            "Hello World from GPU! 0 200\n",
            "Hello World from GPU! 0 201\n",
            "Hello World from GPU! 0 202\n",
            "Hello World from GPU! 0 203\n",
            "Hello World from GPU! 0 204\n",
            "Hello World from GPU! 0 205\n",
            "Hello World from GPU! 0 206\n",
            "Hello World from GPU! 0 207\n",
            "Hello World from GPU! 0 208\n",
            "Hello World from GPU! 0 209\n",
            "Hello World from GPU! 0 210\n",
            "Hello World from GPU! 0 211\n",
            "Hello World from GPU! 0 212\n",
            "Hello World from GPU! 0 213\n",
            "Hello World from GPU! 0 214\n",
            "Hello World from GPU! 0 215\n",
            "Hello World from GPU! 0 216\n",
            "Hello World from GPU! 0 217\n",
            "Hello World from GPU! 0 218\n",
            "Hello World from GPU! 0 219\n",
            "Hello World from GPU! 0 220\n",
            "Hello World from GPU! 0 221\n",
            "Hello World from GPU! 0 222\n",
            "Hello World from GPU! 0 223\n",
            "Hello World from GPU! 0 0\n",
            "Hello World from GPU! 0 1\n",
            "Hello World from GPU! 0 2\n",
            "Hello World from GPU! 0 3\n",
            "Hello World from GPU! 0 4\n",
            "Hello World from GPU! 0 5\n",
            "Hello World from GPU! 0 6\n",
            "Hello World from GPU! 0 7\n",
            "Hello World from GPU! 0 8\n",
            "Hello World from GPU! 0 9\n",
            "Hello World from GPU! 0 10\n",
            "Hello World from GPU! 0 11\n",
            "Hello World from GPU! 0 12\n",
            "Hello World from GPU! 0 13\n",
            "Hello World from GPU! 0 14\n",
            "Hello World from GPU! 0 15\n",
            "Hello World from GPU! 0 16\n",
            "Hello World from GPU! 0 17\n",
            "Hello World from GPU! 0 18\n",
            "Hello World from GPU! 0 19\n",
            "Hello World from GPU! 0 20\n",
            "Hello World from GPU! 0 21\n",
            "Hello World from GPU! 0 22\n",
            "Hello World from GPU! 0 23\n",
            "Hello World from GPU! 0 24\n",
            "Hello World from GPU! 0 25\n",
            "Hello World from GPU! 0 26\n",
            "Hello World from GPU! 0 27\n",
            "Hello World from GPU! 0 28\n",
            "Hello World from GPU! 0 29\n",
            "Hello World from GPU! 0 30\n",
            "Hello World from GPU! 0 31\n",
            "Hello World from GPU! 0 128\n",
            "Hello World from GPU! 0 129\n",
            "Hello World from GPU! 0 130\n",
            "Hello World from GPU! 0 131\n",
            "Hello World from GPU! 0 132\n",
            "Hello World from GPU! 0 133\n",
            "Hello World from GPU! 0 134\n",
            "Hello World from GPU! 0 135\n",
            "Hello World from GPU! 0 136\n",
            "Hello World from GPU! 0 137\n",
            "Hello World from GPU! 0 138\n",
            "Hello World from GPU! 0 139\n",
            "Hello World from GPU! 0 140\n",
            "Hello World from GPU! 0 141\n",
            "Hello World from GPU! 0 142\n",
            "Hello World from GPU! 0 143\n",
            "Hello World from GPU! 0 144\n",
            "Hello World from GPU! 0 145\n",
            "Hello World from GPU! 0 146\n",
            "Hello World from GPU! 0 147\n",
            "Hello World from GPU! 0 148\n",
            "Hello World from GPU! 0 149\n",
            "Hello World from GPU! 0 150\n",
            "Hello World from GPU! 0 151\n",
            "Hello World from GPU! 0 152\n",
            "Hello World from GPU! 0 153\n",
            "Hello World from GPU! 0 154\n",
            "Hello World from GPU! 0 155\n",
            "Hello World from GPU! 0 156\n",
            "Hello World from GPU! 0 157\n",
            "Hello World from GPU! 0 158\n",
            "Hello World from GPU! 0 159\n",
            "Hello World from GPU! 0 960\n",
            "Hello World from GPU! 0 961\n",
            "Hello World from GPU! 0 962\n",
            "Hello World from GPU! 0 963\n",
            "Hello World from GPU! 0 964\n",
            "Hello World from GPU! 0 965\n",
            "Hello World from GPU! 0 966\n",
            "Hello World from GPU! 0 967\n",
            "Hello World from GPU! 0 968\n",
            "Hello World from GPU! 0 969\n",
            "Hello World from GPU! 0 970\n",
            "Hello World from GPU! 0 971\n",
            "Hello World from GPU! 0 972\n",
            "Hello World from GPU! 0 973\n",
            "Hello World from GPU! 0 974\n",
            "Hello World from GPU! 0 975\n",
            "Hello World from GPU! 0 976\n",
            "Hello World from GPU! 0 977\n",
            "Hello World from GPU! 0 978\n",
            "Hello World from GPU! 0 979\n",
            "Hello World from GPU! 0 980\n",
            "Hello World from GPU! 0 981\n",
            "Hello World from GPU! 0 982\n",
            "Hello World from GPU! 0 983\n",
            "Hello World from GPU! 0 984\n",
            "Hello World from GPU! 0 985\n",
            "Hello World from GPU! 0 986\n",
            "Hello World from GPU! 0 987\n",
            "Hello World from GPU! 0 988\n",
            "Hello World from GPU! 0 989\n",
            "Hello World from GPU! 0 990\n",
            "Hello World from GPU! 0 991\n",
            "Hello World from GPU! 0 256\n",
            "Hello World from GPU! 0 257\n",
            "Hello World from GPU! 0 258\n",
            "Hello World from GPU! 0 259\n",
            "Hello World from GPU! 0 260\n",
            "Hello World from GPU! 0 261\n",
            "Hello World from GPU! 0 262\n",
            "Hello World from GPU! 0 263\n",
            "Hello World from GPU! 0 264\n",
            "Hello World from GPU! 0 265\n",
            "Hello World from GPU! 0 266\n",
            "Hello World from GPU! 0 267\n",
            "Hello World from GPU! 0 268\n",
            "Hello World from GPU! 0 269\n",
            "Hello World from GPU! 0 270\n",
            "Hello World from GPU! 0 271\n",
            "Hello World from GPU! 0 272\n",
            "Hello World from GPU! 0 273\n",
            "Hello World from GPU! 0 274\n",
            "Hello World from GPU! 0 275\n",
            "Hello World from GPU! 0 276\n",
            "Hello World from GPU! 0 277\n",
            "Hello World from GPU! 0 278\n",
            "Hello World from GPU! 0 279\n",
            "Hello World from GPU! 0 280\n",
            "Hello World from GPU! 0 281\n",
            "Hello World from GPU! 0 282\n",
            "Hello World from GPU! 0 283\n",
            "Hello World from GPU! 0 284\n",
            "Hello World from GPU! 0 285\n",
            "Hello World from GPU! 0 286\n",
            "Hello World from GPU! 0 287\n",
            "Hello World from GPU! 0 352\n",
            "Hello World from GPU! 0 353\n",
            "Hello World from GPU! 0 354\n",
            "Hello World from GPU! 0 355\n",
            "Hello World from GPU! 0 356\n",
            "Hello World from GPU! 0 357\n",
            "Hello World from GPU! 0 358\n",
            "Hello World from GPU! 0 359\n",
            "Hello World from GPU! 0 360\n",
            "Hello World from GPU! 0 361\n",
            "Hello World from GPU! 0 362\n",
            "Hello World from GPU! 0 363\n",
            "Hello World from GPU! 0 364\n",
            "Hello World from GPU! 0 365\n",
            "Hello World from GPU! 0 366\n",
            "Hello World from GPU! 0 367\n",
            "Hello World from GPU! 0 368\n",
            "Hello World from GPU! 0 369\n",
            "Hello World from GPU! 0 370\n",
            "Hello World from GPU! 0 371\n",
            "Hello World from GPU! 0 372\n",
            "Hello World from GPU! 0 373\n",
            "Hello World from GPU! 0 374\n",
            "Hello World from GPU! 0 375\n",
            "Hello World from GPU! 0 376\n",
            "Hello World from GPU! 0 377\n",
            "Hello World from GPU! 0 378\n",
            "Hello World from GPU! 0 379\n",
            "Hello World from GPU! 0 380\n",
            "Hello World from GPU! 0 381\n",
            "Hello World from GPU! 0 382\n",
            "Hello World from GPU! 0 383\n",
            "Hello World from GPU! 0 480\n",
            "Hello World from GPU! 0 481\n",
            "Hello World from GPU! 0 482\n",
            "Hello World from GPU! 0 483\n",
            "Hello World from GPU! 0 484\n",
            "Hello World from GPU! 0 485\n",
            "Hello World from GPU! 0 486\n",
            "Hello World from GPU! 0 487\n",
            "Hello World from GPU! 0 488\n",
            "Hello World from GPU! 0 489\n",
            "Hello World from GPU! 0 490\n",
            "Hello World from GPU! 0 491\n",
            "Hello World from GPU! 0 492\n",
            "Hello World from GPU! 0 493\n",
            "Hello World from GPU! 0 494\n",
            "Hello World from GPU! 0 495\n",
            "Hello World from GPU! 0 496\n",
            "Hello World from GPU! 0 497\n",
            "Hello World from GPU! 0 498\n",
            "Hello World from GPU! 0 499\n",
            "Hello World from GPU! 0 500\n",
            "Hello World from GPU! 0 501\n",
            "Hello World from GPU! 0 502\n",
            "Hello World from GPU! 0 503\n",
            "Hello World from GPU! 0 504\n",
            "Hello World from GPU! 0 505\n",
            "Hello World from GPU! 0 506\n",
            "Hello World from GPU! 0 507\n",
            "Hello World from GPU! 0 508\n",
            "Hello World from GPU! 0 509\n",
            "Hello World from GPU! 0 510\n",
            "Hello World from GPU! 0 511\n",
            "Hello World from GPU! 0 416\n",
            "Hello World from GPU! 0 417\n",
            "Hello World from GPU! 0 418\n",
            "Hello World from GPU! 0 419\n",
            "Hello World from GPU! 0 420\n",
            "Hello World from GPU! 0 421\n",
            "Hello World from GPU! 0 422\n",
            "Hello World from GPU! 0 423\n",
            "Hello World from GPU! 0 424\n",
            "Hello World from GPU! 0 425\n",
            "Hello World from GPU! 0 426\n",
            "Hello World from GPU! 0 427\n",
            "Hello World from GPU! 0 428\n",
            "Hello World from GPU! 0 429\n",
            "Hello World from GPU! 0 430\n",
            "Hello World from GPU! 0 431\n",
            "Hello World from GPU! 0 432\n",
            "Hello World from GPU! 0 433\n",
            "Hello World from GPU! 0 434\n",
            "Hello World from GPU! 0 435\n",
            "Hello World from GPU! 0 436\n",
            "Hello World from GPU! 0 437\n",
            "Hello World from GPU! 0 438\n",
            "Hello World from GPU! 0 439\n",
            "Hello World from GPU! 0 440\n",
            "Hello World from GPU! 0 441\n",
            "Hello World from GPU! 0 442\n",
            "Hello World from GPU! 0 443\n",
            "Hello World from GPU! 0 444\n",
            "Hello World from GPU! 0 445\n",
            "Hello World from GPU! 0 446\n",
            "Hello World from GPU! 0 447\n",
            "Hello World from GPU! 0 160\n",
            "Hello World from GPU! 0 161\n",
            "Hello World from GPU! 0 162\n",
            "Hello World from GPU! 0 163\n",
            "Hello World from GPU! 0 164\n",
            "Hello World from GPU! 0 165\n",
            "Hello World from GPU! 0 166\n",
            "Hello World from GPU! 0 167\n",
            "Hello World from GPU! 0 168\n",
            "Hello World from GPU! 0 169\n",
            "Hello World from GPU! 0 170\n",
            "Hello World from GPU! 0 171\n",
            "Hello World from GPU! 0 172\n",
            "Hello World from GPU! 0 173\n",
            "Hello World from GPU! 0 174\n",
            "Hello World from GPU! 0 175\n",
            "Hello World from GPU! 0 176\n",
            "Hello World from GPU! 0 177\n",
            "Hello World from GPU! 0 178\n",
            "Hello World from GPU! 0 179\n",
            "Hello World from GPU! 0 180\n",
            "Hello World from GPU! 0 181\n",
            "Hello World from GPU! 0 182\n",
            "Hello World from GPU! 0 183\n",
            "Hello World from GPU! 0 184\n",
            "Hello World from GPU! 0 185\n",
            "Hello World from GPU! 0 186\n",
            "Hello World from GPU! 0 187\n",
            "Hello World from GPU! 0 188\n",
            "Hello World from GPU! 0 189\n",
            "Hello World from GPU! 0 190\n",
            "Hello World from GPU! 0 191\n",
            "Hello World from GPU! 0 864\n",
            "Hello World from GPU! 0 865\n",
            "Hello World from GPU! 0 866\n",
            "Hello World from GPU! 0 867\n",
            "Hello World from GPU! 0 868\n",
            "Hello World from GPU! 0 869\n",
            "Hello World from GPU! 0 870\n",
            "Hello World from GPU! 0 871\n",
            "Hello World from GPU! 0 872\n",
            "Hello World from GPU! 0 873\n",
            "Hello World from GPU! 0 874\n",
            "Hello World from GPU! 0 875\n",
            "Hello World from GPU! 0 876\n",
            "Hello World from GPU! 0 877\n",
            "Hello World from GPU! 0 878\n",
            "Hello World from GPU! 0 879\n",
            "Hello World from GPU! 0 880\n",
            "Hello World from GPU! 0 881\n",
            "Hello World from GPU! 0 882\n",
            "Hello World from GPU! 0 883\n",
            "Hello World from GPU! 0 884\n",
            "Hello World from GPU! 0 885\n",
            "Hello World from GPU! 0 886\n",
            "Hello World from GPU! 0 887\n",
            "Hello World from GPU! 0 888\n",
            "Hello World from GPU! 0 889\n",
            "Hello World from GPU! 0 890\n",
            "Hello World from GPU! 0 891\n",
            "Hello World from GPU! 0 892\n",
            "Hello World from GPU! 0 893\n",
            "Hello World from GPU! 0 894\n",
            "Hello World from GPU! 0 895\n",
            "Hello World from GPU! 0 576\n",
            "Hello World from GPU! 0 577\n",
            "Hello World from GPU! 0 578\n",
            "Hello World from GPU! 0 579\n",
            "Hello World from GPU! 0 580\n",
            "Hello World from GPU! 0 581\n",
            "Hello World from GPU! 0 582\n",
            "Hello World from GPU! 0 583\n",
            "Hello World from GPU! 0 584\n",
            "Hello World from GPU! 0 585\n",
            "Hello World from GPU! 0 586\n",
            "Hello World from GPU! 0 587\n",
            "Hello World from GPU! 0 588\n",
            "Hello World from GPU! 0 589\n",
            "Hello World from GPU! 0 590\n",
            "Hello World from GPU! 0 591\n",
            "Hello World from GPU! 0 592\n",
            "Hello World from GPU! 0 593\n",
            "Hello World from GPU! 0 594\n",
            "Hello World from GPU! 0 595\n",
            "Hello World from GPU! 0 596\n",
            "Hello World from GPU! 0 597\n",
            "Hello World from GPU! 0 598\n",
            "Hello World from GPU! 0 599\n",
            "Hello World from GPU! 0 600\n",
            "Hello World from GPU! 0 601\n",
            "Hello World from GPU! 0 602\n",
            "Hello World from GPU! 0 603\n",
            "Hello World from GPU! 0 604\n",
            "Hello World from GPU! 0 605\n",
            "Hello World from GPU! 0 606\n",
            "Hello World from GPU! 0 607\n",
            "Hello World from GPU! 0 608\n",
            "Hello World from GPU! 0 609\n",
            "Hello World from GPU! 0 610\n",
            "Hello World from GPU! 0 611\n",
            "Hello World from GPU! 0 612\n",
            "Hello World from GPU! 0 613\n",
            "Hello World from GPU! 0 614\n",
            "Hello World from GPU! 0 615\n",
            "Hello World from GPU! 0 616\n",
            "Hello World from GPU! 0 617\n",
            "Hello World from GPU! 0 618\n",
            "Hello World from GPU! 0 619\n",
            "Hello World from GPU! 0 620\n",
            "Hello World from GPU! 0 621\n",
            "Hello World from GPU! 0 622\n",
            "Hello World from GPU! 0 623\n",
            "Hello World from GPU! 0 624\n",
            "Hello World from GPU! 0 625\n",
            "Hello World from GPU! 0 626\n",
            "Hello World from GPU! 0 627\n",
            "Hello World from GPU! 0 628\n",
            "Hello World from GPU! 0 629\n",
            "Hello World from GPU! 0 630\n",
            "Hello World from GPU! 0 631\n",
            "Hello World from GPU! 0 632\n",
            "Hello World from GPU! 0 633\n",
            "Hello World from GPU! 0 634\n",
            "Hello World from GPU! 0 635\n",
            "Hello World from GPU! 0 636\n",
            "Hello World from GPU! 0 637\n",
            "Hello World from GPU! 0 638\n",
            "Hello World from GPU! 0 639\n",
            "Hello World from GPU! 0 288\n",
            "Hello World from GPU! 0 289\n",
            "Hello World from GPU! 0 290\n",
            "Hello World from GPU! 0 291\n",
            "Hello World from GPU! 0 292\n",
            "Hello World from GPU! 0 293\n",
            "Hello World from GPU! 0 294\n",
            "Hello World from GPU! 0 295\n",
            "Hello World from GPU! 0 296\n",
            "Hello World from GPU! 0 297\n",
            "Hello World from GPU! 0 298\n",
            "Hello World from GPU! 0 299\n",
            "Hello World from GPU! 0 300\n",
            "Hello World from GPU! 0 301\n",
            "Hello World from GPU! 0 302\n",
            "Hello World from GPU! 0 303\n",
            "Hello World from GPU! 0 304\n",
            "Hello World from GPU! 0 305\n",
            "Hello World from GPU! 0 306\n",
            "Hello World from GPU! 0 307\n",
            "Hello World from GPU! 0 308\n",
            "Hello World from GPU! 0 309\n",
            "Hello World from GPU! 0 310\n",
            "Hello World from GPU! 0 311\n",
            "Hello World from GPU! 0 312\n",
            "Hello World from GPU! 0 313\n",
            "Hello World from GPU! 0 314\n",
            "Hello World from GPU! 0 315\n",
            "Hello World from GPU! 0 316\n",
            "Hello World from GPU! 0 317\n",
            "Hello World from GPU! 0 318\n",
            "Hello World from GPU! 0 319\n",
            "Hello World from GPU! 0 928\n",
            "Hello World from GPU! 0 929\n",
            "Hello World from GPU! 0 930\n",
            "Hello World from GPU! 0 931\n",
            "Hello World from GPU! 0 932\n",
            "Hello World from GPU! 0 933\n",
            "Hello World from GPU! 0 934\n",
            "Hello World from GPU! 0 935\n",
            "Hello World from GPU! 0 936\n",
            "Hello World from GPU! 0 937\n",
            "Hello World from GPU! 0 938\n",
            "Hello World from GPU! 0 939\n",
            "Hello World from GPU! 0 940\n",
            "Hello World from GPU! 0 941\n",
            "Hello World from GPU! 0 942\n",
            "Hello World from GPU! 0 943\n",
            "Hello World from GPU! 0 944\n",
            "Hello World from GPU! 0 945\n",
            "Hello World from GPU! 0 946\n",
            "Hello World from GPU! 0 947\n",
            "Hello World from GPU! 0 948\n",
            "Hello World from GPU! 0 949\n",
            "Hello World from GPU! 0 950\n",
            "Hello World from GPU! 0 951\n",
            "Hello World from GPU! 0 952\n",
            "Hello World from GPU! 0 953\n",
            "Hello World from GPU! 0 954\n",
            "Hello World from GPU! 0 955\n",
            "Hello World from GPU! 0 956\n",
            "Hello World from GPU! 0 957\n",
            "Hello World from GPU! 0 958\n",
            "Hello World from GPU! 0 959\n",
            "Hello World from GPU! 0 512\n",
            "Hello World from GPU! 0 513\n",
            "Hello World from GPU! 0 514\n",
            "Hello World from GPU! 0 515\n",
            "Hello World from GPU! 0 516\n",
            "Hello World from GPU! 0 517\n",
            "Hello World from GPU! 0 518\n",
            "Hello World from GPU! 0 519\n",
            "Hello World from GPU! 0 520\n",
            "Hello World from GPU! 0 521\n",
            "Hello World from GPU! 0 522\n",
            "Hello World from GPU! 0 523\n",
            "Hello World from GPU! 0 524\n",
            "Hello World from GPU! 0 525\n",
            "Hello World from GPU! 0 526\n",
            "Hello World from GPU! 0 527\n",
            "Hello World from GPU! 0 528\n",
            "Hello World from GPU! 0 529\n",
            "Hello World from GPU! 0 530\n",
            "Hello World from GPU! 0 531\n",
            "Hello World from GPU! 0 532\n",
            "Hello World from GPU! 0 533\n",
            "Hello World from GPU! 0 534\n",
            "Hello World from GPU! 0 535\n",
            "Hello World from GPU! 0 536\n",
            "Hello World from GPU! 0 537\n",
            "Hello World from GPU! 0 538\n",
            "Hello World from GPU! 0 539\n",
            "Hello World from GPU! 0 540\n",
            "Hello World from GPU! 0 541\n",
            "Hello World from GPU! 0 542\n",
            "Hello World from GPU! 0 543\n",
            "Hello World from GPU! 0 640\n",
            "Hello World from GPU! 0 641\n",
            "Hello World from GPU! 0 642\n",
            "Hello World from GPU! 0 643\n",
            "Hello World from GPU! 0 644\n",
            "Hello World from GPU! 0 645\n",
            "Hello World from GPU! 0 646\n",
            "Hello World from GPU! 0 647\n",
            "Hello World from GPU! 0 648\n",
            "Hello World from GPU! 0 649\n",
            "Hello World from GPU! 0 650\n",
            "Hello World from GPU! 0 651\n",
            "Hello World from GPU! 0 652\n",
            "Hello World from GPU! 0 653\n",
            "Hello World from GPU! 0 654\n",
            "Hello World from GPU! 0 655\n",
            "Hello World from GPU! 0 656\n",
            "Hello World from GPU! 0 657\n",
            "Hello World from GPU! 0 658\n",
            "Hello World from GPU! 0 659\n",
            "Hello World from GPU! 0 660\n",
            "Hello World from GPU! 0 661\n",
            "Hello World from GPU! 0 662\n",
            "Hello World from GPU! 0 663\n",
            "Hello World from GPU! 0 664\n",
            "Hello World from GPU! 0 665\n",
            "Hello World from GPU! 0 666\n",
            "Hello World from GPU! 0 667\n",
            "Hello World from GPU! 0 668\n",
            "Hello World from GPU! 0 669\n",
            "Hello World from GPU! 0 670\n",
            "Hello World from GPU! 0 671\n",
            "Hello World from GPU! 0 896\n",
            "Hello World from GPU! 0 897\n",
            "Hello World from GPU! 0 898\n",
            "Hello World from GPU! 0 899\n",
            "Hello World from GPU! 0 900\n",
            "Hello World from GPU! 0 901\n",
            "Hello World from GPU! 0 902\n",
            "Hello World from GPU! 0 903\n",
            "Hello World from GPU! 0 904\n",
            "Hello World from GPU! 0 905\n",
            "Hello World from GPU! 0 906\n",
            "Hello World from GPU! 0 907\n",
            "Hello World from GPU! 0 908\n",
            "Hello World from GPU! 0 909\n",
            "Hello World from GPU! 0 910\n",
            "Hello World from GPU! 0 911\n",
            "Hello World from GPU! 0 912\n",
            "Hello World from GPU! 0 913\n",
            "Hello World from GPU! 0 914\n",
            "Hello World from GPU! 0 915\n",
            "Hello World from GPU! 0 916\n",
            "Hello World from GPU! 0 917\n",
            "Hello World from GPU! 0 918\n",
            "Hello World from GPU! 0 919\n",
            "Hello World from GPU! 0 920\n",
            "Hello World from GPU! 0 921\n",
            "Hello World from GPU! 0 922\n",
            "Hello World from GPU! 0 923\n",
            "Hello World from GPU! 0 924\n",
            "Hello World from GPU! 0 925\n",
            "Hello World from GPU! 0 926\n",
            "Hello World from GPU! 0 927\n",
            "Hello World from GPU! 0 448\n",
            "Hello World from GPU! 0 449\n",
            "Hello World from GPU! 0 450\n",
            "Hello World from GPU! 0 451\n",
            "Hello World from GPU! 0 452\n",
            "Hello World from GPU! 0 453\n",
            "Hello World from GPU! 0 454\n",
            "Hello World from GPU! 0 455\n",
            "Hello World from GPU! 0 456\n",
            "Hello World from GPU! 0 457\n",
            "Hello World from GPU! 0 458\n",
            "Hello World from GPU! 0 459\n",
            "Hello World from GPU! 0 460\n",
            "Hello World from GPU! 0 461\n",
            "Hello World from GPU! 0 462\n",
            "Hello World from GPU! 0 463\n",
            "Hello World from GPU! 0 464\n",
            "Hello World from GPU! 0 465\n",
            "Hello World from GPU! 0 466\n",
            "Hello World from GPU! 0 467\n",
            "Hello World from GPU! 0 468\n",
            "Hello World from GPU! 0 469\n",
            "Hello World from GPU! 0 470\n",
            "Hello World from GPU! 0 471\n",
            "Hello World from GPU! 0 472\n",
            "Hello World from GPU! 0 473\n",
            "Hello World from GPU! 0 474\n",
            "Hello World from GPU! 0 475\n",
            "Hello World from GPU! 0 476\n",
            "Hello World from GPU! 0 477\n",
            "Hello World from GPU! 0 478\n",
            "Hello World from GPU! 0 479\n",
            "Hello World from GPU! 0 384\n",
            "Hello World from GPU! 0 385\n",
            "Hello World from GPU! 0 386\n",
            "Hello World from GPU! 0 387\n",
            "Hello World from GPU! 0 388\n",
            "Hello World from GPU! 0 389\n",
            "Hello World from GPU! 0 390\n",
            "Hello World from GPU! 0 391\n",
            "Hello World from GPU! 0 392\n",
            "Hello World from GPU! 0 393\n",
            "Hello World from GPU! 0 394\n",
            "Hello World from GPU! 0 395\n",
            "Hello World from GPU! 0 396\n",
            "Hello World from GPU! 0 397\n",
            "Hello World from GPU! 0 398\n",
            "Hello World from GPU! 0 399\n",
            "Hello World from GPU! 0 400\n",
            "Hello World from GPU! 0 401\n",
            "Hello World from GPU! 0 402\n",
            "Hello World from GPU! 0 403\n",
            "Hello World from GPU! 0 404\n",
            "Hello World from GPU! 0 405\n",
            "Hello World from GPU! 0 406\n",
            "Hello World from GPU! 0 407\n",
            "Hello World from GPU! 0 408\n",
            "Hello World from GPU! 0 409\n",
            "Hello World from GPU! 0 410\n",
            "Hello World from GPU! 0 411\n",
            "Hello World from GPU! 0 412\n",
            "Hello World from GPU! 0 413\n",
            "Hello World from GPU! 0 414\n",
            "Hello World from GPU! 0 415\n",
            "Hello World from GPU! 0 64\n",
            "Hello World from GPU! 0 65\n",
            "Hello World from GPU! 0 66\n",
            "Hello World from GPU! 0 67\n",
            "Hello World from GPU! 0 68\n",
            "Hello World from GPU! 0 69\n",
            "Hello World from GPU! 0 70\n",
            "Hello World from GPU! 0 71\n",
            "Hello World from GPU! 0 72\n",
            "Hello World from GPU! 0 73\n",
            "Hello World from GPU! 0 74\n",
            "Hello World from GPU! 0 75\n",
            "Hello World from GPU! 0 76\n",
            "Hello World from GPU! 0 77\n",
            "Hello World from GPU! 0 78\n",
            "Hello World from GPU! 0 79\n",
            "Hello World from GPU! 0 80\n",
            "Hello World from GPU! 0 81\n",
            "Hello World from GPU! 0 82\n",
            "Hello World from GPU! 0 83\n",
            "Hello World from GPU! 0 84\n",
            "Hello World from GPU! 0 85\n",
            "Hello World from GPU! 0 86\n",
            "Hello World from GPU! 0 87\n",
            "Hello World from GPU! 0 88\n",
            "Hello World from GPU! 0 89\n",
            "Hello World from GPU! 0 90\n",
            "Hello World from GPU! 0 91\n",
            "Hello World from GPU! 0 92\n",
            "Hello World from GPU! 0 93\n",
            "Hello World from GPU! 0 94\n",
            "Hello World from GPU! 0 95\n",
            "Hello World from GPU! 0 832\n",
            "Hello World from GPU! 0 833\n",
            "Hello World from GPU! 0 834\n",
            "Hello World from GPU! 0 835\n",
            "Hello World from GPU! 0 836\n",
            "Hello World from GPU! 0 837\n",
            "Hello World from GPU! 0 838\n",
            "Hello World from GPU! 0 839\n",
            "Hello World from GPU! 0 840\n",
            "Hello World from GPU! 0 841\n",
            "Hello World from GPU! 0 842\n",
            "Hello World from GPU! 0 843\n",
            "Hello World from GPU! 0 844\n",
            "Hello World from GPU! 0 845\n",
            "Hello World from GPU! 0 846\n",
            "Hello World from GPU! 0 847\n",
            "Hello World from GPU! 0 848\n",
            "Hello World from GPU! 0 849\n",
            "Hello World from GPU! 0 850\n",
            "Hello World from GPU! 0 851\n",
            "Hello World from GPU! 0 852\n",
            "Hello World from GPU! 0 853\n",
            "Hello World from GPU! 0 854\n",
            "Hello World from GPU! 0 855\n",
            "Hello World from GPU! 0 856\n",
            "Hello World from GPU! 0 857\n",
            "Hello World from GPU! 0 858\n",
            "Hello World from GPU! 0 859\n",
            "Hello World from GPU! 0 860\n",
            "Hello World from GPU! 0 861\n",
            "Hello World from GPU! 0 862\n",
            "Hello World from GPU! 0 863\n",
            "Hello World from GPU! 0 800\n",
            "Hello World from GPU! 0 801\n",
            "Hello World from GPU! 0 802\n",
            "Hello World from GPU! 0 803\n",
            "Hello World from GPU! 0 804\n",
            "Hello World from GPU! 0 805\n",
            "Hello World from GPU! 0 806\n",
            "Hello World from GPU! 0 807\n",
            "Hello World from GPU! 0 808\n",
            "Hello World from GPU! 0 809\n",
            "Hello World from GPU! 0 810\n",
            "Hello World from GPU! 0 811\n",
            "Hello World from GPU! 0 812\n",
            "Hello World from GPU! 0 813\n",
            "Hello World from GPU! 0 814\n",
            "Hello World from GPU! 0 815\n",
            "Hello World from GPU! 0 816\n",
            "Hello World from GPU! 0 817\n",
            "Hello World from GPU! 0 818\n",
            "Hello World from GPU! 0 819\n",
            "Hello World from GPU! 0 820\n",
            "Hello World from GPU! 0 821\n",
            "Hello World from GPU! 0 822\n",
            "Hello World from GPU! 0 823\n",
            "Hello World from GPU! 0 824\n",
            "Hello World from GPU! 0 825\n",
            "Hello World from GPU! 0 826\n",
            "Hello World from GPU! 0 827\n",
            "Hello World from GPU! 0 828\n",
            "Hello World from GPU! 0 829\n",
            "Hello World from GPU! 0 830\n",
            "Hello World from GPU! 0 831\n",
            "Hello World from GPU! 0 768\n",
            "Hello World from GPU! 0 769\n",
            "Hello World from GPU! 0 770\n",
            "Hello World from GPU! 0 771\n",
            "Hello World from GPU! 0 772\n",
            "Hello World from GPU! 0 773\n",
            "Hello World from GPU! 0 774\n",
            "Hello World from GPU! 0 775\n",
            "Hello World from GPU! 0 776\n",
            "Hello World from GPU! 0 777\n",
            "Hello World from GPU! 0 778\n",
            "Hello World from GPU! 0 779\n",
            "Hello World from GPU! 0 780\n",
            "Hello World from GPU! 0 781\n",
            "Hello World from GPU! 0 782\n",
            "Hello World from GPU! 0 783\n",
            "Hello World from GPU! 0 784\n",
            "Hello World from GPU! 0 785\n",
            "Hello World from GPU! 0 786\n",
            "Hello World from GPU! 0 787\n",
            "Hello World from GPU! 0 788\n",
            "Hello World from GPU! 0 789\n",
            "Hello World from GPU! 0 790\n",
            "Hello World from GPU! 0 791\n",
            "Hello World from GPU! 0 792\n",
            "Hello World from GPU! 0 793\n",
            "Hello World from GPU! 0 794\n",
            "Hello World from GPU! 0 795\n",
            "Hello World from GPU! 0 796\n",
            "Hello World from GPU! 0 797\n",
            "Hello World from GPU! 0 798\n",
            "Hello World from GPU! 0 799\n",
            "Hello World from GPU! 0 992\n",
            "Hello World from GPU! 0 993\n",
            "Hello World from GPU! 0 994\n",
            "Hello World from GPU! 0 995\n",
            "Hello World from GPU! 0 996\n",
            "Hello World from GPU! 0 997\n",
            "Hello World from GPU! 0 998\n",
            "Hello World from GPU! 0 999\n",
            "Hello World from GPU! 0 1000\n",
            "Hello World from GPU! 0 1001\n",
            "Hello World from GPU! 0 1002\n",
            "Hello World from GPU! 0 1003\n",
            "Hello World from GPU! 0 1004\n",
            "Hello World from GPU! 0 1005\n",
            "Hello World from GPU! 0 1006\n",
            "Hello World from GPU! 0 1007\n",
            "Hello World from GPU! 0 1008\n",
            "Hello World from GPU! 0 1009\n",
            "Hello World from GPU! 0 1010\n",
            "Hello World from GPU! 0 1011\n",
            "Hello World from GPU! 0 1012\n",
            "Hello World from GPU! 0 1013\n",
            "Hello World from GPU! 0 1014\n",
            "Hello World from GPU! 0 1015\n",
            "Hello World from GPU! 0 1016\n",
            "Hello World from GPU! 0 1017\n",
            "Hello World from GPU! 0 1018\n",
            "Hello World from GPU! 0 1019\n",
            "Hello World from GPU! 0 1020\n",
            "Hello World from GPU! 0 1021\n",
            "Hello World from GPU! 0 1022\n",
            "Hello World from GPU! 0 1023\n",
            "Hello World from GPU! 0 672\n",
            "Hello World from GPU! 0 673\n",
            "Hello World from GPU! 0 674\n",
            "Hello World from GPU! 0 675\n",
            "Hello World from GPU! 0 676\n",
            "Hello World from GPU! 0 677\n",
            "Hello World from GPU! 0 678\n",
            "Hello World from GPU! 0 679\n",
            "Hello World from GPU! 0 680\n",
            "Hello World from GPU! 0 681\n",
            "Hello World from GPU! 0 682\n",
            "Hello World from GPU! 0 683\n",
            "Hello World from GPU! 0 684\n",
            "Hello World from GPU! 0 685\n",
            "Hello World from GPU! 0 686\n",
            "Hello World from GPU! 0 687\n",
            "Hello World from GPU! 0 688\n",
            "Hello World from GPU! 0 689\n",
            "Hello World from GPU! 0 690\n",
            "Hello World from GPU! 0 691\n",
            "Hello World from GPU! 0 692\n",
            "Hello World from GPU! 0 693\n",
            "Hello World from GPU! 0 694\n",
            "Hello World from GPU! 0 695\n",
            "Hello World from GPU! 0 696\n",
            "Hello World from GPU! 0 697\n",
            "Hello World from GPU! 0 698\n",
            "Hello World from GPU! 0 699\n",
            "Hello World from GPU! 0 700\n",
            "Hello World from GPU! 0 701\n",
            "Hello World from GPU! 0 702\n",
            "Hello World from GPU! 0 703\n",
            "Hello World from GPU! 0 736\n",
            "Hello World from GPU! 0 737\n",
            "Hello World from GPU! 0 738\n",
            "Hello World from GPU! 0 739\n",
            "Hello World from GPU! 0 740\n",
            "Hello World from GPU! 0 741\n",
            "Hello World from GPU! 0 742\n",
            "Hello World from GPU! 0 743\n",
            "Hello World from GPU! 0 744\n",
            "Hello World from GPU! 0 745\n",
            "Hello World from GPU! 0 746\n",
            "Hello World from GPU! 0 747\n",
            "Hello World from GPU! 0 748\n",
            "Hello World from GPU! 0 749\n",
            "Hello World from GPU! 0 750\n",
            "Hello World from GPU! 0 751\n",
            "Hello World from GPU! 0 752\n",
            "Hello World from GPU! 0 753\n",
            "Hello World from GPU! 0 754\n",
            "Hello World from GPU! 0 755\n",
            "Hello World from GPU! 0 756\n",
            "Hello World from GPU! 0 757\n",
            "Hello World from GPU! 0 758\n",
            "Hello World from GPU! 0 759\n",
            "Hello World from GPU! 0 760\n",
            "Hello World from GPU! 0 761\n",
            "Hello World from GPU! 0 762\n",
            "Hello World from GPU! 0 763\n",
            "Hello World from GPU! 0 764\n",
            "Hello World from GPU! 0 765\n",
            "Hello World from GPU! 0 766\n",
            "Hello World from GPU! 0 767\n",
            "Hello World from GPU! 0 96\n",
            "Hello World from GPU! 0 97\n",
            "Hello World from GPU! 0 98\n",
            "Hello World from GPU! 0 99\n",
            "Hello World from GPU! 0 100\n",
            "Hello World from GPU! 0 101\n",
            "Hello World from GPU! 0 102\n",
            "Hello World from GPU! 0 103\n",
            "Hello World from GPU! 0 104\n",
            "Hello World from GPU! 0 105\n",
            "Hello World from GPU! 0 106\n",
            "Hello World from GPU! 0 107\n",
            "Hello World from GPU! 0 108\n",
            "Hello World from GPU! 0 109\n",
            "Hello World from GPU! 0 110\n",
            "Hello World from GPU! 0 111\n",
            "Hello World from GPU! 0 112\n",
            "Hello World from GPU! 0 113\n",
            "Hello World from GPU! 0 114\n",
            "Hello World from GPU! 0 115\n",
            "Hello World from GPU! 0 116\n",
            "Hello World from GPU! 0 117\n",
            "Hello World from GPU! 0 118\n",
            "Hello World from GPU! 0 119\n",
            "Hello World from GPU! 0 120\n",
            "Hello World from GPU! 0 121\n",
            "Hello World from GPU! 0 122\n",
            "Hello World from GPU! 0 123\n",
            "Hello World from GPU! 0 124\n",
            "Hello World from GPU! 0 125\n",
            "Hello World from GPU! 0 126\n",
            "Hello World from GPU! 0 127\n",
            "Hello World from GPU! 0 224\n",
            "Hello World from GPU! 0 225\n",
            "Hello World from GPU! 0 226\n",
            "Hello World from GPU! 0 227\n",
            "Hello World from GPU! 0 228\n",
            "Hello World from GPU! 0 229\n",
            "Hello World from GPU! 0 230\n",
            "Hello World from GPU! 0 231\n",
            "Hello World from GPU! 0 232\n",
            "Hello World from GPU! 0 233\n",
            "Hello World from GPU! 0 234\n",
            "Hello World from GPU! 0 235\n",
            "Hello World from GPU! 0 236\n",
            "Hello World from GPU! 0 237\n",
            "Hello World from GPU! 0 238\n",
            "Hello World from GPU! 0 239\n",
            "Hello World from GPU! 0 240\n",
            "Hello World from GPU! 0 241\n",
            "Hello World from GPU! 0 242\n",
            "Hello World from GPU! 0 243\n",
            "Hello World from GPU! 0 244\n",
            "Hello World from GPU! 0 245\n",
            "Hello World from GPU! 0 246\n",
            "Hello World from GPU! 0 247\n",
            "Hello World from GPU! 0 248\n",
            "Hello World from GPU! 0 249\n",
            "Hello World from GPU! 0 250\n",
            "Hello World from GPU! 0 251\n",
            "Hello World from GPU! 0 252\n",
            "Hello World from GPU! 0 253\n",
            "Hello World from GPU! 0 254\n",
            "Hello World from GPU! 0 255\n",
            "Hello World from GPU! 0 320\n",
            "Hello World from GPU! 0 321\n",
            "Hello World from GPU! 0 322\n",
            "Hello World from GPU! 0 323\n",
            "Hello World from GPU! 0 324\n",
            "Hello World from GPU! 0 325\n",
            "Hello World from GPU! 0 326\n",
            "Hello World from GPU! 0 327\n",
            "Hello World from GPU! 0 328\n",
            "Hello World from GPU! 0 329\n",
            "Hello World from GPU! 0 330\n",
            "Hello World from GPU! 0 331\n",
            "Hello World from GPU! 0 332\n",
            "Hello World from GPU! 0 333\n",
            "Hello World from GPU! 0 334\n",
            "Hello World from GPU! 0 335\n",
            "Hello World from GPU! 0 336\n",
            "Hello World from GPU! 0 337\n",
            "Hello World from GPU! 0 338\n",
            "Hello World from GPU! 0 339\n",
            "Hello World from GPU! 0 340\n",
            "Hello World from GPU! 0 341\n",
            "Hello World from GPU! 0 342\n",
            "Hello World from GPU! 0 343\n",
            "Hello World from GPU! 0 344\n",
            "Hello World from GPU! 0 345\n",
            "Hello World from GPU! 0 346\n",
            "Hello World from GPU! 0 347\n",
            "Hello World from GPU! 0 348\n",
            "Hello World from GPU! 0 349\n",
            "Hello World from GPU! 0 350\n",
            "Hello World from GPU! 0 351\n",
            "Hello World from GPU! 0 704\n",
            "Hello World from GPU! 0 705\n",
            "Hello World from GPU! 0 706\n",
            "Hello World from GPU! 0 707\n",
            "Hello World from GPU! 0 708\n",
            "Hello World from GPU! 0 709\n",
            "Hello World from GPU! 0 710\n",
            "Hello World from GPU! 0 711\n",
            "Hello World from GPU! 0 712\n",
            "Hello World from GPU! 0 713\n",
            "Hello World from GPU! 0 714\n",
            "Hello World from GPU! 0 715\n",
            "Hello World from GPU! 0 716\n",
            "Hello World from GPU! 0 717\n",
            "Hello World from GPU! 0 718\n",
            "Hello World from GPU! 0 719\n",
            "Hello World from GPU! 0 720\n",
            "Hello World from GPU! 0 721\n",
            "Hello World from GPU! 0 722\n",
            "Hello World from GPU! 0 723\n",
            "Hello World from GPU! 0 724\n",
            "Hello World from GPU! 0 725\n",
            "Hello World from GPU! 0 726\n",
            "Hello World from GPU! 0 727\n",
            "Hello World from GPU! 0 728\n",
            "Hello World from GPU! 0 729\n",
            "Hello World from GPU! 0 730\n",
            "Hello World from GPU! 0 731\n",
            "Hello World from GPU! 0 732\n",
            "Hello World from GPU! 0 733\n",
            "Hello World from GPU! 0 734\n",
            "Hello World from GPU! 0 735\n",
            "Hello World from GPU! 0 32\n",
            "Hello World from GPU! 0 33\n",
            "Hello World from GPU! 0 34\n",
            "Hello World from GPU! 0 35\n",
            "Hello World from GPU! 0 36\n",
            "Hello World from GPU! 0 37\n",
            "Hello World from GPU! 0 38\n",
            "Hello World from GPU! 0 39\n",
            "Hello World from GPU! 0 40\n",
            "Hello World from GPU! 0 41\n",
            "Hello World from GPU! 0 42\n",
            "Hello World from GPU! 0 43\n",
            "Hello World from GPU! 0 44\n",
            "Hello World from GPU! 0 45\n",
            "Hello World from GPU! 0 46\n",
            "Hello World from GPU! 0 47\n",
            "Hello World from GPU! 0 48\n",
            "Hello World from GPU! 0 49\n",
            "Hello World from GPU! 0 50\n",
            "Hello World from GPU! 0 51\n",
            "Hello World from GPU! 0 52\n",
            "Hello World from GPU! 0 53\n",
            "Hello World from GPU! 0 54\n",
            "Hello World from GPU! 0 55\n",
            "Hello World from GPU! 0 56\n",
            "Hello World from GPU! 0 57\n",
            "Hello World from GPU! 0 58\n",
            "Hello World from GPU! 0 59\n",
            "Hello World from GPU! 0 60\n",
            "Hello World from GPU! 0 61\n",
            "Hello World from GPU! 0 62\n",
            "Hello World from GPU! 0 63\n",
            "Hello World from GPU! 0 544\n",
            "Hello World from GPU! 0 545\n",
            "Hello World from GPU! 0 546\n",
            "Hello World from GPU! 0 547\n",
            "Hello World from GPU! 0 548\n",
            "Hello World from GPU! 0 549\n",
            "Hello World from GPU! 0 550\n",
            "Hello World from GPU! 0 551\n",
            "Hello World from GPU! 0 552\n",
            "Hello World from GPU! 0 553\n",
            "Hello World from GPU! 0 554\n",
            "Hello World from GPU! 0 555\n",
            "Hello World from GPU! 0 556\n",
            "Hello World from GPU! 0 557\n",
            "Hello World from GPU! 0 558\n",
            "Hello World from GPU! 0 559\n",
            "Hello World from GPU! 0 560\n",
            "Hello World from GPU! 0 561\n",
            "Hello World from GPU! 0 562\n",
            "Hello World from GPU! 0 563\n",
            "Hello World from GPU! 0 564\n",
            "Hello World from GPU! 0 565\n",
            "Hello World from GPU! 0 566\n",
            "Hello World from GPU! 0 567\n",
            "Hello World from GPU! 0 568\n",
            "Hello World from GPU! 0 569\n",
            "Hello World from GPU! 0 570\n",
            "Hello World from GPU! 0 571\n",
            "Hello World from GPU! 0 572\n",
            "Hello World from GPU! 0 573\n",
            "Hello World from GPU! 0 574\n",
            "Hello World from GPU! 0 575\n",
            "\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "6KH5ZgEbmZD2",
        "colab_type": "text"
      },
      "source": [
        "Resposta. Apresenta grupos de execução sequencial de 32 threads. É possível determinar isso, porém não é possível dizer nada sobre quais grupos executarão primeiro. Provavelmente cada bloco da GPU tem apenas 32 threads, então o compilador alocará vários para rodar todos os 1024 threads solicitados."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "ZOKZlF1-bXUj",
        "colab_type": "text"
      },
      "source": [
        "2. Agora repita o exercício anterior com 1024 blocos com 1 thread cada. Como é a ordem de execução dos threads? Sequencial? Determinística? Agrupada?"
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "8W_U1CnDlgJ2",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 1000
        },
        "outputId": "3ad2d6d3-6e78-4995-ad63-8a48b2aa1455"
      },
      "source": [
        "%%cu\n",
        "#include <stdio.h>\n",
        "#include <stdlib.h>\n",
        "\n",
        "#define CHECK(call) { \\\n",
        "  const cudaError_t error = call; \\\n",
        "  if (error != cudaSuccess) { \\\n",
        "    printf(\"Error: %s:%d, \", __FILE__, __LINE__); \\\n",
        "    printf(\"code:%d, reason: %s\\n\", error, cudaGetErrorString(error)); \\\n",
        "    exit(1); \\\n",
        "  } \\\n",
        "}\n",
        "\n",
        "__global__ void helloFromGPU() {\n",
        "  printf(\"Hello World from GPU! %d %d\\n\", blockIdx.x, threadIdx.x);\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"Hello World from CPU!\\n\");\n",
        "  helloFromGPU<<<1024, 1>>>();\n",
        "  CHECK(cudaDeviceReset());\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Hello World from CPU!\n",
            "Hello World from GPU! 127 0\n",
            "Hello World from GPU! 4 0\n",
            "Hello World from GPU! 117 0\n",
            "Hello World from GPU! 128 0\n",
            "Hello World from GPU! 122 0\n",
            "Hello World from GPU! 9 0\n",
            "Hello World from GPU! 125 0\n",
            "Hello World from GPU! 129 0\n",
            "Hello World from GPU! 6 0\n",
            "Hello World from GPU! 140 0\n",
            "Hello World from GPU! 121 0\n",
            "Hello World from GPU! 143 0\n",
            "Hello World from GPU! 124 0\n",
            "Hello World from GPU! 193 0\n",
            "Hello World from GPU! 18 0\n",
            "Hello World from GPU! 196 0\n",
            "Hello World from GPU! 8 0\n",
            "Hello World from GPU! 25 0\n",
            "Hello World from GPU! 123 0\n",
            "Hello World from GPU! 119 0\n",
            "Hello World from GPU! 126 0\n",
            "Hello World from GPU! 120 0\n",
            "Hello World from GPU! 150 0\n",
            "Hello World from GPU! 154 0\n",
            "Hello World from GPU! 166 0\n",
            "Hello World from GPU! 69 0\n",
            "Hello World from GPU! 195 0\n",
            "Hello World from GPU! 12 0\n",
            "Hello World from GPU! 201 0\n",
            "Hello World from GPU! 200 0\n",
            "Hello World from GPU! 157 0\n",
            "Hello World from GPU! 21 0\n",
            "Hello World from GPU! 202 0\n",
            "Hello World from GPU! 37 0\n",
            "Hello World from GPU! 2 0\n",
            "Hello World from GPU! 139 0\n",
            "Hello World from GPU! 3 0\n",
            "Hello World from GPU! 142 0\n",
            "Hello World from GPU! 71 0\n",
            "Hello World from GPU! 23 0\n",
            "Hello World from GPU! 56 0\n",
            "Hello World from GPU! 26 0\n",
            "Hello World from GPU! 206 0\n",
            "Hello World from GPU! 59 0\n",
            "Hello World from GPU! 161 0\n",
            "Hello World from GPU! 60 0\n",
            "Hello World from GPU! 170 0\n",
            "Hello World from GPU! 19 0\n",
            "Hello World from GPU! 54 0\n",
            "Hello World from GPU! 168 0\n",
            "Hello World from GPU! 48 0\n",
            "Hello World from GPU! 20 0\n",
            "Hello World from GPU! 16 0\n",
            "Hello World from GPU! 36 0\n",
            "Hello World from GPU! 30 0\n",
            "Hello World from GPU! 141 0\n",
            "Hello World from GPU! 194 0\n",
            "Hello World from GPU! 45 0\n",
            "Hello World from GPU! 0 0\n",
            "Hello World from GPU! 14 0\n",
            "Hello World from GPU! 177 0\n",
            "Hello World from GPU! 148 0\n",
            "Hello World from GPU! 33 0\n",
            "Hello World from GPU! 132 0\n",
            "Hello World from GPU! 180 0\n",
            "Hello World from GPU! 22 0\n",
            "Hello World from GPU! 38 0\n",
            "Hello World from GPU! 146 0\n",
            "Hello World from GPU! 153 0\n",
            "Hello World from GPU! 43 0\n",
            "Hello World from GPU! 149 0\n",
            "Hello World from GPU! 78 0\n",
            "Hello World from GPU! 167 0\n",
            "Hello World from GPU! 155 0\n",
            "Hello World from GPU! 57 0\n",
            "Hello World from GPU! 73 0\n",
            "Hello World from GPU! 131 0\n",
            "Hello World from GPU! 77 0\n",
            "Hello World from GPU! 67 0\n",
            "Hello World from GPU! 24 0\n",
            "Hello World from GPU! 81 0\n",
            "Hello World from GPU! 98 0\n",
            "Hello World from GPU! 178 0\n",
            "Hello World from GPU! 75 0\n",
            "Hello World from GPU! 17 0\n",
            "Hello World from GPU! 175 0\n",
            "Hello World from GPU! 46 0\n",
            "Hello World from GPU! 102 0\n",
            "Hello World from GPU! 52 0\n",
            "Hello World from GPU! 44 0\n",
            "Hello World from GPU! 34 0\n",
            "Hello World from GPU! 118 0\n",
            "Hello World from GPU! 32 0\n",
            "Hello World from GPU! 145 0\n",
            "Hello World from GPU! 7 0\n",
            "Hello World from GPU! 55 0\n",
            "Hello World from GPU! 85 0\n",
            "Hello World from GPU! 35 0\n",
            "Hello World from GPU! 63 0\n",
            "Hello World from GPU! 134 0\n",
            "Hello World from GPU! 62 0\n",
            "Hello World from GPU! 13 0\n",
            "Hello World from GPU! 163 0\n",
            "Hello World from GPU! 31 0\n",
            "Hello World from GPU! 151 0\n",
            "Hello World from GPU! 50 0\n",
            "Hello World from GPU! 66 0\n",
            "Hello World from GPU! 15 0\n",
            "Hello World from GPU! 100 0\n",
            "Hello World from GPU! 176 0\n",
            "Hello World from GPU! 11 0\n",
            "Hello World from GPU! 115 0\n",
            "Hello World from GPU! 10 0\n",
            "Hello World from GPU! 107 0\n",
            "Hello World from GPU! 207 0\n",
            "Hello World from GPU! 64 0\n",
            "Hello World from GPU! 89 0\n",
            "Hello World from GPU! 164 0\n",
            "Hello World from GPU! 108 0\n",
            "Hello World from GPU! 68 0\n",
            "Hello World from GPU! 181 0\n",
            "Hello World from GPU! 187 0\n",
            "Hello World from GPU! 53 0\n",
            "Hello World from GPU! 136 0\n",
            "Hello World from GPU! 138 0\n",
            "Hello World from GPU! 28 0\n",
            "Hello World from GPU! 61 0\n",
            "Hello World from GPU! 114 0\n",
            "Hello World from GPU! 147 0\n",
            "Hello World from GPU! 110 0\n",
            "Hello World from GPU! 137 0\n",
            "Hello World from GPU! 172 0\n",
            "Hello World from GPU! 70 0\n",
            "Hello World from GPU! 90 0\n",
            "Hello World from GPU! 58 0\n",
            "Hello World from GPU! 1 0\n",
            "Hello World from GPU! 84 0\n",
            "Hello World from GPU! 169 0\n",
            "Hello World from GPU! 51 0\n",
            "Hello World from GPU! 197 0\n",
            "Hello World from GPU! 72 0\n",
            "Hello World from GPU! 86 0\n",
            "Hello World from GPU! 165 0\n",
            "Hello World from GPU! 199 0\n",
            "Hello World from GPU! 42 0\n",
            "Hello World from GPU! 76 0\n",
            "Hello World from GPU! 179 0\n",
            "Hello World from GPU! 162 0\n",
            "Hello World from GPU! 40 0\n",
            "Hello World from GPU! 96 0\n",
            "Hello World from GPU! 111 0\n",
            "Hello World from GPU! 97 0\n",
            "Hello World from GPU! 106 0\n",
            "Hello World from GPU! 189 0\n",
            "Hello World from GPU! 29 0\n",
            "Hello World from GPU! 91 0\n",
            "Hello World from GPU! 88 0\n",
            "Hello World from GPU! 188 0\n",
            "Hello World from GPU! 204 0\n",
            "Hello World from GPU! 5 0\n",
            "Hello World from GPU! 82 0\n",
            "Hello World from GPU! 113 0\n",
            "Hello World from GPU! 198 0\n",
            "Hello World from GPU! 171 0\n",
            "Hello World from GPU! 185 0\n",
            "Hello World from GPU! 133 0\n",
            "Hello World from GPU! 95 0\n",
            "Hello World from GPU! 130 0\n",
            "Hello World from GPU! 174 0\n",
            "Hello World from GPU! 109 0\n",
            "Hello World from GPU! 152 0\n",
            "Hello World from GPU! 83 0\n",
            "Hello World from GPU! 186 0\n",
            "Hello World from GPU! 87 0\n",
            "Hello World from GPU! 159 0\n",
            "Hello World from GPU! 184 0\n",
            "Hello World from GPU! 203 0\n",
            "Hello World from GPU! 156 0\n",
            "Hello World from GPU! 160 0\n",
            "Hello World from GPU! 104 0\n",
            "Hello World from GPU! 192 0\n",
            "Hello World from GPU! 135 0\n",
            "Hello World from GPU! 173 0\n",
            "Hello World from GPU! 191 0\n",
            "Hello World from GPU! 158 0\n",
            "Hello World from GPU! 190 0\n",
            "Hello World from GPU! 103 0\n",
            "Hello World from GPU! 80 0\n",
            "Hello World from GPU! 105 0\n",
            "Hello World from GPU! 99 0\n",
            "Hello World from GPU! 183 0\n",
            "Hello World from GPU! 182 0\n",
            "Hello World from GPU! 65 0\n",
            "Hello World from GPU! 205 0\n",
            "Hello World from GPU! 92 0\n",
            "Hello World from GPU! 41 0\n",
            "Hello World from GPU! 101 0\n",
            "Hello World from GPU! 74 0\n",
            "Hello World from GPU! 39 0\n",
            "Hello World from GPU! 93 0\n",
            "Hello World from GPU! 27 0\n",
            "Hello World from GPU! 144 0\n",
            "Hello World from GPU! 79 0\n",
            "Hello World from GPU! 112 0\n",
            "Hello World from GPU! 116 0\n",
            "Hello World from GPU! 49 0\n",
            "Hello World from GPU! 94 0\n",
            "Hello World from GPU! 47 0\n",
            "Hello World from GPU! 209 0\n",
            "Hello World from GPU! 208 0\n",
            "Hello World from GPU! 211 0\n",
            "Hello World from GPU! 226 0\n",
            "Hello World from GPU! 220 0\n",
            "Hello World from GPU! 218 0\n",
            "Hello World from GPU! 210 0\n",
            "Hello World from GPU! 216 0\n",
            "Hello World from GPU! 217 0\n",
            "Hello World from GPU! 222 0\n",
            "Hello World from GPU! 221 0\n",
            "Hello World from GPU! 219 0\n",
            "Hello World from GPU! 213 0\n",
            "Hello World from GPU! 227 0\n",
            "Hello World from GPU! 223 0\n",
            "Hello World from GPU! 212 0\n",
            "Hello World from GPU! 228 0\n",
            "Hello World from GPU! 230 0\n",
            "Hello World from GPU! 215 0\n",
            "Hello World from GPU! 214 0\n",
            "Hello World from GPU! 225 0\n",
            "Hello World from GPU! 224 0\n",
            "Hello World from GPU! 229 0\n",
            "Hello World from GPU! 233 0\n",
            "Hello World from GPU! 236 0\n",
            "Hello World from GPU! 234 0\n",
            "Hello World from GPU! 235 0\n",
            "Hello World from GPU! 231 0\n",
            "Hello World from GPU! 243 0\n",
            "Hello World from GPU! 246 0\n",
            "Hello World from GPU! 244 0\n",
            "Hello World from GPU! 237 0\n",
            "Hello World from GPU! 238 0\n",
            "Hello World from GPU! 257 0\n",
            "Hello World from GPU! 245 0\n",
            "Hello World from GPU! 252 0\n",
            "Hello World from GPU! 256 0\n",
            "Hello World from GPU! 240 0\n",
            "Hello World from GPU! 242 0\n",
            "Hello World from GPU! 260 0\n",
            "Hello World from GPU! 241 0\n",
            "Hello World from GPU! 247 0\n",
            "Hello World from GPU! 239 0\n",
            "Hello World from GPU! 249 0\n",
            "Hello World from GPU! 248 0\n",
            "Hello World from GPU! 259 0\n",
            "Hello World from GPU! 255 0\n",
            "Hello World from GPU! 262 0\n",
            "Hello World from GPU! 232 0\n",
            "Hello World from GPU! 250 0\n",
            "Hello World from GPU! 251 0\n",
            "Hello World from GPU! 254 0\n",
            "Hello World from GPU! 264 0\n",
            "Hello World from GPU! 263 0\n",
            "Hello World from GPU! 265 0\n",
            "Hello World from GPU! 266 0\n",
            "Hello World from GPU! 253 0\n",
            "Hello World from GPU! 261 0\n",
            "Hello World from GPU! 267 0\n",
            "Hello World from GPU! 258 0\n",
            "Hello World from GPU! 271 0\n",
            "Hello World from GPU! 274 0\n",
            "Hello World from GPU! 276 0\n",
            "Hello World from GPU! 273 0\n",
            "Hello World from GPU! 277 0\n",
            "Hello World from GPU! 280 0\n",
            "Hello World from GPU! 279 0\n",
            "Hello World from GPU! 268 0\n",
            "Hello World from GPU! 269 0\n",
            "Hello World from GPU! 281 0\n",
            "Hello World from GPU! 287 0\n",
            "Hello World from GPU! 284 0\n",
            "Hello World from GPU! 285 0\n",
            "Hello World from GPU! 289 0\n",
            "Hello World from GPU! 282 0\n",
            "Hello World from GPU! 286 0\n",
            "Hello World from GPU! 283 0\n",
            "Hello World from GPU! 291 0\n",
            "Hello World from GPU! 290 0\n",
            "Hello World from GPU! 270 0\n",
            "Hello World from GPU! 272 0\n",
            "Hello World from GPU! 278 0\n",
            "Hello World from GPU! 275 0\n",
            "Hello World from GPU! 288 0\n",
            "Hello World from GPU! 292 0\n",
            "Hello World from GPU! 295 0\n",
            "Hello World from GPU! 293 0\n",
            "Hello World from GPU! 294 0\n",
            "Hello World from GPU! 296 0\n",
            "Hello World from GPU! 310 0\n",
            "Hello World from GPU! 328 0\n",
            "Hello World from GPU! 299 0\n",
            "Hello World from GPU! 297 0\n",
            "Hello World from GPU! 331 0\n",
            "Hello World from GPU! 340 0\n",
            "Hello World from GPU! 344 0\n",
            "Hello World from GPU! 354 0\n",
            "Hello World from GPU! 338 0\n",
            "Hello World from GPU! 362 0\n",
            "Hello World from GPU! 298 0\n",
            "Hello World from GPU! 367 0\n",
            "Hello World from GPU! 303 0\n",
            "Hello World from GPU! 325 0\n",
            "Hello World from GPU! 312 0\n",
            "Hello World from GPU! 350 0\n",
            "Hello World from GPU! 313 0\n",
            "Hello World from GPU! 306 0\n",
            "Hello World from GPU! 301 0\n",
            "Hello World from GPU! 305 0\n",
            "Hello World from GPU! 386 0\n",
            "Hello World from GPU! 318 0\n",
            "Hello World from GPU! 300 0\n",
            "Hello World from GPU! 311 0\n",
            "Hello World from GPU! 314 0\n",
            "Hello World from GPU! 308 0\n",
            "Hello World from GPU! 302 0\n",
            "Hello World from GPU! 307 0\n",
            "Hello World from GPU! 304 0\n",
            "Hello World from GPU! 392 0\n",
            "Hello World from GPU! 394 0\n",
            "Hello World from GPU! 335 0\n",
            "Hello World from GPU! 341 0\n",
            "Hello World from GPU! 333 0\n",
            "Hello World from GPU! 339 0\n",
            "Hello World from GPU! 329 0\n",
            "Hello World from GPU! 355 0\n",
            "Hello World from GPU! 320 0\n",
            "Hello World from GPU! 327 0\n",
            "Hello World from GPU! 334 0\n",
            "Hello World from GPU! 347 0\n",
            "Hello World from GPU! 337 0\n",
            "Hello World from GPU! 317 0\n",
            "Hello World from GPU! 368 0\n",
            "Hello World from GPU! 324 0\n",
            "Hello World from GPU! 365 0\n",
            "Hello World from GPU! 330 0\n",
            "Hello World from GPU! 374 0\n",
            "Hello World from GPU! 319 0\n",
            "Hello World from GPU! 361 0\n",
            "Hello World from GPU! 358 0\n",
            "Hello World from GPU! 356 0\n",
            "Hello World from GPU! 349 0\n",
            "Hello World from GPU! 372 0\n",
            "Hello World from GPU! 369 0\n",
            "Hello World from GPU! 332 0\n",
            "Hello World from GPU! 360 0\n",
            "Hello World from GPU! 346 0\n",
            "Hello World from GPU! 377 0\n",
            "Hello World from GPU! 373 0\n",
            "Hello World from GPU! 378 0\n",
            "Hello World from GPU! 381 0\n",
            "Hello World from GPU! 382 0\n",
            "Hello World from GPU! 383 0\n",
            "Hello World from GPU! 379 0\n",
            "Hello World from GPU! 371 0\n",
            "Hello World from GPU! 326 0\n",
            "Hello World from GPU! 309 0\n",
            "Hello World from GPU! 376 0\n",
            "Hello World from GPU! 336 0\n",
            "Hello World from GPU! 321 0\n",
            "Hello World from GPU! 375 0\n",
            "Hello World from GPU! 343 0\n",
            "Hello World from GPU! 322 0\n",
            "Hello World from GPU! 380 0\n",
            "Hello World from GPU! 342 0\n",
            "Hello World from GPU! 388 0\n",
            "Hello World from GPU! 345 0\n",
            "Hello World from GPU! 352 0\n",
            "Hello World from GPU! 316 0\n",
            "Hello World from GPU! 359 0\n",
            "Hello World from GPU! 348 0\n",
            "Hello World from GPU! 363 0\n",
            "Hello World from GPU! 357 0\n",
            "Hello World from GPU! 351 0\n",
            "Hello World from GPU! 353 0\n",
            "Hello World from GPU! 364 0\n",
            "Hello World from GPU! 395 0\n",
            "Hello World from GPU! 397 0\n",
            "Hello World from GPU! 401 0\n",
            "Hello World from GPU! 403 0\n",
            "Hello World from GPU! 384 0\n",
            "Hello World from GPU! 390 0\n",
            "Hello World from GPU! 408 0\n",
            "Hello World from GPU! 391 0\n",
            "Hello World from GPU! 396 0\n",
            "Hello World from GPU! 415 0\n",
            "Hello World from GPU! 409 0\n",
            "Hello World from GPU! 370 0\n",
            "Hello World from GPU! 400 0\n",
            "Hello World from GPU! 399 0\n",
            "Hello World from GPU! 389 0\n",
            "Hello World from GPU! 402 0\n",
            "Hello World from GPU! 411 0\n",
            "Hello World from GPU! 412 0\n",
            "Hello World from GPU! 414 0\n",
            "Hello World from GPU! 410 0\n",
            "Hello World from GPU! 385 0\n",
            "Hello World from GPU! 405 0\n",
            "Hello World from GPU! 407 0\n",
            "Hello World from GPU! 393 0\n",
            "Hello World from GPU! 406 0\n",
            "Hello World from GPU! 398 0\n",
            "Hello World from GPU! 323 0\n",
            "Hello World from GPU! 315 0\n",
            "Hello World from GPU! 387 0\n",
            "Hello World from GPU! 413 0\n",
            "Hello World from GPU! 366 0\n",
            "Hello World from GPU! 404 0\n",
            "Hello World from GPU! 417 0\n",
            "Hello World from GPU! 416 0\n",
            "Hello World from GPU! 419 0\n",
            "Hello World from GPU! 420 0\n",
            "Hello World from GPU! 421 0\n",
            "Hello World from GPU! 422 0\n",
            "Hello World from GPU! 423 0\n",
            "Hello World from GPU! 426 0\n",
            "Hello World from GPU! 425 0\n",
            "Hello World from GPU! 424 0\n",
            "Hello World from GPU! 418 0\n",
            "Hello World from GPU! 428 0\n",
            "Hello World from GPU! 427 0\n",
            "Hello World from GPU! 430 0\n",
            "Hello World from GPU! 431 0\n",
            "Hello World from GPU! 429 0\n",
            "Hello World from GPU! 432 0\n",
            "Hello World from GPU! 433 0\n",
            "Hello World from GPU! 442 0\n",
            "Hello World from GPU! 466 0\n",
            "Hello World from GPU! 438 0\n",
            "Hello World from GPU! 470 0\n",
            "Hello World from GPU! 441 0\n",
            "Hello World from GPU! 448 0\n",
            "Hello World from GPU! 463 0\n",
            "Hello World from GPU! 460 0\n",
            "Hello World from GPU! 452 0\n",
            "Hello World from GPU! 437 0\n",
            "Hello World from GPU! 467 0\n",
            "Hello World from GPU! 439 0\n",
            "Hello World from GPU! 462 0\n",
            "Hello World from GPU! 453 0\n",
            "Hello World from GPU! 451 0\n",
            "Hello World from GPU! 447 0\n",
            "Hello World from GPU! 445 0\n",
            "Hello World from GPU! 446 0\n",
            "Hello World from GPU! 444 0\n",
            "Hello World from GPU! 435 0\n",
            "Hello World from GPU! 449 0\n",
            "Hello World from GPU! 455 0\n",
            "Hello World from GPU! 464 0\n",
            "Hello World from GPU! 468 0\n",
            "Hello World from GPU! 471 0\n",
            "Hello World from GPU! 473 0\n",
            "Hello World from GPU! 458 0\n",
            "Hello World from GPU! 454 0\n",
            "Hello World from GPU! 475 0\n",
            "Hello World from GPU! 436 0\n",
            "Hello World from GPU! 477 0\n",
            "Hello World from GPU! 456 0\n",
            "Hello World from GPU! 481 0\n",
            "Hello World from GPU! 504 0\n",
            "Hello World from GPU! 450 0\n",
            "Hello World from GPU! 495 0\n",
            "Hello World from GPU! 476 0\n",
            "Hello World from GPU! 474 0\n",
            "Hello World from GPU! 472 0\n",
            "Hello World from GPU! 480 0\n",
            "Hello World from GPU! 440 0\n",
            "Hello World from GPU! 443 0\n",
            "Hello World from GPU! 506 0\n",
            "Hello World from GPU! 498 0\n",
            "Hello World from GPU! 459 0\n",
            "Hello World from GPU! 519 0\n",
            "Hello World from GPU! 482 0\n",
            "Hello World from GPU! 521 0\n",
            "Hello World from GPU! 494 0\n",
            "Hello World from GPU! 465 0\n",
            "Hello World from GPU! 499 0\n",
            "Hello World from GPU! 469 0\n",
            "Hello World from GPU! 478 0\n",
            "Hello World from GPU! 461 0\n",
            "Hello World from GPU! 488 0\n",
            "Hello World from GPU! 434 0\n",
            "Hello World from GPU! 483 0\n",
            "Hello World from GPU! 490 0\n",
            "Hello World from GPU! 514 0\n",
            "Hello World from GPU! 500 0\n",
            "Hello World from GPU! 497 0\n",
            "Hello World from GPU! 513 0\n",
            "Hello World from GPU! 515 0\n",
            "Hello World from GPU! 508 0\n",
            "Hello World from GPU! 518 0\n",
            "Hello World from GPU! 523 0\n",
            "Hello World from GPU! 520 0\n",
            "Hello World from GPU! 511 0\n",
            "Hello World from GPU! 526 0\n",
            "Hello World from GPU! 457 0\n",
            "Hello World from GPU! 509 0\n",
            "Hello World from GPU! 517 0\n",
            "Hello World from GPU! 485 0\n",
            "Hello World from GPU! 486 0\n",
            "Hello World from GPU! 484 0\n",
            "Hello World from GPU! 527 0\n",
            "Hello World from GPU! 479 0\n",
            "Hello World from GPU! 487 0\n",
            "Hello World from GPU! 510 0\n",
            "Hello World from GPU! 503 0\n",
            "Hello World from GPU! 501 0\n",
            "Hello World from GPU! 532 0\n",
            "Hello World from GPU! 522 0\n",
            "Hello World from GPU! 525 0\n",
            "Hello World from GPU! 537 0\n",
            "Hello World from GPU! 516 0\n",
            "Hello World from GPU! 496 0\n",
            "Hello World from GPU! 492 0\n",
            "Hello World from GPU! 524 0\n",
            "Hello World from GPU! 493 0\n",
            "Hello World from GPU! 512 0\n",
            "Hello World from GPU! 507 0\n",
            "Hello World from GPU! 491 0\n",
            "Hello World from GPU! 502 0\n",
            "Hello World from GPU! 505 0\n",
            "Hello World from GPU! 489 0\n",
            "Hello World from GPU! 530 0\n",
            "Hello World from GPU! 536 0\n",
            "Hello World from GPU! 535 0\n",
            "Hello World from GPU! 538 0\n",
            "Hello World from GPU! 539 0\n",
            "Hello World from GPU! 540 0\n",
            "Hello World from GPU! 543 0\n",
            "Hello World from GPU! 534 0\n",
            "Hello World from GPU! 541 0\n",
            "Hello World from GPU! 545 0\n",
            "Hello World from GPU! 544 0\n",
            "Hello World from GPU! 542 0\n",
            "Hello World from GPU! 547 0\n",
            "Hello World from GPU! 531 0\n",
            "Hello World from GPU! 546 0\n",
            "Hello World from GPU! 528 0\n",
            "Hello World from GPU! 529 0\n",
            "Hello World from GPU! 533 0\n",
            "Hello World from GPU! 549 0\n",
            "Hello World from GPU! 550 0\n",
            "Hello World from GPU! 553 0\n",
            "Hello World from GPU! 554 0\n",
            "Hello World from GPU! 557 0\n",
            "Hello World from GPU! 558 0\n",
            "Hello World from GPU! 561 0\n",
            "Hello World from GPU! 564 0\n",
            "Hello World from GPU! 552 0\n",
            "Hello World from GPU! 572 0\n",
            "Hello World from GPU! 555 0\n",
            "Hello World from GPU! 574 0\n",
            "Hello World from GPU! 573 0\n",
            "Hello World from GPU! 567 0\n",
            "Hello World from GPU! 569 0\n",
            "Hello World from GPU! 548 0\n",
            "Hello World from GPU! 578 0\n",
            "Hello World from GPU! 562 0\n",
            "Hello World from GPU! 576 0\n",
            "Hello World from GPU! 556 0\n",
            "Hello World from GPU! 566 0\n",
            "Hello World from GPU! 580 0\n",
            "Hello World from GPU! 551 0\n",
            "Hello World from GPU! 584 0\n",
            "Hello World from GPU! 563 0\n",
            "Hello World from GPU! 577 0\n",
            "Hello World from GPU! 575 0\n",
            "Hello World from GPU! 565 0\n",
            "Hello World from GPU! 581 0\n",
            "Hello World from GPU! 582 0\n",
            "Hello World from GPU! 585 0\n",
            "Hello World from GPU! 560 0\n",
            "Hello World from GPU! 570 0\n",
            "Hello World from GPU! 586 0\n",
            "Hello World from GPU! 568 0\n",
            "Hello World from GPU! 587 0\n",
            "Hello World from GPU! 591 0\n",
            "Hello World from GPU! 571 0\n",
            "Hello World from GPU! 593 0\n",
            "Hello World from GPU! 588 0\n",
            "Hello World from GPU! 559 0\n",
            "Hello World from GPU! 579 0\n",
            "Hello World from GPU! 583 0\n",
            "Hello World from GPU! 590 0\n",
            "Hello World from GPU! 592 0\n",
            "Hello World from GPU! 589 0\n",
            "Hello World from GPU! 596 0\n",
            "Hello World from GPU! 594 0\n",
            "Hello World from GPU! 595 0\n",
            "Hello World from GPU! 602 0\n",
            "Hello World from GPU! 600 0\n",
            "Hello World from GPU! 604 0\n",
            "Hello World from GPU! 605 0\n",
            "Hello World from GPU! 606 0\n",
            "Hello World from GPU! 603 0\n",
            "Hello World from GPU! 608 0\n",
            "Hello World from GPU! 609 0\n",
            "Hello World from GPU! 599 0\n",
            "Hello World from GPU! 607 0\n",
            "Hello World from GPU! 612 0\n",
            "Hello World from GPU! 614 0\n",
            "Hello World from GPU! 597 0\n",
            "Hello World from GPU! 615 0\n",
            "Hello World from GPU! 619 0\n",
            "Hello World from GPU! 616 0\n",
            "Hello World from GPU! 621 0\n",
            "Hello World from GPU! 620 0\n",
            "Hello World from GPU! 601 0\n",
            "Hello World from GPU! 613 0\n",
            "Hello World from GPU! 611 0\n",
            "Hello World from GPU! 610 0\n",
            "Hello World from GPU! 598 0\n",
            "Hello World from GPU! 617 0\n",
            "Hello World from GPU! 618 0\n",
            "Hello World from GPU! 623 0\n",
            "Hello World from GPU! 622 0\n",
            "Hello World from GPU! 625 0\n",
            "Hello World from GPU! 624 0\n",
            "Hello World from GPU! 627 0\n",
            "Hello World from GPU! 626 0\n",
            "Hello World from GPU! 637 0\n",
            "Hello World from GPU! 628 0\n",
            "Hello World from GPU! 638 0\n",
            "Hello World from GPU! 636 0\n",
            "Hello World from GPU! 639 0\n",
            "Hello World from GPU! 645 0\n",
            "Hello World from GPU! 633 0\n",
            "Hello World from GPU! 646 0\n",
            "Hello World from GPU! 640 0\n",
            "Hello World from GPU! 629 0\n",
            "Hello World from GPU! 644 0\n",
            "Hello World from GPU! 635 0\n",
            "Hello World from GPU! 634 0\n",
            "Hello World from GPU! 642 0\n",
            "Hello World from GPU! 643 0\n",
            "Hello World from GPU! 631 0\n",
            "Hello World from GPU! 641 0\n",
            "Hello World from GPU! 632 0\n",
            "Hello World from GPU! 652 0\n",
            "Hello World from GPU! 653 0\n",
            "Hello World from GPU! 630 0\n",
            "Hello World from GPU! 654 0\n",
            "Hello World from GPU! 661 0\n",
            "Hello World from GPU! 656 0\n",
            "Hello World from GPU! 659 0\n",
            "Hello World from GPU! 651 0\n",
            "Hello World from GPU! 647 0\n",
            "Hello World from GPU! 660 0\n",
            "Hello World from GPU! 658 0\n",
            "Hello World from GPU! 655 0\n",
            "Hello World from GPU! 649 0\n",
            "Hello World from GPU! 650 0\n",
            "Hello World from GPU! 657 0\n",
            "Hello World from GPU! 648 0\n",
            "Hello World from GPU! 667 0\n",
            "Hello World from GPU! 668 0\n",
            "Hello World from GPU! 670 0\n",
            "Hello World from GPU! 669 0\n",
            "Hello World from GPU! 671 0\n",
            "Hello World from GPU! 672 0\n",
            "Hello World from GPU! 686 0\n",
            "Hello World from GPU! 673 0\n",
            "Hello World from GPU! 675 0\n",
            "Hello World from GPU! 678 0\n",
            "Hello World from GPU! 666 0\n",
            "Hello World from GPU! 680 0\n",
            "Hello World from GPU! 683 0\n",
            "Hello World from GPU! 682 0\n",
            "Hello World from GPU! 690 0\n",
            "Hello World from GPU! 714 0\n",
            "Hello World from GPU! 665 0\n",
            "Hello World from GPU! 662 0\n",
            "Hello World from GPU! 693 0\n",
            "Hello World from GPU! 705 0\n",
            "Hello World from GPU! 692 0\n",
            "Hello World from GPU! 700 0\n",
            "Hello World from GPU! 674 0\n",
            "Hello World from GPU! 694 0\n",
            "Hello World from GPU! 684 0\n",
            "Hello World from GPU! 681 0\n",
            "Hello World from GPU! 688 0\n",
            "Hello World from GPU! 704 0\n",
            "Hello World from GPU! 703 0\n",
            "Hello World from GPU! 676 0\n",
            "Hello World from GPU! 664 0\n",
            "Hello World from GPU! 695 0\n",
            "Hello World from GPU! 691 0\n",
            "Hello World from GPU! 663 0\n",
            "Hello World from GPU! 696 0\n",
            "Hello World from GPU! 708 0\n",
            "Hello World from GPU! 685 0\n",
            "Hello World from GPU! 677 0\n",
            "Hello World from GPU! 699 0\n",
            "Hello World from GPU! 706 0\n",
            "Hello World from GPU! 679 0\n",
            "Hello World from GPU! 712 0\n",
            "Hello World from GPU! 719 0\n",
            "Hello World from GPU! 697 0\n",
            "Hello World from GPU! 687 0\n",
            "Hello World from GPU! 716 0\n",
            "Hello World from GPU! 717 0\n",
            "Hello World from GPU! 689 0\n",
            "Hello World from GPU! 709 0\n",
            "Hello World from GPU! 707 0\n",
            "Hello World from GPU! 701 0\n",
            "Hello World from GPU! 710 0\n",
            "Hello World from GPU! 715 0\n",
            "Hello World from GPU! 713 0\n",
            "Hello World from GPU! 698 0\n",
            "Hello World from GPU! 718 0\n",
            "Hello World from GPU! 711 0\n",
            "Hello World from GPU! 702 0\n",
            "Hello World from GPU! 765 0\n",
            "Hello World from GPU! 767 0\n",
            "Hello World from GPU! 769 0\n",
            "Hello World from GPU! 789 0\n",
            "Hello World from GPU! 741 0\n",
            "Hello World from GPU! 758 0\n",
            "Hello World from GPU! 803 0\n",
            "Hello World from GPU! 805 0\n",
            "Hello World from GPU! 751 0\n",
            "Hello World from GPU! 806 0\n",
            "Hello World from GPU! 747 0\n",
            "Hello World from GPU! 749 0\n",
            "Hello World from GPU! 739 0\n",
            "Hello World from GPU! 745 0\n",
            "Hello World from GPU! 736 0\n",
            "Hello World from GPU! 734 0\n",
            "Hello World from GPU! 746 0\n",
            "Hello World from GPU! 748 0\n",
            "Hello World from GPU! 750 0\n",
            "Hello World from GPU! 740 0\n",
            "Hello World from GPU! 743 0\n",
            "Hello World from GPU! 733 0\n",
            "Hello World from GPU! 744 0\n",
            "Hello World from GPU! 752 0\n",
            "Hello World from GPU! 759 0\n",
            "Hello World from GPU! 753 0\n",
            "Hello World from GPU! 756 0\n",
            "Hello World from GPU! 755 0\n",
            "Hello World from GPU! 762 0\n",
            "Hello World from GPU! 754 0\n",
            "Hello World from GPU! 724 0\n",
            "Hello World from GPU! 725 0\n",
            "Hello World from GPU! 742 0\n",
            "Hello World from GPU! 760 0\n",
            "Hello World from GPU! 783 0\n",
            "Hello World from GPU! 730 0\n",
            "Hello World from GPU! 722 0\n",
            "Hello World from GPU! 766 0\n",
            "Hello World from GPU! 770 0\n",
            "Hello World from GPU! 786 0\n",
            "Hello World from GPU! 779 0\n",
            "Hello World from GPU! 771 0\n",
            "Hello World from GPU! 790 0\n",
            "Hello World from GPU! 772 0\n",
            "Hello World from GPU! 738 0\n",
            "Hello World from GPU! 780 0\n",
            "Hello World from GPU! 735 0\n",
            "Hello World from GPU! 788 0\n",
            "Hello World from GPU! 795 0\n",
            "Hello World from GPU! 781 0\n",
            "Hello World from GPU! 774 0\n",
            "Hello World from GPU! 728 0\n",
            "Hello World from GPU! 792 0\n",
            "Hello World from GPU! 793 0\n",
            "Hello World from GPU! 791 0\n",
            "Hello World from GPU! 797 0\n",
            "Hello World from GPU! 723 0\n",
            "Hello World from GPU! 784 0\n",
            "Hello World from GPU! 727 0\n",
            "Hello World from GPU! 801 0\n",
            "Hello World from GPU! 777 0\n",
            "Hello World from GPU! 782 0\n",
            "Hello World from GPU! 768 0\n",
            "Hello World from GPU! 787 0\n",
            "Hello World from GPU! 798 0\n",
            "Hello World from GPU! 778 0\n",
            "Hello World from GPU! 796 0\n",
            "Hello World from GPU! 773 0\n",
            "Hello World from GPU! 757 0\n",
            "Hello World from GPU! 731 0\n",
            "Hello World from GPU! 720 0\n",
            "Hello World from GPU! 764 0\n",
            "Hello World from GPU! 776 0\n",
            "Hello World from GPU! 761 0\n",
            "Hello World from GPU! 785 0\n",
            "Hello World from GPU! 737 0\n",
            "Hello World from GPU! 732 0\n",
            "Hello World from GPU! 763 0\n",
            "Hello World from GPU! 729 0\n",
            "Hello World from GPU! 721 0\n",
            "Hello World from GPU! 800 0\n",
            "Hello World from GPU! 775 0\n",
            "Hello World from GPU! 794 0\n",
            "Hello World from GPU! 726 0\n",
            "Hello World from GPU! 813 0\n",
            "Hello World from GPU! 821 0\n",
            "Hello World from GPU! 827 0\n",
            "Hello World from GPU! 823 0\n",
            "Hello World from GPU! 828 0\n",
            "Hello World from GPU! 812 0\n",
            "Hello World from GPU! 807 0\n",
            "Hello World from GPU! 804 0\n",
            "Hello World from GPU! 811 0\n",
            "Hello World from GPU! 809 0\n",
            "Hello World from GPU! 826 0\n",
            "Hello World from GPU! 820 0\n",
            "Hello World from GPU! 810 0\n",
            "Hello World from GPU! 825 0\n",
            "Hello World from GPU! 819 0\n",
            "Hello World from GPU! 817 0\n",
            "Hello World from GPU! 799 0\n",
            "Hello World from GPU! 818 0\n",
            "Hello World from GPU! 815 0\n",
            "Hello World from GPU! 822 0\n",
            "Hello World from GPU! 802 0\n",
            "Hello World from GPU! 816 0\n",
            "Hello World from GPU! 824 0\n",
            "Hello World from GPU! 814 0\n",
            "Hello World from GPU! 808 0\n",
            "Hello World from GPU! 829 0\n",
            "Hello World from GPU! 830 0\n",
            "Hello World from GPU! 831 0\n",
            "Hello World from GPU! 832 0\n",
            "Hello World from GPU! 835 0\n",
            "Hello World from GPU! 839 0\n",
            "Hello World from GPU! 840 0\n",
            "Hello World from GPU! 833 0\n",
            "Hello World from GPU! 834 0\n",
            "Hello World from GPU! 837 0\n",
            "Hello World from GPU! 841 0\n",
            "Hello World from GPU! 836 0\n",
            "Hello World from GPU! 838 0\n",
            "Hello World from GPU! 846 0\n",
            "Hello World from GPU! 845 0\n",
            "Hello World from GPU! 843 0\n",
            "Hello World from GPU! 844 0\n",
            "Hello World from GPU! 842 0\n",
            "Hello World from GPU! 847 0\n",
            "Hello World from GPU! 849 0\n",
            "Hello World from GPU! 848 0\n",
            "Hello World from GPU! 850 0\n",
            "Hello World from GPU! 853 0\n",
            "Hello World from GPU! 852 0\n",
            "Hello World from GPU! 851 0\n",
            "Hello World from GPU! 855 0\n",
            "Hello World from GPU! 854 0\n",
            "Hello World from GPU! 856 0\n",
            "Hello World from GPU! 857 0\n",
            "Hello World from GPU! 859 0\n",
            "Hello World from GPU! 858 0\n",
            "Hello World from GPU! 861 0\n",
            "Hello World from GPU! 860 0\n",
            "Hello World from GPU! 863 0\n",
            "Hello World from GPU! 862 0\n",
            "Hello World from GPU! 876 0\n",
            "Hello World from GPU! 902 0\n",
            "Hello World from GPU! 882 0\n",
            "Hello World from GPU! 903 0\n",
            "Hello World from GPU! 889 0\n",
            "Hello World from GPU! 884 0\n",
            "Hello World from GPU! 865 0\n",
            "Hello World from GPU! 872 0\n",
            "Hello World from GPU! 891 0\n",
            "Hello World from GPU! 864 0\n",
            "Hello World from GPU! 868 0\n",
            "Hello World from GPU! 880 0\n",
            "Hello World from GPU! 883 0\n",
            "Hello World from GPU! 871 0\n",
            "Hello World from GPU! 885 0\n",
            "Hello World from GPU! 873 0\n",
            "Hello World from GPU! 904 0\n",
            "Hello World from GPU! 897 0\n",
            "Hello World from GPU! 867 0\n",
            "Hello World from GPU! 874 0\n",
            "Hello World from GPU! 881 0\n",
            "Hello World from GPU! 879 0\n",
            "Hello World from GPU! 886 0\n",
            "Hello World from GPU! 895 0\n",
            "Hello World from GPU! 890 0\n",
            "Hello World from GPU! 888 0\n",
            "Hello World from GPU! 892 0\n",
            "Hello World from GPU! 899 0\n",
            "Hello World from GPU! 905 0\n",
            "Hello World from GPU! 866 0\n",
            "Hello World from GPU! 907 0\n",
            "Hello World from GPU! 906 0\n",
            "Hello World from GPU! 875 0\n",
            "Hello World from GPU! 877 0\n",
            "Hello World from GPU! 908 0\n",
            "Hello World from GPU! 896 0\n",
            "Hello World from GPU! 887 0\n",
            "Hello World from GPU! 920 0\n",
            "Hello World from GPU! 878 0\n",
            "Hello World from GPU! 912 0\n",
            "Hello World from GPU! 869 0\n",
            "Hello World from GPU! 893 0\n",
            "Hello World from GPU! 901 0\n",
            "Hello World from GPU! 870 0\n",
            "Hello World from GPU! 923 0\n",
            "Hello World from GPU! 900 0\n",
            "Hello World from GPU! 922 0\n",
            "Hello World from GPU! 894 0\n",
            "Hello World from GPU! 911 0\n",
            "Hello World from GPU! 924 0\n",
            "Hello World from GPU! 918 0\n",
            "Hello World from GPU! 910 0\n",
            "Hello World from GPU! 898 0\n",
            "Hello World from GPU! 917 0\n",
            "Hello World from GPU! 909 0\n",
            "Hello World from GPU! 914 0\n",
            "Hello World from GPU! 915 0\n",
            "Hello World from GPU! 913 0\n",
            "Hello World from GPU! 916 0\n",
            "Hello World from GPU! 919 0\n",
            "Hello World from GPU! 937 0\n",
            "Hello World from GPU! 928 0\n",
            "Hello World from GPU! 939 0\n",
            "Hello World from GPU! 929 0\n",
            "Hello World from GPU! 936 0\n",
            "Hello World from GPU! 941 0\n",
            "Hello World from GPU! 934 0\n",
            "Hello World from GPU! 940 0\n",
            "Hello World from GPU! 942 0\n",
            "Hello World from GPU! 944 0\n",
            "Hello World from GPU! 945 0\n",
            "Hello World from GPU! 938 0\n",
            "Hello World from GPU! 935 0\n",
            "Hello World from GPU! 930 0\n",
            "Hello World from GPU! 933 0\n",
            "Hello World from GPU! 931 0\n",
            "Hello World from GPU! 932 0\n",
            "Hello World from GPU! 943 0\n",
            "Hello World from GPU! 925 0\n",
            "Hello World from GPU! 926 0\n",
            "Hello World from GPU! 921 0\n",
            "Hello World from GPU! 927 0\n",
            "Hello World from GPU! 964 0\n",
            "Hello World from GPU! 982 0\n",
            "Hello World from GPU! 946 0\n",
            "Hello World from GPU! 984 0\n",
            "Hello World from GPU! 961 0\n",
            "Hello World from GPU! 967 0\n",
            "Hello World from GPU! 988 0\n",
            "Hello World from GPU! 995 0\n",
            "Hello World from GPU! 998 0\n",
            "Hello World from GPU! 955 0\n",
            "Hello World from GPU! 954 0\n",
            "Hello World from GPU! 969 0\n",
            "Hello World from GPU! 996 0\n",
            "Hello World from GPU! 972 0\n",
            "Hello World from GPU! 1001 0\n",
            "Hello World from GPU! 997 0\n",
            "Hello World from GPU! 1010 0\n",
            "Hello World from GPU! 958 0\n",
            "Hello World from GPU! 956 0\n",
            "Hello World from GPU! 980 0\n",
            "Hello World from GPU! 965 0\n",
            "Hello World from GPU! 986 0\n",
            "Hello World from GPU! 950 0\n",
            "Hello World from GPU! 974 0\n",
            "Hello World from GPU! 971 0\n",
            "Hello World from GPU! 973 0\n",
            "Hello World from GPU! 985 0\n",
            "Hello World from GPU! 962 0\n",
            "Hello World from GPU! 981 0\n",
            "Hello World from GPU! 1006 0\n",
            "Hello World from GPU! 947 0\n",
            "Hello World from GPU! 975 0\n",
            "Hello World from GPU! 1002 0\n",
            "Hello World from GPU! 1016 0\n",
            "Hello World from GPU! 1015 0\n",
            "Hello World from GPU! 948 0\n",
            "Hello World from GPU! 1004 0\n",
            "Hello World from GPU! 952 0\n",
            "Hello World from GPU! 968 0\n",
            "Hello World from GPU! 999 0\n",
            "Hello World from GPU! 977 0\n",
            "Hello World from GPU! 991 0\n",
            "Hello World from GPU! 966 0\n",
            "Hello World from GPU! 978 0\n",
            "Hello World from GPU! 953 0\n",
            "Hello World from GPU! 970 0\n",
            "Hello World from GPU! 1018 0\n",
            "Hello World from GPU! 957 0\n",
            "Hello World from GPU! 1008 0\n",
            "Hello World from GPU! 1017 0\n",
            "Hello World from GPU! 976 0\n",
            "Hello World from GPU! 1012 0\n",
            "Hello World from GPU! 1019 0\n",
            "Hello World from GPU! 1000 0\n",
            "Hello World from GPU! 1003 0\n",
            "Hello World from GPU! 1023 0\n",
            "Hello World from GPU! 960 0\n",
            "Hello World from GPU! 951 0\n",
            "Hello World from GPU! 959 0\n",
            "Hello World from GPU! 1014 0\n",
            "Hello World from GPU! 990 0\n",
            "Hello World from GPU! 1022 0\n",
            "Hello World from GPU! 1021 0\n",
            "Hello World from GPU! 979 0\n",
            "Hello World from GPU! 989 0\n",
            "Hello World from GPU! 1007 0\n",
            "Hello World from GPU! 992 0\n",
            "Hello World from GPU! 1009 0\n",
            "Hello World from GPU! 993 0\n",
            "Hello World from GPU! 1020 0\n",
            "Hello World from GPU! 949 0\n",
            "Hello World from GPU! 1011 0\n",
            "Hello World from GPU! 983 0\n",
            "Hello World from GPU! 1005 0\n",
            "Hello World from GPU! 963 0\n",
            "Hello World from GPU! 987 0\n",
            "Hello World from GPU! 994 0\n",
            "Hello World from GPU! 1013 0\n",
            "\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "KM_1jViWn-hL",
        "colab_type": "text"
      },
      "source": [
        "Resposta. Executa de forma totalmente imprevisível. Não é sequencial. Não é determinística. Não é agrupada."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "zDgMNO1fbaAE",
        "colab_type": "text"
      },
      "source": [
        "3. Compile com capability 3.5 e gere 2 trilhões de threads - que é o máximo que a GPU poderá executar -, em qualquer combinação de blocos x threads. Use um comando condicional para imprimir apenas alguns threads."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "g78X2njRoZ6l",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 34
        },
        "outputId": "5f7bade1-c065-4330-8cd0-fb0a94b53890"
      },
      "source": [
        "%%writefile hello13.cu\n",
        "#include <stdio.h>\n",
        "#include <stdlib.h>\n",
        "\n",
        "\n",
        "__global__ void helloFromGPU() {\n",
        "  if (blockIdx.x == (1000000000-1) && threadIdx.x == (1000-1))\n",
        "    printf(\"Hello World from GPU! Block X: %d, Thread X: %d\\n\",\n",
        "           blockIdx.x, threadIdx.x);\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"Hello World from CPU!\\n\");\n",
        "  helloFromGPU<<<1000000000, 1000>>>();\n",
        "  cudaDeviceReset();\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Overwriting hello13.cu\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "XT51b4NXo1Ck",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 67
        },
        "outputId": "20d4df91-9bb1-42f6-df2c-ba01089b8d0a"
      },
      "source": [
        "%cd /content/src\n",
        "!nvcc -O2 -arch=sm_35 -o hello13 hello13.cu\n",
        "!./hello13"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "Hello World from CPU!\n",
            "Hello World from GPU! Block X: 999999999, Thread X: 999\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "2Pk9l5UDbdZ-",
        "colab_type": "text"
      },
      "source": [
        "4. Execute alguma operação à sua escolha na GPU, mostrando resultados, e relate-os."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "oYgoz4Brfz54",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 67
        },
        "outputId": "b0fe7aa9-545e-416d-b402-65c72f7c9af8"
      },
      "source": [
        "%%cu\n",
        "#include <stdio.h>\n",
        "#include <stdlib.h>\n",
        "\n",
        "#define CHECK(call) { \\\n",
        "  const cudaError_t error = call; \\\n",
        "  if (error != cudaSuccess) { \\\n",
        "    printf(\"Error: %s:%d, \", __FILE__, __LINE__); \\\n",
        "    printf(\"code:%d, reason: %s\\n\", error, cudaGetErrorString(error)); \\\n",
        "    exit(1); \\\n",
        "  } \\\n",
        "}\n",
        "\n",
        "__global__ void helloFromGPU() {\n",
        "  printf(\"Olá mundo do heavy metal, aqui e a GPU! %d %d\\n\", blockIdx.x, threadIdx.x);\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"Hello World from CPU!\\n\");\n",
        "  helloFromGPU<<<1, 1>>>();\n",
        "  CHECK(cudaDeviceReset());\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Hello World from CPU!\n",
            "Olá mundo do heavy metal, aqui e a GPU! 0 0\n",
            "\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "prtP6LP5bjTm",
        "colab_type": "text"
      },
      "source": [
        "## Exemplo 2 // Soma de Vetores\n",
        "\n",
        "_Referência: Capítulo 2 - CUDA Programming Model_\n",
        "\n",
        "O próximo exemplo já usa o modelo cunhado pela NVIDIA, o _Single Instruction Multiple Thread_ (SIMT), para realizar um cálculo simples em um grande conjunto de dados. Faremos a soma de dois vetores e armazená-la em outro.\n",
        "\n",
        "A forma convencional de se fazer isso na CPU seria\n",
        "``` cpp\n",
        "void sumArraysOnHost(float *A, float *B, float *C, int N) {\n",
        "  for (int i = 0; i < N; i++)\n",
        "    C[i] = A[i] + B[i];\n",
        "}\n",
        "```\n",
        "Porém, com a GPU, vamos desfazer o _loop_ e realizar a soma usando o ID de bloco e thread de cada thread para selecionar qual posição do vetor será somada.\n",
        "\n",
        "Para entender como fazer esse mapeamento entre threads e posições do vetor, é necessário entender a hierarquia das threads. A figura a seguir, extraída do livro, resume bem essa ideia.\n",
        "\n",
        "![Hierarquia de threads](https://docs.google.com/uc?export=download&id=13qryIegnVrDgmaKA7eEcV3P03ddRe3MS)\n",
        "\n",
        "Os blocos são organizados numa grade, e cada bloco contém várias threads. Grades e blocos podem ter até três dimensões, mas geralmente as grades são 2D e os blocos são 3D. Na figura, temos uma grade 2D e um bloco também 2D.\n",
        "\n",
        "Neste exemplo, no entanto, vamos usar uma configuração 1D, pois nossos vetores são unidimensionais. A figura abaixo ilustra essa configuração:\n",
        "\n",
        "![Configuraçãao 1D de threads](https://docs.google.com/uc?export=download&id=1BA6xd8T--b0Y4YaUEa-O5HPiteCceohk)\n",
        "\n",
        "Esses ID's de bloco e de thread são acessados dentro de um _kernel_ com as variáveis vistas no exemplo 1, `blockIdx` e `threadIdx`, respectivamente. Como estamos interessados apenas em uma dimensão, selecionamos apenas a dimensão `x` - as outras dimensões são `y` e `z`. Além dessas, também temos as variáveis `gridDim` e `blockDim`, que informam o tamanho da grade e do bloco, respectivamente, e também tem os campos `x`, `y` e `z`. O código da soma ficaria então:\n",
        "\n",
        "``` cpp\n",
        "__global__ void sumArraysOnGPU(float *A, float *B, float *C) {\n",
        "  int i = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "  C[i] = A[i] + B[i];\n",
        "}\n",
        "```\n",
        "Repare também que no código não há referência ao tamanho dos vetores, já que o valor N é implicitamente definido pela quantidade de blocos e threads usadas para a soma. Por exemplo, se o vetor tiver tamanho 32, você poderia invocar o _kernel_ com qualquer uma das seguintes linhas, dentre outras:\n",
        "``` cpp\n",
        "sumArraysOnGPU<<<1,32>>>(A, B, C);\n",
        "sumArraysOnGPU<<<4,8>>>(A, B, C); // corresponde à figura\n",
        "sumArraysOnGPU<<<16,2>>>(A, B, C);\n",
        "```\n",
        "Para colocar esse _kernel_ em utilização, é necessário usar outras funções da API do CUDA. Serão mostrados comentários explicando em alto nível o que cada uma faz. Note que para usá-las é necessário incluir a biblioteca `cuda_runtime.h`."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "kuTVNh4E-xVK",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 118
        },
        "outputId": "c654e5df-db19-421c-e8bd-6af4734aab7f"
      },
      "source": [
        "%%cu\n",
        "#include <cuda_runtime.h>\n",
        "#include <stdio.h>\n",
        "\n",
        "// Use esta função para conferir se as somas na CPU e na GPU correspondem\n",
        "void checkResult(float *hostRef, float *gpuRef, const int N) {\n",
        "  double epsilon = 1.0E-8;\n",
        "  for (int i=0; i<N; i++) {\n",
        "    if (abs(hostRef[i] - gpuRef[i]) > epsilon) {\n",
        "      printf(\"Arrays do not match!\\n\");\n",
        "      printf(\"host %5.2f gpu %5.2f at current %d\\n\", hostRef[i], gpuRef[i], i);\n",
        "      return;\n",
        "    }\n",
        "  }\n",
        "  printf(\"Arrays match.\\n\\n\");\n",
        "}\n",
        "\n",
        "// Use esta função para atribuir valores iniciais aos vetores de floats\n",
        "void initialData(float *a, int n) {\n",
        "  srand(time(NULL));\n",
        "  for (int i = 0; i < n; i++)\n",
        "    a[i] = (float)(rand() & 0xFF) / 10.0;\n",
        "}\n",
        "\n",
        "// Soma os vetores na CPU\n",
        "void sumArraysOnHost(float *A, float *B, float *C, const int N) {\n",
        "  for (int i = 0; i < N; i++)\n",
        "    C[i] = A[i] + B[i];\n",
        "}\n",
        "\n",
        "// Soma os vetores na GPU\n",
        "__global__ void sumArraysOnGPU(float *A, float *B, float *C) {\n",
        "  int i = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "  C[i] = A[i] + B[i];\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"%s Starting...\\n\", argv[0]);\n",
        "  \n",
        "  // Configura o tamanho dos vetores\n",
        "  int nElem = 2048;\n",
        "  printf(\"Vector size %d\\n\", nElem);\n",
        "  \n",
        "  // Aloca memória no host\n",
        "  size_t nBytes = nElem * sizeof(float);\n",
        "  float *h_A, *h_B, *hostRef, *gpuRef;\n",
        "  h_A = (float *)malloc(nBytes);\n",
        "  h_B = (float *)malloc(nBytes);\n",
        "  hostRef = (float *)malloc(nBytes);\n",
        "  gpuRef = (float *)malloc(nBytes);\n",
        "\n",
        "  // Inicializa os dados ainda no host\n",
        "  initialData(h_A, nElem);\n",
        "  initialData(h_B, nElem);\n",
        "  \n",
        "  // Aloca memória global na GPU (device)\n",
        "  float *d_A, *d_B, *d_C;\n",
        "  cudaMalloc((float **)&d_A, nBytes);\n",
        "  cudaMalloc((float **)&d_B, nBytes);\n",
        "  cudaMalloc((float **)&d_C, nBytes);\n",
        "\n",
        "  // Transfere os dados do host (CPU) para o device (GPU)\n",
        "  cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);\n",
        "  cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);\n",
        "\n",
        "  // Configura os tamanho de:\n",
        "  //   bloco (quantas threads por bloco?)\n",
        "  //   e grade (quantos blocos por grade?)\n",
        "  dim3 block(1024);\n",
        "  dim3 grid(nElem >> 10);\n",
        "\n",
        "  // Executa a soma no device\n",
        "  sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C);\n",
        "  printf(\"Execution configuration <<<%d, %d>>>\\n\", grid.x, block.x);\n",
        "\n",
        "  // Copia o resultado do kernel do device para o host\n",
        "  cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);\n",
        "\n",
        "  // Soma os vetores na CPU para conferir os resultados\n",
        "  sumArraysOnHost(h_A, h_B, hostRef, nElem);\n",
        "  // Compara os resultados\n",
        "  checkResult(hostRef, gpuRef, nElem);\n",
        "\n",
        "  // Libera a memória da GPU\n",
        "  cudaFree(d_A);\n",
        "  cudaFree(d_B);\n",
        "  cudaFree(d_C);\n",
        "\n",
        "  // Libera a memória da CPU\n",
        "  free(h_A);\n",
        "  free(h_B);\n",
        "  free(hostRef);\n",
        "  free(gpuRef);\n",
        "\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/tmp/tmpg1_ooyzm/3057d6bc-4fd9-49e5-a154-1507e8e43700.out Starting...\n",
            "Vector size 2048\n",
            "Execution configuration <<<2, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "lMMsbAYWUXrp",
        "colab_type": "text"
      },
      "source": [
        "Observação: ao modificar este exemplo, tome cuidado para não usar uma configuração de tamanho de bloco que exceda 1024 threads por bloco. Aqui no Google Colab, os kernels não são chamados quando isso acontece. Para usar mais threads, use mais blocos por grade, fazendo os cálculos adequados."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "pioCLy5bPH-k",
        "colab_type": "text"
      },
      "source": [
        "## Tarefa 2 // Soma de Vetores\n",
        "\n",
        "Para realizar esta tarefa, você vai usar a ferramenta `nvprof`. Para usá-la, você precisa compilar seu programa com o `nvcc` - como foi descrito anteriormente - e passar o nome do executável gerado para o `nvprof`. Por exemplo, se você salvar o código fonte do exemplo de soma de vetores como `/content/src/sum.cu`, fica assim o uso do `nvprof`:"
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "4x30OloZhBul",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 34
        },
        "outputId": "0d98b00a-5bb2-46d5-981f-8c1e970504bf"
      },
      "source": [
        "%%writefile /content/src/sum.cu\n",
        "#include <cuda_runtime.h>\n",
        "#include <stdio.h>\n",
        "\n",
        "// Use esta função para conferir se as somas na CPU e na GPU correspondem\n",
        "void checkResult(float *hostRef, float *gpuRef, const int N) {\n",
        "  double epsilon = 1.0E-8;\n",
        "  for (int i=0; i<N; i++) {\n",
        "    if (abs(hostRef[i] - gpuRef[i]) > epsilon) {\n",
        "      printf(\"Arrays do not match!\\n\");\n",
        "      printf(\"host %5.2f gpu %5.2f at current %d\\n\", hostRef[i], gpuRef[i], i);\n",
        "      return;\n",
        "    }\n",
        "  }\n",
        "  printf(\"Arrays match.\\n\\n\");\n",
        "}\n",
        "\n",
        "// Use esta função para atribuir valores iniciais aos vetores de floats\n",
        "void initialData(float *a, int n) {\n",
        "  srand(time(NULL));\n",
        "  for (int i = 0; i < n; i++)\n",
        "    a[i] = (float)(rand() & 0xFF) / 10.0;\n",
        "}\n",
        "\n",
        "// Soma os vetores na CPU\n",
        "void sumArraysOnHost(float *A, float *B, float *C, const int N) {\n",
        "  for (int i = 0; i < N; i++)\n",
        "    C[i] = A[i] + B[i];\n",
        "}\n",
        "\n",
        "// Soma os vetores na GPU\n",
        "__global__ void sumArraysOnGPU(float *A, float *B, float *C) {\n",
        "  int i = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "  C[i] = A[i] + B[i];\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"%s Starting...\\n\", argv[0]);\n",
        "  \n",
        "  // Configura o tamanho dos vetores\n",
        "  int nElem = atoi(argv[1]);\n",
        "  printf(\"Vector size %d\\n\", nElem);\n",
        "  \n",
        "  // Aloca memória no host\n",
        "  size_t nBytes = nElem * sizeof(float);\n",
        "  float *h_A, *h_B, *hostRef, *gpuRef;\n",
        "  h_A = (float *)malloc(nBytes);\n",
        "  h_B = (float *)malloc(nBytes);\n",
        "  hostRef = (float *)malloc(nBytes);\n",
        "  gpuRef = (float *)malloc(nBytes);\n",
        "\n",
        "  // Inicializa os dados ainda no host\n",
        "  initialData(h_A, nElem);\n",
        "  initialData(h_B, nElem);\n",
        "  \n",
        "  // Aloca memória global na GPU (device)\n",
        "  float *d_A, *d_B, *d_C;\n",
        "  cudaMalloc((float **)&d_A, nBytes);\n",
        "  cudaMalloc((float **)&d_B, nBytes);\n",
        "  cudaMalloc((float **)&d_C, nBytes);\n",
        "\n",
        "  // Transfere os dados do host (CPU) para o device (GPU)\n",
        "  cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);\n",
        "  cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);\n",
        "\n",
        "  // Configura os tamanho de:\n",
        "  //   bloco (quantas threads por bloco?)\n",
        "  //   e grade (quantos blocos por grade?)\n",
        "  dim3 block(1024);\n",
        "  dim3 grid(nElem >> 10);\n",
        "\n",
        "  // Executa a soma no device\n",
        "  sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C);\n",
        "  printf(\"Execution configuration <<<%d, %d>>>\\n\", grid.x, block.x);\n",
        "\n",
        "  // Copia o resultado do kernel do device para o host\n",
        "  cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);\n",
        "\n",
        "  // Soma os vetores na CPU para conferir os resultados\n",
        "  sumArraysOnHost(h_A, h_B, hostRef, nElem);\n",
        "  // Compara os resultados\n",
        "  checkResult(hostRef, gpuRef, nElem);\n",
        "\n",
        "  // Libera a memória da GPU\n",
        "  cudaFree(d_A);\n",
        "  cudaFree(d_B);\n",
        "  cudaFree(d_C);\n",
        "\n",
        "  // Libera a memória da CPU\n",
        "  free(h_A);\n",
        "  free(h_B);\n",
        "  free(hostRef);\n",
        "  free(gpuRef);\n",
        "\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Overwriting /content/src/sum.cu\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "jAf6z235UEmw",
        "colab_type": "text"
      },
      "source": [
        "O `nvprof` vai exibir um relatório completo dos tempos de execução de cada chamada à API do CUDA e de cada kernel. Nessa tarefa você vai explorar como avaliar os tempos com o `nvprof`.\n",
        "\n",
        "1. Compile e execute o exemplo de soma de vetores como está, e veja o desempenho com o `nvprof`. Quanto tempo foi gasto em transferências _host to device_ e _device to host_? Qual foi a taxa de transferência, em elementos do vetor por unidade de tempo? Qual foi o tempo de execução do kernel?"
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "_Q7eUxDTTNHb",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 420
        },
        "outputId": "3d1bbaaf-c6bc-4e95-a523-80fea3733548"
      },
      "source": [
        "%cd /content/src\n",
        "!nvcc -O2 -arch=sm_35 -o sum sum.cu\n",
        "!nvprof ./sum 2048"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "./sum Starting...\n",
            "Vector size 2048\n",
            "==13320== NVPROF is profiling process 13320, command: ./sum 2048\n",
            "Execution configuration <<<2, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==13320== Profiling application: ./sum 2048\n",
            "==13320== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   47.47%  6.9120us         2  3.4560us  3.3920us  3.5200us  [CUDA memcpy HtoD]\n",
            "                   26.59%  3.8720us         1  3.8720us  3.8720us  3.8720us  [CUDA memcpy DtoH]\n",
            "                   25.93%  3.7760us         1  3.7760us  3.7760us  3.7760us  sumArraysOnGPU(float*, float*, float*)\n",
            "      API calls:   98.99%  130.23ms         3  43.412ms  4.0660us  130.22ms  cudaMalloc\n",
            "                    0.45%  594.30us         1  594.30us  594.30us  594.30us  cuDeviceTotalMem\n",
            "                    0.26%  341.29us        96  3.5550us     133ns  151.10us  cuDeviceGetAttribute\n",
            "                    0.17%  225.93us         3  75.309us  5.7220us  132.43us  cudaFree\n",
            "                    0.07%  92.332us         3  30.777us  18.664us  51.660us  cudaMemcpy\n",
            "                    0.03%  38.878us         1  38.878us  38.878us  38.878us  cudaLaunchKernel\n",
            "                    0.02%  28.120us         1  28.120us  28.120us  28.120us  cuDeviceGetName\n",
            "                    0.00%  2.9640us         1  2.9640us  2.9640us  2.9640us  cuDeviceGetPCIBusId\n",
            "                    0.00%  1.8610us         3     620ns     197ns  1.0680us  cuDeviceGetCount\n",
            "                    0.00%  1.6560us         2     828ns     258ns  1.3980us  cuDeviceGet\n",
            "                    0.00%     306ns         1     306ns     306ns     306ns  cuDeviceGetUuid\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "1iMZqfGShqWD",
        "colab_type": "text"
      },
      "source": [
        "*   Host para device: 7.2960us\n",
        "*   Device para host: 3.6480us\n",
        "*   Tempo de transferência total: 10.944us\n",
        "*   Elementos do vetor: 2048\n",
        "*   Taxa de transferêcia: 188 elemetos/us\n",
        "*   Tempo de execução do kernel: 14.449us\n",
        "\n",
        "\n",
        "  "
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "KdojjLqFWfkG",
        "colab_type": "text"
      },
      "source": [
        "2. Modifique o tamanho dos vetores para 100 M, 200 M e 300 M elementos. Repita para cada tamanho as análises de tempos e taxas do exercício anterior."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "colab_type": "code",
        "id": "kVgk8QfUjKxy",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 235
        },
        "outputId": "8dad2807-050d-4a1d-955c-9d04bf658a3e"
      },
      "source": [
        "%cd /content/src\n",
        "!nvprof --print-gpu-summary ./sum 128000000"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "./sum Starting...\n",
            "Vector size 128000000\n",
            "==13343== NVPROF is profiling process 13343, command: ./sum 128000000\n",
            "Execution configuration <<<125000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==13343== Profiling application: ./sum 128000000\n",
            "==13343== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   65.49%  305.59ms         1  305.59ms  305.59ms  305.59ms  [CUDA memcpy DtoH]\n",
            "                   31.82%  148.50ms         2  74.248ms  71.194ms  77.302ms  [CUDA memcpy HtoD]\n",
            "                    2.69%  12.533ms         1  12.533ms  12.533ms  12.533ms  sumArraysOnGPU(float*, float*, float*)\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "colab_type": "text",
        "id": "L1HxWpuukCbr"
      },
      "source": [
        "*   Host para device: 250.54ms\n",
        "*   Device para host: 109.06ms\n",
        "*   Tempo de transferência total: 359,6ms\n",
        "*   Elementos do vetor: 100000000\n",
        "*   Taxa de transferêcia: 278087 elementos/ms\n",
        "*   Tempo de execução do kernel: 369.4ms\n",
        "\n",
        "\n",
        "  "
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "colab_type": "code",
        "id": "zeep0TIJkMwG",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 235
        },
        "outputId": "882d1f0a-227c-414d-d4bc-62c190778b15"
      },
      "source": [
        "%cd /content/src\n",
        "!nvprof --print-gpu-summary ./sum 256000000"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "./sum Starting...\n",
            "Vector size 256000000\n",
            "==13366== NVPROF is profiling process 13366, command: ./sum 256000000\n",
            "Execution configuration <<<250000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==13366== Profiling application: ./sum 256000000\n",
            "==13366== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   66.61%  641.91ms         1  641.91ms  641.91ms  641.91ms  [CUDA memcpy DtoH]\n",
            "                   30.78%  296.59ms         2  148.30ms  145.15ms  151.44ms  [CUDA memcpy HtoD]\n",
            "                    2.61%  25.166ms         1  25.166ms  25.166ms  25.166ms  sumArraysOnGPU(float*, float*, float*)\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "colab_type": "text",
        "id": "GdWHyhyHkMwN"
      },
      "source": [
        "*   Host para device: 484.20ms\n",
        "*   Device para host: 231.46ms\n",
        "*   Tempo de transferência total: 715.66ms\n",
        "*   Elementos do vetor: 200000000\n",
        "*   Taxa de transferêcia: 279463 elementos/ms\n",
        "*   Tempo de execução do kernel: 735.29ms\n",
        "\n",
        "\n",
        "  "
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "colab_type": "code",
        "id": "3mf1UHC5kMTB",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 235
        },
        "outputId": "14b7e66c-7bcd-4431-b3eb-89a6b44faa0e"
      },
      "source": [
        "%cd /content/src\n",
        "!nvprof --print-gpu-summary ./sum 512000000"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "./sum Starting...\n",
            "Vector size 512000000\n",
            "==13400== NVPROF is profiling process 13400, command: ./sum 512000000\n",
            "Execution configuration <<<500000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==13400== Profiling application: ./sum 512000000\n",
            "==13400== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   65.50%  1.25264s         1  1.25264s  1.25264s  1.25264s  [CUDA memcpy DtoH]\n",
            "                   31.89%  609.83ms         2  304.91ms  302.92ms  306.90ms  [CUDA memcpy HtoD]\n",
            "                    2.62%  50.019ms         1  50.019ms  50.019ms  50.019ms  sumArraysOnGPU(float*, float*, float*)\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "colab_type": "text",
        "id": "qGBjOK10kMTE"
      },
      "source": [
        "*   Host para device: 736.33ms\n",
        "*   Device para host: 338.87ms\n",
        "*   Tempo de transferência total: 1.075s\n",
        "*   Elementos do vetor: 300000000\n",
        "*   Taxa de transferêcia: 279069 elementos/ms\n",
        "*   Tempo de execução do kernel: 1.104s\n",
        "\n",
        "\n",
        "  "
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "dyUf0uToj5Rb",
        "colab_type": "text"
      },
      "source": [
        ""
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "-HYzK31GY-qY",
        "colab_type": "text"
      },
      "source": [
        "## Exemplo 3 // Unroll\n",
        "\n",
        "_Referência: Capítulo 3 - CUDA Execution Model - Unrolling Loops_\n",
        "\n",
        "Vamos avaliar agora como o desempenho muda quando fazemos com que uma mesma thread calcule mais de um elemento do vetor. Chamamos essa técnica de _unroll_. Se uma thread é responsável por calcular K elementos do vetor, então dizemos que estamos usando _fator de unroll K_.\n",
        "\n",
        "Para facilitar os cálculos, vamos usar o fator de unroll sempre na forma de potência de 2. Assim, por exemplo, para um fator de unroll 8, vamos armazenar o valor 3, pois 2^3 = 8.\n",
        "\n",
        "Se colocarmos o valor do fator de unroll numa variável chamada `unroll`, uma possibilidade de estrutrar o kernel e sua configuração de execução seria da seguinte forma: (perceba que o número total de threads será menor)\n",
        "```cpp\n",
        "__global__ void sumArraysOnGPU(float *A, float *B, float *C, int unroll) {\n",
        "  int i = (blockIdx.x * blockDim.x + threadIdx.x) << unroll;\n",
        "  int k = i + (1 << unroll);\n",
        "  for (; i < k; i++)\n",
        "    C[i] = A[i] + B[i];\n",
        "}\n",
        "\n",
        "// Chamando no main...\n",
        "dim3 block(1024 >> unroll);\n",
        "dim3 grid(nElem >> 10);\n",
        "sumArraysOnGPU<<<grid, block>>>(A, B, C, unroll);\n",
        "```\n",
        "Dessa forma, suponha que configuramos `unroll` para 3 (2^3 = 8) e `nElem` para 4 K elementos. Teremos cada thread processando 8 elementos do vetor. Uma thread de ID 50 num bloco cujo ID é 2, terá seu valor inicial de i dado por `(2 * 4 + 50) << 3 = 58 << 3 = 464`. Essa thread vai então calcular as posições 464, 465, 466, 467, 468, 469, 470 e 471."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "DEbXfFck1Pap",
        "colab_type": "text"
      },
      "source": [
        "## Tarefa 3 // Unroll\n",
        "\n",
        "1. Adapte o código da soma de vetores para usar o fator de unroll. Parametrize o número de elementos no vetor, o número de threads por bloco e o fator de unroll, de forma que você possa passar esses valores como argumentos de linha de comando, como segue:\n",
        "```bash\n",
        "nvprof ./sum <nElem> <nThreads> <unroll>\n",
        "```\n",
        "Use potências de 2 em todos os parâmetros."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "colab_type": "code",
        "id": "jiUICeYrlTZR",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 34
        },
        "outputId": "a656f1f0-9e23-4d8e-a5cd-d6164cb13fb7"
      },
      "source": [
        "%%writefile /content/src/sumUnroll.cu\n",
        "#include <cuda_runtime.h>\n",
        "#include <stdio.h>\n",
        "\n",
        "// Use esta função para conferir se as somas na CPU e na GPU correspondem\n",
        "void checkResult(float *hostRef, float *gpuRef, const int N) {\n",
        "  double epsilon = 1.0E-8;\n",
        "  for (int i=0; i<N; i++) {\n",
        "    if (abs(hostRef[i] - gpuRef[i]) > epsilon) {\n",
        "      printf(\"Arrays do not match!\\n\");\n",
        "      printf(\"host %5.2f gpu %5.2f at current %d\\n\", hostRef[i], gpuRef[i], i);\n",
        "      return;\n",
        "    }\n",
        "  }\n",
        "  printf(\"Arrays match.\\n\\n\");\n",
        "}\n",
        "\n",
        "// Use esta função para atribuir valores iniciais aos vetores de floats\n",
        "void initialData(float *a, int n) {\n",
        "  srand(time(NULL));\n",
        "  for (int i = 0; i < n; i++)\n",
        "    a[i] = (float)(rand() & 0xFF) / 10.0;\n",
        "}\n",
        "\n",
        "// Soma os vetores na CPU\n",
        "void sumArraysOnHost(float *A, float *B, float *C, const int N) {\n",
        "  for (int i = 0; i < N; i++)\n",
        "    C[i] = A[i] + B[i];\n",
        "}\n",
        "\n",
        "__global__ void sumArraysOnGPU(float *A, float *B, float *C, int unroll) {\n",
        "  int i = (blockIdx.x * blockDim.x + threadIdx.x) << unroll;\n",
        "  int k = i + (1 << unroll);\n",
        "  for (; i < k; i++)\n",
        "    C[i] = A[i] + B[i];\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"%s Starting...\\n\", argv[0]);\n",
        "  \n",
        "  // Configura o tamanho dos vetores\n",
        "  int nElem = atoi(argv[1]);\n",
        "  int nThread = atoi(argv[2]);\n",
        "  int unroll = atoi(argv[3]);\n",
        "  printf(\"Vector size %d\\n\", nElem);\n",
        "  \n",
        "  // Aloca memória no host\n",
        "  size_t nBytes = nElem * sizeof(float);\n",
        "  float *h_A, *h_B, *hostRef, *gpuRef;\n",
        "  h_A = (float *)malloc(nBytes);\n",
        "  h_B = (float *)malloc(nBytes);\n",
        "  hostRef = (float *)malloc(nBytes);\n",
        "  gpuRef = (float *)malloc(nBytes);\n",
        "\n",
        "  // Inicializa os dados ainda no host\n",
        "  initialData(h_A, nElem);\n",
        "  initialData(h_B, nElem);\n",
        "  \n",
        "  // Aloca memória global na GPU (device)\n",
        "  float *d_A, *d_B, *d_C;\n",
        "  cudaMalloc((float **)&d_A, nBytes);\n",
        "  cudaMalloc((float **)&d_B, nBytes);\n",
        "  cudaMalloc((float **)&d_C, nBytes);\n",
        "\n",
        "  // Transfere os dados do host (CPU) para o device (GPU)\n",
        "  cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);\n",
        "  cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);\n",
        "\n",
        "  // Configura os tamanho de:\n",
        "  //   bloco (quantas threads por bloco?)\n",
        "  //   e grade (quantos blocos por grade?)\n",
        "  dim3 block(nThread >> unroll);\n",
        "  dim3 grid(nElem >> 10);\n",
        "\n",
        "  // Executa a soma no device\n",
        "  sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C, unroll);\n",
        "  printf(\"Execution configuration <<<%d, %d>>>\\n\", grid.x, block.x);\n",
        "\n",
        "  // Copia o resultado do kernel do device para o host\n",
        "  cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);\n",
        "\n",
        "  // Soma os vetores na CPU para conferir os resultados\n",
        "  sumArraysOnHost(h_A, h_B, hostRef, nElem);\n",
        "  // Compara os resultados\n",
        "  checkResult(hostRef, gpuRef, nElem);\n",
        "\n",
        "  // Libera a memória da GPU\n",
        "  cudaFree(d_A);\n",
        "  cudaFree(d_B);\n",
        "  cudaFree(d_C);\n",
        "\n",
        "  // Libera a memória da CPU\n",
        "  free(h_A);\n",
        "  free(h_B);\n",
        "  free(hostRef);\n",
        "  free(gpuRef);\n",
        "\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Overwriting /content/src/sumUnroll.cu\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "lGrpEA329zEt",
        "colab_type": "text"
      },
      "source": [
        "2. Execute seu código com o `nvprof` para cada combinação de quantidade de elementos (32 M, 64 M, 128 M, 256 M) e fator de unroll (1, 2, 4, 8), deixando a quantidade de threads por bloco fixa em algum valor à sua escolha. Elabore um gráfico que mostre a quantidade de somas por segundo em função do fator de unroll. Nesse gráfico, crie uma curva para cada configuração de tamanho do vetor."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "bJ9puZEAXfFj",
        "colab_type": "text"
      },
      "source": [
        "Gráfico: https://drive.google.com/open?id=1uYZ1Giqjx976i6ovP8J5hE60sT9XZyL6"
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "RCh13k3nxBlF",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 1000
        },
        "outputId": "3441893c-8bae-43c8-9be9-92552a178d52"
      },
      "source": [
        "%cd /content/src\n",
        "!nvcc -O2 -arch=sm_35 -o sumUnroll sumUnroll.cu\n",
        "\n",
        "!nvprof --print-gpu-summary ./sumUnroll 32000000 1024 1\n",
        "!nvprof --print-gpu-summary ./sumUnroll 32000000 1024 2\n",
        "!nvprof --print-gpu-summary ./sumUnroll 32000000 1024 4\n",
        "!nvprof --print-gpu-summary ./sumUnroll 32000000 1024 8\n",
        "!nvprof --print-gpu-summary ./sumUnroll 64000000 1024 1\n",
        "!nvprof --print-gpu-summary ./sumUnroll 64000000 1024 2\n",
        "!nvprof --print-gpu-summary ./sumUnroll 64000000 1024 4\n",
        "!nvprof --print-gpu-summary ./sumUnroll 64000000 1024 8\n",
        "!nvprof --print-gpu-summary ./sumUnroll 128000000 1024 1\n",
        "!nvprof --print-gpu-summary ./sumUnroll 128000000 1024 2\n",
        "!nvprof --print-gpu-summary ./sumUnroll 128000000 1024 4\n",
        "!nvprof --print-gpu-summary ./sumUnroll 128000000 1024 8\n",
        "!nvprof --print-gpu-summary ./sumUnroll 256000000 1024 1\n",
        "!nvprof --print-gpu-summary ./sumUnroll 256000000 1024 2\n",
        "!nvprof --print-gpu-summary ./sumUnroll 256000000 1024 4\n",
        "!nvprof --print-gpu-summary ./sumUnroll 256000000 1024 8"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "./sumUnroll Starting...\n",
            "Vector size 32000000\n",
            "==11054== NVPROF is profiling process 11054, command: ./sumUnroll 32000000 1024 1\n",
            "Execution configuration <<<31250, 512>>>\n",
            "Arrays match.\n",
            "\n",
            "==11054== Profiling application: ./sumUnroll 32000000 1024 1\n",
            "==11054== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   67.71%  81.861ms         1  81.861ms  81.861ms  81.861ms  [CUDA memcpy DtoH]\n",
            "                   29.74%  35.953ms         2  17.976ms  17.789ms  18.164ms  [CUDA memcpy HtoD]\n",
            "                    2.56%  3.0909ms         1  3.0909ms  3.0909ms  3.0909ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 32000000\n",
            "==11068== NVPROF is profiling process 11068, command: ./sumUnroll 32000000 1024 2\n",
            "Execution configuration <<<31250, 256>>>\n",
            "Arrays match.\n",
            "\n",
            "==11068== Profiling application: ./sumUnroll 32000000 1024 2\n",
            "==11068== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   63.17%  83.012ms         1  83.012ms  83.012ms  83.012ms  [CUDA memcpy DtoH]\n",
            "                   33.48%  44.004ms         2  22.002ms  21.834ms  22.169ms  [CUDA memcpy HtoD]\n",
            "                    3.35%  4.3990ms         1  4.3990ms  4.3990ms  4.3990ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 32000000\n",
            "==11079== NVPROF is profiling process 11079, command: ./sumUnroll 32000000 1024 4\n",
            "Execution configuration <<<31250, 64>>>\n",
            "Arrays match.\n",
            "\n",
            "==11079== Profiling application: ./sumUnroll 32000000 1024 4\n",
            "==11079== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   52.36%  80.735ms         1  80.735ms  80.735ms  80.735ms  [CUDA memcpy DtoH]\n",
            "                   23.94%  36.907ms         2  18.453ms  17.736ms  19.170ms  [CUDA memcpy HtoD]\n",
            "                   23.70%  36.550ms         1  36.550ms  36.550ms  36.550ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 32000000\n",
            "==11090== NVPROF is profiling process 11090, command: ./sumUnroll 32000000 1024 8\n",
            "Execution configuration <<<31250, 4>>>\n",
            "Arrays match.\n",
            "\n",
            "==11090== Profiling application: ./sumUnroll 32000000 1024 8\n",
            "==11090== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   53.16%  83.775ms         1  83.775ms  83.775ms  83.775ms  [CUDA memcpy DtoH]\n",
            "                   28.29%  44.577ms         2  22.289ms  19.951ms  24.627ms  [CUDA memcpy HtoD]\n",
            "                   18.55%  29.229ms         1  29.229ms  29.229ms  29.229ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 64000000\n",
            "==11101== NVPROF is profiling process 11101, command: ./sumUnroll 64000000 1024 1\n",
            "Execution configuration <<<62500, 512>>>\n",
            "Arrays match.\n",
            "\n",
            "==11101== Profiling application: ./sumUnroll 64000000 1024 1\n",
            "==11101== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   68.18%  176.87ms         1  176.87ms  176.87ms  176.87ms  [CUDA memcpy DtoH]\n",
            "                   29.44%  76.367ms         2  38.183ms  37.979ms  38.388ms  [CUDA memcpy HtoD]\n",
            "                    2.39%  6.1905ms         1  6.1905ms  6.1905ms  6.1905ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 64000000\n",
            "==11115== NVPROF is profiling process 11115, command: ./sumUnroll 64000000 1024 2\n",
            "Execution configuration <<<62500, 256>>>\n",
            "Arrays match.\n",
            "\n",
            "==11115== Profiling application: ./sumUnroll 64000000 1024 2\n",
            "==11115== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   65.92%  154.70ms         1  154.70ms  154.70ms  154.70ms  [CUDA memcpy DtoH]\n",
            "                   30.34%  71.215ms         2  35.607ms  35.588ms  35.627ms  [CUDA memcpy HtoD]\n",
            "                    3.74%  8.7830ms         1  8.7830ms  8.7830ms  8.7830ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 64000000\n",
            "==11126== NVPROF is profiling process 11126, command: ./sumUnroll 64000000 1024 4\n",
            "Execution configuration <<<62500, 64>>>\n",
            "Arrays match.\n",
            "\n",
            "==11126== Profiling application: ./sumUnroll 64000000 1024 4\n",
            "==11126== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   51.83%  163.58ms         1  163.58ms  163.58ms  163.58ms  [CUDA memcpy DtoH]\n",
            "                   24.98%  78.845ms         2  39.423ms  38.409ms  40.436ms  [CUDA memcpy HtoD]\n",
            "                   23.18%  73.169ms         1  73.169ms  73.169ms  73.169ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 64000000\n",
            "==11137== NVPROF is profiling process 11137, command: ./sumUnroll 64000000 1024 8\n",
            "Execution configuration <<<62500, 4>>>\n",
            "Arrays match.\n",
            "\n",
            "==11137== Profiling application: ./sumUnroll 64000000 1024 8\n",
            "==11137== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   54.25%  162.62ms         1  162.62ms  162.62ms  162.62ms  [CUDA memcpy DtoH]\n",
            "                   26.27%  78.747ms         2  39.374ms  38.105ms  40.642ms  [CUDA memcpy HtoD]\n",
            "                   19.48%  58.376ms         1  58.376ms  58.376ms  58.376ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 128000000\n",
            "==11151== NVPROF is profiling process 11151, command: ./sumUnroll 128000000 1024 1\n",
            "Execution configuration <<<125000, 512>>>\n",
            "Arrays match.\n",
            "\n",
            "==11151== Profiling application: ./sumUnroll 128000000 1024 1\n",
            "==11151== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   66.33%  319.63ms         1  319.63ms  319.63ms  319.63ms  [CUDA memcpy DtoH]\n",
            "                   31.10%  149.87ms         2  74.934ms  72.672ms  77.197ms  [CUDA memcpy HtoD]\n",
            "                    2.57%  12.377ms         1  12.377ms  12.377ms  12.377ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 128000000\n",
            "==11162== NVPROF is profiling process 11162, command: ./sumUnroll 128000000 1024 2\n",
            "Execution configuration <<<125000, 256>>>\n",
            "Arrays match.\n",
            "\n",
            "==11162== Profiling application: ./sumUnroll 128000000 1024 2\n",
            "==11162== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   65.73%  339.45ms         1  339.45ms  339.45ms  339.45ms  [CUDA memcpy DtoH]\n",
            "                   30.87%  159.41ms         2  79.704ms  79.316ms  80.092ms  [CUDA memcpy HtoD]\n",
            "                    3.40%  17.563ms         1  17.563ms  17.563ms  17.563ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 128000000\n",
            "==11176== NVPROF is profiling process 11176, command: ./sumUnroll 128000000 1024 4\n",
            "Execution configuration <<<125000, 64>>>\n",
            "Arrays match.\n",
            "\n",
            "==11176== Profiling application: ./sumUnroll 128000000 1024 4\n",
            "==11176== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   52.16%  320.90ms         1  320.90ms  320.90ms  320.90ms  [CUDA memcpy DtoH]\n",
            "                   24.10%  148.29ms         2  74.144ms  72.542ms  75.746ms  [CUDA memcpy HtoD]\n",
            "                   23.74%  146.03ms         1  146.03ms  146.03ms  146.03ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 128000000\n",
            "==11187== NVPROF is profiling process 11187, command: ./sumUnroll 128000000 1024 8\n",
            "Execution configuration <<<125000, 4>>>\n",
            "Arrays match.\n",
            "\n",
            "==11187== Profiling application: ./sumUnroll 128000000 1024 8\n",
            "==11187== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   55.32%  333.94ms         1  333.94ms  333.94ms  333.94ms  [CUDA memcpy DtoH]\n",
            "                   25.37%  153.12ms         2  76.560ms  76.494ms  76.626ms  [CUDA memcpy HtoD]\n",
            "                   19.31%  116.54ms         1  116.54ms  116.54ms  116.54ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 256000000\n",
            "==11201== NVPROF is profiling process 11201, command: ./sumUnroll 256000000 1024 1\n",
            "Execution configuration <<<250000, 512>>>\n",
            "Arrays match.\n",
            "\n",
            "==11201== Profiling application: ./sumUnroll 256000000 1024 1\n",
            "==11201== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   66.06%  640.86ms         1  640.86ms  640.86ms  640.86ms  [CUDA memcpy DtoH]\n",
            "                   31.38%  304.47ms         2  152.24ms  150.97ms  153.50ms  [CUDA memcpy HtoD]\n",
            "                    2.56%  24.795ms         1  24.795ms  24.795ms  24.795ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 256000000\n",
            "==11215== NVPROF is profiling process 11215, command: ./sumUnroll 256000000 1024 2\n",
            "Execution configuration <<<250000, 256>>>\n",
            "Arrays match.\n",
            "\n",
            "==11215== Profiling application: ./sumUnroll 256000000 1024 2\n",
            "==11215== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   64.98%  636.73ms         1  636.73ms  636.73ms  636.73ms  [CUDA memcpy DtoH]\n",
            "                   31.44%  308.07ms         2  154.04ms  151.77ms  156.30ms  [CUDA memcpy HtoD]\n",
            "                    3.58%  35.094ms         1  35.094ms  35.094ms  35.094ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 256000000\n",
            "==11229== NVPROF is profiling process 11229, command: ./sumUnroll 256000000 1024 4\n",
            "Execution configuration <<<250000, 64>>>\n",
            "Arrays match.\n",
            "\n",
            "==11229== Profiling application: ./sumUnroll 256000000 1024 4\n",
            "==11229== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   51.94%  643.08ms         1  643.08ms  643.08ms  643.08ms  [CUDA memcpy DtoH]\n",
            "                   24.44%  302.53ms         2  151.27ms  150.36ms  152.17ms  [CUDA memcpy HtoD]\n",
            "                   23.62%  292.46ms         1  292.46ms  292.46ms  292.46ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumUnroll Starting...\n",
            "Vector size 256000000\n",
            "==11243== NVPROF is profiling process 11243, command: ./sumUnroll 256000000 1024 8\n",
            "Execution configuration <<<250000, 4>>>\n",
            "Arrays match.\n",
            "\n",
            "==11243== Profiling application: ./sumUnroll 256000000 1024 8\n",
            "==11243== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   53.37%  610.90ms         1  610.90ms  610.90ms  610.90ms  [CUDA memcpy DtoH]\n",
            "                   26.30%  301.03ms         2  150.52ms  146.89ms  154.14ms  [CUDA memcpy HtoD]\n",
            "                   20.34%  232.82ms         1  232.82ms  232.82ms  232.82ms  sumArraysOnGPU(float*, float*, float*, int)\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "L8wTvmMaBb9t",
        "colab_type": "text"
      },
      "source": [
        "3. Discuta sobre o impacto do fator de unroll no desempenho. Melhorou? Piorou? O que você consegue dizer sobre o motivo dessa mudança no desempenho?"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "Yj3-Dx9qULXN",
        "colab_type": "text"
      },
      "source": [
        "\n",
        "\n",
        "> Como esperado, o desempenho com fator unroll elevado se tornou inferior do que com fator unroll reduzido. Isso acontece porque como cada thread tem que executar mais de uma operação, o processo não ocorre paralelamente em relação a todos os elementos.\n"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "pzNrL75lBfUH",
        "colab_type": "text"
      },
      "source": [
        "## Tarefa 4 // Múltiplas Operações\n",
        "\n",
        "Vamos expandir a quantidade de cálculos para cada posição do vetor. Podemos estabelecer uma expressão qualquer que terá M operações. Por exemplo, se tivermos:\n",
        "```cpp\n",
        "C[i] = A[i] + 3 * B[i] - 2 * A[i] * B[i] + 5 * B[i] * B[i]\n",
        "```\n",
        "teremos 8 operações de cálculo para 2 de load e 1 de store. Vamos nos concentrar em alterar a quantidade de operações de cálculo, desconsiderando as de load/store.\n",
        "\n",
        "\n",
        "1. Modifique o código base do Exemplo 2, deixando parametrizada a quantidade de elementos do vetor, a quantidade de threads por bloco, e a quantidade de operações realizadas, de forma que você possa executar o programa como segue:\n",
        "```bash\n",
        "nvprof ./multiops <nElem> <nThreads> <ops>\n",
        "```"
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "aFRztrMU0aPF",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 34
        },
        "outputId": "12ac27e1-9d75-4d1c-f606-096dd6b3b696"
      },
      "source": [
        "%%writefile /content/src/sumOps.cu\n",
        "#include <cuda_runtime.h>\n",
        "#include <stdio.h>\n",
        "\n",
        "// Use esta função para conferir se as somas na CPU e na GPU correspondem\n",
        "void checkResult(float *hostRef, float *gpuRef, const int N) {\n",
        "  double epsilon = 1.0E-8;\n",
        "  for (int i=0; i<N; i++) {\n",
        "    if (abs(hostRef[i] - gpuRef[i]) > epsilon) {\n",
        "      printf(\"Arrays do not match!\\n\");\n",
        "      printf(\"host %5.2f gpu %5.2f at current %d\\n\", hostRef[i], gpuRef[i], i);\n",
        "      return;\n",
        "    }\n",
        "  }\n",
        "  printf(\"Arrays match.\\n\\n\");\n",
        "}\n",
        "\n",
        "// Use esta função para atribuir valores iniciais aos vetores de floats\n",
        "void initialData(float *a, int n) {\n",
        "  srand(time(NULL));\n",
        "  for (int i = 0; i < n; i++)\n",
        "    a[i] = (float)(rand() & 0xFF) / 10.0;\n",
        "}\n",
        "\n",
        "// Soma os vetores na CPU\n",
        "void sumArraysOnHost(float *A, float *B, float *C, const int N) {\n",
        "  for (int i = 0; i < N; i++)\n",
        "    C[i] = A[i] + B[i];\n",
        "}\n",
        "\n",
        "// Soma os vetores na GPU\n",
        "__global__ void sumArraysOnGPU(float *A, float *B, float *C, int ops) {\n",
        "  int i = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "  int a = 0;\n",
        "\n",
        "  for(int j = 0; j < ops; j++) a++;\n",
        "\n",
        "  C[i] = A[i] + B[i];\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"%s Starting...\\n\", argv[0]);\n",
        "  \n",
        "  // Configura o tamanho dos vetores\n",
        "  int nElem = atoi(argv[1]);\n",
        "  int nThread = atoi(argv[2]);\n",
        "  int ops = atoi(argv[3]);\n",
        "\n",
        "  printf(\"Vector size %d\\n\", nElem);\n",
        "  \n",
        "  // Aloca memória no host\n",
        "  size_t nBytes = nElem * sizeof(float);\n",
        "  float *h_A, *h_B, *hostRef, *gpuRef;\n",
        "  h_A = (float *)malloc(nBytes);\n",
        "  h_B = (float *)malloc(nBytes);\n",
        "  hostRef = (float *)malloc(nBytes);\n",
        "  gpuRef = (float *)malloc(nBytes);\n",
        "\n",
        "  // Inicializa os dados ainda no host\n",
        "  initialData(h_A, nElem);\n",
        "  initialData(h_B, nElem);\n",
        "  \n",
        "  // Aloca memória global na GPU (device)\n",
        "  float *d_A, *d_B, *d_C;\n",
        "  cudaMalloc((float **)&d_A, nBytes);\n",
        "  cudaMalloc((float **)&d_B, nBytes);\n",
        "  cudaMalloc((float **)&d_C, nBytes);\n",
        "\n",
        "  // Transfere os dados do host (CPU) para o device (GPU)\n",
        "  cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);\n",
        "  cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);\n",
        "\n",
        "  // Configura os tamanho de:\n",
        "  //   bloco (quantas threads por bloco?)\n",
        "  //   e grade (quantos blocos por grade?)\n",
        "  dim3 block(1024);\n",
        "  dim3 grid(nElem >> 10);\n",
        "\n",
        "  // Executa a soma no device\n",
        "  sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C, ops);\n",
        "  printf(\"Execution configuration <<<%d, %d>>>\\n\", grid.x, block.x);\n",
        "\n",
        "  // Copia o resultado do kernel do device para o host\n",
        "  cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);\n",
        "\n",
        "  // Soma os vetores na CPU para conferir os resultados\n",
        "  sumArraysOnHost(h_A, h_B, hostRef, nElem);\n",
        "  // Compara os resultados\n",
        "  checkResult(hostRef, gpuRef, nElem);\n",
        "\n",
        "  // Libera a memória da GPU\n",
        "  cudaFree(d_A);\n",
        "  cudaFree(d_B);\n",
        "  cudaFree(d_C);\n",
        "\n",
        "  // Libera a memória da CPU\n",
        "  free(h_A);\n",
        "  free(h_B);\n",
        "  free(hostRef);\n",
        "  free(gpuRef);\n",
        "\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Overwriting /content/src/sumOps.cu\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "wkPITOavGnLQ",
        "colab_type": "text"
      },
      "source": [
        "2. Execute seu código com o `nvprof` para cada combinação de quantidade de elementos (como especificado no exercício 2 da Tarefa 3) e quantidade de operações (5, 10 e 15). Elabore um gráfico que mostre o tempo de execução em função da do tamanho do vetor. Faça uma curva para cada configuração de quantidade de operações."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "nfHXK0kMVuNr",
        "colab_type": "text"
      },
      "source": [
        "Gráfico: https://drive.google.com/open?id=1fKbG_PVI7FcRpP5QV64Cq9ZfnhSVSjOP"
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "U7yU5o_y0-Wi",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 1000
        },
        "outputId": "d71ae840-493f-4ede-92fb-c95c0633d6dc"
      },
      "source": [
        "%cd /content/src\n",
        "!nvcc -O2 -arch=sm_35 -o sumOps sumOps.cu\n",
        "\n",
        "!nvprof --print-gpu-summary ./sumOps 32000000 1024 5\n",
        "!nvprof --print-gpu-summary ./sumOps 32000000 1024 10\n",
        "!nvprof --print-gpu-summary ./sumOps 32000000 1024 15\n",
        "!nvprof --print-gpu-summary ./sumOps 64000000 1024 5\n",
        "!nvprof --print-gpu-summary ./sumOps 64000000 1024 10\n",
        "!nvprof --print-gpu-summary ./sumOps 64000000 1024 15\n",
        "!nvprof --print-gpu-summary ./sumOps 128000000 1024 5\n",
        "!nvprof --print-gpu-summary ./sumOps 128000000 1024 10\n",
        "!nvprof --print-gpu-summary ./sumOps 128000000 1024 15\n",
        "!nvprof --print-gpu-summary ./sumOps 256000000 1024 5\n",
        "!nvprof --print-gpu-summary ./sumOps 256000000 1024 10\n",
        "!nvprof --print-gpu-summary ./sumOps 256000000 1024 15"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "./sumOps Starting...\n",
            "Vector size 32000000\n",
            "==11294== NVPROF is profiling process 11294, command: ./sumOps 32000000 1024 5\n",
            "Execution configuration <<<31250, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11294== Profiling application: ./sumOps 32000000 1024 5\n",
            "==11294== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   67.39%  80.610ms         1  80.610ms  80.610ms  80.610ms  [CUDA memcpy DtoH]\n",
            "                   29.98%  35.854ms         2  17.927ms  17.622ms  18.232ms  [CUDA memcpy HtoD]\n",
            "                    2.63%  3.1477ms         1  3.1477ms  3.1477ms  3.1477ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumOps Starting...\n",
            "Vector size 32000000\n",
            "==11305== NVPROF is profiling process 11305, command: ./sumOps 32000000 1024 10\n",
            "Execution configuration <<<31250, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11305== Profiling application: ./sumOps 32000000 1024 10\n",
            "==11305== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   67.12%  83.156ms         1  83.156ms  83.156ms  83.156ms  [CUDA memcpy DtoH]\n",
            "                   30.35%  37.599ms         2  18.800ms  18.731ms  18.868ms  [CUDA memcpy HtoD]\n",
            "                    2.54%  3.1422ms         1  3.1422ms  3.1422ms  3.1422ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumOps Starting...\n",
            "Vector size 32000000\n",
            "==11316== NVPROF is profiling process 11316, command: ./sumOps 32000000 1024 15\n",
            "Execution configuration <<<31250, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11316== Profiling application: ./sumOps 32000000 1024 15\n",
            "==11316== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   64.04%  84.035ms         1  84.035ms  84.035ms  84.035ms  [CUDA memcpy DtoH]\n",
            "                   33.57%  44.059ms         2  22.029ms  21.218ms  22.841ms  [CUDA memcpy HtoD]\n",
            "                    2.39%  3.1389ms         1  3.1389ms  3.1389ms  3.1389ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumOps Starting...\n",
            "Vector size 64000000\n",
            "==11330== NVPROF is profiling process 11330, command: ./sumOps 64000000 1024 5\n",
            "Execution configuration <<<62500, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11330== Profiling application: ./sumOps 64000000 1024 5\n",
            "==11330== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   66.50%  164.44ms         1  164.44ms  164.44ms  164.44ms  [CUDA memcpy DtoH]\n",
            "                   30.96%  76.549ms         2  38.275ms  36.192ms  40.358ms  [CUDA memcpy HtoD]\n",
            "                    2.54%  6.2766ms         1  6.2766ms  6.2766ms  6.2766ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumOps Starting...\n",
            "Vector size 64000000\n",
            "==11341== NVPROF is profiling process 11341, command: ./sumOps 64000000 1024 10\n",
            "Execution configuration <<<62500, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11341== Profiling application: ./sumOps 64000000 1024 10\n",
            "==11341== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   66.38%  163.99ms         1  163.99ms  163.99ms  163.99ms  [CUDA memcpy DtoH]\n",
            "                   31.08%  76.780ms         2  38.390ms  36.235ms  40.544ms  [CUDA memcpy HtoD]\n",
            "                    2.54%  6.2734ms         1  6.2734ms  6.2734ms  6.2734ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumOps Starting...\n",
            "Vector size 64000000\n",
            "==11352== NVPROF is profiling process 11352, command: ./sumOps 64000000 1024 15\n",
            "Execution configuration <<<62500, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11352== Profiling application: ./sumOps 64000000 1024 15\n",
            "==11352== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   66.85%  165.78ms         1  165.78ms  165.78ms  165.78ms  [CUDA memcpy DtoH]\n",
            "                   30.62%  75.934ms         2  37.967ms  37.530ms  38.405ms  [CUDA memcpy HtoD]\n",
            "                    2.53%  6.2667ms         1  6.2667ms  6.2667ms  6.2667ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumOps Starting...\n",
            "Vector size 128000000\n",
            "==11366== NVPROF is profiling process 11366, command: ./sumOps 128000000 1024 5\n",
            "Execution configuration <<<125000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11366== Profiling application: ./sumOps 128000000 1024 5\n",
            "==11366== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   66.45%  334.92ms         1  334.92ms  334.92ms  334.92ms  [CUDA memcpy DtoH]\n",
            "                   31.07%  156.60ms         2  78.298ms  75.885ms  80.711ms  [CUDA memcpy HtoD]\n",
            "                    2.48%  12.524ms         1  12.524ms  12.524ms  12.524ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumOps Starting...\n",
            "Vector size 128000000\n",
            "==11377== NVPROF is profiling process 11377, command: ./sumOps 128000000 1024 10\n",
            "Execution configuration <<<125000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11377== Profiling application: ./sumOps 128000000 1024 10\n",
            "==11377== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   66.74%  326.78ms         1  326.78ms  326.78ms  326.78ms  [CUDA memcpy DtoH]\n",
            "                   30.70%  150.31ms         2  75.157ms  73.105ms  77.210ms  [CUDA memcpy HtoD]\n",
            "                    2.56%  12.515ms         1  12.515ms  12.515ms  12.515ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumOps Starting...\n",
            "Vector size 128000000\n",
            "==11391== NVPROF is profiling process 11391, command: ./sumOps 128000000 1024 15\n",
            "Execution configuration <<<125000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11391== Profiling application: ./sumOps 128000000 1024 15\n",
            "==11391== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   65.84%  310.50ms         1  310.50ms  310.50ms  310.50ms  [CUDA memcpy DtoH]\n",
            "                   31.51%  148.60ms         2  74.298ms  72.627ms  75.968ms  [CUDA memcpy HtoD]\n",
            "                    2.65%  12.498ms         1  12.498ms  12.498ms  12.498ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumOps Starting...\n",
            "Vector size 256000000\n",
            "==11402== NVPROF is profiling process 11402, command: ./sumOps 256000000 1024 5\n",
            "Execution configuration <<<250000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11402== Profiling application: ./sumOps 256000000 1024 5\n",
            "==11402== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   65.91%  628.43ms         1  628.43ms  628.43ms  628.43ms  [CUDA memcpy DtoH]\n",
            "                   31.46%  299.96ms         2  149.98ms  147.36ms  152.60ms  [CUDA memcpy HtoD]\n",
            "                    2.64%  25.143ms         1  25.143ms  25.143ms  25.143ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumOps Starting...\n",
            "Vector size 256000000\n",
            "==11416== NVPROF is profiling process 11416, command: ./sumOps 256000000 1024 10\n",
            "Execution configuration <<<250000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11416== Profiling application: ./sumOps 256000000 1024 10\n",
            "==11416== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   65.25%  636.77ms         1  636.77ms  636.77ms  636.77ms  [CUDA memcpy DtoH]\n",
            "                   32.18%  314.06ms         2  157.03ms  155.03ms  159.03ms  [CUDA memcpy HtoD]\n",
            "                    2.57%  25.130ms         1  25.130ms  25.130ms  25.130ms  sumArraysOnGPU(float*, float*, float*, int)\n",
            "./sumOps Starting...\n",
            "Vector size 256000000\n",
            "==11430== NVPROF is profiling process 11430, command: ./sumOps 256000000 1024 15\n",
            "Execution configuration <<<250000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11430== Profiling application: ./sumOps 256000000 1024 15\n",
            "==11430== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   64.45%  645.89ms         1  645.89ms  645.89ms  645.89ms  [CUDA memcpy DtoH]\n",
            "                   33.04%  331.18ms         2  165.59ms  155.32ms  175.86ms  [CUDA memcpy HtoD]\n",
            "                    2.51%  25.155ms         1  25.155ms  25.155ms  25.155ms  sumArraysOnGPU(float*, float*, float*, int)\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "uYXM48M8Wtfk",
        "colab_type": "code",
        "colab": {}
      },
      "source": [
        ""
      ],
      "execution_count": null,
      "outputs": []
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "4JTD8sQsG8Qw",
        "colab_type": "text"
      },
      "source": [
        "3. Discuta os resultados mostrados pelo gráfico quanto ao tempo de execução."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "8wr1fNNiWA9h",
        "colab_type": "text"
      },
      "source": [
        "> Como é obvio de se esperar, a medida que a quantidade de operações aumenta, independente da quantidade de elementos, o tempo de execução aumenta proporcionalmente quase que em uma linha reta em todos os casos. Isso ocorre porque por mais que emnquanto um thread esteja disponível para executar diferentes operações simultaneamente, o número de operações continua impactando no tempo de execução final."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "2vGauTl5HOZa",
        "colab_type": "text"
      },
      "source": [
        "## Exemplo 5 // Acesso à Memória Global\n",
        "\n",
        "_Referência: Capítulo 4 - Global Memory - Memory Access Patterns_\n",
        "\n",
        "Os acessos à memória global são feitos com a ajuda de caches, e essa memória global é um espaço de memória que você consegue acessar a partir de seus kernels. O padrão ideal de acessos à cache é quando esse acesso é feito de forma _aligned_ (alinhada) e _coalesced_ (agrupada).\n",
        "\n",
        "Acesso alinhado ocorre quando o primeiro endereço acessado numa transferência é um múltiplo par da granularidade de cache usada nessa transferência - 32 B para L2 ou 128 B para L1. Um acesso desalinhado desperdiça os dados carregados.\n",
        "\n",
        "Acesso agrupado occore quando todas as 32 threads em um _warp_ (warp é o conjunto de threads sendo executadas simultaneamente) acessam uma área contígua de memória.\n",
        "\n",
        "![Acesso alinhado e agrupado](https://docs.google.com/uc?export=download&id=116WeX9Ng7J_mgIoCwr5f9mQfTx65F48X)\n"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "TigX_HXaN7a0",
        "colab_type": "text"
      },
      "source": [
        "## Tarefa 5 // Acesso à Memória Global\n",
        "\n",
        "Vamos explorar o que acontece quando quebramos os acessos alinhados e agrupados. \n",
        "\n",
        "Primeiramente, vejamos o que acontece se acessarmos posições mais espaçadas dos vetores do exemplo de soma. Vamos chamar esse espaçamento de pulos, ou _strikes_ no acesso, e seu valor será em potências de 2.\n",
        "\n",
        "Faremos a soma de uma forma diferente: cada posição `C[i]` terá o resultado da soma das posições `A[i << strike] + B[i << strike]`. Por exemplo, se tivermos o _strike_ igual a 3 (2^3 = 8), teremos o seguinte calculo `C[6] = A[48] + B[48]`. Na próxima posição, teremos `C[7] = A[56] + B[56]`. Note que muitas posições dos vetores A e B foram puladas, e note também que o vetor C será menor que os vetores A e B. Especificamente, seu tamanho será `N >> strike`.\n",
        "\n",
        "1. Adapte o código de soma de vetores considerando o _strike_. Deixe parametrizado o tamanho dos vetores A e B, o número de threads por bloco, e o valor do _strike_. "
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "6qDEyz9OWmz0",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 34
        },
        "outputId": "c320fdae-2f20-4d96-c985-dbf040013d77"
      },
      "source": [
        "%%writefile /content/src/sumStrike.cu\n",
        "#include <cuda_runtime.h>\n",
        "#include <stdio.h>\n",
        "\n",
        "// Use esta função para conferir se as somas na CPU e na GPU correspondem\n",
        "void checkResult(float *hostRef, float *gpuRef, const int N) {\n",
        "  double epsilon = 1.0E-8;\n",
        "  for (int i=0; i<N; i++) {\n",
        "    if (abs(hostRef[i] - gpuRef[i]) > epsilon) {\n",
        "      printf(\"Arrays do not match!\\n\");\n",
        "      printf(\"host %5.2f gpu %5.2f at current %d\\n\", hostRef[i], gpuRef[i], i);\n",
        "      return;\n",
        "    }\n",
        "  }\n",
        "  printf(\"Arrays match.\\n\\n\");\n",
        "}\n",
        "\n",
        "// Use esta função para atribuir valores iniciais aos vetores de floats\n",
        "void initialData(float *a, int n) {\n",
        "  srand(time(NULL));\n",
        "  for (int i = 0; i < n; i++)\n",
        "    a[i] = (float)(rand() & 0xFF) / 10.0;\n",
        "}\n",
        "\n",
        "// Soma os vetores na GPU\n",
        "__global__ void sumArraysOnGPU(float *A, float *B, float *C, int strike, int n) {\n",
        "  int i = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "  \n",
        "  int j = i << strike;\n",
        "\n",
        "  if(j>=n){\n",
        "    C[i] = A[i] + B[i];\n",
        "  } else {\n",
        "    C[i] = A[j] + B[j];\n",
        "  }\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"%s Starting...\\n\", argv[0]);\n",
        "  \n",
        "  // Configura o tamanho dos vetores\n",
        "  int nElem = atoi(argv[1]);\n",
        "  int nThread = atoi(argv[2]);\n",
        "  int strike = atoi(argv[3]);\n",
        "\n",
        "  printf(\"Vector size %d\\n\", nElem);\n",
        "  \n",
        "  // Aloca memória no host\n",
        "  size_t nBytes = nElem * sizeof(float);\n",
        "  float *h_A, *h_B, *hostRef, *gpuRef;\n",
        "  h_A = (float *)malloc(nBytes);\n",
        "  h_B = (float *)malloc(nBytes);\n",
        "  hostRef = (float *)malloc(nBytes);\n",
        "  gpuRef = (float *)malloc(nBytes);\n",
        "\n",
        "  // Inicializa os dados ainda no host\n",
        "  initialData(h_A, nElem);\n",
        "  initialData(h_B, nElem);\n",
        "  \n",
        "  // Aloca memória global na GPU (device)\n",
        "  float *d_A, *d_B, *d_C;\n",
        "  cudaMalloc((float **)&d_A, nBytes);\n",
        "  cudaMalloc((float **)&d_B, nBytes);\n",
        "  cudaMalloc((float **)&d_C, nBytes);\n",
        "\n",
        "  // Transfere os dados do host (CPU) para o device (GPU)\n",
        "  cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);\n",
        "  cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);\n",
        "\n",
        "  // Configura os tamanho de:\n",
        "  //   bloco (quantas threads por bloco?)\n",
        "  //   e grade (quantos blocos por grade?)\n",
        "  dim3 block(1024);\n",
        "  dim3 grid(nElem >> 10);\n",
        "\n",
        "  // Executa a soma no device\n",
        "  sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C, strike, nElem);\n",
        "  printf(\"Execution configuration <<<%d, %d>>>\\n\", grid.x, block.x);\n",
        "\n",
        "  // Copia o resultado do kernel do device para o host\n",
        "  cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);\n",
        "\n",
        "  // Libera a memória da GPU\n",
        "  cudaFree(d_A);\n",
        "  cudaFree(d_B);\n",
        "  cudaFree(d_C);\n",
        "\n",
        "  // Libera a memória da CPU\n",
        "  free(h_A);\n",
        "  free(h_B);\n",
        "  free(hostRef);\n",
        "  free(gpuRef);\n",
        "\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Overwriting /content/src/sumStrike.cu\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "8Iqgbfco2ZOC",
        "colab_type": "text"
      },
      "source": [
        "2. Execute seu programa com o `nvprof` usando diferentes combinações de quantidade de elementos para os vetores A e B (os mesmos valores do exercício 2 da Tarefa 3) e diferentes valores de _strike_ (4, 8, 16, 32, 64 e 128). Elabore um gráfico que mostre o tempo de execução em função do tamanho dos vetores, fazendo uma curva para cada valor de _strike_. Use todos os parâmetros como potências de 2."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "HTt6_vxaXqBt",
        "colab_type": "text"
      },
      "source": [
        "Gráfico: https://drive.google.com/open?id=12wfYKBUK7RiaTc_EWNfi90kfQr9hB3LP"
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "6loHITik2cR5",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 1000
        },
        "outputId": "ff4bc05c-c291-46d4-9de4-826ca73a3740"
      },
      "source": [
        "# Reduzi o tamanho do vetor em 10000 vezes, pois iria demorar tempo demais para\n",
        "# executar\n",
        "\n",
        "%cd /content/src\n",
        "!nvcc -O2 -arch=sm_35 -o sumStrike sumStrike.cu\n",
        "\n",
        "!nvprof --print-gpu-summary ./sumStrike 3200 1024 4\n",
        "!nvprof --print-gpu-summary ./sumStrike 3200 1024 8\n",
        "!nvprof --print-gpu-summary ./sumStrike 3200 1024 16\n",
        "!nvprof --print-gpu-summary ./sumStrike 3200 1024 32\n",
        "!nvprof --print-gpu-summary ./sumStrike 3200 1024 64\n",
        "!nvprof --print-gpu-summary ./sumStrike 3200 1024 128\n",
        "!nvprof --print-gpu-summary ./sumStrike 6400 1024 4\n",
        "!nvprof --print-gpu-summary ./sumStrike 6400 1024 8\n",
        "!nvprof --print-gpu-summary ./sumStrike 6400 1024 16\n",
        "!nvprof --print-gpu-summary ./sumStrike 6400 1024 32\n",
        "!nvprof --print-gpu-summary ./sumStrike 6400 1024 64\n",
        "!nvprof --print-gpu-summary ./sumStrike 6400 1024 128\n",
        "!nvprof --print-gpu-summary ./sumStrike 12800 1024 4\n",
        "!nvprof --print-gpu-summary ./sumStrike 12800 1024 8\n",
        "!nvprof --print-gpu-summary ./sumStrike 12800 1024 16\n",
        "!nvprof --print-gpu-summary ./sumStrike 12800 1024 32\n",
        "!nvprof --print-gpu-summary ./sumStrike 12800 1024 64\n",
        "!nvprof --print-gpu-summary ./sumStrike 12800 1024 128\n",
        "!nvprof --print-gpu-summary ./sumStrike 25600 1024 4\n",
        "!nvprof --print-gpu-summary ./sumStrike 25600 1024 8\n",
        "!nvprof --print-gpu-summary ./sumStrike 25600 1024 16\n",
        "!nvprof --print-gpu-summary ./sumStrike 25600 1024 32\n",
        "!nvprof --print-gpu-summary ./sumStrike 25600 1024 64\n",
        "!nvprof --print-gpu-summary ./sumStrike 25600 1024 128"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "./sumStrike Starting...\n",
            "Vector size 3200\n",
            "==13655== NVPROF is profiling process 13655, command: ./sumStrike 3200 1024 4\n",
            "Execution configuration <<<3, 1024>>>\n",
            "==13655== Profiling application: ./sumStrike 3200 1024 4\n",
            "==13655== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   48.44%  8.9280us         2  4.4640us  4.3840us  4.5440us  [CUDA memcpy HtoD]\n",
            "                   26.39%  4.8640us         1  4.8640us  4.8640us  4.8640us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "                   25.17%  4.6400us         1  4.6400us  4.6400us  4.6400us  [CUDA memcpy DtoH]\n",
            "./sumStrike Starting...\n",
            "Vector size 3200\n",
            "==13666== NVPROF is profiling process 13666, command: ./sumStrike 3200 1024 8\n",
            "Execution configuration <<<3, 1024>>>\n",
            "==13666== Profiling application: ./sumStrike 3200 1024 8\n",
            "==13666== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   50.00%  8.8960us         2  4.4480us  4.3520us  4.5440us  [CUDA memcpy HtoD]\n",
            "                   26.44%  4.7040us         1  4.7040us  4.7040us  4.7040us  [CUDA memcpy DtoH]\n",
            "                   23.56%  4.1920us         1  4.1920us  4.1920us  4.1920us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 3200\n",
            "==13677== NVPROF is profiling process 13677, command: ./sumStrike 3200 1024 16\n",
            "Execution configuration <<<3, 1024>>>\n",
            "==13677== Profiling application: ./sumStrike 3200 1024 16\n",
            "==13677== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   52.12%  9.0560us         2  4.5280us  4.3840us  4.6720us  [CUDA memcpy HtoD]\n",
            "                   24.31%  4.2240us         1  4.2240us  4.2240us  4.2240us  [CUDA memcpy DtoH]\n",
            "                   23.57%  4.0960us         1  4.0960us  4.0960us  4.0960us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 3200\n",
            "==13691== NVPROF is profiling process 13691, command: ./sumStrike 3200 1024 32\n",
            "Execution configuration <<<3, 1024>>>\n",
            "==13691== Profiling application: ./sumStrike 3200 1024 32\n",
            "==13691== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   51.96%  8.9280us         2  4.4640us  4.3840us  4.5440us  [CUDA memcpy HtoD]\n",
            "                   26.07%  4.4800us         1  4.4800us  4.4800us  4.4800us  [CUDA memcpy DtoH]\n",
            "                   21.97%  3.7760us         1  3.7760us  3.7760us  3.7760us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 3200\n",
            "==13702== NVPROF is profiling process 13702, command: ./sumStrike 3200 1024 64\n",
            "Execution configuration <<<3, 1024>>>\n",
            "==13702== Profiling application: ./sumStrike 3200 1024 64\n",
            "==13702== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   47.35%  11.424us         1  11.424us  11.424us  11.424us  [CUDA memcpy DtoH]\n",
            "                   37.00%  8.9280us         2  4.4640us  4.3840us  4.5440us  [CUDA memcpy HtoD]\n",
            "                   15.65%  3.7760us         1  3.7760us  3.7760us  3.7760us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 3200\n",
            "==13713== NVPROF is profiling process 13713, command: ./sumStrike 3200 1024 128\n",
            "Execution configuration <<<3, 1024>>>\n",
            "==13713== Profiling application: ./sumStrike 3200 1024 128\n",
            "==13713== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   53.01%  9.0240us         2  4.5120us  4.3840us  4.6400us  [CUDA memcpy HtoD]\n",
            "                   25.00%  4.2560us         1  4.2560us  4.2560us  4.2560us  [CUDA memcpy DtoH]\n",
            "                   21.99%  3.7440us         1  3.7440us  3.7440us  3.7440us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 6400\n",
            "==13724== NVPROF is profiling process 13724, command: ./sumStrike 6400 1024 4\n",
            "Execution configuration <<<6, 1024>>>\n",
            "==13724== Profiling application: ./sumStrike 6400 1024 4\n",
            "==13724== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   56.37%  14.592us         2  7.2960us  7.2320us  7.3600us  [CUDA memcpy HtoD]\n",
            "                   22.62%  5.8560us         1  5.8560us  5.8560us  5.8560us  [CUDA memcpy DtoH]\n",
            "                   21.01%  5.4400us         1  5.4400us  5.4400us  5.4400us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 6400\n",
            "==13735== NVPROF is profiling process 13735, command: ./sumStrike 6400 1024 8\n",
            "Execution configuration <<<6, 1024>>>\n",
            "==13735== Profiling application: ./sumStrike 6400 1024 8\n",
            "==13735== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   51.76%  14.592us         2  7.2960us  7.2320us  7.3600us  [CUDA memcpy HtoD]\n",
            "                   27.47%  7.7440us         1  7.7440us  7.7440us  7.7440us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "                   20.77%  5.8560us         1  5.8560us  5.8560us  5.8560us  [CUDA memcpy DtoH]\n",
            "./sumStrike Starting...\n",
            "Vector size 6400\n",
            "==13746== NVPROF is profiling process 13746, command: ./sumStrike 6400 1024 16\n",
            "Execution configuration <<<6, 1024>>>\n",
            "==13746== Profiling application: ./sumStrike 6400 1024 16\n",
            "==13746== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   59.45%  14.687us         2  7.3430us  7.2310us  7.4560us  [CUDA memcpy HtoD]\n",
            "                   24.22%  5.9840us         1  5.9840us  5.9840us  5.9840us  [CUDA memcpy DtoH]\n",
            "                   16.32%  4.0320us         1  4.0320us  4.0320us  4.0320us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 6400\n",
            "==13757== NVPROF is profiling process 13757, command: ./sumStrike 6400 1024 32\n",
            "Execution configuration <<<6, 1024>>>\n",
            "==13757== Profiling application: ./sumStrike 6400 1024 32\n",
            "==13757== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   58.64%  14.656us         2  7.3280us  7.2320us  7.4240us  [CUDA memcpy HtoD]\n",
            "                   25.86%  6.4640us         1  6.4640us  6.4640us  6.4640us  [CUDA memcpy DtoH]\n",
            "                   15.49%  3.8720us         1  3.8720us  3.8720us  3.8720us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 6400\n",
            "==13768== NVPROF is profiling process 13768, command: ./sumStrike 6400 1024 64\n",
            "Execution configuration <<<6, 1024>>>\n",
            "==13768== Profiling application: ./sumStrike 6400 1024 64\n",
            "==13768== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   60.31%  15.168us         2  7.5840us  7.2320us  7.9360us  [CUDA memcpy HtoD]\n",
            "                   24.05%  6.0480us         1  6.0480us  6.0480us  6.0480us  [CUDA memcpy DtoH]\n",
            "                   15.65%  3.9360us         1  3.9360us  3.9360us  3.9360us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 6400\n",
            "==13779== NVPROF is profiling process 13779, command: ./sumStrike 6400 1024 128\n",
            "Execution configuration <<<6, 1024>>>\n",
            "==13779== Profiling application: ./sumStrike 6400 1024 128\n",
            "==13779== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   68.16%  22.336us         2  11.168us  7.2640us  15.072us  [CUDA memcpy HtoD]\n",
            "                   20.02%  6.5600us         1  6.5600us  6.5600us  6.5600us  [CUDA memcpy DtoH]\n",
            "                   11.82%  3.8720us         1  3.8720us  3.8720us  3.8720us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 12800\n",
            "==13793== NVPROF is profiling process 13793, command: ./sumStrike 12800 1024 4\n",
            "Execution configuration <<<12, 1024>>>\n",
            "==13793== Profiling application: ./sumStrike 12800 1024 4\n",
            "==13793== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   61.45%  26.112us         2  13.056us  12.928us  13.184us  [CUDA memcpy HtoD]\n",
            "                   22.06%  9.3760us         1  9.3760us  9.3760us  9.3760us  [CUDA memcpy DtoH]\n",
            "                   16.49%  7.0080us         1  7.0080us  7.0080us  7.0080us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 12800\n",
            "==13804== NVPROF is profiling process 13804, command: ./sumStrike 12800 1024 8\n",
            "Execution configuration <<<12, 1024>>>\n",
            "==13804== Profiling application: ./sumStrike 12800 1024 8\n",
            "==13804== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   62.52%  26.208us         2  13.104us  12.928us  13.280us  [CUDA memcpy HtoD]\n",
            "                   23.51%  9.8560us         1  9.8560us  9.8560us  9.8560us  [CUDA memcpy DtoH]\n",
            "                   13.97%  5.8560us         1  5.8560us  5.8560us  5.8560us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 12800\n",
            "==13815== NVPROF is profiling process 13815, command: ./sumStrike 12800 1024 16\n",
            "Execution configuration <<<12, 1024>>>\n",
            "==13815== Profiling application: ./sumStrike 12800 1024 16\n",
            "==13815== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   72.89%  35.360us         2  17.680us  13.728us  21.632us  [CUDA memcpy HtoD]\n",
            "                   18.73%  9.0880us         1  9.0880us  9.0880us  9.0880us  [CUDA memcpy DtoH]\n",
            "                    8.38%  4.0640us         1  4.0640us  4.0640us  4.0640us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 12800\n",
            "==13826== NVPROF is profiling process 13826, command: ./sumStrike 12800 1024 32\n",
            "Execution configuration <<<12, 1024>>>\n",
            "==13826== Profiling application: ./sumStrike 12800 1024 32\n",
            "==13826== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   71.79%  33.952us         2  16.976us  13.088us  20.864us  [CUDA memcpy HtoD]\n",
            "                   19.28%  9.1200us         1  9.1200us  9.1200us  9.1200us  [CUDA memcpy DtoH]\n",
            "                    8.93%  4.2240us         1  4.2240us  4.2240us  4.2240us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 12800\n",
            "==13837== NVPROF is profiling process 13837, command: ./sumStrike 12800 1024 64\n",
            "Execution configuration <<<12, 1024>>>\n",
            "==13837== Profiling application: ./sumStrike 12800 1024 64\n",
            "==13837== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   72.11%  34.912us         2  17.456us  13.152us  21.760us  [CUDA memcpy HtoD]\n",
            "                   19.10%  9.2480us         1  9.2480us  9.2480us  9.2480us  [CUDA memcpy DtoH]\n",
            "                    8.79%  4.2560us         1  4.2560us  4.2560us  4.2560us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 12800\n",
            "==13848== NVPROF is profiling process 13848, command: ./sumStrike 12800 1024 128\n",
            "Execution configuration <<<12, 1024>>>\n",
            "==13848== Profiling application: ./sumStrike 12800 1024 128\n",
            "==13848== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   66.40%  26.496us         2  13.248us  12.960us  13.536us  [CUDA memcpy HtoD]\n",
            "                   22.77%  9.0880us         1  9.0880us  9.0880us  9.0880us  [CUDA memcpy DtoH]\n",
            "                   10.83%  4.3200us         1  4.3200us  4.3200us  4.3200us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 25600\n",
            "==13859== NVPROF is profiling process 13859, command: ./sumStrike 25600 1024 4\n",
            "Execution configuration <<<25, 1024>>>\n",
            "==13859== Profiling application: ./sumStrike 25600 1024 4\n",
            "==13859== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   59.73%  34.976us         2  17.488us  16.864us  18.112us  [CUDA memcpy HtoD]\n",
            "                   26.56%  15.552us         1  15.552us  15.552us  15.552us  [CUDA memcpy DtoH]\n",
            "                   13.72%  8.0320us         1  8.0320us  8.0320us  8.0320us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 25600\n",
            "==13870== NVPROF is profiling process 13870, command: ./sumStrike 25600 1024 8\n",
            "Execution configuration <<<25, 1024>>>\n",
            "==13870== Profiling application: ./sumStrike 25600 1024 8\n",
            "==13870== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   62.49%  35.392us         2  17.696us  16.960us  18.432us  [CUDA memcpy HtoD]\n",
            "                   27.40%  15.519us         1  15.519us  15.519us  15.519us  [CUDA memcpy DtoH]\n",
            "                   10.11%  5.7280us         1  5.7280us  5.7280us  5.7280us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 25600\n",
            "==13881== NVPROF is profiling process 13881, command: ./sumStrike 25600 1024 16\n",
            "Execution configuration <<<25, 1024>>>\n",
            "==13881== Profiling application: ./sumStrike 25600 1024 16\n",
            "==13881== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   63.14%  34.528us         2  17.264us  16.576us  17.952us  [CUDA memcpy HtoD]\n",
            "                   28.38%  15.520us         1  15.520us  15.520us  15.520us  [CUDA memcpy DtoH]\n",
            "                    8.48%  4.6400us         1  4.6400us  4.6400us  4.6400us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 25600\n",
            "==13895== NVPROF is profiling process 13895, command: ./sumStrike 25600 1024 32\n",
            "Execution configuration <<<25, 1024>>>\n",
            "==13895== Profiling application: ./sumStrike 25600 1024 32\n",
            "==13895== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   62.67%  34.815us         2  17.407us  16.383us  18.432us  [CUDA memcpy HtoD]\n",
            "                   28.00%  15.552us         1  15.552us  15.552us  15.552us  [CUDA memcpy DtoH]\n",
            "                    9.33%  5.1840us         1  5.1840us  5.1840us  5.1840us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 25600\n",
            "==13906== NVPROF is profiling process 13906, command: ./sumStrike 25600 1024 64\n",
            "Execution configuration <<<25, 1024>>>\n",
            "==13906== Profiling application: ./sumStrike 25600 1024 64\n",
            "==13906== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   62.63%  35.552us         2  17.776us  17.056us  18.496us  [CUDA memcpy HtoD]\n",
            "                   28.24%  16.032us         1  16.032us  16.032us  16.032us  [CUDA memcpy DtoH]\n",
            "                    9.13%  5.1840us         1  5.1840us  5.1840us  5.1840us  sumArraysOnGPU(float*, float*, float*, int, int)\n",
            "./sumStrike Starting...\n",
            "Vector size 25600\n",
            "==13917== NVPROF is profiling process 13917, command: ./sumStrike 25600 1024 128\n",
            "Execution configuration <<<25, 1024>>>\n",
            "==13917== Profiling application: ./sumStrike 25600 1024 128\n",
            "==13917== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   62.40%  34.463us         2  17.231us  16.383us  18.080us  [CUDA memcpy HtoD]\n",
            "                   28.10%  15.520us         1  15.520us  15.520us  15.520us  [CUDA memcpy DtoH]\n",
            "                    9.50%  5.2480us         1  5.2480us  5.2480us  5.2480us  sumArraysOnGPU(float*, float*, float*, int, int)\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "LTTd3eND4LqS",
        "colab_type": "text"
      },
      "source": [
        "3. Discuta o impacto do _strike_ no desempenho."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "iFm85fCmXvv6",
        "colab_type": "text"
      },
      "source": [
        "> Como é possível notar no gráfico já apresentado, o tempo de execução aumenta bastante com o aumento do strike. Embora que não esteja sendo apresentado no gráfico, o tempo de execução das somas continua sendo o mesmo (ou quase o mesmo) do mesmo algoritmo sem strike, no entanto, como as posições dos vetores não são exploradas de forma contígua, o princípio da localidade acaba sendo cada vez menos explorado a medida que o número do strike aumenta. Portanto, o fato de o novo tempo de execução ser superior é provocado pela taxa de transferência de dados entre CPU e GPU."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "KGDS3jJS4SXJ",
        "colab_type": "text"
      },
      "source": [
        "Para os próximos exercícios desta tarefa, vamos simular um acesso aleatório nas posições do vetor, usando novamente a soma de vetores como base, com is vetores A, B e C de mesmo tamanho. Para isso, vamos usar um vetor a mais, com o mesmo tamanho dos outros, chamado R. Este vai mapear uma posição em outra, aleatória. Você deve montar esse vetor de forma que o conteúdo de cada um de seus elementos seja um índice, entre 0 e N - 1, porém esses valores devem ser \"sorteados\" no vetor.\n",
        "\n",
        "Fazendo isso, sempre que você acessar alguma posição dos vetores A, B ou C, ao invés de usar diretamente o índice `i`, você vai usar o índice `R[i]`:\n",
        "```cpp\n",
        "// No corpo do for da função sumArraysOnHost\n",
        "//  ou no corpo do kernel sumArraysOnGPU\n",
        "C[R[i]] = A[R[i]] + B[R[i]];\n",
        "```\n",
        "4. Altere o código para que seja criado esse vetor `R`, e deixe parametrizado: o número de elementos nos vetores, o número de threads por bloco e uma opção para se usar ou não o vetor `R`."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "AbPtUDdb5cBz",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 34
        },
        "outputId": "9b0ad8b1-cc9f-4a05-9e50-0beee554ce93"
      },
      "source": [
        "%%writefile /content/src/sumR.cu\n",
        "#include <cuda_runtime.h>\n",
        "#include <stdio.h>\n",
        "\n",
        "// Use esta função para conferir se as somas na CPU e na GPU correspondem\n",
        "void checkResult(float *hostRef, float *gpuRef, const int N) {\n",
        "  double epsilon = 1.0E-8;\n",
        "  for (int i=0; i<N; i++) {\n",
        "    if (abs(hostRef[i] - gpuRef[i]) > epsilon) {\n",
        "      printf(\"Arrays do not match!\\n\");\n",
        "      printf(\"host %5.2f gpu %5.2f at current %d\\n\", hostRef[i], gpuRef[i], i);\n",
        "      return;\n",
        "    }\n",
        "  }\n",
        "  printf(\"Arrays match.\\n\\n\");\n",
        "}\n",
        "\n",
        "// Use esta função para atribuir valores iniciais aos vetores de floats\n",
        "void initialData(float *a, int n) {\n",
        "  srand(time(NULL));\n",
        "  for (int i = 0; i < n; i++)\n",
        "    a[i] = (float)(rand() & 0xFF) / 10.0;\n",
        "}\n",
        "\n",
        "// Use esta função para atribuir valores iniciais aos vetores de floats\n",
        "void initialDataInt(int *a, int n) {\n",
        "  srand(time(NULL));\n",
        "  for (int i = 0; i < n; i++)\n",
        "    a[i] = (int)(rand() % n);\n",
        "}\n",
        "\n",
        "// Soma os vetores na CPU\n",
        "void sumArraysOnHost(float *A, float *B, float *C, const int N, int * R, int use) {\n",
        "  for (int i = 0; i < N; i++)\n",
        "    {\n",
        "        if(use==0) {\n",
        "          C[i] = A[i] + B[i];\n",
        "        } else {\n",
        "          C[i] = A[R[i]] + B[R[i]];\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "// Soma os vetores na GPU\n",
        "__global__ void sumArraysOnGPU(float *A, float *B, float *C, int *R, int use) {\n",
        "  int i = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "\n",
        "  if(use==0) {\n",
        "    C[i] = A[i] + B[i];\n",
        "  } else {\n",
        "    C[i] = A[R[i]] + B[R[i]];\n",
        "  }\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"%s Starting...\\n\", argv[0]);\n",
        "  \n",
        "  // Configura o tamanho dos vetores\n",
        "  int nElem = atoi(argv[1]);\n",
        "  int nThread = atoi(argv[2]);\n",
        "  int use = atoi(argv[3]);\n",
        "\n",
        "  printf(\"Vector size %d\\n\", nElem);\n",
        "  \n",
        "  // Aloca memória no host\n",
        "  size_t nBytes = nElem * sizeof(float);\n",
        "  float *h_A, *h_B, *hostRef, *gpuRef;\n",
        "  int *h_R;\n",
        "  h_A = (float *)malloc(nBytes);\n",
        "  h_B = (float *)malloc(nBytes);\n",
        "  h_R = (int *)malloc(nElem * sizeof(int));\n",
        "\n",
        "  hostRef = (float *)malloc(nBytes);\n",
        "  gpuRef = (float *)malloc(nBytes);\n",
        "\n",
        "  // Inicializa os dados ainda no host\n",
        "  initialData(h_A, nElem);\n",
        "  initialData(h_B, nElem);\n",
        "  \n",
        "  // Aloca memória global na GPU (device)\n",
        "  float *d_A, *d_B, *d_C;\n",
        "  int *d_R;\n",
        "\n",
        "  cudaMalloc((float **)&d_A, nBytes);\n",
        "  cudaMalloc((float **)&d_B, nBytes);\n",
        "  cudaMalloc((float **)&d_C, nBytes);\n",
        "  cudaMalloc((int **)&d_R, nElem * sizeof(int));\n",
        "\n",
        "  // Transfere os dados do host (CPU) para o device (GPU)\n",
        "  cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);\n",
        "  cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);\n",
        "\n",
        "  cudaMemcpy(d_R, h_R, nElem * sizeof(int), cudaMemcpyHostToDevice);\n",
        "\n",
        "  // Configura os tamanho de:\n",
        "  //   bloco (quantas threads por bloco?)\n",
        "  //   e grade (quantos blocos por grade?)\n",
        "  dim3 block(1024);\n",
        "  dim3 grid(nElem >> 10);\n",
        "\n",
        "  // Executa a soma no device\n",
        "  sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C, d_R, use);\n",
        "  printf(\"Execution configuration <<<%d, %d>>>\\n\", grid.x, block.x);\n",
        "\n",
        "  // Copia o resultado do kernel do device para o host\n",
        "  cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);\n",
        "\n",
        "  // Soma os vetores na CPU para conferir os resultados\n",
        "  sumArraysOnHost(h_A, h_B, hostRef, nElem, h_R, use);\n",
        "  // Compara os resultados\n",
        "  checkResult(hostRef, gpuRef, nElem);\n",
        "\n",
        "  // Libera a memória da GPU\n",
        "  cudaFree(d_A);\n",
        "  cudaFree(d_B);\n",
        "  cudaFree(d_C);\n",
        "\n",
        "  // Libera a memória da CPU\n",
        "  free(h_A);\n",
        "  free(h_B);\n",
        "  free(hostRef);\n",
        "  free(gpuRef);\n",
        "\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Overwriting /content/src/sumR.cu\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "PZL1IB8E8lu4",
        "colab_type": "text"
      },
      "source": [
        "5. Execute com o `nvprof` com cada combinação de quantidade de elementos (mesmos valores do exercício 2 da Tarefa 3) e opção de usar ou não o vetor `R`. Elabore um gráfico do tempo de execução em função do tamanho do vetor, fazendo uma curva para o acesso direto dos índices e outra para o acesso via o vetor aleatório `R`."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "GOHMI_NnikKI",
        "colab_type": "text"
      },
      "source": [
        "Gráfico: https://drive.google.com/open?id=1XXiba0nASpDntFe_iAabOSWmlOimwX2h"
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "iyLp0rWu7EF9",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 1000
        },
        "outputId": "006336d3-48da-42df-c67b-8be1ca25be78"
      },
      "source": [
        "%cd /content/src\n",
        "!nvcc -O2 -arch=sm_35 -o sumR sumR.cu\n",
        "\n",
        "!nvprof --print-gpu-summary ./sumR 32000000 1024 0\n",
        "!nvprof --print-gpu-summary ./sumR 32000000 1024 1\n",
        "!nvprof --print-gpu-summary ./sumR 64000000 1024 0\n",
        "!nvprof --print-gpu-summary ./sumR 64000000 1024 1\n",
        "!nvprof --print-gpu-summary ./sumR 128000000 1024 0\n",
        "!nvprof --print-gpu-summary ./sumR 128000000 1024 1\n",
        "!nvprof --print-gpu-summary ./sumR 256000000 1024 0\n",
        "!nvprof --print-gpu-summary ./sumR 256000000 1024 1"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "./sumR Starting...\n",
            "Vector size 32000000\n",
            "==11521== NVPROF is profiling process 11521, command: ./sumR 32000000 1024 0\n",
            "Execution configuration <<<31250, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11521== Profiling application: ./sumR 32000000 1024 0\n",
            "==11521== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   55.87%  86.927ms         1  86.927ms  86.927ms  86.927ms  [CUDA memcpy DtoH]\n",
            "                   42.09%  65.486ms         3  21.829ms  18.234ms  29.006ms  [CUDA memcpy HtoD]\n",
            "                    2.05%  3.1877ms         1  3.1877ms  3.1877ms  3.1877ms  sumArraysOnGPU(float*, float*, float*, int*, int)\n",
            "./sumR Starting...\n",
            "Vector size 32000000\n",
            "==11532== NVPROF is profiling process 11532, command: ./sumR 32000000 1024 1\n",
            "Execution configuration <<<31250, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11532== Profiling application: ./sumR 32000000 1024 1\n",
            "==11532== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   54.85%  80.221ms         1  80.221ms  80.221ms  80.221ms  [CUDA memcpy DtoH]\n",
            "                   43.13%  63.083ms         3  21.028ms  18.145ms  26.471ms  [CUDA memcpy HtoD]\n",
            "                    2.02%  2.9546ms         1  2.9546ms  2.9546ms  2.9546ms  sumArraysOnGPU(float*, float*, float*, int*, int)\n",
            "./sumR Starting...\n",
            "Vector size 64000000\n",
            "==11543== NVPROF is profiling process 11543, command: ./sumR 64000000 1024 0\n",
            "Execution configuration <<<62500, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11543== Profiling application: ./sumR 64000000 1024 0\n",
            "==11543== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   53.81%  162.60ms         1  162.60ms  162.60ms  162.60ms  [CUDA memcpy DtoH]\n",
            "                   44.08%  133.19ms         3  44.398ms  36.132ms  60.769ms  [CUDA memcpy HtoD]\n",
            "                    2.11%  6.3755ms         1  6.3755ms  6.3755ms  6.3755ms  sumArraysOnGPU(float*, float*, float*, int*, int)\n",
            "./sumR Starting...\n",
            "Vector size 64000000\n",
            "==11554== NVPROF is profiling process 11554, command: ./sumR 64000000 1024 1\n",
            "Execution configuration <<<62500, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11554== Profiling application: ./sumR 64000000 1024 1\n",
            "==11554== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   53.25%  163.49ms         1  163.49ms  163.49ms  163.49ms  [CUDA memcpy DtoH]\n",
            "                   44.82%  137.60ms         3  45.868ms  37.035ms  62.904ms  [CUDA memcpy HtoD]\n",
            "                    1.92%  5.9008ms         1  5.9008ms  5.9008ms  5.9008ms  sumArraysOnGPU(float*, float*, float*, int*, int)\n",
            "./sumR Starting...\n",
            "Vector size 128000000\n",
            "==11568== NVPROF is profiling process 11568, command: ./sumR 128000000 1024 0\n",
            "Execution configuration <<<125000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11568== Profiling application: ./sumR 128000000 1024 0\n",
            "==11568== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   52.86%  323.53ms         1  323.53ms  323.53ms  323.53ms  [CUDA memcpy DtoH]\n",
            "                   45.06%  275.76ms         3  91.920ms  77.352ms  113.64ms  [CUDA memcpy HtoD]\n",
            "                    2.08%  12.705ms         1  12.705ms  12.705ms  12.705ms  sumArraysOnGPU(float*, float*, float*, int*, int)\n",
            "./sumR Starting...\n",
            "Vector size 128000000\n",
            "==11579== NVPROF is profiling process 11579, command: ./sumR 128000000 1024 1\n",
            "Execution configuration <<<125000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11579== Profiling application: ./sumR 128000000 1024 1\n",
            "==11579== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   53.22%  317.46ms         1  317.46ms  317.46ms  317.46ms  [CUDA memcpy DtoH]\n",
            "                   44.80%  267.23ms         3  89.078ms  74.257ms  115.69ms  [CUDA memcpy HtoD]\n",
            "                    1.98%  11.813ms         1  11.813ms  11.813ms  11.813ms  sumArraysOnGPU(float*, float*, float*, int*, int)\n",
            "./sumR Starting...\n",
            "Vector size 256000000\n",
            "==11593== NVPROF is profiling process 11593, command: ./sumR 256000000 1024 0\n",
            "Execution configuration <<<250000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11593== Profiling application: ./sumR 256000000 1024 0\n",
            "==11593== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   53.97%  638.56ms         1  638.56ms  638.56ms  638.56ms  [CUDA memcpy DtoH]\n",
            "                   43.88%  519.15ms         3  173.05ms  145.98ms  222.57ms  [CUDA memcpy HtoD]\n",
            "                    2.15%  25.495ms         1  25.495ms  25.495ms  25.495ms  sumArraysOnGPU(float*, float*, float*, int*, int)\n",
            "./sumR Starting...\n",
            "Vector size 256000000\n",
            "==11607== NVPROF is profiling process 11607, command: ./sumR 256000000 1024 1\n",
            "Execution configuration <<<250000, 1024>>>\n",
            "Arrays match.\n",
            "\n",
            "==11607== Profiling application: ./sumR 256000000 1024 1\n",
            "==11607== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   52.94%  638.07ms         1  638.07ms  638.07ms  638.07ms  [CUDA memcpy DtoH]\n",
            "                   45.10%  543.59ms         3  181.20ms  153.66ms  235.25ms  [CUDA memcpy HtoD]\n",
            "                    1.96%  23.618ms         1  23.618ms  23.618ms  23.618ms  sumArraysOnGPU(float*, float*, float*, int*, int)\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "vaIniQMi9MII",
        "colab_type": "text"
      },
      "source": [
        "6. Discuta sobre o impacto do acesso aleatório no desempenho."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "tls0981RioKG",
        "colab_type": "text"
      },
      "source": [
        "> É levemente perceptível no gráfico já apresentado que a linha indicadora do tempo de execução do algoritmo sem vetor de números aleatórios está acima da linha indicadora do algoritmo com vetor de números aleatórios. Isso ocorre porque quando não se calcula previamente os valores aleatórios, a GPU deixa de executar totalmente em paralelo, já que tem que despender parte de seus recursos para calcular o número aleatório almejado. No entanto, comparado qualquer um dos dois algoritmos com um que não faça uso de acesso aleatório de qualquer espécie, ambos os dois algoritmos demorarão muito mais para serem executados, já que eles fazem muito mal aproveitamento do princípio de localidade da cache."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "5NpPufDM9QP-",
        "colab_type": "text"
      },
      "source": [
        "## Exemplo 6 // Memória Compartilhada\n",
        "\n",
        "_Referência: Capítulo 5 - Shared Memory and Constant Memory_\n",
        "\n",
        "A memória compartilhada - _shared memory_ - pode ser vista como uma memória cache gerenciada pelo programador, e geralmente serve também como canal de comunicação intrabloco para as threads, ou mesmo uma memória _scratchpad_ para auxiliar na transformação de dados.\n",
        "\n",
        "Para atingir maiores taxas de transferência, essa memória é organizada em 32 módulos de memória de mesmo tamanho que podem ser acessados simultaneamente, chamados bancos. São 32 porque há 32 threads em um _warp_. Se uma operação de load ou store lançada num _warp_ não acessa mais de um endereço de memória por banco, a operação é servida com apenas uma transferência de memória. Do contrário, as transferências são serializadas, diminuindo o aproveitamento da taxa de transferência.\n",
        "\n",
        "![Acessos na memória compartilhada](https://docs.google.com/uc?export=download&id=1aSip-o4c5J9WorgJV3Lar9ftnqlpRfia)\n",
        "\n",
        "Na figura, vemos primeiro um caso onde há apenas um acesso num mesmo banco, e depois um caso de muitos acessos em um mesmo banco."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "xV0vn8R4CvCA",
        "colab_type": "text"
      },
      "source": [
        "## Tarefa 6 // Memória Compartilhada\n",
        "\n",
        "Nesta tarefa, estamos interessados apenas no impacto do desempenho causado pelo conflito de bancos, não na operação em si que será realizada na GPU. Portanto, você pode usar a código de soma de vetores como base, mas não usaremos a função que verifica os resultados, nem precisamos repetir os cálculos na CPU, já que não faremos comparações.\n",
        "\n",
        "Vejamos como usar a memória compartilhada. Dentro de um kernel, se cada thread acessar um banco diferente, como segue, não haverá conflito de bancos:\n",
        "```cpp\n",
        "__shared__ float smem[512];\n",
        "smem[threadIdx.x] += 1.0;\n",
        "```\n",
        "Mas se fizermos com que múltiplas threads acessem um mesmo banco, em posições diferentes, teremos conflito:\n",
        "```cpp\n",
        "__shared__ float smem[512];\n",
        "smem[threadIdx.x * 2] += 1.0;\n",
        "```\n",
        "Nesse caso, se lançarmos o kernel com 32 threads, a thread 0 e a thread 16 serão mapeadas para as posições 0 e 32, respectivamente, que estão no mesmo banco de memória compartilhada, o banco 0.\n",
        "\n",
        "1. Elabore um código em que você possa testar o conflito nos bancos de forma parametrizada (pode ser com base na soma de vetores ou não). Você pode deixar fixa a quantidade de threads em 32."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "3plql7MhHQYe",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 34
        },
        "outputId": "a837eb7b-ae93-4890-9cc3-a8afd575a9ea"
      },
      "source": [
        "%%writefile /content/src/sumConflit.cu\n",
        "#include <cuda_runtime.h>\n",
        "#include <stdio.h>\n",
        "\n",
        " __shared__ float smem[32];\n",
        "\n",
        "// Soma os vetores na GPU\n",
        "__global__ void confliGPU(int conflit) {\n",
        "  float a = 3.878 + 2.76;\n",
        "  smem[(threadIdx.x * conflit) % 32] += a;\n",
        "\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"%s Starting...\\n\", argv[0]);\n",
        "  \n",
        "  int nThread = atoi(argv[1]);\n",
        "  int conflit = atoi(argv[2]);\n",
        "\n",
        "  // Executa a soma no device\n",
        "  confliGPU<<<1, 32>>>(conflit);\n",
        "\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Overwriting /content/src/sumConflit.cu\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "HA0SkVeJmZPR",
        "colab_type": "text"
      },
      "source": [
        "2. Execute seu código com o `nvprof` para alguns valores de quantidade de conflitos (2, 4, 8 e 16). Elabore um gráfico que mostre o tempo de execução em função da quantidade de conflitos. "
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "gCTEgyRGZ6oX",
        "colab_type": "text"
      },
      "source": [
        "Gráfico: https://drive.google.com/open?id=1TCeosStLEdLJ_M9GuCZDxKdW3QF3nNMy"
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "KochAYdSH4Rt",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 538
        },
        "outputId": "3e8ed72b-8bdd-40fb-eae6-8283404cbd06"
      },
      "source": [
        "%cd /content/src\n",
        "!nvcc -O2 -arch=sm_35 -o sumConflit sumConflit.cu\n",
        "\n",
        "!nvprof --print-gpu-summary ./sumConflit 32 1\n",
        "!nvprof --print-gpu-summary ./sumConflit 32 2\n",
        "!nvprof --print-gpu-summary ./sumConflit 32 4\n",
        "!nvprof --print-gpu-summary ./sumConflit 32 8\n",
        "!nvprof --print-gpu-summary ./sumConflit 32 16"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "./sumConflit Starting...\n",
            "==16549== NVPROF is profiling process 16549, command: ./sumConflit 32 1\n",
            "==16549== Profiling application: ./sumConflit 32 1\n",
            "==16549== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:  100.00%  2.8160us         1  2.8160us  2.8160us  2.8160us  confliGPU(int)\n",
            "./sumConflit Starting...\n",
            "==16563== NVPROF is profiling process 16563, command: ./sumConflit 32 2\n",
            "==16563== Profiling application: ./sumConflit 32 2\n",
            "==16563== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:  100.00%  2.8160us         1  2.8160us  2.8160us  2.8160us  confliGPU(int)\n",
            "./sumConflit Starting...\n",
            "==16574== NVPROF is profiling process 16574, command: ./sumConflit 32 4\n",
            "==16574== Profiling application: ./sumConflit 32 4\n",
            "==16574== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:  100.00%  2.9120us         1  2.9120us  2.9120us  2.9120us  confliGPU(int)\n",
            "./sumConflit Starting...\n",
            "==16585== NVPROF is profiling process 16585, command: ./sumConflit 32 8\n",
            "==16585== Profiling application: ./sumConflit 32 8\n",
            "==16585== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:  100.00%  2.7840us         1  2.7840us  2.7840us  2.7840us  confliGPU(int)\n",
            "./sumConflit Starting...\n",
            "==16596== NVPROF is profiling process 16596, command: ./sumConflit 32 16\n",
            "==16596== Profiling application: ./sumConflit 32 16\n",
            "==16596== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:  100.00%  2.7840us         1  2.7840us  2.7840us  2.7840us  confliGPU(int)\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "Ewys6WNonRjH",
        "colab_type": "text"
      },
      "source": [
        "3. Discuta sobre o impacto no desempenho com o aumento dos conflitos."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "MFLjlJ9-aCWx",
        "colab_type": "text"
      },
      "source": [
        "> Com os conflitos, o tempo de execução tendem a não se reduzir ou até mesmo \n",
        "aumetarem. Isso acontece porque quando um conflito acontece, o hardware é reposnsável de trata-los evitando que algum erro de baixo nível aconteça. No entanto, essse tratamento dispende tempo, ou seja, quanto mais conflitos, mais tempo é gasto para resolvê-los."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "k94mg7IwezO3",
        "colab_type": "text"
      },
      "source": [
        "## Tarefa 7 // Registros\n",
        "\n",
        "As variáveis locais e as variáveis de \"controle\" - `threadIdx`, `blockIdx`, etc. - do kernel são preferencialmente alocadas em registradores. Caso se tenha uma demanda maior que a quantidade de registradores alocados para uma thread, essas variáveis podem ser alocadas na memória global de forma transparente para o programador, porém isso pode gerar perda de desempenho, mesmo com a ajuda das caches L1 e L2.\n",
        "\n",
        "Vamos ver esse impacto criando mais variáveis locais no kernel. Suponha que temos um vetor local, como segue:\n",
        "```cpp\n",
        "__global__ void sumArraysOnGPU(float *A, float *B, float *C, int N) {\n",
        "  int i = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "  float local[4];\n",
        "  for (int j = 0; j < 4; j++)\n",
        "    local[j] = 2 * A[(i + j) % N];\n",
        "  C[i] = A[i] + B[i] + local[i % 4];\n",
        "}\n",
        "```\n",
        "Note que, assim como na tarefa anterior, não estamos mais interessados na operação em si realizada, vamos apenas ver as alterações no desempenho.\n",
        "\n",
        "1. Modifique o código da soma de vetores para usar o vetor local como mostrado. Faça também uma outra versão do kernel, no mesmo código, com outro nome, e que use quatro variáveis locais ao invés de um vetor de quatro posições. Deixe seu código parametrizado de forma que se possa escolher qual desses dois kernels será chamado."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "1ByFL65ON_G3",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 34
        },
        "outputId": "8e279d9f-75d1-4324-87e4-f19019c89dab"
      },
      "source": [
        "%%writefile /content/src/sumVariableVector.cu\n",
        "#include <cuda_runtime.h>\n",
        "#include <stdio.h>\n",
        "\n",
        "// Use esta função para atribuir valores iniciais aos vetores de floats\n",
        "void initialData(float *a, int n) {\n",
        "  srand(time(NULL));\n",
        "  for (int i = 0; i < n; i++)\n",
        "    a[i] = (float)(rand() & 0xFF) / 10.0;\n",
        "}\n",
        "\n",
        "__global__ void sumArraysOnGPUVector(float *A, float *B, float *C, int N) {\n",
        "  int i = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "  float local[4];\n",
        "  for (int j = 0; j < 4; j++)\n",
        "    local[j] = 2 * A[(i + j) % N];\n",
        "  C[i] = A[i] + B[i] + local[i % 4];\n",
        "}\n",
        "\n",
        "__global__ void sumArraysOnGPUVariables(float *A, float *B, float *C, int N) {\n",
        "  int i = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "  float local0, local1, local2, local3;\n",
        "\n",
        "  local0 = 2 * A[(i + 0) % N];\n",
        "  local1 = 2 * A[(i + 1) % N];\n",
        "  local2 = 2 * A[(i + 2) % N];\n",
        "  local3 = 2 * A[(i + 3) % N];\n",
        "\n",
        "  C[i] = A[i] + B[i] + local0 + local1 + local2+ local3;\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"%s Starting...\\n\", argv[0]);\n",
        "  \n",
        "  // Configura o tamanho dos vetores\n",
        "  int nElem = atoi(argv[1]);\n",
        "  int nThread = atoi(argv[2]);\n",
        "  int use = atoi(argv[3]);\n",
        "\n",
        "  printf(\"Vector size %d\\n\", nElem);\n",
        "  \n",
        "  // Aloca memória no host\n",
        "  size_t nBytes = nElem * sizeof(float);\n",
        "  float *h_A, *h_B, *hostRef, *gpuRef;\n",
        "  h_A = (float *)malloc(nBytes);\n",
        "  h_B = (float *)malloc(nBytes);\n",
        "  hostRef = (float *)malloc(nBytes);\n",
        "  gpuRef = (float *)malloc(nBytes);\n",
        "\n",
        "  // Inicializa os dados ainda no host\n",
        "  initialData(h_A, nElem);\n",
        "  initialData(h_B, nElem);\n",
        "  \n",
        "  // Aloca memória global na GPU (device)\n",
        "  float *d_A, *d_B, *d_C;\n",
        "  cudaMalloc((float **)&d_A, nBytes);\n",
        "  cudaMalloc((float **)&d_B, nBytes);\n",
        "  cudaMalloc((float **)&d_C, nBytes);\n",
        "\n",
        "  // Transfere os dados do host (CPU) para o device (GPU)\n",
        "  cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);\n",
        "  cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);\n",
        "\n",
        "  // Configura os tamanho de:\n",
        "  //   bloco (quantas threads por bloco?)\n",
        "  //   e grade (quantos blocos por grade?)\n",
        "  dim3 block(1024);\n",
        "  dim3 grid(nElem >> 10);\n",
        "\n",
        "  // Executa a soma no device\n",
        "  if(use==1) {\n",
        "      sumArraysOnGPUVector<<<grid, block>>>(d_A, d_B, d_C, nElem);\n",
        "      printf(\"Vector allocated in GPU\\n\");\n",
        "  } else {\n",
        "      printf(\"Variables allocated in GPU\\n\");\n",
        "      sumArraysOnGPUVariables<<<grid, block>>>(d_A, d_B, d_C, nElem);\n",
        "  }\n",
        "\n",
        "  printf(\"Execution configuration <<<%d, %d>>>\\n\", grid.x, block.x);\n",
        "\n",
        "  // Copia o resultado do kernel do device para o host\n",
        "  cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);\n",
        "\n",
        "  // Libera a memória da GPU\n",
        "  cudaFree(d_A);\n",
        "  cudaFree(d_B);\n",
        "  cudaFree(d_C);\n",
        "\n",
        "  // Libera a memória da CPU\n",
        "  free(h_A);\n",
        "  free(h_B);\n",
        "  free(hostRef);\n",
        "  free(gpuRef);\n",
        "\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Overwriting /content/src/sumVariableVector.cu\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "wXWNmLmgkczr",
        "colab_type": "text"
      },
      "source": [
        "2. Escolha um tamanho de vetor, de grade, e de bloco, e execute seu código duas vezes, um para cada kernel. Compare os tempos de execução e discuta sobre."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "XFe6-GGrmaz2",
        "colab_type": "text"
      },
      "source": [
        "> É possível notar que quando utilizadas 4 varáveis ao invés de um vetor de 4 posições o desempenho foi bem pior. Isso foi causado tanto pelo demorado acesso a memoria global quanto pela não exploração do princípio de localidade."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "bY2O4xnL4fGI",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 403
        },
        "outputId": "e20a5ce1-66fc-498a-8f25-6cfd9eeeb5ff"
      },
      "source": [
        "%cd /content/src\n",
        "!nvcc -O2 -arch=sm_35 -o sumVariableVector sumVariableVector.cu\n",
        "\n",
        "!nvprof --print-gpu-summary ./sumVariableVector 1024 32 1\n",
        "!nvprof --print-gpu-summary ./sumVariableVector 1024 32 0"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "./sumVariableVector Starting...\n",
            "Vector size 1024\n",
            "==323== NVPROF is profiling process 323, command: ./sumVariableVector 1024 32 1\n",
            "Vector allocated in GPU\n",
            "Execution configuration <<<1, 1024>>>\n",
            "==323== Profiling application: ./sumVariableVector 1024 32 1\n",
            "==323== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   43.02%  6.1120us         1  6.1120us  6.1120us  6.1120us  sumArraysOnGPUVector(float*, float*, float*, int)\n",
            "                   36.49%  5.1840us         2  2.5920us  2.4320us  2.7520us  [CUDA memcpy HtoD]\n",
            "                   20.50%  2.9120us         1  2.9120us  2.9120us  2.9120us  [CUDA memcpy DtoH]\n",
            "./sumVariableVector Starting...\n",
            "Vector size 1024\n",
            "==334== NVPROF is profiling process 334, command: ./sumVariableVector 1024 32 0\n",
            "Variables allocated in GPU\n",
            "Execution configuration <<<1, 1024>>>\n",
            "==334== Profiling application: ./sumVariableVector 1024 32 0\n",
            "==334== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   56.92%  14.080us         1  14.080us  14.080us  14.080us  [CUDA memcpy DtoH]\n",
            "                   21.99%  5.4400us         1  5.4400us  5.4400us  5.4400us  sumArraysOnGPUVariables(float*, float*, float*, int)\n",
            "                   21.09%  5.2160us         2  2.6080us  2.4320us  2.7840us  [CUDA memcpy HtoD]\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "wq_5h0lwkfJw",
        "colab_type": "text"
      },
      "source": [
        "3. Modifique novamente o código, dessa vez alocando um vetor de quatro posições no `main` e passando-o como argumento para o kernel."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "mZd-5WJz5GM9",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 34
        },
        "outputId": "f52c3a58-6914-46c4-b0cd-9826812135ed"
      },
      "source": [
        "%%writefile /content/src/sumVectorMain.cu\n",
        "#include <cuda_runtime.h>\n",
        "#include <stdio.h>\n",
        "\n",
        "// Use esta função para atribuir valores iniciais aos vetores de floats\n",
        "void initialData(float *a, int n) {\n",
        "  srand(time(NULL));\n",
        "  for (int i = 0; i < n; i++)\n",
        "    a[i] = (float)(rand() & 0xFF) / 10.0;\n",
        "}\n",
        "\n",
        "__global__ void sumArraysOnGPUVector(float *A, float *B, float *C, float *local, int N) {\n",
        "  int i = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "  \n",
        "  for (int j = 0; j < 4; j++)\n",
        "    local[j] = 2 * A[(i + j) % N];\n",
        "  C[i] = A[i] + B[i] + local[i % 4];\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"%s Starting...\\n\", argv[0]);\n",
        "  \n",
        "  // Configura o tamanho dos vetores\n",
        "  int nElem = atoi(argv[1]);\n",
        "  int nThread = atoi(argv[2]);\n",
        "\n",
        "  printf(\"Vector size %d\\n\", nElem);\n",
        "  \n",
        "  // Aloca memória no host\n",
        "  size_t nBytes = nElem * sizeof(float);\n",
        "  float *h_A, *h_B, *h_local, *hostRef, *gpuRef;\n",
        "  h_A = (float *)malloc(nBytes);\n",
        "  h_B = (float *)malloc(nBytes);\n",
        "\n",
        "  // Vector for question\n",
        "  h_local = (float *) malloc(4 * sizeof(float));\n",
        "\n",
        "  hostRef = (float *)malloc(nBytes);\n",
        "  gpuRef = (float *)malloc(nBytes);\n",
        "\n",
        "  // Inicializa os dados ainda no host\n",
        "  initialData(h_A, nElem);\n",
        "  initialData(h_B, nElem);\n",
        "  \n",
        "  // Aloca memória global na GPU (device)\n",
        "  float *d_A, *d_B, *d_C, *d_local;\n",
        "  cudaMalloc((float **)&d_A, nBytes);\n",
        "  cudaMalloc((float **)&d_B, nBytes);\n",
        "  cudaMalloc((float **)&d_C, nBytes);\n",
        "\n",
        "  // Vector for question\n",
        "  cudaMalloc((float **)&d_local, 4 * sizeof(float));\n",
        "\n",
        "  // Transfere os dados do host (CPU) para o device (GPU)\n",
        "  cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);\n",
        "  cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);\n",
        "\n",
        "  // Vector for question\n",
        "  cudaMemcpy(d_local, h_local, 4 * sizeof(float), cudaMemcpyHostToDevice);\n",
        "\n",
        "  // Configura os tamanho de:\n",
        "  //   bloco (quantas threads por bloco?)\n",
        "  //   e grade (quantos blocos por grade?)\n",
        "  dim3 block(1024);\n",
        "  dim3 grid(nElem >> 10);\n",
        "\n",
        "  // Executa a soma no device\n",
        "  printf(\"Vector allocated in Main Code\\n\");\n",
        "  sumArraysOnGPUVector<<<grid, block>>>(d_A, d_B, d_C, d_local, nElem);\n",
        "\n",
        "  printf(\"Execution configuration <<<%d, %d>>>\\n\", grid.x, block.x);\n",
        "\n",
        "  // Copia o resultado do kernel do device para o host\n",
        "  cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);\n",
        "\n",
        "  // Libera a memória da GPU\n",
        "  cudaFree(d_A);\n",
        "  cudaFree(d_B);\n",
        "  cudaFree(d_C);\n",
        "\n",
        "  // Libera a memória da CPU\n",
        "  free(h_A);\n",
        "  free(h_B);\n",
        "  free(hostRef);\n",
        "  free(gpuRef);\n",
        "\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Overwriting /content/src/sumVectorMain.cu\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "ZpOt9WaGkiLj",
        "colab_type": "text"
      },
      "source": [
        "4. Execute com a mesma configuração de tamanhos no exercício 2 e compare o tempo de execução com os tempos previamente obtidos."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "jNwqvZ8BnY-l",
        "colab_type": "text"
      },
      "source": [
        "> Embora que tenha sido necessário enviar para o Kernel uma quantidade maior de dados, o tempo de execução foi tão bom quanto do melhor caso explorado no exercício 2. Isso provavelmente ocorreu porque não foi necessário um acesso a memória global e o princípio de localidade foi explorado."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "UVmZo8VZ583N",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 218
        },
        "outputId": "acba109f-56d3-40a8-ec2f-1908f03a61db"
      },
      "source": [
        "%cd /content/src\n",
        "!nvcc -O2 -arch=sm_35 -o sumVectorMain sumVectorMain.cu\n",
        "\n",
        "!nvprof --print-gpu-summary ./sumVectorMain 1024 32"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "./sumVectorMain Starting...\n",
            "Vector size 1024\n",
            "==520== NVPROF is profiling process 520, command: ./sumVectorMain 1024 32\n",
            "Vector allocated in Main Code\n",
            "Execution configuration <<<1, 1024>>>\n",
            "==520== Profiling application: ./sumVectorMain 1024 32\n",
            "==520== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   42.08%  6.7200us         3  2.2400us  1.5360us  2.7520us  [CUDA memcpy HtoD]\n",
            "                   40.08%  6.4000us         1  6.4000us  6.4000us  6.4000us  sumArraysOnGPUVector(float*, float*, float*, float*, int)\n",
            "                   17.84%  2.8480us         1  2.8480us  2.8480us  2.8480us  [CUDA memcpy DtoH]\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "yAkrfEokknMo",
        "colab_type": "text"
      },
      "source": [
        "## Tarefa 8 // Shuffle\n",
        "\n",
        "A GPU oferece também a instrução _shuffle_ para troca de valores entre threads num mesmo _warp_. Vamos comparar o desempenho do _shuffle_ com os outros recursos de memória como _shared memory_.\n",
        "\n",
        "O kernel a seguir usa a função `__shfl` para passar o valor da posição `A[i]` para a posição `B[i+1]` (perceba a soma `+1` no segundo argumento), fazendo um `loop` a cada 32 elementos (terceiro argumento).\n",
        "```cpp\n",
        "__global__ void shuffle(int *A, int *B) {\n",
        "  int i = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "  int val = A[i];\n",
        "  val = __shfl(val, threadIdx.x + 1, 32);\n",
        "  B[i] = val;\n",
        "}\n",
        "```\n",
        "Assim, supondo que os vetores A e B sejam de tamanho 64, esses seriam seus valores após a execução do kernel:\n",
        "```\n",
        "A = 0, 1, 2, ..., 30, 31, 32, 33, ..., 62, 63\n",
        "B = 1, 2, 3, ..., 31,  0, 33, 34, ..., 63, 32\n",
        "```\n",
        "\n",
        "1. Elabore o código completo que usa o kernel mostrado."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "gxWXnfrw7cJr",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 34
        },
        "outputId": "651ab909-9ea3-4db7-b476-e7a5b8415565"
      },
      "source": [
        "%%writefile /content/src/shuffleWarp.cu\n",
        "#include <cuda_runtime.h>\n",
        "#include <stdio.h>\n",
        "\n",
        "// Use esta função para atribuir valores iniciais aos vetores de ints\n",
        "void initialData(int *a, int n) {\n",
        "  srand(time(NULL));\n",
        "  for (int i = 0; i < n; i++)\n",
        "    a[i] = (int)(rand()%100);\n",
        "}\n",
        "\n",
        "__global__ void shuffle(int *A, int *B) {\n",
        "  int i = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "  int val = A[i];\n",
        "  val = __shfl_sync(val, threadIdx.x + 1, 32);\n",
        "  B[i] = val;\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv) {\n",
        "  printf(\"%s Starting...\\n\", argv[0]);\n",
        "  \n",
        "  // Configura o tamanho dos vetores\n",
        "  int nElem = atoi(argv[1]);\n",
        "  int nThread = atoi(argv[2]);\n",
        "\n",
        "  printf(\"Running shuffle shared memory\\n\");\n",
        "  printf(\"Vector size %d\\n\", nElem);\n",
        "  \n",
        "  // Aloca memória no host\n",
        "  size_t nBytes = nElem * sizeof(int);\n",
        "  int *h_A, *h_B;\n",
        "  h_A = (int *)malloc(nBytes);\n",
        "  h_B = (int *)malloc(nBytes);\n",
        "\n",
        "  // Inicializa os dados ainda no host\n",
        "  initialData(h_A, nElem);\n",
        "  initialData(h_B, nElem);\n",
        "  \n",
        "  // Aloca memória global na GPU (device)\n",
        "  int *d_A, *d_B;\n",
        "  cudaMalloc((int **)&d_A, nBytes);\n",
        "  cudaMalloc((int **)&d_B, nBytes);\n",
        "\n",
        "  // Transfere os dados do host (CPU) para o device (GPU)\n",
        "  cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);\n",
        "  cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);\n",
        "\n",
        "  // Configura os tamanho de:\n",
        "  //   bloco (quantas threads por bloco?)\n",
        "  //   e grade (quantos blocos por grade?)\n",
        "  dim3 block(1024);\n",
        "  dim3 grid(nElem >> 10);\n",
        "\n",
        "  shuffle<<<grid, block>>>(d_A, d_B);\n",
        "\n",
        "  printf(\"Execution configuration <<<%d, %d>>>\\n\", grid.x, block.x);\n",
        "\n",
        "  // Copia o resultado do kernel do device para o host\n",
        "  //cudaMemcpy(h_A, d_A, nBytes, cudaMemcpyDeviceToHost);\n",
        "  //cudaMemcpy(h_B, d_B, nBytes, cudaMemcpyDeviceToHost);\n",
        "\n",
        "  // Libera a memória da GPU\n",
        "  cudaFree(d_A);\n",
        "  cudaFree(d_B);\n",
        "\n",
        "  // Libera a memória da CPU\n",
        "  free(h_A);\n",
        "  free(h_B);\n",
        "\n",
        "  return 0;\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Overwriting /content/src/shuffleWarp.cu\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "TH6GRlXwxMs5",
        "colab_type": "text"
      },
      "source": [
        "2. Elabore o código que implementa a mesma funcionalidade, porém usando memória compartilhada."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "K2iMaDco-KRI",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 34
        },
        "outputId": "acf30e32-c589-4cef-c19a-927837c0876b"
      },
      "source": [
        "%%writefile /content/src/shuffleSharedMemory.cu\n",
        "#include \"cuda_runtime.h\"\n",
        "#include \"device_launch_parameters.h\"\n",
        "#include <cuda_profiler_api.h>\n",
        "#include <stdio.h>\n",
        "#include <stdio.h>\n",
        "#include <stdlib.h>\n",
        "#include <cuda_runtime.h>\n",
        "\n",
        "\n",
        "__inline__ __device__\n",
        "int warpReduceSum(int val) {\n",
        "    for (int offset = 16; offset > 0; offset /= 2)\n",
        "        val += __shfl_down(val, offset);\n",
        "    return val;\n",
        "}\n",
        "\n",
        "__inline__ __device__\n",
        "int blockReduceSum(int val) {\n",
        "    static __shared__ int shared[32];\n",
        "    int lane = threadIdx.x%32;\n",
        "    int wid = threadIdx.x / 32;\n",
        "    val = warpReduceSum(val);\n",
        "\n",
        "    //write reduced value to shared memory\n",
        "    if (lane == 0) shared[wid] = val;\n",
        "    __syncthreads();\n",
        "\n",
        "    //ensure we only grab a value from shared memory if that warp existed\n",
        "    val = (threadIdx.x<blockDim.x / 32) ? shared[lane] : int(0);\n",
        "    if (wid == 0) val = warpReduceSum(val);\n",
        "\n",
        "    return val;\n",
        "}\n",
        "\n",
        "__global__ void device_reduce_stable_kernel(int *in, int* out, int N) {\n",
        "    int sum = int(0);\n",
        "    //printf(\"value = %d \", blockDim.x*gridDim.x);\n",
        "    for (int i = blockIdx.x*blockDim.x + threadIdx.x; i<N; i += blockDim.x*gridDim.x) {\n",
        "        sum += in[i];\n",
        "    }\n",
        "    sum = blockReduceSum(sum);\n",
        "    if (threadIdx.x == 0)\n",
        "        out[blockIdx.x] = sum;\n",
        "}\n",
        "\n",
        "void device_reduce_stable(int *in, int* out, int N) {\n",
        "    //int threads = 512;\n",
        "    //int blocks = min((N + threads - 1) / threads, 1024);\n",
        "    const int maxThreadsPerBlock = 1024;\n",
        "    int threads = maxThreadsPerBlock;\n",
        "    int blocks = N / maxThreadsPerBlock;\n",
        "    device_reduce_stable_kernel << <blocks, threads >> >(in, out, N);\n",
        "    cudaError_t err = cudaGetLastError();\n",
        "    if (err != cudaSuccess)\n",
        "        printf(\"Error: %s\\n\", cudaGetErrorString(err));\n",
        "\n",
        "    device_reduce_stable_kernel << <1, 1024 >> >(out, out, blocks);\n",
        "    //cudaError_t err = cudaGetLastError();\n",
        "    if (err != cudaSuccess)\n",
        "        printf(\"Error: %s\\n\", cudaGetErrorString(err));\n",
        "}\n",
        "\n",
        "__global__ void global_reduce_kernel(int * d_out, int * d_in)\n",
        "{\n",
        "    int myId = threadIdx.x + blockDim.x * blockIdx.x;\n",
        "    int tid = threadIdx.x;\n",
        "\n",
        "    // do reduction in global mem\n",
        "    for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1)\n",
        "    {\n",
        "        if (tid < s)\n",
        "        {\n",
        "            d_in[myId] += d_in[myId + s];\n",
        "        }\n",
        "        __syncthreads();        // make sure all adds at one stage are done!\n",
        "    }\n",
        "\n",
        "    // only thread 0 writes result for this block back to global mem\n",
        "    if (tid == 0)\n",
        "    {\n",
        "        d_out[blockIdx.x] = d_in[myId];\n",
        "    }\n",
        "}\n",
        "\n",
        "__global__ void shmem_reduce_kernel(int * d_out, const int * d_in)\n",
        "{\n",
        "    // sdata is allocated in the kernel call: 3rd arg to <<<b, t, shmem>>>\n",
        "    extern __shared__ int sdata[];\n",
        "\n",
        "    int myId = threadIdx.x + blockDim.x * blockIdx.x;\n",
        "    int tid = threadIdx.x;\n",
        "\n",
        "    // load shared mem from global mem\n",
        "    sdata[tid] = d_in[myId];\n",
        "    __syncthreads();            // make sure entire block is loaded!\n",
        "\n",
        "    // do reduction in shared mem\n",
        "    for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1)\n",
        "    {\n",
        "        if (tid < s)\n",
        "        {\n",
        "            sdata[tid] += sdata[tid + s];\n",
        "        }\n",
        "        __syncthreads();        // make sure all adds at one stage are done!\n",
        "    }\n",
        "\n",
        "    // only thread 0 writes result for this block back to global mem\n",
        "    if (tid == 0)\n",
        "    {\n",
        "        d_out[blockIdx.x] = sdata[0];\n",
        "    }\n",
        "}\n",
        "\n",
        "void reduce(int * d_out, int * d_intermediate, int * d_in,\n",
        "    int size, bool usesSharedMemory)\n",
        "{\n",
        "    // assumes that size is not greater than maxThreadsPerBlock^2\n",
        "    // and that size is a multiple of maxThreadsPerBlock\n",
        "    const int maxThreadsPerBlock = 1024;\n",
        "    int threads = maxThreadsPerBlock;\n",
        "    int blocks = size / maxThreadsPerBlock;\n",
        "    if (usesSharedMemory)\n",
        "    {\n",
        "        shmem_reduce_kernel << <blocks, threads, threads * sizeof(int) >> >\n",
        "            (d_intermediate, d_in);\n",
        "        cudaError_t err = cudaGetLastError();\n",
        "        if (err != cudaSuccess)\n",
        "            printf(\"Error: %s\\n\", cudaGetErrorString(err));\n",
        "    }\n",
        "    else\n",
        "    {\n",
        "        global_reduce_kernel << <blocks, threads >> >\n",
        "            (d_intermediate, d_in);\n",
        "        cudaError_t err = cudaGetLastError();\n",
        "        if (err != cudaSuccess)\n",
        "            printf(\"Error: %s\\n\", cudaGetErrorString(err));\n",
        "    }\n",
        "    // now we're down to one block left, so reduce it\n",
        "    threads = blocks; // launch one thread for each block in prev step\n",
        "    blocks = 1;\n",
        "    if (usesSharedMemory)\n",
        "    {\n",
        "        shmem_reduce_kernel << <blocks, threads, threads * sizeof(int) >> >\n",
        "            (d_out, d_intermediate);\n",
        "        cudaError_t err = cudaGetLastError();\n",
        "        if (err != cudaSuccess)\n",
        "            printf(\"Error: %s\\n\", cudaGetErrorString(err));\n",
        "    }\n",
        "    else\n",
        "    {\n",
        "        global_reduce_kernel << <blocks, threads >> >\n",
        "            (d_out, d_intermediate);\n",
        "        cudaError_t err = cudaGetLastError();\n",
        "        if (err != cudaSuccess)\n",
        "            printf(\"Error: %s\\n\", cudaGetErrorString(err));\n",
        "    }\n",
        "}\n",
        "\n",
        "int main(int argc, char **argv)\n",
        "{\n",
        "    const int ARRAY_SIZE = atoi(argv[1]);\n",
        "    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(int);\n",
        "    printf(\"%s Starting...\\n\", argv[0]);\n",
        "\n",
        "    // generate the input array on the host\n",
        "    int h_in[ARRAY_SIZE];\n",
        "    int sum = 0.0f;\n",
        "    for (int i = 0; i < ARRAY_SIZE; i++) {\n",
        "        // generate random int in [-1.0f, 1.0f]\n",
        "        h_in[i] = i;\n",
        "        sum += h_in[i];\n",
        "    }\n",
        "\n",
        "    // declare GPU memory pointers\n",
        "    int * d_in, *d_intermediate, *d_out;\n",
        "\n",
        "    // allocate GPU memory\n",
        "    cudaMalloc((void **)&d_in, ARRAY_BYTES);\n",
        "    cudaMalloc((void **)&d_intermediate, ARRAY_BYTES); // overallocated\n",
        "    cudaMalloc((void **)&d_out, sizeof(int));\n",
        "\n",
        "    // transfer the input array to the GPU\n",
        "    cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);\n",
        "\n",
        "    cudaEvent_t start, stop;\n",
        "    cudaEventCreate(&start);\n",
        "    cudaEventCreate(&stop);\n",
        "    // launch the kernel\n",
        "    cudaProfilerStart();\n",
        "    \n",
        "    printf(\"Running shuffle shared memory\\n\");\n",
        "    printf(\"Vector size %d\\n\", ARRAY_SIZE);\n",
        "    cudaEventRecord(start, 0);\n",
        "    \n",
        "    reduce(d_out, d_intermediate, d_in, ARRAY_SIZE, true);\n",
        "    \n",
        "    cudaEventRecord(stop, 0);\n",
        "       \n",
        "    cudaProfilerStop();\n",
        "    cudaEventSynchronize(stop);\n",
        "    float elapsedTime;\n",
        "    cudaEventElapsedTime(&elapsedTime, start, stop);\n",
        "    elapsedTime /= 100.0f;      // 100 trials\n",
        "\n",
        "    // copy back the sum from GPU\n",
        "    //int h_out;\n",
        "    //cudaMemcpy(&h_out, d_out, sizeof(int), cudaMemcpyDeviceToHost);\n",
        "\n",
        "    //printf(\"average time elapsed: %f\\n\", elapsedTime);\n",
        "\n",
        "    // free GPU memory allocation\n",
        "    cudaFree(d_in);\n",
        "    cudaFree(d_intermediate);\n",
        "    cudaFree(d_out);\n",
        "\n",
        "    return 0;\n",
        "\n",
        "}"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "Overwriting /content/src/shuffleSharedMemory.cu\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "dIRBCUsUxgep",
        "colab_type": "text"
      },
      "source": [
        "3. Execute os dois programas com o `nvprof` com diferentes configurações de conflito para o código da memória compartilhada: sem conflitos, 2, 4 e 8 conflitos. Fixe a quantidade de threads por bloco em 32 e a quantidade de elementos no vetor a um valor à sua escolha (no mínimo 4096). Discuta sobre o desempenho de cada execução."
      ]
    },
    {
      "cell_type": "code",
      "metadata": {
        "id": "XTmA_9s181xa",
        "colab_type": "code",
        "colab": {
          "base_uri": "https://localhost:8080/",
          "height": 608
        },
        "outputId": "b67982ec-abb8-40c4-8188-0718c4eee4d6"
      },
      "source": [
        "%cd /content/src\n",
        "!nvcc -O2 -arch=sm_35 -o shuffleWarp shuffleWarp.cu\n",
        "!nvcc -O2 -arch=sm_35 -o shuffleSharedMemory shuffleSharedMemory.cu\n",
        "\n",
        "!echo \"----------------------------------------------------------------------------------------------------------------\"\n",
        "\n",
        "!nvprof --print-gpu-summary ./shuffleWarp 4096 1024\n",
        "!nvprof --print-gpu-summary ./shuffleSharedMemory 4096\n",
        "\n",
        "# Referencia: https://stackoverflow.com/questions/44278317/cuda-shuffle-instruction-reduction-slower-than-shared-memory-reduction"
      ],
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "text": [
            "/content/src\n",
            "shuffleSharedMemory.cu(13): warning: function \"__shfl_down(int, unsigned int, int)\"\n",
            "/usr/local/cuda/bin/../targets/x86_64-linux/include/sm_30_intrinsics.hpp(195): here was declared deprecated (\"__shfl_down() is deprecated in favor of __shfl_down_sync() and may be removed in a future release (Use -Wno-deprecated-declarations to suppress this warning).\")\n",
            "\n",
            "ptxas /tmp/tmpxft_00000a02_00000000-5_shuffleSharedMemory.ptx, line 58; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version\n",
            "ptxas /tmp/tmpxft_00000a02_00000000-5_shuffleSharedMemory.ptx, line 63; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version\n",
            "ptxas /tmp/tmpxft_00000a02_00000000-5_shuffleSharedMemory.ptx, line 68; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version\n",
            "ptxas /tmp/tmpxft_00000a02_00000000-5_shuffleSharedMemory.ptx, line 73; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version\n",
            "ptxas /tmp/tmpxft_00000a02_00000000-5_shuffleSharedMemory.ptx, line 78; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version\n",
            "ptxas /tmp/tmpxft_00000a02_00000000-5_shuffleSharedMemory.ptx, line 108; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version\n",
            "ptxas /tmp/tmpxft_00000a02_00000000-5_shuffleSharedMemory.ptx, line 112; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version\n",
            "ptxas /tmp/tmpxft_00000a02_00000000-5_shuffleSharedMemory.ptx, line 116; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version\n",
            "ptxas /tmp/tmpxft_00000a02_00000000-5_shuffleSharedMemory.ptx, line 120; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version\n",
            "ptxas /tmp/tmpxft_00000a02_00000000-5_shuffleSharedMemory.ptx, line 124; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version\n",
            "----------------------------------------------------------------------------------------------------------------\n",
            "./shuffleWarp Starting...\n",
            "Running shuffle shared memory\n",
            "Vector size 4096\n",
            "==2603== NVPROF is profiling process 2603, command: ./shuffleWarp 4096 1024\n",
            "Execution configuration <<<4, 1024>>>\n",
            "==2603== Profiling application: ./shuffleWarp 4096 1024\n",
            "==2603== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   79.28%  10.528us         2  5.2640us  5.1520us  5.3760us  [CUDA memcpy HtoD]\n",
            "                   20.72%  2.7520us         1  2.7520us  2.7520us  2.7520us  shuffle(int*, int*)\n",
            "./shuffleSharedMemory Starting...\n",
            "==2617== NVPROF is profiling process 2617, command: ./shuffleSharedMemory 4096\n",
            "Running shuffle shared memory\n",
            "Vector size 4096\n",
            "==2617== Profiling application: ./shuffleSharedMemory 4096\n",
            "==2617== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   66.86%  10.848us         2  5.4240us  3.5520us  7.2960us  shmem_reduce_kernel(int*, int const *)\n",
            "                   33.14%  5.3760us         1  5.3760us  5.3760us  5.3760us  [CUDA memcpy HtoD]\n"
          ],
          "name": "stdout"
        }
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "uW8LFdFYkjoM",
        "colab_type": "text"
      },
      "source": [
        "> É notável que o Warp apresenta desempenho bem superior ao SharedMemory, pórém isso não é uma surpresa, pois, enquanto o SharedMemory é uma implementação manual da solução do problema acima citado e, por isso, acaba consumindo mais recursos por estar níveis acima do nível que o warp está impleementado, pois é um recurso disponível nas GPU's. Sendo assim, alguns tratamentos que já são automaticamente tratados utilizando warp, terão de ser tratados de forma convencional, o que torna a execução do SharedMemory bem mais lenta."
      ]
    }
  ]
}