L1 - Introduction.ipynb

{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {
    "tags": []
   },
   "source": [
    "# Introduction to OpenCL\n",
    "\n",
    "<figure style=\"float:right; width:30%;\">\n",
    "    <img src=\"images/OpenCL_RGB_Apr20.svg\" alt=\"OpenCL logo\"/>\n",
    "    <figcaption style= \"text-align:lower; margin:1em; float:bottom; vertical-align:bottom;\">OpenCL and the OpenCL logo are trademarks of Apple Inc. used by permission by Khronos.</figcaption>\n",
    "</figure>\n",
    "\n",
    "OpenCL (short for Open Computing Language) is an open standard for running compute workloads on many different kinds of compute hardware (e.g CPUs, GPU's). The OpenCL trademark is held by Apple, and the standard is developed and released by the [Khronos](https://www.khronos.org) group, a non-for-profit organisation that provides a focal point for the development of royalty-free standards such as OpenGL. The OpenCL specification itself is just a document, and can be downloaded from the Khronos website [here](https://www.khronos.org/registry/OpenCL/specs/). It is then the task of compute hardware vendors to produce software implementations of OpenCL that best make use of their compute devices."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## How does OpenCL work?\n",
    "\n",
    "In order to answer how an OpenCL implementation works, we need to start thinking about hardware. In every compute device such as a CPU or GPU there are a number of cores on which software can be run. In OpenCL terminology these cores are called **Compute Units**. Each Compute Unit makes available to the operating system a number of hardware threads that can run software. In OpenCL terminology we call these hardware threads **Processing Elements**. For example, an NVIDIA GP102 die is shown below. Each die contains 30 compute units, shown contained by the orange squares. Each compute unit provides 128 processing elements (CUDA cores), so in this example there are $30\\times128 = 3840$ processing elements available for use in compute applications. \n",
    "\n",
    "<figure style=\"margin: 1em; margin-left:auto; margin-right:auto; width:70%;\">\n",
    "    <img src=\"images/compute_units.svg\">\n",
    "    <figcaption style= \"text-align:lower; margin:1em; float:bottom; vertical-align:bottom;\">NVIDIA GP102 die with compute units highlighted in orange. Image credit: <a href=\"https://www.flickr.com/photos/130561288@N04/46079430302/\")>Fritzchens Fritz</a></figcaption>\n",
    "</figure>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "During execution of an OpenCL program, processing elements each run an instance of a user-specified piece of compiled code called a **kernel**. Below is an example OpenCL C kernel that takes the absolute value of a single element of an array.\n",
    "\n",
    "```C\n",
    "__kernel void vec_fabs(\n",
    "        // Memory allocations that are on the compute device\n",
    "        __global float *src, \n",
    "        __global float *dst,\n",
    "        // Number of elements in the memory allocations\n",
    "        int length) {\n",
    "\n",
    "    // Get our position in the array\n",
    "    size_t gid0 = get_global_id(0);\n",
    "\n",
    "    // Get the absolute value of \n",
    "    if (gid0 < length) {\n",
    "        dst[gid0] = fabs(src[gid0]);\n",
    "    }\n",
    "}\n",
    "```\n",
    "\n",
    "We want to run a kernel instance for every element of the array. An OpenCL implementation is a way to run kernel instances on processing elements as they become available. The implementation also provides the means to upload and download memory to and from compute devices. We specify how many kernel instances we want at runtime by defining a 3D execution space called a **Grid** and specifying its size at kernel launch. Every point in the Grid is called a **work-item** and represents a unique invocation of the kernel. A work-item is equivalent to a single kernel invocation. This is much like defining an execution space using nested loops, however with OpenCL there are no guarantees on the order in which work items are completed.\n",
    "\n",
    "<figure style=\"margin-left:auto; margin-right:auto; width:70%;\">\n",
    "    <img style=\"vertical-align:middle\" src=\"images/grid.svg\">\n",
    "    <figcaption style= \"text-align:lower; margin:1em; float:bottom; vertical-align:bottom;\">Three-dimensional Grid with work-items and work-groups.</figcaption>\n",
    "</figure>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Work-items are executed in teams called work-groups. In the example above, the grid is of global size (10, 8, 2) and each work-group is of size (5,4,1). The the number of work-groups in each dimension is then (2,2,2). Every work item has access to device memory that it can use exclusively (**private memory**), access to memory the team can use (**local memory**), and access to memory that other teams use (**global** and **constant** memory). Every kernel invocation or work-item can query its location within the **Grid** and use that position as a reference to access allocated memory on the compute device at an appropriately calculated offset."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "<figure style=\"margin-left:auto; margin-right:auto; width:70%;\">\n",
    "    <img style=\"vertical-\n",
    "                align:middle\" src=\"images/mem_access.svg\">\n",
    "    <figcaption style= \"text-align:lower; margin:1em; float:bottom; vertical-align:bottom;\">Using the location within the Grid to access memory within a memory allocation on a GPU compute device.</figcaption>\n",
    "</figure>\n",
    "\n",
    "The above concepts form the core ideas surrounding OpenCL. Everything that follows in this course is supporting information on how to prepare compute devices, memory allocations, kernel invocations, and how best to use these concepts together to get the best performance out of your compute devices. "
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Elements of an accelerated application\n",
    "\n",
    "In every accelerated application there is the concept of a host computer with one or more attached compute devices. The host usually has the largest memory space available and the compute device usually has the most compute power and memory bandwidth. This is why we say the application is \"accelerated\" by the compute device.\n",
    "\n",
    "At runtime, the host executes the application and compiles kernels for execution on the compute device. The host manages memory allocations and submits kernels to the compute device for execution. For instances where the compute device is a CPU, the host CPU and the compute device are the same.\n",
    "\n",
    "Every accelerated application follows the same logical progression of steps: \n",
    "\n",
    "1. Compute devices discovered\n",
    "1. Kernels prepared for compute devices\n",
    "1. Memory allocated on the compute device\n",
    "1. Memory copied to the compute device\n",
    "1. Kernels run on the compute device\n",
    "1. Wait for kernels to finish\n",
    "1. Memory copied back from the computed device to the host\n",
    "1. Repeat steps 3 - 8 as many times as necessary\n",
    "1. Clean up resources and exit\n",
    "\n",
    "We now discuss the OpenCL components that make these steps possible."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Taxonomy of an OpenCL application\n",
    "\n",
    "Below is a representation of the core software components that are available to an OpenCL application.\n",
    "\n",
    "<figure style=\"margin-left:auto; margin-right:auto; width:50%;\">\n",
    "    <img style=\"vertical-\n",
    "                align:middle\" src=\"images/opencl_components.svg\">\n",
    "    <figcaption style= \"text-align:lower; margin:1em; float:bottom; vertical-align:bottom;\">Components of an OpenCL application.</figcaption>\n",
    "</figure>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "The first is the **Platform**. This is a software representation of the vendor. A platform provides access to all **devices** that the platform supports. During device discovery, available platforms must be queried before anything else. A platform provides access to one or more compute devices and possibly even a mixture of accelerator devices from the same vendor.\n",
    "\n",
    "A **Device** provides a way to query the capabilites of the compute device and provides a foundation to build a context.\n",
    "\n",
    "Surrounding the devices is a **Context**. A Context is like a registry that keeps track of everything (i.e kernel executions and memory allocations) that are happening on the compute device/s. A context is constructed on using both a platform and one or more devices on the platform. There are some benefits (such as memory copies) that could be obtained by encapsulating one or more devices under the same context, however this assumes that devices must belong to the same platform - an assumption which may not be true. A simpler and more general design is to create a unique context for every compute device.\n",
    "\n",
    "Within the context are **Buffers**. Buffers are memory allocations managed under the context, and may exist on either the host or the compute device. At runtime memory is migrated to where is needed, but you can have some control over where the buffer lives. \n",
    "\n",
    "At runtime, source code for the kernels is collated into a **Program**, and the program is compiled for every device in a context. There must be a program for every context, and every program must be compiled with knowledge of the associated devices under the context.\n",
    "\n",
    "Once a context has been created and devices are known, then one can create one or more **Command queue/s** for each device. A command queue is a place to submit work, such as kernel invocations and memory copies. \n",
    "\n",
    "A **Kernel** is a component of a compiled **Program**. At runtime we set the arguments of compiled kernels and then submit kernels to command queues for execution. We can keep track of the status of a command submitted to the command queue using an **Event**.\n",
    "\n",
    "In summary we have the following components:\n",
    "\n",
    "* **Platform**: provides access to devices\n",
    "* **Device**: represents a way to access the compute device and to query device capabilities\n",
    "* **Context**: provides a way to create Buffers and keep track of what is happening on compute devices\n",
    "* **Buffer**: provides a way to allocate memory on devices\n",
    "* **Program**: provides a way to compile kernels for each device\n",
    "* **Command queue**: provides a place to send work such as memory copy commands and kernel executions\n",
    "* **Kernel**: provides a way to do work on a compute device\n",
    "* **Event**: provides a way to keep track of work submitted to a command queue"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Specification Roadmap\n",
    "\n",
    "From [Wikipedia](https://en.wikipedia.org/wiki/OpenCL) OpenCL was originally designed by Apple, who developed a proposal to submit to the Khronos group and holds the trademarks to OpenCL. The first specification, OpenCL 1.0, was ratified on November 18, 2008 and the first public release of the standard was on December 2008. Since then a number of different versions of the standard have been released. \n",
    "\n",
    "**Version 1.1** introduced thread safety so that calls to most OpenCL functions from different threads didn't introduce race conditions. If memory allocations in buffers are used to represent 2D and 3D arrays, then Version 1.1 introduced routines to copy rectangular regions of those buffers to and from the host. \n",
    "\n",
    "**Version 1.2** is probably the most significant release of OpenCL. It remained the defacto OpenCL standard for at least 10 years. Abilities such as being able to divide the processing elements of a compute device into sub-devices that share a common cache and offline compilation of kernels were useful. Having math operations conform to the IEEE754 precision standard meant consistent results across heterogeneous compute architectures.\n",
    "\n",
    "**Version 2.0** introduced support for Shared Virtual Memory (SVM). Implementation of SVM meant we no longer needed to qualify which space (i.e global, local..) a memory allocation belonged to, and memory could be transferred to and from devices transparently to the user. This was too much for some vendors to implement however, and a few vendor implementations remained at 1.2 for a number of years.\n",
    "\n",
    "**Version 2.1** brought the SPIR-V (Standard Portable Intermediate Representation) language to OpenCL. During compilation a open-source compiler can take C or C++ kernel code and emit a compiled program as SPIR-V intermediate code. At runtime this program is loaded by the application and passed to the vendor driver for further compilation to binary code that can run on the compute device. This is a significant advancement, because if a vendor can implement support for SPIR-V then it dramatically reduces the number of intermediate representations the vendor compiler must support. It also offloads support for kernel language advancements to the open source compiler and provides a measure of security against intellectual property theft.\n",
    "\n",
    "**Version 2.2** allowed kernels to be produced using a subset of the C++14 standard. It also updated support for SPIR-V to version 1.2. The combination of shared virtual memory, C++ kernels, and SPIR-V support meant that very few vendors actually succeeded in producing viable implementations of OpenCL 2.2, and OpenCL stagnated for a period of 5 years. \n",
    "\n",
    "**Version 3.0** addressed the issue of stagnation by making Version 1.2 standard and all the other improvements in Version 2.x optional. This gave vendors freedom to implement what they wanted for customers and gave the standard some breathing room. Version 3.0 also introduced a new C++ language for kernels (called C++ for OpenCL) that uses a subset of the C++17 standard. The Clang compiler supports compilation of C++ for OpenCL kernels into SPIR-V format.\n",
    "\n",
    "Below is a summary of major features implemented with each release:"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "<table>\n",
    "    \n",
    "<tr>\n",
    "<th>Specification</th>\n",
    "<th>Release year</th>\n",
    "<th>Specifics</th>\n",
    "</tr>\n",
    "\n",
    "<tr>\n",
    "    <td>1.0</td>\n",
    "    <td>2008</td>\n",
    "    <td>Initial implementation</td>\n",
    "</tr>\n",
    "\n",
    "<tr>\n",
    "    <td>1.1</td>\n",
    "    <td>2010</td>\n",
    "    <td> \n",
    "        <ul> \n",
    "            <li> Copy rectangular sections to and from Buffers </li>\n",
    "            <li> User-defined Events </li>\n",
    "            <li> 3-component vector types </li>\n",
    "            <li> Support for making Buffers from Buffers (sub-Buffers) </li>\n",
    "            <li> Thread safety for all functions except setting kernel arguments </li>\n",
    "        </ul> \n",
    "    </td>\n",
    "</tr>\n",
    "    \n",
    "<tr>\n",
    "    <td>1.2</td>\n",
    "    <td>2011</td>\n",
    "    <td> \n",
    "        <ul> \n",
    "            <li> The ability to divide a compute device into sub-devices, that share a cache for example </li>\n",
    "            <li> Offline compilation of kernels </li>\n",
    "            <li> Support for built-in kernels (i.e for FPGA's) </li>\n",
    "            <li> IEEE754 compliance for consistent math across devices </li>\n",
    "            <li> Enabling double precision math </li>\n",
    "            <li> Support for using printf in kernels for debugging </li>\n",
    "        </ul> \n",
    "    </td>\n",
    "</tr>\n",
    "\n",
    "<tr>\n",
    "    <td>2.0</td>\n",
    "    <td>2013</td>\n",
    "    <td> \n",
    "        <ul> \n",
    "            <li> Support for Shared Virtual Memory (SVM) </li>\n",
    "            <li> The ability to run kernels from kernels </li>\n",
    "            <li> Enhanced support for Images (specialised Buffers) </li>\n",
    "            <li> Simplified atomics </li>\n",
    "            <li> Pipe storage </li>\n",
    "            <li> Double precision IEEE754 operations </li>\n",
    "        </ul> \n",
    "    </td>\n",
    "</tr>\n",
    "    \n",
    "<tr>\n",
    "    <td>2.1</td>\n",
    "    <td>2015</td>\n",
    "    <td> \n",
    "        <ul> \n",
    "            <li> Low latency timers for profiling events </li>\n",
    "            <li> The ability to make some command queues have higher priority than others </li>\n",
    "            <li> Introduces the SPIR-V 1.1 intermediate language for compiled kernels </li>\n",
    "            <li> The ability to clone kernels </li>\n",
    "        </ul> \n",
    "    </td>\n",
    "</tr>\n",
    "    \n",
    "<tr>\n",
    "    <td>2.2</td>\n",
    "    <td>2015</td>\n",
    "    <td> \n",
    "        <ul> \n",
    "            <li> A subset of C++14 for Kernels </li>\n",
    "            <li> The ability to make some command queues have higher priority than others </li>\n",
    "            <li> Updates the SPIR-V intermediate language to version 1.2 </li>\n",
    "        </ul> \n",
    "    </td>\n",
    "</tr>\n",
    "    \n",
    "<tr>\n",
    "    <td>3.0</td>\n",
    "    <td>2020</td>\n",
    "    <td> \n",
    "        <ul> \n",
    "            <li> Version 1.2 is canon, everything else is an option </li>\n",
    "            <li> The ability to call a function when a context is destroyed </li>\n",
    "            <li> Introduces C++ for OpenCL </li>\n",
    "        </ul> \n",
    "    </td>\n",
    "</tr>\n",
    "    \n",
    "</table>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Vendor implementations\n",
    "\n",
    "All of the major vendors have OpenCL implementations at varying levels of support for the OpenCL specification. The table below shows the latest known level of support for each version of the specification, along with links to the vendor's OpenCL developer page.\n",
    "\n",
    "|Vendor| 1.2 | 2.0 | 2.1 | 2.2 | 3.0 |\n",
    "| :- | :- | :- | :- | :- | :- |\n",
    "| [AMD](https://rocmdocs.amd.com/en/latest/Programming_Guides/Opencl-programming-guide.html) | Y | Y | Y | Some | N |\n",
    "| [Apple](https://developer.apple.com/opencl) | Y | N | N | N | N |\n",
    "| [ARM](https://developer.arm.com/solutions/graphics-and-gaming/apis/opencl) | Y | Y | Y | N | Y |\n",
    "| [Intel](https://www.intel.com/content/www/us/en/developer/tools/opencl-sdk/overview.html) | Y | Y | Y | Some | Y |\n",
    "| [NVIDIA](https://developer.nvidia.com/opencl) | Y | N | N | N | Y |\n",
    "| [Portable OpenCL](http://portablecl.org) | Y | Some | N | N | N |\n",
    "\n",
    "**[Apple](https://developer.apple.com/opencl)** was the original vendor for OpenCL and it comes baked into the MacOS operating system. However the company has since moved on to their proprietary framework **Metal** and they haven't invested in OpenCL beyond specification 1.2. Support for OpenCL is built in to **[NVIDIA](https://developer.nvidia.com/opencl)'s** CUDA toolkit, though after an initial flurry of development activity up to version 1.2, development stalled until version 3.0. Support for OpenCL with **[AMD](https://rocmdocs.amd.com/en/latest/Programming_Guides/Opencl-programming-guide.html)** is part of the **[ROCM](https://rocmdocs.amd.com/en/latest/Programming_Guides/Opencl-programming-guide.html)** suite. **[Intel](https://www.intel.com/content/www/us/en/developer/tools/opencl-sdk/overview.html)** strongly supports OpenCL development for CPU's and GPU's with its NEO implementation. **[ARM](https://developer.arm.com/solutions/graphics-and-gaming/apis/opencl)** has solid support for OpenCL on its Mali GPU's. The open source Portable OpenCL implementation has a CPU implementation as well as support for OpenCL on CUDA and OpenCL on MacOS.\n",
    "\n",
    "#### Conformance\n",
    "\n",
    "A conformant OpenCL implementation is an implementation of OpenCL that has passed Khronos' [test suite](https://github.com/KhronosGroup/OpenCL-CTS). The number of vendors with conformant implementations is an evolving list, click [here](https://www.khronos.org/conformance/adopters/conformant-products/opencl) to see the latest conformant implementations."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Getting help for OpenCL\n",
    "\n",
    "The best source of help for OpenCL is [Khronos OpenCL registry](https://www.khronos.org/registry/OpenCL/). There you can find excellent documentation on  the latest specification that your vendor supports. As an exercise, download the latest API specification in PDF format and have it ready as reference material."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "tags": []
   },
   "source": [
    "## Is OpenCL right for you?\n",
    "\n",
    "This is sometimes a difficult question to answer. Researchers often have diverse computing environments, in such cases OpenCL is a good fit as it will provide a solid foundation for your research tools. However if you are looking for the best possible performance and can live with vendor lock-in, then using vendor-specific tools will help with that. \n",
    "\n",
    "**Drawbacks to using OpenCL**\n",
    "\n",
    "* Can't readily utilise device-specific hardware (i.e tensor or matrix cores)?\n",
    "* When vendors have their own accelerator libraries it creates a financial incentive to prioritise development and performance of their libraries over their OpenCL implementation.\n",
    "* Buffer allocations are sometimes limited to $1/4$ or more of available device memory (vendor specific)\n",
    "* Lots of code required to set up the computation, increased potential for error\n",
    "* Paucity of vendor-supported tools for debugging and profiling\n",
    "\n",
    "**Benefits of using OpenCL**\n",
    "\n",
    "* Straightforward well-defined C API with good documentation\n",
    "* Ability to use a wide variety of hardware\n",
    "* Data types to facilitate consistent precision across implementations\n",
    "* Consistent math across implementations\n",
    "* Support for vectors of up to 16 elements\n",
    "* Open standard - the standard is not (explicitly) contingent on the wellbeing of a single vendor\n",
    "* Mature, production quality OpenCL implementations"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Compiling OpenCL programs\n",
    "\n",
    "Just to avoid confusion there are two compilation steps for OpenCL applications: \n",
    "\n",
    "1. Compiling the application before execution\n",
    "2. Compiling kernels during execution\n",
    "\n",
    "Thankfully, when compiling an OpenCL application before execution (Step 1) we don't need to link against every available implementation. We just need to link against a single library file called the **Installable Client Driver (ICD)** that may be provided by any vendor. The ICD has the name (**opencl.dll**) on Windows and (**libOpenCL.so**) on Linux. Accompanying the ICD are header files (**opencl.h** for C and **cl.hpp** for C++) that must be \"included\" from the C/C++ source code. The ICD takes care of intercepting all library calls and routing them to the appropriate vendor implementation. The routing process happens transparently to the user. "
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "tags": []
   },
   "source": [
    "## Exercise: compiling your first OpenCL application\n",
    "\n",
    "At the location [src/L1_Introduction/hello_devices.cpp](src/L1_Introduction/hello_devices.cpp) is a complete OpenCL application to obtain the size of on-device memory and the maximum Buffer size that is possible within that memory. \n",
    "\n",
    "* **Step 1.** From the Jupyter launcher start a Terminal and use cd to navigate to the src/L1_Introduction directory in the course material\n",
    "\n",
    "```bash\n",
    "cd src/L1_Introduction\n",
    "```\n",
    "\n",
    "* **Step 2.** You need to know where the OpenCL ICD loader and OpenCL header files are located. For this particular example the locations are as follows:\n",
    "\n",
    "| File | Directory |\n",
    "| :--- | :--- |\n",
    "| ICD loader (libOpenCL.so) | /usr/local/cuda/lib64 |\n",
    "| OpenCL C++ headers directory (CL) | /usr/local/cuda/include |\n",
    "\n",
    "In the Terminal use **ls** to list the contents of these directories and locate the **CL** directory in which the OpenCL header files are located. \n",
    "\n",
    "* **Step 3.** Compile the application **hello_devices.cxx** using the **g++** compiler. The compilation command should look like this:\n",
    "\n",
    "```bash\n",
    "g++ -g -O2 -I/usr/local/cuda/include -L/usr/local/cuda/lib64 hello_devices.cpp\\\n",
    "    -o hello_devices.exe -lOpenCL\n",
    "```\n",
    " \n",
    "* **Step 4.** Now run the application\n",
    "\n",
    "```bash\n",
    "./hello_devices.exe\n",
    "```\n",
    "\n",
    "You should see at least one device printed with the name and memory sizes. Now that you know how to let the compiler know about OpenCL you can use the **make** command within that directory to compile the example. "
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 8,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "rm -r *.exe\n",
      "g++ -g -O2 -fopenmp -I/usr/local/cuda/include -I../include -L/usr/local/cuda/lib64 hello_devices.cpp\\\n",
      "\t-o hello_devices.exe \"-lOpenCL\"\n",
      "In file included from \u001b[01m\u001b[Khello_devices.cpp:8\u001b[m\u001b[K:\n",
      "\u001b[01m\u001b[K/usr/local/cuda/include/CL/cl.hpp:5085:28:\u001b[m\u001b[K \u001b[01;35m\u001b[Kwarning: \u001b[m\u001b[Kignoring attributes on template argument ‘\u001b[01m\u001b[Kcl_int\u001b[m\u001b[K’ {aka ‘\u001b[01m\u001b[Kint\u001b[m\u001b[K’} [\u001b[01;35m\u001b[K\u001b]8;;https://gcc.gnu.org/onlinedocs/gcc/Warning-Options.html#index-Wignored-attributes\u0007-Wignored-attributes\u001b]8;;\u0007\u001b[m\u001b[K]\n",
      " 5085 |         VECTOR_CLASS<cl_int\u001b[01;35m\u001b[K>\u001b[m\u001b[K* binaryStatus = NULL,\n",
      "      |                            \u001b[01;35m\u001b[K^\u001b[m\u001b[K\n",
      "\u001b[01m\u001b[Khello_devices.cpp:14:29:\u001b[m\u001b[K \u001b[01;35m\u001b[Kwarning: \u001b[m\u001b[Kignoring attributes on template argument ‘\u001b[01m\u001b[Kcl_int\u001b[m\u001b[K’ {aka ‘\u001b[01m\u001b[Kint\u001b[m\u001b[K’} [\u001b[01;35m\u001b[K\u001b]8;;https://gcc.gnu.org/onlinedocs/gcc/Warning-Options.html#index-Wignored-attributes\u0007-Wignored-attributes\u001b]8;;\u0007\u001b[m\u001b[K]\n",
      "   14 | std::map<cl_int, std::string\u001b[01;35m\u001b[K>\u001b[m\u001b[K error_codes {\n",
      "      |                             \u001b[01;35m\u001b[K^\u001b[m\u001b[K\n"
     ]
    }
   ],
   "source": [
    "!cd src/L1_Introduction; make clean; make"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "tags": []
   },
   "source": [
    "This application is rather rudimentary, however there is a far more sophisticated OpenCL query application called **clinfo**. You can use it to query a great deal on information on the available devices."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 6,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Number of platforms                               3\n",
      "  Platform Name                                   NVIDIA CUDA\n",
      "  Platform Vendor                                 NVIDIA Corporation\n",
      "  Platform Version                                OpenCL 3.0 CUDA 11.4.136\n",
      "  Platform Profile                                FULL_PROFILE\n",
      "  Platform Extensions                             cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_fp64 cl_khr_3d_image_writes cl_khr_byte_addressable_store cl_khr_icd cl_khr_gl_sharing cl_nv_compiler_options cl_nv_device_attribute_query cl_nv_pragma_unroll cl_nv_copy_opts cl_nv_create_buffer cl_khr_int64_base_atomics cl_khr_int64_extended_atomics cl_khr_device_uuid cl_khr_pci_bus_info\n",
      "  Platform Extensions with Version                cl_khr_global_int32_base_atomics                                 0x400000 (1.0.0)\n",
      "                                                  cl_khr_global_int32_extended_atomics                             0x400000 (1.0.0)\n",
      "                                                  cl_khr_local_int32_base_atomics                                  0x400000 (1.0.0)\n",
      "                                                  cl_khr_local_int32_extended_atomics                              0x400000 (1.0.0)\n",
      "                                                  cl_khr_fp64                                                      0x400000 (1.0.0)\n",
      "                                                  cl_khr_3d_image_writes                                           0x400000 (1.0.0)\n",
      "                                                  cl_khr_byte_addressable_store                                    0x400000 (1.0.0)\n",
      "                                                  cl_khr_icd                                                       0x400000 (1.0.0)\n",
      "                                                  cl_khr_gl_sharing                                                0x400000 (1.0.0)\n",
      "                                                  cl_nv_compiler_options                                           0x400000 (1.0.0)\n",
      "                                                  cl_nv_device_attribute_query                                     0x400000 (1.0.0)\n",
      "                                                  cl_nv_pragma_unroll                                              0x400000 (1.0.0)\n",
      "                                                  cl_nv_copy_opts                                                  0x400000 (1.0.0)\n",
      "                                                  cl_nv_create_buffer                                              0x400000 (1.0.0)\n",
      "                                                  cl_khr_int64_base_atomics                                        0x400000 (1.0.0)\n",
      "                                                  cl_khr_int64_extended_atomics                                    0x400000 (1.0.0)\n",
      "                                                  cl_khr_device_uuid                                               0x400000 (1.0.0)\n",
      "                                                  cl_khr_pci_bus_info                                              0x400000 (1.0.0)\n",
      "  Platform Numeric Version                        0xc00000 (3.0.0)\n",
      "  Platform Extensions function suffix             NV\n",
      "  Platform Host timer resolution                  0ns\n",
      "\n",
      "  Platform Name                                   Portable Computing Language\n",
      "  Platform Vendor                                 The pocl project\n",
      "  Platform Version                                OpenCL 2.0 pocl 1.8  Linux, RelWithDebInfo, RELOC, LLVM 13.0.0, SLEEF, DISTRO, POCL_DEBUG\n",
      "  Platform Profile                                FULL_PROFILE\n",
      "  Platform Extensions                             cl_khr_icd cl_pocl_content_size\n",
      "  Platform Extensions function suffix             POCL\n",
      "\n",
      "  Platform Name                                   AMD Accelerated Parallel Processing\n",
      "  Platform Vendor                                 Advanced Micro Devices, Inc.\n",
      "  Platform Version                                OpenCL 2.2 AMD-APP (3361.0)\n",
      "  Platform Profile                                FULL_PROFILE\n",
      "  Platform Extensions                             cl_khr_icd cl_amd_event_callback \n",
      "  Platform Extensions function suffix             AMD\n",
      "  Platform Host timer resolution                  1ns\n",
      "\n",
      "  Platform Name                                   NVIDIA CUDA\n",
      "Number of devices                                 1\n",
      "  Device Name                                     NVIDIA GeForce RTX 3060\n",
      "  Device Vendor                                   NVIDIA Corporation\n",
      "  Device Vendor ID                                0x10de\n",
      "  Device Version                                  OpenCL 3.0 CUDA\n",
      "  Device UUID                                     b695a982-d4d7-8731-3bc6-36c596b643eb\n",
      "  Driver UUID                                     b695a982-d4d7-8731-3bc6-36c596b643eb\n",
      "  Valid Device LUID                               No\n",
      "  Device LUID                                     6d69-637300000000\n",
      "  Device Node Mask                                0\n",
      "  Device Numeric Version                          0xc00000 (3.0.0)\n",
      "  Driver Version                                  470.74\n",
      "  Device OpenCL C Version                         OpenCL C 1.2 \n",
      "  Device OpenCL C all versions                    OpenCL C                                                         0x400000 (1.0.0)\n",
      "                                                  OpenCL C                                                         0x401000 (1.1.0)\n",
      "                                                  OpenCL C                                                         0x402000 (1.2.0)\n",
      "                                                  OpenCL C                                                         0xc00000 (3.0.0)\n",
      "  Device OpenCL C features                        __opencl_c_fp64                                                  0xc00000 (3.0.0)\n",
      "                                                  __opencl_c_images                                                0xc00000 (3.0.0)\n",
      "                                                  __opencl_c_int64                                                 0xc00000 (3.0.0)\n",
      "                                                  __opencl_c_3d_image_writes                                       0xc00000 (3.0.0)\n",
      "  Latest comfornace test passed                   v2021-02-01-00\n",
      "  Device Type                                     GPU\n",
      "  Device Topology (NV)                            PCI-E, 0000:43:00.0\n",
      "  Device Profile                                  FULL_PROFILE\n",
      "  Device Available                                Yes\n",
      "  Compiler Available                              Yes\n",
      "  Linker Available                                Yes\n",
      "  Max compute units                               28\n",
      "  Max clock frequency                             1837MHz\n",
      "  Compute Capability (NV)                         8.6\n",
      "  Device Partition                                (core)\n",
      "    Max number of sub-devices                     1\n",
      "    Supported partition types                     None\n",
      "    Supported affinity domains                    (n/a)\n",
      "  Max work item dimensions                        3\n",
      "  Max work item sizes                             1024x1024x64\n",
      "  Max work group size                             1024\n",
      "  Preferred work group size multiple (device)     32\n",
      "  Preferred work group size multiple (kernel)     32\n",
      "  Warp size (NV)                                  32\n",
      "  Max sub-groups per work group                   0\n",
      "  Preferred / native vector sizes                 \n",
      "    char                                                 1 / 1       \n",
      "    short                                                1 / 1       \n",
      "    int                                                  1 / 1       \n",
      "    long                                                 1 / 1       \n",
      "    half                                                 0 / 0        (n/a)\n",
      "    float                                                1 / 1       \n",
      "    double                                               1 / 1        (cl_khr_fp64)\n",
      "  Half-precision Floating-point support           (n/a)\n",
      "  Single-precision Floating-point support         (core)\n",
      "    Denormals                                     Yes\n",
      "    Infinity and NANs                             Yes\n",
      "    Round to nearest                              Yes\n",
      "    Round to zero                                 Yes\n",
      "    Round to infinity                             Yes\n",
      "    IEEE754-2008 fused multiply-add               Yes\n",
      "    Support is emulated in software               No\n",
      "    Correctly-rounded divide and sqrt operations  Yes\n",
      "  Double-precision Floating-point support         (cl_khr_fp64)\n",
      "    Denormals                                     Yes\n",
      "    Infinity and NANs                             Yes\n",
      "    Round to nearest                              Yes\n",
      "    Round to zero                                 Yes\n",
      "    Round to infinity                             Yes\n",
      "    IEEE754-2008 fused multiply-add               Yes\n",
      "    Support is emulated in software               No\n",
      "  Address bits                                    64, Little-Endian\n",
      "  Global memory size                              12636061696 (11.77GiB)\n",
      "  Error Correction support                        No\n",
      "  Max memory allocation                           3159015424 (2.942GiB)\n",
      "  Unified memory for Host and Device              No\n",
      "  Integrated memory (NV)                          No\n",
      "  Shared Virtual Memory (SVM) capabilities        (core)\n",
      "    Coarse-grained buffer sharing                 Yes\n",
      "    Fine-grained buffer sharing                   No\n",
      "    Fine-grained system sharing                   No\n",
      "    Atomics                                       No\n",
      "  Minimum alignment for any data type             128 bytes\n",
      "  Alignment of base address                       4096 bits (512 bytes)\n",
      "  Preferred alignment for atomics                 \n",
      "    SVM                                           0 bytes\n",
      "    Global                                        0 bytes\n",
      "    Local                                         0 bytes\n",
      "  Atomic memory capabilities                      relaxed, work-group scope\n",
      "  Atomic fence capabilities                       relaxed, acquire/release, work-group scope\n",
      "  Max size for global variable                    0\n",
      "  Preferred total size of global vars             0\n",
      "  Global Memory cache type                        Read/Write\n",
      "  Global Memory cache size                        802816 (784KiB)\n",
      "  Global Memory cache line size                   128 bytes\n",
      "  Image support                                   Yes\n",
      "    Max number of samplers per kernel             32\n",
      "    Max size for 1D images from buffer            268435456 pixels\n",
      "    Max 1D or 2D image array size                 2048 images\n",
      "    Max 2D image size                             32768x32768 pixels\n",
      "    Max 3D image size                             16384x16384x16384 pixels\n",
      "    Max number of read image args                 256\n",
      "    Max number of write image args                32\n",
      "    Max number of read/write image args           0\n",
      "  Pipe support                                    No\n",
      "  Max number of pipe args                         0\n",
      "  Max active pipe reservations                    0\n",
      "  Max pipe packet size                            0\n",
      "  Local memory type                               Local\n",
      "  Local memory size                               49152 (48KiB)\n",
      "  Registers per block (NV)                        65536\n",
      "  Max number of constant args                     9\n",
      "  Max constant buffer size                        65536 (64KiB)\n",
      "  Generic address space support                   No\n",
      "  Max size of kernel argument                     4352 (4.25KiB)\n",
      "  Queue properties (on host)                      \n",
      "    Out-of-order execution                        Yes\n",
      "    Profiling                                     Yes\n",
      "  Device enqueue capabilities                     (n/a)\n",
      "  Queue properties (on device)                    \n",
      "    Out-of-order execution                        No\n",
      "    Profiling                                     No\n",
      "    Preferred size                                0\n",
      "    Max size                                      0\n",
      "  Max queues on device                            0\n",
      "  Max events on device                            0\n",
      "  Prefer user sync for interop                    No\n",
      "  Profiling timer resolution                      1000ns\n",
      "  Execution capabilities                          \n",
      "    Run OpenCL kernels                            Yes\n",
      "    Run native kernels                            No\n",
      "    Non-uniform work-groups                       No\n",
      "    Work-group collective functions               No\n",
      "    Sub-group independent forward progress        No\n",
      "    Kernel execution timeout (NV)                 Yes\n",
      "  Concurrent copy and kernel execution (NV)       Yes\n",
      "    Number of async copy engines                  2\n",
      "    IL version                                    (n/a)\n",
      "    ILs with version                              <printDeviceInfo:186: get CL_DEVICE_ILS_WITH_VERSION : error -30>\n",
      "  printf() buffer size                            1048576 (1024KiB)\n",
      "  Built-in kernels                                (n/a)\n",
      "  Built-in kernels with version                   <printDeviceInfo:190: get CL_DEVICE_BUILT_IN_KERNELS_WITH_VERSION : error -30>\n",
      "  Device Extensions                               cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_fp64 cl_khr_3d_image_writes cl_khr_byte_addressable_store cl_khr_icd cl_khr_gl_sharing cl_nv_compiler_options cl_nv_device_attribute_query cl_nv_pragma_unroll cl_nv_copy_opts cl_nv_create_buffer cl_khr_int64_base_atomics cl_khr_int64_extended_atomics cl_khr_device_uuid cl_khr_pci_bus_info\n",
      "  Device Extensions with Version                  cl_khr_global_int32_base_atomics                                 0x400000 (1.0.0)\n",
      "                                                  cl_khr_global_int32_extended_atomics                             0x400000 (1.0.0)\n",
      "                                                  cl_khr_local_int32_base_atomics                                  0x400000 (1.0.0)\n",
      "                                                  cl_khr_local_int32_extended_atomics                              0x400000 (1.0.0)\n",
      "                                                  cl_khr_fp64                                                      0x400000 (1.0.0)\n",
      "                                                  cl_khr_3d_image_writes                                           0x400000 (1.0.0)\n",
      "                                                  cl_khr_byte_addressable_store                                    0x400000 (1.0.0)\n",
      "                                                  cl_khr_icd                                                       0x400000 (1.0.0)\n",
      "                                                  cl_khr_gl_sharing                                                0x400000 (1.0.0)\n",
      "                                                  cl_nv_compiler_options                                           0x400000 (1.0.0)\n",
      "                                                  cl_nv_device_attribute_query                                     0x400000 (1.0.0)\n",
      "                                                  cl_nv_pragma_unroll                                              0x400000 (1.0.0)\n",
      "                                                  cl_nv_copy_opts                                                  0x400000 (1.0.0)\n",
      "                                                  cl_nv_create_buffer                                              0x400000 (1.0.0)\n",
      "                                                  cl_khr_int64_base_atomics                                        0x400000 (1.0.0)\n",
      "                                                  cl_khr_int64_extended_atomics                                    0x400000 (1.0.0)\n",
      "                                                  cl_khr_device_uuid                                               0x400000 (1.0.0)\n",
      "                                                  cl_khr_pci_bus_info                                              0x400000 (1.0.0)\n",
      "\n",
      "  Platform Name                                   Portable Computing Language\n",
      "Number of devices                                 1\n",
      "  Device Name                                     pthread-AMD Ryzen Threadripper 2950X 16-Core Processor\n",
      "  Device Vendor                                   AuthenticAMD\n",
      "  Device Vendor ID                                0x1022\n",
      "  Device Version                                  OpenCL 1.2 pocl HSTR: pthread-x86_64-unknown-linux-gnu-znver1\n",
      "  Driver Version                                  1.8\n",
      "  Device OpenCL C Version                         OpenCL C 1.2 pocl\n",
      "  Device Type                                     CPU\n",
      "  Device Profile                                  FULL_PROFILE\n",
      "  Device Available                                Yes\n",
      "  Compiler Available                              Yes\n",
      "  Linker Available                                Yes\n",
      "  Max compute units                               32\n",
      "  Max clock frequency                             3500MHz\n",
      "  Device Partition                                (core)\n",
      "    Max number of sub-devices                     32\n",
      "    Supported partition types                     equally, by counts\n",
      "    Supported affinity domains                    (n/a)\n",
      "  Max work item dimensions                        3\n",
      "  Max work item sizes                             4096x4096x4096\n",
      "  Max work group size                             4096\n",
      "  Preferred work group size multiple (kernel)     8\n",
      "  Preferred / native vector sizes                 \n",
      "    char                                                16 / 16      \n",
      "    short                                               16 / 16      \n",
      "    int                                                  8 / 8       \n",
      "    long                                                 4 / 4       \n",
      "    half                                                 0 / 0        (n/a)\n",
      "    float                                                8 / 8       \n",
      "    double                                               4 / 4        (cl_khr_fp64)\n",
      "  Half-precision Floating-point support           (n/a)\n",
      "  Single-precision Floating-point support         (core)\n",
      "    Denormals                                     Yes\n",
      "    Infinity and NANs                             Yes\n",
      "    Round to nearest                              Yes\n",
      "    Round to zero                                 Yes\n",
      "    Round to infinity                             Yes\n",
      "    IEEE754-2008 fused multiply-add               Yes\n",
      "    Support is emulated in software               No\n",
      "    Correctly-rounded divide and sqrt operations  Yes\n",
      "  Double-precision Floating-point support         (cl_khr_fp64)\n",
      "    Denormals                                     Yes\n",
      "    Infinity and NANs                             Yes\n",
      "    Round to nearest                              Yes\n",
      "    Round to zero                                 Yes\n",
      "    Round to infinity                             Yes\n",
      "    IEEE754-2008 fused multiply-add               Yes\n",
      "    Support is emulated in software               No\n",
      "  Address bits                                    64, Little-Endian\n",
      "  Global memory size                              132854779904 (123.7GiB)\n",
      "  Error Correction support                        No\n",
      "  Max memory allocation                           34359738368 (32GiB)\n",
      "  Unified memory for Host and Device              Yes\n",
      "  Minimum alignment for any data type             128 bytes\n",
      "  Alignment of base address                       1024 bits (128 bytes)\n",
      "  Global Memory cache type                        Read/Write\n",
      "  Global Memory cache size                        8388608 (8MiB)\n",
      "  Global Memory cache line size                   64 bytes\n",
      "  Image support                                   Yes\n",
      "    Max number of samplers per kernel             16\n",
      "    Max size for 1D images from buffer            2147483648 pixels\n",
      "    Max 1D or 2D image array size                 2048 images\n",
      "    Max 2D image size                             32768x32768 pixels\n",
      "    Max 3D image size                             2048x2048x2048 pixels\n",
      "    Max number of read image args                 128\n",
      "    Max number of write image args                128\n",
      "  Local memory type                               Global\n",
      "  Local memory size                               524288 (512KiB)\n",
      "  Max number of constant args                     8\n",
      "  Max constant buffer size                        524288 (512KiB)\n",
      "  Max size of kernel argument                     1024\n",
      "  Queue properties                                \n",
      "    Out-of-order execution                        Yes\n",
      "    Profiling                                     Yes\n",
      "  Prefer user sync for interop                    Yes\n",
      "  Profiling timer resolution                      1ns\n",
      "  Execution capabilities                          \n",
      "    Run OpenCL kernels                            Yes\n",
      "    Run native kernels                            Yes\n",
      "  printf() buffer size                            16777216 (16MiB)\n",
      "  Built-in kernels                                (n/a)\n",
      "  Device Extensions                               cl_khr_byte_addressable_store cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_3d_image_writes cl_khr_fp64 cl_khr_int64_base_atomics cl_khr_int64_extended_atomics cl_khr_fp64\n",
      "\n",
      "  Platform Name                                   AMD Accelerated Parallel Processing\n",
      "Number of devices                                 1\n",
      "  Device Name                                     gfx906:sramecc+:xnack-\n",
      "  Device Vendor                                   Advanced Micro Devices, Inc.\n",
      "  Device Vendor ID                                0x1002\n",
      "  Device Version                                  OpenCL 2.0 \n",
      "  Driver Version                                  3361.0 (HSA1.1,LC)\n",
      "  Device OpenCL C Version                         OpenCL C 2.0 \n",
      "  Device Type                                     GPU\n",
      "  Device Board Name (AMD)                         AMD Radeon VII\n",
      "  Device PCI-e ID (AMD)                           0x66af\n",
      "  Device Topology (AMD)                           PCI-E, 0000:0a:00.0\n",
      "  Device Profile                                  FULL_PROFILE\n",
      "  Device Available                                Yes\n",
      "  Compiler Available                              Yes\n",
      "  Linker Available                                Yes\n",
      "  Max compute units                               60\n",
      "  SIMD per compute unit (AMD)                     4\n",
      "  SIMD width (AMD)                                16\n",
      "  SIMD instruction width (AMD)                    1\n",
      "  Max clock frequency                             1801MHz\n",
      "  Graphics IP (AMD)                               9.0\n",
      "  Device Partition                                (core)\n",
      "    Max number of sub-devices                     60\n",
      "    Supported partition types                     None\n",
      "    Supported affinity domains                    (n/a)\n",
      "  Max work item dimensions                        3\n",
      "  Max work item sizes                             1024x1024x1024\n",
      "  Max work group size                             256\n",
      "  Preferred work group size (AMD)                 256\n",
      "  Max work group size (AMD)                       1024\n",
      "  Preferred work group size multiple (kernel)     64\n",
      "  Wavefront width (AMD)                           64\n",
      "  Preferred / native vector sizes                 \n",
      "    char                                                 4 / 4       \n",
      "    short                                                2 / 2       \n",
      "    int                                                  1 / 1       \n",
      "    long                                                 1 / 1       \n",
      "    half                                                 1 / 1        (cl_khr_fp16)\n",
      "    float                                                1 / 1       \n",
      "    double                                               1 / 1        (cl_khr_fp64)\n",
      "  Half-precision Floating-point support           (cl_khr_fp16)\n",
      "    Denormals                                     No\n",
      "    Infinity and NANs                             No\n",
      "    Round to nearest                              No\n",
      "    Round to zero                                 No\n",
      "    Round to infinity                             No\n",
      "    IEEE754-2008 fused multiply-add               No\n",
      "    Support is emulated in software               No\n",
      "  Single-precision Floating-point support         (core)\n",
      "    Denormals                                     Yes\n",
      "    Infinity and NANs                             Yes\n",
      "    Round to nearest                              Yes\n",
      "    Round to zero                                 Yes\n",
      "    Round to infinity                             Yes\n",
      "    IEEE754-2008 fused multiply-add               Yes\n",
      "    Support is emulated in software               No\n",
      "    Correctly-rounded divide and sqrt operations  Yes\n",
      "  Double-precision Floating-point support         (cl_khr_fp64)\n",
      "    Denormals                                     Yes\n",
      "    Infinity and NANs                             Yes\n",
      "    Round to nearest                              Yes\n",
      "    Round to zero                                 Yes\n",
      "    Round to infinity                             Yes\n",
      "    IEEE754-2008 fused multiply-add               Yes\n",
      "    Support is emulated in software               No\n",
      "  Address bits                                    64, Little-Endian\n",
      "  Global memory size                              17163091968 (15.98GiB)\n",
      "  Global free memory (AMD)                        16760832 (15.98GiB) 16760832 (15.98GiB)\n",
      "  Global memory channels (AMD)                    128\n",
      "  Global memory banks per channel (AMD)           4\n",
      "  Global memory bank width (AMD)                  256 bytes\n",
      "  Error Correction support                        No\n",
      "  Max memory allocation                           14588628168 (13.59GiB)\n",
      "  Unified memory for Host and Device              No\n",
      "  Shared Virtual Memory (SVM) capabilities        (core)\n",
      "    Coarse-grained buffer sharing                 Yes\n",
      "    Fine-grained buffer sharing                   Yes\n",
      "    Fine-grained system sharing                   No\n",
      "    Atomics                                       No\n",
      "  Minimum alignment for any data type             128 bytes\n",
      "  Alignment of base address                       1024 bits (128 bytes)\n",
      "  Preferred alignment for atomics                 \n",
      "    SVM                                           0 bytes\n",
      "    Global                                        0 bytes\n",
      "    Local                                         0 bytes\n",
      "  Max size for global variable                    14588628168 (13.59GiB)\n",
      "  Preferred total size of global vars             17163091968 (15.98GiB)\n",
      "  Global Memory cache type                        Read/Write\n",
      "  Global Memory cache size                        16384 (16KiB)\n",
      "  Global Memory cache line size                   64 bytes\n",
      "  Image support                                   Yes\n",
      "    Max number of samplers per kernel             26287\n",
      "    Max size for 1D images from buffer            134217728 pixels\n",
      "    Max 1D or 2D image array size                 8192 images\n",
      "    Base address alignment for 2D image buffers   256 bytes\n",
      "    Pitch alignment for 2D image buffers          256 pixels\n",
      "    Max 2D image size                             16384x16384 pixels\n",
      "    Max 3D image size                             16384x16384x8192 pixels\n",
      "    Max number of read image args                 128\n",
      "    Max number of write image args                8\n",
      "    Max number of read/write image args           64\n",
      "  Max number of pipe args                         16\n",
      "  Max active pipe reservations                    16\n",
      "  Max pipe packet size                            1703726280 (1.587GiB)\n",
      "  Local memory type                               Local\n",
      "  Local memory size                               65536 (64KiB)\n",
      "  Local memory size per CU (AMD)                  65536 (64KiB)\n",
      "  Local memory banks (AMD)                        32\n",
      "  Max number of constant args                     8\n",
      "  Max constant buffer size                        14588628168 (13.59GiB)\n",
      "  Preferred constant buffer size (AMD)            16384 (16KiB)\n",
      "  Max size of kernel argument                     1024\n",
      "  Queue properties (on host)                      \n",
      "    Out-of-order execution                        No\n",
      "    Profiling                                     Yes\n",
      "  Queue properties (on device)                    \n",
      "    Out-of-order execution                        Yes\n",
      "    Profiling                                     Yes\n",
      "    Preferred size                                262144 (256KiB)\n",
      "    Max size                                      8388608 (8MiB)\n",
      "  Max queues on device                            1\n",
      "  Max events on device                            1024\n",
      "  Prefer user sync for interop                    Yes\n",
      "  Number of P2P devices (AMD)                     0\n",
      "  Profiling timer resolution                      1ns\n",
      "  Profiling timer offset since Epoch (AMD)        0ns (Thu Jan  1 08:00:00 1970)\n",
      "  Execution capabilities                          \n",
      "    Run OpenCL kernels                            Yes\n",
      "    Run native kernels                            No\n",
      "    Thread trace supported (AMD)                  No\n",
      "    Number of async queues (AMD)                  8\n",
      "    Max real-time compute queues (AMD)            8\n",
      "    Max real-time compute units (AMD)             60\n",
      "  printf() buffer size                            4194304 (4MiB)\n",
      "  Built-in kernels                                (n/a)\n",
      "  Device Extensions                               cl_khr_fp64 cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_int64_base_atomics cl_khr_int64_extended_atomics cl_khr_3d_image_writes cl_khr_byte_addressable_store cl_khr_fp16 cl_khr_gl_sharing cl_amd_device_attribute_query cl_amd_media_ops cl_amd_media_ops2 cl_khr_image2d_from_buffer cl_khr_subgroups cl_khr_depth_images cl_amd_copy_buffer_p2p cl_amd_assembly_program \n",
      "\n",
      "\n",
      "NULL platform behavior\n",
      "  clGetPlatformInfo(NULL, CL_PLATFORM_NAME, ...)  No platform\n",
      "  clGetDeviceIDs(NULL, CL_DEVICE_TYPE_ALL, ...)   No platform\n",
      "  clCreateContext(NULL, ...) [default]            No platform\n",
      "  clCreateContext(NULL, ...) [other]              Success [NV]\n",
      "  clCreateContextFromType(NULL, CL_DEVICE_TYPE_DEFAULT)  No platform\n",
      "  clCreateContextFromType(NULL, CL_DEVICE_TYPE_CPU)  No devices found in platform\n",
      "  clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU)  No platform\n",
      "  clCreateContextFromType(NULL, CL_DEVICE_TYPE_ACCELERATOR)  No devices found in platform\n",
      "  clCreateContextFromType(NULL, CL_DEVICE_TYPE_CUSTOM)  Invalid device type for platform\n",
      "  clCreateContextFromType(NULL, CL_DEVICE_TYPE_ALL)  No platform\n",
      "\tNOTE:\tyour OpenCL library only supports OpenCL 2.0,\n",
      "\t\tbut some installed platforms support OpenCL 3.0.\n",
      "\t\tPrograms using 3.0 features may crash\n",
      "\t\tor behave unexpectedly\n"
     ]
    }
   ],
   "source": [
    "!clinfo --human"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "tags": []
   },
   "source": [
    "## Resources"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "<address>\n",
    "Written by Dr. Toby Potter of Pelagos Consulting and Education for the Pawsey Supercomputing Centre<br>\n",
    "Visit us at: <a href=\"https://www.pelagos-consulting.com\">www.pelagos-consulting.com</a><br>\n",
    "</address>"
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 3 (ipykernel)",
   "language": "python",
   "name": "python3"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 3
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython3",
   "version": "3.8.5"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 4
}