cuda 0.12.8.0 → 0.13.0.0

raw patch · 27 files changed

+371/−2020 lines, 27 filessetup-changedPVP ok

version bump matches the API change (PVP)

API changes (from Hackage documentation)

Files

• CHANGELOG.md +11/−2
• README.md +8/−5
• Setup.hs +125/−32
• cbits/stubs.c +21/−18
• cbits/stubs.h +5/−6
• cuda.cabal +19/−37
• examples/src/deviceQueryDrv/DeviceQuery.hs +22/−9
• src/Foreign/CUDA.hs +0/−19
• src/Foreign/CUDA/Analysis/Device.chs +15/−3
• src/Foreign/CUDA/Driver.hs +1/−1
• src/Foreign/CUDA/Driver/Device.chs +11/−1
• src/Foreign/CUDA/Driver/Error.chs +1/−1
• src/Foreign/CUDA/Driver/Event.chs +7/−0
• src/Foreign/CUDA/Driver/Graph/Build.chs +62/−0
• src/Foreign/CUDA/Driver/Graph/Capture.chs +25/−5
• src/Foreign/CUDA/Driver/Marshal.chs +10/−0
• src/Foreign/CUDA/Driver/Unified.chs +10/−1
• src/Foreign/CUDA/Internal/Describe.hs +18/−0
• src/Foreign/CUDA/Runtime.hs +0/−28
• src/Foreign/CUDA/Runtime/Device.chs +0/−440
• src/Foreign/CUDA/Runtime/Error.chs +0/−123
• src/Foreign/CUDA/Runtime/Event.chs +0/−149
• src/Foreign/CUDA/Runtime/Exec.chs +0/−294
• src/Foreign/CUDA/Runtime/Marshal.chs +0/−648
• src/Foreign/CUDA/Runtime/Stream.chs +0/−117
• src/Foreign/CUDA/Runtime/Utils.chs +0/−63
• src/Text/Show/Describe.hs +0/−18

@@ -5,12 +5,21 @@ 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+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.  -## [0.12.8.0] - ???+## [0.13.0.0] - 2026-03-30+### Added+  * Support for CUDA-13++### Removed+  * Support for the runtime API (`Foreign.CUDA.RUntime`). There is an+    experimental cuda-runtime package in the Git repository; contact us if you+    depend on this.++## [0.12.8.0] - 2025-08-21 ### Added   * Support for CUDA-12       - Thanks to @noahmartinwilliams on GitHub for helping out!

@@ -25,10 +25,10 @@  ## Missing functionality -_This library is currently in **maintenance mode**. While we plan to release-updates to keep the existing interface working with newer CUDA versions (as-long as the underlying APIs remain available), no binding of new features is-planned at the moment. Get in touch if you want to contribute._+This library is currently in **maintenance mode**. While we are happy to+release updates to keep the existing interface working with newer CUDA versions+(as long as the underlying APIs remain available), no binding of new features+is planned at the moment. Get in touch if you want to contribute.  Here is an incomplete historical list of missing bindings. Pull requests welcome! @@ -150,9 +150,12 @@ - cuGraphMemAllocNodeGetParams - cuGraphMemFreeNodeGetParams -### CUDA-12+### CUDA >= 12  A lot. PRs welcome.++- CUDA-12.3+  - Edge data in the driver Graph API (`cuGraphAddDependencies_v2` etc.)   # Old compatibility notes

@@ -1,13 +1,22 @@+-- Decouple from GHC's default language setting, so that it's easier+-- to maintain compatibility with old GHCs.+{-# LANGUAGE Haskell2010     #-}+{-# OPTIONS_GHC -Wall        #-}++{-# LANGUAGE ConstraintKinds #-} {-# LANGUAGE CPP             #-}+{-# LANGUAGE DataKinds       #-}+{-# LANGUAGE KindSignatures  #-} {-# LANGUAGE QuasiQuotes     #-} {-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TupleSections   #-}  -- The MIN_VERSION_Cabal macro was introduced with Cabal-1.24 (??) #ifndef MIN_VERSION_Cabal #define MIN_VERSION_Cabal(major1,major2,minor) 0 #endif -import Distribution.PackageDescription+import Distribution.PackageDescription                              hiding ( Flag ) import Distribution.Simple import Distribution.Simple.BuildPaths import Distribution.Simple.Command@@ -33,6 +42,18 @@ #if MIN_VERSION_Cabal(3,8,0) import Distribution.Simple.PackageDescription #endif+#if MIN_VERSION_Cabal(3,14,0)+-- Note [Cabal 3.14]+--+-- If you change any path stuff, either test that the package still works with+-- Cabal 3.12 or stop declaring support for it in cuda.cabal. (If you do the+-- latter, also remove all of the other conditionals in this file.)+-- Note that supporting old versions of Cabal is useful for being able to run+-- e.g. Accelerate on old GPU clusters, which is nice.+import Distribution.Utils.Path (SymbolicPath, FileOrDir(File, Dir), Lib, Include, Pkg, CWD, makeSymbolicPath, interpretSymbolicPath, makeRelativePathEx)+import qualified Distribution.Types.LocalBuildConfig as LBC+#else+#endif  import Control.Exception import Control.Monad@@ -40,6 +61,7 @@ import Data.Function import Data.List import Data.Maybe+import Data.String (fromString) import System.Directory import System.Environment import System.FilePath@@ -67,8 +89,9 @@ main :: IO () main = defaultMainWithHooks customHooks   where+    -- Be careful changing flags/paths stuff here; see Note [Cabal 3.14].     readHook get_verbosity a flags = do-        getHookedBuildInfo (fromFlag (get_verbosity flags))+        getHookedBuildInfo (flagToMaybe (workingDirFlag flags)) (fromFlag (get_verbosity flags))      preprocessors = hookedPreProcessors simpleUserHooks @@ -87,14 +110,16 @@         , preReg              = readHook regVerbosity         , preUnreg            = readHook regVerbosity         , postConf            = postConfHook-        , hookedPreProcessors = ("chs", ppC2hs) : filter (\x -> fst x /= "chs") preprocessors+        , hookedPreProcessors = (fromString "chs", ppC2hs) : filter (\x -> fst x /= fromString "chs") preprocessors         }      -- The hook just loads the HookedBuildInfo generated by postConfHook,     -- unless there is user-provided info that overwrites it.     --     preBuildHook :: Args -> BuildFlags -> IO HookedBuildInfo-    preBuildHook _ flags = getHookedBuildInfo $ fromFlag $ buildVerbosity flags+    preBuildHook _ flags = getHookedBuildInfo cwd verbosity+      where cwd = flagToMaybe (workingDirFlag flags)+            verbosity = fromFlag (buildVerbosity flags)      -- The hook scans system in search for CUDA Toolkit. If the toolkit is not     -- found, an error is raised. Otherwise the toolkit location is used to@@ -103,12 +128,14 @@     postConfHook :: Args -> ConfigFlags -> PackageDescription -> LocalBuildInfo -> IO ()     postConfHook args flags pkg_descr lbi = do       let+          cwd             = flagToMaybe (workingDirFlag flags)           verbosity       = fromFlagOrDefault normal (configVerbosity flags)           profile         = fromFlagOrDefault False  (configProfLib flags)           currentPlatform = hostPlatform lbi           compilerId_     = compilerId (compiler lbi)       --       generateAndStoreBuildInfo+          cwd           verbosity           profile           currentPlatform@@ -118,7 +145,7 @@           generatedBuildInfoFilePath       validateLinker verbosity currentPlatform $ withPrograms lbi       ---      actualBuildInfoToUse <- getHookedBuildInfo verbosity+      actualBuildInfoToUse <- getHookedBuildInfo cwd verbosity       let pkg_descr' = updatePackageDescription actualBuildInfoToUse pkg_descr       postConf simpleUserHooks args flags pkg_descr' lbi @@ -131,27 +158,29 @@ -- visible to underlying build tools. -- libraryBuildInfo-    :: Verbosity+    :: Maybe CWDPath+    -> Verbosity     -> Bool     -> FilePath     -> Platform     -> Version-    -> [FilePath]-    -> [FilePath]+    -> [ExtraLibsPath]+    -> [ExtraIncludesPath]     -> IO HookedBuildInfo-libraryBuildInfo verbosity profile installPath platform@(Platform arch os) ghcVersion extraLibs extraIncludes = do+libraryBuildInfo cwd verbosity profile installPath platform@(Platform arch os) ghcVersion extraLibs extraIncludes = do   let-      libraryPaths      = cudaLibraryPaths platform installPath ++ extraLibs-      includePaths      = cudaIncludePath platform installPath : extraIncludes+      -- Be careful changing flags/paths stuff here; see Note [Cabal 3.14].+      libraryPaths      = map makeSymbolicPath (cudaLibraryPaths platform installPath) ++ extraLibs+      includePaths      = makeSymbolicPath (cudaIncludePath platform installPath) : extraIncludes        takeFirstExisting paths = do-          existing <- filterM doesDirectoryExist libraryPaths+          existing <- filterM (doesDirectoryExist . interpretSymbolicPath cwd) libraryPaths           case existing of                (p0:_) -> return p0                _      -> die' verbosity $ "Could not find path: " ++ show paths    -- This can only be defined once, so take the first path which exists-  canonicalLibraryPath <- takeFirstExisting libraryPaths+  canonicalLibraryPath <- interpretSymbolicPath cwd <$> takeFirstExisting libraryPaths    let       -- OS-specific escaping for -D path defines@@ -163,16 +192,16 @@       extraLibDirs'     = libraryPaths       ccOptions'        = [ "-DCUDA_INSTALL_PATH=\"" ++ escDefPath installPath ++ "\""                           , "-DCUDA_LIBRARY_PATH=\"" ++ escDefPath canonicalLibraryPath ++ "\""-                          ] ++ map ("-I" ++) includePaths-      ldOptions'        = map ("-L" ++) libraryPaths+                          ] ++ map (("-I" ++) . interpretSymbolicPath cwd) includePaths+      ldOptions'        = map (("-L" ++) . interpretSymbolicPath cwd) libraryPaths       ghcOptions        = map ("-optc"++) ccOptions'                        ++ map ("-optl"++) ldOptions'                        ++ if os /= Windows && not profile-                            then map ("-optl-Wl,-rpath,"++) extraLibDirs'+                            then map (("-optl-Wl,-rpath," ++) . interpretSymbolicPath cwd) extraLibDirs'                             else []       extraLibs'        = cudaLibraries platform-      frameworks'       = [ "CUDA" | os == OSX ]-      frameworkDirs'    = [ "/Library/Frameworks" | os == OSX ]+      frameworks'       = [ makeRelativePathEx "CUDA" | os == OSX ]+      frameworkDirs'    = [ makeSymbolicPath "/Library/Frameworks" | os == OSX ]        -- options or c2hs       archFlag          = case arch of@@ -227,7 +256,9 @@         (Windows, X86_64)  -> ["lib/x64"]         (OSX,     _)       -> ["lib"]    -- MacOS does not distinguish 32- vs. 64-bit paths         (_,       X86_64)  -> ["lib64", "lib"]  -- prefer lib64 for 64-bit systems+#if MIN_VERSION_Cabal(2,4,0)         (_,       AArch64) -> ["lib64", "lib"]+#endif         _                  -> ["lib"]           -- otherwise  @@ -238,7 +269,7 @@ cudaLibraries (Platform _ os) =   case os of     OSX -> ["cudadevrt", "cudart_static"]-    _   -> ["cudart", "cuda"]+    _   -> ["cuda"]  cudaGHCiLibraries     :: Platform@@ -248,7 +279,7 @@ cudaGHCiLibraries platform@(Platform _ os) installPath libraries =   case os of     Windows -> cudaGhciLibrariesWindows platform installPath libraries-    OSX     -> return ["cudart"]+    OSX     -> return ["cuda"]     _       -> return []  -- Windows compatibility function.@@ -398,7 +429,7 @@   windowsHelpPage :: String-windowsHelpPage = "https://github.com/tmcdonell/cuda/blob/master/WINDOWS.markdown"+windowsHelpPage = "https://github.com/tmcdonell/cuda/blob/master/cuda/WINDOWS.md"  windowsLinkerBugMsg :: FilePath -> String windowsLinkerBugMsg ldPath = printf (unlines msg) windowsHelpPage ldPath@@ -427,17 +458,18 @@ -- Runs CUDA detection procedure and stores .buildinfo to a file. -- generateAndStoreBuildInfo-    :: Verbosity+    :: Maybe CWDPath+    -> Verbosity     -> Bool     -> Platform     -> CompilerId-    -> [FilePath]-    -> [FilePath]+    -> [ExtraLibsPath]+    -> [ExtraIncludesPath]     -> FilePath     -> IO ()-generateAndStoreBuildInfo verbosity profile platform (CompilerId _ghcFlavor ghcVersion) extraLibs extraIncludes path = do+generateAndStoreBuildInfo cwd verbosity profile platform (CompilerId _ghcFlavor ghcVersion) extraLibs extraIncludes path = do   installPath <- findCUDAInstallPath verbosity platform-  hbi         <- libraryBuildInfo verbosity profile installPath platform ghcVersion extraLibs extraIncludes+  hbi         <- libraryBuildInfo cwd verbosity profile installPath platform ghcVersion extraLibs extraIncludes   storeHookedBuildInfo verbosity path hbi  storeHookedBuildInfo@@ -622,21 +654,22 @@ -- (generated one should be always present, as it is created in the post-conf step) -- getHookedBuildInfo-    :: Verbosity+    :: Maybe CWDPath+    -> Verbosity     -> IO HookedBuildInfo-getHookedBuildInfo verbosity = do-  doesCustomBuildInfoExists <- doesFileExist customBuildInfoFilePath+getHookedBuildInfo cwd verbosity = do+  doesCustomBuildInfoExists <- doesFileExist (customBuildInfoFilePath)   if doesCustomBuildInfoExists     then do       notice verbosity $ printf "The user-provided buildinfo from file %s will be used. To use default settings, delete this file.\n" customBuildInfoFilePath-      readHookedBuildInfo verbosity customBuildInfoFilePath+      readHookedBuildInfoWithCWD verbosity cwd (makeSymbolicPath customBuildInfoFilePath)     else do       doesGeneratedBuildInfoExists <- doesFileExist generatedBuildInfoFilePath       if doesGeneratedBuildInfoExists         then do           notice verbosity $ printf "Using build information from '%s'.\n" generatedBuildInfoFilePath           notice verbosity $ printf "Provide a '%s' file to override this behaviour.\n" customBuildInfoFilePath-          readHookedBuildInfo verbosity generatedBuildInfoFilePath+          readHookedBuildInfoWithCWD verbosity cwd (makeSymbolicPath generatedBuildInfoFilePath)         else           die' verbosity $ printf "Unexpected failure. Neither the default %s nor custom %s exist.\n" generatedBuildInfoFilePath customBuildInfoFilePath @@ -672,7 +705,7 @@ getCppOptions :: BuildInfo -> LocalBuildInfo -> [String] getCppOptions bi lbi     = hcDefines (compiler lbi)-   ++ ["-I" ++ dir | dir <- includeDirs bi]+   ++ ["-I" ++ interpretSymbolicPath (lbiCWD lbi) dir | dir <- includeDirs bi]    ++ [opt | opt@('-':c:_) <- ccOptions bi, c `elem` "DIU"]  hcDefines :: Compiler -> [String]@@ -704,5 +737,65 @@ #if !MIN_VERSION_Cabal(2,0,0) die' :: Verbosity -> String -> IO a die' _ = die+#endif+++-- Compatibility across Cabal 3.14 symbolic paths.+-- If we want to drop pre-Cabal-3.14 compatibility at some point, this should all be merged in above.++lbiCWD :: LocalBuildInfo -> Maybe CWDPath++#if MIN_VERSION_Cabal(3,14,0)+type ExtraLibsPath = SymbolicPath Pkg ('Dir Lib)+type ExtraIncludesPath = SymbolicPath Pkg ('Dir Include)+type CWDPath = SymbolicPath CWD ('Dir Pkg)++regVerbosity :: RegisterFlags -> Flag Verbosity+regVerbosity = setupVerbosity . registerCommonFlags++workingDirFlag :: HasCommonFlags flags => flags -> Flag CWDPath+workingDirFlag = setupWorkingDir . getCommonFlags++lbiCWD = flagToMaybe . setupWorkingDir . configCommonFlags . LBC.configFlags . LBC.packageBuildDescr . localBuildDescr++-- makeSymbolicPath is an actual useful function in Cabal 3.14+-- makeRelativePathEx is an actual useful function in Cabal 3.14+-- interpretSymbolicPath is an actual useful function in Cabal 3.14++class HasCommonFlags flags where getCommonFlags :: flags -> CommonSetupFlags+instance HasCommonFlags BuildFlags where getCommonFlags = buildCommonFlags+instance HasCommonFlags CleanFlags where getCommonFlags = cleanCommonFlags+instance HasCommonFlags ConfigFlags where getCommonFlags = configCommonFlags+instance HasCommonFlags CopyFlags where getCommonFlags = copyCommonFlags+instance HasCommonFlags InstallFlags where getCommonFlags = installCommonFlags+instance HasCommonFlags HscolourFlags where getCommonFlags = hscolourCommonFlags+instance HasCommonFlags HaddockFlags where getCommonFlags = haddockCommonFlags+instance HasCommonFlags RegisterFlags where getCommonFlags = registerCommonFlags++readHookedBuildInfoWithCWD :: Verbosity -> Maybe CWDPath -> SymbolicPath Pkg 'File -> IO HookedBuildInfo+readHookedBuildInfoWithCWD = readHookedBuildInfo+#else+type ExtraLibsPath = FilePath+type ExtraIncludesPath = FilePath+type CWDPath = ()++-- regVerbosity is still present as an actual field in Cabal 3.12++workingDirFlag :: flags -> Flag CWDPath+workingDirFlag _ = NoFlag++lbiCWD _ = Nothing++makeSymbolicPath :: FilePath -> FilePath+makeSymbolicPath = id++makeRelativePathEx :: FilePath -> FilePath+makeRelativePathEx = id++interpretSymbolicPath :: Maybe CWDPath -> FilePath -> FilePath+interpretSymbolicPath _ = id++readHookedBuildInfoWithCWD :: Verbosity -> Maybe CWDPath -> FilePath -> IO HookedBuildInfo+readHookedBuildInfoWithCWD verb _ path = readHookedBuildInfo verb path #endif 

@@ -3,24 +3,7 @@  */  #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};--    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+#include <string.h>  // memset  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) {@@ -196,7 +179,13 @@  CUresult CUDAAPI cuCtxCreate(CUcontext *pctx, unsigned int flags, CUdevice dev) {+#if CUDA_VERSION >= 13000+    CUctxCreateParams params;+    memset(&params, 0, sizeof params);+    return cuCtxCreate_v4(pctx, &params, flags, dev);+#else     return cuCtxCreate_v2(pctx, flags, dev);+#endif }  CUresult CUDAAPI cuModuleGetGlobal(CUdeviceptr *dptr, size_t *bytes, CUmodule hmod, const char *name)@@ -424,3 +413,17 @@ } #endif +#if CUDA_VERSION >= 13000+// This is the signature of the CUDA <=12 version; much easier to shim here than in Haskell.+CUresult cuMemAdvise_device(CUdeviceptr dptr, size_t count, CUmem_advise advice, CUdevice device)+{+  return cuMemAdvise(dptr, count, advice, (CUmemLocation){.id = device, .type = CU_MEM_LOCATION_TYPE_DEVICE});+}++// This is the signature of the CUDA <=12 version; much easier to shim here than in Haskell.+CUresult cuMemPrefetchAsync_device(CUdeviceptr dptr, size_t count, CUdevice device, CUstream hStream)+{+  // flags is reserved and must be 0 in CUDA 13+  return cuMemPrefetchAsync(dptr, count, (CUmemLocation){.id = device, .type = CU_MEM_LOCATION_TYPE_DEVICE}, 0, hStream);+}+#endif

@@ -21,12 +21,6 @@  void enable_constructors(); -#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);@@ -182,6 +176,11 @@ #if CUDA_VERSION >= 11010 #undef cuIpcOpenMemHandle CUresult CUDAAPI cuIpcOpenMemHandle(CUdeviceptr *pdptr, CUipcMemHandle handle, unsigned int Flags);+#endif++#if CUDA_VERSION >= 13000+CUresult cuMemAdvise_device(CUdeviceptr dptr, size_t count, CUmem_advise advice, CUdevice device);+CUresult cuMemPrefetchAsync_device(CUdeviceptr dptr, size_t count, CUdevice device, CUstream hStream); #endif  #ifdef __cplusplus

@@ -1,8 +1,8 @@ cabal-version:          1.24  Name:                   cuda-Version:                0.12.8.0-Synopsis:               FFI binding to the CUDA interface for programming NVIDIA GPUs+Version:                0.13.0.0+Synopsis:               FFI binding to the CUDA driver interface for programming NVIDIA GPUs Description:     The CUDA library provides a direct, general purpose C-like SPMD programming     model for NVIDIA graphics cards (G8x series onwards). This is a collection@@ -21,18 +21,13 @@     .       3. Checking at @\/usr\/local\/cuda@     .-      4. @CUDA_PATH_Vx_y@ environment variable, for recent CUDA toolkit versions x.y-    .-    This library provides bindings to both the CUDA Driver and Runtime APIs. To-    get started, see one of:-    .-    * "Foreign.CUDA.Driver" (a short tutorial is available here)+      4. Environment variables of the form @CUDA_PATH_Vx_y@ (deprecated)     .-    * "Foreign.CUDA.Runtime"+    This library provides bindings to the CUDA Driver API, not the Runtime API.+    To get started, see "Foreign.CUDA.Driver"; a short tutorial is available+    there.     .-    Tested with library versions up to CUDA-12.8. See also the-    <https://travis-ci.org/tmcdonell/cuda travis-ci.org> build matrix for-    version compatibility.+    Tested with library versions up to and including CUDA-13.0.     .     [/NOTES:/]     .@@ -40,16 +35,18 @@     .     * <https://github.com/tmcdonell/cuda/blob/master/WINDOWS.md>     .-    This library is currently in __maintenance mode__. While we plan to release-    updates to keep the existing interface working with newer CUDA versions (as-    long as the underlying APIs remain available), no binding of new features is-    planned at the moment. Get in touch if you want to contribute.+    This library is currently in __maintenance mode__. While we are happy to+    release updates to keep the existing interface working with newer CUDA+    versions (as long as the underlying APIs remain available), no binding of+    new features is planned at the moment. Get in touch if you want to+    contribute.  License:                BSD3 License-file:           LICENSE Copyright:              Copyright (c) [2009..2023]. Trevor L. McDonell <trevor.mcdonell@gmail.com> Author:                 Trevor L. McDonell <trevor.mcdonell@gmail.com>-Maintainer:             Trevor L. McDonell <trevor.mcdonell@gmail.com>+Maintainer:             Trevor L. McDonell <trevor.mcdonell@gmail.com>,+                        Tom Smeding <tom@tomsmeding.com> Homepage:               https://github.com/tmcdonell/cuda Bug-reports:            https://github.com/tmcdonell/cuda/issues Category:               Foreign@@ -61,6 +58,8 @@  Extra-source-files:   cbits/stubs.h++Extra-doc-files:   CHANGELOG.md   README.md   WINDOWS.md@@ -68,14 +67,13 @@ custom-setup   setup-depends:       base              >= 4.7  && < 5-    , Cabal             >= 1.24 && < 3.11+    , Cabal             >= 1.24 && < 3.17     , directory         >= 1.0     , filepath          >= 1.0  Library   hs-source-dirs:       src   exposed-modules:-      Foreign.CUDA       Foreign.CUDA.Path       Foreign.CUDA.Ptr @@ -111,22 +109,10 @@       Foreign.CUDA.Driver.Unified       Foreign.CUDA.Driver.Utils -      -- Runtime API-      Foreign.CUDA.Runtime-      Foreign.CUDA.Runtime.Device-      Foreign.CUDA.Runtime.Error-      Foreign.CUDA.Runtime.Event-      Foreign.CUDA.Runtime.Exec-      Foreign.CUDA.Runtime.Marshal-      Foreign.CUDA.Runtime.Stream-      Foreign.CUDA.Runtime.Utils-       -- Extras       Foreign.C.Extra--  other-modules:       Foreign.CUDA.Internal.C2HS-      Text.Show.Describe+      Foreign.CUDA.Internal.Describe    include-dirs:         .   c-sources:            cbits/stubs.c@@ -150,10 +136,6 @@       -fwarn-tabs       -fno-warn-unused-imports -  ghc-prof-options:-      -fprof-auto-      -fprof-cafs-   if impl(ghc == 8.0.1)     cpp-options:        -DCUDA_PRELOAD @@ -177,6 +159,6 @@ source-repository this     type:               git     location:           https://github.com/tmcdonell/cuda-    tag:                v0.12.8.0+    tag:                v0.13.0.0  -- vim: nospell

@@ -5,6 +5,7 @@ module Main where  import Control.Monad+import Foreign.Marshal.Utils                            ( toBool ) import Numeric import Prelude                                          hiding ( (<>) ) import Text.PrettyPrint@@ -69,22 +70,34 @@         ,("  2D:",                                      grid maxTextureDim2D)         ,("  3D:",                                      cube maxTextureDim3D)         ,("Texture alignment:",                         text $ showBytes textureAlignment)-        ,("Maximum memory pitch:",                      text $ showBytes memPitch)-        ,("Concurrent kernel execution:",               bool concurrentKernels)-        ,("Concurrent copy and execution:",             bool deviceOverlap <> text (printf ", with %d copy engine%s" asyncEngineCount (if asyncEngineCount > 1 then "s" else "")))-        ,("Runtime limit on kernel execution:",         bool kernelExecTimeoutEnabled)+        ,("Maximum memory pitch:",                      text $ showBytes memPitch)]++++        $(if CUDA.libraryVersion >= 13000 then [|+        [("Concurrent copy and kernel execution:",      bool concurrentKernels <> text (printf " with %d copy engine%s" asyncEngineCount (if asyncEngineCount > 1 then "s" else "")))]+        |] else [|+        [("Concurrent kernel execution:",               bool concurrentKernels)+        ,("Concurrent copy and execution:",             bool deviceOverlap <> text (printf ", with %d copy engine%s" asyncEngineCount (if asyncEngineCount > 1 then "s" else "")))]+        |])++++        [("Runtime limit on kernel execution:",         bool kernelExecTimeoutEnabled)         ,("Integrated GPU sharing host memory:",        bool integrated)         ,("Host page-locked memory mapping:",           bool canMapHostMemory)         ,("ECC memory support:",                        bool eccEnabled)         ,("Unified addressing (UVA):",                  bool unifiedAddressing)]++-#if __GLASGOW_HASKELL__ > 710+         $(if CUDA.libraryVersion >= 8000 then [|         [("Single to double precision performance:",    text $ printf "%d : 1" singleToDoublePerfRatio)-        ,("Supports compute pre-emption:",              bool preemption)]|] else [|[]|])+++        ,("Supports compute pre-emption:",              bool preemption)]+        |] else [|[]|])+++         $(if CUDA.libraryVersion >= 9000 then [|-        [("Supports cooperative launch:",               bool cooperativeLaunch)-        ,("Supports multi-device cooperative launch:",  bool cooperativeLaunchMultiDevice)]|] else [|[]|])++-#endif+        [("Supports cooperative launch:",               bool cooperativeLaunch)]+        |] else [|[]|])++++        $(if CUDA.libraryVersion >= 9000 && CUDA.libraryVersion < 13000 then [|+        [("Supports multi-device cooperative launch:",  bool cooperativeLaunchMultiDevice)]+        |] else [|[]|])+++         [("PCI bus/location:",                          int (busID pciInfo) <> char '/' <> int (deviceID pciInfo))         ,("Compute mode:",                              text (show computeMode))         ]

@@ -1,19 +0,0 @@------------------------------------------------------------------------------------ |--- Module    : Foreign.CUDA--- Copyright : [2009..2023] Trevor L. McDonell--- License   : BSD------ Top level bindings. By default, expose the C-for-CUDA runtime API bindings,--- as they are slightly more user friendly.--------------------------------------------------------------------------------------module Foreign.CUDA (--  module Foreign.CUDA.Runtime--) where--import Foreign.CUDA.Runtime-

@@ -23,7 +23,7 @@ import Data.Set (Set) import Data.Int import Data.IORef-import Text.Show.Describe+import Foreign.CUDA.Internal.Describe import System.IO.Unsafe  import Debug.Trace@@ -32,13 +32,19 @@ -- | -- The compute mode the device is currently in --+#if CUDA_VERSION < 13000 {# enum CUcomputemode as ComputeMode     { underscoreToCase }     with prefix="CU_COMPUTEMODE" deriving (Eq, Show) #}+#else+{# enum cudaComputeMode as ComputeMode+    { }+    with prefix="cudaComputeMode" deriving (Eq, Show) #}+#endif  instance Describe ComputeMode where   describe Default          = "Multiple contexts are allowed on the device simultaneously"-#if CUDA_VERSION < 8000+#if CUDA_VERSION < 8000 || CUDA_VERSION >= 13000   describe Exclusive        = "Only one context used by a single thread can be present on this device at a time" #endif   describe Prohibited       = "No contexts can be created on this device at this time"@@ -69,7 +75,13 @@ --}  -- |--- The properties of a compute device+-- Some properties of a compute device. Originally, this mirrored+-- @struct cudaDeviceProp@ in CUDA, but since CUDA 13 some fields have been+-- removed from that struct, and this data type keeps them for backwards+-- compatibility.+--+-- There are more device "properties" than those listed in this data type; use+-- 'Foreign.CUDA.Driver.Device.attribute' to query them. -- data DeviceProperties = DeviceProperties   {

@@ -31,7 +31,7 @@ -- Next, we must select a GPU that we will execute operations on. Each GPU -- is assigned a unique identifier (beginning at zero). We can get a handle -- to a compute device at a given ordinal using the 'device' operation.--- Given a device handle, we can query the properties of that device using+-- Given a device handle, we can query some properties of that device using -- 'props'. The number of available CUDA-capable devices is given via -- 'count'. For example: --

@@ -246,8 +246,11 @@     peekS s _ = peekCString s  --- | Returns a UUID for the device+-- | Returns a UUID for the device. --+-- Since CUDA-13: If the device is in MIG mode, this function returns its MIG+-- UUID which uniquely identifies the subscribed MIG compute instance.+-- -- Requires CUDA-9.2 -- -- <https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__DEVICE.html#group__CUDA__DEVICE_1g987b46b884c101ed5be414ab4d9e60e4>@@ -265,10 +268,17 @@     unpack ptr   where     {-# INLINE cuDeviceGetUuid #-}+#if CUDA_VERSION < 13000     {# fun unsafe cuDeviceGetUuid       {           `Ptr ()'       , useDevice `Device'       } -> `()' checkStatus*- #}+#else+    {# fun unsafe cuDeviceGetUuid_v2 as cuDeviceGetUuid+      {           `Ptr ()'+      , useDevice `Device'+      } -> `()' checkStatus*- #}+#endif      {-# INLINE unpack #-}     unpack :: Ptr () -> IO UUID

@@ -23,7 +23,7 @@  -- Friends import Foreign.CUDA.Internal.C2HS-import Text.Show.Describe+import Foreign.CUDA.Internal.Describe  -- System import Control.Exception

@@ -114,10 +114,17 @@ elapsedTime !ev1 !ev2 = resultIfOk =<< cuEventElapsedTime ev1 ev2  {-# INLINE cuEventElapsedTime #-}+#if CUDA_VERSION < 13000 {# fun unsafe cuEventElapsedTime   { alloca-  `Float' peekFloatConv*   , useEvent `Event'   , useEvent `Event'                } -> `Status' cToEnum #}+#else+{# fun unsafe cuEventElapsedTime_v2 as cuEventElapsedTime+  { alloca-  `Float' peekFloatConv*+  , useEvent `Event'+  , useEvent `Event'                } -> `Status' cToEnum #}+#endif   -- |

@@ -204,13 +204,24 @@   where     (from, to) = unzip deps +#if CUDA_VERSION < 13000     {# fun unsafe cuGraphAddDependencies       { useGraph          `Graph'       , withNodeArray*    `[Node]'       , withNodeArrayLen* `[Node]'&       }       -> `()' checkStatus*- #}+#else+    cuGraphAddDependencies g' from' to' = cuGraphAddDependencies_v2 g' from' to' (length deps)+    {# fun unsafe cuGraphAddDependencies_v2+      { useGraph          `Graph'+      , withNodeArray*    `[Node]'+      , withNodeArray*    `[Node]'+      , withNullEdgeDataLen* `Int'&+      }+      -> `()' checkStatus*- #} #endif+#endif   -- | Remove dependency edges from the graph@@ -230,13 +241,24 @@   where     (from, to) = unzip deps +#if CUDA_VERSION < 13000     {# fun unsafe cuGraphRemoveDependencies       { useGraph          `Graph'       , withNodeArray*    `[Node]'       , withNodeArrayLen* `[Node]'&       }       -> `()' checkStatus*- #}+#else+    cuGraphRemoveDependencies g' from' to' = cuGraphRemoveDependencies_v2 g' from' to' (length deps)+    {# fun unsafe cuGraphRemoveDependencies_v2+      { useGraph          `Graph'+      , withNodeArray*    `[Node]'+      , withNodeArray*    `[Node]'+      , withNullEdgeDataLen* `Int'&+      }+      -> `()' checkStatus*- #} #endif+#endif   -- | Create an empty node and add it to the graph.@@ -507,6 +529,7 @@        to   <- peekArray count p_to        return $ zip from to   where+#if CUDA_VERSION < 13000     {# fun unsafe cuGraphGetEdges       { useGraph     `Graph'       , castPtr      `Ptr Node'@@ -514,7 +537,18 @@       , id           `Ptr CULong'       }       -> `()' checkStatus*- #}+#else+    cuGraphGetEdges g' f t c = cuGraphGetEdges_v2 g' f t nullPtr c+    {# fun unsafe cuGraphGetEdges_v2+      { useGraph     `Graph'+      , castPtr      `Ptr Node'+      , castPtr      `Ptr Node'+      , castPtr      `Ptr edgeData'+      , id           `Ptr CULong'+      }+      -> `()' checkStatus*- #} #endif+#endif   -- | Return a graph's nodes@@ -598,13 +632,24 @@       cuGraphNodeGetDependencies n p_deps p_count       peekArray count p_deps   where+#if CUDA_VERSION < 13000     {# fun unsafe cuGraphNodeGetDependencies       { useNode `Node'       , castPtr `Ptr Node'       , id      `Ptr CULong'       }       -> `()' checkStatus*- #}+#else+    cuGraphNodeGetDependencies n' d c = cuGraphNodeGetDependencies_v2 n' d nullPtr c+    {# fun unsafe cuGraphNodeGetDependencies_v2+      { useNode `Node'+      , castPtr `Ptr Node'+      , castPtr `Ptr edgeData'+      , id      `Ptr CULong'+      }+      -> `()' checkStatus*- #} #endif+#endif   -- | Return a node's dependent nodes@@ -628,13 +673,24 @@       cuGraphNodeGetDependentNodes n p_deps p_count       peekArray count p_deps   where+#if CUDA_VERSION < 13000     {# fun unsafe cuGraphNodeGetDependentNodes       { useNode `Node'       , castPtr `Ptr Node'       , id      `Ptr CULong'       }       -> `()' checkStatus*- #}+#else+    cuGraphNodeGetDependentNodes n' d c = cuGraphNodeGetDependentNodes_v2 n' d nullPtr c+    {# fun unsafe cuGraphNodeGetDependentNodes_v2+      { useNode `Node'+      , castPtr `Ptr Node'+      , castPtr `Ptr edgeData'+      , id      `Ptr CULong'+      }+      -> `()' checkStatus*- #} #endif+#endif   -- | Find a cloned version of a node@@ -679,5 +735,11 @@ {-# INLINE withNodeArrayLen #-} withNodeArrayLen :: [Node] -> ((Ptr {# type CUgraphNode #}, CULong) -> IO a) -> IO a withNodeArrayLen ns f = withArrayLen ns $ \i p -> f (castPtr p, cIntConv i)+#endif++#if CUDA_VERSION >= 13000+{-# INLINE withNullEdgeDataLen #-}+withNullEdgeDataLen :: Int -> ((Ptr (), CULong) -> IO a) -> IO a+withNullEdgeDataLen len f = f (nullPtr, cIntConv len) #endif 

@@ -143,6 +143,10 @@ -- | Query the capture status of a stream and get an id for the capture -- sequence, which is unique over the lifetime of the process. --+-- Since CUDA-13, "edge data" can be associated with an edge in the graph. This+-- function assumes no such data is present (if there is, a CUDA error+-- (@CUDA_ERROR_LOSSY_QUERY@) will be raised).+-- -- Requires CUDA-10.1 -- -- <https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__STREAM.html#group__CUDA__STREAM_1g13145ece1d79a1d79a1d22abb9663216>@@ -152,24 +156,38 @@ #if CUDA_VERSION < 10010 info :: Stream -> IO (Status, Int64) info = requireSDK 'info 10.1-#elif CUDA_VERSION < 12000+-- Not another elif because c2hs seems to be buggy+#else+#if CUDA_VERSION < 12000 {# fun unsafe cuStreamGetCaptureInfo as info   { useStream `Stream'   , alloca-   `Status' peekEnum*   , alloca-   `Int64'  peekIntConv*   }   -> `()' checkStatus*- #}-#else+#elif CUDA_VERSION < 13000 {# fun unsafe cuStreamGetCaptureInfo_v2 as info   { useStream `Stream'   , alloca-   `Status' peekEnum*   , alloca-   `Int64'  peekIntConv*-  , alloca-   `Graph'-  , alloca-   `Node'-  , alloca-   `CSize'+  , withNullPtr- `Graph'+  , withNullPtr- `Node'+  , withNullPtr- `CSize'   }   -> `()' checkStatus*- #}+#else+{# fun unsafe cuStreamGetCaptureInfo_v3 as info+  { useStream `Stream'+  , alloca-   `Status' peekEnum*+  , alloca-   `Int64'  peekIntConv*+  , withNullPtr- `Graph'+  , withNullPtr- `()'+  , withNullPtr- `Node'+  , withNullPtr- `CSize'+  }+  -> `()' checkStatus*- #} #endif+#endif   -- | Set the stream capture interaction mode for this thread. Return the previous value.@@ -201,3 +219,5 @@ peekGraph = liftM Graph . peek #endif +withNullPtr :: (Ptr a -> r) -> r+withNullPtr f = f nullPtr

@@ -358,6 +358,7 @@     go x _ = nothingIfOk =<< cuMemPrefetchAsync ptr (n * sizeOf x) (maybe (-1) useDevice mdev) (fromMaybe defaultStream mst)  {-# INLINE cuMemPrefetchAsync #-}+#if CUDA_VERSION < 13000 {# fun unsafe cuMemPrefetchAsync   { useDeviceHandle `DevicePtr a'   ,                 `Int'@@ -365,6 +366,15 @@   , useStream       `Stream'   }   -> `Status' cToEnum #}+#else+{# fun unsafe cuMemPrefetchAsync_device as cuMemPrefetchAsync+  { useDeviceHandle `DevicePtr a'+  ,                 `Int'+  , id              `CInt'+  , useStream       `Stream'+  }+  -> `Status' cToEnum #}+#endif #endif  

@@ -258,6 +258,7 @@     go x _ = nothingIfOk =<< cuMemAdvise ptr (n * sizeOf x) a (maybe (-1) useDevice mdev)  {-# INLINE cuMemAdvise #-}+#if CUDA_VERSION < 13000 {# fun unsafe cuMemAdvise   { useHandle `Ptr a'   ,           `Int'@@ -265,7 +266,15 @@   ,           `CInt'   }   -> `Status' cToEnum #}+#else+{# fun unsafe cuMemAdvise_device as cuMemAdvise+  { useHandle `Ptr a'+  ,           `Int'+  , cFromEnum `Advice'+  ,           `CInt'+  }+  -> `Status' cToEnum #}+#endif   where     useHandle = fromIntegral . ptrToIntPtr #endif-

@@ -0,0 +1,18 @@+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Internal.Describe+-- Copyright : [2016..2023] Trevor L. McDonell+-- License   : BSD+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Internal.Describe+  where+++-- | Like 'Text.Show.Show', but focuses on providing a more detailed description+-- of the value rather than a 'Text.Read.read'able representation.+--+class Describe a where+    describe :: a -> String+

@@ -1,28 +0,0 @@------------------------------------------------------------------------------------ |--- Module    : Foreign.CUDA.Runtime--- Copyright : [2009..2023] Trevor L. McDonell--- License   : BSD------ Top level bindings to the C-for-CUDA runtime API--------------------------------------------------------------------------------------module Foreign.CUDA.Runtime (--  module Foreign.CUDA.Ptr,-  module Foreign.CUDA.Runtime.Device,-  module Foreign.CUDA.Runtime.Error,-  module Foreign.CUDA.Runtime.Exec,-  module Foreign.CUDA.Runtime.Marshal,-  module Foreign.CUDA.Runtime.Utils--) where--import Foreign.CUDA.Ptr-import Foreign.CUDA.Runtime.Device-import Foreign.CUDA.Runtime.Error-import Foreign.CUDA.Runtime.Exec-import Foreign.CUDA.Runtime.Marshal-import Foreign.CUDA.Runtime.Utils-

@@ -1,440 +0,0 @@-{-# LANGUAGE BangPatterns             #-}-{-# LANGUAGE CPP                      #-}-{-# LANGUAGE EmptyDataDecls           #-}-{-# LANGUAGE ForeignFunctionInterface #-}-{-# LANGUAGE RecordWildCards          #-}-{-# LANGUAGE TemplateHaskell          #-}-{-# OPTIONS_GHC -fno-warn-orphans #-}-#ifdef USE_EMPTY_CASE-{-# LANGUAGE EmptyCase                #-}-#endif------------------------------------------------------------------------------------ |--- Module    : Foreign.CUDA.Runtime.Device--- Copyright : [2009..2023] Trevor L. McDonell--- License   : BSD------ Device management routines--------------------------------------------------------------------------------------module Foreign.CUDA.Runtime.Device (--  -- * Device Management-  Device, DeviceFlag(..), DeviceProperties(..), Compute(..), ComputeMode(..),-  choose, get, count, props, set, setFlags, setOrder, reset, sync,--  -- * Peer Access-  PeerFlag,-  accessible, add, remove,--  -- * Cache Configuration-  Limit(..),-  getLimit, setLimit--) where--#include "cbits/stubs.h"-{# context lib="cudart" #}---- Friends-import Foreign.CUDA.Analysis.Device-import Foreign.CUDA.Runtime.Error-import Foreign.CUDA.Internal.C2HS---- System-import Control.Applicative-import Foreign-import Foreign.C-import Prelude--#c-typedef struct cudaDeviceProp   cudaDeviceProp;--typedef enum-{-    cudaDeviceFlagScheduleAuto    = cudaDeviceScheduleAuto,-    cudaDeviceFlagScheduleSpin    = cudaDeviceScheduleSpin,-    cudaDeviceFlagScheduleYield   = cudaDeviceScheduleYield,-    cudaDeviceFlagBlockingSync    = cudaDeviceBlockingSync,-    cudaDeviceFlagMapHost         = cudaDeviceMapHost,-#if CUDART_VERSION >= 3000-    cudaDeviceFlagLMemResizeToMax = cudaDeviceLmemResizeToMax-#endif-} cudaDeviceFlags;-#endc-------------------------------------------------------------------------------------- Data Types------------------------------------------------------------------------------------- |--- A device identifier----type Device = Int--{# pointer *cudaDeviceProp as ^ foreign -> DeviceProperties nocode #}---- |--- Device execution flags----{# enum cudaDeviceFlags as DeviceFlag { }-    with prefix="cudaDeviceFlag" deriving (Eq, Show, Bounded) #}---instance Storable DeviceProperties where-  sizeOf _    = {#sizeof cudaDeviceProp#}-  alignment _ = alignment (undefined :: Ptr ())--  poke _ _    = error "no instance for Foreign.Storable.poke DeviceProperties"-  peek p      = do-    deviceName                    <- peekCString =<< {#get cudaDeviceProp.name#} p-    computeCapability             <- Compute <$> (fromIntegral <$> {#get cudaDeviceProp.major#} p)-                                             <*> (fromIntegral <$> {#get cudaDeviceProp.minor#} p)-    totalGlobalMem                <- cIntConv <$> {#get cudaDeviceProp.totalGlobalMem#} p-    sharedMemPerBlock             <- cIntConv <$> {#get cudaDeviceProp.sharedMemPerBlock#} p-    regsPerBlock                  <- cIntConv <$> {#get cudaDeviceProp.regsPerBlock#} p-    warpSize                      <- cIntConv <$> {#get cudaDeviceProp.warpSize#} p-    memPitch                      <- cIntConv <$> {#get cudaDeviceProp.memPitch#} p-    maxThreadsPerBlock            <- cIntConv <$> {#get cudaDeviceProp.maxThreadsPerBlock#} p-    clockRate                     <- cIntConv <$> {#get cudaDeviceProp.clockRate#} p-    totalConstMem                 <- cIntConv <$> {#get cudaDeviceProp.totalConstMem#} p-    textureAlignment              <- cIntConv <$> {#get cudaDeviceProp.textureAlignment#} p-    deviceOverlap                 <- cToBool  <$> {#get cudaDeviceProp.deviceOverlap#} p-    multiProcessorCount           <- cIntConv <$> {#get cudaDeviceProp.multiProcessorCount#} p-    kernelExecTimeoutEnabled      <- cToBool  <$> {#get cudaDeviceProp.kernelExecTimeoutEnabled#} p-    integrated                    <- cToBool  <$> {#get cudaDeviceProp.integrated#} p-    canMapHostMemory              <- cToBool  <$> {#get cudaDeviceProp.canMapHostMemory#} p-    computeMode                   <- cToEnum  <$> {#get cudaDeviceProp.computeMode#} p-#if CUDART_VERSION >= 3000-    concurrentKernels             <- cToBool  <$> {#get cudaDeviceProp.concurrentKernels#} p-    maxTextureDim1D               <- cIntConv <$> {#get cudaDeviceProp.maxTexture1D#} p-#endif-#if CUDART_VERSION >= 3000 && CUDART_VERSION < 3010-    -- XXX: not visible from CUDA runtime API 3.0 (only accessible from the driver API)-    let eccEnabled = False-#endif-#if CUDART_VERSION >= 3010-    eccEnabled                    <- cToBool  <$> {#get cudaDeviceProp.ECCEnabled#} p-#endif-#if CUDART_VERSION >= 4000-    asyncEngineCount              <- cIntConv <$> {#get cudaDeviceProp.asyncEngineCount#} p-    cacheMemL2                    <- cIntConv <$> {#get cudaDeviceProp.l2CacheSize#} p-    maxThreadsPerMultiProcessor   <- cIntConv <$> {#get cudaDeviceProp.maxThreadsPerMultiProcessor#} p-    memBusWidth                   <- cIntConv <$> {#get cudaDeviceProp.memoryBusWidth#} p-    memClockRate                  <- cIntConv <$> {#get cudaDeviceProp.memoryClockRate#} p-    pciInfo                       <- PCI <$> (cIntConv <$> {#get cudaDeviceProp.pciBusID#} p)-                                         <*> (cIntConv <$> {#get cudaDeviceProp.pciDeviceID#} p)-                                         <*> (cIntConv <$> {#get cudaDeviceProp.pciDomainID#} p)-    tccDriverEnabled              <- cToBool <$> {#get cudaDeviceProp.tccDriver#} p-    unifiedAddressing             <- cToBool <$> {#get cudaDeviceProp.unifiedAddressing#} p-#endif-    [t1,t2,t3]                    <- peekArrayWith cIntConv 3 =<< {#get cudaDeviceProp.maxThreadsDim#} p-    [g1,g2,g3]                    <- peekArrayWith cIntConv 3 =<< {#get cudaDeviceProp.maxGridSize#} p-    let maxBlockSize = (t1,t2,t3)-        maxGridSize  = (g1,g2,g3)-#if CUDART_VERSION >= 3000-    [u21,u22]                     <- peekArrayWith cIntConv 2 =<< {#get cudaDeviceProp.maxTexture2D#} p-    [u31,u32,u33]                 <- peekArrayWith cIntConv 3 =<< {#get cudaDeviceProp.maxTexture3D#} p-    let maxTextureDim2D = (u21,u22)-        maxTextureDim3D = (u31,u32,u33)-#endif-#if CUDART_VERSION >= 5050-    streamPriorities              <- cToBool  <$> {#get cudaDeviceProp.streamPrioritiesSupported#} p-#endif-#if CUDART_VERSION >= 6000-    globalL1Cache                 <- cToBool  <$> {#get cudaDeviceProp.globalL1CacheSupported#} p-    localL1Cache                  <- cToBool  <$> {#get cudaDeviceProp.localL1CacheSupported#} p-    managedMemory                 <- cToBool  <$> {#get cudaDeviceProp.managedMemory#} p-    multiGPUBoard                 <- cToBool  <$> {#get cudaDeviceProp.isMultiGpuBoard#} p-    multiGPUBoardGroupID          <- cIntConv <$> {#get cudaDeviceProp.multiGpuBoardGroupID#} p-#endif-#if CUDART_VERSION >= 8000-#if CUDART_VERSION == 8000-    -- XXX: Not visible from the CUDA runtime API 8.0 (only accessible from the driver API)-    let preemption = False-#else-    preemption                    <- cToBool  <$> {#get cudaDeviceProp.computePreemptionSupported#} p-#endif-    singleToDoublePerfRatio       <- cIntConv <$> {#get cudaDeviceProp.singleToDoublePrecisionPerfRatio#} p-#endif-#if CUDART_VERSION >= 9000-    cooperativeLaunch             <- cToBool <$> {#get cudaDeviceProp.cooperativeLaunch#} p-    cooperativeLaunchMultiDevice  <- cToBool <$> {#get cudaDeviceProp.cooperativeMultiDeviceLaunch#} p-#endif--    return DeviceProperties{..}-------------------------------------------------------------------------------------- Device Management------------------------------------------------------------------------------------- |--- Select the compute device which best matches the given criteria----{-# INLINEABLE choose #-}-choose :: DeviceProperties -> IO Device-choose !dev = resultIfOk =<< cudaChooseDevice dev--{-# INLINE cudaChooseDevice #-}-{# fun unsafe cudaChooseDevice-  { alloca-      `Int'              peekIntConv*-  , withDevProp* `DeviceProperties'              } -> `Status' cToEnum #}-  where-      withDevProp = with----- |--- Returns which device is currently being used----{-# INLINEABLE get #-}-get :: IO Device-get = resultIfOk =<< cudaGetDevice--{-# INLINE cudaGetDevice #-}-{# fun unsafe cudaGetDevice-  { alloca- `Int' peekIntConv* } -> `Status' cToEnum #}----- |--- Returns the number of devices available for execution, with compute--- capability >= 1.0----{-# INLINEABLE count #-}-count :: IO Int-count = resultIfOk =<< cudaGetDeviceCount--{-# INLINE cudaGetDeviceCount #-}-{# fun unsafe cudaGetDeviceCount-  { alloca- `Int' peekIntConv* } -> `Status' cToEnum #}----- |--- Return information about the selected compute device----{-# INLINEABLE props #-}-props :: Device -> IO DeviceProperties-props !n = resultIfOk =<< cudaGetDeviceProperties n--{-# INLINE cudaGetDeviceProperties #-}-#if CUDA_VERSION < 12000-{# fun unsafe cudaGetDeviceProperties-  { alloca- `DeviceProperties' peek*-  ,         `Int'                    } -> `Status' cToEnum #}-#else-{# fun unsafe cudaGetDeviceProperties_v2 as cudaGetDeviceProperties-  { alloca- `DeviceProperties' peek*-  ,         `Int'                    } -> `Status' cToEnum #}-#endif----- |--- Set device to be used for GPU execution----{-# INLINEABLE set #-}-set :: Device -> IO ()-set !n = nothingIfOk =<< cudaSetDevice n--{-# INLINE cudaSetDevice #-}-{# fun unsafe cudaSetDevice-  { `Int' } -> `Status' cToEnum #}----- |--- Set flags to be used for device executions----{-# INLINEABLE setFlags #-}-setFlags :: [DeviceFlag] -> IO ()-setFlags !f = nothingIfOk =<< cudaSetDeviceFlags (combineBitMasks f)--{-# INLINE cudaSetDeviceFlags #-}-{# fun unsafe cudaSetDeviceFlags-  { `Int' } -> `Status' cToEnum #}----- |--- Set list of devices for CUDA execution in priority order----{-# INLINEABLE setOrder #-}-setOrder :: [Device] -> IO ()-setOrder !l = nothingIfOk =<< cudaSetValidDevices l (length l)--{-# INLINE cudaSetValidDevices #-}-{# fun unsafe cudaSetValidDevices-  { withArrayIntConv* `[Int]'-  ,                   `Int'   } -> `Status' cToEnum #}-  where-      withArrayIntConv = withArray . map cIntConv---- |--- Block until the device has completed all preceding requested tasks. Returns--- an error if one of the tasks fails.----{-# INLINEABLE sync #-}-sync :: IO ()-#if CUDART_VERSION < 4000-{-# INLINE cudaThreadSynchronize #-}-sync = nothingIfOk =<< cudaThreadSynchronize-{# fun cudaThreadSynchronize { } -> `Status' cToEnum #}-#else-{-# INLINE cudaDeviceSynchronize #-}-sync = nothingIfOk =<< cudaDeviceSynchronize-{# fun cudaDeviceSynchronize { } -> `Status' cToEnum #}-#endif---- |--- Explicitly destroys and cleans up all runtime resources associated with the--- current device in the current process. Any subsequent API call will--- reinitialise the device.------ Note that this function will reset the device immediately. It is the caller’s--- responsibility to ensure that the device is not being accessed by any other--- host threads from the process when this function is called.----{-# INLINEABLE reset #-}-reset :: IO ()-#if CUDART_VERSION >= 4000-{-# INLINE cudaDeviceReset #-}-reset = nothingIfOk =<< cudaDeviceReset-{# fun unsafe cudaDeviceReset { } -> `Status' cToEnum #}-#else-{-# INLINE cudaThreadExit #-}-reset = nothingIfOk =<< cudaThreadExit-{# fun unsafe cudaThreadExit  { } -> `Status' cToEnum #}-#endif-------------------------------------------------------------------------------------- Peer Access------------------------------------------------------------------------------------- |--- Possible option values for direct peer memory access----data PeerFlag-instance Enum PeerFlag where-#ifdef USE_EMPTY_CASE-  toEnum   x = error ("PeerFlag.toEnum: Cannot match " ++ show x)-  fromEnum x = case x of {}-#endif---- |--- Queries if the first device can directly access the memory of the second. If--- direct access is possible, it can then be enabled with 'add'. Requires--- cuda-4.0.----{-# INLINEABLE accessible #-}-accessible :: Device -> Device -> IO Bool-#if CUDART_VERSION < 4000-accessible _ _        = requireSDK 'accessible 4.0-#else-accessible !dev !peer = resultIfOk =<< cudaDeviceCanAccessPeer dev peer--{-# INLINE cudaDeviceCanAccessPeer #-}-{# fun unsafe cudaDeviceCanAccessPeer-  { alloca-  `Bool'   peekBool*-  , cIntConv `Device'-  , cIntConv `Device'           } -> `Status' cToEnum #}-#endif---- |--- If the devices of both the current and supplied contexts support unified--- addressing, then enable allocations in the supplied context to be accessible--- by the current context. Requires cuda-4.0.----{-# INLINEABLE add #-}-add :: Device -> [PeerFlag] -> IO ()-#if CUDART_VERSION < 4000-add _ _         = requireSDK 'add 4.0-#else-add !dev !flags = nothingIfOk =<< cudaDeviceEnablePeerAccess dev flags--{-# INLINE cudaDeviceEnablePeerAccess #-}-{# fun unsafe cudaDeviceEnablePeerAccess-  { cIntConv        `Device'-  , combineBitMasks `[PeerFlag]' } -> `Status' cToEnum #}-#endif----- |--- Disable direct memory access from the current context to the supplied--- context. Requires cuda-4.0.----{-# INLINEABLE remove #-}-remove :: Device -> IO ()-#if CUDART_VERSION < 4000-remove _    = requireSDK 'remove 4.0-#else-remove !dev = nothingIfOk =<< cudaDeviceDisablePeerAccess dev--{-# INLINE cudaDeviceDisablePeerAccess #-}-{# fun unsafe cudaDeviceDisablePeerAccess-  { cIntConv `Device' } -> `Status' cToEnum #}-#endif-------------------------------------------------------------------------------------- Cache Configuration------------------------------------------------------------------------------------- |--- Device limit flags----#if CUDART_VERSION < 3010-data Limit-#else-{# enum cudaLimit as Limit-    { underscoreToCase }-    with prefix="cudaLimit" deriving (Eq, Show) #}-#endif----- |--- Query compute 2.0 call stack limits. Requires cuda-3.1.----{-# INLINEABLE getLimit #-}-getLimit :: Limit -> IO Int-#if   CUDART_VERSION < 3010-getLimit _  = requireSDK 'getLimit 3.1-#elif CUDART_VERSION < 4000-getLimit !l = resultIfOk =<< cudaThreadGetLimit l--{-# INLINE cudaThreadGetLimit #-}-{# fun unsafe cudaThreadGetLimit-  { alloca-   `Int' peekIntConv*-  , cFromEnum `Limit'            } -> `Status' cToEnum #}-#else-getLimit !l = resultIfOk =<< cudaDeviceGetLimit l--{-# INLINE cudaDeviceGetLimit #-}-{# fun unsafe cudaDeviceGetLimit-  { alloca-   `Int' peekIntConv*-  , cFromEnum `Limit'            } -> `Status' cToEnum #}-#endif----- |--- Set compute 2.0 call stack limits. Requires cuda-3.1.----{-# INLINEABLE setLimit #-}-setLimit :: Limit -> Int -> IO ()-#if   CUDART_VERSION < 3010-setLimit _ _   = requireSDK 'setLimit 3.1-#elif CUDART_VERSION < 4000-setLimit !l !n = nothingIfOk =<< cudaThreadSetLimit l n--{-# INLINE cudaThreadSetLimit #-}-{# fun unsafe cudaThreadSetLimit-  { cFromEnum `Limit'-  , cIntConv  `Int'   } -> `Status' cToEnum #}-#else-setLimit !l !n = nothingIfOk =<< cudaDeviceSetLimit l n--{-# INLINE cudaDeviceSetLimit #-}-{# fun unsafe cudaDeviceSetLimit-  { cFromEnum `Limit'-  , cIntConv  `Int'   } -> `Status' cToEnum #}-#endif-

@@ -1,123 +0,0 @@-{-# LANGUAGE BangPatterns             #-}-{-# LANGUAGE DeriveDataTypeable       #-}-{-# LANGUAGE ForeignFunctionInterface #-}------------------------------------------------------------------------------------ |--- Module    : Foreign.CUDA.Runtime.Error--- Copyright : [2009..2023] Trevor L. McDonell--- License   : BSD------ Error handling functions--------------------------------------------------------------------------------------module Foreign.CUDA.Runtime.Error (--  Status(..), CUDAException(..),--  cudaError, describe, requireSDK,-  resultIfOk, nothingIfOk, checkStatus,--) where---- Friends-import Foreign.CUDA.Internal.C2HS-import Text.Show.Describe---- System-import Control.Exception-import Data.Typeable-import Foreign.C-import Foreign.Ptr-import Language.Haskell.TH-import System.IO.Unsafe-import Text.Printf--#include "cbits/stubs.h"-{# context lib="cudart" #}-------------------------------------------------------------------------------------- Return Status------------------------------------------------------------------------------------- |--- Return codes from API functions----{# enum cudaError as Status-    { cudaSuccess as Success }-    with prefix="cudaError" deriving (Eq, Show) #}------------------------------------------------------------------------------------- Exceptions-----------------------------------------------------------------------------------data CUDAException-  = ExitCode  Status-  | UserError String-  deriving Typeable--instance Exception CUDAException--instance Show CUDAException where-  showsPrec _ (ExitCode  s) = showString ("CUDA Exception: " ++ describe s)-  showsPrec _ (UserError s) = showString ("CUDA Exception: " ++ s)----- |--- Raise a 'CUDAException' in the IO Monad----cudaError :: String -> IO a-cudaError s = throwIO (UserError s)---- |--- A specially formatted error message----requireSDK :: Name -> Double -> IO a-requireSDK n v = cudaError $ printf "'%s' requires at least cuda-%3.1f\n" (show n) v-------------------------------------------------------------------------------------- Helper Functions------------------------------------------------------------------------------------- |--- Return the descriptive string associated with a particular error code----instance Describe Status where-    describe = cudaGetErrorString---- Logically, this must be a pure function, returning a pointer to a statically--- defined string constant.----{# fun pure unsafe cudaGetErrorString-    { cFromEnum `Status' } -> `String' #}----- |--- Return the results of a function on successful execution, otherwise return--- the error string associated with the return code----{-# INLINE resultIfOk #-}-resultIfOk :: (Status, a) -> IO a-resultIfOk (status, !result) =-    case status of-        Success -> return  result-        _       -> throwIO (ExitCode status)----- |--- Return the error string associated with an unsuccessful return code,--- otherwise Nothing----{-# INLINE nothingIfOk #-}-nothingIfOk :: Status -> IO ()-nothingIfOk status =-    case status of-        Success -> return  ()-        _       -> throwIO (ExitCode status)--{-# INLINE checkStatus #-}-checkStatus :: CInt -> IO ()-checkStatus = nothingIfOk . cToEnum-

@@ -1,149 +0,0 @@-{-# LANGUAGE BangPatterns             #-}-{-# LANGUAGE CPP                      #-}-{-# LANGUAGE EmptyDataDecls           #-}-{-# LANGUAGE ForeignFunctionInterface #-}-{-# LANGUAGE TemplateHaskell          #-}------------------------------------------------------------------------------------ |--- Module    : Foreign.CUDA.Driver.Event--- Copyright : [2009..2023] Trevor L. McDonell--- License   : BSD------ Event management for C-for-CUDA runtime environment--------------------------------------------------------------------------------------module Foreign.CUDA.Runtime.Event (--  -- * Event Management-  Event, EventFlag(..), WaitFlag,-  create, destroy, elapsedTime, query, record, wait, block--) where--#include "cbits/stubs.h"-{# context lib="cudart" #}---- Friends-import Foreign.CUDA.Driver.Event                          ( Event(..), EventFlag(..), WaitFlag )-import Foreign.CUDA.Driver.Stream                         ( Stream(..), defaultStream )-import Foreign.CUDA.Internal.C2HS-import Foreign.CUDA.Runtime.Error---- System-import Foreign-import Foreign.C-import Control.Monad                                      ( liftM )-import Control.Exception                                  ( throwIO )-import Data.Maybe                                         ( fromMaybe )-------------------------------------------------------------------------------------- Event management------------------------------------------------------------------------------------- |--- Create a new event----{-# INLINEABLE create #-}-create :: [EventFlag] -> IO Event-create !flags = resultIfOk =<< cudaEventCreateWithFlags flags--{-# INLINE cudaEventCreateWithFlags #-}-{# fun unsafe cudaEventCreateWithFlags-  { alloca-         `Event'       peekEvt*-  , combineBitMasks `[EventFlag]'          } -> `Status' cToEnum #}-  where peekEvt = liftM Event . peek----- |--- Destroy an event----{-# INLINEABLE destroy #-}-destroy :: Event -> IO ()-destroy !ev = nothingIfOk =<< cudaEventDestroy ev--{-# INLINE cudaEventDestroy #-}-{# fun unsafe cudaEventDestroy-  { useEvent `Event' } -> `Status' cToEnum #}----- |--- Determine the elapsed time (in milliseconds) between two events----{-# INLINEABLE elapsedTime #-}-elapsedTime :: Event -> Event -> IO Float-elapsedTime !ev1 !ev2 = resultIfOk =<< cudaEventElapsedTime ev1 ev2--{-# INLINE cudaEventElapsedTime #-}-{# fun unsafe cudaEventElapsedTime-  { alloca-  `Float' peekFloatConv*-  , useEvent `Event'-  , useEvent `Event'                } -> `Status' cToEnum #}----- |--- Determines if a event has actually been recorded----{-# INLINEABLE query #-}-query :: Event -> IO Bool-query !ev =-  cudaEventQuery ev >>= \rv ->-  case rv of-    Success  -> return True-    NotReady -> return False-    _        -> throwIO (ExitCode rv)--{-# INLINE cudaEventQuery #-}-{# fun unsafe cudaEventQuery-  { useEvent `Event' } -> `Status' cToEnum #}----- |--- Record an event once all operations in the current context (or optionally--- specified stream) have completed. This operation is asynchronous.----{-# INLINEABLE record #-}-record :: Event -> Maybe Stream -> IO ()-record !ev !mst =-  nothingIfOk =<< cudaEventRecord ev (maybe defaultStream id mst)--{-# INLINE cudaEventRecord #-}-{# fun unsafe cudaEventRecord-  { useEvent  `Event'-  , useStream `Stream' } -> `Status' cToEnum #}----- |--- Makes all future work submitted to the (optional) stream wait until the given--- event reports completion before beginning execution. Synchronisation is--- performed on the device, including when the event and stream are from--- different device contexts. Requires cuda-3.2.----{-# INLINEABLE wait #-}-wait :: Event -> Maybe Stream -> [WaitFlag] -> IO ()-#if CUDART_VERSION < 3020-wait _ _ _           = requireSDK 'wait 3.2-#else-wait !ev !mst !flags =-  let st = fromMaybe defaultStream mst-  in  nothingIfOk =<< cudaStreamWaitEvent st ev flags--{-# INLINE cudaStreamWaitEvent #-}-{# fun unsafe cudaStreamWaitEvent-  { useStream       `Stream'-  , useEvent        `Event'-  , combineBitMasks `[WaitFlag]' } -> `Status' cToEnum #}-#endif---- |--- Wait until the event has been recorded----{-# INLINEABLE block #-}-block :: Event -> IO ()-block !ev = nothingIfOk =<< cudaEventSynchronize ev--{-# INLINE cudaEventSynchronize #-}-{# fun cudaEventSynchronize-  { useEvent `Event' } -> `Status' cToEnum #}-

@@ -1,294 +0,0 @@-{-# LANGUAGE BangPatterns             #-}-{-# LANGUAGE CPP                      #-}-{-# LANGUAGE ForeignFunctionInterface #-}-{-# LANGUAGE GADTs                    #-}-{-# LANGUAGE TemplateHaskell          #-}------------------------------------------------------------------------------------ |--- Module    : Foreign.CUDA.Runtime.Exec--- Copyright : [2009..2023] Trevor L. McDonell--- License   : BSD------ Kernel execution control for C-for-CUDA runtime interface--------------------------------------------------------------------------------------module Foreign.CUDA.Runtime.Exec (--  -- * Kernel Execution-  Fun, FunAttributes(..), FunParam(..), CacheConfig(..),-  attributes, setCacheConfig, launchKernel,--) where--#include "cbits/stubs.h"-{# context lib="cudart" #}---- Friends-import Foreign.CUDA.Runtime.Stream                      ( Stream(..), defaultStream )-import Foreign.CUDA.Runtime.Error-import Foreign.CUDA.Internal.C2HS---- System-import Foreign-import Foreign.C-import Control.Monad-import Data.Maybe--#c-typedef struct cudaFuncAttributes cudaFuncAttributes;-#endc-------------------------------------------------------------------------------------- Data Types------------------------------------------------------------------------------------- |--- A @__global__@ device function.------ Note that the use of a string naming a function was deprecated in CUDA 4.1--- and removed in CUDA 5.0.----#if CUDART_VERSION >= 5000-type Fun = FunPtr ()-#else-type Fun = String-#endif-------- Function Attributes----{# pointer *cudaFuncAttributes as ^ foreign -> FunAttributes nocode #}--data FunAttributes = FunAttributes-  {-    constSizeBytes           :: !Int64,-    localSizeBytes           :: !Int64,-    sharedSizeBytes          :: !Int64,-    maxKernelThreadsPerBlock :: !Int,   -- ^ maximum block size that can be successively launched (based on register usage)-    numRegs                  :: !Int    -- ^ number of registers required for each thread-  }-  deriving (Show)--instance Storable FunAttributes where-  sizeOf _    = {# sizeof cudaFuncAttributes #}-  alignment _ = alignment (undefined :: Ptr ())--  poke _ _    = error "Can not poke Foreign.CUDA.Runtime.FunAttributes"-  peek p      = do-    cs <- cIntConv `fmap` {#get cudaFuncAttributes.constSizeBytes#} p-    ls <- cIntConv `fmap` {#get cudaFuncAttributes.localSizeBytes#} p-    ss <- cIntConv `fmap` {#get cudaFuncAttributes.sharedSizeBytes#} p-    tb <- cIntConv `fmap` {#get cudaFuncAttributes.maxThreadsPerBlock#} p-    nr <- cIntConv `fmap` {#get cudaFuncAttributes.numRegs#} p--    return FunAttributes-      {-        constSizeBytes           = cs,-        localSizeBytes           = ls,-        sharedSizeBytes          = ss,-        maxKernelThreadsPerBlock = tb,-        numRegs                  = nr-      }--#if CUDART_VERSION < 3000-data CacheConfig-#else--- |--- Cache configuration preference----{# enum cudaFuncCache as CacheConfig-    { }-    with prefix="cudaFuncCachePrefer" deriving (Eq, Show) #}-#endif---- |--- Kernel function parameters. Doubles will be converted to an internal float--- representation on devices that do not support doubles natively.----data FunParam where-  IArg :: !Int             -> FunParam-  FArg :: !Float           -> FunParam-  DArg :: !Double          -> FunParam-  VArg :: Storable a => !a -> FunParam-------------------------------------------------------------------------------------- Execution Control------------------------------------------------------------------------------------- |--- Obtain the attributes of the named @__global__@ device function. This--- itemises the requirements to successfully launch the given kernel.----{-# INLINEABLE attributes #-}-attributes :: Fun -> IO FunAttributes-attributes !fn = resultIfOk =<< cudaFuncGetAttributes fn--{-# INLINE cudaFuncGetAttributes #-}-{# fun unsafe cudaFuncGetAttributes-  { alloca-  `FunAttributes' peek*-  , withFun* `Fun'                 } -> `Status' cToEnum #}----- |--- 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 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 setConfig #-}-{# fun unsafe cudaConfigureCall_simple as setConfig-  {           `Int', `Int'-  ,           `Int', `Int', `Int'-  , cIntConv  `Int64'-  , useStream `Stream'-  }-  -> `()' checkStatus*- #}---- |--- Set the argument parameters that will be passed to the next kernel--- invocation. This is used in conjunction with 'setConfig' to control kernel--- execution.----{-# INLINEABLE setParams #-}-setParams :: [FunParam] -> IO ()-setParams = foldM_ k 0-  where-    k !offset !arg = do-      let s = size arg-      set arg s offset-      return (offset + s)--    size (IArg _) = sizeOf (undefined :: Int)-    size (FArg _) = sizeOf (undefined :: Float)-    size (DArg _) = sizeOf (undefined :: Double)-    size (VArg a) = sizeOf a--    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 = do-      d <- cudaSetDoubleForDevice v-      cudaSetupArgument d s o--{-# INLINE cudaSetupArgument #-}-{# fun unsafe cudaSetupArgument-  `Storable a' =>-  { with'* `a'-  ,        `Int'-  ,        `Int'-  }-  -> `()' checkStatus*- #}-  where-    with' v a = with v $ \p -> a (castPtr p)--{-# INLINE cudaSetDoubleForDevice #-}-{# fun unsafe cudaSetDoubleForDevice-  { with'* `Double' peek'* } -> `()' checkStatus*- #}-  where-    with' v a = with v $ \p -> a (castPtr p)-    peek'     = peek . castPtr-#endif------------------------------------------------------------------------------------- Internals------------------------------------------------------------------------------------- CUDA 5.0 changed the type of a kernel function from char* to void*----#if CUDART_VERSION >= 5000-withFun :: Fun -> (Ptr a -> IO b) -> IO b-withFun fn action = action (castFunPtrToPtr fn)-#else-withFun :: Fun -> (Ptr CChar -> IO a) -> IO a-withFun           = withCString-#endif-

@@ -1,648 +0,0 @@-{-# LANGUAGE BangPatterns             #-}-{-# LANGUAGE EmptyDataDecls           #-}-{-# LANGUAGE ForeignFunctionInterface #-}-{-# LANGUAGE TemplateHaskell          #-}------------------------------------------------------------------------------------ |--- Module    : Foreign.CUDA.Runtime.Marshal--- Copyright : [2009..2023] Trevor L. McDonell--- License   : BSD------ Memory management for CUDA devices--------------------------------------------------------------------------------------module Foreign.CUDA.Runtime.Marshal (--  -- * Host Allocation-  AllocFlag(..),-  mallocHostArray, freeHost,--  -- * Device Allocation-  mallocArray, allocaArray, free,--  -- * Unified Memory Allocation-  AttachFlag(..),-  mallocManagedArray,--  -- * Marshalling-  peekArray, peekArrayAsync, peekArray2D, peekArray2DAsync, peekListArray,-  pokeArray, pokeArrayAsync, pokeArray2D, pokeArray2DAsync, pokeListArray,-  copyArray, copyArrayAsync, copyArray2D, copyArray2DAsync,--  -- * Combined Allocation and Marshalling-  newListArray,  newListArrayLen,-  withListArray, withListArrayLen,--  -- * Utility-  memset--) where--#include "cbits/stubs.h"-{# context lib="cudart" #}---- Friends-import Foreign.CUDA.Ptr-import Foreign.CUDA.Runtime.Error-import Foreign.CUDA.Runtime.Stream-import Foreign.CUDA.Internal.C2HS---- System-import Data.Int-import Data.Maybe-import Control.Exception--import Foreign.C-import Foreign.Ptr-import Foreign.Storable-import qualified Foreign.Marshal as F--#c-typedef enum cudaMemHostAlloc_option_enum {-//  CUDA_MEMHOSTALLOC_OPTION_DEFAULT        = cudaHostAllocDefault,-    CUDA_MEMHOSTALLOC_OPTION_DEVICE_MAPPED  = cudaHostAllocMapped,-    CUDA_MEMHOSTALLOC_OPTION_PORTABLE       = cudaHostAllocPortable,-    CUDA_MEMHOSTALLOC_OPTION_WRITE_COMBINED = cudaHostAllocWriteCombined-} cudaMemHostAlloc_option;-#endc--#if CUDART_VERSION >= 6000-#c-typedef enum cudaMemAttachFlags_option_enum {-    CUDA_MEM_ATTACH_OPTION_GLOBAL = cudaMemAttachGlobal,-    CUDA_MEM_ATTACH_OPTION_HOST   = cudaMemAttachHost,-    CUDA_MEM_ATTACH_OPTION_SINGLE = cudaMemAttachSingle-} cudaMemAttachFlags_option;-#endc-#endif-------------------------------------------------------------------------------------- Host Allocation------------------------------------------------------------------------------------- |--- Options for host allocation----{# enum cudaMemHostAlloc_option as AllocFlag-    { underscoreToCase }-    with prefix="CUDA_MEMHOSTALLOC_OPTION" deriving (Eq, Show, Bounded) #}----- |--- Allocate a section of linear memory on the host which is page-locked and--- directly accessible from the device. The storage is sufficient to hold the--- given number of elements of a storable type. The runtime system automatically--- accelerates calls to functions such as 'peekArrayAsync' and 'pokeArrayAsync'--- that refer to page-locked memory.------ Note that since the amount of pageable memory is thusly reduced, overall--- system performance may suffer. This is best used sparingly to allocate--- staging areas for data exchange----{-# INLINEABLE mallocHostArray #-}-mallocHostArray :: Storable a => [AllocFlag] -> Int -> IO (HostPtr a)-mallocHostArray !flags = doMalloc undefined-  where-    doMalloc :: Storable a' => a' -> Int -> IO (HostPtr a')-    doMalloc x !n = resultIfOk =<< cudaHostAlloc (fromIntegral n * fromIntegral (sizeOf x)) flags--{-# INLINE cudaHostAlloc #-}-{# fun unsafe cudaHostAlloc-  { alloca'-        `HostPtr a' hptr*-  , cIntConv        `Int64'-  , combineBitMasks `[AllocFlag]'     } -> `Status' cToEnum #}-  where-    alloca' !f = F.alloca $ \ !p -> poke p nullPtr >> f (castPtr p)-    hptr !p    = (HostPtr . castPtr) `fmap` peek p----- |--- Free page-locked host memory previously allocated with 'mallecHost'----{-# INLINEABLE freeHost #-}-freeHost :: HostPtr a -> IO ()-freeHost !p = nothingIfOk =<< cudaFreeHost p--{-# INLINE cudaFreeHost #-}-{# fun unsafe cudaFreeHost-  { hptr `HostPtr a' } -> `Status' cToEnum #}-  where hptr = castPtr . useHostPtr-------------------------------------------------------------------------------------- Device Allocation------------------------------------------------------------------------------------- |--- Allocate a section of linear memory on the device, and return a reference to--- it. The memory is sufficient to hold the given number of elements of storable--- type. It is suitable aligned, and not cleared.----{-# INLINEABLE mallocArray #-}-mallocArray :: Storable a => Int -> IO (DevicePtr a)-mallocArray = doMalloc undefined-  where-    doMalloc :: Storable a' => a' -> Int -> IO (DevicePtr a')-    doMalloc x !n = resultIfOk =<< cudaMalloc (fromIntegral n * fromIntegral (sizeOf x))--{-# INLINE cudaMalloc #-}-{# fun unsafe cudaMalloc-  { alloca'- `DevicePtr a' dptr*-  , cIntConv `Int64'             } -> `Status' cToEnum #}-  where-    -- C-> Haskell doesn't like qualified imports in marshaller specifications-    alloca' !f = F.alloca $ \ !p -> poke p nullPtr >> f (castPtr p)-    dptr !p    = (castDevPtr . DevicePtr) `fmap` peek p----- |--- Execute a computation, passing a pointer to a temporarily allocated block of--- memory sufficient to hold the given number of elements of storable type. The--- memory is freed when the computation terminates (normally or via an--- exception), so the pointer must not be used after this.------ Note that kernel launches can be asynchronous, so you may need to add a--- synchronisation point at the end of the computation.----{-# INLINEABLE allocaArray #-}-allocaArray :: Storable a => Int -> (DevicePtr a -> IO b) -> IO b-allocaArray n = bracket (mallocArray n) free----- |--- Free previously allocated memory on the device----{-# INLINEABLE free #-}-free :: DevicePtr a -> IO ()-free !p = nothingIfOk =<< cudaFree p--{-# INLINE cudaFree #-}-{# fun unsafe cudaFree-  { dptr `DevicePtr a' } -> `Status' cToEnum #}-  where-    dptr = useDevicePtr . castDevPtr-------------------------------------------------------------------------------------- Unified memory allocation------------------------------------------------------------------------------------- |--- Options for unified memory allocations----#if CUDART_VERSION >= 6000-{# enum cudaMemAttachFlags_option as AttachFlag-    { underscoreToCase }-    with prefix="CUDA_MEM_ATTACH_OPTION" deriving (Eq, Show, Bounded) #}-#else-data AttachFlag-#endif---- |--- Allocates memory that will be automatically managed by the Unified Memory--- system----{-# INLINEABLE mallocManagedArray #-}-mallocManagedArray :: Storable a => [AttachFlag] -> Int -> IO (DevicePtr a)-#if CUDART_VERSION < 6000-mallocManagedArray _ _    = requireSDK 'mallocManagedArray 6.0-#else-mallocManagedArray !flags = doMalloc undefined-  where-    doMalloc :: Storable a' => a' -> Int -> IO (DevicePtr a')-    doMalloc x !n = resultIfOk =<< cudaMallocManaged (fromIntegral n * fromIntegral (sizeOf x)) flags--{-# INLINE cudaMallocManaged #-}-{# fun unsafe cudaMallocManaged-  { alloca'-        `DevicePtr a' dptr*-  , cIntConv        `Int64'-  , combineBitMasks `[AttachFlag]'      } -> `Status' cToEnum #}-  where-    alloca' !f = F.alloca $ \ !p -> poke p nullPtr >> f (castPtr p)-    dptr !p    = (castDevPtr . DevicePtr) `fmap` peek p-#endif-------------------------------------------------------------------------------------- Marshalling------------------------------------------------------------------------------------- |--- Copy a number of elements from the device to host memory. This is a--- synchronous operation.----{-# INLINEABLE peekArray #-}-peekArray :: Storable a => Int -> DevicePtr a -> Ptr a -> IO ()-peekArray !n !dptr !hptr = memcpy hptr (useDevicePtr dptr) n DeviceToHost----- |--- Copy memory from the device asynchronously, possibly associated with a--- particular stream. The destination memory must be page locked.----{-# INLINEABLE peekArrayAsync #-}-peekArrayAsync :: Storable a => Int -> DevicePtr a -> HostPtr a -> Maybe Stream -> IO ()-peekArrayAsync !n !dptr !hptr !mst =-  memcpyAsync (useHostPtr hptr) (useDevicePtr dptr) n DeviceToHost mst----- |--- Copy a 2D memory area from the device to the host. This is a synchronous--- operation.----{-# INLINEABLE peekArray2D #-}-peekArray2D-    :: Storable a-    => Int                      -- ^ width to copy (elements)-    -> Int                      -- ^ height to copy (elements)-    -> DevicePtr a              -- ^ source array-    -> Int                      -- ^ source array width-    -> Ptr a                    -- ^ destination array-    -> Int                      -- ^ destination array width-    -> IO ()-peekArray2D !w !h !dptr !dw !hptr !hw =-  memcpy2D hptr hw (useDevicePtr dptr) dw w h DeviceToHost----- |--- Copy a 2D memory area from the device to the host asynchronously, possibly--- associated with a particular stream. The destination array must be page--- locked.----{-# INLINEABLE peekArray2DAsync #-}-peekArray2DAsync-    :: Storable a-    => Int                      -- ^ width to copy (elements)-    -> Int                      -- ^ height to copy (elements)-    -> DevicePtr a              -- ^ source array-    -> Int                      -- ^ source array width-    -> HostPtr a                -- ^ destination array-    -> Int                      -- ^ destination array width-    -> Maybe Stream-    -> IO ()-peekArray2DAsync !w !h !dptr !dw !hptr !hw !mst =-  memcpy2DAsync (useHostPtr hptr) hw (useDevicePtr dptr) dw w h DeviceToHost mst----- |--- Copy a number of elements from the device into a new Haskell list. Note that--- this requires two memory copies: firstly from the device into a heap--- allocated array, and from there marshalled into a list----{-# INLINEABLE peekListArray #-}-peekListArray :: Storable a => Int -> DevicePtr a -> IO [a]-peekListArray !n !dptr =-  F.allocaArray n $ \p -> do-    peekArray   n dptr p-    F.peekArray n p----- |--- Copy a number of elements onto the device. This is a synchronous operation.----{-# INLINEABLE pokeArray #-}-pokeArray :: Storable a => Int -> Ptr a -> DevicePtr a -> IO ()-pokeArray !n !hptr !dptr = memcpy (useDevicePtr dptr) hptr n HostToDevice----- |--- Copy memory onto the device asynchronously, possibly associated with a--- particular stream. The source memory must be page-locked.----{-# INLINEABLE pokeArrayAsync #-}-pokeArrayAsync :: Storable a => Int -> HostPtr a -> DevicePtr a -> Maybe Stream -> IO ()-pokeArrayAsync !n !hptr !dptr !mst =-  memcpyAsync (useDevicePtr dptr) (useHostPtr hptr) n HostToDevice mst----- |--- Copy a 2D memory area onto the device. This is a synchronous operation.----{-# INLINEABLE pokeArray2D #-}-pokeArray2D-    :: Storable a-    => Int                      -- ^ width to copy (elements)-    -> Int                      -- ^ height to copy (elements)-    -> Ptr a                    -- ^ source array-    -> Int                      -- ^ source array width-    -> DevicePtr a              -- ^ destination array-    -> Int                      -- ^ destination array width-    -> IO ()-pokeArray2D !w !h !hptr !dw !dptr !hw =-  memcpy2D (useDevicePtr dptr) dw hptr hw w h HostToDevice----- |--- Copy a 2D memory area onto the device asynchronously, possibly associated--- with a particular stream. The source array must be page locked.----{-# INLINEABLE pokeArray2DAsync #-}-pokeArray2DAsync-    :: Storable a-    => Int                      -- ^ width to copy (elements)-    -> Int                      -- ^ height to copy (elements)-    -> HostPtr a                -- ^ source array-    -> Int                      -- ^ source array width-    -> DevicePtr a              -- ^ destination array-    -> Int                      -- ^ destination array width-    -> Maybe Stream-    -> IO ()-pokeArray2DAsync !w !h !hptr !hw !dptr !dw !mst =-  memcpy2DAsync (useDevicePtr dptr) dw (useHostPtr hptr) hw w h HostToDevice mst---- |--- Write a list of storable elements into a device array. The array must be--- sufficiently large to hold the entire list. This requires two marshalling--- operations----{-# INLINEABLE pokeListArray #-}-pokeListArray :: Storable a => [a] -> DevicePtr a -> IO ()-pokeListArray !xs !dptr = F.withArrayLen xs $ \len p -> pokeArray len p dptr----- |--- Copy the given number of elements from the first device array (source) to the--- second (destination). The copied areas may not overlap. This operation is--- asynchronous with respect to host, but will not overlap other device--- operations.----{-# INLINEABLE copyArray #-}-copyArray :: Storable a => Int -> DevicePtr a -> DevicePtr a -> IO ()-copyArray !n !src !dst = memcpy (useDevicePtr dst) (useDevicePtr src) n DeviceToDevice----- |--- Copy the given number of elements from the first device array (source) to the--- second (destination). The copied areas may not overlap. This operation is--- asynchronous with respect to the host, and may be associated with a--- particular stream.----{-# INLINEABLE copyArrayAsync #-}-copyArrayAsync :: Storable a => Int -> DevicePtr a -> DevicePtr a -> Maybe Stream -> IO ()-copyArrayAsync !n !src !dst !mst =-  memcpyAsync (useDevicePtr dst) (useDevicePtr src) n DeviceToDevice mst----- |--- Copy a 2D memory area from the first device array (source) to the second--- (destination). The copied areas may not overlap. This operation is--- asynchronous with respect to the host, but will not overlap other device--- operations.----{-# INLINEABLE copyArray2D #-}-copyArray2D-    :: Storable a-    => Int                      -- ^ width to copy (elements)-    -> Int                      -- ^ height to copy (elements)-    -> DevicePtr a              -- ^ source array-    -> Int                      -- ^ source array width-    -> DevicePtr a              -- ^ destination array-    -> Int                      -- ^ destination array width-    -> IO ()-copyArray2D !w !h !src !sw !dst !dw =-  memcpy2D (useDevicePtr dst) dw (useDevicePtr src) sw w h DeviceToDevice----- |--- Copy a 2D memory area from the first device array (source) to the second--- device array (destination). The copied areas may not overlay. This operation--- is asynchronous with respect to the host, and may be associated with a--- particular stream.----{-# INLINEABLE copyArray2DAsync #-}-copyArray2DAsync-    :: Storable a-    => Int                      -- ^ width to copy (elements)-    -> Int                      -- ^ height to copy (elements)-    -> DevicePtr a              -- ^ source array-    -> Int                      -- ^ source array width-    -> DevicePtr a              -- ^ destination array-    -> Int                      -- ^ destination array width-    -> Maybe Stream-    -> IO ()-copyArray2DAsync !w !h !src !sw !dst !dw !mst =-  memcpy2DAsync (useDevicePtr dst) dw (useDevicePtr src) sw w h DeviceToDevice mst-------- Memory copy kind----{# enum cudaMemcpyKind as CopyDirection {}-    with prefix="cudaMemcpy" deriving (Eq, Show) #}---- |--- Copy data between host and device. This is a synchronous operation.----{-# INLINEABLE memcpy #-}-memcpy :: Storable a-       => Ptr a                 -- ^ destination-       -> Ptr a                 -- ^ source-       -> Int                   -- ^ number of elements-       -> CopyDirection-       -> IO ()-memcpy !dst !src !n !dir = doMemcpy undefined dst-  where-    doMemcpy :: Storable a' => a' -> Ptr a' -> IO ()-    doMemcpy x _ =-      nothingIfOk =<< cudaMemcpy dst src (fromIntegral n * fromIntegral (sizeOf x)) dir--{-# INLINE cudaMemcpy #-}-{# fun cudaMemcpy-  { castPtr   `Ptr a'-  , castPtr   `Ptr a'-  , cIntConv  `Int64'-  , cFromEnum `CopyDirection' } -> `Status' cToEnum #}----- |--- Copy data between the host and device asynchronously, possibly associated--- with a particular stream. The host-side memory must be page-locked (allocated--- with 'mallocHostArray').----{-# INLINEABLE memcpyAsync #-}-memcpyAsync :: Storable a-            => Ptr a            -- ^ destination-            -> Ptr a            -- ^ source-            -> Int              -- ^ number of elements-            -> CopyDirection-            -> Maybe Stream-            -> IO ()-memcpyAsync !dst !src !n !kind !mst = doMemcpy undefined dst-  where-    doMemcpy :: Storable a' => a' -> Ptr a' -> IO ()-    doMemcpy x _ =-      let bytes = fromIntegral n * fromIntegral (sizeOf x) in-      nothingIfOk =<< cudaMemcpyAsync dst src bytes kind (fromMaybe defaultStream mst)--{-# INLINE cudaMemcpyAsync #-}-{# fun cudaMemcpyAsync-  { castPtr   `Ptr a'-  , castPtr   `Ptr a'-  , cIntConv  `Int64'-  , cFromEnum `CopyDirection'-  , useStream `Stream'        } -> `Status' cToEnum #}----- |--- Copy a 2D memory area between the host and device. This is a synchronous--- operation.----{-# INLINEABLE memcpy2D #-}-memcpy2D :: Storable a-         => Ptr a               -- ^ destination-         -> Int                 -- ^ width of destination array-         -> Ptr a               -- ^ source-         -> Int                 -- ^ width of source array-         -> Int                 -- ^ width to copy-         -> Int                 -- ^ height to copy-         -> CopyDirection-         -> IO ()-memcpy2D !dst !dw !src !sw !w !h !kind = doCopy undefined dst-  where-    doCopy :: Storable a' => a' -> Ptr a' -> IO ()-    doCopy x _ =-      let bytes = fromIntegral (sizeOf x)-          dw'   = fromIntegral dw * bytes-          sw'   = fromIntegral sw * bytes-          w'    = fromIntegral w  * bytes-          h'    = fromIntegral h-      in-      nothingIfOk =<< cudaMemcpy2D dst dw' src sw' w' h' kind--{-# INLINE cudaMemcpy2D #-}-{# fun cudaMemcpy2D-  { castPtr   `Ptr a'-  ,           `Int64'-  , castPtr   `Ptr a'-  ,           `Int64'-  ,           `Int64'-  ,           `Int64'-  , cFromEnum `CopyDirection'-  }-  -> `Status' cToEnum #}----- |--- Copy a 2D memory area between the host and device asynchronously, possibly--- associated with a particular stream. The host-side memory must be--- page-locked.----{-# INLINEABLE memcpy2DAsync #-}-memcpy2DAsync :: Storable a-              => Ptr a          -- ^ destination-              -> Int            -- ^ width of destination array-              -> Ptr a          -- ^ source-              -> Int            -- ^ width of source array-              -> Int            -- ^ width to copy-              -> Int            -- ^ height to copy-              -> CopyDirection-              -> Maybe Stream-              -> IO ()-memcpy2DAsync !dst !dw !src !sw !w !h !kind !mst = doCopy undefined dst-  where-    doCopy :: Storable a' => a' -> Ptr a' -> IO ()-    doCopy x _ =-      let bytes = fromIntegral (sizeOf x)-          dw'   = fromIntegral dw * bytes-          sw'   = fromIntegral sw * bytes-          w'    = fromIntegral w  * bytes-          h'    = fromIntegral h-          st    = fromMaybe defaultStream mst-      in-      nothingIfOk =<< cudaMemcpy2DAsync dst dw' src sw' w' h' kind st--{-# INLINE cudaMemcpy2DAsync #-}-{# fun cudaMemcpy2DAsync-  { castPtr   `Ptr a'-  ,           `Int64'-  , castPtr   `Ptr a'-  ,           `Int64'-  ,           `Int64'-  ,           `Int64'-  , cFromEnum `CopyDirection'-  , useStream `Stream'-  }-  -> `Status' cToEnum #}-------------------------------------------------------------------------------------- Combined Allocation and Marshalling------------------------------------------------------------------------------------- |--- Write a list of storable elements into a newly allocated device array,--- returning the device pointer together with the number of elements that were--- written. Note that this requires two copy operations: firstly from a Haskell--- list into a heap-allocated array, and from there into device memory. The--- array should be 'free'd when no longer required.----{-# INLINEABLE newListArrayLen #-}-newListArrayLen :: Storable a => [a] -> IO (DevicePtr a, Int)-newListArrayLen !xs =-  F.withArrayLen xs                     $ \len p ->-  bracketOnError (mallocArray len) free $ \d_xs  -> do-    pokeArray len p d_xs-    return (d_xs, len)----- |--- Write a list of storable elements into a newly allocated device array. This--- is 'newListArrayLen' composed with 'fst'.----{-# INLINEABLE newListArray #-}-newListArray :: Storable a => [a] -> IO (DevicePtr a)-newListArray !xs = fst `fmap` newListArrayLen xs----- |--- Temporarily store a list of elements into a newly allocated device array. An--- IO action is applied to the array, the result of which is returned. Similar--- to 'newListArray', this requires two marshalling operations of the data.------ As with 'allocaArray', the memory is freed once the action completes, so you--- should not return the pointer from the action, and be sure that any--- asynchronous operations (such as kernel execution) have completed.----{-# INLINEABLE withListArray #-}-withListArray :: Storable a => [a] -> (DevicePtr a -> IO b) -> IO b-withListArray !xs = withListArrayLen xs . const----- |--- A variant of 'withListArray' which also supplies the number of elements in--- the array to the applied function----{-# INLINEABLE withListArrayLen #-}-withListArrayLen :: Storable a => [a] -> (Int -> DevicePtr a -> IO b) -> IO b-withListArrayLen !xs !f =-  bracket (newListArrayLen xs) (free . fst) (uncurry . flip $ f)------ XXX: Will this attempt to double-free the device array on error (together--- with newListArrayLen)?----------------------------------------------------------------------------------------- Utility------------------------------------------------------------------------------------- |--- Initialise device memory to a given 8-bit value----{-# INLINEABLE memset #-}-memset :: DevicePtr a                   -- ^ The device memory-       -> Int64                         -- ^ Number of bytes-       -> Int8                          -- ^ Value to set for each byte-       -> IO ()-memset !dptr !bytes !symbol = nothingIfOk =<< cudaMemset dptr symbol bytes--{-# INLINE cudaMemset #-}-{# fun unsafe cudaMemset-  { dptr     `DevicePtr a'-  , cIntConv `Int8'-  , cIntConv `Int64'       } -> `Status' cToEnum #}-  where-    dptr = useDevicePtr . castDevPtr-

@@ -1,117 +0,0 @@-{-# LANGUAGE BangPatterns             #-}-{-# LANGUAGE ForeignFunctionInterface #-}------------------------------------------------------------------------------------ |--- Module    : Foreign.CUDA.Runtime.Stream--- Copyright : [2009..2023] Trevor L. McDonell--- License   : BSD------ Stream management routines--------------------------------------------------------------------------------------module Foreign.CUDA.Runtime.Stream (--  -- * Stream Management-  Stream(..),-  create, destroy, finished, block,--  defaultStream, defaultStreamLegacy, defaultStreamPerThread,--) where--#include "cbits/stubs.h"-{# context lib="cudart" #}---- Friends-import Foreign.CUDA.Driver.Stream                         ( Stream(..), defaultStream, defaultStreamLegacy, defaultStreamPerThread )-import Foreign.CUDA.Internal.C2HS-import Foreign.CUDA.Runtime.Error---- System-import Foreign-import Foreign.C-import Control.Monad                                      ( liftM )-import Control.Exception                                  ( throwIO )-------------------------------------------------------------------------------------- Functions------------------------------------------------------------------------------------- |--- Create a new asynchronous stream----{-# INLINEABLE create #-}-create :: IO Stream-create = resultIfOk =<< cudaStreamCreate--{-# INLINE cudaStreamCreate #-}-{# fun unsafe cudaStreamCreate-  { alloca- `Stream' peekStream* } -> `Status' cToEnum #}----- |--- Destroy and clean up an asynchronous stream----{-# INLINEABLE destroy #-}-destroy :: Stream -> IO ()-destroy !s = nothingIfOk =<< cudaStreamDestroy s--{-# INLINE cudaStreamDestroy #-}-{# fun unsafe cudaStreamDestroy-  { useStream `Stream' } -> `Status' cToEnum #}----- |--- Determine if all operations in a stream have completed----{-# INLINEABLE finished #-}-finished :: Stream -> IO Bool-finished !s =-  cudaStreamQuery s >>= \rv -> do-  case rv of-      Success  -> return True-      NotReady -> return False-      _        -> throwIO (ExitCode rv)--{-# INLINE cudaStreamQuery #-}-{# fun unsafe cudaStreamQuery-  { useStream `Stream' } -> `Status' cToEnum #}----- |--- Block until all operations in a Stream have been completed----{-# INLINEABLE block #-}-block :: Stream -> IO ()-block !s = nothingIfOk =<< cudaStreamSynchronize s--{-# INLINE cudaStreamSynchronize #-}-{# fun cudaStreamSynchronize-  { useStream `Stream' } -> `Status' cToEnum #}----- |--- The main execution stream (0)------ {-# INLINE defaultStream #-}--- defaultStream :: Stream--- #if CUDART_VERSION < 3010--- defaultStream = Stream 0--- #else--- defaultStream = Stream nullPtr--- #endif------------------------------------------------------------------------------------- Internal-----------------------------------------------------------------------------------{-# INLINE peekStream #-}-peekStream :: Ptr {#type cudaStream_t#} -> IO Stream-#if CUDART_VERSION < 3010-peekStream = liftM Stream . peekIntConv-#else-peekStream = liftM Stream . peek-#endif-

@@ -1,63 +0,0 @@-{-# LANGUAGE ForeignFunctionInterface #-}------------------------------------------------------------------------------------ |--- Module    : Foreign.CUDA.Runtime.Utils--- Copyright : [2009..2023] Trevor L. McDonell--- License   : BSD------ Utility functions--------------------------------------------------------------------------------------module Foreign.CUDA.Runtime.Utils (--  runtimeVersion,-  driverVersion,-  libraryVersion,--) where--#include "cbits/stubs.h"-{# context lib="cudart" #}---- Friends-import Foreign.CUDA.Runtime.Error-import Foreign.CUDA.Internal.C2HS---- System-import Foreign-import Foreign.C----- |--- Return the version number of the installed CUDA driver----{-# INLINEABLE runtimeVersion #-}-runtimeVersion :: IO Int-runtimeVersion =  resultIfOk =<< cudaRuntimeGetVersion--{-# INLINE cudaRuntimeGetVersion #-}-{# fun unsafe cudaRuntimeGetVersion-  { alloca- `Int' peekIntConv* } -> `Status' cToEnum #}----- |--- Return the version number of the installed CUDA runtime----{-# INLINEABLE driverVersion #-}-driverVersion :: IO Int-driverVersion =  resultIfOk =<< cudaDriverGetVersion--{-# INLINE cudaDriverGetVersion #-}-{# fun unsafe cudaDriverGetVersion-  { alloca- `Int' peekIntConv* } -> `Status' cToEnum #}----- |--- Return the version number of the CUDA library (API) that this package was--- compiled against.----{-# INLINEABLE libraryVersion #-}-libraryVersion :: Int-libraryVersion = {#const CUDART_VERSION #}-

@@ -1,18 +0,0 @@------------------------------------------------------------------------------------ |--- Module    : Text.Show.Describe--- Copyright : [2016..2023] Trevor L. McDonell--- License   : BSD--------------------------------------------------------------------------------------module Text.Show.Describe-  where----- | Like 'Text.Show.Show', but focuses on providing a more detailed description--- of the value rather than a 'Text.Read.read'able representation.----class Describe a where-    describe :: a -> String-