cuda 0.10.0.0 → 0.10.1.0

raw patch · 16 files changed

+600/−491 lines, 16 filesPVP: major bump suggested

API removals or changes: PVP suggests a major version bump

API changes (from Hackage documentation)

Files

• CHANGELOG.markdown +0/−162
• CHANGELOG.md +184/−0
• README.markdown +0/−24
• README.md +24/−0
• WINDOWS.markdown +0/−54
• WINDOWS.md +54/−0
• cbits/stubs.c +41/−70
• cbits/stubs.h +19/−77
• cuda.cabal +10/−9
• src/Foreign/CUDA/Analysis/Device.chs +23/−3
• src/Foreign/CUDA/Driver.hs +10/−4
• src/Foreign/CUDA/Driver/Graph/Capture.chs +68/−3
• src/Foreign/CUDA/Driver/Graph/Exec.chs +63/−3
• src/Foreign/CUDA/Driver/Texture.chs +3/−3
• src/Foreign/CUDA/Runtime/Error.chs +5/−1
• src/Foreign/CUDA/Runtime/Exec.chs +96/−78

@@ -1,162 +0,0 @@-# Change Log--Notable changes to the project will be documented in this file.--The format is based on [Keep a Changelog](http://keepachangelog.com/).--## [0.10.0.0] - 2018-10-02-### Added-  * Device properties for SM7-  * Functions from CUDA-9.2-    * `Device.uuid`-    * `Stream.getContext`--  * Functions from CUDA-10.0-    * `Foreign.CUDA.Driver.Graph*`--  * Additional bindings from older CUDA releases--### Changed-  * Replace uses of `String` with `ShortByteString`--### Removed-  * Support for ghc-7.6--## [0.9.0.3] - 2018-03-12-### Fixed-  * Build fix for Cabal-2.2 (ghc-8.4)--## [0.9.0.2] - 2018-03-07-### Fixed-  * Build fix for Nix ([#53])--## [0.9.0.1] - 2018-02-16-### Fixed-  * Build fix for macOS High Sierra (10.13)--## [0.9.0.0] - 2017-11-15-### Fixed-  * Build fixes for CUDA-9--### Added-  * `Peer.getAttribute`-  * `Exec.launchKernelCooperative`--### Changed-  * Changed type of `Stream.wait` and `Stream.write` to support 64-bit values--## [0.8.0.1] - 2017-10-24-### Fixed-  * Escape backslashes used in -D flags on Windows ([#50])--## [0.8.0.0] - 2017-08-24-### Changed-  * Tested with CUDA toolkit 8.0--### Added-  * Add operations for unified addressing in the device API-  * Add `write` and `wait` operations for streams in the device API-  * (internals) The paths this module was configured against are exposed by the-    module `Foreign.CUDA.Paths`.--## [0.7.5.3] - 2017-03-23-### Fixed-  * Bug fix in occupancy calculator--## [0.7.5.2] - 2017-01-06-### Fixed-  * Build fails with library profiling ([#43])-  * On Windows, the Cabal installer is looking in the wrong place ([#45])-  * Windows install fix ([#47])--## [0.7.5.1] - 2016-10-21-### Fixed-  * Re-enable support for Cabal-1.22-  * Unknown CUDA device compute capability 6.1 ([#40])-  * Compilation fails for CUDA-8 [was: ghc 7.10.3 fail to install] ([#44])--## [0.7.5.0] - 2016-10-07-### Changed-  * Tested with CUDA toolkit 7.5--### Added-  * Add functions from CUDA-7.5-  * Add profiler control functions-  * Add function `mallocHostForeignPtr`--## [0.7.0.0] - 2015-11-30-### Changed-  * Add support for operations from CUDA-7.0-  * Add support for online linking-  * Add support for inter-process communication-  * Bug fixes, extra documentation, improve library coverage.-  * Mac OS X no longer requires the DYLD_LIBRARY_PATH environment variable in-    order to compile or run programs that use this package.--## [0.6.7.0] - 2015-09-12-### Added-  * Add support for building on Windows (thanks to @mwu-tow)--## [0.6.6.2] - 2015-04-04-### Fixed-  * Build fix--## [0.6.6.1] - 2015-04-04 [YANKED]-### Fixed-  * Build fixes for ghc-7.6 and ghc-7.10--## [0.6.6.0] - 2015-03-10-### Added-  * Add compute-capability data for 3.7, 5.2 devices.--### Changed-  * Combine the definition of the 'Event' and 'Stream' data types. As of-    CUDA-3.1 these data structures are equivalent, and can be safely shared-    between runtime and driver API calls and libraries.--  * Mark FFI imports of potentially long-running API functions as safe. This-    allows them to be safely called from Haskell threads without blocking the-    entire HEC.--### Removed-  * Drop support for CUDA 3.0 and older.--## [0.6.5.1] - 2014-12-02-### Fixed-  * Build fix for Mac OS X 10.10 (Yosemite)--## [0.6.5.0] - 2014-09-03-### Changed-  * Tested with CUDA toolkit 6.5--### Added-  * Add functions from CUDA-6.5---[0.10.0.0]:   https://github.com/tmcdonell/cuda/compare/0.9.0.3...v0.10.0.0-[0.9.0.3]:    https://github.com/tmcdonell/cuda/compare/0.9.0.2...0.9.0.3-[0.9.0.2]:    https://github.com/tmcdonell/cuda/compare/0.9.0.1...0.9.0.2-[0.9.0.1]:    https://github.com/tmcdonell/cuda/compare/0.9.0.0...0.9.0.1-[0.9.0.0]:    https://github.com/tmcdonell/cuda/compare/0.8.0.1...0.9.0.0-[0.8.0.1]:    https://github.com/tmcdonell/cuda/compare/0.8.0.0...0.8.0.1-[0.8.0.0]:    https://github.com/tmcdonell/cuda/compare/0.7.5.3...0.8.0.0-[0.7.5.3]:    https://github.com/tmcdonell/cuda/compare/0.7.5.2...0.7.5.3-[0.7.5.2]:    https://github.com/tmcdonell/cuda/compare/0.7.5.1...0.7.5.2-[0.7.5.1]:    https://github.com/tmcdonell/cuda/compare/0.7.5.0...0.7.5.1-[0.7.5.0]:    https://github.com/tmcdonell/cuda/compare/0.7.0.0...0.7.5.0-[0.7.0.0]:    https://github.com/tmcdonell/cuda/compare/0.6.7.0...0.7.0.0-[0.6.7.0]:    https://github.com/tmcdonell/cuda/compare/0.6.6.2...0.6.7.0-[0.6.6.2]:    https://github.com/tmcdonell/cuda/compare/0.6.6.1...0.6.6.2-[0.6.6.1]:    https://github.com/tmcdonell/cuda/compare/0.6.6.0...0.6.6.1-[0.6.6.0]:    https://github.com/tmcdonell/cuda/compare/0.6.5.1...0.6.6.0-[0.6.5.1]:    https://github.com/tmcdonell/cuda/compare/0.6.5.0...0.6.5.1-[0.6.5.0]:    https://github.com/tmcdonell/cuda/compare/0.6.0.1...0.6.5.0--[#40]:        https://github.com/tmcdonell/cuda/issues/40-[#43]:        https://github.com/tmcdonell/cuda/issues/43-[#44]:        https://github.com/tmcdonell/cuda/issues/44-[#45]:        https://github.com/tmcdonell/cuda/issues/45-[#47]:        https://github.com/tmcdonell/cuda/pull/47-[#50]:        https://github.com/tmcdonell/cuda/pull/50-[#53]:        https://github.com/tmcdonell/cuda/pull/53-

@@ -0,0 +1,184 @@+# Change Log++Notable changes to the project will be documented in this file.++The format is based on [Keep a Changelog](http://keepachangelog.com/).++**NOTE:** The version numbers of this package roughly align to the latest+version of the CUDA API this package is built against This means that this+package _DOES NOT_ follow the PVP, or indeed any sensible version scheme,+because NVIDIA are A-OK introducing breaking changes in minor updates.+++## [next]+### Added+  * Add support for CUDA-10.1++### Changed+  * The function `Foreign.CUDA.Driver.Graph.Capture.start` has an extra+    parameter to specify the capture mode++### Removed+  * The following functions have been deprecated (since at least CUDA-8) and are+    finally removed:+      - `Foreign.CUDA.Runtime.Exec.launch`+      - `Foreign.CUDA.Runtime.Exec.setParams`+      - `Foreign.CUDA.Runtime.Exec.setConfig`++## [0.10.0.0] - 2018-10-02+### Added+  * Device properties for SM7+  * Functions from CUDA-9.2+    * `Device.uuid`+    * `Stream.getContext`++  * Functions from CUDA-10.0+    * `Foreign.CUDA.Driver.Graph*`++  * Additional bindings from older CUDA releases++### Changed+  * Replace uses of `String` with `ShortByteString`++### Removed+  * Support for ghc-7.6++## [0.9.0.3] - 2018-03-12+### Fixed+  * Build fix for Cabal-2.2 (ghc-8.4)++## [0.9.0.2] - 2018-03-07+### Fixed+  * Build fix for Nix ([#53])++## [0.9.0.1] - 2018-02-16+### Fixed+  * Build fix for macOS High Sierra (10.13)++## [0.9.0.0] - 2017-11-15+### Fixed+  * Build fixes for CUDA-9++### Added+  * `Peer.getAttribute`+  * `Exec.launchKernelCooperative`++### Changed+  * Changed type of `Stream.wait` and `Stream.write` to support 64-bit values++## [0.8.0.1] - 2017-10-24+### Fixed+  * Escape backslashes used in -D flags on Windows ([#50])++## [0.8.0.0] - 2017-08-24+### Changed+  * Tested with CUDA toolkit 8.0++### Added+  * Add operations for unified addressing in the device API+  * Add `write` and `wait` operations for streams in the device API+  * (internals) The paths this module was configured against are exposed by the+    module `Foreign.CUDA.Paths`.++## [0.7.5.3] - 2017-03-23+### Fixed+  * Bug fix in occupancy calculator++## [0.7.5.2] - 2017-01-06+### Fixed+  * Build fails with library profiling ([#43])+  * On Windows, the Cabal installer is looking in the wrong place ([#45])+  * Windows install fix ([#47])++## [0.7.5.1] - 2016-10-21+### Fixed+  * Re-enable support for Cabal-1.22+  * Unknown CUDA device compute capability 6.1 ([#40])+  * Compilation fails for CUDA-8 [was: ghc 7.10.3 fail to install] ([#44])++## [0.7.5.0] - 2016-10-07+### Changed+  * Tested with CUDA toolkit 7.5++### Added+  * Add functions from CUDA-7.5+  * Add profiler control functions+  * Add function `mallocHostForeignPtr`++## [0.7.0.0] - 2015-11-30+### Changed+  * Add support for operations from CUDA-7.0+  * Add support for online linking+  * Add support for inter-process communication+  * Bug fixes, extra documentation, improve library coverage.+  * Mac OS X no longer requires the DYLD_LIBRARY_PATH environment variable in+    order to compile or run programs that use this package.++## [0.6.7.0] - 2015-09-12+### Added+  * Add support for building on Windows (thanks to @mwu-tow)++## [0.6.6.2] - 2015-04-04+### Fixed+  * Build fix++## [0.6.6.1] - 2015-04-04 [YANKED]+### Fixed+  * Build fixes for ghc-7.6 and ghc-7.10++## [0.6.6.0] - 2015-03-10+### Added+  * Add compute-capability data for 3.7, 5.2 devices.++### Changed+  * Combine the definition of the 'Event' and 'Stream' data types. As of+    CUDA-3.1 these data structures are equivalent, and can be safely shared+    between runtime and driver API calls and libraries.++  * Mark FFI imports of potentially long-running API functions as safe. This+    allows them to be safely called from Haskell threads without blocking the+    entire HEC.++### Removed+  * Drop support for CUDA 3.0 and older.++## [0.6.5.1] - 2014-12-02+### Fixed+  * Build fix for Mac OS X 10.10 (Yosemite)++## [0.6.5.0] - 2014-09-03+### Changed+  * Tested with CUDA toolkit 6.5++### Added+  * Add functions from CUDA-6.5+++[next]:       https://github.com/tmcdonell/cuda/compare/v0.10.0.0...HEAD+[0.10.0.0]:   https://github.com/tmcdonell/cuda/compare/0.9.0.3...v0.10.0.0+[0.9.0.3]:    https://github.com/tmcdonell/cuda/compare/0.9.0.2...0.9.0.3+[0.9.0.2]:    https://github.com/tmcdonell/cuda/compare/0.9.0.1...0.9.0.2+[0.9.0.1]:    https://github.com/tmcdonell/cuda/compare/0.9.0.0...0.9.0.1+[0.9.0.0]:    https://github.com/tmcdonell/cuda/compare/0.8.0.1...0.9.0.0+[0.8.0.1]:    https://github.com/tmcdonell/cuda/compare/0.8.0.0...0.8.0.1+[0.8.0.0]:    https://github.com/tmcdonell/cuda/compare/0.7.5.3...0.8.0.0+[0.7.5.3]:    https://github.com/tmcdonell/cuda/compare/0.7.5.2...0.7.5.3+[0.7.5.2]:    https://github.com/tmcdonell/cuda/compare/0.7.5.1...0.7.5.2+[0.7.5.1]:    https://github.com/tmcdonell/cuda/compare/0.7.5.0...0.7.5.1+[0.7.5.0]:    https://github.com/tmcdonell/cuda/compare/0.7.0.0...0.7.5.0+[0.7.0.0]:    https://github.com/tmcdonell/cuda/compare/0.6.7.0...0.7.0.0+[0.6.7.0]:    https://github.com/tmcdonell/cuda/compare/0.6.6.2...0.6.7.0+[0.6.6.2]:    https://github.com/tmcdonell/cuda/compare/0.6.6.1...0.6.6.2+[0.6.6.1]:    https://github.com/tmcdonell/cuda/compare/0.6.6.0...0.6.6.1+[0.6.6.0]:    https://github.com/tmcdonell/cuda/compare/0.6.5.1...0.6.6.0+[0.6.5.1]:    https://github.com/tmcdonell/cuda/compare/0.6.5.0...0.6.5.1+[0.6.5.0]:    https://github.com/tmcdonell/cuda/compare/0.6.0.1...0.6.5.0++[#40]:        https://github.com/tmcdonell/cuda/issues/40+[#43]:        https://github.com/tmcdonell/cuda/issues/43+[#44]:        https://github.com/tmcdonell/cuda/issues/44+[#45]:        https://github.com/tmcdonell/cuda/issues/45+[#47]:        https://github.com/tmcdonell/cuda/pull/47+[#50]:        https://github.com/tmcdonell/cuda/pull/50+[#53]:        https://github.com/tmcdonell/cuda/pull/53+

@@ -1,24 +0,0 @@-Haskell FFI Bindings to CUDA-============================--[![Travis build status](https://img.shields.io/travis/tmcdonell/cuda/master.svg?label=linux)](https://travis-ci.org/tmcdonell/cuda)-[![AppVeyor build status](https://img.shields.io/appveyor/ci/tmcdonell/cuda/master.svg?label=windows)](https://ci.appveyor.com/project/tmcdonell/cuda)-[![Stackage LTS](https://stackage.org/package/cuda/badge/lts)](https://stackage.org/lts/package/cuda)-[![Stackage Nightly](https://stackage.org/package/cuda/badge/nightly)](https://stackage.org/nightly/package/cuda)-[![Hackage](https://img.shields.io/hackage/v/cuda.svg)](https://hackage.haskell.org/package/cuda)--The CUDA library provides a direct, general purpose C-like SPMD programming-model for NVIDIA graphics cards (G8x series onwards). This is a collection of-bindings to allow you to call and control, although not write, such functions-from Haskell-land. You will need to install the CUDA driver and developer-toolkit.--  <http://developer.nvidia.com/object/cuda.html>--The configure step will look for your CUDA installation in the standard places,-and if the `nvcc` compiler is found in your `PATH`, relative to that.--For important information on installing on Windows, see:--  <https://github.com/tmcdonell/cuda/blob/master/WINDOWS.markdown>-

@@ -0,0 +1,24 @@+Haskell FFI Bindings to CUDA+============================++[![Travis build status](https://img.shields.io/travis/tmcdonell/cuda/master.svg?label=linux)](https://travis-ci.org/tmcdonell/cuda)+[![AppVeyor build status](https://img.shields.io/appveyor/ci/tmcdonell/cuda/master.svg?label=windows)](https://ci.appveyor.com/project/tmcdonell/cuda)+[![Stackage LTS](https://stackage.org/package/cuda/badge/lts)](https://stackage.org/lts/package/cuda)+[![Stackage Nightly](https://stackage.org/package/cuda/badge/nightly)](https://stackage.org/nightly/package/cuda)+[![Hackage](https://img.shields.io/hackage/v/cuda.svg)](https://hackage.haskell.org/package/cuda)++The CUDA library provides a direct, general purpose C-like SPMD programming+model for NVIDIA graphics cards (G8x series onwards). This is a collection of+bindings to allow you to call and control, although not write, such functions+from Haskell-land. You will need to install the CUDA driver and developer+toolkit.++  <http://developer.nvidia.com/object/cuda.html>++The configure step will look for your CUDA installation in the standard places,+and if the `nvcc` compiler is found in your `PATH`, relative to that.++For important information on installing on Windows, see:++  <https://github.com/tmcdonell/cuda/blob/master/WINDOWS.markdown>+

@@ -1,54 +0,0 @@-Using the CUDA package on Windows-=================================--The CUDA package works on Windows and is actively maintained. If you encounter-any other issues, please report them.--Note that if you build your applications for the Windows 64-bit architecture,-you'll need to update your `ld.exe` as described below.---Windows 64-bit-----------------There is a known issue with the version of `ld.exe` that ships with the 64-bit-versions of (at least) GHC-7.8.4 and GHC-7.10.2. The version of `ld.exe` that-ships with these GHC distributions does not properly link against MS-style-dynamic libraries (such as those that ship with the CUDA toolkit), causing the-application to crash at runtime once those library routines are called. The-configure step will fail if it detects an old version of `ld.exe` (< 2.25.1),-which are known to be broken.--If you are using the 64-bit GHC distributions mentioned above, you will need to-apply the following steps. This bug does not affect 32-bit GHC distributions.-The bug has been fixed in MinGW binutils `ld.exe` >= 2.25.1, so it is expected-that newer releases of GHC will not have this issue.--The problem is fixed by replacing the linker binary `ld.exe` with the newer-(patched) version, available as part of the MSys2 binutils package here:--> <http://repo.msys2.org/mingw/x86_64/mingw-w64-x86_64-binutils-2.25.1-1-any.pkg.tar.xz>--The updated `ld.exe` binary must replace the version at the path:--> `GHC_PATH\mingw\x86_64-w64-mingw32\bin\`--Note that there is another copy of `ld.exe` located at `GHC_PATH\mingw\bin\`,-but this version does not seem to be used, so replacing it as well is not-necessary. It is not sufficient to replace whatever version of `ld.exe` appears-first in your `PATH`.--Please note that having another MinGW installation in `PATH` before the one-shipped with GHC may break things, particularly if you mix 32/64-bit-distributions of MinGW and GHC.--For further discussion of the bug, see:--  * [CUDA package issue][cuda31]-  * [GHC issue][ghc10885]-  * [binutils issue][binutils16598]-- [cuda31]:              https://github.com/tmcdonell/cuda/issues/31- [ghc10885]:            https://ghc.haskell.org/trac/ghc/ticket/10885- [binutils16598]:       https://sourceware.org/bugzilla/show_bug.cgi?id=16598-

@@ -0,0 +1,54 @@+Using the CUDA package on Windows+=================================++The CUDA package works on Windows and is actively maintained. If you encounter+any other issues, please report them.++Note that if you build your applications for the Windows 64-bit architecture,+you'll need to update your `ld.exe` as described below.+++Windows 64-bit+--------------++There is a known issue with the version of `ld.exe` that ships with the 64-bit+versions of (at least) GHC-7.8.4 and GHC-7.10.2. The version of `ld.exe` that+ships with these GHC distributions does not properly link against MS-style+dynamic libraries (such as those that ship with the CUDA toolkit), causing the+application to crash at runtime once those library routines are called. The+configure step will fail if it detects an old version of `ld.exe` (< 2.25.1),+which are known to be broken.++If you are using the 64-bit GHC distributions mentioned above, you will need to+apply the following steps. This bug does not affect 32-bit GHC distributions.+The bug has been fixed in MinGW binutils `ld.exe` >= 2.25.1, so it is expected+that newer releases of GHC will not have this issue.++The problem is fixed by replacing the linker binary `ld.exe` with the newer+(patched) version, available as part of the MSys2 binutils package here:++> <http://repo.msys2.org/mingw/x86_64/mingw-w64-x86_64-binutils-2.25.1-1-any.pkg.tar.xz>++The updated `ld.exe` binary must replace the version at the path:++> `GHC_PATH\mingw\x86_64-w64-mingw32\bin\`++Note that there is another copy of `ld.exe` located at `GHC_PATH\mingw\bin\`,+but this version does not seem to be used, so replacing it as well is not+necessary. It is not sufficient to replace whatever version of `ld.exe` appears+first in your `PATH`.++Please note that having another MinGW installation in `PATH` before the one+shipped with GHC may break things, particularly if you mix 32/64-bit+distributions of MinGW and GHC.++For further discussion of the bug, see:++  * [CUDA package issue][cuda31]+  * [GHC issue][ghc10885]+  * [binutils issue][binutils16598]++ [cuda31]:              https://github.com/tmcdonell/cuda/issues/31+ [ghc10885]:            https://ghc.haskell.org/trac/ghc/ticket/10885+ [binutils16598]:       https://sourceware.org/bugzilla/show_bug.cgi?id=16598+

@@ -4,33 +4,25 @@  #include "cbits/stubs.h" +#if CUDART_VERSION >= 7000+cudaError_t cudaLaunchKernel_simple(const void *func, unsigned int gridX, unsigned int gridY, unsigned int gridZ, unsigned int blockX, unsigned int blockY, unsigned int blockZ, void **args, size_t sharedMem, cudaStream_t stream)+{+    dim3 gridDim  = {gridX, gridY, gridZ};+    dim3 blockDim = {blockX, blockY, blockZ}; -cudaError_t-cudaConfigureCallSimple-(-    int gridX,  int gridY,-    int blockX, int blockY, int blockZ,-    size_t sharedMem,-    cudaStream_t stream-)+    return cudaLaunchKernel(func, gridDim, blockDim, args, sharedMem, stream);+}+#else+cudaError_t cudaConfigureCall_simple(unsigned int gridX, unsigned int gridY, unsigned int blockX, unsigned int blockY, unsigned int blockZ, size_t sharedMem, cudaStream_t stream) {     dim3 gridDim  = {gridX, gridY, 1};     dim3 blockDim = {blockX,blockY,blockZ};      return cudaConfigureCall(gridDim, blockDim, sharedMem, stream); }+#endif -CUresult-cuTexRefSetAddress2DSimple-(-    CUtexref tex,-    CUarray_format format,-    unsigned int numChannels,-    CUdeviceptr dptr,-    size_t width,-    size_t height,-    size_t pitch-)+CUresult cuTexRefSetAddress2D_simple(CUtexref tex, CUarray_format format, unsigned int numChannels, CUdeviceptr dptr, size_t width, size_t height, size_t pitch) {     CUDA_ARRAY_DESCRIPTOR desc;     desc.Format      = format;@@ -41,14 +33,7 @@     return cuTexRefSetAddress2D(tex, &desc, dptr, pitch); } -CUresult-cuMemcpy2DHtoD-(-    CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY,-    void* srcHost,         unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY,-    unsigned int widthInBytes,-    unsigned int height-)+CUresult cuMemcpy2DHtoD(CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY, void* srcHost, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY, unsigned int widthInBytes, unsigned int height) {     CUDA_MEMCPY2D desc; @@ -75,15 +60,7 @@     return cuMemcpy2D(&desc); } -CUresult-cuMemcpy2DHtoDAsync-(-    CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY,-    void* srcHost,         unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY,-    unsigned int widthInBytes,-    unsigned int height,-    CUstream hStream-)+CUresult cuMemcpy2DHtoDAsync(CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY, void* srcHost, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY, unsigned int widthInBytes, unsigned int height, CUstream hStream) {     CUDA_MEMCPY2D desc; @@ -110,14 +87,7 @@     return cuMemcpy2DAsync(&desc, hStream); } -CUresult-cuMemcpy2DDtoH-(-    void* dstHost,         unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY,-    CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY,-    unsigned int widthInBytes,-    unsigned int height-)+CUresult cuMemcpy2DDtoH(void* dstHost, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY, CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY, unsigned int widthInBytes, unsigned int height) {     CUDA_MEMCPY2D desc; @@ -144,15 +114,7 @@     return cuMemcpy2D(&desc); } -CUresult-cuMemcpy2DDtoHAsync-(-    void* dstHost,         unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY,-    CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY,-    unsigned int widthInBytes,-    unsigned int height,-    CUstream hStream-)+CUresult cuMemcpy2DDtoHAsync(void* dstHost, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY, CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY, unsigned int widthInBytes, unsigned int height, CUstream hStream) {     CUDA_MEMCPY2D desc; @@ -179,14 +141,7 @@     return cuMemcpy2DAsync(&desc, hStream); } -CUresult-cuMemcpy2DDtoD-(-    CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY,-    CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY,-    unsigned int widthInBytes,-    unsigned int height-)+CUresult cuMemcpy2DDtoD(CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY, CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY, unsigned int widthInBytes, unsigned int height) {     CUDA_MEMCPY2D desc; @@ -213,15 +168,7 @@     return cuMemcpy2D(&desc); } -CUresult-cuMemcpy2DDtoDAsync-(-    CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY,-    CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY,-    unsigned int widthInBytes,-    unsigned int height,-    CUstream hStream-)+CUresult cuMemcpy2DDtoDAsync(CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY, CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY, unsigned int widthInBytes, unsigned int height, CUstream hStream) {     CUDA_MEMCPY2D desc; @@ -466,6 +413,30 @@     memsetParams.width       = width;      return cuGraphAddMemsetNode(phGraphNode, hGraph, dependencies, numDependencies, &memsetParams, ctx);+}+#endif++#if CUDA_VERSION >= 10010+CUresult CUDAAPI cuStreamBeginCapture(CUstream hStream, CUstreamCaptureMode mode)+{+    return cuStreamBeginCapture_v2(hStream, mode);+}++CUresult CUDAAPI cuGraphExecKernelNodeSetParams_simple(CUgraphExec hGraphExec, CUgraphNode hNode, CUfunction func, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, void** kernelParams)+{+    CUDA_KERNEL_NODE_PARAMS nodeParams;+    nodeParams.func           = func;+    nodeParams.gridDimX       = gridDimX;+    nodeParams.gridDimY       = gridDimY;+    nodeParams.gridDimZ       = gridDimZ;+    nodeParams.blockDimX      = blockDimX;+    nodeParams.blockDimY      = blockDimY;+    nodeParams.blockDimZ      = blockDimZ;+    nodeParams.sharedMemBytes = sharedMemBytes;+    nodeParams.kernelParams   = kernelParams;+    nodeParams.extra          = NULL;++    return cuGraphExecKernelNodeSetParams(hGraphExec, hNode, &nodeParams); } #endif 

@@ -21,84 +21,19 @@  void enable_constructors(); -cudaError_t-cudaConfigureCallSimple-(-    int gridX,  int gridY,-    int blockX, int blockY, int blockZ,-    size_t sharedMem,-    cudaStream_t stream-);--CUresult-cuTexRefSetAddress2DSimple-(-    CUtexref            tex,-    CUarray_format      format,-    unsigned int        numChannels,-    CUdeviceptr         dptr,-    size_t              width,-    size_t              height,-    size_t              pitch-);--CUresult-cuMemcpy2DHtoD-(-    CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY,-    void* srcHost,         unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY,-    unsigned int widthInBytes,-    unsigned int height-);--CUresult-cuMemcpy2DHtoDAsync-(-    CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY,-    void* srcHost,         unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY,-    unsigned int widthInBytes,-    unsigned int height,-    CUstream hStream-);--CUresult-cuMemcpy2DDtoH-(-    void* dstHost,         unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY,-    CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY,-    unsigned int widthInBytes,-    unsigned int height-);--CUresult-cuMemcpy2DDtoHAsync-(-    void* dstHost,         unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY,-    CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY,-    unsigned int widthInBytes,-    unsigned int height,-    CUstream hStream-);--CUresult-cuMemcpy2DDtoD-(-    CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY,-    CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY,-    unsigned int widthInBytes,-    unsigned int height-);--CUresult-cuMemcpy2DDtoDAsync-(-    CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY,-    CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY,-    unsigned int widthInBytes,-    unsigned int height,-    CUstream hStream-);+#if CUDART_VERSION >= 7000+cudaError_t cudaLaunchKernel_simple(const void *func, unsigned int gridX, unsigned int gridY, unsigned int gridZ, unsigned int blockX, unsigned int blockY, unsigned int blockZ, void **args, size_t sharedMem, cudaStream_t stream);+#else+cudaError_t cudaConfigureCall_simple(unsigned int gridX, unsigned int gridY, unsigned int blockX, unsigned int blockY, unsigned int blockZ, size_t sharedMem, cudaStream_t stream);+#endif +CUresult cuTexRefSetAddress2D_simple(CUtexref tex, CUarray_format format, unsigned int numChannels, CUdeviceptr dptr, size_t width, size_t height, size_t pitch);+CUresult cuMemcpy2DHtoD(CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY, void* srcHost, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY, unsigned int widthInBytes, unsigned int height);+CUresult cuMemcpy2DHtoDAsync(CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY, void* srcHost, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY, unsigned int widthInBytes, unsigned int height, CUstream hStream);+CUresult cuMemcpy2DDtoH(void* dstHost, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY, CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY, unsigned int widthInBytes, unsigned int height);+CUresult cuMemcpy2DDtoHAsync(void* dstHost, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY, CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY, unsigned int widthInBytes, unsigned int height, CUstream hStream);+CUresult cuMemcpy2DDtoD(CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY, CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY, unsigned int widthInBytes, unsigned int height);+CUresult cuMemcpy2DDtoDAsync(CUdeviceptr dstDevice, unsigned int dstPitch, unsigned int dstXInBytes, unsigned int dstY, CUdeviceptr srcDevice, unsigned int srcPitch, unsigned int srcXInBytes, unsigned int srcY, unsigned int widthInBytes, unsigned int height, CUstream hStream);  /*  * Need to re-export some symbols as they are now generated by #defines, which@@ -223,6 +158,13 @@ CUresult CUDAAPI cuGraphAddKernelNode_simple(CUgraphNode *phGraphNode, CUgraph hGraph, CUgraphNode *dependencies, size_t numDependencies, CUfunction func, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, void** kernelParams); CUresult CUDAAPI cuGraphAddMemcpyNode_simple(CUgraphNode *phGraphNode, CUgraph hGraph, CUgraphNode *dependencies, size_t numDependencies, CUcontext ctx, size_t srcXInBytes, size_t srcY, size_t srcZ, size_t srcLOD, CUmemorytype srcMemoryType, const void *srcPtr, size_t srcPitch, size_t srcHeight, size_t dstXInBytes, size_t dstY, size_t dstZ, size_t dstLOD, CUmemorytype dstMemoryType, void *dstPtr, size_t dstPitch, size_t dstHeight, size_t widthInBytes, size_t height, size_t depth); CUresult CUDAAPI cuGraphAddMemsetNode_simple(CUgraphNode *phGraphNode, CUgraph hGraph, CUgraphNode *dependencies, size_t numDependencies, CUcontext ctx, CUdeviceptr dst, unsigned int elementSize, size_t height, size_t pitch, unsigned int value, size_t width);+#endif++#if CUDA_VERSION >= 10010+#undef cuStreamBeginCapture++CUresult CUDAAPI cuStreamBeginCapture(CUstream hStream, CUstreamCaptureMode mode);+CUresult CUDAAPI cuGraphExecKernelNodeSetParams_simple(CUgraphExec hGraphExec, CUgraphNode hNode, CUfunction func, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, void** kernelParams); #endif  #ifdef __cplusplus

@@ -1,5 +1,5 @@ Name:                   cuda-Version:                0.10.0.0+Version:                0.10.1.0 Synopsis:               FFI binding to the CUDA interface for programming NVIDIA GPUs Description:     The CUDA library provides a direct, general purpose C-like SPMD programming@@ -53,7 +53,7 @@     .     For additional notes on installing on Windows, see:     .-    * <https://github.com/tmcdonell/cuda/blob/master/WINDOWS.markdown>+    * <https://github.com/tmcdonell/cuda/blob/master/WINDOWS.md>     .  License:                BSD3@@ -66,15 +66,16 @@ Category:               Foreign Cabal-version:          >= 1.24 Tested-with:            GHC >= 7.8- Build-type:             Custom -Extra-tmp-files:        cuda.buildinfo.generated+Extra-tmp-files:+  cuda.buildinfo.generated -Extra-source-files:     cbits/stubs.h-                        CHANGELOG.markdown-                        README.markdown-                        WINDOWS.markdown+Extra-source-files:+  cbits/stubs.h+  CHANGELOG.md+  README.md+  WINDOWS.md  custom-setup   setup-depends:@@ -189,6 +190,6 @@ source-repository this     type:               git     location:           https://github.com/tmcdonell/cuda-    tag:                v0.10.0.0+    tag:                v0.10.1.0  -- vim: nospell

@@ -322,14 +322,34 @@         , maxGridsPerDevice     = 16         } -      Compute 7 _ -> DeviceResources              -- Volta GV100+      Compute 7 0 -> resources (Compute 7 2)      -- Volta GV100+      Compute 7 2 -> DeviceResources         { threadsPerWarp        = 32         , coresPerMP            = 64         , warpsPerMP            = 64         , threadsPerMP          = 2048         , threadBlocksPerMP     = 32-        , sharedMemPerMP        = 98304-        , maxSharedMemPerBlock  = 49152           -- XXX: or 96KB?+        , sharedMemPerMP        = 98304           -- of 128KB+        , maxSharedMemPerBlock  = 98304+        , regFileSizePerMP      = 65536+        , maxRegPerBlock        = 65536+        , regAllocUnit          = 256+        , regAllocationStyle    = Warp+        , maxRegPerThread       = 255+        , sharedMemAllocUnit    = 256+        , warpAllocUnit         = 4+        , warpRegAllocUnit      = 256+        , maxGridsPerDevice     = 128+        }++      Compute 7 5 -> DeviceResources              -- Turing TU1xx+        { threadsPerWarp        = 32+        , coresPerMP            = 64+        , warpsPerMP            = 32+        , threadsPerMP          = 1024+        , threadBlocksPerMP     = 16+        , sharedMemPerMP        = 65536           -- of 96KB+        , maxSharedMemPerBlock  = 65536         , regFileSizePerMP      = 65536         , maxRegPerBlock        = 65536         , regAllocUnit          = 256

@@ -228,12 +228,18 @@  ) where +-- If we don't import everything from a module below, the generated haddocks+-- will not have the convenient link to a module, but instead enumerate+-- documentation for everything which was imported. This, unfortunately, looks+-- like garbage, so it would be nice to find a better solution.+--    -- TLM 2018-11-18+-- import Foreign.CUDA.Ptr-import Foreign.CUDA.Driver.Context      hiding ( useContext, device )-import Foreign.CUDA.Driver.Device       hiding ( useDevice )+import Foreign.CUDA.Driver.Context                        hiding ( device )+import Foreign.CUDA.Driver.Device import Foreign.CUDA.Driver.Error-import Foreign.CUDA.Driver.Exec         hiding ( useFun )-import Foreign.CUDA.Driver.Marshal      hiding ( useDeviceHandle, peekDeviceHandle )+import Foreign.CUDA.Driver.Exec+import Foreign.CUDA.Driver.Marshal import Foreign.CUDA.Driver.Module import Foreign.CUDA.Driver.Unified import Foreign.CUDA.Driver.Utils

@@ -18,8 +18,8 @@  module Foreign.CUDA.Driver.Graph.Capture ( -  Status(..),-  start, stop, status,+  Status(..), Mode(..),+  start, stop, status, info, mode,  ) where @@ -52,7 +52,22 @@   with prefix="CU_STREAM_CAPTURE_STATUS" deriving (Eq, Show, Bounded) #} #endif +#if CUDA_VERSION < 10010+data Mode +instance Enum Mode where+#ifdef USE_EMPTY_CASE+  toEnum   x = error ("Mode.toEnum: Cannot match " ++ show x)+  fromEnum x = case x of {}+#endif++#else+{# enum CUstreamCaptureMode as Mode+  { underscoreToCase }+  with prefix="CU_STREAM_CAPTURE_MODE" deriving (Eq, Show, Bounded) #}+#endif++ -------------------------------------------------------------------------------- -- Graph capture --------------------------------------------------------------------------------@@ -66,11 +81,20 @@ -- @since 0.10.0.0 -- #if CUDA_VERSION < 10000-start :: Stream -> IO ()+start :: Stream -> Mode -> IO () start = requireSDK 'start 10.0+#elif CUDA_VERSION < 10010+start :: Stream -> Mode -> IO ()+start s _ = cuStreamBeginCapture s+  where+    {# fun unsafe cuStreamBeginCapture+      { useStream `Stream'+      }+      -> `()' checkStatus*- #} #else {# fun unsafe cuStreamBeginCapture as start   { useStream `Stream'+  ,           `Mode'   }   -> `()' checkStatus*- #} #endif@@ -111,6 +135,47 @@ {# fun unsafe cuStreamIsCapturing as status   { useStream `Stream'   , alloca-   `Status' peekEnum*+  }+  -> `()' checkStatus*- #}+#endif+++-- | Query the capture status of a stream and get an id for the capture+-- sequence, which is unique over the lifetime of the process.+--+-- Requires CUDA-10.1+--+-- <https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__STREAM.html#group__CUDA__STREAM_1g13145ece1d79a1d79a1d22abb9663216>+--+-- @since 0.10.1.0+--+#if CUDA_VERSION < 10010+info :: Stream -> IO (Status, Int64)+info = requireSDK 'info 10.1+#else+{# fun unsafe cuStreamGetCaptureInfo as info+  { useStream `Stream'+  , alloca-   `Status' peekEnum*+  , alloca-   `Int64'  peekIntConv*+  }+  -> `()' checkStatus*- #}+#endif+++-- | Set the stream capture interaction mode for this thread. Return the previous value.+--+-- Requires CUDA-10.1+--+-- <https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__STREAM.html#group__CUDA__STREAM_1g378135b262f02a43a7caeab239ae493d>+--+-- @since 0.10.1.0+--+#if CUDA_VERSION < 10010+mode :: Mode -> IO Mode+mode = requireSDK 'mode 10.1+#else+{# fun unsafe cuThreadExchangeStreamCaptureMode as mode+  { withEnum* `Mode' peekEnum*   }   -> `()' checkStatus*- #} #endif

@@ -22,6 +22,7 @@   launch,   instantiate,   destroy,+  setKernel,  ) where @@ -29,6 +30,7 @@ {# context lib="cuda" #}  import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Driver.Exec                           ( Fun(..), FunParam(..) ) import Foreign.CUDA.Driver.Graph.Base import Foreign.CUDA.Driver.Stream                         ( Stream(..) ) import Foreign.CUDA.Internal.C2HS@@ -59,7 +61,7 @@ {-# INLINEABLE launch #-} #if CUDA_VERSION < 10000 launch :: Executable -> Stream -> IO ()-launch _ _ = requireSDK 'launch 10.0+launch = requireSDK 'launch 10.0 #else {# fun unsafe cuGraphLaunch as launch   { useExecutable `Executable'@@ -81,7 +83,7 @@ {-# INLINEABLE instantiate #-} instantiate :: Graph -> IO Executable #if CUDA_VERSION < 10000-instantiate _ = requireSDK 'instantiate 10.0+instantiate = requireSDK 'instantiate 10.0 #else instantiate !g = do   let logSize = 2048@@ -111,6 +113,64 @@ #endif  +-- | Update the parameters for a kernel node in the given executable graph+--+-- Requires CUDA-10.1+--+-- <https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__GRAPH.html#group__CUDA__GRAPH_1gd84243569e4c3d6356b9f2eea20ed48c>+--+-- @since 0.10.1.0+--+setKernel+    :: Executable+    -> Node+    -> Fun+    -> (Int, Int, Int)  -- ^ grid dimension+    -> (Int, Int, Int)  -- ^ thread block dimensions+    -> Int              -- ^ shared memory (bytes)+    -> [FunParam]+    -> IO ()+#if CUDA_VERSION < 10010+setKernel = requireSDK 'setKernel 10.1+#else+setKernel !exe !n !fun (!gx,!gy,!gz) (!tx,!ty,!tz) !sm !args+  = withMany withFP args+  $ \pa -> withArray pa+  $ \pp -> cuGraphExecKernelNodeSetParams_simple exe n fun gx gy gz tx ty tz sm pp+  where+    withFP :: FunParam -> (Ptr () -> IO b) -> IO b+    withFP !p !f = case p of+      IArg v -> with' v (f . castPtr)+      FArg v -> with' v (f . castPtr)+      VArg v -> with' v (f . castPtr)++    -- can't use the standard 'with' because 'alloca' will pass an undefined+    -- dummy argument when determining 'sizeOf' and 'alignment', but sometimes+    -- instances in Accelerate need to evaluate this argument.+    --+    with' :: Storable a => a -> (Ptr a -> IO b) -> IO b+    with' !val !f =+      allocaBytes (sizeOf val) $ \ptr -> do+        poke ptr val+        f ptr++    {# fun unsafe cuGraphExecKernelNodeSetParams_simple+      { useExecutable `Executable'+      , useNode       `Node'+      , useFun        `Fun'+      ,               `Int'+      ,               `Int'+      ,               `Int'+      ,               `Int'+      ,               `Int'+      ,               `Int'+      ,               `Int'+      , id            `Ptr (Ptr ())'+      }+      -> `()' checkStatus*- #}+#endif++ -- | Destroy an executable graph -- -- Requires CUDA-10.0@@ -122,7 +182,7 @@ {-# INLINEABLE destroy #-} #if CUDA_VERSION < 10000 destroy :: Executable -> IO ()-destroy _ = requireSDK 'destroy 10.0+destroy = requireSDK 'destroy 10.0 #else {# fun unsafe cuGraphExecDestroy as destroy   { useExecutable `Executable'

@@ -170,10 +170,10 @@ {-# INLINEABLE bind2D #-} bind2D :: Texture -> Format -> Int -> DevicePtr a -> (Int,Int) -> Int64 -> IO () bind2D !tex !fmt !chn !dptr (!width,!height) !pitch =-  nothingIfOk =<< cuTexRefSetAddress2DSimple tex fmt chn dptr width height pitch+  nothingIfOk =<< cuTexRefSetAddress2D_simple tex fmt chn dptr width height pitch -{-# INLINE cuTexRefSetAddress2DSimple #-}-{# fun unsafe cuTexRefSetAddress2DSimple+{-# INLINE cuTexRefSetAddress2D_simple #-}+{# fun unsafe cuTexRefSetAddress2D_simple   { useTexture      `Texture'   , cFromEnum       `Format'   ,                 `Int'

@@ -16,7 +16,7 @@   Status(..), CUDAException(..),    cudaError, describe, requireSDK,-  resultIfOk, nothingIfOk+  resultIfOk, nothingIfOk, checkStatus,  ) where @@ -116,4 +116,8 @@     case status of         Success -> return  ()         _       -> throwIO (ExitCode status)++{-# INLINE checkStatus #-}+checkStatus :: CInt -> IO ()+checkStatus = nothingIfOk . cToEnum 

@@ -17,7 +17,7 @@    -- * Kernel Execution   Fun, FunAttributes(..), FunParam(..), CacheConfig(..),-  attributes, setConfig, setParams, setCacheConfig, launch, launchKernel,+  attributes, setCacheConfig, launchKernel,  ) where @@ -33,6 +33,7 @@ import Foreign import Foreign.C import Control.Monad+import Data.Maybe  #c typedef struct cudaFuncAttributes cudaFuncAttributes;@@ -133,33 +134,104 @@   -- |--- Specify the grid and block dimensions for a device call. Used in conjunction--- with 'setParams', this pushes data onto the execution stack that will be--- popped when a function is 'launch'ed.+-- On devices where the L1 cache and shared memory use the same hardware+-- resources, this sets the preferred cache configuration for the given device+-- function. This is only a preference; the driver is free to choose a different+-- configuration as required to execute the function. ---{-# INLINEABLE setConfig #-}-setConfig :: (Int,Int)          -- ^ grid dimensions-          -> (Int,Int,Int)      -- ^ block dimensions-          -> Int64              -- ^ shared memory per block (bytes)-          -> Maybe Stream       -- ^ associated processing stream-          -> IO ()-setConfig (!gx,!gy) (!bx,!by,!bz) !sharedMem !mst =-  nothingIfOk =<<-    cudaConfigureCallSimple gx gy bx by bz sharedMem (maybe defaultStream id mst)+-- Switching between configuration modes may insert a device-side+-- synchronisation point for streamed kernel launches+--+{-# INLINEABLE setCacheConfig #-}+setCacheConfig :: Fun -> CacheConfig -> IO ()+#if CUDART_VERSION < 3000+setCacheConfig _ _       = requireSDK 'setCacheConfig 3.0+#else+setCacheConfig !fn !pref = nothingIfOk =<< cudaFuncSetCacheConfig fn pref +{-# INLINE cudaFuncSetCacheConfig #-}+{# fun unsafe cudaFuncSetCacheConfig+  { withFun*  `Fun'+  , cFromEnum `CacheConfig' } -> `Status' cToEnum #}+#endif ++-- |+-- Invoke a kernel on a @(gx * gy)@ grid of blocks, where each block contains+-- @(tx * ty * tz)@ threads and has access to a given number of bytes of shared+-- memory. The launch may also be associated with a specific 'Stream'. --+{-# INLINEABLE launchKernel #-}+launchKernel+    :: Fun              -- ^ Device function symbol+    -> (Int,Int)        -- ^ grid dimensions+    -> (Int,Int,Int)    -- ^ thread block shape+    -> Int64            -- ^ shared memory per block (bytes)+    -> Maybe Stream     -- ^ (optional) execution stream+    -> [FunParam]+    -> IO ()+#if CUDART_VERSION >= 7000+launchKernel !fn (!gx,!gy) (!bx,!by,!bz) !sm !mst !args+  = (=<<) nothingIfOk+  $ withMany withFP args+  $ \pa -> withArray pa+  $ \pp -> cudaLaunchKernel_simple fn gx gy 1 bx by bz pp sm (fromMaybe defaultStream mst)+  where+    withFP :: FunParam -> (Ptr FunParam -> IO b) -> IO b+    withFP p f = case p of+      IArg v -> with' v (f . castPtr)+      FArg v -> with' v (f . castPtr)+      DArg v -> with' v (f . castPtr)+      VArg v -> with' v (f . castPtr)++    with' :: Storable a => a -> (Ptr a -> IO b) -> IO b+    with' !val !f =+      allocaBytes (sizeOf val) $ \ptr -> do+        poke ptr val+        f ptr++{-# INLINE cudaLaunchKernel_simple #-}+{# fun unsafe cudaLaunchKernel_simple+  { withFun*  `Fun'+  ,           `Int', `Int', `Int'+  ,           `Int', `Int', `Int'+  , castPtr   `Ptr (Ptr FunParam)'+  ,           `Int64'+  , useStream `Stream'+  }+  -> `Status' cToEnum #}++#else+launchKernel !fn (!gx,!gy) (!bx,!by,!bz) !sm !mst !args = do+  setConfig gx gy bx by bz sm (fromMaybe defaultStream mst)+  setParams args+  launch fn+++-- Invoke the @__global__@ kernel function on the device. This must be preceded+-- by a call to 'setConfig' and (if appropriate) 'setParams'.+--+{-# INLINE launch #-}+{# fun unsafe cudaLaunch as launch+  { withFun* `Fun' } -> `()' checkStatus*- #}+++-- Specify the grid and block dimensions for a device call. Used in conjunction+-- with 'setParams', this pushes data onto the execution stack that will be+-- popped when a function is 'launch'ed.+-- -- The FFI does not support passing deferenced structures to C functions, as -- this is highly platform/compiler dependent. Wrap our own function stub -- accepting plain integers. ---{-# INLINE cudaConfigureCallSimple #-}-{# fun unsafe cudaConfigureCallSimple+{-# INLINE setConfig #-}+{# fun unsafe cudaConfigureCall_simple as setConfig   {           `Int', `Int'   ,           `Int', `Int', `Int'   , cIntConv  `Int64'-  , useStream `Stream'            } -> `Status' cToEnum #}-+  , useStream `Stream'+  }+  -> `()' checkStatus*- #}  -- | -- Set the argument parameters that will be passed to the next kernel@@ -172,7 +244,7 @@   where     k !offset !arg = do       let s = size arg-      set arg s offset >>= nothingIfOk+      set arg s offset       return (offset + s)      size (IArg _) = sizeOf (undefined :: Int)@@ -183,82 +255,28 @@     set (IArg v) s o = cudaSetupArgument v s o     set (FArg v) s o = cudaSetupArgument v s o     set (VArg v) s o = cudaSetupArgument v s o-    set (DArg v) s o =-      cudaSetDoubleForDevice v >>= resultIfOk >>= \d ->+    set (DArg v) s o = do+      d <- cudaSetDoubleForDevice v       cudaSetupArgument d s o - {-# INLINE cudaSetupArgument #-} {# fun unsafe cudaSetupArgument   `Storable a' =>   { with'* `a'   ,        `Int'-  ,        `Int'   } -> `Status' cToEnum #}+  ,        `Int'+  }+  -> `()' checkStatus*- #}   where     with' v a = with v $ \p -> a (castPtr p)  {-# INLINE cudaSetDoubleForDevice #-} {# fun unsafe cudaSetDoubleForDevice-  { with'* `Double' peek'* } -> `Status' cToEnum #}+  { with'* `Double' peek'* } -> `()' checkStatus*- #}   where     with' v a = with v $ \p -> a (castPtr p)     peek'     = peek . castPtr----- |--- On devices where the L1 cache and shared memory use the same hardware--- resources, this sets the preferred cache configuration for the given device--- function. This is only a preference; the driver is free to choose a different--- configuration as required to execute the function.------ Switching between configuration modes may insert a device-side--- synchronisation point for streamed kernel launches----{-# INLINEABLE setCacheConfig #-}-setCacheConfig :: Fun -> CacheConfig -> IO ()-#if CUDART_VERSION < 3000-setCacheConfig _ _       = requireSDK 'setCacheConfig 3.0-#else-setCacheConfig !fn !pref = nothingIfOk =<< cudaFuncSetCacheConfig fn pref--{-# INLINE cudaFuncSetCacheConfig #-}-{# fun unsafe cudaFuncSetCacheConfig-  { withFun*  `Fun'-  , cFromEnum `CacheConfig' } -> `Status' cToEnum #} #endif----- |--- Invoke the @__global__@ kernel function on the device. This must be preceded--- by a call to 'setConfig' and (if appropriate) 'setParams'.----{-# INLINEABLE launch #-}-launch :: Fun -> IO ()-launch !fn = nothingIfOk =<< cudaLaunch fn--{-# INLINE cudaLaunch #-}-{# fun unsafe cudaLaunch-  { withFun* `Fun' } -> `Status' cToEnum #}----- |--- Invoke a kernel on a @(gx * gy)@ grid of blocks, where each block contains--- @(tx * ty * tz)@ threads and has access to a given number of bytes of shared--- memory. The launch may also be associated with a specific 'Stream'.----{-# INLINEABLE launchKernel #-}-launchKernel-    :: Fun              -- ^ Device function symbol-    -> (Int,Int)        -- ^ grid dimensions-    -> (Int,Int,Int)    -- ^ thread block shape-    -> Int64            -- ^ shared memory per block (bytes)-    -> Maybe Stream     -- ^ (optional) execution stream-    -> [FunParam]-    -> IO ()-launchKernel !fn !grid !block !sm !mst !args = do-  setConfig grid block sm mst-  setParams args-  launch fn  -------------------------------------------------------------------------------- -- Internals