diff --git a/CHANGELOG.md b/CHANGELOG.md --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -7,6 +7,33 @@ Policy (PVP)](https://pvp.haskell.org) +## [1.4.0.0] - ? +### Changed + * Support for LLVM-16 to 22. + * Use shuffle instructions for faster warp-level folds and scans. + * Support new compute capabilities. + +### Fixed + * Device memory was deallocated too early or overwritten [accelerate-llvm#111] + * CUDA Context double free [accelerate-llvm#113] + * Execution of stencil boundaries were not properly synchronised + * Scans sometimes gave incorrect results + * Various other small correctness issues + +### Contributors + +Special thanks to those who contributed patches as part of this release: + + * Trevor L. McDonell (@tmcdonell) + * Tom Smeding (@tomsmeding) + * David van Balen (@dpvanbalen) + * Ivo Gabe de Wolff (@ivogabe) + * Michael Swam (@michael-swan) + * Travis Whitaker (@TravisWhitaker) + * Noah Williams (@noahmartinwilliams) + * Robbert van der Helm (@robbert-vdh) + * Jann Müller (@j-mueller) + ## [1.3.0.0] - 2018-08-27 ### Changed * Code generation improvements for stencil operations @@ -88,4 +115,5 @@ [#436]: https://github.com/AccelerateHS/accelerate/issues/436 [accelerate-llvm#17]: https://github.com/AccelerateHS/accelerate-llvm/issues/17 [accelerate#423]: https://github.com/AccelerateHS/accelerate/issues/423 - +[accelerate-llvm#111]: https://github.com/AccelerateHS/accelerate-llvm/pull/111 +[accelerate-llvm#113]: https://github.com/AccelerateHS/accelerate-llvm/pull/113 diff --git a/README.md b/README.md --- a/README.md +++ b/README.md @@ -3,21 +3,16 @@ # LLVM backends for the Accelerate array language -[![GitHub CI](https://github.com/tmcdonell/accelerate-llvm/workflows/CI/badge.svg)](https://github.com/tmcdonell/accelerate-llvm/actions) +[![CI](https://github.com/AccelerateHS/accelerate-llvm/actions/workflows/ci.yml/badge.svg)](https://github.com/AccelerateHS/accelerate-llvm/actions/workflows/ci.yml) [![Gitter](https://img.shields.io/gitter/room/nwjs/nw.js.svg)](https://gitter.im/AccelerateHS/Lobby) -
-[![Stackage LTS](https://stackage.org/package/accelerate-llvm/badge/lts)](https://stackage.org/lts/package/accelerate-llvm) -[![Stackage Nightly](https://stackage.org/package/accelerate-llvm/badge/nightly)](https://stackage.org/nightly/package/accelerate-llvm) [![Hackage](https://img.shields.io/hackage/v/accelerate-llvm.svg)](https://hackage.haskell.org/package/accelerate-llvm) -
-[![Docker Automated build](https://img.shields.io/docker/automated/tmcdonell/accelerate-llvm.svg)](https://hub.docker.com/r/tmcdonell/accelerate-llvm/) -[![Docker status](https://images.microbadger.com/badges/image/tmcdonell/accelerate-llvm.svg)](https://microbadger.com/images/tmcdonell/accelerate-llvm) This package compiles Accelerate code to LLVM IR, and executes that code on -multicore CPUs as well as NVIDIA GPUs. This avoids the need to go through `nvcc` -or `clang`. For details on Accelerate, refer to the [main repository][GitHub]. +multicore CPUs as well as NVIDIA GPUs. This avoids the need to go through +`nvcc` or write C++ code. For details on Accelerate, refer to the [main +repository][GitHub]. We love all kinds of contributions, so feel free to open issues for missing features as well as report (or fix!) bugs on the [issue tracker][Issues]. @@ -27,179 +22,143 @@ * [Dependencies](#dependencies) - * [Docker](#docker) - * [Installing LLVM](#installing-llvm) - * [Homebrew](#homebrew) + * [macOS](#macos) * [Debian/Ubuntu](#debianubuntu) - * [Building from source](#building-from-source) - * [Installing Accelerate-LLVM](#installing-accelerate-llvm) - * [libNVVM](#libNVVM) + * [Arch Linux](#archlinux) + * [Windows](#windows) Dependencies ------------ -Haskell dependencies are available from Hackage, but there are several external +Haskell dependencies are available from Hackage, but there are some external library dependencies that you will need to install as well: - * [`LLVM`](http://llvm.org) - * [`libFFI`](http://sourceware.org/libffi/) (if using the `accelerate-llvm-native` backend for multicore CPUs) - * [`CUDA`](https://developer.nvidia.com/cuda-downloads) (if using the `accelerate-llvm-ptx` backend for NVIDIA GPUs) - - -Docker ------- - -A [docker](https://www.docker.com) container is provided with this package -preinstalled (via stack) at `/opt/accelerate-llvm`. Note that if you wish to use -the `accelerate-llvm-ptx` GPU backend, you will need to install the [NVIDIA -docker](https://github.com/NVIDIA/nvidia-docker) plugin; see that page for more -information. - -```sh -$ docker run -it tmcdonell/accelerate-llvm -``` - - -Installing LLVM ---------------- - -When installing LLVM, make sure that it includes the `libLLVM` shared library. -If you want to use the GPU targeting `accelerate-llvm-ptx` backend, make sure -you install (or build) LLVM with the 'nvptx' target. +- if using `accelerate-llvm-native` for multicore CPU: + [`libFFI`](http://sourceware.org/libffi/) +- if using `accelerate-llvm-ptx` for GPU: + [`CUDA`](https://developer.nvidia.com/cuda-downloads); + [Note that not all versions of CUDA support all NVIDIA GPUs](https://en.wikipedia.org/wiki/CUDA#GPUs_supported) +- [`clang`](https://clang.llvm.org/) (if using `accelerate-llvm-ptx`: version + 16 or higher, built with support for the `nvptx` backend). `accelerate-llvm` + uses the command-line tool as a way to be compatible with many different LLVM + versions, not to compile C code. (Accelerate passes LLVM IR to `clang`.) -## Homebrew +Below are some OS-specific instructions. If anything here is wrong or out of +date, please file an issue. -Example using [Homebrew](http://brew.sh) on macOS: +## macOS -```sh -$ brew install llvm-hs/llvm/llvm-9 -``` +To get `libFFI`, run `brew install libffi`. `clang` is already provided with +macOS (you may need to `xcode-select --install`), and CUDA is not supported on +macOS. ## Debian/Ubuntu -For Debian/Ubuntu based Linux distributions, the LLVM.org website provides -binary distribution packages. Check [apt.llvm.org](http://apt.llvm.org) for -instructions for adding the correct package database for your OS version, and -then: - -```sh -$ apt-get install llvm-9-dev -``` - -## Building from source - -If your OS does not have an appropriate LLVM distribution available, you can also build from source. Detailed build instructions are available on the [LLVM.org website](http://releases.llvm.org/6.0.0/docs/CMake.html). Note that you will require at least [CMake 3.4.3](http://www.cmake.org/cmake/resources/software.html) and a recent C++ compiler; at least Clang 3.1, GCC 4.8, or Visual Studio 2015 (update 3). - - 1. Download and unpack the [LLVM-9 source code](https://github.com/llvm/llvm-project/releases/download/llvmorg-9.0.1/llvm-9.0.1.src.tar.xz). We'll refer to - the path that the source tree was unpacked to as `LLVM_SRC`. Only the main - LLVM source tree is required, but you can optionally add other components - such as the Clang compiler or Polly loop optimiser. See the [LLVM releases](http://releases.llvm.org/download.html#9.0.1) - page for the complete list. - - 2. Create a temporary build directory and `cd` into it, for example: - ```sh - $ mkdir /tmp/build - $ cd /tmp/build - ``` - - 3. Execute the following to configure the build. Here `INSTALL_PREFIX` is - where LLVM is to be installed, for example `/usr/local` or - `$HOME/opt/llvm`: - ```sh - $ cmake $LLVM_SRC -DCMAKE_INSTALL_PREFIX=$INSTALL_PREFIX -DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_ASSERTIONS=ON -DLLVM_BUILD_LLVM_DYLIB=ON -DLLVM_LINK_LLVM_DYLIB=ON - ``` - See [options and variables](http://llvm.org/docs/CMake.html#options-and-variables) - for a list of additional build parameters you can specify. - - 4. Build and install: - ```sh - $ cmake --build . - $ cmake --build . --target install - ``` - - 5. For macOS only, some additional steps are useful to work around issues related - to [System Integrity Protection](https://en.wikipedia.org/wiki/System_Integrity_Protection): - ```sh - cd $INSTALL_PREFIX/lib - ln -s libLLVM.dylib libLLVM-9.dylib - install_name_tool -id $PWD/libLTO.dylib libLTO.dylib - install_name_tool -id $PWD/libLLVM.dylib libLLVM.dylib - install_name_tool -change '@rpath/libLLVM.dylib' $PWD/libLLVM.dylib libLTO.dylib - ``` - +For `clang`: +- On Ubuntu 24.04 (noble) / Debian trixie or higher: `sudo apt install clang`. +- Otherwise, if you need only the CPU backend (`accelerate-llvm-native`): + `sudo apt install clang` will give you an old version of `clang`, but the CPU + backend is likely to work fine. +- If you are on an older distro and need the GPU backend + (`accelerate-llvm-ptx`): `clang` version 16 or higher is required. + Add the apt source from [apt.llvm.org](https://apt.llvm.org/). The neatest + way to do this is to create a file `/etc/apt/sources.list.d/llvm.list` (the + precise file name does not matter) and put in it, for Ubuntu (change "jammy" + as appropriate): -Installing Accelerate-LLVM --------------------------- + deb http://apt.llvm.org/jammy/ llvm-toolchain-jammy main + deb-src http://apt.llvm.org/jammy/ llvm-toolchain-jammy main -Once the dependencies are installed, we are ready to install `accelerate-llvm`. + or for Debian (change "bookworm" as appropriate): -For example, installation using [`stack`](http://docs.haskellstack.org/en/stable/README.html) -just requires you to point it to the appropriate configuration file: -```sh -$ ln -s stack-8.8.yaml stack.yaml -$ stack setup -$ stack install -``` + deb http://apt.llvm.org/bookworm/ llvm-toolchain-bookworm main + deb-src http://apt.llvm.org/bookworm/ llvm-toolchain-bookworm main -Note that the version of [`llvm-hs`](https://hackage.haskell.org/package/llvm-hs) -used must match the installed version of LLVM, which is currently 9.0. + and `sudo apt update; sudo apt install clang`. This gets you the latest + version of `clang`; different sources are also available for specific + versions (see [apt.llvm.org](https://apt.llvm.org)). +To use the CPU backend (`accelerate-llvm-native`), install `libFFI` using +`sudo apt install libffi-dev`. -## libNVVM +To use the GPU backend (`accelerate-llvm-ptx`), install CUDA from +[here](https://developer.nvidia.com/cuda-downloads?target_os=Linux) +("deb (network)" is smoother than the "deb (local)" option). -The `accelerate-llvm-ptx` backend can optionally be compiled to generate GPU -code using the `libNVVM` library, rather than LLVM's inbuilt NVPTX code -generator. `libNVVM` is a closed-source library distributed as part of the -NVIDIA CUDA toolkit, and is what the `nvcc` compiler itself uses internally when -compiling CUDA C code. +## Arch Linux -Using `libNVVM` _may_ improve GPU performance compared to the code generator -built in to LLVM. One difficulty with using it however is that since `libNVVM` -is also based on LLVM, and typically lags LLVM by several releases, you must -install `accelerate-llvm` with a "compatible" version of LLVM, which will depend -on the version of the CUDA toolkit you have installed. The following table shows -combinations which have been tested: +Run `sudo pacman -S clang`. To use the CPU backend (`accelerate-llvm-native`), +additionally run `sudo pacman -S libffi`. To use the GPU backend +(`accelerate-llvm-ptx`), additionally run `sudo pacman -S cuda`. -| | LLVM-3.3 | LLVM-3.4 | LLVM-3.5 | LLVM-3.8 | LLVM-3.9 | LLVM-4.0 | LLVM-5.0 | LLVM-6.0 | LLVM-7 | LLVM-8 | LLVM-9 | -| ------------- | :------: | :------: | :------: | :------: | :------: | :------: | :------: | :------: | :----: | :----: | :----: | -| **CUDA-7.0** | ⭕ | ❌ | | | | | | | | | | -| **CUDA-7.5** | | ⭕ | ⭕ | ❌ | | | | | | | | -| **CUDA-8.0** | | | ⭕ | ⭕ | ❌ | ❌ | | | | | | -| **CUDA-9.0** | | | | | | ❌ | ❌ | | | | | -| **CUDA-9.1** | | | | | | | | | | | | -| **CUDA-9.2** | | | | | | | | | | | | -| **CUDA-10.0** | | | | | | | | | | | | -| **CUDA-10.1** | | | | | | | | | | | | +## Windows -Where ⭕ = Works, and ❌ = Does not work. +We recommend WSL2 (not WSL1, WSL2!) and following the Ubuntu instructions +above. The remainder of this text attemps to give you a working system on +Windows native. -The above table is incomplete! If you try a particular combination and find that -it does or does not work, please let us know! +Install `clang`; you have two options: +1. Using + [WinGet](https://learn.microsoft.com/en-us/windows/package-manager/winget/): + `winget install LLVM.LLVM` +2. By downloading the installer directly (WinGet just runs the same installer) + from [here](https://github.com/llvm/llvm-project/releases) (choose + "LLVM-<version>-win64.exe" from the latest release; you may need to click + "Show all 57 assets"). +This will also give you `libFFI`. -Note that the above restrictions on CUDA and LLVM version exist _only_ if you -want to use the NVVM component. Otherwise, you should be free to use any -combination of CUDA and LLVM. +
Optionally, add clang (and more) to your system path. Click to see how. -Also note that `accelerate-llvm-ptx` itself currently requires at least LLVM-4.0. +Accelerate should be able to find `clang` automatically even if you do not do +this. However, for easy access to `clang` and all other LLVM executables, add +`C:\Program Files\LLVM\bin` to the system path as follows: +1. Search for "environment variables" in the start menu +2. Click "Edit the system environment variables" +3. Click on "Environment Variables..." +4. Double-click on the user variable called "Path" +5. And add a new entry containing `C:\Program Files\LLVM\bin`. -Using `stack`, either edit the `stack.yaml` and add the following section: +Note that if you add an entry here manually, it is a good idea to clean it up +again if you uninstall LLVM/clang. (Leaving it there is not very harmful, +however.) -```yaml -flags: - accelerate-llvm-ptx: - nvvm: true -``` +You may find that the LLVM/clang installer has already added the Path entry +automatically (it did not for us); if so, no need to add a second entry. -Or install using the following option on the command line: +—— +
-```sh -$ stack install accelerate-llvm-ptx --flag accelerate-llvm-ptx:nvvm -``` +You may additionally need the VS Build Tools, if you have not yet installed and +set up Visual Studio otherwise. You need this if `clang` complains that it is +`unable to find a Visual Studio installation; try running Clang from a developer command prompt`. -If installing via `cabal`: +1. If you already have the Visual Studio Installer on your system, open it and + check if you already have Visual Studio (Community) installed. Note that + this is completely unrelated to VS _Code_. + - If you already have VS (Community): inside the Visual Studio Installer, + click on "Modify" in the VS (Community) box. This should get you a screen + with "workloads" you can select. + - If you do not yet have VS (Community), install the VS Build Tools: go to + https://visualstudio.microsoft.com/downloads, scroll down to "All + Downloads", open "Tools for Visual Studio", and select "Build Tools for + Visual Studio". If you run the installer, you should get a screen with + "workloads" you can select. +2. Under the Workloads tab, choose the "Desktop development with C++" workload. + If you want to save a bit of disk space (not much), keep only the following + two options selected: + - "MSVC v143 - VE 2022 C++ x64/x86 build tools (Latest)" + - "Windows 11 SDK (…)" (choose the latest option). The attentive reader may + note that the wizard also offers Clang; we recommend a separate Clang + install for Accelerate because the one from VS somehow doesn’t seem to + work properly with Accelerate. If you find out why, please let us know. +3. Install that. This takes a while. -```sh -$ cabal install accelerate-llvm-ptx -fnvvm -``` +It turns out that having both Visual Studio and the Build Tools installed +results in Clang getting confused between the two (it appears that Visual +Studio is 64-bit (x64) and the Build Tools are 32-bit (x86)). If Clang +complains about the bit-ness of your system libraries, double-check that you +haven’t installed both simultaneously. +The GPU backend (`accelerate-llvm-ptx`) probably doesn't work on Windows; in +any case, it is untested. diff --git a/Setup.hs b/Setup.hs deleted file mode 100644 --- a/Setup.hs +++ /dev/null @@ -1,2 +0,0 @@ -import Distribution.Simple -main = defaultMain diff --git a/accelerate-llvm-ptx.cabal b/accelerate-llvm-ptx.cabal --- a/accelerate-llvm-ptx.cabal +++ b/accelerate-llvm-ptx.cabal @@ -1,72 +1,39 @@ +cabal-version: 2.2 + name: accelerate-llvm-ptx -version: 1.3.0.0 -cabal-version: >= 1.10 -tested-with: GHC >= 8.6 +version: 1.4.0.0 +tested-with: GHC >= 9.4 build-type: Simple synopsis: Accelerate backend for NVIDIA GPUs description: This library implements a backend for the /Accelerate/ language which - generates LLVM-IR targeting CUDA capable GPUs. For further information, + generates LLVM IR targeting CUDA capable GPUs. For further information, refer to the main package. . [/Dependencies/] . Haskell dependencies are available from Hackage. The following external - libraries are alse required: - . - * - . - * - . - [/Installing LLVM/] - . - /Homebrew/ - . - Example using Homebrew on macOS: - . - > brew install llvm-hs/llvm/llvm-9 - . - /Debian & Ubuntu/ - . - For Debian/Ubuntu based Linux distributions, the LLVM.org website provides - binary distribution packages. Check for - instructions for adding the correct package database for your OS version, - and then: - . - > apt-get install llvm-9-dev - . - /Building from source/ + dependencies are also required: . - If your OS does not have an appropriate LLVM distribution available, you can - also build from source. Detailed build instructions are available on - . Make sure to - include the cmake build options - @-DLLVM_BUILD_LLVM_DYLIB=ON -DLLVM_LINK_LLVM_DYLIB=ON@ so that the @libLLVM@ - shared library will be built. Also ensure that the @LLVM_TARGETS_TO_BUILD@ - option includes the @NVPTX@ target (if not specified all targets are built). + * (not used to compile C code, but to compile generated LLVM IR via a mostly LLVM-version-independent interface) + * . + For installation instructions, see the . -license: BSD3 +license: BSD-3-Clause license-file: LICENSE author: Trevor L. McDonell maintainer: Trevor L. McDonell bug-reports: https://github.com/AccelerateHS/accelerate/issues category: Accelerate, Compilers/Interpreters, Concurrency, Data, Parallelism -extra-source-files: +extra-doc-files: CHANGELOG.md README.md --- Configuration flags --- ------------------- - -Flag nvvm - Default: False - Description: Use the NVVM library to generate optimised PTX - -- Build configuration -- ------------------- @@ -111,9 +78,7 @@ Data.Array.Accelerate.LLVM.PTX.Compile Data.Array.Accelerate.LLVM.PTX.Compile.Cache - Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.Load - Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.TH Data.Array.Accelerate.LLVM.PTX.Link Data.Array.Accelerate.LLVM.PTX.Link.Cache @@ -126,26 +91,32 @@ Paths_accelerate_llvm_ptx + autogen-modules: + Paths_accelerate_llvm_ptx + build-depends: base >= 4.10 && < 5 - , accelerate == 1.3.* - , accelerate-llvm == 1.3.* + , accelerate == 1.4.* + , accelerate-llvm == 1.4.* , bytestring >= 0.10.4 - , containers >= 0.5 && < 0.7 + , containers >= 0.5 && < 0.9 , cuda >= 0.10 , deepseq >= 1.3 , directory >= 1.0 , dlist >= 0.6 , file-embed >= 0.0.8 , filepath >= 1.0 + , formatting >= 7.0 , hashable >= 1.2 - , llvm-hs >= 4.1 && < 9.1 - , llvm-hs-pure >= 4.1 && < 9.1 + -- , llvm-pretty >= 0.12 , mtl >= 2.2.1 - , nvvm >= 0.9 - , pretty >= 1.1 + -- only used to render llvm-pretty output + , pretty + , prettyprinter >= 1.7 + , primitive >= 0.7 , process >= 1.4.3 , template-haskell + , text >= 1.2 , unordered-containers >= 0.2 if impl(ghc >= 8.6) @@ -156,9 +127,6 @@ ghc , ghci - cpp-options: - -DUNBOXED_TUPLES=1 - hs-source-dirs: src @@ -170,28 +138,20 @@ -Wall -fwarn-tabs - ghc-prof-options: - -caf-all - -auto-all - if impl(ghc >= 8.0) - ghc-options: - -Wmissed-specialisations - - if flag(nvvm) - cpp-options: - -DACCELERATE_USE_NVVM - - test-suite nofib-llvm-ptx type: exitcode-stdio-1.0 hs-source-dirs: test/nofib main-is: Main.hs + other-modules: + Data.Array.Accelerate.LLVM.PTX.NoFib.RunQ build-depends: base >= 4.10 , accelerate , accelerate-llvm-ptx + , tasty + , tasty-hunit default-language: Haskell2010 @@ -212,7 +172,7 @@ source-repository this type: git - tag: v1.3.0.0 + tag: v1.4.0.0 location: https://github.com/AccelerateHS/accelerate-llvm.git -- vim: nospell diff --git a/src/Data/Array/Accelerate/LLVM/PTX.hs b/src/Data/Array/Accelerate/LLVM/PTX.hs --- a/src/Data/Array/Accelerate/LLVM/PTX.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX.hs @@ -1,13 +1,14 @@ {-# LANGUAGE AllowAmbiguousTypes #-} {-# LANGUAGE BangPatterns #-} -{-# LANGUAGE CPP #-} {-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE GADTs #-} +{-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE RankNTypes #-} {-# LANGUAGE ScopedTypeVariables #-} {-# LANGUAGE TemplateHaskell #-} {-# LANGUAGE TypeApplications #-} {-# LANGUAGE TypeFamilies #-} +{-# LANGUAGE TypeOperators #-} {-# LANGUAGE TypeSynonymInstances #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX @@ -22,6 +23,37 @@ -- NVPTX for execution on NVIDIA GPUs. Expressions are on-line translated into -- LLVM code, which is just-in-time executed in parallel on the GPU. -- +-- * Concurrent kernel execution +-- +-- Accelerate enables copies to and from device memory to occur concurrently +-- with kernel execution. However, in order for kernels to be executed +-- concurrently, you must create your own 'PTX' execution target and pass it to +-- the 'runWith'-style functions. For example: +-- +-- > import Data.Array.Accelerate as A +-- > import qualified Data.Array.Accelerate.LLVM.PTX as PTX +-- > import qualified Foreign.CUDA.Driver as C +-- > +-- > dotp :: Acc (Vector Float) -> Acc (Vector Float) -> Acc (Scalar Float) +-- > dotp xs ys = A.fold (+) 0 (A.zipWith (*) xs ys) +-- > +-- > runDotP = do +-- > C.initialise [] +-- > -- Assuming just one CUDA device +-- > dev0 <- C.device 0 +-- > dev0Props <- C.props dev0 +-- > ptx <- PTX.createTargetForDevice dev0 dev0Props [] +-- > -- Using runAsyncWith +-- > let xs = fromList (Z:.10) [0..] :: Vector Float +-- > ys = fromList (Z:.10) [1,3..] :: Vector Float +-- > asyncDotP <- PTX.runAsyncWith ptx (dotp (use +-- > ... +-- > dotP <- PTX.wait asyncDotP +-- +-- The CUDA scheduler will only execute kernels concurrently if there are +-- sufficient resources available. Concurrency is limited by the number of +-- available threads, blocks, warps and registers; a kernel occupancy of \(n\) +-- does not always imply that \(1-n\) device resources are available. module Data.Array.Accelerate.LLVM.PTX ( @@ -58,7 +90,6 @@ import Data.Array.Accelerate.AST.LeftHandSide ( lhsToTupR ) import Data.Array.Accelerate.Array.Data import Data.Array.Accelerate.Async ( Async, asyncBound, wait, poll, cancel ) -import Data.Array.Accelerate.Debug import Data.Array.Accelerate.Error import Data.Array.Accelerate.Representation.Array ( liftArraysR ) import Data.Array.Accelerate.Smart ( Acc ) @@ -71,6 +102,7 @@ import Data.Array.Accelerate.LLVM.PTX.Array.Data import Data.Array.Accelerate.LLVM.PTX.Compile import Data.Array.Accelerate.LLVM.PTX.Context +import Data.Array.Accelerate.LLVM.PTX.Debug import Data.Array.Accelerate.LLVM.PTX.Embed import Data.Array.Accelerate.LLVM.PTX.Execute import Data.Array.Accelerate.LLVM.PTX.Execute.Async ( Par, evalPar, getArrays ) @@ -84,10 +116,9 @@ import Control.Exception import Control.Monad.Trans import Data.Maybe +import Formatting ( shown ) import System.IO.Unsafe -import Text.Printf -import qualified Language.Haskell.TH as TH -import qualified Language.Haskell.TH.Syntax as TH +import qualified Data.Array.Accelerate.TH.Compat as TH -- Accelerate: LLVM backend for NVIDIA GPUs @@ -143,9 +174,9 @@ execute = do dumpGraph acc evalPTX target $ do - build <- phase "compile" (compileAcc acc) >>= dumpStats - exec <- phase "link" (linkAcc build) - res <- phase "execute" (evalPar (executeAcc exec >>= copyToHostLazy (Sugar.arraysR @a))) + build <- phase Compile (compileAcc acc) >>= dumpStats + exec <- phase Link (linkAcc build) + res <- phase Execute (evalPar (executeAcc exec >>= copyToHostLazy (Sugar.arraysR @a))) return $ toArr res @@ -241,8 +272,8 @@ !afun = unsafePerformIO $ do dumpGraph acc evalPTX target $ do - build <- phase "compile" (compileAfun acc) >>= dumpStats - link <- phase "link" (linkAfun build) + build <- phase Compile (compileAfun acc) >>= dumpStats + link <- phase Link (linkAfun build) return link go :: forall aenv t r trepr. @@ -255,7 +286,7 @@ a <- useRemoteAsync (lhsToTupR lhs) $ fromArr arrs return (aenv `push` (lhs, a)) in go repr l k' - go AfunctionReprBody (Abody b) k = unsafePerformIO . phase "execute" . evalPTX target . evalPar $ do + go AfunctionReprBody (Abody b) k = unsafePerformIO . phase Execute . evalPTX target . evalPar $ do aenv <- k fut <- executeOpenAcc b aenv toArr <$> copyToHostLazy (Sugar.arraysR @r) fut @@ -300,8 +331,8 @@ !afun = unsafePerformIO $ do dumpGraph acc evalPTX target $ do - build <- phase "compile" (compileAfun acc) >>= dumpStats - link <- phase "link" (linkAfun build) + build <- phase Compile (compileAfun acc) >>= dumpStats + link <- phase Link (linkAfun build) return link !exec = runAsync' target afun (return Empty) @@ -321,7 +352,7 @@ instance Arrays b => RunAsync (IO (Async b)) where type RunAsyncR (IO (Async b)) = ArraysR b runAsync' _ Alam{} _ = error "runAsync: function not fully applied" - runAsync' target (Abody b) k = asyncBound . phase "execute" . evalPTX target . evalPar $ do + runAsync' target (Abody b) k = asyncBound . phase Execute . evalPTX target . evalPar $ do aenv <- k ans <- executeOpenAcc b aenv arrs <- getArrays (arraysR b) ans @@ -464,14 +495,14 @@ in TH.runIO $ do dumpGraph acc evalPTX defaultTarget $ - phase "compile" (compileAfun acc) >>= dumpStats + phase Compile (compileAfun acc) >>= dumpStats let go :: CompiledOpenAfun PTX aenv t -> [TH.PatQ] -> [TH.ExpQ] -> [TH.StmtQ] -> TH.ExpQ go (Alam lhs l) xs as stmts = do x <- TH.newName "x" -- lambda bound variable a <- TH.newName "a" -- local array name - s <- TH.bindS (TH.varP a) [| useRemoteAsync $(TH.unTypeQ $ liftArraysR (lhsToTupR lhs)) (fromArr $(TH.varE x)) |] - go l (TH.bangP (TH.varP x) : xs) ([| ($(TH.unTypeQ $ liftALeftHandSide lhs), $(TH.varE a)) |] : as) (return s : stmts) + s <- TH.bindS (TH.varP a) [| useRemoteAsync $(TH.unTypeCode $ liftArraysR (lhsToTupR lhs)) (fromArr $(TH.varE x)) |] + go l (TH.bangP (TH.varP x) : xs) ([| ($(TH.unTypeCode $ liftALeftHandSide lhs), $(TH.varE a)) |] : as) (return s : stmts) go (Abody b) xs as stmts = do r <- TH.newName "r" -- result @@ -481,10 +512,10 @@ body = embedOpenAcc defaultTarget b -- TH.lamE (reverse xs) - [| $using (phase "execute" $(k ( + [| $using (phase Execute $(k ( TH.doE ( reverse stmts ++ - [ TH.bindS (TH.varP r) [| executeOpenAcc $(TH.unTypeQ body) $aenv |] - , TH.bindS (TH.varP s) [| copyToHostLazy $(TH.unTypeQ (liftArraysR (arraysR b))) $(TH.varE r) |] + [ TH.bindS (TH.varP r) [| executeOpenAcc $(TH.unTypeCode body) $aenv |] + , TH.bindS (TH.varP s) [| copyToHostLazy $(TH.unTypeCode (liftArraysR (arraysR b))) $(TH.varE r) |] , TH.noBindS [| return $ toArr $(TH.varE s) |] ])))) |] @@ -519,7 +550,7 @@ registerForeignPtrAllocator $ \bytes -> withContext (ptxContext target) (CUDA.mallocHostForeignPtr [] bytes) `catch` - \e -> internalError (show (e :: CUDAException)) + \e -> internalError shown (e :: CUDAException) -- Debugging @@ -527,7 +558,4 @@ dumpStats :: MonadIO m => a -> m a dumpStats x = liftIO dumpSimplStats >> return x - -phase :: MonadIO m => String -> m a -> m a -phase n go = timed dump_phases (\wall cpu -> printf "phase %s: %s" n (elapsed wall cpu)) go diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs b/src/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs @@ -19,7 +19,7 @@ ) where import Foreign.CUDA.Analysis as CUDA -import Language.Haskell.TH +import Data.Array.Accelerate.TH.Compat -- | Given information about the resource usage of the compiled kernel, @@ -33,7 +33,7 @@ , Int -- thread block size , Int -> Int -- grid size required to process the given input size , Int -- #bytes dynamic shared memory - , Q (TExp (Int -> Int)) + , CodeQ (Int -> Int) ) -- | Analytics for a simple kernel which requires no additional shared memory or @@ -51,7 +51,7 @@ -> [Int] -- ^ Thread block sizes to consider -> (Int -> Int) -- ^ Shared memory (#bytes) as a function of thread block size -> (Int -> Int -> Int) -- ^ Determine grid size for input size 'n' (first arg) over thread blocks of size 'm' (second arg) - -> Q (TExp (Int -> Int -> Int)) + -> CodeQ (Int -> Int -> Int) -> LaunchConfig launchConfig dev candidates dynamic_smem grid_size grid_sizeQ maxThreads registers static_smem = let @@ -69,6 +69,6 @@ multipleOf :: Int -> Int -> Int multipleOf x y = ((x + y - 1) `quot` y) -multipleOfQ :: Q (TExp (Int -> Int -> Int)) +multipleOfQ :: CodeQ (Int -> Int -> Int) multipleOfQ = [|| multipleOf ||] diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Array/Data.hs b/src/Data/Array/Accelerate/LLVM/PTX/Array/Data.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Array/Data.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Array/Data.hs @@ -1,13 +1,13 @@ {-# LANGUAGE BangPatterns #-} -{-# LANGUAGE CPP #-} {-# LANGUAGE GADTs #-} {-# LANGUAGE MagicHash #-} +{-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE RankNTypes #-} {-# LANGUAGE RecordWildCards #-} {-# LANGUAGE ScopedTypeVariables #-} {-# LANGUAGE TemplateHaskell #-} {-# LANGUAGE UnboxedTuples #-} -{-# OPTIONS_GHC -fno-warn-orphans #-} +{-# OPTIONS_GHC -Wno-orphans #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Array.Data -- Copyright : [2014..2020] The Accelerate Team @@ -43,8 +43,9 @@ import qualified Data.Array.Accelerate.LLVM.PTX.Array.Prim as Prim import Control.Applicative -import Control.Monad.Reader -import Control.Monad.State ( gets ) +import Control.Monad +import Control.Monad.IO.Class ( liftIO ) +import Control.Monad.Reader ( asks ) import System.IO.Unsafe import Prelude @@ -98,7 +99,7 @@ a2 <- copyToHostLazy r2 f2 return (a1, a2) copyToHostLazy (TupRsingle (ArrayR shr tp)) future = do - ptx <- gets llvmTarget + ptx <- asks llvmTarget liftIO $ do Array sh adata <- wait future diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs b/src/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs @@ -1,10 +1,12 @@ {-# LANGUAGE BangPatterns #-} -{-# LANGUAGE CPP #-} {-# LANGUAGE DataKinds #-} {-# LANGUAGE GADTs #-} {-# LANGUAGE MagicHash #-} +{-# LANGUAGE MagicHash #-} +{-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE RecordWildCards #-} {-# LANGUAGE ScopedTypeVariables #-} +{-# LANGUAGE TypeApplications #-} {-# LANGUAGE TypeOperators #-} {-# LANGUAGE UnboxedTuples #-} -- | @@ -34,7 +36,7 @@ import Data.Array.Accelerate.Array.Data import Data.Array.Accelerate.Array.Unique import Data.Array.Accelerate.Error -import Data.Array.Accelerate.Lifetime hiding ( withLifetime ) +import Data.Array.Accelerate.Lifetime hiding ( withLifetime ) import Data.Array.Accelerate.Representation.Elt import Data.Array.Accelerate.Representation.Type import Data.Array.Accelerate.Type @@ -45,20 +47,23 @@ import Data.Array.Accelerate.LLVM.PTX.Execute.Async import Data.Array.Accelerate.LLVM.PTX.Execute.Event import Data.Array.Accelerate.LLVM.PTX.Execute.Stream -import Data.Array.Accelerate.LLVM.PTX.Array.Remote as Remote -import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug +import Data.Array.Accelerate.LLVM.PTX.Array.Remote as Remote +import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug -import qualified Foreign.CUDA.Driver as CUDA -import qualified Foreign.CUDA.Driver.Stream as CUDA +import qualified Foreign.CUDA.Driver as CUDA +import qualified Foreign.CUDA.Driver.Stream as CUDA import Control.Monad import Control.Monad.Reader import Data.IORef -import GHC.Base -import Text.Printf +import Data.Text.Lazy.Builder +import Formatting hiding ( bytes ) +import qualified Formatting as F import Prelude +import GHC.Base ( IO(..), Int(..), Double(..), touch#, int2Double# ) + -- | Allocate a device-side array associated with the given host array. If the -- allocation fails due to a memory error, we attempt some last-ditch memory -- cleanup before trying again. If it still fails; error. @@ -71,7 +76,7 @@ -> ArrayData e -> LLVM PTX () mallocArray !t !n !ad = do - message ("mallocArray: " ++ showBytes (n * bytesElt (TupRsingle (SingleScalarType t)))) + message ("mallocArray: " % formatBytes) (n * bytesElt (TupRsingle (SingleScalarType t))) void $ Remote.malloc t ad n False @@ -88,7 +93,7 @@ -> ArrayData e -> Par PTX (Future (ArrayData e)) useArrayAsync !t !n !ad = do - message ("useArrayAsync: " ++ showBytes (n * bytesElt (TupRsingle (SingleScalarType t)))) + message ("useArrayAsync: " % formatBytes) (n * bytesElt (TupRsingle (SingleScalarType t))) alloc <- liftPar $ Remote.malloc t ad n True if alloc then pokeArrayAsync t n ad @@ -111,14 +116,14 @@ let !src = CUDA.HostPtr (unsafeUniqueArrayPtr ad) !bytes = n * bytesElt (TupRsingle (SingleScalarType t)) -- - stream <- asks ptxStream + stream <- asksParState ptxStream result <- liftPar $ withLifetime stream $ \st -> withDevicePtr t ad $ \dst -> nonblocking stream $ do transfer "pokeArray" bytes (Just st) $ do CUDA.pokeArrayAsync n src dst (Just st) - Debug.didCopyBytesToRemote (fromIntegral bytes) + Debug.memcpy_to_remote bytes return ad -- return result @@ -145,14 +150,14 @@ let !bytes = n * bytesElt (TupRsingle (SingleScalarType t)) !dst = CUDA.HostPtr (unsafeUniqueArrayPtr ad_dst) -- - stream <- asks ptxStream + stream <- asksParState ptxStream result <- liftPar $ withLifetime stream $ \st -> withDevicePtr t ad_src $ \src -> nonblocking stream $ do transfer "indexArray" bytes (Just st) $ do CUDA.peekArrayAsync n (src `CUDA.advanceDevPtr` (i*n)) dst (Just st) - Debug.didCopyBytesFromRemote (fromIntegral bytes) + Debug.memcpy_from_remote bytes return ad_dst -- return result @@ -174,14 +179,14 @@ let !bytes = n * bytesElt (TupRsingle (SingleScalarType t)) !dst = CUDA.HostPtr (unsafeUniqueArrayPtr ad) -- - stream <- asks ptxStream + stream <- asksParState ptxStream result <- liftPar $ withLifetime stream $ \st -> withDevicePtr t ad $ \src -> nonblocking stream $ do transfer "peekArray" bytes (Just st) $ do CUDA.peekArrayAsync n src dst (Just st) - Debug.didCopyBytesFromRemote (fromIntegral bytes) + Debug.memcpy_from_remote bytes return ad -- return result @@ -203,7 +208,7 @@ = do let !bytes = n * bytesElt (TupRsingle (SingleScalarType t)) -- - stream <- asks ptxStream + stream <- asksParState ptxStream result <- liftPar $ withLifetime stream $ \st -> withDevicePtr t ad_src $ \src -> @@ -282,7 +287,7 @@ = do let !bytes = n * bytesElt (TupRsingle (SingleScalarType t)) -- - stream <- asks ptxStream + stream <- asksParState ptxStream result <- liftPar $ withLifetime stream $ \st -> withDevicePtr t ad $ \ptr -> @@ -345,7 +350,7 @@ return (Nothing, future) else do - future <- Future <$> liftIO (newIORef (Pending event Nothing result)) + future <- Future <$> liftIO (newIORef (Pending event (return ()) result)) return (Just event, future) {-# INLINE withLifetime #-} @@ -363,27 +368,24 @@ -- Debug -- ----- -{-# INLINE showBytes #-} -showBytes :: Int -> String -showBytes x = Debug.showFFloatSIBase (Just 0) 1024 (fromIntegral x :: Double) "B" +{-# INLINE double #-} +double :: Int -> Double +double (I# i#) = D# (int2Double# i#) -{-# INLINE trace #-} -trace :: MonadIO m => String -> m a -> m a -trace msg next = liftIO (Debug.traceIO Debug.dump_gc ("gc: " ++ msg)) >> next +{-# INLINE formatBytes #-} +formatBytes :: Format r (Int -> r) +formatBytes = F.bytes @Double shortest {-# INLINE message #-} -message :: MonadIO m => String -> m () -message s = s `trace` return () +message :: MonadIO m => Format (m ()) a -> a +message fmt = Debug.traceM Debug.dump_gc ("gc: " % fmt) {-# INLINE transfer #-} -transfer :: MonadIO m => String -> Int -> Maybe CUDA.Stream -> IO () -> m () +transfer :: MonadIO m => Builder -> Int -> Maybe CUDA.Stream -> IO () -> m () transfer name bytes stream action = - let showRate x t = Debug.showFFloatSIBase (Just 3) 1024 (fromIntegral x / t) "B/s" - msg wall cpu gpu = printf "gc: %s: %s @ %s, %s" - name - (showBytes bytes) - (showRate bytes wall) - (Debug.elapsed wall cpu gpu) + let fmt wall cpu gpu = + message (builder % ": " % F.bytes @Double shortest % " @ " % Debug.formatSIBase (Just 3) 1024 % "B/s, " % Debug.elapsed) + name bytes (double bytes / wall) wall cpu gpu in - liftIO (Debug.timed Debug.dump_gc msg stream action) + liftIO (Debug.timed Debug.dump_gc fmt stream action) diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs b/src/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs @@ -1,6 +1,7 @@ {-# LANGUAGE BangPatterns #-} {-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE MagicHash #-} +{-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE RecordWildCards #-} {-# LANGUAGE ScopedTypeVariables #-} {-# LANGUAGE TypeApplications #-} @@ -33,20 +34,21 @@ import Data.Array.Accelerate.Representation.Elt import Data.Array.Accelerate.Representation.Type import Data.Array.Accelerate.Type -import qualified Data.Array.Accelerate.Array.Remote as Remote -import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug +import qualified Data.Array.Accelerate.Array.Remote as Remote +import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug import Foreign.CUDA.Driver.Error -import qualified Foreign.CUDA.Ptr as CUDA -import qualified Foreign.CUDA.Driver as CUDA -import qualified Foreign.CUDA.Driver.Stream as CUDA +import qualified Foreign.CUDA.Ptr as CUDA +import qualified Foreign.CUDA.Driver as CUDA +import qualified Foreign.CUDA.Driver.Stream as CUDA import Control.Exception -import Control.Monad.State -import Text.Printf +import Control.Monad.Reader +import Data.Text.Lazy.Builder +import Formatting hiding ( bytes ) +import qualified Formatting as F -import GHC.Base -import GHC.Int +import GHC.Base ( Int(..), Double(..), int2Double# ) -- Events signal once a computation has completed @@ -60,14 +62,16 @@ -- mallocRemote n | n <= 0 = return (Just CUDA.nullDevPtr) - | otherwise = liftIO $ do - ep <- try (CUDA.mallocArray n) - case ep of - Right p -> do liftIO (Debug.didAllocateBytesRemote (i64 n)) - return (Just p) - Left (ExitCode OutOfMemory) -> do return Nothing - Left e -> do message ("malloc failed with error: " ++ show e) - throwIO e + | otherwise = do + name <- asks ptxDeviceName + liftIO $ do + ep <- try (CUDA.mallocArray n) + case ep of + Right p -> do Debug.remote_memory_alloc name (CUDA.useDevicePtr p) n + return (Just p) + Left (ExitCode OutOfMemory) -> do return Nothing + Left e -> do message ("malloc failed with error: " % shown) e + throwIO e peekRemote t n src ad | SingleArrayDict <- singleArrayDict t @@ -77,7 +81,7 @@ in blocking $ \stream -> withLifetime stream $ \st -> do - Debug.didCopyBytesFromRemote (i64 bytes) + Debug.memcpy_from_remote bytes transfer "peekRemote" bytes (Just st) $ CUDA.peekArrayAsync n src dst (Just st) pokeRemote t n dst ad @@ -88,7 +92,7 @@ in blocking $ \stream -> withLifetime stream $ \st -> do - Debug.didCopyBytesToRemote (i64 bytes) + Debug.memcpy_to_remote bytes transfer "pokeRemote" bytes (Just st) $ CUDA.pokeArrayAsync n src dst (Just st) castRemotePtr = CUDA.castDevPtr @@ -110,7 +114,7 @@ -> Bool -> LLVM PTX Bool malloc !tp !ad !n !frozen = do - PTX{..} <- gets llvmTarget + PTX{..} <- asks llvmTarget Remote.malloc ptxMemoryTable tp ad frozen n @@ -123,7 +127,7 @@ -> (CUDA.DevicePtr (ScalarArrayDataR e) -> LLVM PTX (Maybe Event, r)) -> LLVM PTX (Maybe r) withRemote !tp !ad !f = do - PTX{..} <- gets llvmTarget + PTX{..} <- asks llvmTarget Remote.withRemote ptxMemoryTable tp ad f @@ -140,10 +144,6 @@ liftIO $ block e return r -{-# INLINE i64 #-} -i64 :: Int -> Int64 -i64 (I# i#) = I64# i# - {-# INLINE double #-} double :: Int -> Double double (I# i#) = D# (int2Double# i#) @@ -152,27 +152,16 @@ -- Debugging -- --------- -{-# INLINE showBytes #-} -showBytes :: Int -> String -showBytes x = Debug.showFFloatSIBase (Just 0) 1024 (double x) "B" - -{-# INLINE trace #-} -trace :: String -> IO a -> IO a -trace msg next = Debug.traceIO Debug.dump_gc ("gc: " ++ msg) >> next - {-# INLINE message #-} -message :: String -> IO () -message s = s `trace` return () +message :: Format (IO ()) a -> a +message fmt = Debug.traceM Debug.dump_gc ("gc: " % fmt) {-# INLINE transfer #-} -transfer :: String -> Int -> Maybe CUDA.Stream -> IO () -> IO () -transfer name bytes stream action - = let showRate x t = Debug.showFFloatSIBase (Just 3) 1024 (double x / t) "B/s" - msg wall cpu gpu = printf "gc: %s: %s bytes @ %s, %s" - name - (showBytes bytes) - (showRate bytes wall) - (Debug.elapsed wall cpu gpu) - in - Debug.timed Debug.dump_gc msg stream action +transfer :: Builder -> Int -> Maybe CUDA.Stream -> IO () -> IO () +transfer name bytes stream action = + let fmt wall cpu gpu = + message (builder % ": " % F.bytes @Double shortest % " @ " % Debug.formatSIBase (Just 3) 1024 % "B/s, " % Debug.elapsed) + name bytes (double bytes / wall) wall cpu gpu + in + Debug.timed Debug.dump_gc fmt stream action diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Array/Table.hs b/src/Data/Array/Accelerate/LLVM/PTX/Array/Table.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Array/Table.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Array/Table.hs @@ -1,4 +1,5 @@ -{-# LANGUAGE BangPatterns #-} +{-# LANGUAGE BangPatterns #-} +{-# LANGUAGE OverloadedStrings #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Array.Table -- Copyright : [2014..2020] The Accelerate Team @@ -17,6 +18,7 @@ ) where import Data.Array.Accelerate.LLVM.PTX.Context ( Context, withContext ) +import qualified Data.Array.Accelerate.LLVM.PTX.Context as Context import qualified Data.Array.Accelerate.Array.Remote as Remote import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug import {-# SOURCE #-} Data.Array.Accelerate.LLVM.PTX.Execute.Event @@ -24,7 +26,7 @@ import qualified Foreign.CUDA.Ptr as CUDA import qualified Foreign.CUDA.Driver as CUDA -import Text.Printf +import Formatting -- Remote memory tables. This builds upon the LRU-cached memory tables provided @@ -42,18 +44,15 @@ where freeRemote :: CUDA.DevicePtr a -> IO () freeRemote !ptr = do - message (printf "freeRemote %s" (show ptr)) - withContext ctx (CUDA.free ptr) + message ("freeRemote " % shown) ptr + Context.contextFinalizeResource ctx $ + withContext ctx (CUDA.free ptr) -- Debugging -- --------- -{-# INLINE trace #-} -trace :: String -> IO a -> IO a -trace msg next = Debug.traceIO Debug.dump_gc ("gc: " ++ msg) >> next - {-# INLINE message #-} -message :: String -> IO () -message s = s `trace` return () +message :: Format (IO ()) a -> a +message fmt = Debug.traceM Debug.dump_gc ("gc: " % fmt) diff --git a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs @@ -32,14 +32,14 @@ instance Skeleton PTX where - map _ = mkMap - generate _ = mkGenerate - transform _ = mkTransform - fold _ = mkFold - foldSeg _ = mkFoldSeg - scan _ = mkScan - scan' _ = mkScan' - permute _ = mkPermute - stencil1 _ = mkStencil1 - stencil2 _ = mkStencil2 + map = mkMap + generate = mkGenerate + transform = mkTransform + fold = mkFold + foldSeg = mkFoldSeg + scan = mkScan + scan' = mkScan' + permute = mkPermute + stencil1 = mkStencil1 + stencil2 = mkStencil2 diff --git a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs @@ -1,11 +1,13 @@ -{-# LANGUAGE CPP #-} +{-# LANGUAGE DataKinds #-} {-# LANGUAGE GADTs #-} +{-# LANGUAGE MultiWayIf #-} {-# LANGUAGE OverloadedStrings #-} +{-# LANGUAGE RankNTypes #-} {-# LANGUAGE ScopedTypeVariables #-} -{-# LANGUAGE TemplateHaskell #-} {-# LANGUAGE TypeApplications #-} {-# LANGUAGE TypeFamilies #-} {-# LANGUAGE ViewPatterns #-} +{-# OPTIONS_GHC -Wno-orphans #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.Base -- Copyright : [2014..2020] The Accelerate Team @@ -36,10 +38,14 @@ __syncwarp, __syncwarp_mask, __threadfence_block, __threadfence_grid, + -- Warp shuffle instructions + __shfl_up, __shfl_down, __shfl_idx, __broadcast, + -- Shared memory staticSharedMem, + sharedMemorySizeAdd, dynamicSharedMem, - sharedMemAddrSpace, + sharedMemAddrSpace, sharedMemVolatility, -- Kernel definitions (+++), @@ -47,56 +53,54 @@ ) where +import Data.Primitive.Vec import Data.Array.Accelerate.Error import Data.Array.Accelerate.LLVM.CodeGen.Arithmetic as A import Data.Array.Accelerate.LLVM.CodeGen.Base -import Data.Array.Accelerate.LLVM.CodeGen.Constant import Data.Array.Accelerate.LLVM.CodeGen.IR import Data.Array.Accelerate.LLVM.CodeGen.Module import Data.Array.Accelerate.LLVM.CodeGen.Monad import Data.Array.Accelerate.LLVM.CodeGen.Ptr import Data.Array.Accelerate.LLVM.CodeGen.Sugar +import Data.Array.Accelerate.LLVM.Compile.Cache import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch import Data.Array.Accelerate.LLVM.PTX.Target import Data.Array.Accelerate.Representation.Array import Data.Array.Accelerate.Representation.Elt import Data.Array.Accelerate.Representation.Shape import Data.Array.Accelerate.Representation.Type +import qualified Data.Array.Accelerate.LLVM.CodeGen.Constant as A import Foreign.CUDA.Analysis ( Compute(..), computeCapability ) import qualified Foreign.CUDA.Analysis as CUDA -import LLVM.AST.Type.AddrSpace import LLVM.AST.Type.Constant import LLVM.AST.Type.Downcast import LLVM.AST.Type.Function -import LLVM.AST.Type.InlineAssembly +import LLVM.AST.Type.GetElementPtr import LLVM.AST.Type.Instruction +import LLVM.AST.Type.Instruction.Atomic +import qualified LLVM.AST.Type.Instruction.RMW as RMW import LLVM.AST.Type.Instruction.Volatile import LLVM.AST.Type.Metadata import LLVM.AST.Type.Name import LLVM.AST.Type.Operand import LLVM.AST.Type.Representation -import qualified LLVM.AST.Constant as LLVM hiding ( type' ) -import qualified LLVM.AST.Global as LLVM -import qualified LLVM.AST.Linkage as LLVM -import qualified LLVM.AST.Name as LLVM -import qualified LLVM.AST.Type as LLVM +import qualified Data.Array.Accelerate.LLVM.Internal.LLVMPretty as LP import Control.Applicative import Control.Monad ( void ) -import Control.Monad.State ( gets ) +import Control.Monad.Reader ( asks ) +import Data.Bits +import Data.Proxy +import Data.String +import Foreign.Storable import Prelude as P -#if MIN_VERSION_llvm_hs(10,0,0) -import qualified LLVM.AST.Type.Instruction.RMW as RMW -import LLVM.AST.Type.Instruction.Atomic -#elif !MIN_VERSION_llvm_hs(9,0,0) -import Data.String -import Text.Printf -#endif +import GHC.TypeLits + -- Thread identifiers -- ------------------ @@ -135,7 +139,7 @@ -- warpId :: CodeGen PTX (Operands Int32) warpId = do - dev <- liftCodeGen $ gets ptxDeviceProperties + dev <- liftCodeGen $ asks ptxDeviceProperties tid <- threadIdx A.quot integralType tid (A.liftInt32 (P.fromIntegral (CUDA.warpSize dev))) @@ -240,15 +244,12 @@ -- __syncwarp_mask :: HasCallStack => Operands Word32 -> CodeGen PTX () __syncwarp_mask mask = do - dev <- liftCodeGen $ gets ptxDeviceProperties - if computeCapability dev < Compute 7 0 - then return () - else -#if !MIN_VERSION_llvm_hs(6,0,0) - internalError "LLVM-6.0 or above is required for Volta devices and later" -#else - void $ call (Lam primType (op primType mask) (Body VoidType (Just Tail) "llvm.nvvm.bar.warp.sync")) [NoUnwind, NoDuplicate, Convergent] -#endif + llvmver <- getLLVMversion + dev <- liftCodeGen $ asks ptxDeviceProperties + case (computeCapability dev >= Compute 7 0, llvmver >= 6) of + (True, True) -> void $ call (Lam primType (op primType mask) (Body VoidType (Just Tail) "llvm.nvvm.bar.warp.sync")) [NoUnwind, NoDuplicate, Convergent] + (True, False) -> internalError "LLVM-6.0 or above is required for Volta devices and later" + (False, _) -> return () -- | Ensure that all writes to shared and global memory before the call to @@ -289,42 +290,274 @@ -- -- atomicAdd_f :: HasCallStack => FloatingType a -> Operand (Ptr a) -> Operand a -> CodeGen PTX () -atomicAdd_f t addr val = -#if MIN_VERSION_llvm_hs(10,0,0) - void . instr' $ AtomicRMW (FloatingNumType t) NonVolatile RMW.FAdd addr val (CrossThread, AcquireRelease) -#else +atomicAdd_f t addr val = do + llvmver <- getLLVMversion + if | llvmver >= 10 -> + void . instr' $ AtomicRMW (FloatingNumType t) NonVolatile RMW.Add addr val (CrossThread, AcquireRelease) + + | otherwise -> + internalError "LLVM < 10 not supported" + + +-- Warp shuffle functions +-- ---------------------- +-- +-- Exchange a variable between threads within a warp. Requires compute +-- capability 3.0 or higher. +-- +-- +-- +-- Each of the shfl primitives also exists in ".p" form. This version +-- returns, alongside the normal value, a boolean that returns whether the +-- source lane was in range. This could be useful when doing bounds checks +-- in folds and scans. +-- + +data ShuffleOp + = Idx -- ^ Direct copy from indexed lane + | Up -- ^ Copy from a lane with lower ID relative to the caller + | Down -- ^ Copy from a lane with higher ID relative to the caller + | XOR -- ^ Copy from a lane based on bitwise XOR of own lane ID + +-- | Each thread gets the value provided by lower threads +-- +__shfl_up :: TypeR a -> Operands a -> Operands Word32 -> CodeGen PTX (Operands a) +__shfl_up = shfl Up + +-- | Each thread gets the value provided by higher threads +-- +__shfl_down :: TypeR a -> Operands a -> Operands Word32 -> CodeGen PTX (Operands a) +__shfl_down = shfl Down + +-- | shfl_idx takes an argument representing the source lane index. +-- +__shfl_idx :: TypeR a -> Operands a -> Operands Word32 -> CodeGen PTX (Operands a) +__shfl_idx = shfl Idx + +-- | Distribute the value from lane 0 across the warp +-- +__broadcast :: TypeR a -> Operands a -> CodeGen PTX (Operands a) +__broadcast aR a = __shfl_idx aR a (liftWord32 0) + + +shfl :: ShuffleOp + -> TypeR a + -> Operands a + -> Operands Word32 + -> CodeGen PTX (Operands a) +shfl sop tR val delta = go tR val + where + delta' :: Operand Word32 + delta' = op integralType delta + + go :: TypeR s -> Operands s -> CodeGen PTX (Operands s) + go TupRunit OP_Unit = return OP_Unit + go (TupRpair aR bR) (OP_Pair a b) = OP_Pair <$> go aR a <*> go bR b + go (TupRsingle t) a = scalar t a + + scalar :: ScalarType s -> Operands s -> CodeGen PTX (Operands s) + scalar (SingleScalarType t) = single t + scalar (VectorScalarType t) = vector t + + single :: SingleType s -> Operands s -> CodeGen PTX (Operands s) + single (NumSingleType t) = num t + + vector :: forall n s. VectorType (Vec n s) -> Operands (Vec n s) -> CodeGen PTX (Operands (Vec n s)) + vector v@(VectorType w t) a + | SingleDict <- singleDict t + = let bytes = sizeOf (undefined::s) + (m,r) = P.quotRem (w * bytes) 4 + + withSomeNat :: Int -> (forall m. KnownNat m => Proxy m -> b) -> b + withSomeNat n k = + case someNatVal (toInteger n) of + Nothing -> error "Welcome to overthinkers club. The first rule of overthinkers club is: yet to be decided." + Just (SomeNat p) -> k p + in + if r == 0 + -- bitcast into a vector + -- special case for a single element vector + then + if m == 1 + then do + b <- A.bitcast (VectorScalarType v) (scalarType @Int32) a + c <- integral (integralType @Int32) b + d <- A.bitcast scalarType (VectorScalarType v) c + return d + + else + let + vec :: forall m. KnownNat m => Proxy m -> CodeGen PTX (Operands (Vec n s)) + vec _ = do + let v' = VectorType m (singleType @Int32) + + b <- A.bitcast (VectorScalarType v) (VectorScalarType v') a + + let c = op v' b + + repack :: Int32 -> CodeGen PTX (Operands (Vec m Int32)) + repack 0 = return $ ir v' (A.undef (VectorScalarType v')) + repack i = do + d <- instr $ ExtractElement (i-1) c + e <- integral integralType d + f <- repack (i-1) + g <- instr $ InsertElement (i-1) (op v' f) (op integralType e) + return g + + h <- repack (P.fromIntegral m) + i <- A.bitcast (VectorScalarType v') (VectorScalarType v) h + return i + in + withSomeNat m vec + + -- Round up to the next multiple of 32: + -- + -- 1. bitcast to an integer of the same number of bits: e.g. bitcast <3 x i16> i48 + -- 2. extend that to the next multiple of 32: e.g. zext i48 i64 + -- 3. bitcast to : e.g. bitcast i64 <2 x i32> + -- + else + let raw :: LP.Type -> LP.Instr -> CodeGen PTX (LP.Typed LP.Value) + raw ty ins = do + name <- freshLocalName + instr_ (LP.Result (nameToPrettyI name) ins [] []) + return (LP.Typed ty (LP.ValIdent (nameToPrettyI name))) + + rawUp :: Type u -> LP.Instr -> CodeGen PTX (Operand u) + rawUp ty ins = do + name <- freshLocalName + instr_ (LP.Result (nameToPrettyI name) ins [] []) + return (LocalReference ty name) + + + vec :: forall m. KnownNat m => Proxy m -> CodeGen PTX (Operands (Vec n s)) + vec _ = do + let t0Up :: Type (Vec n s) + t0Up = PrimType (ScalarPrimType (VectorScalarType v)) + t0 = downcast t0Up + + t1 = LP.PrimType (LP.Integer (P.fromIntegral ((w*bytes) * 8))) + t2 = LP.PrimType (LP.Integer (P.fromIntegral ((m+1) * 4 * 8))) + + v' :: VectorType (Vec m Int32) + v' = VectorType (m+1) (singleType @Int32) + t3Up :: Type (Vec m Int32) + t3Up = PrimType (ScalarPrimType (VectorScalarType v')) + t3 = downcast t3Up + + b <- raw t1 (LP.Conv LP.BitCast (downcast (op v a)) t1) + c <- raw t2 (LP.Conv (LP.ZExt False) b t2) + d <- rawUp t3Up (LP.Conv LP.BitCast c t3) + e <- vector v' (ir v' d) + f <- raw t2 (LP.Conv LP.BitCast (downcast (op v' e)) t2) + g <- raw t1 (LP.Conv (LP.Trunc False False) f t1) + h <- rawUp t0Up (LP.Conv LP.BitCast g t0) + return (ir v h) + in + withSomeNat (m+1) vec + + num :: NumType s -> Operands s -> CodeGen PTX (Operands s) + num (IntegralNumType t) = integral t + num (FloatingNumType t) = floating t + + integral :: forall s. IntegralType s -> Operands s -> CodeGen PTX (Operands s) + integral TypeInt32 a = shfl_op sop ShuffleInt32 delta' a + integral t a + | IntegralDict <- integralDict t + = case finiteBitSize (undefined::s) of + 64 -> do + let ta = SingleScalarType (NumSingleType (IntegralNumType t)) + tb = scalarType @(Vec 2 Int32) + -- + b <- A.bitcast ta tb a + c <- vector (VectorType 2 singleType) b + d <- A.bitcast tb ta c + return d + + _ -> do + b <- A.fromIntegral t (numType @Int32) a + c <- integral integralType b + d <- A.fromIntegral integralType (IntegralNumType t) c + return d + + floating :: FloatingType s -> Operands s -> CodeGen PTX (Operands s) + floating TypeFloat a = shfl_op sop ShuffleFloat delta' a + floating TypeDouble a = do + b <- A.bitcast scalarType (scalarType @(Vec 2 Int32)) a + c <- vector (VectorType 2 singleType) b + d <- A.bitcast scalarType (scalarType @Double) c + return d + floating TypeHalf a = do + b <- A.bitcast scalarType (scalarType @Int16) a + c <- integral integralType b + d <- A.bitcast scalarType (scalarType @Half) c + return d + + +data ShuffleType a where + ShuffleInt32 :: ShuffleType Int32 + ShuffleFloat :: ShuffleType Float + +shfl_op + :: forall a. + ShuffleOp + -> ShuffleType a + -> Operand Word32 -- delta + -> Operands a -- value to give + -> CodeGen PTX (Operands a) -- value received +shfl_op sop t delta val = do + dev <- liftCodeGen $ asks ptxDeviceProperties + let - _width :: Int - _width = - case t of - TypeHalf -> 16 - TypeFloat -> 32 - TypeDouble -> 64 + -- The CUDA __shfl* instruction take an optional final parameter + -- which is the warp size. Setting this value to something (always + -- a power-of-two) other than 32 emulates the shfl behaviour at that + -- warp size. Behind the scenes, a bunch of instructions happen with + -- this width parameter before they get passed into the actual shfl + -- instruction. Here, we have to directly set them to the 'actual' + -- width parameter. The formula that clang compiles to is in the + -- comments + -- + width :: Operand Int32 + width = case sop of + Up -> A.integral integralType 0 -- ((32 - warpSize) `shiftL` 8) + Down -> A.integral integralType 31 -- ((32 - warpSize) `shiftL` 8) `or` 31 + Idx -> A.integral integralType 31 -- ((32 - warpSize) `shiftL` 8) `or` 31 + XOR -> A.integral integralType 31 -- ((32 - warpSize) `shiftL` 8) `or` 31 - (t_addr, t_val, _addrspace) = - case typeOf addr of - PrimType ta@(PtrPrimType (ScalarPrimType tv) (AddrSpace as)) - -> (ta, tv, as) - _ -> internalError "unexpected operand type" + -- Starting CUDA 9.0, the normal `shfl` primitives are deprecated in + -- favour of the newer `shfl_sync` primitives. They behave the same + -- way, except they start with a 'mask' argument specifying which + -- threads participate in the shuffle. + -- + mask :: Operand Int32 + mask = A.integral integralType (-1) -- all threads participate - t_ret = PrimType (ScalarPrimType t_val) -#if MIN_VERSION_llvm_hs(9,0,0) || !MIN_VERSION_llvm_hs(6,0,0) - asm = - case t of - -- assuming .address_size 64 - TypeHalf -> InlineAssembly "atom.add.noftz.f16 $0, [$1], $2;" "=c,l,c" True False ATTDialect - TypeFloat -> InlineAssembly "atom.global.add.f32 $0, [$1], $2;" "=f,l,f" True False ATTDialect - TypeDouble -> InlineAssembly "atom.global.add.f64 $0, [$1], $2;" "=d,l,d" True False ATTDialect - in - void $ instr (Call (Lam t_addr addr (Lam (ScalarPrimType t_val) val (Body t_ret (Just Tail) (Left asm)))) [Right NoUnwind]) -#else - fun = fromString $ printf "llvm.nvvm.atomic.load.add.f%d.p%df%d" _width (_addrspace :: Word32) _width - in - void $ call (Lam t_addr addr (Lam (ScalarPrimType t_val) val (Body t_ret (Just Tail) fun))) [NoUnwind] -#endif -#endif + useSyncShfl = CUDA.computeCapability dev >= Compute 7 0 + call' = if useSyncShfl + then call . Lam primType mask + else call + sync = if useSyncShfl then "sync." else "" + asm = "llvm.nvvm.shfl." + <> sync + <> case sop of + Idx -> "idx." + Up -> "up." + Down -> "down." + XOR -> "bfly." + <> case t of + ShuffleInt32 -> "i32" + ShuffleFloat -> "f32" + + t_val = case t of + ShuffleInt32 -> primType :: PrimType Int32 + ShuffleFloat -> primType :: PrimType Float + + call' (Lam t_val (op t_val val) (Lam primType delta (Lam primType width (Body (PrimType t_val) (Just Tail) asm)))) [Convergent, NoUnwind, InaccessibleMemOnly] + + -- Shared memory -- ------------- @@ -338,6 +571,28 @@ -- Declare a new statically allocated array in the __shared__ memory address -- space, with enough storage to contain the given number of elements. -- +-- Previously, like initialiseDynamicSharedMemory, this function declared an +-- external global, e.g. for 1 i64: +-- @sdata = external addrspace(3) global [1 x i64], align 8 +-- This would correspond to the following CUDA source: +-- extern __shared__ int64_t sdata[1]; +-- +-- But this CUDA C++ is rejected by Clang. When LLVM is fed LLVM IR, however, +-- things are more subtle; in the old llvm-hs backend where we linked against +-- LLVM, with LLVM 15, the above IR (defining @0) was accepted. However, +-- passing this same IR to Clang 18 with the llvm-pretty backend (yes I'm aware +-- the clang version is also changing here), clang first calls ptxas and then +-- nvlink; nvlink complains: +-- Undefined reference to 'sdata' in '/tmp/test-409abe.cubin' +-- When linking against LLVM 15, nvlink is never invoked, but instead ptxas is +-- _not_ given the -c flag and it immediately produces a SASS file. +-- +-- Because Clang doesn't even accept the corresponding C++ code, but does +-- accept this: +-- __shared__ int64_t sdata[1]; +-- the global created in this function was changed to be of internal linkage +-- instead. The assigned value is 'undef', just like what Clang generates for +-- the internal sdata C++ declaration. staticSharedMem :: TypeR e -> Word64 @@ -345,7 +600,7 @@ staticSharedMem tp n = do ad <- go tp return $ IRArray { irArrayRepr = ArrayR dim1 tp - , irArrayShape = OP_Pair OP_Unit $ OP_Int $ integral integralType $ P.fromIntegral n + , irArrayShape = OP_Pair OP_Unit $ OP_Int $ A.integral integralType $ P.fromIntegral n , irArrayData = ad , irArrayAddrSpace = sharedMemAddrSpace , irArrayVolatility = sharedMemVolatility @@ -357,20 +612,28 @@ go tt@(TupRsingle t) = do -- Declare a new global reference for the statically allocated array -- located in the __shared__ memory space. - nm <- freshName - sm <- return $ ConstantOperand $ GlobalReference (PrimType (PtrPrimType (ArrayPrimType n t) sharedMemAddrSpace)) nm - declare $ LLVM.globalVariableDefaults - { LLVM.addrSpace = sharedMemAddrSpace - , LLVM.type' = LLVM.ArrayType n (downcast t) - , LLVM.linkage = LLVM.External - , LLVM.name = downcast nm - , LLVM.alignment = 4 `P.max` P.fromIntegral (bytesElt tt) + nm <- freshGlobalName + let arrt = ArrayPrimType n t + ptrarrt = PrimType (PtrPrimType arrt sharedMemAddrSpace) + sm <- return $ ConstantOperand $ GlobalReference ptrarrt nm + declareGlobalVar $ LP.Global + { LP.globalSym = nameToPrettyS nm + , LP.globalAttrs = LP.GlobalAttrs + { LP.gaLinkage = Just LP.Internal + , LP.gaVisibility = Nothing + , LP.gaAddrSpace = sharedMemAddrSpace + , LP.gaConstant = False } + , LP.globalType = LP.Array n (downcast t) + , LP.globalValue = Just LP.ValUndef + , LP.globalAlign = Just (4 `P.max` P.fromIntegral (bytesElt tt)) + , LP.globalMetadata = mempty } -- Return a pointer to the first element of the __shared__ memory array. -- We do this rather than just returning the global reference directly due -- to how __shared__ memory needs to be indexed with the GEP instruction. - p <- instr' $ GetElementPtr sm [num numType 0, num numType 0 :: Operand Int32] + p <- instr' $ GetElementPtr (GEP (PrimType arrt) sm (A.num numType 0 :: Operand Int32) + (GEPArray (A.num numType 0 :: Operand Int32) (GEPEmpty (ScalarPrimType t)))) q <- instr' $ PtrCast (PtrPrimType (ScalarPrimType t) sharedMemAddrSpace) p return $ ir t (unPtr q) @@ -382,19 +645,46 @@ -- -- > @__shared__ = external addrspace(3) global [0 x i8] -- -initialiseDynamicSharedMemory :: CodeGen PTX (Operand (Ptr Word8)) +initialiseDynamicSharedMemory :: CodeGen PTX (Operand (Ptr Int8)) initialiseDynamicSharedMemory = do - declare $ LLVM.globalVariableDefaults - { LLVM.addrSpace = sharedMemAddrSpace - , LLVM.type' = LLVM.ArrayType 0 (LLVM.IntegerType 8) - , LLVM.linkage = LLVM.External - , LLVM.name = LLVM.Name "__shared__" - , LLVM.alignment = 4 + declareGlobalVar $ LP.Global + { LP.globalSym = LP.Symbol "__shared__" + , LP.globalAttrs = LP.GlobalAttrs + { LP.gaLinkage = Just LP.External + , LP.gaVisibility = Nothing + , LP.gaAddrSpace = sharedMemAddrSpace + , LP.gaConstant = False } + , LP.globalType = LP.Array 0 (LP.PrimType (LP.Integer 8)) + , LP.globalValue = Nothing + , LP.globalAlign = Nothing + , LP.globalMetadata = mempty } - return $ ConstantOperand $ GlobalReference (PrimType (PtrPrimType (ArrayPrimType 0 scalarType) sharedMemAddrSpace)) "__shared__" + return $ ConstantOperand + $ ConstantGetElementPtr (GEP (PrimType (ArrayPrimType 0 (scalarType @Int8))) + (GlobalReference (PrimType (PtrPrimType (ArrayPrimType 0 scalarType) sharedMemAddrSpace)) "__shared__") + (ScalarConstant (scalarType @Int32) 0) + (GEPArray (ScalarConstant (scalarType @Int32) 0) (GEPEmpty primType))) +sharedMemorySizeAdd + :: TypeR e + -> Int -- number of array elements + -> Int -- #bytes of shared memory the have already been allocated + -> Int +sharedMemorySizeAdd tp n i = case tp of + TupRunit -> i + TupRpair t2 t1 -> + -- First handle the second element of the tuple, then the first, + -- to match the behaviour of dynamicSharedMem + sharedMemorySizeAdd t2 n $ sharedMemorySizeAdd t1 n i + TupRsingle t -> + let + bytes = bytesElt tp + -- Align 'i' to the alignment of t + aligned = alignTo (scalarAlignment t) i + in + aligned + bytes * n --- Declared a new dynamically allocated array in the __shared__ memory space +-- Declare a new dynamically allocated array in the __shared__ memory space -- with enough space to contain the given number of elements. -- dynamicSharedMem @@ -402,7 +692,7 @@ TypeR e -> IntegralType int -> Operands int -- number of array elements - -> Operands int -- #bytes of shared memory the have already been allocated + -> Operands int -- #bytes of shared memory that have already been allocated -> CodeGen PTX (IRArray (Vector e)) dynamicSharedMem tp int n@(op int -> m) (op int -> offset) | IntegralDict <- integralDict int = do @@ -417,10 +707,14 @@ (i2, p2) <- go t2 i1 return $ (i2, OP_Pair p2 p1) go (TupRsingle t) i = do - p <- instr' $ GetElementPtr smem [num numTp 0, i] -- TLM: note initial zero index!! + let bytes = bytesElt (TupRsingle t) + let align = scalarAlignment t + i' <- instr' $ Add numTp i (A.integral int $ P.fromIntegral $ align - 1) + aligned <- instr' $ BAnd int i' (A.integral int $ P.fromIntegral $ Data.Bits.complement $ align - 1) + p <- instr' $ GetElementPtr (GEP1 scalarType smem aligned) q <- instr' $ PtrCast (PtrPrimType (ScalarPrimType t) sharedMemAddrSpace) p - a <- instr' $ Mul numTp m (integral int (P.fromIntegral (bytesElt (TupRsingle t)))) - b <- instr' $ Add numTp i a + a <- instr' $ Mul numTp m (A.integral int (P.fromIntegral bytes)) + b <- instr' $ Add numTp aligned a return (b, ir t (unPtr q)) -- (_, ad) <- go tp offset @@ -440,12 +734,12 @@ -- Requires compute capability >= 7.0 -- nanosleep :: Operands Int32 -> CodeGen PTX () -nanosleep ns = - let - attrs = [NoUnwind, Convergent] - asm = InlineAssembly "nanosleep.u32 $0;" "r" True False ATTDialect - in - void $ instr (Call (Lam primType (op integralType ns) (Body VoidType (Just Tail) (Left asm))) (map Right attrs)) +nanosleep ns = do + -- This is an acc prelude function because it requires inline assembly, and + -- llvm-pretty does not yet support caling inline assembly snippets. Thus we + -- manually wrap the assembly in an inlineable function and call that. + let label = makeAccPreludeLabel "nanosleep" + void $ instr (Call (Lam primType (op integralType ns) (Body VoidType (Just Tail) (Right label)))) -- Global kernel definitions @@ -462,24 +756,26 @@ -- | Create a single kernel program with the default launch configuration. -- makeOpenAcc - :: Label - -> [LLVM.Parameter] + :: UID + -> Label + -> [LP.Typed LP.Ident] -> CodeGen PTX () -> CodeGen PTX (IROpenAcc PTX aenv a) -makeOpenAcc name param kernel = do - dev <- liftCodeGen $ gets ptxDeviceProperties - makeOpenAccWith (simpleLaunchConfig dev) name param kernel +makeOpenAcc uid name param kernel = do + dev <- liftCodeGen $ asks ptxDeviceProperties + makeOpenAccWith (simpleLaunchConfig dev) uid name param kernel -- | Create a single kernel program with the given launch analysis information. -- makeOpenAccWith :: LaunchConfig + -> UID -> Label - -> [LLVM.Parameter] + -> [LP.Typed LP.Ident] -> CodeGen PTX () -> CodeGen PTX (IROpenAcc PTX aenv a) -makeOpenAccWith config name param kernel = do - body <- makeKernel config name param kernel +makeOpenAccWith config uid name param kernel = do + body <- makeKernel config (name <> fromString ('_' : show uid)) param kernel return $ IROpenAcc [body] -- | Create a complete kernel function by running the code generation process @@ -488,24 +784,42 @@ makeKernel :: LaunchConfig -> Label - -> [LLVM.Parameter] + -> [LP.Typed LP.Ident] -> CodeGen PTX () -> CodeGen PTX (Kernel PTX aenv a) -makeKernel config name@(Label l) param kernel = do +makeKernel config name param kernel = do _ <- kernel code <- createBlocks + let define = LP.Define + { LP.defLinkage = Nothing + , LP.defVisibility = Nothing + , LP.defRetType = LP.PrimType LP.Void + , LP.defName = labelToPrettyS name + , LP.defArgs = param + , LP.defVarArgs = False + , LP.defAttrs = [] + , LP.defSection = Nothing + , LP.defGC = Nothing + , LP.defBody = code + , LP.defMetadata = mempty + , LP.defComdat = Nothing + } addMetadata "nvvm.annotations" - [ Just . MetadataConstantOperand $ LLVM.GlobalReference (LLVM.PointerType (LLVM.FunctionType LLVM.VoidType [ t | LLVM.Parameter t _ _ <- param ] False) (AddrSpace 0)) (LLVM.Name l) + [ Just . MetadataConstantOperand + $ LP.Typed (LP.defFunType define) (LP.ValSymbol (labelToPrettyS name)) , Just . MetadataStringOperand $ "kernel" - , Just . MetadataConstantOperand $ LLVM.Int 32 1 + , Just . MetadataConstantOperand $ LP.Typed (LP.PrimType (LP.Integer 32)) (LP.ValInteger 1) ] return $ Kernel { kernelMetadata = KM_PTX config - , unKernel = LLVM.functionDefaults - { LLVM.returnType = LLVM.VoidType - , LLVM.name = downcast name - , LLVM.parameters = (param, False) - , LLVM.basicBlocks = code - } + , unKernel = define } +scalarAlignment :: ScalarType t -> Int +scalarAlignment t@(SingleScalarType _) = bytesElt (TupRsingle t) +scalarAlignment (VectorScalarType (VectorType _ t)) = bytesElt (TupRsingle $ SingleScalarType t) + +-- Align 'ptr' to the given alignment. +-- Assumes 'align' is a power of 2. +alignTo :: Int -> Int -> Int +alignTo align ptr = (ptr + align - 1) .&. Data.Bits.complement (align - 1) diff --git a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs @@ -21,7 +21,6 @@ where import Data.Array.Accelerate.Representation.Array -import Data.Array.Accelerate.Representation.Elt import Data.Array.Accelerate.Representation.Shape hiding ( size ) import Data.Array.Accelerate.Representation.Type @@ -35,6 +34,7 @@ import Data.Array.Accelerate.LLVM.CodeGen.Loop as Loop import Data.Array.Accelerate.LLVM.CodeGen.Monad import Data.Array.Accelerate.LLVM.CodeGen.Sugar +import Data.Array.Accelerate.LLVM.Compile.Cache import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base @@ -46,7 +46,7 @@ import qualified Foreign.CUDA.Analysis as CUDA import Control.Monad ( (>=>) ) -import Control.Monad.State ( gets ) +import Control.Monad.Reader ( asks ) import Data.String ( fromString ) import Data.Bits as P import Prelude as P @@ -61,19 +61,20 @@ -- mkFold :: forall aenv sh e. - Gamma aenv + UID + -> Gamma aenv -> ArrayR (Array sh e) -> IRFun2 PTX aenv (e -> e -> e) -> Maybe (IRExp PTX aenv e) -> MIRDelayed PTX aenv (Array (sh, Int) e) -> CodeGen PTX (IROpenAcc PTX aenv (Array sh e)) -mkFold aenv repr f z acc = case z of - Just z' -> (+++) <$> codeFold <*> mkFoldFill aenv repr z' +mkFold uid aenv repr f z acc = case z of + Just z' -> (+++) <$> codeFold <*> mkFoldFill uid aenv repr z' Nothing -> codeFold where codeFold = case repr of - ArrayR ShapeRz tp -> mkFoldAll aenv tp f z acc - _ -> mkFoldDim aenv repr f z acc + ArrayR ShapeRz tp -> mkFoldAll uid aenv tp f z acc + _ -> mkFoldDim uid aenv repr f z acc -- Reduce an array to a single element. @@ -96,17 +97,18 @@ -- mkFoldAll :: forall aenv e. - Gamma aenv -- ^ array environment + UID + -> Gamma aenv -- ^ array environment -> TypeR e -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function -> MIRExp PTX aenv e -- ^ (optional) initial element for exclusive reductions -> MIRDelayed PTX aenv (Vector e) -- ^ input data -> CodeGen PTX (IROpenAcc PTX aenv (Scalar e)) -mkFoldAll aenv tp combine mseed macc = do - dev <- liftCodeGen $ gets ptxDeviceProperties - foldr1 (+++) <$> sequence [ mkFoldAllS dev aenv tp combine mseed macc - , mkFoldAllM1 dev aenv tp combine macc - , mkFoldAllM2 dev aenv tp combine mseed +mkFoldAll uid aenv tp combine mseed macc = do + dev <- liftCodeGen $ asks ptxDeviceProperties + foldr1 (+++) <$> sequence [ mkFoldAllS uid dev aenv tp combine mseed macc + , mkFoldAllM1 uid dev aenv tp combine macc + , mkFoldAllM2 uid dev aenv tp combine mseed ] @@ -115,28 +117,27 @@ -- mkFoldAllS :: forall aenv e. - DeviceProperties -- ^ properties of the target GPU + UID + -> DeviceProperties -- ^ properties of the target GPU -> Gamma aenv -- ^ array environment -> TypeR e -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function -> MIRExp PTX aenv e -- ^ (optional) initial element for exclusive reductions -> MIRDelayed PTX aenv (Vector e) -- ^ input data -> CodeGen PTX (IROpenAcc PTX aenv (Scalar e)) -mkFoldAllS dev aenv tp combine mseed marr = +mkFoldAllS uid dev aenv tp combine mseed marr = let (arrOut, paramOut) = mutableArray (ArrayR dim0 tp) "out" (arrIn, paramIn) = delayedArray "in" marr paramEnv = envParam aenv -- config = launchConfig dev (CUDA.incWarp dev) smem multipleOf multipleOfQ - smem n = warps * (1 + per_warp) * bytes + smem n = sharedMemorySizeAdd tp warps 0 where ws = CUDA.warpSize dev warps = n `P.quot` ws - per_warp = ws + ws `P.quot` 2 - bytes = bytesElt tp in - makeOpenAccWith config "foldAllS" (paramOut ++ paramIn ++ paramEnv) $ do + makeOpenAccWith config uid "foldAllS" (paramOut ++ paramIn ++ paramEnv) $ do tid <- threadIdx bd <- blockDim @@ -155,8 +156,8 @@ -- reduction. x0 <- app1 (delayedLinearIndex arrIn) =<< int i0 r0 <- if (tp, A.eq singleType sz bd) - then reduceBlockSMem dev tp combine Nothing x0 - else reduceBlockSMem dev tp combine (Just sz) x0 + then reduceBlock dev tp combine Nothing x0 + else reduceBlock dev tp combine (Just sz) x0 when (A.eq singleType tid (liftInt32 0)) $ writeArray TypeInt32 arrOut tid =<< @@ -173,13 +174,14 @@ -- mkFoldAllM1 :: forall aenv e. - DeviceProperties -- ^ properties of the target GPU + UID + -> DeviceProperties -- ^ properties of the target GPU -> Gamma aenv -- ^ array environment -> TypeR e -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function -> MIRDelayed PTX aenv (Vector e) -- ^ input data -> CodeGen PTX (IROpenAcc PTX aenv (Scalar e)) -mkFoldAllM1 dev aenv tp combine marr = +mkFoldAllM1 uid dev aenv tp combine marr = let (arrTmp, paramTmp) = mutableArray (ArrayR dim1 tp) "tmp" (arrIn, paramIn) = delayedArray "in" marr @@ -187,14 +189,12 @@ start = liftInt 0 -- config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] - smem n = warps * (1 + per_warp) * bytes + smem n = sharedMemorySizeAdd tp warps 0 where ws = CUDA.warpSize dev warps = n `P.quot` ws - per_warp = ws + ws `P.quot` 2 - bytes = bytesElt tp in - makeOpenAccWith config "foldAllM1" (paramTmp ++ paramIn ++ paramEnv) $ do + makeOpenAccWith config uid "foldAllM1" (paramTmp ++ paramIn ++ paramEnv) $ do -- Each thread block cooperatively reduces a stripe of the input and stores -- that value into a temporary array at a corresponding index. Since the @@ -229,13 +229,14 @@ -- mkFoldAllM2 :: forall aenv e. - DeviceProperties + UID + -> DeviceProperties -> Gamma aenv -> TypeR e -> IRFun2 PTX aenv (e -> e -> e) -> MIRExp PTX aenv e -> CodeGen PTX (IROpenAcc PTX aenv (Scalar e)) -mkFoldAllM2 dev aenv tp combine mseed = +mkFoldAllM2 uid dev aenv tp combine mseed = let (arrTmp, paramTmp) = mutableArray (ArrayR dim1 tp) "tmp" (arrOut, paramOut) = mutableArray (ArrayR dim1 tp) "out" @@ -243,14 +244,12 @@ start = liftInt 0 -- config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] - smem n = warps * (1 + per_warp) * bytes + smem n = sharedMemorySizeAdd tp warps 0 where ws = CUDA.warpSize dev warps = n `P.quot` ws - per_warp = ws + ws `P.quot` 2 - bytes = bytesElt tp in - makeOpenAccWith config "foldAllM2" (paramTmp ++ paramOut ++ paramEnv) $ do + makeOpenAccWith config uid "foldAllM2" (paramTmp ++ paramOut ++ paramEnv) $ do -- Threads cooperatively reduce a stripe of the input (temporary) array -- output from the first phase, storing the results into another temporary. @@ -296,14 +295,15 @@ -- mkFoldDim :: forall aenv sh e. - Gamma aenv -- ^ array environment + UID + -> Gamma aenv -- ^ array environment -> ArrayR (Array sh e) -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function -> MIRExp PTX aenv e -- ^ (optional) seed element, if this is an exclusive reduction -> MIRDelayed PTX aenv (Array (sh, Int) e) -- ^ input data -> CodeGen PTX (IROpenAcc PTX aenv (Array sh e)) -mkFoldDim aenv repr@(ArrayR shr tp) combine mseed marr = do - dev <- liftCodeGen $ gets ptxDeviceProperties +mkFoldDim uid aenv repr@(ArrayR shr tp) combine mseed marr = do + dev <- liftCodeGen $ asks ptxDeviceProperties -- let (arrOut, paramOut) = mutableArray repr "out" @@ -311,14 +311,12 @@ paramEnv = envParam aenv -- config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] - smem n = warps * (1 + per_warp) * bytes + smem n = sharedMemorySizeAdd tp warps 0 where ws = CUDA.warpSize dev warps = n `P.quot` ws - per_warp = ws + ws `P.quot` 2 - bytes = bytesElt tp -- - makeOpenAccWith config "fold" (paramOut ++ paramIn ++ paramEnv) $ do + makeOpenAccWith config uid "fold" (paramOut ++ paramIn ++ paramEnv) $ do -- If the innermost dimension is smaller than the number of threads in the -- block, those threads will never contribute to the output. @@ -349,8 +347,8 @@ x0 <- app1 (delayedLinearIndex arrIn) i0 bd <- blockDim r0 <- if (tp, A.gte singleType sz' bd) - then reduceBlockSMem dev tp combine Nothing x0 - else reduceBlockSMem dev tp combine (Just sz') x0 + then reduceBlock dev tp combine Nothing x0 + else reduceBlock dev tp combine (Just sz') x0 -- Step 2: keep walking over the input bd' <- int bd @@ -368,7 +366,7 @@ -- bounds checks. then do x <- app1 (delayedLinearIndex arrIn) i - y <- reduceBlockSMem dev tp combine Nothing x + y <- reduceBlock dev tp combine Nothing x return y -- Otherwise, we require bounds checks when reading the input @@ -388,7 +386,7 @@ return $ go tp v <- i32 v' - y <- reduceBlockSMem dev tp combine (Just v) x + y <- reduceBlock dev tp combine (Just v) x return y if (tp, A.eq singleType tid (liftInt32 0)) @@ -411,12 +409,13 @@ -- dimensions with the initial element. -- mkFoldFill - :: Gamma aenv + :: UID + -> Gamma aenv -> ArrayR (Array sh e) -> IRExp PTX aenv e -> CodeGen PTX (IROpenAcc PTX aenv (Array sh e)) -mkFoldFill aenv repr seed = - mkGenerate aenv repr (IRFun1 (const seed)) +mkFoldFill uid aenv repr seed = + mkGenerate uid aenv repr (IRFun1 (const seed)) -- Efficient threadblock-wide reduction using the specified operator. The @@ -427,38 +426,28 @@ -- -- Example: https://github.com/NVlabs/cub/blob/1.5.2/cub/block/specializations/block_reduce_warp_reductions.cuh -- -reduceBlockSMem +reduceBlock :: forall aenv e. DeviceProperties -- ^ properties of the target device -> TypeR e -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function - -> Maybe (Operands Int32) -- ^ number of valid elements (may be less than block size) - -> Operands e -- ^ calling thread's input element - -> CodeGen PTX (Operands e) -- ^ thread-block-wide reduction using the specified operator (lane 0 only) -reduceBlockSMem dev tp combine size = warpReduce >=> warpAggregate + -> Maybe (Operands Int32) -- ^ number of valid elements (may be less than block size) + -> Operands e -- ^ calling thread's input element + -> CodeGen PTX (Operands e) -- ^ thread-block-wide reduction using the specified operator (lane 0 only) +reduceBlock dev tp combine size = warpReduce >=> warpAggregate where int32 :: Integral a => a -> Operands Int32 int32 = liftInt32 . P.fromIntegral - -- Temporary storage required for each warp - bytes = bytesElt tp - warp_smem_elems = CUDA.warpSize dev + (CUDA.warpSize dev `P.quot` 2) - -- Step 1: Reduction in every warp -- warpReduce :: Operands e -> CodeGen PTX (Operands e) warpReduce input = do - -- Allocate (1.5 * warpSize) elements of shared memory for each warp wid <- warpId - skip <- A.mul numType wid (int32 (warp_smem_elems * bytes)) - smem <- dynamicSharedMem tp TypeInt32 (int32 warp_smem_elems) skip - -- Are we doing bounds checking for this warp? - -- case size of -- The entire thread block is valid, so skip bounds checks. - Nothing -> - reduceWarpSMem dev tp combine smem Nothing input + Nothing -> reduceWarp dev tp combine Nothing input -- Otherwise check how many elements are valid for this warp. If it is -- full then we can still skip bounds checks for it. @@ -466,8 +455,8 @@ offset <- A.mul numType wid (int32 (CUDA.warpSize dev)) valid <- A.sub numType n offset if (tp, A.gte singleType valid (int32 (CUDA.warpSize dev))) - then reduceWarpSMem dev tp combine smem Nothing input - else reduceWarpSMem dev tp combine smem (Just valid) input + then reduceWarp dev tp combine Nothing input + else reduceWarp dev tp combine (Just valid) input -- Step 2: Aggregate per-warp reductions -- @@ -476,8 +465,7 @@ -- Allocate #warps elements of shared memory bd <- blockDim warps <- A.quot integralType bd (int32 (CUDA.warpSize dev)) - skip <- A.mul numType warps (int32 (warp_smem_elems * bytes)) - smem <- dynamicSharedMem tp TypeInt32 warps skip + smem <- dynamicSharedMem tp TypeInt32 warps (liftInt32 0) -- Share the per-lane aggregates wid <- warpId @@ -488,9 +476,18 @@ -- Wait for each warp to finish its local reduction __syncthreads + -- -- Now, warp 0 will reduce all the warp-wide results. + -- -- TODO: this means that the block can only consist of 32 warps, + -- -- reducing 1024 elements total. Check if this gets handled by callers! + -- if (tp, A.eq singleType wid (liftInt32 0)) + -- then warpReduce input + -- else + -- return input + -- Update the total aggregate. Thread 0 just does this sequentially (as is -- done in CUB), but we could also do this cooperatively (better for -- larger thread blocks?) + -- tid <- threadIdx if (tp, A.eq singleType tid (liftInt32 0)) then do @@ -515,16 +512,14 @@ -- -- Example: https://github.com/NVlabs/cub/blob/1.5.2/cub/warp/specializations/warp_reduce_smem.cuh#L128 -- -reduceWarpSMem - :: forall aenv e. - DeviceProperties -- ^ properties of the target device +reduceWarp + :: forall e aenv. DeviceProperties -> TypeR e - -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function - -> IRArray (Vector e) -- ^ temporary storage array in shared memory (1.5 warp size elements) + -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function -> Maybe (Operands Int32) -- ^ number of items that will be reduced by this warp, otherwise all lanes are valid - -> Operands e -- ^ calling thread's input element - -> CodeGen PTX (Operands e) -- ^ warp-wide reduction using the specified operator (lane 0 only) -reduceWarpSMem dev tp combine smem size = reduce 0 + -> Operands e -- ^ this thread's input value + -> CodeGen PTX (Operands e) -- ^ final result +reduceWarp dev typer combine size = reduce 0 where log2 :: Double -> Double log2 = P.logBase 2 @@ -532,49 +527,30 @@ -- Number steps required to reduce warp steps = P.floor . log2 . P.fromIntegral . CUDA.warpSize $ dev - -- Return whether the index is valid. Assume that constant branches are - -- optimised away. - valid i = + valid offset = do + lane <- laneId + i <- A.add numType lane (liftInt32 offset) case size of - Nothing -> return (liftBool True) + Nothing -> A.lt singleType i (liftInt32 (P.fromIntegral (CUDA.warpSize dev))) Just n -> A.lt singleType i n -- Unfold the reduction as a recursive code generation function. reduce :: Int -> Operands e -> CodeGen PTX (Operands e) reduce step x - | step >= steps = return x + | step > steps = return x | otherwise = do - let offset = liftInt32 (1 `P.shiftL` step) + let + offset :: (Bits i, Integral i) => i + offset = 1 `P.shiftL` step - -- share input through buffer - lane <- laneId - writeArray TypeInt32 smem lane x + y <- __shfl_down typer x (liftWord32 offset) + x' <- if (typer, valid offset) + then app2 combine x y + else return x + reduce (step + 1) x' - __syncwarp - -- update input if in range - i <- A.add numType lane offset - x' <- if (tp, valid i) - then app2 combine x =<< readArray TypeInt32 smem i - else return x - __syncwarp - - reduce (step+1) x' - - --- Efficient warp reduction using __shfl_up instruction (compute >= 3.0) --- --- Example: https://github.com/NVlabs/cub/blob/1.5.2/cub/warp/specializations/warp_reduce_shfl.cuh#L310 --- --- reduceWarpShfl --- :: IRFun2 PTX aenv (e -> e -> e) -- ^ combination function --- -> Operands e -- ^ this thread's input value --- -> CodeGen (Operands e) -- ^ final result --- reduceWarpShfl combine input = --- error "TODO: PTX.reduceWarpShfl" - - -- Reduction loops -- --------------- @@ -583,8 +559,8 @@ -> TypeR a -> Operands Int -- ^ starting index -> Operands Int -- ^ final index (exclusive) - -> (IRFun2 PTX aenv (a -> a -> a)) -- ^ combination function - -> (Operands Int -> CodeGen PTX (Operands a)) -- ^ function to retrieve element at index + -> (IRFun2 PTX aenv (a -> a -> a)) -- ^ combination function + -> (Operands Int -> CodeGen PTX (Operands a)) -- ^ function to retrieve element at index -> (Operands a -> CodeGen PTX ()) -- ^ what to do with the value -> CodeGen PTX () reduceFromTo dev tp from to combine get set = do @@ -600,7 +576,7 @@ -- All threads in the block will participate in the reduction, so -- we can avoid bounds checks x <- get i - r <- reduceBlockSMem dev tp combine Nothing x + r <- reduceBlock dev tp combine Nothing x set r return (lift TupRunit ()) @@ -610,7 +586,7 @@ when (A.lt singleType i to) $ do x <- get i v <- i32 valid - r <- reduceBlockSMem dev tp combine (Just v) x + r <- reduceBlock dev tp combine (Just v) x set r return (lift TupRunit ()) diff --git a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs @@ -20,7 +20,6 @@ where import Data.Array.Accelerate.Representation.Array -import Data.Array.Accelerate.Representation.Elt import Data.Array.Accelerate.Representation.Shape import Data.Array.Accelerate.Representation.Type @@ -34,10 +33,11 @@ import Data.Array.Accelerate.LLVM.CodeGen.Loop as Loop import Data.Array.Accelerate.LLVM.CodeGen.Monad import Data.Array.Accelerate.LLVM.CodeGen.Sugar +import Data.Array.Accelerate.LLVM.Compile.Cache import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base -import Data.Array.Accelerate.LLVM.PTX.CodeGen.Fold ( reduceBlockSMem, reduceWarpSMem, imapFromTo ) +import qualified Data.Array.Accelerate.LLVM.PTX.CodeGen.Fold as Fold ( reduceBlock, reduceWarp, imapFromTo ) import Data.Array.Accelerate.LLVM.PTX.Target import LLVM.AST.Type.Representation @@ -45,7 +45,7 @@ import qualified Foreign.CUDA.Analysis as CUDA import Control.Monad ( void ) -import Control.Monad.State ( gets ) +import Control.Monad.Reader ( asks ) import Data.String ( fromString ) import Prelude as P @@ -55,7 +55,8 @@ -- mkFoldSeg :: forall aenv sh i e. - Gamma aenv + UID + -> Gamma aenv -> ArrayR (Array (sh, Int) e) -> IntegralType i -> IRFun2 PTX aenv (e -> e -> e) @@ -63,9 +64,9 @@ -> MIRDelayed PTX aenv (Array (sh, Int) e) -> MIRDelayed PTX aenv (Segments i) -> CodeGen PTX (IROpenAcc PTX aenv (Array (sh, Int) e)) -mkFoldSeg aenv repr intTp combine seed arr seg = - (+++) <$> mkFoldSegP_block aenv repr intTp combine seed arr seg - <*> mkFoldSegP_warp aenv repr intTp combine seed arr seg +mkFoldSeg uid aenv repr intTp combine seed arr seg = + (+++) <$> mkFoldSegP_block uid aenv repr intTp combine seed arr seg + <*> mkFoldSegP_warp uid aenv repr intTp combine seed arr seg -- This implementation assumes that the segments array represents the offset @@ -77,7 +78,8 @@ -- mkFoldSegP_block :: forall aenv sh i e. - Gamma aenv + UID + -> Gamma aenv -> ArrayR (Array (sh, Int) e) -> IntegralType i -> IRFun2 PTX aenv (e -> e -> e) @@ -85,8 +87,8 @@ -> MIRDelayed PTX aenv (Array (sh, Int) e) -> MIRDelayed PTX aenv (Segments i) -> CodeGen PTX (IROpenAcc PTX aenv (Array (sh, Int) e)) -mkFoldSegP_block aenv repr@(ArrayR shr tp) intTp combine mseed marr mseg = do - dev <- liftCodeGen $ gets ptxDeviceProperties +mkFoldSegP_block uid aenv repr@(ArrayR shr tp) intTp combine mseed marr mseg = do + dev <- liftCodeGen $ asks ptxDeviceProperties -- let (arrOut, paramOut) = mutableArray repr "out" @@ -95,14 +97,12 @@ paramEnv = envParam aenv -- config = launchConfig dev (CUDA.decWarp dev) dsmem const [|| const ||] - dsmem n = warps * (1 + per_warp) * bytes + dsmem n = sharedMemorySizeAdd tp warps 0 where ws = CUDA.warpSize dev warps = n `P.quot` ws - per_warp = ws + ws `P.quot` 2 - bytes = bytesElt tp -- - makeOpenAccWith config "foldSeg_block" (paramOut ++ paramIn ++ paramSeg ++ paramEnv) $ do + makeOpenAccWith config uid "foldSeg_block" (paramOut ++ paramIn ++ paramSeg ++ paramEnv) $ do -- We use a dynamically scheduled work queue in order to evenly distribute -- the uneven workload, due to the variable length of each segment, over the @@ -137,7 +137,7 @@ start <- return (liftInt 0) end <- shapeSize shr (irArrayShape arrOut) - imapFromTo start end $ \s -> do + Fold.imapFromTo start end $ \s -> do -- The first two threads of the block determine the indices of the -- segments array that we will reduce between and distribute those values @@ -205,8 +205,8 @@ v0 <- A.sub numType sup inf v0' <- i32 v0 r0 <- if (tp, A.gte singleType v0 bd) - then reduceBlockSMem dev tp combine Nothing x0 - else reduceBlockSMem dev tp combine (Just v0') x0 + then Fold.reduceBlock dev tp combine Nothing x0 + else Fold.reduceBlock dev tp combine (Just v0') x0 -- Step 2: keep walking over the input nxt <- A.add numType inf bd @@ -222,7 +222,7 @@ -- can avoid bounds checks. then do x <- app1 (delayedLinearIndex arrIn) i' - y <- reduceBlockSMem dev tp combine Nothing x + y <- Fold.reduceBlock dev tp combine Nothing x return y -- Not all threads are valid. Note that we still @@ -241,7 +241,7 @@ return $ go tp z <- i32 v' - y <- reduceBlockSMem dev tp combine (Just z) x + y <- Fold.reduceBlock dev tp combine (Just z) x return y -- first thread incorporates the result from the previous @@ -273,7 +273,8 @@ -- mkFoldSegP_warp :: forall aenv sh i e. - Gamma aenv + UID + -> Gamma aenv -> ArrayR (Array (sh, Int) e) -> IntegralType i -> IRFun2 PTX aenv (e -> e -> e) @@ -281,8 +282,8 @@ -> MIRDelayed PTX aenv (Array (sh, Int) e) -> MIRDelayed PTX aenv (Segments i) -> CodeGen PTX (IROpenAcc PTX aenv (Array (sh, Int) e)) -mkFoldSegP_warp aenv repr@(ArrayR shr tp) intTp combine mseed marr mseg = do - dev <- liftCodeGen $ gets ptxDeviceProperties +mkFoldSegP_warp uid aenv repr@(ArrayR shr tp) intTp combine mseed marr mseg = do + dev <- liftCodeGen $ asks ptxDeviceProperties -- let (arrOut, paramOut) = mutableArray repr "out" @@ -291,21 +292,17 @@ paramEnv = envParam aenv -- config = launchConfig dev (CUDA.decWarp dev) dsmem grid gridQ - dsmem n = warps * per_warp_bytes - where - warps = (n + ws - 1) `P.quot` ws + where dsmem _n = 0 -- grid n m = multipleOf n (m `P.quot` ws) gridQ = [|| \n m -> $$multipleOfQ n (m `P.quot` ws) ||] -- - per_warp_bytes = (per_warp_elems * bytesElt tp) `P.max` (2 * bytesElt tp) - per_warp_elems = ws + (ws `P.quot` 2) ws = CUDA.warpSize dev int32 :: Integral a => a -> Operands Int32 int32 = liftInt32 . P.fromIntegral -- - makeOpenAccWith config "foldSeg_warp" (paramOut ++ paramIn ++ paramSeg ++ paramEnv) $ do + makeOpenAccWith config uid "foldSeg_warp" (paramOut ++ paramIn ++ paramSeg ++ paramEnv) $ do -- Each warp works independently. -- Determine the ID of this warp within the thread block. @@ -322,33 +319,6 @@ b <- A.add numType wid a return b - -- All threads in the warp need to know what the start and end indices of - -- this segment are in order to participate in the reduction. We use - -- variables in __shared__ memory to communicate these values between - -- threads. Furthermore, by using a 2-element array, we can have the first - -- two threads of the warp read the start and end indices as a single - -- coalesced read, as these elements will be adjacent in the segment-offset - -- array. - -- - -- Note that this is aliased with the memory used to communicate reduction - -- values within the warp. - -- - lim <- do - a <- A.mul numType wid (int32 per_warp_bytes) - b <- dynamicSharedMem (TupRsingle scalarTypeInt) TypeInt32 (liftInt32 2) a - return b - - -- Allocate (1.5 * warpSize) elements of shared memory for each warp to - -- communicate reduction values. - -- - -- Note that this is aliased with the memory used to communicate the start - -- and end indices of this segment. - -- - smem <- do - a <- A.mul numType wid (int32 per_warp_bytes) - b <- dynamicSharedMem tp TypeInt32 (int32 per_warp_elems) a - return b - -- Compute the number of segments and size of the innermost dimension. These -- are required if we are reducing a rank-2 or higher array, to properly -- compute the start and end indices of the portion of the array this warp @@ -374,21 +344,26 @@ -- array that we will reduce between and distribute those values to the -- other threads in the warp lane <- laneId - when (A.lt singleType lane (liftInt32 2)) $ do - a <- case shr of - ShapeRsnoc ShapeRz -> return s - _ -> A.rem integralType s ss - b <- A.add numType a =<< int lane - c <- app1 (delayedLinearIndex arrSeg) b - writeArray TypeInt32 lim lane =<< A.fromIntegral intTp numType c + idx <- if (TupRsingle scalarTypeInt, A.lt singleType lane (liftInt32 2)) + then do + a <- case shr of + ShapeRsnoc ShapeRz -> return s + _ -> A.rem integralType s ss + b <- A.add numType a =<< int lane + c <- app1 (delayedLinearIndex arrSeg) b + d <- A.fromIntegral intTp numType c + return d + else + return (ir integralType (undef scalarType)) __syncwarp -- Determine the index range of the input array we will reduce over. -- Necessary for multidimensional segmented reduction. (inf,sup) <- do - u <- readArray TypeInt32 lim (liftInt32 0) - v <- readArray TypeInt32 lim (liftInt32 1) + u <- __shfl_idx (TupRsingle scalarTypeInt) idx (liftWord32 0) + + v <- __shfl_idx (TupRsingle scalarTypeInt) idx (liftWord32 1) A.unpair <$> case shr of ShapeRsnoc ShapeRz -> return (A.pair u v) _ -> do q <- A.quot integralType s ss @@ -429,8 +404,8 @@ v0 <- A.sub numType sup inf v0' <- i32 v0 r0 <- if (tp, A.gte singleType v0 (liftInt ws)) - then reduceWarpSMem dev tp combine smem Nothing x0 - else reduceWarpSMem dev tp combine smem (Just v0') x0 + then reduceWarp dev tp combine Nothing x0 + else reduceWarp dev tp combine (Just v0') x0 -- Step 2: Keep walking over the rest of the segment nx <- A.add numType inf (liftInt ws) @@ -445,7 +420,7 @@ then do -- All lanes are in bounds, so avoid bounds checks x <- app1 (delayedLinearIndex arrIn) i' - y <- reduceWarpSMem dev tp combine smem Nothing x + y <- reduceWarp dev tp combine Nothing x return y else do @@ -460,7 +435,7 @@ return $ go tp z <- i32 v' - y <- reduceWarpSMem dev tp combine smem (Just z) x + y <- reduceWarp dev tp combine (Just z) x return y -- The first lane incorporates the result from the previous @@ -488,4 +463,14 @@ int :: IsIntegral i => Operands i -> CodeGen PTX (Operands Int) int = A.fromIntegral integralType numType + +reduceWarp + :: forall aenv e. + DeviceProperties -- ^ properties of the target device + -> TypeR e + -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function + -> Maybe (Operands Int32) -- ^ number of items that will be reduced by this warp, otherwise all lanes are valid + -> Operands e -- ^ calling thread's input element + -> CodeGen PTX (Operands e) -- ^ warp-wide reduction using the specified operator (lane 0 only) +reduceWarp dev t c = Fold.reduceWarp dev t c diff --git a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Generate.hs b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Generate.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Generate.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Generate.hs @@ -27,6 +27,7 @@ import Data.Array.Accelerate.LLVM.CodeGen.Exp import Data.Array.Accelerate.LLVM.CodeGen.Monad import Data.Array.Accelerate.LLVM.CodeGen.Sugar +import Data.Array.Accelerate.LLVM.Compile.Cache import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base import Data.Array.Accelerate.LLVM.PTX.CodeGen.Loop @@ -37,16 +38,17 @@ -- processes multiple adjacent elements. -- mkGenerate - :: Gamma aenv + :: UID + -> Gamma aenv -> ArrayR (Array sh e) -> IRFun1 PTX aenv (sh -> e) -> CodeGen PTX (IROpenAcc PTX aenv (Array sh e)) -mkGenerate aenv repr@(ArrayR shr _) apply = +mkGenerate uid aenv repr@(ArrayR shr _) apply = let (arrOut, paramOut) = mutableArray repr "out" paramEnv = envParam aenv in - makeOpenAcc "generate" (paramOut ++ paramEnv) $ do + makeOpenAcc uid "generate" (paramOut ++ paramEnv) $ do start <- return (liftInt 0) end <- shapeSize shr (irArrayShape arrOut) diff --git a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Map.hs b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Map.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Map.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Map.hs @@ -27,6 +27,7 @@ import Data.Array.Accelerate.LLVM.CodeGen.Exp import Data.Array.Accelerate.LLVM.CodeGen.Monad import Data.Array.Accelerate.LLVM.CodeGen.Sugar +import Data.Array.Accelerate.LLVM.Compile.Cache import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base import Data.Array.Accelerate.LLVM.PTX.CodeGen.Loop @@ -36,18 +37,19 @@ -- Apply a unary function to each element of an array. Each thread processes -- multiple elements, striding the array by the grid size. -- -mkMap :: Gamma aenv +mkMap :: UID + -> Gamma aenv -> ArrayR (Array sh a) -> TypeR b -> IRFun1 PTX aenv (a -> b) -> CodeGen PTX (IROpenAcc PTX aenv (Array sh b)) -mkMap aenv repr@(ArrayR shr _) tp' apply = +mkMap uid aenv repr@(ArrayR shr _) tp' apply = let (arrOut, paramOut) = mutableArray (ArrayR shr tp') "out" (arrIn, paramIn) = mutableArray repr "in" paramEnv = envParam aenv in - makeOpenAcc "map" (paramOut ++ paramIn ++ paramEnv) $ do + makeOpenAcc uid "map" (paramOut ++ paramIn ++ paramEnv) $ do start <- return (liftInt 0) end <- shapeSize shr (irArrayShape arrIn) diff --git a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Permute.hs b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Permute.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Permute.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Permute.hs @@ -1,4 +1,3 @@ -{-# LANGUAGE CPP #-} {-# LANGUAGE GADTs #-} {-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE RecordWildCards #-} @@ -39,23 +38,24 @@ import Data.Array.Accelerate.LLVM.CodeGen.Permute import Data.Array.Accelerate.LLVM.CodeGen.Ptr import Data.Array.Accelerate.LLVM.CodeGen.Sugar +import Data.Array.Accelerate.LLVM.Compile.Cache import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base import Data.Array.Accelerate.LLVM.PTX.CodeGen.Loop import Data.Array.Accelerate.LLVM.PTX.Target -import LLVM.AST.Type.AddrSpace import LLVM.AST.Type.Instruction import LLVM.AST.Type.Instruction.Atomic import LLVM.AST.Type.Instruction.RMW as RMW import LLVM.AST.Type.Instruction.Volatile +import LLVM.AST.Type.GetElementPtr import LLVM.AST.Type.Operand import LLVM.AST.Type.Representation import Foreign.CUDA.Analysis import Control.Monad ( void ) -import Control.Monad.State ( gets ) +import Control.Monad.Reader ( asks ) import Prelude @@ -79,17 +79,18 @@ -- mkPermute :: HasCallStack - => Gamma aenv + => UID + -> Gamma aenv -> ArrayR (Array sh e) -> ShapeR sh' -> IRPermuteFun PTX aenv (e -> e -> e) -> IRFun1 PTX aenv (sh -> PrimMaybe sh') -> MIRDelayed PTX aenv (Array sh e) -> CodeGen PTX (IROpenAcc PTX aenv (Array sh' e)) -mkPermute aenv repr shr' IRPermuteFun{..} project arr = +mkPermute uid aenv repr shr' IRPermuteFun{..} project arr = case atomicRMW of - Just (rmw, f) -> mkPermute_rmw aenv repr shr' rmw f project arr - _ -> mkPermute_mutex aenv repr shr' combine project arr + Just (rmw, f) -> mkPermute_rmw uid aenv repr shr' rmw f project arr + _ -> mkPermute_mutex uid aenv repr shr' combine project arr -- Parallel forward permutation function which uses atomic instructions to @@ -115,7 +116,8 @@ -- mkPermute_rmw :: HasCallStack - => Gamma aenv + => UID + -> Gamma aenv -> ArrayR (Array sh e) -> ShapeR sh' -> RMWOperation @@ -123,8 +125,8 @@ -> IRFun1 PTX aenv (sh -> PrimMaybe sh') -> MIRDelayed PTX aenv (Array sh e) -> CodeGen PTX (IROpenAcc PTX aenv (Array sh' e)) -mkPermute_rmw aenv (ArrayR shr tp) shr' rmw update project marr = do - dev <- liftCodeGen $ gets ptxDeviceProperties +mkPermute_rmw uid aenv (ArrayR shr tp) shr' rmw update project marr = do + dev <- liftCodeGen $ asks ptxDeviceProperties -- let outR = ArrayR shr' tp @@ -139,7 +141,7 @@ compute60 = Compute 6 0 compute70 = Compute 7 0 -- - makeOpenAcc "permute_rmw" (paramOut ++ paramIn ++ paramEnv) $ do + makeOpenAcc uid "permute_rmw" (paramOut ++ paramIn ++ paramEnv) $ do shIn <- delayedExtent arrIn end <- shapeSize shr shIn @@ -161,7 +163,7 @@ _ | TupRsingle (SingleScalarType s) <- tp , adata <- irArrayData arrOut -> do - addr <- instr' $ GetElementPtr (asPtr defaultAddrSpace (op s adata)) [op integralType j] + addr <- instr' $ GetElementPtr $ GEP1 (SingleScalarType s) (asPtr defaultAddrSpace (op s adata)) (op integralType j) -- let rmw_integral :: IntegralType t -> Operand (Ptr t) -> Operand t -> CodeGen PTX () @@ -215,14 +217,15 @@ -- a mutex before updating the value at that location. -- mkPermute_mutex - :: Gamma aenv + :: UID + -> Gamma aenv -> ArrayR (Array sh e) -> ShapeR sh' -> IRFun2 PTX aenv (e -> e -> e) -> IRFun1 PTX aenv (sh -> PrimMaybe sh') -> MIRDelayed PTX aenv (Array sh e) -> CodeGen PTX (IROpenAcc PTX aenv (Array sh' e)) -mkPermute_mutex aenv (ArrayR shr tp) shr' combine project marr = +mkPermute_mutex uid aenv (ArrayR shr tp) shr' combine project marr = let outR = ArrayR shr' tp lockR = ArrayR (ShapeRsnoc ShapeRz) (TupRsingle scalarTypeWord32) @@ -232,7 +235,7 @@ paramEnv = envParam aenv start = liftInt 0 in - makeOpenAcc "permute_mutex" (paramOut ++ paramLock ++ paramIn ++ paramEnv) $ do + makeOpenAcc uid "permute_mutex" (paramOut ++ paramLock ++ paramIn ++ paramEnv) $ do shIn <- delayedExtent arrIn end <- shapeSize shr shIn @@ -264,7 +267,7 @@ -> CodeGen PTX a -> CodeGen PTX a atomically barriers i action = do - dev <- liftCodeGen $ gets ptxDeviceProperties + dev <- liftCodeGen $ asks ptxDeviceProperties if computeCapability dev >= Compute 7 0 then atomically_thread barriers i action else atomically_warp barriers i action @@ -305,7 +308,7 @@ exit <- newBlock "spinlock.exit" ns <- fresh i32 - addr <- instr' $ GetElementPtr (asPtr defaultAddrSpace (op integralType (irArrayData barriers))) [op integralType i] + addr <- instr' $ GetElementPtr $ GEP1 scalarType (asPtr defaultAddrSpace (op integralType (irArrayData barriers))) (op integralType i) top <- br entry -- Loop until this thread has completed its critical section. If the slot @@ -403,7 +406,7 @@ end <- newBlock "spinlock.critical-end" exit <- newBlock "spinlock.exit" - addr <- instr' $ GetElementPtr (asPtr defaultAddrSpace (op integralType (irArrayData barriers))) [op integralType i] + addr <- instr' $ GetElementPtr $ GEP1 scalarType (asPtr defaultAddrSpace (op integralType (irArrayData barriers))) (op integralType i) _ <- br entry -- Loop until this thread has completed its critical section. If the slot was @@ -425,7 +428,7 @@ -- incoming edge the thread arrived at this block from determines whether they -- have completed their critical section. setBlock end - res <- freshName + res <- freshLocalName done <- phi1 end res [(boolean True, yes), (boolean False, no)] __syncthreads diff --git a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs @@ -26,7 +26,6 @@ import Data.Array.Accelerate.AST ( Direction(..) ) import Data.Array.Accelerate.Representation.Array -import Data.Array.Accelerate.Representation.Elt import Data.Array.Accelerate.Representation.Shape import Data.Array.Accelerate.Representation.Type @@ -40,6 +39,7 @@ import Data.Array.Accelerate.LLVM.CodeGen.Loop import Data.Array.Accelerate.LLVM.CodeGen.Monad import Data.Array.Accelerate.LLVM.CodeGen.Sugar +import Data.Array.Accelerate.LLVM.Compile.Cache import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base import Data.Array.Accelerate.LLVM.PTX.CodeGen.Generate @@ -51,7 +51,7 @@ import Control.Applicative import Control.Monad ( (>=>), void ) -import Control.Monad.State ( gets ) +import Control.Monad.Reader ( asks ) import Data.String ( fromString ) import Data.Coerce as Safe import Data.Bits as P @@ -69,28 +69,30 @@ -- mkScan :: forall aenv sh e. - Gamma aenv + UID + -> Gamma aenv -> ArrayR (Array (sh, Int) e) -> Direction -> IRFun2 PTX aenv (e -> e -> e) -> Maybe (IRExp PTX aenv e) -> MIRDelayed PTX aenv (Array (sh, Int) e) -> CodeGen PTX (IROpenAcc PTX aenv (Array (sh, Int) e)) -mkScan aenv repr dir combine seed arr +mkScan uid aenv repr dir combine seed arr = foldr1 (+++) <$> sequence (codeScan ++ codeFill) where codeScan = case repr of - ArrayR (ShapeRsnoc ShapeRz) tp -> [ mkScanAllP1 dir aenv tp combine seed arr - , mkScanAllP2 dir aenv tp combine - , mkScanAllP3 dir aenv tp combine seed + ArrayR (ShapeRsnoc ShapeRz) tp -> [ mkScanAllP1 dir uid aenv tp combine seed arr + , mkScanAllP2 dir uid aenv tp combine + , mkScanAllP3 dir uid aenv tp combine seed ] - _ -> [ mkScanDim dir aenv repr combine seed arr + _ -> [ mkScanDim dir uid aenv repr combine seed arr ] codeFill = case seed of - Just s -> [ mkScanFill aenv repr s ] + Just s -> [ mkScanFill uid aenv repr s ] Nothing -> [] + -- Variant of 'scanl' where the final result is returned in a separate array. -- -- > scanr' (+) 10 (use $ fromList (Z :. 10) [0..]) @@ -101,24 +103,25 @@ -- mkScan' :: forall aenv sh e. - Gamma aenv + UID + -> Gamma aenv -> ArrayR (Array (sh, Int) e) -> Direction -> IRFun2 PTX aenv (e -> e -> e) -> IRExp PTX aenv e -> MIRDelayed PTX aenv (Array (sh, Int) e) -> CodeGen PTX (IROpenAcc PTX aenv (Array (sh, Int) e, Array sh e)) -mkScan' aenv repr dir combine seed arr +mkScan' uid aenv repr dir combine seed arr | ArrayR (ShapeRsnoc ShapeRz) tp <- repr - = foldr1 (+++) <$> sequence [ mkScan'AllP1 dir aenv tp combine seed arr - , mkScan'AllP2 dir aenv tp combine - , mkScan'AllP3 dir aenv tp combine - , mkScan'Fill aenv repr seed + = foldr1 (+++) <$> sequence [ mkScan'AllP1 dir uid aenv tp combine seed arr + , mkScan'AllP2 dir uid aenv tp combine + , mkScan'AllP3 dir uid aenv tp combine + , mkScan'Fill uid aenv repr seed ] -- | otherwise - = (+++) <$> mkScan'Dim dir aenv repr combine seed arr - <*> mkScan'Fill aenv repr seed + = (+++) <$> mkScan'Dim dir uid aenv repr combine seed arr + <*> mkScan'Fill uid aenv repr seed -- Device wide scans @@ -138,14 +141,15 @@ mkScanAllP1 :: forall aenv e. Direction + -> UID -> Gamma aenv -- ^ array environment -> TypeR e -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function -> MIRExp PTX aenv e -- ^ seed element, if this is an exclusive scan -> MIRDelayed PTX aenv (Vector e) -- ^ input data -> CodeGen PTX (IROpenAcc PTX aenv (Vector e)) -mkScanAllP1 dir aenv tp combine mseed marr = do - dev <- liftCodeGen $ gets ptxDeviceProperties +mkScanAllP1 dir uid aenv tp combine mseed marr = do + dev <- liftCodeGen $ asks ptxDeviceProperties -- let (arrOut, paramOut) = mutableArray (ArrayR dim1 tp) "out" @@ -154,15 +158,9 @@ end = indexHead (irArrayShape arrTmp) paramEnv = envParam aenv -- - config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] - smem n = warps * (1 + per_warp) * bytes - where - ws = CUDA.warpSize dev - warps = n `P.quot` ws - per_warp = ws + ws `P.quot` 2 - bytes = bytesElt tp + config = launchConfig dev (CUDA.incWarp dev) (scanSMemSize dev tp) const [|| const ||] -- - makeOpenAccWith config "scanP1" (paramTmp ++ paramOut ++ paramIn ++ paramEnv) $ do + makeOpenAccWith config uid "scanP1" (paramTmp ++ paramOut ++ paramIn ++ paramEnv) $ do -- Size of the input array sz <- indexHead <$> delayedExtent arrIn @@ -180,8 +178,15 @@ s0 <- int bid -- iterating over thread-block-wide segments + -- Note that 'end' is a multiple of the gd', and the control flow is thus uniform in the loop. + -- This is set in scanAllOp in Data.Array.Accelerate.LLVM.PTX.Execute. + -- Hence we can run __syncthreads safely. imapFromStepTo s0 gd' end $ \chunk -> do + -- Make sure all threads have finished previous iterations, + -- so we can reuse (and overwrite) shared memory. + __syncthreads + bd <- blockDim bd' <- int bd inf <- A.mul numType chunk bd' @@ -228,8 +233,8 @@ n <- A.sub numType sz inf n' <- i32 n x2 <- if (tp, A.gte singleType n bd') - then scanBlockSMem dir dev tp combine Nothing x1 - else scanBlockSMem dir dev tp combine (Just n') x1 + then scanBlock dir dev tp combine Nothing x1 + else scanBlock dir dev tp combine (Just n') x1 -- Write this thread's scan result to memory writeArray TypeInt arrOut j0 x2 @@ -258,12 +263,13 @@ mkScanAllP2 :: forall aenv e. Direction + -> UID -> Gamma aenv -- ^ array environment -> TypeR e -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function -> CodeGen PTX (IROpenAcc PTX aenv (Vector e)) -mkScanAllP2 dir aenv tp combine = do - dev <- liftCodeGen $ gets ptxDeviceProperties +mkScanAllP2 dir uid aenv tp combine = do + dev <- liftCodeGen $ asks ptxDeviceProperties -- let (arrTmp, paramTmp) = mutableArray (ArrayR dim1 tp) "tmp" @@ -271,17 +277,11 @@ start = liftInt 0 end = indexHead (irArrayShape arrTmp) -- - config = launchConfig dev (CUDA.incWarp dev) smem grid gridQ + config = launchConfig dev (CUDA.incWarp dev) (scanSMemSize dev tp) grid gridQ grid _ _ = 1 gridQ = [|| \_ _ -> 1 ||] - smem n = warps * (1 + per_warp) * bytes - where - ws = CUDA.warpSize dev - warps = n `P.quot` ws - per_warp = ws + ws `P.quot` 2 - bytes = bytesElt tp -- - makeOpenAccWith config "scanP2" (paramTmp ++ paramEnv) $ do + makeOpenAccWith config uid "scanP2" (paramTmp ++ paramEnv) $ do -- The first and last threads of the block need to communicate the -- block-wide aggregate as a carry-in value across iterations. @@ -310,11 +310,10 @@ LeftToRight -> A.lt singleType i end RightToLeft -> A.gte singleType i start - when (valid i0) $ do - - -- wait for the carry-in value to be updated - __syncthreads + -- wait for the carry-in value to be updated + __syncthreads + when (valid i0) $ do x0 <- readArray TypeInt arrTmp i0 x1 <- if (tp, A.gt singleType offset (liftInt 0) `land'` A.eq singleType tid (liftInt32 0)) then do @@ -328,8 +327,8 @@ n <- A.sub numType end offset n' <- i32 n x2 <- if (tp, A.gte singleType n bd') - then scanBlockSMem dir dev tp combine Nothing x1 - else scanBlockSMem dir dev tp combine (Just n') x1 + then scanBlock dir dev tp combine Nothing x1 + else scanBlock dir dev tp combine (Just n') x1 -- Update the temporary array with this thread's result writeArray TypeInt arrTmp i0 x2 @@ -351,13 +350,14 @@ mkScanAllP3 :: forall aenv e. Direction + -> UID -> Gamma aenv -- ^ array environment -> TypeR e -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function -> MIRExp PTX aenv e -- ^ seed element, if this is an exclusive scan -> CodeGen PTX (IROpenAcc PTX aenv (Vector e)) -mkScanAllP3 dir aenv tp combine mseed = do - dev <- liftCodeGen $ gets ptxDeviceProperties +mkScanAllP3 dir uid aenv tp combine mseed = do + dev <- liftCodeGen $ asks ptxDeviceProperties -- let (arrOut, paramOut) = mutableArray (ArrayR dim1 tp) "out" @@ -369,7 +369,7 @@ -- config = launchConfig dev (CUDA.incWarp dev) (const 0) const [|| const ||] -- - makeOpenAccWith config "scanP3" (paramTmp ++ paramOut ++ paramStride ++ paramEnv) $ do + makeOpenAccWith config uid "scanP3" (paramTmp ++ paramOut ++ paramStride ++ paramEnv) $ do sz <- return $ indexHead (irArrayShape arrOut) tid <- int =<< threadIdx @@ -448,14 +448,15 @@ mkScan'AllP1 :: forall aenv e. Direction + -> UID -> Gamma aenv -> TypeR e -> IRFun2 PTX aenv (e -> e -> e) -> IRExp PTX aenv e -> MIRDelayed PTX aenv (Vector e) -> CodeGen PTX (IROpenAcc PTX aenv (Vector e, Scalar e)) -mkScan'AllP1 dir aenv tp combine seed marr = do - dev <- liftCodeGen $ gets ptxDeviceProperties +mkScan'AllP1 dir uid aenv tp combine seed marr = do + dev <- liftCodeGen $ asks ptxDeviceProperties -- let (arrOut, paramOut) = mutableArray (ArrayR dim1 tp) "out" @@ -464,15 +465,9 @@ end = indexHead (irArrayShape arrTmp) paramEnv = envParam aenv -- - config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] - smem n = warps * (1 + per_warp) * bytes - where - ws = CUDA.warpSize dev - warps = n `P.quot` ws - per_warp = ws + ws `P.quot` 2 - bytes = bytesElt tp + config = launchConfig dev (CUDA.incWarp dev) (scanSMemSize dev tp) const [|| const ||] -- - makeOpenAccWith config "scanP1" (paramTmp ++ paramOut ++ paramIn ++ paramEnv) $ do + makeOpenAccWith config uid "scanP1" (paramTmp ++ paramOut ++ paramIn ++ paramEnv) $ do -- Size of the input array sz <- indexHead <$> delayedExtent arrIn @@ -483,8 +478,15 @@ gd <- int =<< gridDim -- iterate over thread-block wide segments + -- Note that 'end' is a multiple of the gd', and the control flow is thus uniform in the loop. + -- This is set in scan'AllOp in Data.Array.Accelerate.LLVM.PTX.Execute. + -- Hence we can run __syncthreads safely. imapFromStepTo bid gd end $ \seg -> do + -- Make sure all threads have finished previous iterations, + -- so we can reuse (and overwrite) shared memory. + __syncthreads + bd <- int =<< blockDim inf <- A.mul numType seg bd @@ -528,8 +530,8 @@ n <- A.sub numType sz inf n' <- i32 n x2 <- if (tp, A.gte singleType n bd) - then scanBlockSMem dir dev tp combine Nothing x1 - else scanBlockSMem dir dev tp combine (Just n') x1 + then scanBlock dir dev tp combine Nothing x1 + else scanBlock dir dev tp combine (Just n') x1 -- Write this thread's scan result to memory. Recall that we had to make -- space for the initial element, so the very last thread does not store @@ -561,12 +563,13 @@ mkScan'AllP2 :: forall aenv e. Direction + -> UID -> Gamma aenv -> TypeR e -> IRFun2 PTX aenv (e -> e -> e) -> CodeGen PTX (IROpenAcc PTX aenv (Vector e, Scalar e)) -mkScan'AllP2 dir aenv tp combine = do - dev <- liftCodeGen $ gets ptxDeviceProperties +mkScan'AllP2 dir uid aenv tp combine = do + dev <- liftCodeGen $ asks ptxDeviceProperties -- let (arrTmp, paramTmp) = mutableArray (ArrayR dim1 tp) "tmp" @@ -575,17 +578,11 @@ start = liftInt 0 end = indexHead (irArrayShape arrTmp) -- - config = launchConfig dev (CUDA.incWarp dev) smem grid gridQ + config = launchConfig dev (CUDA.incWarp dev) (scanSMemSize dev tp) grid gridQ grid _ _ = 1 gridQ = [|| \_ _ -> 1 ||] - smem n = warps * (1 + per_warp) * bytes - where - ws = CUDA.warpSize dev - warps = n `P.quot` ws - per_warp = ws + ws `P.quot` 2 - bytes = bytesElt tp -- - makeOpenAccWith config "scanP2" (paramTmp ++ paramSum ++ paramEnv) $ do + makeOpenAccWith config uid "scanP2" (paramTmp ++ paramSum ++ paramEnv) $ do -- The first and last threads of the block need to communicate the -- block-wide aggregate as a carry-in value across iterations. @@ -616,12 +613,7 @@ x0 <- if (tp, valid i0) then readArray TypeInt arrTmp i0 else - let go :: TypeR a -> Operands a - go TupRunit = OP_Unit - go (TupRpair a b) = OP_Pair (go a) (go b) - go (TupRsingle t) = ir t (undef t) - in - return $ go tp + return $ tupUndef tp x1 <- if (tp, A.gt singleType offset (liftInt 0) `A.land'` A.eq singleType tid (liftInt32 0)) then do @@ -635,8 +627,8 @@ n <- A.sub numType end offset n' <- i32 n x2 <- if (tp, A.gte singleType n bd) - then scanBlockSMem dir dev tp combine Nothing x1 - else scanBlockSMem dir dev tp combine (Just n') x1 + then scanBlock dir dev tp combine Nothing x1 + else scanBlock dir dev tp combine (Just n') x1 -- Update the partial results array when (valid i0) $ @@ -668,12 +660,13 @@ mkScan'AllP3 :: forall aenv e. Direction + -> UID -> Gamma aenv -- ^ array environment -> TypeR e -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function -> CodeGen PTX (IROpenAcc PTX aenv (Vector e, Scalar e)) -mkScan'AllP3 dir aenv tp combine = do - dev <- liftCodeGen $ gets ptxDeviceProperties +mkScan'AllP3 dir uid aenv tp combine = do + dev <- liftCodeGen $ asks ptxDeviceProperties -- let (arrOut, paramOut) = mutableArray (ArrayR dim1 tp) "out" @@ -685,7 +678,7 @@ -- config = launchConfig dev (CUDA.incWarp dev) (const 0) const [|| const ||] -- - makeOpenAccWith config "scanP3" (paramTmp ++ paramOut ++ paramStride ++ paramEnv) $ do + makeOpenAccWith config uid "scanP3" (paramTmp ++ paramOut ++ paramStride ++ paramEnv) $ do sz <- return $ indexHead (irArrayShape arrOut) tid <- int =<< threadIdx @@ -753,29 +746,24 @@ mkScanDim :: forall aenv sh e. Direction + -> UID -> Gamma aenv -- ^ array environment -> ArrayR (Array (sh, Int) e) -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function -> MIRExp PTX aenv e -- ^ seed element, if this is an exclusive scan -> MIRDelayed PTX aenv (Array (sh, Int) e) -- ^ input data -> CodeGen PTX (IROpenAcc PTX aenv (Array (sh, Int) e)) -mkScanDim dir aenv repr@(ArrayR (ShapeRsnoc shr) tp) combine mseed marr = do - dev <- liftCodeGen $ gets ptxDeviceProperties +mkScanDim dir uid aenv repr@(ArrayR (ShapeRsnoc shr) tp) combine mseed marr = do + dev <- liftCodeGen $ asks ptxDeviceProperties -- let (arrOut, paramOut) = mutableArray repr "out" (arrIn, paramIn) = delayedArray "in" marr paramEnv = envParam aenv -- - config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] - smem n = warps * (1 + per_warp) * bytes - where - ws = CUDA.warpSize dev - warps = n `P.quot` ws - per_warp = ws + ws `P.quot` 2 - bytes = bytesElt tp + config = launchConfig dev (CUDA.incWarp dev) (scanSMemSize dev tp) const [|| const ||] -- - makeOpenAccWith config "scan" (paramOut ++ paramIn ++ paramEnv) $ do + makeOpenAccWith config uid "scan" (paramOut ++ paramIn ++ paramEnv) $ do -- The first and last threads of the block need to communicate the -- block-wide aggregate as a carry-in value across iterations. @@ -796,8 +784,18 @@ gd <- int =<< gridDim end <- shapeSize shr (indexTail (irArrayShape arrOut)) + -- Iterate over the outer dimensions (all but the innermost dimension). + -- Within this loop we perform a scan over a row (innermost dimension) of + -- the input. + -- + -- Since 'bid', 'gd' and 'end' are uniform, the control flow within this + -- loop is also uniform. We can thus perform __syncthreads in the loop. imapFromStepTo bid gd end $ \seg -> do + -- Make sure all threads have finished previous iterations, + -- so we can reuse (and overwrite) shared memory. + __syncthreads + -- Index this thread reads from tid <- threadIdx tid' <- int tid @@ -856,12 +854,19 @@ return $ A.trip sz i0 j1 Nothing -> do + -- We cannot call scanBlock under non-uniform control flow. + -- Instead, we conditionally read the input, and then + -- unconditionally call scanBlock. + x0 <- if (tp, A.lt singleType tid' sz) + then app1 (delayedLinearIndex arrIn) i0 + else return $ tupUndef tp + n' <- i32 sz + + r0 <- if (tp, A.gte singleType sz bd') + then scanBlock dir dev tp combine Nothing x0 + else scanBlock dir dev tp combine (Just n') x0 + when (A.lt singleType tid' sz) $ do - n' <- i32 sz - x0 <- app1 (delayedLinearIndex arrIn) i0 - r0 <- if (tp, A.gte singleType sz bd') - then scanBlockSMem dir dev tp combine Nothing x0 - else scanBlockSMem dir dev tp combine (Just n') x0 writeArray TypeInt arrOut j0 r0 ll <- A.sub numType bd (liftInt32 1) @@ -874,6 +879,11 @@ return $ A.trip n1 i1 j1 -- Iterate over the remaining elements in this segment + -- The loop condition uses the first triple of the state, 'n'. + -- This variable is uniform (the same for all threads in the thread + -- block), since it is initialized as 'sz' or 'sz - bd', and lowered by + -- 'bd' each iteration. Hence the control flow in this loop is uniform, + -- and we can thus call __syncthreads within the loop. void $ while (TupRunit `TupRpair` TupRsingle scalarTypeInt `TupRpair` TupRsingle scalarTypeInt `TupRpair` TupRsingle scalarTypeInt) (\(A.fst3 -> n) -> A.gt singleType n (liftInt 0)) @@ -891,13 +901,7 @@ -- x <- if (tp, A.lt singleType tid' n) then app1 (delayedLinearIndex arrIn) i - else let - go :: TypeR a -> Operands a - go TupRunit = OP_Unit - go (TupRpair a b) = OP_Pair (go a) (go b) - go (TupRsingle t) = ir t (undef t) - in - return $ go tp + else return $ tupUndef tp -- Thread zero incorporates the carry-in element y <- if (tp, A.eq singleType tid (liftInt32 0)) @@ -912,8 +916,8 @@ -- Perform the scan and write the result to memory m <- i32 n z <- if (tp, A.gte singleType n bd') - then scanBlockSMem dir dev tp combine Nothing y - else scanBlockSMem dir dev tp combine (Just m) y + then scanBlock dir dev tp combine Nothing y + else scanBlock dir dev tp combine (Just m) y when (A.lt singleType tid' n) $ do writeArray TypeInt arrOut j z @@ -950,14 +954,15 @@ mkScan'Dim :: forall aenv sh e. Direction + -> UID -> Gamma aenv -- ^ array environment -> ArrayR (Array (sh, Int) e) -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function -> IRExp PTX aenv e -- ^ seed element -> MIRDelayed PTX aenv (Array (sh, Int) e) -- ^ input data -> CodeGen PTX (IROpenAcc PTX aenv (Array (sh, Int) e, Array sh e)) -mkScan'Dim dir aenv repr@(ArrayR (ShapeRsnoc shr) tp) combine seed marr = do - dev <- liftCodeGen $ gets ptxDeviceProperties +mkScan'Dim dir uid aenv repr@(ArrayR (ShapeRsnoc shr) tp) combine seed marr = do + dev <- liftCodeGen $ asks ptxDeviceProperties -- let (arrSum, paramSum) = mutableArray (reduceRank repr) "sum" @@ -965,15 +970,9 @@ (arrIn, paramIn) = delayedArray "in" marr paramEnv = envParam aenv -- - config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] - smem n = warps * (1 + per_warp) * bytes - where - ws = CUDA.warpSize dev - warps = n `P.quot` ws - per_warp = ws + ws `P.quot` 2 - bytes = bytesElt tp + config = launchConfig dev (CUDA.incWarp dev) (scanSMemSize dev tp) const [|| const ||] -- - makeOpenAccWith config "scan" (paramOut ++ paramSum ++ paramIn ++ paramEnv) $ do + makeOpenAccWith config uid "scan" (paramOut ++ paramSum ++ paramIn ++ paramEnv) $ do -- The first and last threads of the block need to communicate the -- block-wide aggregate as a carry-in value across iterations. @@ -1066,7 +1065,7 @@ RightToLeft -> app2 combine x c else return x - z <- scanBlockSMem dir dev tp combine Nothing y + z <- scanBlock dir dev tp combine Nothing y -- Write results to the output array. Note that if we -- align directly on the boundary of the array this is not @@ -1103,16 +1102,10 @@ return x return y else - let - go :: TypeR a -> Operands a - go TupRunit = OP_Unit - go (TupRpair a b) = OP_Pair (go a) (go b) - go (TupRsingle t) = ir t (undef t) - in - return $ go tp + return $ tupUndef tp l <- i32 n - y <- scanBlockSMem dir dev tp combine (Just l) x + y <- scanBlock dir dev tp combine (Just l) x m <- A.sub numType l (liftInt32 1) when (A.lt singleType tid m) $ writeArray TypeInt arrOut j y @@ -1140,21 +1133,28 @@ -- the seed element. -- mkScanFill - :: Gamma aenv + :: UID + -> Gamma aenv -> ArrayR (Array sh e) -> IRExp PTX aenv e -> CodeGen PTX (IROpenAcc PTX aenv (Array sh e)) -mkScanFill aenv repr seed = - mkGenerate aenv repr (IRFun1 (const seed)) +mkScanFill uid aenv repr seed = + mkGenerate uid aenv repr (IRFun1 (const seed)) mkScan'Fill - :: Gamma aenv + :: UID + -> Gamma aenv -> ArrayR (Array (sh, Int) e) -> IRExp PTX aenv e -> CodeGen PTX (IROpenAcc PTX aenv (Array (sh, Int) e, Array sh e)) -mkScan'Fill aenv repr seed = - Safe.coerce <$> mkGenerate aenv (reduceRank repr) (IRFun1 (const seed)) +mkScan'Fill uid aenv repr seed = + Safe.coerce <$> mkGenerate uid aenv (reduceRank repr) (IRFun1 (const seed)) +scanSMemSize :: DeviceProperties -> TypeR e -> Int -> Int +scanSMemSize dev tp n = sharedMemorySizeAdd tp warps 0 + where + ws = CUDA.warpSize dev + warps = n `P.quot` ws -- Block wide scan -- --------------- @@ -1166,17 +1166,9 @@ -- -- Example: https://github.com/NVlabs/cub/blob/1.5.4/cub/block/specializations/block_scan_warp_scans.cuh -- --- NOTE: [Synchronisation problems with SM_70 and greater] --- --- This operation uses thread synchronisation. When calling this operation, it --- is important that all active (that is, non-exited) threads of the thread --- block participate. It seems that sm_70+ (devices with independent thread --- scheduling) are stricter about the requirement that all non-existed threads --- participate in every barrier. --- --- See: https://github.com/AccelerateHS/accelerate/issues/436 --- -scanBlockSMem +-- Must be called under uniform control flow within a thread block +-- (as this function may use __syncthreads) +scanBlock :: forall aenv e. Direction -> DeviceProperties -- ^ properties of the target device @@ -1185,24 +1177,18 @@ -> Maybe (Operands Int32) -- ^ number of valid elements (may be less than block size) -> Operands e -- ^ calling thread's input element -> CodeGen PTX (Operands e) -scanBlockSMem dir dev tp combine nelem = warpScan >=> warpPrefix +scanBlock dir dev tp combine nelem = warpScan >=> warpPrefix where int32 :: Integral a => a -> Operands Int32 int32 = liftInt32 . P.fromIntegral -- Temporary storage required for each warp - warp_smem_elems = CUDA.warpSize dev + (CUDA.warpSize dev `P.quot` 2) - warp_smem_bytes = warp_smem_elems * bytesElt tp + -- warp_smem_elems = CUDA.warpSize dev + (CUDA.warpSize dev `P.quot` 2) + -- warp_smem_bytes = warp_smem_elems * bytesElt tp -- Step 1: Scan in every warp warpScan :: Operands e -> CodeGen PTX (Operands e) - warpScan input = do - -- Allocate (1.5 * warpSize) elements of shared memory for each warp - -- (individually addressable by each warp) - wid <- warpId - skip <- A.mul numType wid (int32 warp_smem_bytes) - smem <- dynamicSharedMem tp TypeInt32 (int32 warp_smem_elems) skip - scanWarpSMem dir dev tp combine smem input + warpScan = scanWarp dir dev tp combine -- Step 2: Collect the aggregate results of each warp to compute the prefix -- values for each warp and combine with the partial result to compute each @@ -1212,8 +1198,7 @@ -- Allocate #warps elements of shared memory bd <- blockDim warps <- A.quot integralType bd (int32 (CUDA.warpSize dev)) - skip <- A.mul numType warps (int32 warp_smem_bytes) - smem <- dynamicSharedMem tp TypeInt32 warps skip + smem <- dynamicSharedMem tp TypeInt32 warps (liftInt32 0) -- Share warp aggregates wid <- warpId @@ -1233,7 +1218,7 @@ -- their prefix value. We do this sequentially, but could also have -- warp 0 do it cooperatively if we limit thread block sizes to -- (warp size ^ 2). - steps <- case nelem of + steps <- case nelem of Nothing -> return wid Just n -> A.min singleType wid =<< A.quot integralType n (int32 (CUDA.warpSize dev)) @@ -1260,54 +1245,49 @@ -- -- Example: https://github.com/NVlabs/cub/blob/1.5.4/cub/warp/specializations/warp_scan_smem.cuh -- -scanWarpSMem +scanWarp :: forall aenv e. Direction -> DeviceProperties -- ^ properties of the target device -> TypeR e -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function - -> IRArray (Vector e) -- ^ temporary storage array in shared memory (1.5 x warp size elements) -> Operands e -- ^ calling thread's input element -> CodeGen PTX (Operands e) -scanWarpSMem dir dev tp combine smem = scan 0 +scanWarp dir dev tp combine = scan 0 where log2 :: Double -> Double log2 = P.logBase 2 -- Number of steps required to scan warp - steps = P.floor (log2 (P.fromIntegral (CUDA.warpSize dev))) - halfWarp = P.fromIntegral (CUDA.warpSize dev `P.quot` 2) + steps = P.floor (log2 (P.fromIntegral (CUDA.warpSize dev))) -- Unfold the scan as a recursive code generation function scan :: Int -> Operands e -> CodeGen PTX (Operands e) scan step x | step >= steps = return x | otherwise = do - let offset = liftInt32 (1 `P.shiftL` step) + let offset = 1 `P.shiftL` step -- share partial result through shared memory buffer + y <- __shfl_up tp x (liftWord32 offset) lane <- laneId - i <- A.add numType lane (liftInt32 halfWarp) - writeArray TypeInt32 smem i x - __syncwarp - -- update partial result if in range - x' <- if (tp, A.gte singleType lane offset) + x' <- if (tp, A.gte singleType lane (liftInt32 . P.fromIntegral $ offset)) then do - i' <- A.sub numType i offset -- lane + HALF_WARP - offset - x' <- readArray TypeInt32 smem i' case dir of - LeftToRight -> app2 combine x' x - RightToLeft -> app2 combine x x' + LeftToRight -> app2 combine y x + RightToLeft -> app2 combine x y else return x - __syncwarp - scan (step+1) x' +tupUndef :: TypeR a -> Operands a +tupUndef TupRunit = OP_Unit +tupUndef (TupRpair a b) = OP_Pair (tupUndef a) (tupUndef b) +tupUndef (TupRsingle t) = ir t (undef t) -- Utilities -- --------- diff --git a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Stencil.hs b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Stencil.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Stencil.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Stencil.hs @@ -35,12 +35,13 @@ import Data.Array.Accelerate.LLVM.CodeGen.Monad import Data.Array.Accelerate.LLVM.CodeGen.Stencil import Data.Array.Accelerate.LLVM.CodeGen.Sugar +import Data.Array.Accelerate.LLVM.Compile.Cache import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base import Data.Array.Accelerate.LLVM.PTX.CodeGen.Loop import Data.Array.Accelerate.LLVM.PTX.Target ( PTX ) -import qualified LLVM.AST.Global as LLVM +import qualified Data.Array.Accelerate.LLVM.Internal.LLVMPretty.AST as LP import Control.Monad @@ -56,24 +57,26 @@ -- * stencil_border: applies boundary condition check to each array access -- mkStencil1 - :: Gamma aenv + :: UID + -> Gamma aenv -> StencilR sh a stencil -> TypeR b -> IRFun1 PTX aenv (stencil -> b) -> IRBoundary PTX aenv (Array sh a) -> MIRDelayed PTX aenv (Array sh a) -> CodeGen PTX (IROpenAcc PTX aenv (Array sh b)) -mkStencil1 aenv stencil tp fun bnd marr = +mkStencil1 uid aenv stencil tp fun bnd marr = let repr = ArrayR shr tp (shr, halo) = stencilHalo stencil (arrIn, paramIn) = delayedArray "in" marr in - (+++) <$> mkInside aenv repr halo (IRFun1 $ app1 fun <=< stencilAccess stencil Nothing arrIn) paramIn - <*> mkBorder aenv repr (IRFun1 $ app1 fun <=< stencilAccess stencil (Just bnd) arrIn) paramIn + (+++) <$> mkInside uid aenv repr halo (IRFun1 $ app1 fun <=< stencilAccess stencil Nothing arrIn) paramIn + <*> mkBorder uid aenv repr (IRFun1 $ app1 fun <=< stencilAccess stencil (Just bnd) arrIn) paramIn mkStencil2 - :: Gamma aenv + :: UID + -> Gamma aenv -> StencilR sh a stencil1 -> StencilR sh b stencil2 -> TypeR c @@ -83,7 +86,7 @@ -> IRBoundary PTX aenv (Array sh b) -> MIRDelayed PTX aenv (Array sh b) -> CodeGen PTX (IROpenAcc PTX aenv (Array sh c)) -mkStencil2 aenv stencil1 stencil2 tp f bnd1 marr1 bnd2 marr2 = +mkStencil2 uid aenv stencil1 stencil2 tp f bnd1 marr1 bnd2 marr2 = let repr = ArrayR shr tp (arrIn1, paramIn1) = delayedArray "in1" marr1 @@ -103,27 +106,27 @@ (_, halo2) = stencilHalo stencil2 halo = union shr halo1 halo2 in - (+++) <$> mkInside aenv repr halo inside (paramIn1 ++ paramIn2) - <*> mkBorder aenv repr border (paramIn1 ++ paramIn2) + (+++) <$> mkInside uid aenv repr halo inside (paramIn1 ++ paramIn2) + <*> mkBorder uid aenv repr border (paramIn1 ++ paramIn2) mkInside - :: Gamma aenv + :: UID + -> Gamma aenv -> ArrayR (Array sh e) -> sh -> IRFun1 PTX aenv (sh -> e) - -> [LLVM.Parameter] + -> [LP.Typed LP.Ident] -> CodeGen PTX (IROpenAcc PTX aenv (Array sh e)) -mkInside aenv repr@(ArrayR shr _) halo apply paramIn = +mkInside uid aenv repr@(ArrayR shr _) halo apply paramIn = let (arrOut, paramOut) = mutableArray repr "out" paramInside = parameter (shapeType shr) "shInside" shInside = local (shapeType shr) "shInside" shOut = irArrayShape arrOut paramEnv = envParam aenv - -- in - makeOpenAcc "stencil_inside" (paramInside ++ paramOut ++ paramIn ++ paramEnv) $ do + makeOpenAcc uid "stencil_inside" (paramInside ++ paramOut ++ paramIn ++ paramEnv) $ do start <- return (liftInt 0) end <- shapeSize shr shInside @@ -142,12 +145,13 @@ mkBorder - :: Gamma aenv + :: UID + -> Gamma aenv -> ArrayR (Array sh e) -> IRFun1 PTX aenv (sh -> e) - -> [LLVM.Parameter] + -> [LP.Typed LP.Ident] -> CodeGen PTX (IROpenAcc PTX aenv (Array sh e)) -mkBorder aenv repr@(ArrayR shr _) apply paramIn = +mkBorder uid aenv repr@(ArrayR shr _) apply paramIn = let (arrOut, paramOut) = mutableArray repr "out" paramFrom = parameter (shapeType shr) "shFrom" @@ -156,9 +160,8 @@ shInside = local (shapeType shr) "shInside" shOut = irArrayShape arrOut paramEnv = envParam aenv - -- in - makeOpenAcc "stencil_border" (paramFrom ++ paramInside ++ paramOut ++ paramIn ++ paramEnv) $ do + makeOpenAcc uid "stencil_border" (paramFrom ++ paramInside ++ paramOut ++ paramIn ++ paramEnv) $ do start <- return (liftInt 0) end <- shapeSize shr shInside diff --git a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Transform.hs b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Transform.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Transform.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Transform.hs @@ -26,6 +26,7 @@ import Data.Array.Accelerate.LLVM.CodeGen.Exp import Data.Array.Accelerate.LLVM.CodeGen.Monad import Data.Array.Accelerate.LLVM.CodeGen.Sugar +import Data.Array.Accelerate.LLVM.Compile.Cache import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base import Data.Array.Accelerate.LLVM.PTX.CodeGen.Loop @@ -36,19 +37,20 @@ -- multiple elements, striding the array by the grid size. -- mkTransform - :: Gamma aenv + :: UID + -> Gamma aenv -> ArrayR (Array sh a) -> ArrayR (Array sh' b) -> IRFun1 PTX aenv (sh' -> sh) -> IRFun1 PTX aenv (a -> b) -> CodeGen PTX (IROpenAcc PTX aenv (Array sh' b)) -mkTransform aenv repr@(ArrayR shr _) repr'@(ArrayR shr' _) p f = +mkTransform uid aenv repr@(ArrayR shr _) repr'@(ArrayR shr' _) p f = let (arrOut, paramOut) = mutableArray repr' "out" (arrIn, paramIn) = mutableArray repr "in" paramEnv = envParam aenv in - makeOpenAcc "transform" (paramOut ++ paramIn ++ paramEnv) $ do + makeOpenAcc uid "transform" (paramOut ++ paramIn ++ paramEnv) $ do let start = liftInt 0 end <- shapeSize shr' (irArrayShape arrOut) diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Compile.hs b/src/Data/Array/Accelerate/LLVM/PTX/Compile.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Compile.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Compile.hs @@ -1,9 +1,8 @@ -{-# LANGUAGE CPP #-} {-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE RecordWildCards #-} {-# LANGUAGE TemplateHaskell #-} {-# LANGUAGE TypeFamilies #-} -{-# OPTIONS_GHC -fno-warn-orphans #-} +{-# OPTIONS_GHC -Wno-orphans #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Compile -- Copyright : [2014..2020] The Accelerate Team @@ -25,65 +24,54 @@ import Data.Array.Accelerate.Error import Data.Array.Accelerate.Trafo.Delayed -import Data.Array.Accelerate.LLVM.CodeGen +import Data.Array.Accelerate.LLVM.CodeGen ( llvmOfPreOpenAcc ) import Data.Array.Accelerate.LLVM.CodeGen.Environment ( Gamma ) import Data.Array.Accelerate.LLVM.CodeGen.Module ( Module(..) ) import Data.Array.Accelerate.LLVM.Compile -import Data.Array.Accelerate.LLVM.Extra import Data.Array.Accelerate.LLVM.State +import Data.Array.Accelerate.LLVM.Target.ClangInfo ( hostLLVMVersion, llvmverFromTuple, clangExePath, clangExePathEnvironment ) import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch import Data.Array.Accelerate.LLVM.PTX.CodeGen import Data.Array.Accelerate.LLVM.PTX.Compile.Cache -import Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice +import Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.Load import Data.Array.Accelerate.LLVM.PTX.Foreign ( ) import Data.Array.Accelerate.LLVM.PTX.Target import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug -import Foreign.CUDA.Path +import Foreign.CUDA.Path ( cudaInstallPath ) import qualified Foreign.CUDA.Analysis as CUDA -import qualified Foreign.NVVM as NVVM -import qualified LLVM.AST as AST -import qualified LLVM.AST.Name as LLVM -import qualified LLVM.Context as LLVM -import qualified LLVM.Module as LLVM -import qualified LLVM.PassManager as LLVM -import qualified LLVM.Target as LLVM -import qualified LLVM.Internal.Module as LLVM.Internal -import qualified LLVM.Internal.FFI.LLVMCTypes as LLVM.Internal.FFI -import qualified LLVM.Analysis as LLVM +import qualified LLVM.AST.Type.Name as LLVM -import Control.DeepSeq -import Control.Exception -import Control.Monad.Except -import Control.Monad.State -import Data.ByteString ( ByteString ) +import qualified Data.Array.Accelerate.LLVM.Internal.LLVMPretty as LP +import qualified Data.Array.Accelerate.LLVM.Internal.LLVMPretty.PP as LP +import qualified Text.PrettyPrint as Pretty + +import Control.Monad ( when ) +import Control.Monad.Reader import Data.ByteString.Short ( ShortByteString ) -import Data.Maybe -import Data.Word -import Foreign.ForeignPtr -import Foreign.Ptr -import Foreign.Storable +import Data.List ( intercalate ) +import qualified Data.List.NonEmpty as NE +import Data.Foldable ( toList ) +import GHC.IO.Exception ( IOErrorType(OtherError) ) +import Formatting import System.Directory -import System.Exit -import System.FilePath -import System.IO +import System.Exit ( ExitCode(..) ) +import System.IO ( hPutStrLn, stderr ) +import System.IO.Error ( mkIOError ) import System.IO.Unsafe import System.Process -import System.Process.Extra import Text.Printf ( printf ) -import qualified Data.Map as Map -import qualified Data.ByteString as B -import qualified Data.ByteString.Char8 as B8 -import qualified Data.ByteString.Internal as B -import Prelude as P +import qualified Data.ByteString.Short.Char8 as SBS8 +import qualified Data.Map.Strict as Map instance Compile PTX where data ObjectR PTX = ObjectR { objId :: {-# UNPACK #-} !UID - , ptxConfig :: ![(ShortByteString, LaunchConfig)] - , objData :: {- LAZY -} ByteString + , -- | Config for each exported kernel (symbol) + ptxConfig :: ![(ShortByteString, LaunchConfig)] + , objPath :: {- LAZY -} FilePath } compileForTarget = compile @@ -98,11 +86,19 @@ -- Generate code for this Acc operation -- - dev <- gets ptxDeviceProperties - (uid, cacheFile) <- cacheOfPreOpenAcc pacc - Module ast md <- llvmOfPreOpenAcc uid pacc aenv - let config = [ (f,x) | (LLVM.Name f, KM_PTX x) <- Map.toList md ] + dev <- asks ptxDeviceProperties + let CUDA.Compute m n = CUDA.computeCapability dev + let arch = printf "sm_%d%d" m n + (uid, cacheFile) <- cacheOfPreOpenAcc pacc + Module ast md <- llvmOfPreOpenAcc uid pacc aenv + let config = [ (SBS8.pack f, x) | (LP.Symbol f, KM_PTX x) <- Map.toList md ] + libdevice_bc <- liftIO libdeviceBitcodePath + + case isDeviceSupported (CUDA.computeCapability dev) of + Nothing -> return () -- all fine + Just err -> internalError string err + -- Lower the generated LLVM into a CUBIN object code. -- -- The 'objData' field is lazily evaluated since the object code might have @@ -114,171 +110,174 @@ recomp <- if Debug.debuggingIsEnabled then Debug.getFlag Debug.force_recomp else return False if exists && not recomp then do - Debug.traceIO Debug.dump_cc (printf "cc: found cached object code %s" (show uid)) - B.readFile cacheFile + Debug.traceM Debug.dump_cc ("cc: found cached object code " % shown) uid - else - LLVM.withContext $ \ctx -> do - ptx <- compilePTX dev ctx ast - cubin <- compileCUBIN dev cacheFile ptx - return cubin + else do + -- Detect LLVM version + -- Note: this LLVM version is incorporated in the cache path, so we're safe detecting it at runtime. + let prettyHostLLVMVersion = intercalate "." (map show (toList hostLLVMVersion)) + llvmver <- case llvmverFromTuple hostLLVMVersion of + Just llvmver -> return llvmver + Nothing -> internalError ("accelerate-llvm-ptx: Unsupported LLVM version: " % string) + prettyHostLLVMVersion + Debug.traceM Debug.dump_cc ("Using Clang at " % string % " version " % shown) clangExePath prettyHostLLVMVersion - return $! ObjectR uid config cubin + when (NE.head hostLLVMVersion < 16) $ + case clangExePathEnvironment of + Nothing -> do + hPutStrLn stderr $ + "[accelerate-llvm-ptx] Clang version 16 or newer is required for the Nvidia PTX " ++ + "backend, but version " ++ prettyHostLLVMVersion ++ " was found at '" ++ + clangExePath ++ "'. To override this choice, set the ACCELERATE_LLVM_CLANG_PATH " ++ + "environment variable to point to the desired clang executable." + -- not an IOError because we're in unsafePerformIO, somewhere up the call chain + errorWithoutStackTrace $ + "accelerate-llvm-ptx: Clang version " ++ prettyHostLLVMVersion ++ + " found but >=16 required (set ACCELERATE_LLVM_CLANG_PATH to override)" + Just{} -> -- If an explicit path was given, let's just continue and see what happens. + return () + -- Convert module to llvm-pretty format so that we can print it + let unoptimisedText = Pretty.renderStyle + Pretty.style { Pretty.lineLength = maxBound `div` 2 } + (LP.ppLLVM llvmver (LP.ppModule ast)) + ++ "\n\n" ++ accPreludePTX + Debug.when Debug.verbose $ do + Debug.traceM Debug.dump_cc ("Unoptimised LLVM IR:\n" % string) unoptimisedText --- | Compile the LLVM module to PTX assembly. This is done either by the --- closed-source libNVVM library, or via the standard NVPTX backend (which is --- the default). --- -compilePTX :: CUDA.DeviceProperties -> LLVM.Context -> AST.Module -> IO ByteString -compilePTX dev ctx ast = do -#ifdef ACCELERATE_USE_NVVM - ptx <- withLibdeviceNVVM dev ctx ast (_compileModuleNVVM dev (AST.moduleName ast)) -#else - ptx <- withLibdeviceNVPTX dev ctx ast (_compileModuleNVPTX dev) -#endif - Debug.when Debug.dump_asm $ Debug.traceIO Debug.verbose (B8.unpack ptx) - return ptx + isVerboseFlagSet <- Debug.getFlag Debug.verbose + let clangArgs = ["-O3", "--target=nvptx64-nvidia-cuda", "-march=" ++ arch + ,"-o", cacheFile + ,"-Wno-override-module" + ,"--cuda-path=" ++ cudaInstallPath + ,"-x", "ir", "-" + -- See Note [Internalizing Libdevice] + -- TODO: only link in libdevice if we're actually using __nv_ functions! + ,"-Xclang", "-mlink-builtin-bitcode", "-Xclang", libdevice_bc] + ++ (if isVerboseFlagSet then ["-v"] else []) + Debug.traceM Debug.dump_cc ("Arguments to clang: " % shown) clangArgs --- | Compile the given PTX assembly to a CUBIN file (SASS object code). The --- compiled code will be stored at the given FilePath. --- -compileCUBIN :: HasCallStack => CUDA.DeviceProperties -> FilePath -> ByteString -> IO ByteString -compileCUBIN dev sass ptx = do - _verbose <- if Debug.debuggingIsEnabled then Debug.getFlag Debug.verbose else return False - _debug <- if Debug.debuggingIsEnabled then Debug.getFlag Debug.debug else return False - -- - let verboseFlag = if _verbose then [ "-v" ] else [] - debugFlag = if _debug then [ "-g", "-lineinfo" ] else [] - arch = printf "-arch=sm_%d%d" m n - CUDA.Compute m n = CUDA.computeCapability dev - flags = "-" : "-o" : sass : arch : verboseFlag ++ debugFlag - -- - cp = (proc (cudaBinPath "ptxas") flags) - { std_in = CreatePipe - , std_out = NoStream - , std_err = CreatePipe - } + -- Remove some diagnostics from clang (and subprocesses) output that we + -- know are fine. See filterClangStderr. Unfortunately, System.Process + -- does not have a combinator for "give me stdout and stderr but throw + -- exception on ExitFailure", so we do it manually. + (clangEC, clangOut, clangErr) <- readProcessWithExitCode clangExePath clangArgs unoptimisedText + putStr clangOut + putStr (filterClangStderr clangErr) + case clangEC of + ExitSuccess -> return () + ExitFailure code -> do + let msg = "clang returned non-zero exit code: " ++ show code ++ + " (invocation: " ++ show (clangExePath : clangArgs) ++ ")" + ioError $ mkIOError OtherError msg Nothing Nothing - -- Invoke the 'ptxas' executable to compile the generated PTX into SASS (GPU - -- object code). The output is written directly to the final cache location. - -- - withCreateProcess cp $ \(Just inh) Nothing (Just errh) ph -> do + Debug.traceM Debug.dump_cc ("Written PTX to: " % string) cacheFile - -- fork off a thread to start consuming stderr - info <- hGetContents errh - withForkWait (evaluate (rnf info)) $ \waitErr -> do + return cacheFile - -- write the PTX to the input handle - -- closing the handle performs an implicit flush, thus may trigger SIGPIPE - ignoreSIGPIPE $ B.hPut inh ptx - ignoreSIGPIPE $ hClose inh + return $! ObjectR uid config cubin - -- wait on the output - waitErr - hClose errh - -- wait on the process - ex <- waitForProcess ph - case ex of - ExitFailure r -> internalError (printf "ptxas %s (exit %d)\n%s" (unwords flags) r info) - ExitSuccess -> return () - - when _verbose $ - unless (null info) $ - Debug.traceIO Debug.dump_cc (printf "ptx: compiled entry function(s)\n%s" info) - - -- Read back the results - B.readFile sass - +{- Note [Internalizing Libdevice] --- Compile and optimise the module to PTX using the (closed source) NVVM --- library. This _may_ produce faster object code than the LLVM NVPTX compiler. --- -_compileModuleNVVM :: HasCallStack => CUDA.DeviceProperties -> ShortByteString -> [(ShortByteString, ByteString)] -> LLVM.Module -> IO ByteString -_compileModuleNVVM dev name libdevice mdl = do - _debug <- if Debug.debuggingIsEnabled then Debug.getFlag Debug.debug else return False - -- - let arch = CUDA.computeCapability dev - verbose = if _debug then [ NVVM.GenerateDebugInfo ] else [] - flags = NVVM.Target arch : verbose +"Libdevice" refers to $CUDAPATH/nvvm/libdevice/libdevice.XX.bc, an LLVM bitcode +file that (reportedly) contains definitions of various math functions for use +in NVIDIA PTX. Most interesting primitive arithmetic operations on +floating-point numbers get compiled to calls to functions from libdevice, so it +is essential that we link it into any kernel that we create (or at least, any +kernel that references functions from libdevice). - -- Note: [NVVM and target datalayout] - -- - -- The NVVM library does not correctly parse the target datalayout field, - -- instead doing a (very dodgy) string compare against exactly two - -- expected values. This means that it is sensitive to, e.g. the ordering - -- of the fields, and changes to the representation in each LLVM release. - -- - -- We get around this by only specifying the data layout in a separate - -- (otherwise empty) module that we additionally link against. - -- - header = case bitSize (undefined::Int) of - 32 -> "target triple = \"nvptx-nvidia-cuda\"\ntarget datalayout = \"e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64\"" - 64 -> "target triple = \"nvptx64-nvidia-cuda\"\ntarget datalayout = \"e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64\"" - _ -> internalError "I don't know what architecture I am" +However, libdevice is quite large; it is 473 KB of LLVM bitcode for cuda 12.6 +on my machine, and clang takes >1 second to compile it on my (5 GHz Intel) +machine. Indeed, the LLVM NVPTX usage guide [1] recommends _internalizing_ the +symbols from libdevice after linking it with the kernel module; more precisely, +it recommends to first link the kernel module with libdevice, and subsequently +internalize all functions that we don't explicitly want exported (the public +kernel functions). - Debug.when Debug.dump_cc $ do - Debug.when Debug.verbose $ do - ll <- LLVM.moduleLLVMAssembly mdl -- TLM: unfortunate to do the lowering twice in debug mode - Debug.traceIO Debug.verbose (B8.unpack ll) +Clang doesn't have a command-line option to internalize symbols. Indeed, it +would be somewhat ambiguous when in the compilation process to do said +internalization. The LLVM command-line tool that _can_ do internalization is +`llvm-link`, the tool for linking LLVM modules together (and doing little +else). So translating the recommended [1] strategy to command-line tools +(because linking with LLVM through bindings is a version nightmare -- been +there, done that, not again), we get the following sensible procedure: - -- Lower the generated module to bitcode, then compile and link together with - -- the shim header and libdevice library (if necessary) - bc <- LLVM.moduleBitcode mdl - ptx <- NVVM.compileModules (("",header) : (name,bc) : libdevice) flags +$ llvm-link --internalize kernel.ll libdevice.bc -o kernel-linked.bc +$ clang --target=... kernel-linked.bc -o kernel.sass - unless (B.null (NVVM.compileLog ptx)) $ do - Debug.traceIO Debug.dump_cc $ "llvm: " ++ B8.unpack (NVVM.compileLog ptx) +However, llvm-link is not clang, and we'd very much like to depend _only_ on +clang, not on the full LLVM suite of tools. Especially not for this vexingly +small bit of functionality! But clang is huge, and surely it can do +internalization somehow? - -- Return the generated binary code - return (NVVM.compileResult ptx) +It turns out it can, but they did their absolute best to hide it. (All +references in this paragraph are to LLVM HEAD on 2024-12-04: 7954a0514ba7de.) +The workhorse function, called from `llvm-link`, is internalizeModule(). This +function is also called from clang in BackendConsumer::LinkInModules() in +clang/lib/CodeGen/CodeGenAction.cpp, but only if .Internalize is set on the +CodeGenAction::LinkModule in question. In CompilerInvocation::ParseCodeGenArgs +(clang/lib/Frontend/CompilerInvocation.cpp), we see that _some_ field called +"Internalize" is set on _something_ (not a LinkModule, but whatever?) if the +OPT_mlink_builtin_bitcode flag is set. Of course, no documentation anywhere +explains what this option does; the only mention I could find anywhere is here +[2], as well as some mailing list posts / issue tracker comments mentioning it. +How do we use the option? Well, it's not a clang option, it's actually a (I +think!) cc1 option, so you have to do: +$ clang -Xclang -mlink-builtin-bitcode -Xclang libdevice.bc --- Compiling with the NVPTX backend uses LLVM-3.3 and above --- -_compileModuleNVPTX :: CUDA.DeviceProperties -> LLVM.Module -> IO ByteString -_compileModuleNVPTX dev mdl = - withPTXTargetMachine dev $ \nvptx -> do +This makes clang internalize everything in that module that is not globally +exported (I think), which is what we want. - when Debug.internalChecksAreEnabled $ LLVM.verify mdl +[1]: https://releases.llvm.org/19.1.0/docs/NVPTXUsage.html#linking-with-libdevice +[2]: https://clang.llvm.org/docs/OffloadingDesign.html#offload-device-compilation +-} - -- Run the standard optimisation pass +filterClangStderr :: String -> String +filterClangStderr = unlines . filter (not . isShflSyncWarn) . lines + where + -- ptxas warns about use of shfl instructions without the .sync suffix on + -- CC 6.0, because such non-sync shuffles are deprecated (and indeed + -- removed in CC 7.0). We still use them in CC 6.0 (and not any more in CC + -- 7.0) because the shfl.sync in CC 6.0 has restrictions: -- - let pss = LLVM.defaultCuratedPassSetSpec { LLVM.optLevel = Just 3 } - LLVM.withPassManager pss $ \pm -> do - - b1 <- LLVM.runPassManager pm mdl - - -- debug printout - Debug.when Debug.dump_cc $ do - Debug.traceIO Debug.dump_cc $ printf "llvm: optimisation did work? %s" (show b1) - Debug.traceIO Debug.verbose . B8.unpack =<< LLVM.moduleLLVMAssembly mdl - - -- Lower the LLVM module into target assembly (PTX) - moduleTargetAssembly nvptx mdl + -- > For .target `sm_6x` or below, all threads in `membermask` must execute + -- > the same `shfl.sync` instruction in convergence, and only threads + -- > belonging to some `membermask` can be active when the `shfl.sync` + -- > instruction is executed. Otherwise, the behavior is undefined. + -- (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-shfl-sync) + -- + -- Perhaps we do use the shuffles in convergence, but we don't want to risk + -- it. Hence in CC 6.0, we still use non-sync shuffles. + -- + -- The ptxas warning cannot be turned off, however, and is **incredibly** + -- noisy (there's a warning for every single shfl instruction). Hence we + -- filter them out here. + -- + -- > ptxas /tmp/--f8a421.s, line 119; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version + isShflSyncWarn line = + let (presemi, postsemi) = break (== ';') line + in takeWhile (/= ' ') presemi == "ptxas" && + postsemi == "; warning : Instruction 'shfl' without '.sync' is deprecated since " ++ + "PTX ISA version 6.0 and will be discontinued in a future PTX ISA version" +-- | Returns a human-readable error message in case the device is unsupported, +-- and Nothing if everything is alright. +isDeviceSupported :: CUDA.Compute -> Maybe String +isDeviceSupported cc@(CUDA.Compute m _) + -- We require shfl instructions which are available only from CC 3.0. + | m >= 3 = Nothing + | otherwise = Just $ + "Your GPU has compute capability " ++ show cc ++ ", but only >= 3.0 is supported." --- | Produce target specific assembly as a 'ByteString'. --- -moduleTargetAssembly :: LLVM.TargetMachine -> LLVM.Module -> IO ByteString -moduleTargetAssembly tm m = unsafe0 =<< LLVM.Internal.emitToByteString LLVM.Internal.FFI.codeGenFileTypeAssembly tm m +accPreludePTX :: String +accPreludePTX = unlines + -- see Data.Array.Accelerate.LLVM.PTX.CodeGen.Base.nanosleep for why this is a hand-written function + ["define private void @" ++ name_nanosleep ++ "(i32 noundef %0) alwaysinline convergent nounwind {" + ," tail call void asm sideeffect \"nanosleep.u32 $0;\", \"r\"(i32 %0)" + ," ret void" + ,"}"] where - -- Ensure that the ByteString is NULL-terminated, so that it can be passed - -- directly to C. This will unsafely mutate the underlying ForeignPtr if the - -- string is not NULL-terminated but the last character is a whitespace - -- character (there are usually a few blank lines at the end). - -- - unsafe0 :: ByteString -> IO ByteString - unsafe0 bs@(B.PS fp s l) = - liftIO . withForeignPtr fp $ \p -> do - let p' :: Ptr Word8 - p' = p `plusPtr` (s+l-1) - -- - x <- peek p' - case x of - 0 -> return bs - _ | B.isSpaceWord8 x -> poke p' 0 >> return bs - _ -> return (B.snoc bs 0) - + name_nanosleep = let LLVM.Label name = LLVM.makeAccPreludeLabel "nanosleep" in SBS8.unpack name diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Compile/Cache.hs b/src/Data/Array/Accelerate/LLVM/PTX/Compile/Cache.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Compile/Cache.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Compile/Cache.hs @@ -1,5 +1,4 @@ -{-# LANGUAGE CPP #-} -{-# OPTIONS_GHC -fno-warn-orphans #-} +{-# OPTIONS_GHC -Wno-orphans #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Compile.Cache -- Copyright : [2017..2020] The Accelerate Team @@ -18,8 +17,11 @@ import Data.Array.Accelerate.LLVM.Compile.Cache import Data.Array.Accelerate.LLVM.PTX.Target +import Data.Array.Accelerate.LLVM.Target.ClangInfo ( hostLLVMVersion ) -import Control.Monad.State +import Control.Monad.Reader +import Data.Foldable ( toList ) +import Data.List ( intercalate ) import Data.Version import Foreign.CUDA.Analysis import System.FilePath @@ -31,9 +33,9 @@ instance Persistent PTX where targetCacheTemplate = do - Compute m n <- gets (computeCapability . ptxDeviceProperties) + Compute m n <- asks (computeCapability . ptxDeviceProperties) return $ "accelerate-llvm-ptx-" ++ showVersion version - "llvm-hs-" ++ VERSION_llvm_hs + "llvmpr-" ++ intercalate "." (map show (toList hostLLVMVersion)) S8.unpack ptxTargetTriple printf "sm%d%d" m n "morp.sass" diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs b/src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs deleted file mode 100644 --- a/src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs +++ /dev/null @@ -1,176 +0,0 @@ -{-# LANGUAGE OverloadedStrings #-} -{-# LANGUAGE RecordWildCards #-} -{-# LANGUAGE TemplateHaskell #-} -{-# LANGUAGE ViewPatterns #-} --- | --- Module : Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice --- Copyright : [2014..2020] The Accelerate Team --- License : BSD3 --- --- Maintainer : Trevor L. McDonell --- Stability : experimental --- Portability : non-portable (GHC extensions) --- - -module Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice ( - - withLibdeviceNVVM, - withLibdeviceNVPTX, - -) where - --- llvm-hs -import LLVM.Context -import qualified LLVM.Module as LLVM - -import LLVM.AST as AST -import LLVM.AST.Global as G -import LLVM.AST.Linkage - --- accelerate -import Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.Load -import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug - --- cuda -import Foreign.CUDA.Analysis - --- standard library -import Control.Monad -import Data.ByteString ( ByteString ) -import Data.ByteString.Short.Char8 ( ShortByteString ) -import Data.HashSet ( HashSet ) -import Data.List -import Data.Maybe -import Text.Printf -import qualified Data.ByteString.Short.Char8 as S8 -import qualified Data.ByteString.Short.Extra as BS -import qualified Data.HashSet as Set - - --- | Lower an LLVM AST to C++ objects and link it against the libdevice module, --- iff any libdevice functions are referenced from the base module. --- --- Note: [Linking with libdevice] --- --- The CUDA toolkit comes with an LLVM bitcode library called 'libdevice' that --- implements many common mathematical functions. The library can be used as a --- high performance math library for targets of the LLVM NVPTX backend, such as --- this one. To link a module 'foo' with libdevice, the following compilation --- pipeline is recommended: --- --- 1. Save all external functions in module 'foo' --- --- 2. Link module 'foo' with the appropriate 'libdevice_compute_XX.YY.bc' --- --- 3. Internalise all functions not in the list from (1) --- --- 4. Eliminate all unused internal functions --- --- 5. Run the NVVMReflect pass (see note: [NVVM Reflect Pass]) --- --- 6. Run the standard optimisation pipeline --- -withLibdeviceNVPTX - :: DeviceProperties - -> Context - -> Module - -> (LLVM.Module -> IO a) - -> IO a -withLibdeviceNVPTX dev ctx ast next = - case Set.null externs of - True -> LLVM.withModuleFromAST ctx ast next - False -> - LLVM.withModuleFromAST ctx ast $ \mdl -> - LLVM.withModuleFromAST ctx nvvmReflect $ \refl -> - LLVM.withModuleFromAST ctx (internalise externs libdev) $ \libd -> do - LLVM.linkModules mdl refl - LLVM.linkModules mdl libd - Debug.traceIO Debug.dump_cc msg - next mdl - where - -- Replace the target triple and datalayout from the libdevice.bc module - -- with those of the generated code. This avoids warnings such as "linking - -- two modules of different target triples..." - libdev = (libdevice arch) { moduleTargetTriple = moduleTargetTriple ast - , moduleDataLayout = moduleDataLayout ast - } - externs = analyse ast - arch = computeCapability dev - - msg = printf "cc: linking with libdevice: %s" - $ intercalate ", " - $ map S8.unpack - $ Set.toList externs - - --- | Lower an LLVM AST to C++ objects and prepare it for linking against --- libdevice using the nvvm bindings, iff any libdevice functions are referenced --- from the base module. --- --- Rather than internalise and strip any unused functions ourselves, allow the --- nvvm library to do so when linking the two modules together. --- --- TLM: This really should work with the above method, however for some reason --- we get a "CUDA Exception: function named symbol not found" error, even though --- the function is clearly visible in the generated code. hmm... --- -withLibdeviceNVVM - :: DeviceProperties - -> Context - -> Module - -> ([(ShortByteString, ByteString)] -> LLVM.Module -> IO a) - -> IO a -withLibdeviceNVVM dev ctx ast next = - LLVM.withModuleFromAST ctx ast $ \mdl -> do - when withlib $ Debug.traceIO Debug.dump_cc msg - next lib mdl - where - externs = analyse ast - withlib = not (Set.null externs) - lib | withlib = [ nvvmReflect, libdevice arch ] - | otherwise = [] - - arch = computeCapability dev - - msg = printf "cc: linking with libdevice: %s" - $ intercalate ", " - $ map S8.unpack - $ Set.toList externs - - --- | Analyse the LLVM AST module and determine if any of the external --- declarations are intrinsics implemented by libdevice. The set of such --- functions is returned, and will be used when determining which functions from --- libdevice to internalise. --- -analyse :: Module -> HashSet ShortByteString -analyse Module{..} = - let intrinsic (GlobalDefinition Function{..}) - | null basicBlocks - , Name n <- name - , "__nv_" <- BS.take 5 n - = Just n - - intrinsic _ - = Nothing - in - Set.fromList (mapMaybe intrinsic moduleDefinitions) - - --- | Mark all definitions in the module as internal linkage. This means that --- unused definitions can be removed as dead code. Be careful to leave any --- declarations as external. --- -internalise :: HashSet ShortByteString -> Module -> Module -internalise externals Module{..} = - let internal (GlobalDefinition Function{..}) - | Name n <- name - , not (Set.member n externals) -- we don't call this function directly; and - , not (null basicBlocks) -- it is not an external declaration - = GlobalDefinition Function { linkage=Internal, .. } - - internal x - = x - in - Module { moduleDefinitions = map internal moduleDefinitions, .. } - diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/Load.hs b/src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/Load.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/Load.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/Load.hs @@ -1,8 +1,5 @@ -{-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE OverloadedStrings #-} -{-# LANGUAGE RecordWildCards #-} -{-# LANGUAGE TemplateHaskell #-} -{-# LANGUAGE TupleSections #-} + -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.Load -- Copyright : [2014..2020] The Accelerate Team @@ -15,41 +12,21 @@ module Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.Load ( - nvvmReflect, libdevice, + libdeviceBitcodePath, ) where -import LLVM.Context -import LLVM.Module as LLVM -import LLVM.AST as AST ( Module(..) ) - import Data.Array.Accelerate.Error -import Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.TH import Data.Array.Accelerate.LLVM.PTX.Execute.Event ( ) -- GHC#1012 import Data.Array.Accelerate.LLVM.PTX.Execute.Stream ( ) -- GHC#1012 -import Foreign.CUDA.Analysis import qualified Foreign.CUDA.Driver as CUDA -import System.IO.Unsafe -import Data.ByteString ( ByteString ) -import Data.ByteString.Short.Char8 ( ShortByteString ) -import qualified Data.ByteString.Short.Char8 as S8 -import qualified Language.Haskell.TH as TH -import qualified Language.Haskell.TH.Syntax as TH - - --- NVVM Reflect --- ------------ - -class NVVMReflect a where - nvvmReflect :: a - -instance NVVMReflect AST.Module where - nvvmReflect = nvvmReflectModule +import Foreign.CUDA.Path -instance NVVMReflect (ShortByteString, ByteString) where - nvvmReflect = $$( nvvmReflectBitcode nvvmReflectModule ) +import Data.List ( isPrefixOf, sortBy ) +import System.Directory +import System.FilePath -- libdevice @@ -58,69 +35,26 @@ -- Compatible version of libdevice for a given compute capability should be -- listed here: -- --- https://github.com/llvm-mirror/llvm/blob/master/lib/Target/NVPTX/NVPTX.td#L72 --- -class Libdevice a where - libdevice :: Compute -> a - --- Load the libdevice bitcode files as an LLVM AST module. The top-level --- unsafePerformIO ensures that the data is only read from disk once per --- program execution. --- --- As of CUDA-9.0, libdevice is no longer split into multiple files --- depending on the target compute architecture. --- -$( let - libdeviceModule :: TH.ExpQ - libdeviceModule = [| \(name, bc) -> - unsafePerformIO $ - withContext $ \ctx -> - withModuleFromBitcode ctx (S8.unpack name, bc) moduleAST |] - in - if CUDA.libraryVersion < 9000 - then - [d| {-# NOINLINE libdevice_20_mdl #-} - {-# NOINLINE libdevice_30_mdl #-} - {-# NOINLINE libdevice_35_mdl #-} - {-# NOINLINE libdevice_50_mdl #-} - libdevice_20_mdl, libdevice_30_mdl, libdevice_35_mdl, libdevice_50_mdl :: AST.Module - libdevice_20_mdl = $libdeviceModule libdevice_20_bc - libdevice_30_mdl = $libdeviceModule libdevice_30_bc - libdevice_35_mdl = $libdeviceModule libdevice_35_bc - libdevice_50_mdl = $libdeviceModule libdevice_50_bc - - libdevice_20_bc, libdevice_30_bc, libdevice_35_bc, libdevice_50_bc :: (ShortByteString,ByteString) - libdevice_20_bc = $( TH.unTypeQ $ libdeviceBitcode (Compute 2 0) ) - libdevice_30_bc = $( TH.unTypeQ $ libdeviceBitcode (Compute 3 0) ) - libdevice_35_bc = $( TH.unTypeQ $ libdeviceBitcode (Compute 3 5) ) - libdevice_50_bc = $( TH.unTypeQ $ libdeviceBitcode (Compute 5 0) ) - - instance Libdevice AST.Module where - libdevice compute = - case compute of - Compute 2 _ -> libdevice_20_mdl -- 2.0, 2.1 - Compute 3 x - | x < 5 -> libdevice_30_mdl -- 3.0, 3.2 - | otherwise -> libdevice_35_mdl -- 3.5, 3.7 - Compute 5 _ -> libdevice_50_mdl -- 5.x - _ -> internalError - $ unlines [ "This device (compute capability " ++ show compute ++ ") is not supported by this version of the CUDA toolkit (" ++ show CUDA.libraryVersion ++ ")" - , "Please upgrade to the latest version of the CUDA toolkit and reinstall the 'cuda' package." - ] - |] - else - [d| {-# NOINLINE libdevice_mdl #-} - libdevice_mdl :: AST.Module - libdevice_mdl = $libdeviceModule libdevice_bc - - libdevice_bc :: (ShortByteString,ByteString) - libdevice_bc = $( TH.unTypeQ $ libdeviceBitcode undefined ) +-- https://github.com/llvm/llvm-project/blob/master/lib/Target/NVPTX/NVPTX.td - instance Libdevice AST.Module where - libdevice _ = libdevice_mdl +-- | Find the libdevice bitcode file for the given compute architecture. The name +-- of the bitcode file follows the format @libdevice.XX.bc@, where XX +-- represents a version(?). We search the libdevice path for all files of the +-- appropriate compute capability and load the "most recent" (by sort order). +libdeviceBitcodePath :: HasCallStack => IO FilePath +libdeviceBitcodePath + | CUDA.libraryVersion < 9000 = + -- There is some support code for cuda < 9 in an earlier version of these + -- files; in particular, look at commit + -- 2b5d69448557e89002c0179ea1aaf59bb757a6e3 (2023-08-22) + -- for original llvm-hs code. + internalError "Cuda < 9 is unsupported." + | otherwise = do + let nvvm = cudaInstallPath "nvvm" "libdevice" - instance Libdevice (ShortByteString,ByteString) where - libdevice _ = libdevice_bc - |] - ) + files <- getDirectoryContents nvvm + let matches f = "libdevice" `isPrefixOf` f && takeExtension f == ".bc" + return $ case sortBy (flip compare) (filter matches files) of + name : _ -> nvvm name + [] -> internalError "not found: libdevice.XX.bc" diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/TH.hs b/src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/TH.hs deleted file mode 100644 --- a/src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/TH.hs +++ /dev/null @@ -1,141 +0,0 @@ -{-# LANGUAGE CPP #-} -{-# LANGUAGE OverloadedStrings #-} -{-# LANGUAGE TemplateHaskell #-} -{-# LANGUAGE TypeApplications #-} --- | --- Module : Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.TH --- Copyright : [2017..2020] The Accelerate Team --- License : BSD3 --- --- Maintainer : Trevor L. McDonell --- Stability : experimental --- Portability : non-portable (GHC extensions) --- - -module Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.TH ( - - nvvmReflectModule, nvvmReflectBitcode, - libdeviceBitcode, - -) where - -import qualified LLVM.AST as AST -import qualified LLVM.AST.Attribute as AST -import qualified LLVM.AST.Global as AST.G -import qualified LLVM.Context as LLVM -import qualified LLVM.Module as LLVM - -import LLVM.AST.Type.Downcast -import LLVM.AST.Type.Representation - -import Data.Array.Accelerate.Error -import Data.Array.Accelerate.LLVM.CodeGen.Base -import Data.Array.Accelerate.LLVM.PTX.Target - -import Foreign.CUDA.Analysis -import qualified Foreign.CUDA.Driver as CUDA -#if MIN_VERSION_nvvm(0,10,0) -import Foreign.NVVM.Path -#else -import Foreign.CUDA.Path -#endif - -import Data.ByteString ( ByteString ) -import Data.ByteString.Short ( ShortByteString ) -import Data.FileEmbed -import Data.List -import Data.Maybe -import Language.Haskell.TH.Syntax ( Q, TExp ) -import System.Directory -import System.FilePath -import Text.Printf -import qualified Data.ByteString.Short as BS -import qualified Data.ByteString.Short.Char8 as S8 -import qualified Data.ByteString.Short.Extra as BS -import qualified Language.Haskell.TH as TH -import qualified Language.Haskell.TH.Syntax as TH - - --- This is a hacky module that can be linked against in order to provide the --- same functionality as running the NVVMReflect pass. --- --- Note: [NVVM Reflect Pass] --- --- To accommodate various math-related compiler flags that can affect code --- generation of libdevice code, the library code depends on a special LLVM IR --- pass (NVVMReflect) to handle conditional compilation within LLVM IR. This --- pass looks for calls to the @__nvvm_reflect function and replaces them with --- constants based on the defined reflection parameters. --- --- libdevice currently uses the following reflection parameters to control code --- generation: --- --- * __CUDA_FTZ={0,1} fast math that flushes denormals to zero --- --- Since this is currently the only reflection parameter supported, and that we --- prefer correct results over pure speed, we do not flush denormals to zero. If --- the list of supported parameters ever changes, we may need to re-evaluate --- this implementation. --- -nvvmReflectModule :: AST.Module -nvvmReflectModule = - AST.Module - { AST.moduleName = "nvvm-reflect" - , AST.moduleSourceFileName = BS.empty - , AST.moduleDataLayout = targetDataLayout @PTX - , AST.moduleTargetTriple = targetTriple @PTX - , AST.moduleDefinitions = [AST.GlobalDefinition $ AST.G.functionDefaults - { AST.G.name = AST.Name "__nvvm_reflect" - , AST.G.returnType = downcast (integralType :: IntegralType Int32) - , AST.G.parameters = ( [ptrParameter scalarType (UnName 0 :: Name (Ptr Int8))], False ) - , AST.G.functionAttributes = map Right [AST.NoUnwind, AST.ReadNone, AST.AlwaysInline] - , AST.G.basicBlocks = [] - }] - } - - --- Lower the given NVVM Reflect module into bitcode. --- -nvvmReflectBitcode :: AST.Module -> Q (TExp (ShortByteString, ByteString)) -nvvmReflectBitcode mdl = do - let name = "__nvvm_reflect" - -- - bs <- TH.runIO $ LLVM.withContext $ \ctx -> - LLVM.withModuleFromAST ctx mdl LLVM.moduleLLVMAssembly - TH.unsafeTExpCoerce $ TH.tupE [ TH.unTypeQ (BS.liftSBS name) - , bsToExp bs ] - - --- Load the libdevice bitcode file for the given compute architecture. The name --- of the bitcode files follows the format: --- --- libdevice.compute_XX.YY.bc --- --- Where XX represents the compute capability, and YY represents a version(?) We --- search the libdevice PATH for all files of the appropriate compute capability --- and load the "most recent" (by sort order). --- -libdeviceBitcode :: HasCallStack => Compute -> Q (TExp (ShortByteString, ByteString)) -libdeviceBitcode compute = do - let basename - | CUDA.libraryVersion < 9000 - , Compute m n <- compute = printf "libdevice.compute_%d%d" m n - | otherwise = "libdevice" - -- - err = internalError (printf "not found: %s.YY.bc" basename) - best f = basename `isPrefixOf` f && takeExtension f == ".bc" -#if MIN_VERSION_nvvm(0,10,0) - base = nvvmDeviceLibraryPath -#else - base = cudaInstallPath "nvvm" "libdevice" -#endif - - -- - files <- TH.runIO $ getDirectoryContents base - -- - let name = fromMaybe err . listToMaybe . sortBy (flip compare) $ filter best files - path = base name - -- - TH.unsafeTExpCoerce $ TH.tupE [ TH.unTypeQ (BS.liftSBS (S8.pack name)) - , embedFile path ] - diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Context.hs b/src/Data/Array/Accelerate/LLVM/PTX/Context.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Context.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Context.hs @@ -1,5 +1,8 @@ -{-# LANGUAGE MagicHash #-} -{-# LANGUAGE RecordWildCards #-} +{-# LANGUAGE BangPatterns #-} +{-# LANGUAGE LambdaCase #-} +{-# LANGUAGE MagicHash #-} +{-# LANGUAGE OverloadedStrings #-} +{-# LANGUAGE RecordWildCards #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Context -- Copyright : [2014..2020] The Accelerate Team @@ -14,25 +17,37 @@ Context(..), new, raw, withContext, + contextFinalizeResource, ) where import Data.Array.Accelerate.Lifetime import Data.Array.Accelerate.LLVM.PTX.Analysis.Device -import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug +import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug -import qualified Foreign.CUDA.Driver.Device as CUDA -import qualified Foreign.CUDA.Driver.Context as CUDA +import qualified Foreign.CUDA.Driver.Device as CUDA +import qualified Foreign.CUDA.Driver.Context as CUDA import Control.Concurrent import Control.Exception import Control.Monad +import Data.Char import Data.Hashable -import Text.PrettyPrint -import Prelude hiding ( (<>) ) +import Data.Int +import Data.IORef +import Data.Primitive.ByteArray +import Data.Text.Lazy.Builder +import Data.Word +import Formatting +import Prettyprinter +import Prettyprinter.Internal +import Prettyprinter.Render.Util.Panic +import Text.Printf +import qualified Data.Text.Lazy.Builder as TLB +import Prelude hiding ( (<>) ) -import GHC.Base ( Int(..), addr2Int#, ) -import GHC.Ptr ( Ptr(..) ) +import GHC.Base ( Int(..), addr2Int#, ) +import GHC.Ptr ( Ptr(..) ) -- | An execution context, which is tied to a specific device and CUDA execution @@ -40,7 +55,13 @@ -- data Context = Context { deviceProperties :: {-# UNPACK #-} !CUDA.DeviceProperties -- information on hardware resources + , deviceName :: {-# UNPACK #-} !ByteArray -- device name, used for profiling , deviceContext :: {-# UNPACK #-} !(Lifetime CUDA.Context) -- device execution context + + -- | The number of finalizers currently using the context to free resources, + -- plus 1 if the Context finalizer does not yet want to destroy the context + -- itself. See Note: [Finalizing a CUDA Context]. + , deviceFinalizerRefcount :: {-# UNPACK #-} !(IORef Int) } instance Eq Context where @@ -73,10 +94,7 @@ -> CUDA.Context -> IO Context raw dev prp ctx = do - lft <- newLifetime ctx - addFinalizer lft $ do - message $ "finalise context " ++ showContext ctx - CUDA.destroy ctx + refcountVar <- newIORef 1 -- there is one user of the context: the context itself -- The kernels don't use much shared memory, so for devices that support it -- prefer using those memory banks as an L1 cache. @@ -89,12 +107,27 @@ when (CUDA.computeCapability prp >= CUDA.Compute 2 0) (CUDA.setCache CUDA.PreferL1) + -- Generate the context name + let str = printf "[%d] %s\0" (fromIntegral (CUDA.useDevice dev) :: Int) (CUDA.deviceName prp) + mba <- newPinnedByteArray (length str) + + let go !_ [] = unsafeFreezeByteArray mba + go !i (x:xs) = do + writeByteArray mba i (fromIntegral (ord x) :: Word8) + go (i+1) xs + + nm <- go 0 str + -- Display information about the selected device - Debug.traceIO Debug.dump_phases (deviceInfo dev prp) + Debug.traceM Debug.dump_phases builder (deviceInfo dev prp) - return $! Context prp lft + lft <- newLifetime ctx -- on the CUDA context + let !result = Context prp nm lft refcountVar + addFinalizer lft $ decrementContext result + return result + -- | Push the context onto the CPUs thread stack of current contexts and execute -- some operation. -- @@ -108,16 +141,44 @@ {-# INLINE push #-} push :: CUDA.Context -> IO () push ctx = do - message $ "push context: " ++ showContext ctx + message ("push context: " % formatContext) ctx CUDA.push ctx {-# INLINE pop #-} pop :: IO () pop = do ctx <- CUDA.pop - message $ "pop context: " ++ showContext ctx + message ("pop context: " % formatContext) ctx +decrementContext :: Context -> IO () +decrementContext ctx = do + newCount <- atomicModifyIORef' (deviceFinalizerRefcount ctx) (\i -> (i - 1, i - 1)) + message ("decrement context " % formatContext % " to " % shown) (unsafeGetValue (deviceContext ctx)) newCount + when (newCount == 0) $ CUDA.destroy (unsafeGetValue (deviceContext ctx)) +-- | If the underlying CUDA context is already slated for destruction entirely +-- (or has already been destroyed) by its finalizer, this function does +-- nothing. If the CUDA context will live on (for now), the passed @IO@ action +-- is invoked with a lock held so that the CUDA context will not be destroyed +-- while your action runs. Use this in finalizers of CUDA resources such as +-- arrays and linked modules. +-- +-- The context is not automatically pushed in the action; if you need to +-- 'withContext', do it yourself. +contextFinalizeResource :: Context -> IO () -> IO () +contextFinalizeResource ctx action = + -- See Note: [Finalizing a CUDA Context] + bracket + (do newCount <- atomicModifyIORef' (deviceFinalizerRefcount ctx) $ \i -> + if i == 0 then (0, 0) else (i + 1, i + 1) + message ("increment context " % formatContext % " to " % shown) (unsafeGetValue (deviceContext ctx)) newCount + return (newCount > 0)) + (\contextStillLive -> + when contextStillLive (decrementContext ctx)) + (\contextStillLive -> + when contextStillLive action) + + -- Debugging -- --------- @@ -125,36 +186,84 @@ -- -- Device 0: GeForce 9600M GT (compute capability 1.1), 4 multiprocessors @ 1.25GHz (32 cores), 512MB global memory -- -deviceInfo :: CUDA.Device -> CUDA.DeviceProperties -> String -deviceInfo dev prp = render $ +deviceInfo :: CUDA.Device -> CUDA.DeviceProperties -> Builder +deviceInfo dev prp = go $ layoutPretty defaultLayoutOptions $ devID <> colon <+> name <+> parens compute - <> comma <+> processors <+> at <+> text clock <+> parens cores + <> comma <+> processors <+> at <+> pretty clock <+> parens cores <> comma <+> memory where - name = text (CUDA.deviceName prp) - compute = text "compute capability" <+> text (show $ CUDA.computeCapability prp) - devID = text "device" <+> int (fromIntegral $ CUDA.useDevice dev) - processors = int (CUDA.multiProcessorCount prp) <+> text "multiprocessors" - cores = int (CUDA.multiProcessorCount prp * coresPerMultiProcessor prp) <+> text "cores" - memory = text mem <+> text "global memory" - -- - clock = Debug.showFFloatSIBase (Just 2) 1000 (fromIntegral $ CUDA.clockRate prp * 1000 :: Double) "Hz" - mem = Debug.showFFloatSIBase (Just 0) 1024 (fromIntegral $ CUDA.totalGlobalMem prp :: Double) "B" - at = char '@' - -- reset = zeroWidthText "\r" + name = pretty (CUDA.deviceName prp) + compute = "compute capability" <+> unsafeViaShow (CUDA.computeCapability prp) + devID = "device" <+> unsafeViaShow (CUDA.useDevice dev) + processors = pretty (CUDA.multiProcessorCount prp) <+> "multiprocessors" + cores = pretty (CUDA.multiProcessorCount prp * coresPerMultiProcessor prp) <+> "cores" + memory = pretty mem <+> "global memory" + ---- + clock = toLazyText $ Debug.showFFloatSIBase (Just 2) 1000 (fromIntegral $ CUDA.clockRate prp * 1000 :: Double) "Hz" + mem = toLazyText $ Debug.showFFloatSIBase (Just 0) 1024 (fromIntegral $ CUDA.totalGlobalMem prp :: Double) "B" + at = pretty '@' + go = \case + SFail -> panicUncaughtFail + SEmpty -> mempty + SChar c rest -> TLB.singleton c <> go rest + SText _l t rest -> TLB.fromText t <> go rest + SLine i rest -> TLB.singleton '\n' <> (TLB.fromText (textSpaces i) <> go rest) + SAnnPush _ann rest -> go rest + SAnnPop rest -> go rest -{-# INLINE trace #-} -trace :: String -> IO a -> IO a -trace msg next = do - Debug.traceIO Debug.dump_gc ("gc: " ++ msg) - next {-# INLINE message #-} -message :: String -> IO () -message s = s `trace` return () +message :: Format (IO ()) a -> a +message fmt = Debug.traceM Debug.dump_gc ("gc: " % fmt) -{-# INLINE showContext #-} -showContext :: CUDA.Context -> String -showContext (CUDA.Context c) = show c +{-# INLINE formatContext #-} +formatContext :: Format r (CUDA.Context -> r) +formatContext = later $ \(CUDA.Context c) -> bformat shown c + + +-- Note: [Finalizing a CUDA Context] +-- +-- Both a CUDA context and the resources we allocate within such a context +-- (currently, arrays, executable modules and events) are freed with +-- finalizers; these are invoked by the GC when it detects (after a GC pass) +-- that the Haskell heap objects are no longer reachable. +-- +-- In our case, finalizers are attached to 'Lifetime' objects. The problem is +-- that even if a finalizer for Lifetime 1 refers to Lifetime 2, the GC does +-- not guarantee that the finalizer for Lifetime 1 runs to completion before the +-- finalizer of Lifetime 2 starts. (See the documentation for 'touchForeignPtr' +-- in base.) This is a problem for us because we are in this situation: to free +-- a resource within a CUDA context, we need a reference to that context and +-- the context needs to be alive. Thus finalizers for e.g. arrays reference the +-- Lifetime for the CUDA context. +-- +-- If we just leave GHC to do finalization as it wishes, this means that a CUDA +-- context may well be destroyed before the resourcees in it are finalized, +-- leading to use-after-free errors and segfaults. We have two choices here: +-- 1. either we let the finalizer for a Context wait until the other finalizers +-- have run, or +-- 2. we free the Context when first we can, and let resource finalizers that +-- come later, do nothing. +-- We choose option 2 because destroying a CUDA context already frees the +-- resources in it, so there is no need to do meticulous manual cleanup here. +-- +-- To accomplish this, we use reference counting. The context itself being +-- alive (precisely: its finalizer not yet having run) counts for one "use"; +-- the only other "uses" are the finalizers for the CUDA resources in the +-- context. (There is no need to track anything before we start finalizing, so +-- this is enough.) The Context finalizer does nothing more than a decrement on +-- the refcount to release the "use" of the CUDA context by the Context object +-- itself. +-- +-- To complete the picture: +-- - When decrementing, if we decremented to zero, we destroy the CUDA context. +-- - When incrementing, we are apparently in a resource finalizer, because that +-- is the only place where we increment. If here we see that the refcount is +-- already zero, we simply do nothing (and *skip* the resource finalizer +-- entirely): the CUDA context has already been destroyed, taking this +-- resource with it, so nothing more needs to be done. +-- +-- The resource finalizers use 'contextFinalizeResource' to access this +-- functionality. diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Debug.hs b/src/Data/Array/Accelerate/LLVM/PTX/Debug.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Debug.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Debug.hs @@ -1,4 +1,6 @@ {-# LANGUAGE ForeignFunctionInterface #-} +{-# LANGUAGE LambdaCase #-} +{-# LANGUAGE OverloadedStrings #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Debug -- Copyright : [2014..2020] The Accelerate Team @@ -11,20 +13,22 @@ module Data.Array.Accelerate.LLVM.PTX.Debug ( - module Data.Array.Accelerate.Debug, + module Data.Array.Accelerate.Debug.Internal, module Data.Array.Accelerate.LLVM.PTX.Debug, ) where -import Data.Array.Accelerate.Debug hiding ( timed, elapsed ) +import Data.Array.Accelerate.Debug.Internal hiding ( timed, elapsed ) +import qualified Data.Array.Accelerate.Debug.Internal as D -import Foreign.CUDA.Driver.Stream ( Stream ) -import qualified Foreign.CUDA.Driver.Event as Event +import Foreign.CUDA.Driver.Stream ( Stream ) +import qualified Foreign.CUDA.Driver.Event as Event -import Control.Monad.Trans import Control.Concurrent +import Control.Monad.Trans +import Data.Text.Lazy.Builder +import Formatting import System.CPUTime -import Text.Printf import GHC.Float @@ -33,14 +37,15 @@ -- to format the output string given elapsed GPU and CPU time respectively -- timed - :: Flag - -> (Double -> Double -> Double -> String) + :: MonadIO m + => Flag + -> (Double -> Double -> Double -> IO ()) -> Maybe Stream - -> IO () - -> IO () + -> m a + -> m a {-# INLINE timed #-} -timed f msg = - monitorProcTime (getFlag f) (\t1 t2 t3 -> traceIO f (msg t1 t2 t3)) +timed f fmt = + monitorProcTime (getFlag f) fmt monitorProcTime :: MonadIO m @@ -87,13 +92,22 @@ {-# INLINE elapsed #-} -elapsed :: Double -> Double -> Double -> String -elapsed wallTime cpuTime gpuTime = - printf "%s (wall), %s (cpu), %s (gpu)" - (showFFloatSIBase (Just 3) 1000 wallTime "s") - (showFFloatSIBase (Just 3) 1000 cpuTime "s") - (showFFloatSIBase (Just 3) 1000 gpuTime "s") +elapsed :: Format r (Double -> Double -> Double -> r) +elapsed = formatSIBase (Just 3) 1000 % "s (wall), " + % formatSIBase (Just 3) 1000 % "s (cpu), " + % formatSIBase (Just 3) 1000 % "s (gpu)" -- accelerate/cbits/clock.c foreign import ccall unsafe "clock_gettime_monotonic_seconds" getMonotonicTime :: IO Double + +data Phase = Compile | Link | Execute + +buildPhase :: Phase -> Builder +buildPhase = \case + Compile -> "compile" + Link -> "link" + Execute -> "execute" + +phase :: MonadIO m => Phase -> m a -> m a +phase p = D.timed dump_phases (now ("phase " <> buildPhase p <> ": ") % D.elapsed) diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Embed.hs b/src/Data/Array/Accelerate/LLVM/PTX/Embed.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Embed.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Embed.hs @@ -1,7 +1,6 @@ -{-# LANGUAGE CPP #-} {-# LANGUAGE QuasiQuotes #-} {-# LANGUAGE TemplateHaskell #-} -{-# OPTIONS_GHC -fno-warn-orphans #-} +{-# OPTIONS_GHC -Wno-orphans #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Embed -- Copyright : [2017..2020] The Accelerate Team @@ -32,14 +31,14 @@ import qualified Foreign.CUDA.Driver as CUDA +import Control.Monad.IO.Class ( liftIO ) import Foreign.Ptr import GHC.Ptr ( Ptr(..) ) -import Language.Haskell.TH ( Q, TExp ) +import Data.Array.Accelerate.TH.Compat ( CodeQ ) import System.IO.Unsafe import qualified Data.ByteString as B import qualified Data.ByteString.Unsafe as B -import qualified Language.Haskell.TH as TH -import qualified Language.Haskell.TH.Syntax as TH +import qualified Data.Array.Accelerate.TH.Compat as TH instance Embed PTX where @@ -47,32 +46,38 @@ -- Embed the given object code and set up to be reloaded at execution time. -- -embed :: PTX -> ObjectR PTX -> Q (TExp (ExecutableR PTX)) -embed target (ObjectR _ cfg obj) = do - -- Load the module to recover information such as number of registers and - -- bytes of shared memory. It may be possible to do this without requiring an - -- active CUDA context. - kmd <- TH.runIO $ withContext (ptxContext target) $ do - jit <- B.unsafeUseAsCString obj $ \p -> CUDA.loadDataFromPtrEx (castPtr p) [] - ks <- mapM (uncurry (linkFunctionQ (CUDA.jitModule jit))) cfg - CUDA.unload (CUDA.jitModule jit) - return ks - +embed :: PTX -> ObjectR PTX -> CodeQ (ExecutableR PTX) +embed target (ObjectR _ cfg objFname) = do -- Generate the embedded kernel executable. This will load the embedded object -- code into the current (at execution time) context. - [|| unsafePerformIO $ do - jit <- CUDA.loadDataFromPtrEx $$( TH.unsafeTExpCoerce [| Ptr $(TH.litE (TH.StringPrimL (B.unpack obj))) |] ) [] - fun <- newLifetime (FunctionTable $$(listE (map (linkQ 'jit) kmd))) - return $ PTXR fun - ||] + loadQ `TH.bindCode` \kmd -> + [|| unsafePerformIO $ do + jit <- CUDA.loadDataFromPtrEx + $$( liftIO (B.readFile objFname) `TH.bindCode` \obj -> + TH.unsafeCodeCoerce [| Ptr $(TH.litE (TH.StringPrimL (B.unpack obj))) |] ) + [] + fun <- newLifetime (FunctionTable $$(listE (map (linkQ 'jit) kmd))) + return $ PTXR fun + ||] where - linkQ :: TH.Name -> (Kernel, Q (TExp (Int -> Int))) -> Q (TExp Kernel) + -- Load the module to recover information such as number of registers + -- and bytes of shared memory. It may be possible to do this without + -- requiring an active CUDA context. + loadQ :: TH.Q [(Kernel, CodeQ (Int -> Int))] + loadQ = TH.runIO $ withContext (ptxContext target) $ do + obj <- B.readFile objFname + jit <- B.unsafeUseAsCString obj $ \p -> CUDA.loadDataFromPtrEx (castPtr p) [] + ks <- mapM (uncurry (linkFunctionQ (CUDA.jitModule jit))) cfg + CUDA.unload (CUDA.jitModule jit) + return ks + + linkQ :: TH.Name -> (Kernel, CodeQ (Int -> Int)) -> CodeQ Kernel linkQ jit (Kernel name _ dsmem cta _, grid) = [|| unsafePerformIO $ do - f <- CUDA.getFun (CUDA.jitModule $$(TH.unsafeTExpCoerce (TH.varE jit))) $$(liftSBS name) + f <- CUDA.getFun (CUDA.jitModule $$(TH.unsafeCodeCoerce (TH.varE jit))) $$(liftSBS name) return $ Kernel $$(liftSBS name) f dsmem cta $$grid ||] - listE :: [Q (TExp a)] -> Q (TExp [a]) - listE xs = TH.unsafeTExpCoerce (TH.listE (map TH.unTypeQ xs)) + listE :: [CodeQ a] -> CodeQ [a] + listE xs = TH.unsafeCodeCoerce (TH.listE (map TH.unTypeCode xs)) diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Execute.hs b/src/Data/Array/Accelerate/LLVM/PTX/Execute.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Execute.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Execute.hs @@ -1,4 +1,3 @@ -{-# LANGUAGE CPP #-} {-# LANGUAGE FlexibleContexts #-} {-# LANGUAGE GADTs #-} {-# LANGUAGE OverloadedStrings #-} @@ -8,7 +7,7 @@ {-# LANGUAGE TypeApplications #-} {-# LANGUAGE TypeOperators #-} {-# LANGUAGE ViewPatterns #-} -{-# OPTIONS_GHC -fno-warn-orphans #-} +{-# OPTIONS_GHC -Wno-orphans #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Execute -- Copyright : [2014..2020] The Accelerate Team @@ -36,29 +35,30 @@ import Data.Array.Accelerate.LLVM.Execute -import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch ( multipleOf ) +import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch ( multipleOf ) import Data.Array.Accelerate.LLVM.PTX.Array.Data -import Data.Array.Accelerate.LLVM.PTX.Array.Prim ( memsetArrayAsync ) +import Data.Array.Accelerate.LLVM.PTX.Array.Prim ( memsetArrayAsync ) import Data.Array.Accelerate.LLVM.PTX.Execute.Async import Data.Array.Accelerate.LLVM.PTX.Execute.Environment import Data.Array.Accelerate.LLVM.PTX.Execute.Marshal -import Data.Array.Accelerate.LLVM.PTX.Execute.Stream ( Stream ) +import Data.Array.Accelerate.LLVM.PTX.Execute.Stream ( Stream ) import Data.Array.Accelerate.LLVM.PTX.Link import Data.Array.Accelerate.LLVM.PTX.Target -import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug -import qualified Data.Array.Accelerate.LLVM.PTX.Execute.Event as Event +import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug +import qualified Data.Array.Accelerate.LLVM.PTX.Execute.Event as Event -import qualified Foreign.CUDA.Driver as CUDA +import qualified Foreign.CUDA.Driver as CUDA -import Control.Monad ( when, forM_ ) -import Control.Monad.Reader ( asks, local ) -import Control.Monad.State ( liftIO ) -import Data.ByteString.Short.Char8 ( ShortByteString, unpack ) -import qualified Data.DList as DL -import Data.List ( find ) -import Data.Maybe ( fromMaybe ) -import Text.Printf ( printf ) -import Prelude hiding ( exp, map, sum, scanl, scanr ) +import Control.Monad ( forM_ ) +import Control.Monad.State ( liftIO ) +import Data.ByteString.Short.Char8 ( ShortByteString, unpack ) +import Data.List ( find ) +import Data.Maybe ( fromMaybe ) +import Formatting +import Prelude hiding ( exp, map, sum, scanl, scanr ) +import qualified Data.ByteString.Short as S +import qualified Data.ByteString.Short.Extra as SE +import qualified Data.DList as DL {-# SPECIALISE INLINE executeAcc :: ExecAcc PTX a -> Par PTX (FutureArraysR PTX a) #-} @@ -130,8 +130,8 @@ result <- allocateRemote repr sh -- let paramR = TupRsingle $ ParamRarray repr - executeOp (ptxExecutable !# name) gamma aenv (arrayRshape repr) sh paramR result - put future result + cleanup <- executeOp (ptxExecutable !# name) gamma aenv (arrayRshape repr) sh paramR result + putCleanup future cleanup result return future -- Mapping over an array can ignore the dimensionality of the array and @@ -157,8 +157,8 @@ Nothing -> allocateRemote reprOut sh -- let paramsR = TupRsingle (ParamRarray reprOut) `TupRpair` TupRsingle (ParamRarray repr) - executeOp (ptxExecutable !# "map") gamma aenv (arrayRshape repr) sh paramsR (result, input) - put future result + cleanup <- executeOp (ptxExecutable !# "map") gamma aenv (arrayRshape repr) sh paramsR (result, input) + putCleanup future cleanup result return future {-# INLINE generateOp #-} @@ -188,8 +188,8 @@ future <- new result <- allocateRemote repr' sh' let paramsR = TupRsingle (ParamRarray repr') `TupRpair` TupRsingle (ParamRarray repr) - executeOp (ptxExecutable !# "transform") gamma aenv (arrayRshape repr') sh' paramsR (result, input) - put future result + cleanup <- executeOp (ptxExecutable !# "transform") gamma aenv (arrayRshape repr') sh' paramsR (result, input) + putCleanup future cleanup result return future {-# INLINE backpermuteOp #-} @@ -290,29 +290,30 @@ -- The array is small enough that we can compute it in a single step result <- allocateRemote (ArrayR dim0 tp) () let paramsR = paramsRdim0 `TupRpair` paramsRinput - executeOp ks gamma aenv dim1 sh paramsR (result, manifest input) - put future result + cleanup <- executeOp ks gamma aenv dim1 sh paramsR (result, manifest input) + putCleanup future cleanup result else do -- Multi-kernel reduction to a single element. The first kernel integrates -- any delayed elements, and the second is called recursively until -- reaching a single element. + -- The cleanup function is accumulated. let - rec :: Vector e -> Par PTX () - rec tmp@(Array ((),m) adata) - | m <= 1 = put future (Array () adata) + rec :: Vector e -> IO () -> Par PTX () + rec tmp@(Array ((),m) adata) cleanup + | m <= 1 = putCleanup future cleanup (Array () adata) | otherwise = do let sh' = ((), m `multipleOf` kernelThreadBlockSize km2) out <- allocateRemote (ArrayR dim1 tp) sh' let paramsR2 = paramsRdim1 `TupRpair` paramsRdim1 - executeOp km2 gamma aenv dim1 sh' paramsR2 (tmp, out) - rec out + cleanup2 <- executeOp km2 gamma aenv dim1 sh' paramsR2 (tmp, out) + rec out (cleanup >> cleanup2) -- let sh' = ((), n `multipleOf` kernelThreadBlockSize km1) tmp <- allocateRemote (ArrayR dim1 tp) sh' let paramsR1 = paramsRdim1 `TupRpair` paramsRinput - executeOp km1 gamma aenv dim1 sh' paramsR1 (tmp, manifest input) - rec tmp + cleanup <- executeOp km1 gamma aenv dim1 sh' paramsR1 (tmp, manifest input) + rec tmp cleanup -- return future @@ -334,8 +335,8 @@ result <- allocateRemote repr sh -- let paramsR = TupRsingle (ParamRarray repr) `TupRpair` TupRsingle (ParamRmaybe $ ParamRarray $ ArrayR (ShapeRsnoc shr) tp) - executeOp (ptxExecutable !# "fold") gamma aenv shr sh paramsR (result, manifest input) - put future result + cleanup <- executeOp (ptxExecutable !# "fold") gamma aenv shr sh paramsR (result, manifest input) + putCleanup future cleanup result return future @@ -368,8 +369,8 @@ future <- new result <- allocateRemote repr (sh, n) let paramsR = TupRsingle (ParamRarray repr) `TupRpair` TupRsingle (ParamRmaybe $ ParamRarray repr) `TupRpair` TupRsingle (ParamRmaybe $ ParamRarray reprSeg) - executeOp foldseg gamma aenv dim1 ((), m) paramsR ((result, manifest input), manifest segments) - put future result + cleanup <- executeOp foldseg gamma aenv dim1 ((), m) paramsR ((result, manifest input), manifest segments) + putCleanup future cleanup result return future @@ -396,9 +397,10 @@ -> Val aenv -> Delayed (Array (sh, Int) e) -> Par PTX (Future (Array (sh, Int) e)) -scan1Op repr exe gamma aenv input@(delayedShape -> (_, n)) - = boundsCheck "empty array" (n > 0) - $ scanCore repr exe gamma aenv n input +scan1Op repr exe gamma aenv input@(delayedShape -> sh@(_, n)) = + case n of + 0 -> newFull =<< allocateRemote repr sh + _ -> scanCore repr exe gamma aenv n input {-# INLINE scanCore #-} scanCore @@ -449,15 +451,20 @@ -- which can be computed by a single thread block will require no -- additional work. tmp <- allocateRemote repr ((), s) - executeOp k1 gamma aenv dim1 ((), s) paramsR1 ((tmp, result), manifest input) + cleanup1 <- executeOp k1 gamma aenv dim1 ((), s) paramsR1 ((tmp, result), manifest input) -- Step 2: Multi-block reductions need to compute the per-block prefix, -- then apply those values to the partial results. - when (s > 1) $ do - executeOp k2 gamma aenv dim1 ((), s) paramR tmp - executeOp k3 gamma aenv dim1 ((), s-1) paramsR3 ((tmp, result), c) + cleanup2 <- + if s > 1 + then do + cleanup2a <- executeOp k2 gamma aenv dim1 ((), s) paramR tmp + cleanup2b <- executeOp k3 gamma aenv dim1 ((), s-1) paramsR3 ((tmp, result), c) + return (cleanup2a >> cleanup2b) + else + return (return ()) - put future result + putCleanup future (cleanup1 >> cleanup2) result return future {-# INLINE scanDimOp #-} @@ -476,8 +483,8 @@ future <- new result <- allocateRemote repr (sz, m) let paramsR = TupRsingle (ParamRarray repr) `TupRpair` TupRsingle (ParamRmaybe $ ParamRarray repr) - executeOp (ptxExecutable !# "scan") gamma aenv dim1 ((), size shr' sz) paramsR (result, manifest input) - put future result + cleanup <- executeOp (ptxExecutable !# "scan") gamma aenv dim1 ((), size shr' sz) paramsR (result, manifest input) + putCleanup future cleanup result return future @@ -547,7 +554,7 @@ -- Step 1: independent thread-block-wide scans. Each block stores its partial -- sum to a temporary array. let paramsR1 = paramRdim1 `TupRpair` paramRdim1 `TupRpair` TupRsingle (ParamRmaybe $ ParamRarray repr) - executeOp k1 gamma aenv dim1 ((), s) paramsR1 ((tmp, result), manifest input) + cleanup1 <- executeOp k1 gamma aenv dim1 ((), s) paramsR1 ((tmp, result), manifest input) -- If this was a small array that was processed by a single thread block then -- we are done, otherwise compute the per-block prefix and apply those values @@ -555,15 +562,15 @@ if s == 1 then case tmp of - Array _ ad -> put future (result, Array () ad) + Array _ ad -> putCleanup future cleanup1 (result, Array () ad) else do sums <- allocateRemote (ArrayR dim0 tp) () let paramsR2 = paramRdim1 `TupRpair` paramRdim0 let paramsR3 = paramRdim1 `TupRpair` paramRdim1 `TupRpair` TupRsingle ParamRint - executeOp k2 gamma aenv dim1 ((), s) paramsR2 (tmp, sums) - executeOp k3 gamma aenv dim1 ((), s-1) paramsR3 ((tmp, result), c) - put future (result, sums) + cleanup2 <- executeOp k2 gamma aenv dim1 ((), s) paramsR2 (tmp, sums) + cleanup3 <- executeOp k3 gamma aenv dim1 ((), s-1) paramsR3 ((tmp, result), c) + putCleanup future (cleanup1 >> cleanup2 >> cleanup3) (result, sums) -- return future @@ -582,8 +589,8 @@ result <- allocateRemote repr sh sums <- allocateRemote (reduceRank repr) sz let paramsR = TupRsingle (ParamRarray repr) `TupRpair` TupRsingle (ParamRarray $ reduceRank repr) `TupRpair` TupRsingle (ParamRmaybe $ ParamRarray repr) - executeOp (ptxExecutable !# "scan") gamma aenv dim1 ((), size shr' sz) paramsR ((result, sums), manifest input) - put future (result, sums) + cleanup <- executeOp (ptxExecutable !# "scan") gamma aenv dim1 ((), size shr' sz) paramsR ((result, sums), manifest input) + putCleanup future cleanup (result, sums) return future @@ -617,7 +624,10 @@ then Debug.trace Debug.dump_exec "exec: permute/inplace" $ return defaults else Debug.trace Debug.dump_exec "exec: permute/clone" $ get =<< cloneArrayAsync repr' defaults -- - case kernelName kernel of + let kernelName' = + let kn = kernelName kernel + in SE.take (S.length kn - 65) kn + cleanup <- case kernelName' of -- execute directly using atomic operations "permute_rmw" -> let paramsR = paramR' `TupRpair` paramR @@ -635,7 +645,7 @@ _ -> internalError "unexpected kernel image" -- - put future result + putCleanup future cleanup result return future @@ -703,32 +713,38 @@ -- future <- new result <- allocateRemote repr shOut - parent <- asks ptxStream + parent <- asksParState ptxStream + parentStartPoint <- liftPar (Event.waypoint parent) -- interior (no bounds checking) let paramsRinside = TupRsingle (ParamRshape shr) `TupRpair` TupRsingle (ParamRarray repr) `TupRpair` paramsR - executeOp inside gamma aenv shr shIn paramsRinside ((shIn, result), params) + cleanup1 <- executeOp inside gamma aenv shr shIn paramsRinside ((shIn, result), params) -- halo regions (bounds checking) -- executed in separate streams so that they might overlap the main stencil -- and each other, as individually they will not saturate the device forM_ (stencilBorders (arrayRshape repr) shOut halo) $ \(u, v) -> fork $ do + -- synchronise with start of stencil computation, so that the arguments + -- are available + child <- asksParState ptxStream + liftIO (Event.after parentStartPoint child) + -- launch in a separate stream let sh = trav (-) v u let paramsRborder = TupRsingle (ParamRshape shr) `TupRpair` TupRsingle (ParamRshape shr) `TupRpair` TupRsingle (ParamRarray repr) `TupRpair` paramsR - executeOp border gamma aenv shr sh paramsRborder (((u, sh), result), params) + cleanup2 <- executeOp border gamma aenv shr sh paramsRborder (((u, sh), result), params) + addCleanup future cleanup2 - -- synchronisation with main stream - child <- asks ptxStream + -- make remainder of the parent stream depend on the border results event <- liftPar (Event.waypoint child) ready <- liftIO (Event.query event) if ready then return () else liftIO (Event.after event parent) - put future result + putCleanup future cleanup1 result return future -- Compute the stencil border regions, where we may need to evaluate the @@ -772,17 +788,14 @@ -> as -> Par PTX (Future bs) aforeignOp name _ _ asm arr = do - stream <- asks ptxStream + stream <- asksParState ptxStream Debug.monitorProcTime query msg (Just (unsafeGetValue stream)) (asm arr) where - query = if Debug.monitoringIsEnabled + msg = Debug.traceM Debug.dump_exec ("exec: " % string % " " % Debug.elapsed) name + query = if Debug.debuggingIsEnabled then return True else liftIO $ Debug.getFlag Debug.dump_exec - msg wall cpu gpu = do - Debug.addProcessorTime Debug.PTX gpu - Debug.traceIO Debug.dump_exec $ - printf "exec: %s %s" name (Debug.elapsed wall cpu gpu) -- Skeleton execution @@ -792,12 +805,12 @@ -- (!#) :: HasCallStack => FunctionTable -> ShortByteString -> Kernel (!#) exe name - = fromMaybe (internalError ("function not found: " ++ unpack name)) + = fromMaybe (internalError ("function not found: " % string) (unpack name)) $ lookupKernel name exe lookupKernel :: ShortByteString -> FunctionTable -> Maybe Kernel lookupKernel name ptxExecutable = - find (\k -> kernelName k == name) (functionTable ptxExecutable) + find (\k -> let n = kernelName k in SE.take (S.length n - 65) n == name) (functionTable ptxExecutable) delayedShape :: Delayed (Array sh e) -> sh delayedShape (Delayed sh) = sh @@ -811,7 +824,7 @@ -- withExecutable :: HasCallStack => ExecutableR PTX -> (FunctionTable -> Par PTX b) -> Par PTX b withExecutable PTXR{..} f = - local (\(s,_) -> (s,Just ptxExecutable)) $ do + localParState (\(s,_) -> (s,Just ptxExecutable)) $ do r <- f (unsafeGetValue ptxExecutable) liftIO $ touchLifetime ptxExecutable return r @@ -828,13 +841,17 @@ -> sh -> ParamsR PTX params -> params - -> Par PTX () + -> Par PTX (IO ()) executeOp kernel gamma aenv shr sh paramsR params = let n = size shr sh - in when (n > 0) $ do - stream <- asks ptxStream - argv <- marshalParams' @PTX (paramsR `TupRpair` TupRsingle (ParamRenv gamma)) (params, aenv) - liftIO $ launch kernel stream n $ DL.toList argv + in if n > 0 + then do + stream <- asksParState ptxStream + (argv, cleanup) <- marshalParams' @PTX (paramsR `TupRpair` TupRsingle (ParamRenv gamma)) (params, aenv) + liftIO $ launch kernel stream n $ DL.toList argv + return cleanup + else + return (return ()) -- Execute a device function with the given thread configuration and function @@ -851,14 +868,15 @@ smem = kernelSharedMemBytes -- Debugging/monitoring support - query = if Debug.monitoringIsEnabled + query = if Debug.debuggingIsEnabled then return True else Debug.getFlag Debug.dump_exec fst3 (x,_,_) = x msg wall cpu gpu = do - Debug.addProcessorTime Debug.PTX gpu - Debug.traceIO Debug.dump_exec $ - printf "exec: %s <<< %d, %d, %d >>> %s" - (unpack kernelName) (fst3 grid) (fst3 cta) smem (Debug.elapsed wall cpu gpu) + verbose <- Debug.getFlag Debug.verbose + let kernelName' | verbose = kernelName + | otherwise = SE.take (S.length kernelName - 65) kernelName + Debug.traceM Debug.dump_exec ("exec: " % string % " <<< " % int % ", " % int % ", " % int % " >>> " % Debug.elapsed) + (unpack kernelName') (fst3 grid) (fst3 cta) smem wall cpu gpu diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs b/src/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs @@ -2,6 +2,7 @@ {-# LANGUAGE GeneralizedNewtypeDeriving #-} {-# LANGUAGE LambdaCase #-} {-# LANGUAGE MultiParamTypeClasses #-} +{-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE TypeFamilies #-} {-# LANGUAGE TypeSynonymInstances #-} {-# OPTIONS_GHC -fno-warn-orphans #-} @@ -22,7 +23,6 @@ ) where --- accelerate import Data.Array.Accelerate.Error import Data.Array.Accelerate.Lifetime @@ -36,8 +36,6 @@ import qualified Data.Array.Accelerate.LLVM.PTX.Execute.Event as Event import qualified Data.Array.Accelerate.LLVM.PTX.Execute.Stream as Stream --- standard library -import Control.Monad.State import Control.Monad.Reader import Data.IORef @@ -68,32 +66,45 @@ data IVar a = Full !a - | Pending {-# UNPACK #-} !Event !(Maybe (Lifetime FunctionTable)) !a - | Empty + | Pending {-# UNPACK #-} !Event !(IO ()) !a + | Empty !(IO ()) +askParState :: Par PTX ParState +askParState = Par ask + +asksParState :: (ParState -> a) -> Par PTX a +asksParState f = Par (asks f) + +localParState :: (ParState -> ParState) -> Par PTX a -> Par PTX a +localParState f (Par m) = Par (local f m) + +instance MonadReader PTX (Par PTX) where + ask = Par (lift ask) + local f (Par (ReaderT g)) = Par (ReaderT (\parstate -> local f (g parstate))) + instance Async PTX where type FutureR PTX = Future newtype Par PTX a = Par { runPar :: ReaderT ParState (LLVM PTX) a } - deriving ( Functor, Applicative, Monad, MonadIO, MonadReader ParState, MonadState PTX ) + deriving ( Functor, Applicative, Monad, MonadIO ) {-# INLINEABLE new #-} {-# INLINEABLE newFull #-} - new = Future <$> liftIO (newIORef Empty) + new = Future <$> liftIO (newIORef (Empty (return ()))) newFull v = Future <$> liftIO (newIORef (Full v)) {-# INLINEABLE spawn #-} spawn m = do s' <- liftPar Stream.create - r <- local (const (s', Nothing)) m + r <- localParState (const (s', Nothing)) m liftIO (Stream.destroy s') return r {-# INLINEABLE fork #-} fork m = do s' <- liftPar (Stream.create) - () <- local (const (s', Nothing)) m + () <- localParState (const (s', Nothing)) m liftIO (Stream.destroy s') -- When we call 'put' the actual work may not have been evaluated yet; get @@ -102,13 +113,16 @@ -- {-# INLINEABLE put #-} put (Future ref) v = do - stream <- asks ptxStream - kernel <- asks ptxKernel + stream <- asksParState ptxStream + kernel <- asksParState ptxKernel event <- liftPar (Event.waypoint stream) ready <- liftIO (Event.query event) - liftIO . modifyIORef' ref $ \case - Empty -> if ready then Full v - else Pending event kernel v + let cleanupK = case kernel of + Just k -> touchLifetime k + Nothing -> return () + liftIO . atomicModifyIORef' ref $ \case + Empty cleanup -> if ready then (Full v, ()) + else (Pending event (cleanup >> cleanupK) v, ()) _ -> internalError "multiple put" -- Get the value of Future. Since the actual cross-stream synchronisation @@ -118,23 +132,21 @@ -- {-# INLINEABLE get #-} get (Future ref) = do - stream <- asks ptxStream + stream <- asksParState ptxStream liftIO $ do ivar <- readIORef ref case ivar of Full v -> return v - Pending event k v -> do + Pending event cleanup v -> do ready <- Event.query event if ready then do writeIORef ref (Full v) - case k of - Just f -> touchLifetime f - Nothing -> return () + cleanup else Event.after event stream return v - Empty -> internalError "blocked on an IVar" + Empty _ -> internalError "blocked on an IVar" {-# INLINEABLE block #-} block = liftIO . wait @@ -152,12 +164,34 @@ ivar <- readIORef ref case ivar of Full v -> return v - Pending event k v -> do + Pending event cleanup v -> do Event.block event writeIORef ref (Full v) - case k of - Just f -> touchLifetime f - Nothing -> return () + cleanup return v - Empty -> internalError "blocked on an IVar" + Empty _ -> internalError "blocked on an IVar" + +{-# INLINEABLE putCleanup #-} +putCleanup :: HasCallStack => FutureR PTX a -> IO () -> a -> Par PTX () +putCleanup (Future ref) cleanup v = do + stream <- asksParState ptxStream + kernel <- asksParState ptxKernel + event <- liftPar (Event.waypoint stream) + ready <- liftIO (Event.query event) + let cleanupK = case kernel of + Just k -> touchLifetime k + Nothing -> return () + liftIO . atomicModifyIORef' ref $ \case + Empty cleanup2 -> if ready then (Full v, ()) + else (Pending event (cleanup2 >> cleanup >> cleanupK) v, ()) + _ -> internalError "multiple put" + +{-# INLINEABLE addCleanup #-} +addCleanup :: HasCallStack => FutureR PTX a -> IO () -> Par PTX () +addCleanup (Future ref) cleanup = liftIO $ do + toRunNow <- atomicModifyIORef' ref $ \case + Full v -> (Full v, cleanup) + Pending event cleanup2 v -> (Pending event (cleanup2 >> cleanup) v, return ()) + Empty cleanup2 -> (Empty (cleanup2 >> cleanup), return ()) + toRunNow diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs b/src/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs @@ -1,4 +1,5 @@ -{-# LANGUAGE NamedFieldPuns #-} +{-# LANGUAGE NamedFieldPuns #-} +{-# LANGUAGE OverloadedStrings #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Event -- Copyright : [2014..2020] The Accelerate Team @@ -16,23 +17,25 @@ ) where --- accelerate import Data.Array.Accelerate.Lifetime import qualified Data.Array.Accelerate.Array.Remote.LRU as Remote import Data.Array.Accelerate.LLVM.PTX.Array.Remote ( ) +import qualified Data.Array.Accelerate.LLVM.PTX.Context as Context import Data.Array.Accelerate.LLVM.PTX.Target ( PTX(..) ) import Data.Array.Accelerate.LLVM.State import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug import {-# SOURCE #-} Data.Array.Accelerate.LLVM.PTX.Execute.Stream --- cuda import Foreign.CUDA.Driver.Error import qualified Foreign.CUDA.Driver.Event as Event import qualified Foreign.CUDA.Driver.Stream as Stream import Control.Exception -import Control.Monad.State +import Control.Monad +import Control.Monad.Reader +import Data.Text.Lazy.Builder +import Formatting -- | Events can be used for efficient device-side synchronisation between @@ -47,15 +50,18 @@ {-# INLINEABLE create #-} create :: LLVM PTX Event create = do + ctx <- asks ptxContext e <- create' event <- liftIO $ newLifetime e - liftIO $ addFinalizer event $ do message $ "destroy " ++ showEvent e - Event.destroy e + liftIO $ addFinalizer event $ do + message ("destroy " % formatEvent) e + Context.contextFinalizeResource ctx $ + Event.destroy e return event create' :: LLVM PTX Event.Event create' = do - PTX{ptxMemoryTable} <- gets llvmTarget + PTX{ptxMemoryTable} <- asks llvmTarget me <- attempt "create/new" (liftIO . catchOOM $ Event.create [Event.DisableTiming]) `orElse` do Remote.reclaim ptxMemoryTable @@ -75,12 +81,12 @@ ExitCode OutOfMemory -> return Nothing _ -> throwIO e - attempt :: MonadIO m => String -> m (Maybe a) -> m (Maybe a) + attempt :: MonadIO m => Builder -> m (Maybe a) -> m (Maybe a) attempt msg ea = do ma <- ea case ma of Nothing -> return Nothing - Just a -> do liftIO (message msg) + Just a -> do message builder msg return (Just a) orElse :: MonadIO m => m (Maybe a) -> m (Maybe a) -> m (Maybe a) @@ -107,7 +113,7 @@ liftIO $ withLifetime stream $ \s -> do withLifetime event $ \e -> do - message $ "add waypoint " ++ showEvent e ++ " in stream " ++ showStream s + message ("add waypoint " % formatEvent % " in stream " % formatStream) e s Event.record e (Just s) return event @@ -119,7 +125,7 @@ after event stream = withLifetime stream $ \s -> withLifetime event $ \e -> do - message $ "after " ++ showEvent e ++ " in stream " ++ showStream s + message ("after " % formatEvent % " in stream " % formatStream) e s Event.wait e (Just s) [] -- | Block the calling thread until the event is recorded @@ -128,7 +134,7 @@ block :: Event -> IO () block event = withLifetime event $ \e -> do - message $ "blocked on event " ++ showEvent e + message ("blocked on event " % formatEvent) e Event.block e -- | Test whether an event has completed @@ -141,21 +147,15 @@ -- Debug -- ----- -{-# INLINE trace #-} -trace :: String -> IO a -> IO a -trace msg next = do - Debug.traceIO Debug.dump_sched ("event: " ++ msg) - next - {-# INLINE message #-} -message :: String -> IO () -message s = s `trace` return () +message :: MonadIO m => Format (m ()) a -> a +message fmt = Debug.traceM Debug.dump_sched ("event: " % fmt) -{-# INLINE showEvent #-} -showEvent :: Event.Event -> String -showEvent (Event.Event e) = show e +{-# INLINE formatEvent #-} +formatEvent :: Format r (Event.Event -> r) +formatEvent = later $ \(Event.Event e) -> bformat shown e -{-# INLINE showStream #-} -showStream :: Stream.Stream -> String -showStream (Stream.Stream s) = show s +{-# INLINE formatStream #-} +formatStream :: Format r (Stream.Stream -> r) +formatStream = later $ \(Stream.Stream s) -> bformat shown s diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs b/src/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs @@ -37,27 +37,72 @@ import qualified Foreign.CUDA.Driver as CUDA +import Control.Concurrent +import Control.Monad.IO.Class (liftIO) import qualified Data.DList as DL instance Marshal PTX where type ArgR PTX = CUDA.FunParam + type MarshalCleanup PTX = IO () marshalInt = CUDA.VArg marshalScalarData' t | SingleArrayDict <- singleArrayDict t - = liftPar . fmap (DL.singleton . CUDA.VArg) . unsafeGetDevicePtr t + = liftPar . fmap (\(ptr, cleanup) -> (DL.singleton (CUDA.VArg ptr), cleanup)) . getCudaDevicePtr t --- TODO FIXME !!! +-- | Return the CUDA device pointer corresponding to the given array, as well +-- as a cleanup IO action that __MUST__ be run once you are done with the +-- pointer (i.e. the GPU kernel has completed). Not calling the cleanup action +-- will result in leaked memory and resources. Calling the action twice will +-- block indefinitely on an MVar. -- --- We will probably need to change marshal to be a bracketed function, so that --- the garbage collector does not try to evict the array in the middle of --- a computation. +-- This function is a hack. Prim.withDevicePtr is intended to be a wrapping +-- function that retains the resource while the callback is running and +-- releases it when the callback returns. This is all nice, but since the PTX +-- Accelerate runtime is asynchronous, uses of withDevicePtr would not all be +-- neatly nested: the actual array lifetimes are haphazard intervals during +-- program execution. -- -unsafeGetDevicePtr +-- Originally, this function just gave up and extracted the DevicePtr by +-- calling withDevicePtr with a trivial body that simply leaks p; this is +-- unsound (which was acknowledged by a 'fixme' comment...) and appears to have +-- been the cause of silent incorrect results (!) on a GTX 1050 Ti on the +-- adbench-gmmgrad test in accelerate-tests [1]. +-- +-- [1]: https://github.com/tomsmeding/accelerate-tests/blob/master/src/Data/Array/Accelerate/Tests/Prog/ADBenchGMMGrad.hs +-- +-- Fortunately, it turns out that the MemoryTable implementation underlying +-- withDevicePtr does not in fact assume lexical nesting of array usages. Thus +-- we can use a hack to let the callback of withDevicePtr live for the correct +-- amount of time without needing to rearchitect the entire PTX backend: let +-- the call run in a forkIO thread and use MVars to communicate when it should +-- return. This means that we now have the possibility to return a +-- self-contained "cleanup" handler from getCudaDevicePtr that does nothing but +-- signal to the withDevicePtr callback that the array's lifetime has ended and +-- the scope can close. All this is possible because the 'LLVM' monad is just a +-- reader monad over IO, so we can unlift it into IO. +-- +-- As a final, questionable improvement, we let the "cleanup" handler wait +-- until withDevicePtr has properly returned so that we know the array's +-- refcount has been properly decremented and memory has been released if +-- possible. +getCudaDevicePtr :: SingleType e -> ArrayData e - -> LLVM PTX (CUDA.DevicePtr (ScalarArrayDataR e)) -unsafeGetDevicePtr !t !ad = - Prim.withDevicePtr t ad (\p -> return (Nothing, p)) + -> LLVM PTX (CUDA.DevicePtr (ScalarArrayDataR e), IO ()) +getCudaDevicePtr !t !ad = do + ptrVar <- liftIO newEmptyMVar + doneVar <- liftIO newEmptyMVar + releasedVar <- liftIO newEmptyMVar + + _ <- unliftIOLLVM $ \inLLVM -> forkIO $ inLLVM $ do + Prim.withDevicePtr t ad $ \p -> liftIO $ do + putMVar ptrVar p + takeMVar doneVar + return (Nothing, ()) + liftIO $ putMVar releasedVar () + + ptr <- liftIO $ readMVar ptrVar + return (ptr, putMVar doneVar () >> readMVar releasedVar) diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs b/src/Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs @@ -1,6 +1,7 @@ -{-# LANGUAGE BangPatterns #-} -{-# LANGUAGE MagicHash #-} -{-# LANGUAGE RecordWildCards #-} +{-# LANGUAGE BangPatterns #-} +{-# LANGUAGE MagicHash #-} +{-# LANGUAGE OverloadedStrings #-} +{-# LANGUAGE RecordWildCards #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Stream -- Copyright : [2014..2020] The Accelerate Team @@ -18,7 +19,6 @@ ) where --- accelerate import Data.Array.Accelerate.Lifetime import qualified Data.Array.Accelerate.Array.Remote.LRU as Remote @@ -30,13 +30,14 @@ import qualified Data.Array.Accelerate.LLVM.PTX.Execute.Event as Event import Data.Array.Accelerate.LLVM.PTX.Execute.Stream.Reservoir as RSV --- cuda import Foreign.CUDA.Driver.Error import qualified Foreign.CUDA.Driver.Stream as Stream --- standard library import Control.Exception -import Control.Monad.State +import Control.Monad +import Control.Monad.Reader +import Data.Text.Lazy.Builder +import Formatting -- | A 'Stream' represents an independent sequence of computations executed on @@ -110,7 +111,7 @@ {-# INLINEABLE create #-} create :: LLVM PTX Stream create = do - PTX{..} <- gets llvmTarget + PTX{..} <- asks llvmTarget s <- create' stream <- liftIO $ newLifetime s liftIO $ addFinalizer stream (RSV.insert ptxStreamReservoir s) @@ -118,7 +119,7 @@ create' :: LLVM PTX Stream.Stream create' = do - PTX{..} <- gets llvmTarget + PTX{..} <- asks llvmTarget ms <- attempt "create/reservoir" (liftIO $ RSV.malloc ptxStreamReservoir) `orElse` attempt "create/new" (liftIO . catchOOM $ Stream.create []) @@ -140,12 +141,12 @@ ExitCode OutOfMemory -> return Nothing _ -> throwIO e - attempt :: MonadIO m => String -> m (Maybe a) -> m (Maybe a) + attempt :: MonadIO m => Builder -> m (Maybe a) -> m (Maybe a) attempt msg ea = do ma <- ea case ma of Nothing -> return Nothing - Just a -> do liftIO (message msg) + Just a -> do message builder msg return (Just a) orElse :: MonadIO m => m (Maybe a) -> m (Maybe a) -> m (Maybe a) @@ -167,13 +168,7 @@ -- Debug -- ----- -{-# INLINE trace #-} -trace :: String -> IO a -> IO a -trace msg next = do - Debug.traceIO Debug.dump_sched ("stream: " ++ msg) - next - {-# INLINE message #-} -message :: String -> IO () -message s = s `trace` return () +message :: MonadIO m => Format (m ()) a -> a +message fmt = Debug.traceM Debug.dump_sched ("stream: " % fmt) diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs b/src/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs @@ -1,4 +1,5 @@ -{-# LANGUAGE BangPatterns #-} +{-# LANGUAGE BangPatterns #-} +{-# LANGUAGE OverloadedStrings #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Stream.Reservoir -- Copyright : [2016..2020] The Accelerate Team @@ -21,6 +22,7 @@ import Control.Concurrent.MVar import Data.Sequence ( Seq ) +import Formatting import qualified Data.Sequence as Seq import qualified Foreign.CUDA.Driver.Stream as Stream @@ -79,24 +81,18 @@ {-# INLINEABLE insert #-} insert :: Reservoir -> Stream.Stream -> IO () insert !ref !stream = do - message ("stash stream " ++ showStream stream) + message ("stash stream " % formatStream) stream modifyMVar_ ref $ \rsv -> return (rsv Seq.|> stream) -- Debug -- ----- -{-# INLINE trace #-} -trace :: String -> IO a -> IO a -trace msg next = do - Debug.traceIO Debug.dump_sched ("stream: " ++ msg) - next - {-# INLINE message #-} -message :: String -> IO () -message s = s `trace` return () +message :: Format (IO ()) a -> a +message fmt = Debug.traceM Debug.dump_sched ("stream: " % fmt) -{-# INLINE showStream #-} -showStream :: Stream.Stream -> String -showStream (Stream.Stream s) = show s +{-# INLINE formatStream #-} +formatStream :: Format r (Stream.Stream -> r) +formatStream = later $ \(Stream.Stream s) -> bformat shown s diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Link.hs b/src/Data/Array/Accelerate/LLVM/PTX/Link.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Link.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Link.hs @@ -1,7 +1,7 @@ -{-# LANGUAGE CPP #-} -{-# LANGUAGE RecordWildCards #-} -{-# LANGUAGE TypeFamilies #-} -{-# OPTIONS_GHC -fno-warn-orphans #-} +{-# LANGUAGE OverloadedStrings #-} +{-# LANGUAGE RecordWildCards #-} +{-# LANGUAGE TypeFamilies #-} +{-# OPTIONS_GHC -Wno-orphans #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Link -- Copyright : [2017..2020] The Accelerate Team @@ -33,16 +33,15 @@ import Data.Array.Accelerate.LLVM.PTX.Target import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug --- cuda import qualified Foreign.CUDA.Analysis as CUDA import qualified Foreign.CUDA.Driver as CUDA --- standard library -import Control.Monad.State +import Control.Monad.Reader import Data.ByteString.Short.Char8 ( ShortByteString, unpack ) +import Formatting import Foreign.Ptr -import Language.Haskell.TH -import Text.Printf ( printf ) +import Data.Array.Accelerate.TH.Compat +import qualified Data.ByteString as B import qualified Data.ByteString.Unsafe as B import Prelude as P hiding ( lookup ) @@ -56,11 +55,12 @@ -- | Load the generated object code into the current CUDA context. -- link :: ObjectR PTX -> LLVM PTX (ExecutableR PTX) -link (ObjectR uid cfg obj) = do - target <- gets llvmTarget - cache <- gets ptxKernelTable +link (ObjectR uid cfg objFname) = do + target <- asks llvmTarget + cache <- asks ptxKernelTable funs <- liftIO $ dlsym uid cache $ do -- Load the SASS object code into the current CUDA context + obj <- B.readFile objFname jit <- B.unsafeUseAsCString obj $ \p -> CUDA.loadDataFromPtrEx (castPtr p) [] let mdl = CUDA.jitModule jit @@ -70,8 +70,9 @@ -- Finalise the module by unloading it from the CUDA context addFinalizer oc $ do - Debug.traceIO Debug.dump_ld ("ld: unload module: " ++ show nm) - withContext (ptxContext target) (CUDA.unload mdl) + Debug.traceM Debug.dump_ld ("ld: unload module: " % formatFunctionTable) nm + contextFinalizeResource (ptxContext target) $ + withContext (ptxContext target) (CUDA.unload mdl) return (nm, oc) -- @@ -95,7 +96,7 @@ :: CUDA.Module -> ShortByteString -> LaunchConfig - -> IO (Kernel, Q (TExp (Int -> Int))) + -> IO (Kernel, CodeQ (Int -> Int)) linkFunctionQ mdl name configure = do f <- CUDA.getFun mdl name regs <- CUDA.requires f CUDA.NumRegs @@ -107,17 +108,16 @@ let (occ, cta, grid, dsmem, gridQ) = configure maxt regs ssmem - msg1, msg2 :: String - msg1 = printf "kernel function '%s' used %d registers, %d bytes smem, %d bytes lmem, %d bytes cmem" + msg1 = bformat ("kernel function " % squoted string % " used " % int % " registers, " % int % " bytes smem, " % int % " bytes lmem, " % int % " bytes cmem") (unpack name) regs (ssmem + dsmem) lmem cmem - msg2 = printf "multiprocessor occupancy %.1f %% : %d threads over %d warps in %d blocks" + msg2 = bformat ("multiprocessor occupancy " % fixed 1 % "% : " % int % " threads over " % int % " warps in " % int % " blocks") (CUDA.occupancy100 occ) (CUDA.activeThreads occ) (CUDA.activeWarps occ) (CUDA.activeThreadBlocks occ) - Debug.traceIO Debug.dump_cc (printf "cc: %s\n %s" msg1 msg2) + Debug.traceM Debug.dump_cc ("cc: " % builder % "\n " % builder) msg1 msg2 return (Kernel name f dsmem cta grid, gridQ) diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Link/Object.hs b/src/Data/Array/Accelerate/LLVM/PTX/Link/Object.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Link/Object.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Link/Object.hs @@ -14,6 +14,7 @@ import Data.Array.Accelerate.Lifetime import Data.ByteString.Short.Char8 ( ShortByteString, unpack ) import Data.List +import Formatting import qualified Foreign.CUDA.Driver as CUDA @@ -34,6 +35,10 @@ = showString "<<" . showString (intercalate "," [ unpack (kernelName k) | k <- functionTable f ]) . showString ">>" + +formatFunctionTable :: Format r (FunctionTable -> r) +formatFunctionTable = later $ \f -> + bformat (angled (angled (commaSep string))) [ unpack (kernelName k) | k <- functionTable f ] -- | Object code consists of executable code in the device address space -- diff --git a/src/Data/Array/Accelerate/LLVM/PTX/State.hs b/src/Data/Array/Accelerate/LLVM/PTX/State.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/State.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/State.hs @@ -1,4 +1,4 @@ -{-# LANGUAGE CPP #-} +{-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE ScopedTypeVariables #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.State @@ -33,15 +33,17 @@ import qualified Data.Array.Accelerate.LLVM.PTX.Link.Cache as LC import qualified Data.Array.Accelerate.LLVM.PTX.Pool as Pool +import Foreign.CUDA.Driver.Error +import qualified Foreign.CUDA.Driver as CUDA +import qualified Foreign.CUDA.Driver.Context as Context + +import Control.Concurrent import Control.Exception ( try, catch ) import Data.Maybe ( fromMaybe, catMaybes ) +import Formatting import System.Environment ( lookupEnv ) import System.IO.Unsafe ( unsafePerformIO, unsafeInterleaveIO ) -import Text.Printf ( printf ) import Text.Read ( readMaybe ) -import Foreign.CUDA.Driver.Error -import qualified Foreign.CUDA.Driver as CUDA -import qualified Foreign.CUDA.Driver.Context as Context -- | Execute a PTX computation @@ -50,7 +52,7 @@ evalPTX ptx acc = CT.withContext (ptxContext ptx) (evalLLVM ptx acc) `catch` - \e -> internalError (show (e :: CUDAException)) + \e -> internalError shown (e :: CUDAException) -- | Create a new PTX execution target for the given device @@ -125,7 +127,9 @@ -- {-# NOINLINE defaultTarget #-} defaultTarget :: PTX -defaultTarget = head (unmanaged defaultTargetPool) +defaultTarget = case unmanaged defaultTargetPool of + ptx : _ -> ptx + _ -> error "impossible" -- ensured by defaultTargetPool -- | Create a shared resource pool of the available CUDA devices. -- @@ -137,7 +141,7 @@ {-# NOINLINE defaultTargetPool #-} defaultTargetPool :: Pool PTX defaultTargetPool = unsafePerformIO $! do - Debug.traceIO Debug.dump_gc "gc: initialise default PTX pool" + Debug.traceM Debug.dump_gc "gc: initialise default PTX pool" CUDA.initialise [] -- Figure out which GPUs we should put into the execution pool @@ -150,14 +154,14 @@ -- Spin up the GPU at the given ordinal. -- boot :: Int -> IO (Maybe PTX) - boot i = unsafeInterleaveIO $ do + boot i = unsafeInterleaveIO $ runInBoundThread $ do dev <- CUDA.device i prp <- CUDA.props dev r <- try $ createTargetForDevice dev prp [CUDA.SchedAuto] case r of Right ptx -> return (Just ptx) Left (e::CUDAException) -> do - Debug.traceIO Debug.dump_gc (printf "gc: failed to initialise device %d: %s" i (show e)) + Debug.traceM Debug.dump_gc ("gc: failed to initialise device " % int % ": " % shown) i e return Nothing -- Create the pool from the available devices, which get spun-up lazily as diff --git a/src/Data/Array/Accelerate/LLVM/PTX/Target.hs b/src/Data/Array/Accelerate/LLVM/PTX/Target.hs --- a/src/Data/Array/Accelerate/LLVM/PTX/Target.hs +++ b/src/Data/Array/Accelerate/LLVM/PTX/Target.hs @@ -1,4 +1,3 @@ -{-# LANGUAGE CPP #-} {-# LANGUAGE EmptyDataDecls #-} {-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE TypeApplications #-} @@ -20,15 +19,6 @@ ) where --- llvm-hs -import LLVM.AST.AddrSpace -import LLVM.AST.DataLayout -import LLVM.Target hiding ( Target ) -import qualified LLVM.Target as LLVM -import qualified LLVM.Relocation as R -import qualified LLVM.CodeModel as CM -import qualified LLVM.CodeGenOpt as CGO - -- accelerate import Data.Array.Accelerate.Error @@ -36,22 +26,18 @@ import Data.Array.Accelerate.LLVM.Target import Data.Array.Accelerate.LLVM.PTX.Array.Table ( MemoryTable ) -import Data.Array.Accelerate.LLVM.PTX.Context ( Context, deviceProperties ) +import Data.Array.Accelerate.LLVM.PTX.Context ( Context, deviceProperties, deviceName ) import Data.Array.Accelerate.LLVM.PTX.Execute.Stream.Reservoir ( Reservoir ) import Data.Array.Accelerate.LLVM.PTX.Link.Cache ( KernelTable ) -- CUDA -import Foreign.CUDA.Analysis.Device +import Foreign.CUDA.Analysis.Device ( DeviceProperties ) -- standard library -import Data.ByteString ( ByteString ) import Data.ByteString.Short ( ShortByteString ) -import Data.String -import Debug.Trace -import System.IO.Unsafe -import Text.Printf -import qualified Data.Map as Map -import qualified Data.Set as Set +import Data.Primitive.ByteArray +import Foreign.C.String +import Foreign.Ptr -- | The PTX execution target for NVIDIA GPUs. @@ -72,11 +58,7 @@ instance Target PTX where targetTriple = Just ptxTargetTriple -#if ACCELERATE_USE_NVVM - targetDataLayout = Nothing -- see note: [NVVM and target data layout] -#else - targetDataLayout = Just ptxDataLayout -#endif + targetDataLayout = Nothing -- | Extract the properties of the device the current PTX execution state is @@ -85,36 +67,10 @@ ptxDeviceProperties :: PTX -> DeviceProperties ptxDeviceProperties = deviceProperties . ptxContext - --- | A description of the various data layout properties that may be used during --- optimisation. For CUDA the following data layouts are supported: --- --- 32-bit: --- e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64 --- --- 64-bit: --- e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64 --- --- Thus, only the size of the pointer layout changes depending on the host --- architecture. +-- | Extract the name of the device of the current execution context -- -ptxDataLayout :: DataLayout -ptxDataLayout = DataLayout - { endianness = LittleEndian - , mangling = Nothing - , aggregateLayout = AlignmentInfo 0 64 - , stackAlignment = Nothing - , pointerLayouts = Map.fromList - [ (AddrSpace 0, (wordSize, AlignmentInfo wordSize wordSize)) ] - , typeLayouts = Map.fromList $ - [ ((IntegerAlign, 1), AlignmentInfo 8 8) ] ++ - [ ((IntegerAlign, i), AlignmentInfo i i) | i <- [8,16,32,64]] ++ - [ ((VectorAlign, v), AlignmentInfo v v) | v <- [16,32,64,128]] ++ - [ ((FloatAlign, f), AlignmentInfo f f) | f <- [32,64] ] - , nativeSizes = Just $ Set.fromList [ 16,32,64 ] - } - where - wordSize = bitSize (undefined :: Int) +ptxDeviceName :: PTX -> CString +ptxDeviceName = castPtr . byteArrayContents . deviceName . ptxContext -- | String that describes the target host. @@ -125,83 +81,3 @@ 32 -> "nvptx-nvidia-cuda" 64 -> "nvptx64-nvidia-cuda" _ -> internalError "I don't know what architecture I am" - - --- | Bracket creation and destruction of the NVVM TargetMachine. --- -withPTXTargetMachine - :: HasCallStack - => DeviceProperties - -> (TargetMachine -> IO a) - -> IO a -withPTXTargetMachine dev go = - let (sm, isa) = ptxTargetVersion (computeCapability dev) - in - withTargetOptions $ \options -> do - withTargetMachine - ptxTarget - ptxTargetTriple - sm -- CPU - (Map.singleton (CPUFeature isa) True) -- CPU features - options -- target options - R.Default -- relocation model - CM.Default -- code model - CGO.Default -- optimisation level - go - --- Compile using the earliest version of the SM target PTX ISA supported by --- the given compute device and this version of LLVM. --- --- Note that we require at least ptx40 for some libnvvm device functions. --- --- See table NVPTX supported processors: --- --- https://github.com/llvm-mirror/llvm/blob/master/lib/Target/NVPTX/NVPTX.td --- --- PTX ISA verison history: --- --- https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#release-notes --- -ptxTargetVersion :: Compute -> (ByteString, ByteString) -ptxTargetVersion compute@(Compute m n) -#if MIN_VERSION_llvm_hs(8,0,0) - | m >= 7 && n >= 5 = ("sm_75", "ptx63") -#endif -#if MIN_VERSION_llvm_hs(7,0,0) - | m >= 7 && n >= 2 = ("sm_72", "ptx61") -#endif -#if MIN_VERSION_llvm_hs(6,0,0) - | m >= 7 = ("sm_70", "ptx60") -#endif - | m > 6 = ("sm_62", "ptx50") -- fallthrough - -- - | m == 6 && n == 2 = ("sm_62", "ptx50") - | m == 6 && n == 1 = ("sm_61", "ptx50") - | m == 6 = ("sm_60", "ptx50") - | m == 5 && n == 3 = ("sm_53", "ptx42") - | m == 5 && n == 2 = ("sm_52", "ptx41") - | m == 5 = ("sm_50", "ptx40") - | m == 3 && n == 7 = ("sm_37", "ptx41") - | m == 3 && n == 5 = ("sm_35", "ptx40") - | m == 3 && n == 2 = ("sm_32", "ptx40") - | m == 3 = ("sm_30", "ptx40") - | m == 2 && n == 1 = ("sm_21", "ptx40") - | m == 2 = ("sm_20", "ptx40") - -- - | otherwise - = trace warning (fromString (printf "sm_%d%d" m n), "ptx40") - where - warning = unlines [ "*** Warning: Unhandled CUDA device compute capability: " ++ show compute - , "*** Please submit a bug report at https://github.com/AccelerateHS/accelerate/issues" ] - --- | The NVPTX target for this host. --- --- The top-level 'unsafePerformIO' is so that 'initializeAllTargets' is run once --- per program execution (although that might not be necessary?) --- -{-# NOINLINE ptxTarget #-} -ptxTarget :: LLVM.Target -ptxTarget = unsafePerformIO $ do - initializeAllTargets - fst `fmap` lookupTarget Nothing ptxTargetTriple - diff --git a/test/nofib/Data/Array/Accelerate/LLVM/PTX/NoFib/RunQ.hs b/test/nofib/Data/Array/Accelerate/LLVM/PTX/NoFib/RunQ.hs new file mode 100644 --- /dev/null +++ b/test/nofib/Data/Array/Accelerate/LLVM/PTX/NoFib/RunQ.hs @@ -0,0 +1,31 @@ +{-# LANGUAGE BangPatterns #-} +{-# LANGUAGE TemplateHaskell #-} +module Data.Array.Accelerate.LLVM.PTX.NoFib.RunQ where + +import qualified Data.Array.Accelerate as A +import qualified Data.Array.Accelerate.LLVM.PTX as GPU + +import Test.Tasty +import Test.Tasty.HUnit + + +-- WARNING: This module is duplicated (apart from Native/PTX) between the +-- accelerate-llvm-native and accelerate-llvm-ptx backends. This code is not +-- included in the main Accelerate nofib testsuite because of staging issues: +-- the test can only be defined after runQ is known, and runQ is only built +-- after the 'accelerate' package has already finished building. It would be +-- possible to deduplicate the little Accelerate program in there, but that was +-- not deemed worth the effort. + + +test_runq :: TestTree +test_runq = + testGroup "runQ" + [ testCase "simple" test_simple ] + +test_simple :: Assertion +test_simple = do + let prog :: A.Vector Int -> A.Scalar Int + !prog = $(GPU.runQ $ \a -> A.sum (A.map (+1) (a :: A.Acc (A.Vector Int)))) + let n = 10000 + prog (A.fromList (A.Z A.:. 10000) [1..]) @=? A.fromList A.Z [n * (n + 1) `div` 2 + n] diff --git a/test/nofib/Main.hs b/test/nofib/Main.hs --- a/test/nofib/Main.hs +++ b/test/nofib/Main.hs @@ -12,10 +12,8 @@ import Data.Array.Accelerate.Test.NoFib import Data.Array.Accelerate.LLVM.PTX -import Data.Array.Accelerate.Debug +import Data.Array.Accelerate.LLVM.PTX.NoFib.RunQ main :: IO () -main = do - beginMonitoring - nofib runN +main = nofib runN test_runq