accelerate-llvm-ptx 1.3.0.0 → 1.4.0.0
raw patch · 40 files changed
+1859/−1880 lines, 40 filesdep +formattingdep +prettyprinterdep +primitivedep −llvm-hsdep −llvm-hs-puredep −nvvmdep ~acceleratedep ~accelerate-llvmdep ~containerssetup-changednew-uploaderPVP ok
version bump matches the API change (PVP)
Dependencies added: formatting, prettyprinter, primitive, tasty, tasty-hunit, text
Dependencies removed: llvm-hs, llvm-hs-pure, nvvm
Dependency ranges changed: accelerate, accelerate-llvm, containers, pretty
API changes (from Hackage documentation)
+ Data.Array.Accelerate.LLVM.PTX.Foreign: [deviceFinalizerRefcount] :: Context -> {-# UNPACK #-} !IORef Int
+ Data.Array.Accelerate.LLVM.PTX.Foreign: [deviceName] :: Context -> {-# UNPACK #-} !ByteArray
+ Data.Array.Accelerate.LLVM.PTX.Foreign: addCleanup :: HasCallStack => FutureR PTX a -> IO () -> Par PTX ()
+ Data.Array.Accelerate.LLVM.PTX.Foreign: askParState :: Par PTX ParState
+ Data.Array.Accelerate.LLVM.PTX.Foreign: asksParState :: (ParState -> a) -> Par PTX a
+ Data.Array.Accelerate.LLVM.PTX.Foreign: localParState :: (ParState -> ParState) -> Par PTX a -> Par PTX a
+ Data.Array.Accelerate.LLVM.PTX.Foreign: putCleanup :: HasCallStack => FutureR PTX a -> IO () -> a -> Par PTX ()
- Data.Array.Accelerate.LLVM.PTX: class Afunction f
+ Data.Array.Accelerate.LLVM.PTX: class () => Afunction f
- Data.Array.Accelerate.LLVM.PTX: class Arrays a
+ Data.Array.Accelerate.LLVM.PTX: class () => Arrays a
- Data.Array.Accelerate.LLVM.PTX: data Acc a
+ Data.Array.Accelerate.LLVM.PTX: data () => Acc a
- Data.Array.Accelerate.LLVM.PTX: data Async a
+ Data.Array.Accelerate.LLVM.PTX: data () => Async a
- Data.Array.Accelerate.LLVM.PTX.Foreign: Context :: {-# UNPACK #-} !DeviceProperties -> {-# UNPACK #-} !Lifetime Context -> Context
+ Data.Array.Accelerate.LLVM.PTX.Foreign: Context :: {-# UNPACK #-} !DeviceProperties -> {-# UNPACK #-} !ByteArray -> {-# UNPACK #-} !Lifetime Context -> {-# UNPACK #-} !IORef Int -> Context
- Data.Array.Accelerate.LLVM.PTX.Foreign: Empty :: IVar a
+ Data.Array.Accelerate.LLVM.PTX.Foreign: Empty :: !IO () -> IVar a
- Data.Array.Accelerate.LLVM.PTX.Foreign: Pending :: {-# UNPACK #-} !Event -> !Maybe (Lifetime FunctionTable) -> !a -> IVar a
+ Data.Array.Accelerate.LLVM.PTX.Foreign: Pending :: {-# UNPACK #-} !Event -> !IO () -> !a -> IVar a
- Data.Array.Accelerate.LLVM.PTX.Foreign: data LLVM target a
+ Data.Array.Accelerate.LLVM.PTX.Foreign: data () => LLVM target a
Files
- CHANGELOG.md +29/−1
- README.md +110/−151
- Setup.hs +0/−2
- accelerate-llvm-ptx.cabal +29/−69
- src/Data/Array/Accelerate/LLVM/PTX.hs +52/−24
- src/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs +4/−4
- src/Data/Array/Accelerate/LLVM/PTX/Array/Data.hs +6/−5
- src/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs +37/−35
- src/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs +34/−45
- src/Data/Array/Accelerate/LLVM/PTX/Array/Table.hs +9/−10
- src/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs +10/−10
- src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs +421/−107
- src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs +87/−111
- src/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs +53/−68
- src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Generate.hs +5/−3
- src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Map.hs +5/−3
- src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Permute.hs +22/−19
- src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs +151/−171
- src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Stencil.hs +22/−19
- src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Transform.hs +5/−3
- src/Data/Array/Accelerate/LLVM/PTX/Compile.hs +183/−184
- src/Data/Array/Accelerate/LLVM/PTX/Compile/Cache.hs +7/−5
- src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs +0/−176
- src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/Load.hs +26/−92
- src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/TH.hs +0/−141
- src/Data/Array/Accelerate/LLVM/PTX/Context.hs +150/−41
- src/Data/Array/Accelerate/LLVM/PTX/Debug.hs +32/−18
- src/Data/Array/Accelerate/LLVM/PTX/Embed.hs +30/−25
- src/Data/Array/Accelerate/LLVM/PTX/Execute.hs +98/−80
- src/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs +59/−25
- src/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs +26/−26
- src/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs +54/−9
- src/Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs +14/−19
- src/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs +9/−13
- src/Data/Array/Accelerate/LLVM/PTX/Link.hs +19/−19
- src/Data/Array/Accelerate/LLVM/PTX/Link/Object.hs +5/−0
- src/Data/Array/Accelerate/LLVM/PTX/State.hs +14/−10
- src/Data/Array/Accelerate/LLVM/PTX/Target.hs +9/−133
- test/nofib/Data/Array/Accelerate/LLVM/PTX/NoFib/RunQ.hs +31/−0
- test/nofib/Main.hs +2/−4
CHANGELOG.md view
@@ -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
README.md view
@@ -3,21 +3,16 @@ # LLVM backends for the Accelerate array language -[](https://github.com/tmcdonell/accelerate-llvm/actions)+[](https://github.com/AccelerateHS/accelerate-llvm/actions/workflows/ci.yml) [](https://gitter.im/AccelerateHS/Lobby)-<br>-[](https://stackage.org/lts/package/accelerate-llvm)-[](https://stackage.org/nightly/package/accelerate-llvm) [](https://hackage.haskell.org/package/accelerate-llvm)-<br>-[](https://hub.docker.com/r/tmcdonell/accelerate-llvm/)-[](https://microbadger.com/images/tmcdonell/accelerate-llvm) </div> 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.+<details><summary>Optionally, add <code>clang</code> (and more) to your system path. Click to see how.</summary> -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:+——+</details> -```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.
− Setup.hs
@@ -1,2 +0,0 @@-import Distribution.Simple-main = defaultMain
accelerate-llvm-ptx.cabal view
@@ -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 <http://hackage.haskell.org/package/accelerate accelerate> package. . [/Dependencies/] . Haskell dependencies are available from Hackage. The following external- libraries are alse required:- .- * <http://llvm.org LLVM>- .- * <https://developer.nvidia.com/cuda-downloads CUDA>- .- [/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 <http://apt.llvm.org apt.llvm.org> 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- <http://releases.llvm.org/9.0.0/docs/CMake.html LLVM.org>. 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).+ * <https://clang.llvm.org/ clang> (not used to compile C code, but to compile generated LLVM IR via a mostly LLVM-version-independent interface)+ * <https://developer.nvidia.com/cuda-downloads CUDA> .+ For installation instructions, see the <https://github.com/AccelerateHS/accelerate-llvm#readme README>. -license: BSD3+license: BSD-3-Clause license-file: LICENSE author: Trevor L. McDonell maintainer: Trevor L. McDonell <trevor.mcdonell@gmail.com> 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
src/Data/Array/Accelerate/LLVM/PTX.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs view
@@ -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 ||]
src/Data/Array/Accelerate/LLVM/PTX/Array/Data.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs view
@@ -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)
src/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/Array/Table.hs view
@@ -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)
src/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs view
@@ -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 @@ -- <https://github.com/AccelerateHS/accelerate/issues/363> -- 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.+--+-- <https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#warp-shuffle-functions>+--+-- 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 <m x i32> 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 <m+1 x i32>: 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)
src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs view
@@ -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 ())
src/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Generate.hs view
@@ -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)
src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Map.hs view
@@ -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)
src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Permute.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs view
@@ -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 -- ---------
src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Stencil.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/CodeGen/Transform.hs view
@@ -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)
src/Data/Array/Accelerate/LLVM/PTX/Compile.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/Compile/Cache.hs view
@@ -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"
− src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs
@@ -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 <trevor.mcdonell@gmail.com>--- 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, .. }-
src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/Load.hs view
@@ -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"
− src/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/TH.hs
@@ -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 <trevor.mcdonell@gmail.com>--- 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 ]-
src/Data/Array/Accelerate/LLVM/PTX/Context.hs view
@@ -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.
src/Data/Array/Accelerate/LLVM/PTX/Debug.hs view
@@ -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)
src/Data/Array/Accelerate/LLVM/PTX/Embed.hs view
@@ -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))
src/Data/Array/Accelerate/LLVM/PTX/Execute.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs view
@@ -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)
src/Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs view
@@ -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)
src/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/Link.hs view
@@ -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)
src/Data/Array/Accelerate/LLVM/PTX/Link/Object.hs view
@@ -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 --
src/Data/Array/Accelerate/LLVM/PTX/State.hs view
@@ -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
src/Data/Array/Accelerate/LLVM/PTX/Target.hs view
@@ -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-
+ test/nofib/Data/Array/Accelerate/LLVM/PTX/NoFib/RunQ.hs view
@@ -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]
test/nofib/Main.hs view
@@ -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