diff --git a/CHANGELOG.md b/CHANGELOG.md --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -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! diff --git a/README.md b/README.md --- a/README.md +++ b/README.md @@ -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 diff --git a/Setup.hs b/Setup.hs --- a/Setup.hs +++ b/Setup.hs @@ -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 diff --git a/cbits/stubs.c b/cbits/stubs.c --- a/cbits/stubs.c +++ b/cbits/stubs.c @@ -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 // 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(¶ms, 0, sizeof params); + return cuCtxCreate_v4(pctx, ¶ms, 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 diff --git a/cbits/stubs.h b/cbits/stubs.h --- a/cbits/stubs.h +++ b/cbits/stubs.h @@ -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 diff --git a/cuda.cabal b/cuda.cabal --- a/cuda.cabal +++ b/cuda.cabal @@ -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 - build matrix for - version compatibility. + Tested with library versions up to and including CUDA-13.0. . [/NOTES:/] . @@ -40,16 +35,18 @@ . * . - 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 Author: Trevor L. McDonell -Maintainer: Trevor L. McDonell +Maintainer: Trevor L. McDonell , + Tom Smeding 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 diff --git a/examples/src/deviceQueryDrv/DeviceQuery.hs b/examples/src/deviceQueryDrv/DeviceQuery.hs --- a/examples/src/deviceQueryDrv/DeviceQuery.hs +++ b/examples/src/deviceQueryDrv/DeviceQuery.hs @@ -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)) ] diff --git a/src/Foreign/CUDA.hs b/src/Foreign/CUDA.hs deleted file mode 100644 --- a/src/Foreign/CUDA.hs +++ /dev/null @@ -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 - diff --git a/src/Foreign/CUDA/Analysis/Device.chs b/src/Foreign/CUDA/Analysis/Device.chs --- a/src/Foreign/CUDA/Analysis/Device.chs +++ b/src/Foreign/CUDA/Analysis/Device.chs @@ -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 { diff --git a/src/Foreign/CUDA/Driver.hs b/src/Foreign/CUDA/Driver.hs --- a/src/Foreign/CUDA/Driver.hs +++ b/src/Foreign/CUDA/Driver.hs @@ -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: -- diff --git a/src/Foreign/CUDA/Driver/Device.chs b/src/Foreign/CUDA/Driver/Device.chs --- a/src/Foreign/CUDA/Driver/Device.chs +++ b/src/Foreign/CUDA/Driver/Device.chs @@ -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 -- -- @@ -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 diff --git a/src/Foreign/CUDA/Driver/Error.chs b/src/Foreign/CUDA/Driver/Error.chs --- a/src/Foreign/CUDA/Driver/Error.chs +++ b/src/Foreign/CUDA/Driver/Error.chs @@ -23,7 +23,7 @@ -- Friends import Foreign.CUDA.Internal.C2HS -import Text.Show.Describe +import Foreign.CUDA.Internal.Describe -- System import Control.Exception diff --git a/src/Foreign/CUDA/Driver/Event.chs b/src/Foreign/CUDA/Driver/Event.chs --- a/src/Foreign/CUDA/Driver/Event.chs +++ b/src/Foreign/CUDA/Driver/Event.chs @@ -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 -- | diff --git a/src/Foreign/CUDA/Driver/Graph/Build.chs b/src/Foreign/CUDA/Driver/Graph/Build.chs --- a/src/Foreign/CUDA/Driver/Graph/Build.chs +++ b/src/Foreign/CUDA/Driver/Graph/Build.chs @@ -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 diff --git a/src/Foreign/CUDA/Driver/Graph/Capture.chs b/src/Foreign/CUDA/Driver/Graph/Capture.chs --- a/src/Foreign/CUDA/Driver/Graph/Capture.chs +++ b/src/Foreign/CUDA/Driver/Graph/Capture.chs @@ -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 -- -- @@ -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 diff --git a/src/Foreign/CUDA/Driver/Marshal.chs b/src/Foreign/CUDA/Driver/Marshal.chs --- a/src/Foreign/CUDA/Driver/Marshal.chs +++ b/src/Foreign/CUDA/Driver/Marshal.chs @@ -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 diff --git a/src/Foreign/CUDA/Driver/Unified.chs b/src/Foreign/CUDA/Driver/Unified.chs --- a/src/Foreign/CUDA/Driver/Unified.chs +++ b/src/Foreign/CUDA/Driver/Unified.chs @@ -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 - diff --git a/src/Foreign/CUDA/Internal/Describe.hs b/src/Foreign/CUDA/Internal/Describe.hs new file mode 100644 --- /dev/null +++ b/src/Foreign/CUDA/Internal/Describe.hs @@ -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 + diff --git a/src/Foreign/CUDA/Runtime.hs b/src/Foreign/CUDA/Runtime.hs deleted file mode 100644 --- a/src/Foreign/CUDA/Runtime.hs +++ /dev/null @@ -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 - diff --git a/src/Foreign/CUDA/Runtime/Device.chs b/src/Foreign/CUDA/Runtime/Device.chs deleted file mode 100644 --- a/src/Foreign/CUDA/Runtime/Device.chs +++ /dev/null @@ -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 - diff --git a/src/Foreign/CUDA/Runtime/Error.chs b/src/Foreign/CUDA/Runtime/Error.chs deleted file mode 100644 --- a/src/Foreign/CUDA/Runtime/Error.chs +++ /dev/null @@ -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 - diff --git a/src/Foreign/CUDA/Runtime/Event.chs b/src/Foreign/CUDA/Runtime/Event.chs deleted file mode 100644 --- a/src/Foreign/CUDA/Runtime/Event.chs +++ /dev/null @@ -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 #} - diff --git a/src/Foreign/CUDA/Runtime/Exec.chs b/src/Foreign/CUDA/Runtime/Exec.chs deleted file mode 100644 --- a/src/Foreign/CUDA/Runtime/Exec.chs +++ /dev/null @@ -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 - diff --git a/src/Foreign/CUDA/Runtime/Marshal.chs b/src/Foreign/CUDA/Runtime/Marshal.chs deleted file mode 100644 --- a/src/Foreign/CUDA/Runtime/Marshal.chs +++ /dev/null @@ -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 - diff --git a/src/Foreign/CUDA/Runtime/Stream.chs b/src/Foreign/CUDA/Runtime/Stream.chs deleted file mode 100644 --- a/src/Foreign/CUDA/Runtime/Stream.chs +++ /dev/null @@ -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 - diff --git a/src/Foreign/CUDA/Runtime/Utils.chs b/src/Foreign/CUDA/Runtime/Utils.chs deleted file mode 100644 --- a/src/Foreign/CUDA/Runtime/Utils.chs +++ /dev/null @@ -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 #} - diff --git a/src/Text/Show/Describe.hs b/src/Text/Show/Describe.hs deleted file mode 100644 --- a/src/Text/Show/Describe.hs +++ /dev/null @@ -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 -