cuda 0.6.7.0 → 0.7.0.0
raw patch · 37 files changed
+3077/−1069 lines, 37 filesdep +template-haskelldep −Cabalsetup-changed
Dependencies added: template-haskell
Dependencies removed: Cabal
Files
- CHANGELOG.markdown +47/−0
- Foreign/CUDA/Analysis/Device.chs +42/−32
- Foreign/CUDA/Driver.hs +210/−4
- Foreign/CUDA/Driver/Context.chs +0/−362
- Foreign/CUDA/Driver/Context.hs +23/−0
- Foreign/CUDA/Driver/Context/Base.chs +241/−0
- Foreign/CUDA/Driver/Context/Config.chs +297/−0
- Foreign/CUDA/Driver/Context/Peer.chs +131/−0
- Foreign/CUDA/Driver/Context/Primary.chs +170/−0
- Foreign/CUDA/Driver/Device.chs +127/−48
- Foreign/CUDA/Driver/Error.chs +16/−6
- Foreign/CUDA/Driver/Event.chs +19/−2
- Foreign/CUDA/Driver/Exec.chs +110/−50
- Foreign/CUDA/Driver/IPC/Event.chs +136/−0
- Foreign/CUDA/Driver/IPC/Marshal.chs +168/−0
- Foreign/CUDA/Driver/Marshal.chs +112/−20
- Foreign/CUDA/Driver/Module.chs +0/−331
- Foreign/CUDA/Driver/Module.hs +20/−0
- Foreign/CUDA/Driver/Module/Base.chs +311/−0
- Foreign/CUDA/Driver/Module/Link.chs +219/−0
- Foreign/CUDA/Driver/Module/Query.chs +108/−0
- Foreign/CUDA/Driver/Stream.chs +78/−5
- Foreign/CUDA/Driver/Texture.chs +32/−8
- Foreign/CUDA/Driver/Utils.chs +1/−1
- Foreign/CUDA/Runtime/Device.chs +61/−41
- Foreign/CUDA/Runtime/Error.chs +6/−4
- Foreign/CUDA/Runtime/Event.chs +2/−1
- Foreign/CUDA/Runtime/Exec.chs +4/−2
- Foreign/CUDA/Runtime/Marshal.chs +4/−3
- Foreign/CUDA/Types.chs +10/−2
- README.markdown +20/−0
- Setup.hs +248/−90
- WINDOWS.markdown +54/−0
- cbits/stubs.c +15/−0
- cbits/stubs.h +6/−13
- changelog.md +0/−34
- cuda.cabal +29/−10
+ CHANGELOG.markdown view
@@ -0,0 +1,47 @@+0.7.0.0++ * Add support for operations from CUDA-7.0++ * Add support for online linking++ * Add support for inter-process communication++ * Bug fixes, extra documentation, improve library coverage.++ * Mac OS X no longer requires the DYLD_LIBRARY_PATH environment variable in+ order to compile or run programs that use this package.++0.6.7.0++ * Add support for building on Windows (thanks to mwu-tow)++0.6.6.2++ * Build fix++0.6.6.1++ * Build fixes for ghc-7.6 and ghc-7.10++0.6.6.0++ * Drop support for CUDA 3.0 and older.++ * Combine the definition of the 'Event' and 'Stream' data types. As of+ CUDA-3.1 these data structures are equivalent, and can be safely shared+ between runtime and driver API calls and libraries.++ * Mark FFI imports of potentially long-running API functions as safe. This+ allows them to be safely called from Haskell threads without blocking the+ entire HEC.++ * Add compute-capability data for 3.7, 5.2 devices.++0.6.5.1++ * Build fix for Mac OS X 10.10 (Yosemite)++0.6.5.0++ * Add support for the CUDA-6.5 release+
Foreign/CUDA/Analysis/Device.chs view
@@ -57,49 +57,59 @@ -- data DeviceProperties = DeviceProperties {- deviceName :: !String, -- ^ Identifier- computeCapability :: !Compute, -- ^ Supported compute capability- totalGlobalMem :: !Int64, -- ^ Available global memory on the device in bytes- totalConstMem :: !Int64, -- ^ Available constant memory on the device in bytes- sharedMemPerBlock :: !Int64, -- ^ Available shared memory per block in bytes- regsPerBlock :: !Int, -- ^ 32-bit registers per block- warpSize :: !Int, -- ^ Warp size in threads (SIMD width)- maxThreadsPerBlock :: !Int, -- ^ Max number of threads per block+ deviceName :: !String -- ^ Identifier+ , computeCapability :: !Compute -- ^ Supported compute capability+ , totalGlobalMem :: !Int64 -- ^ Available global memory on the device in bytes+ , totalConstMem :: !Int64 -- ^ Available constant memory on the device in bytes+ , sharedMemPerBlock :: !Int64 -- ^ Available shared memory per block in bytes+ , regsPerBlock :: !Int -- ^ 32-bit registers per block+ , warpSize :: !Int -- ^ Warp size in threads (SIMD width)+ , maxThreadsPerBlock :: !Int -- ^ Maximum number of threads per block #if CUDA_VERSION >= 4000- maxThreadsPerMultiProcessor :: !Int, -- ^ Max number of threads per multiprocessor+ , maxThreadsPerMultiProcessor :: !Int -- ^ Maximum number of threads per multiprocessor #endif- maxBlockSize :: !(Int,Int,Int), -- ^ Max size of each dimension of a block- maxGridSize :: !(Int,Int,Int), -- ^ Max size of each dimension of a grid+ , maxBlockSize :: !(Int,Int,Int) -- ^ Maximum size of each dimension of a block+ , maxGridSize :: !(Int,Int,Int) -- ^ Maximum size of each dimension of a grid #if CUDA_VERSION >= 3000- maxTextureDim1D :: !Int, -- ^ Maximum texture dimensions- maxTextureDim2D :: !(Int,Int),- maxTextureDim3D :: !(Int,Int,Int),+ , maxTextureDim1D :: !Int -- ^ Maximum texture dimensions+ , maxTextureDim2D :: !(Int,Int)+ , maxTextureDim3D :: !(Int,Int,Int) #endif- clockRate :: !Int, -- ^ Clock frequency in kilohertz- multiProcessorCount :: !Int, -- ^ Number of multiprocessors on the device- memPitch :: !Int64, -- ^ Max pitch in bytes allowed by memory copies+ , clockRate :: !Int -- ^ Clock frequency in kilohertz+ , multiProcessorCount :: !Int -- ^ Number of multiprocessors on the device+ , memPitch :: !Int64 -- ^ Maximum pitch in bytes allowed by memory copies #if CUDA_VERSION >= 4000- memBusWidth :: !Int, -- ^ Global memory bus width in bits- memClockRate :: !Int, -- ^ Peak memory clock frequency in kilohertz+ , memBusWidth :: !Int -- ^ Global memory bus width in bits+ , memClockRate :: !Int -- ^ Peak memory clock frequency in kilohertz #endif- textureAlignment :: !Int64, -- ^ Alignment requirement for textures- computeMode :: !ComputeMode,- deviceOverlap :: !Bool, -- ^ Device can concurrently copy memory and execute a kernel+ , textureAlignment :: !Int64 -- ^ Alignment requirement for textures+ , computeMode :: !ComputeMode+ , deviceOverlap :: !Bool -- ^ Device can concurrently copy memory and execute a kernel #if CUDA_VERSION >= 3000- concurrentKernels :: !Bool, -- ^ Device can possibly execute multiple kernels concurrently- eccEnabled :: !Bool, -- ^ Device supports and has enabled error correction+ , concurrentKernels :: !Bool -- ^ Device can possibly execute multiple kernels concurrently+ , eccEnabled :: !Bool -- ^ Device supports and has enabled error correction #endif #if CUDA_VERSION >= 4000- asyncEngineCount :: !Int, -- ^ Number of asynchronous engines- cacheMemL2 :: !Int, -- ^ Size of the L2 cache in bytes- tccDriverEnabled :: !Bool, -- ^ Whether this is a Tesla device using the TCC driver- pciInfo :: !PCI, -- ^ PCI device information for the device+ , asyncEngineCount :: !Int -- ^ Number of asynchronous engines+ , cacheMemL2 :: !Int -- ^ Size of the L2 cache in bytes+ , pciInfo :: !PCI -- ^ PCI device information for the device+ , tccDriverEnabled :: !Bool -- ^ Whether this is a Tesla device using the TCC driver #endif- kernelExecTimeoutEnabled :: !Bool, -- ^ Whether there is a runtime limit on kernels- integrated :: !Bool, -- ^ As opposed to discrete- canMapHostMemory :: !Bool, -- ^ Device can use pinned memory+ , kernelExecTimeoutEnabled :: !Bool -- ^ Whether there is a runtime limit on kernels+ , integrated :: !Bool -- ^ As opposed to discrete+ , canMapHostMemory :: !Bool -- ^ Device can use pinned memory #if CUDA_VERSION >= 4000- unifiedAddressing :: !Bool -- ^ Device shares a unified address space with the host+ , unifiedAddressing :: !Bool -- ^ Device shares a unified address space with the host+#endif+#if CUDA_VERSION >= 5050+ , streamPriorities :: !Bool -- ^ Device supports stream priorities+#endif+#if CUDA_VERSION >= 6000+ , globalL1Cache :: !Bool -- ^ Device supports caching globals in L1 cache+ , localL1Cache :: !Bool -- ^ Device supports caching locals in L1 cache+ , managedMemory :: !Bool -- ^ Device supports allocating managed memory on this system+ , multiGPUBoard :: !Bool -- ^ Device is on a multi-GPU board+ , multiGPUBoardGroupID :: !Int -- ^ Unique identifier for a group of devices associated with the same board #endif } deriving (Show)
Foreign/CUDA/Driver.hs view
@@ -1,11 +1,217 @@ -------------------------------------------------------------------------------- -- | -- Module : Foreign.CUDA.Driver--- Copyright : [2009..2014] Trevor L. McDonell+-- Copyright : [2009..2015] Trevor L. McDonell -- License : BSD ----- Top level bindings to CUDA driver API+-- This module defines an interface to the CUDA driver API. The Driver API+-- is a lower-level interface to CUDA devices than that provided by the+-- Runtime API. Using the Driver API, the programmer must deal explicitly+-- with operations such as initialisation, context management, and loading+-- (kernel) modules. Although more difficult to use initially, the Driver+-- API provides more control over how CUDA is used. Furthermore, since it+-- does not require compiling and linking the program with 'nvcc', the+-- Driver API provides better inter-language compatibility.+ --+-- The following is a short tutorial on using the Driver API. The steps can+-- be copied into a file, or run directly in `ghci`, in which case `ghci`+-- should be launched with the option `-fno-ghci-sandbox`. This is because+-- CUDA maintains CPU-local state, so operations should always be run from+-- a bound thread. --+--+-- [/Using the Driver API/]+--+-- Before any operation can be performed, the Driver API must be+-- initialised:+--+-- >>> import Foreign.CUDA.Driver+-- >>> initialise []+--+-- 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+-- 'props'. The number of available CUDA-capable devices is given via+-- 'count'. For example:+--+-- >>> count+-- 1+-- >>> dev0 <- device 0+-- >>> props dev0+-- DeviceProperties {deviceName = "GeForce GT 650M", computeCapability = 3.0, ...}+--+-- This package also includes the executable 'nvidia-device-query', which when+-- executed displays the key properties of all available devices. See+-- "Foreign.CUDA.Driver.Device" for additional operations to query the+-- capabilities or status of a device.+--+-- Once you have chosen a device to use, the next step is to create a CUDA+-- context. A context is associated with a particular device, and all+-- operations, such as memory allocation and kernel execution, take place+-- with respect to that context. For example, to 'create' a new execution+-- context on CUDA device 0:+--+-- >>> ctx <- create dev0 []+--+-- The second argument is a set of 'ContextFlag's which control how the+-- context behaves in various situations, for example, whether or not the+-- CPU should actively spin when waiting for results from the GPU+-- ('SchedSpin'), or to yield control to other threads instead+-- ('SchedYield').+--+-- The newly created context is now the /active/ context, and all+-- subsequent operations take place within that context. More than one+-- context can be created per device, but resources, such as memory+-- allocated in the GPU, are unique to each context. The module+-- "Foreign.CUDA.Driver.Context" contains operations for managing multiple+-- contexts. Some devices allow data to be shared between contexts without+-- copying, see "Foreign.CUDA.Driver.Context.Peer" for more information.+--+-- Once the context is no longer needed, it should be 'destroy'ed in order+-- to free up any resources that were allocated to it.+--+-- >>> destroy ctx+--+-- Each device also has a unique context which is used by the Runtime API.+-- This context can be accessed with the module+-- "Foreign.CUDA.Driver.Context.Primary".+--+--+-- [/Executing kernels onto the GPU/]+--+-- Once the Driver API is initialised and an execution context is created+-- on the GPU, we can begin to interact with it.+--+-- At an example, we'll step through executing the CUDA equivalent of the+-- following Haskell function, which element-wise adds the elements of two+-- arrays:+--+-- >>> vecAdd xs ys = zipWith (+) xs ys+--+-- The following CUDA kernel can be used to implement this on the GPU:+--+-- > extern "C" __global__ void vecAdd(float *xs, float *ys, float *zs, int N)+-- > {+-- > int ix = blockIdx.x * blockDim.x + threadIdx.x;+-- >+-- > if ( ix < N ) {+-- > zs[ix] = xs[ix] + ys[ix];+-- > }+-- > }+--+-- Here, the `__global__` keyword marks the function as a kernel that+-- should be computed on the GPU in data parallel. When we execute this+-- function on the GPU, (at least) /N/ threads will execute /N/ individual+-- instances of the kernel function `vecAdd`. Each thread will operate on+-- a single element of each input array to create a single value in the+-- result. See the CUDA programming guide for more details.+--+-- We can save this to a file `vector_add.cu`, and compile it using `nvcc`+-- into a form that we can then load onto the GPU and execute:+--+-- > $ nvcc --ptx vector_add.cu+--+-- The module "Foreign.CUDA.Driver.Module" contains functions for loading+-- the resulting .ptx file (or .cubin files) into the running program.+--+-- >>> mdl <- loadFile "vector_add.ptx"+--+-- Once finished with the module, it is also a good idea to 'unload' it.+--+-- Modules may export kernel functions, global variables, and texture+-- references. Before we can execute our function, we need to look it up in+-- the module by name.+--+-- >>> vecAdd <- getFun mdl "vecAdd"+--+-- Given this reference to our kernel function, we are almost ready to+-- execute it on the device using 'launchKernel', but first, we must create+-- some data that we can execute the function on.+--+--+-- [/Transferring data to and from the GPU/]+--+-- GPUs typically have their own memory which is separate from the CPU's+-- memory, and we need to explicitly copy data back and forth between these+-- two regions. The module "Foreign.CUDA.Driver.Marshal" provides functions+-- for allocating memory on the GPU, and copying data between the CPU and+-- GPU, as well as directly between multiple GPUs.+--+-- For simplicity, we'll use standard Haskell lists for our input and+-- output data structure. Note however that this will have significantly+-- lower effective bandwidth than reading a single contiguous region of+-- memory, so for most practical purposes you will want to use some kind of+-- unboxed array.+--+-- >>> let xs = [1..1024] :: [Float]+-- >>> let ys = [2,4..2048] :: [Float]+--+-- In CUDA, like C, all memory management is explicit, and arrays on the+-- device must be explicitly allocated and freed. As mentioned previously,+-- data transfer is also explicit. However, we do provide convenience+-- functions for combined allocation and marshalling, as well as bracketed+-- operations.+--+-- >>> xs_dev <- newListArray xs+-- >>> ys_dev <- newListArray ys+-- >>> zs_dev <- mallocArray 1024 :: IO (DevicePtr Float)+--+-- After executing the kernel (see next section), we transfer the result+-- back to the host, and free the memory that was allocated on the GPU.+--+-- >>> zs <- peekListArray 1024 zs_dev+-- >>> free xs_dev+-- >>> free ys_dev+-- >>> free zs_dev+--+--+-- [/Piecing it all together/]+--+-- Finally, we have everything in place to execute our operation on the+-- GPU. Launching a kernel on the GPU consists of creating many threads on+-- the GPU which all execute the same function, and each thread has+-- a unique identifier in the grid/block hierarchy which can be used to+-- identify exactly which element this thread should process (the+-- `blockIdx` and `threadIdx` parameters that we saw earlier,+-- respectively).+--+-- To execute our function, we will use a grid of 4 blocks, each containing+-- 256 threads. Thus, a total of 1024 threads will be launched, which will+-- each compute a single element of the output array (recall that our input+-- arrays each have 1024 elements). The module+-- "Foreign.CUDA.Analysis.Occupancy" contains functions to help determine+-- the ideal thread block size for a given kernel and GPU combination.+--+-- >>> launchKernel vecAdd (4,1,1) (256,1,1) 0 Nothing [VArg xs_dev, VArg ys_dev, VArg zs_dev, IArg 1024]+--+-- Note that kernel execution is asynchronous, so we should also wait for+-- the operation to complete before attempting to read the results back.+--+-- >>> sync+--+-- And that's it!+--+--+-- [/Next steps/]+--+-- As mentioned at the end of the previous section, kernels on the GPU are+-- executed asynchronously with respect to the host, and other operations+-- such as data transfers can also be executed asynchronously. This allows+-- the CPU to continue doing other work while the GPU is busy.+-- 'Foreign.CUDA.Driver.Event.Event's can be used to check whether an+-- operation has completed yet.+--+-- It is also possible to execute multiple kernels or data transfers+-- concurrently with each other, by assigning those operations to different+-- execution 'Foreign.CUDA.Driver.Stream.Stream's. Used in conjunction with+-- 'Foreign.CUDA.Driver.Event.Event's, operations will be scheduled+-- efficiently only once all dependencies (in the form of+-- 'Foreign.CUDA.Driver.Event.Event's) have been cleared.+--+-- See "Foreign.CUDA.Driver.Event" and "Foreign.CUDA.Driver.Stream" for+-- more information on this topic.+-- -------------------------------------------------------------------------------- module Foreign.CUDA.Driver (@@ -22,8 +228,8 @@ ) where import Foreign.CUDA.Ptr-import Foreign.CUDA.Driver.Context hiding ( device, useContext )-import Foreign.CUDA.Driver.Device+import Foreign.CUDA.Driver.Context hiding ( useContext, device )+import Foreign.CUDA.Driver.Device hiding ( useDevice ) import Foreign.CUDA.Driver.Error import Foreign.CUDA.Driver.Exec import Foreign.CUDA.Driver.Marshal hiding ( useDeviceHandle, peekDeviceHandle )
− Foreign/CUDA/Driver/Context.chs
@@ -1,362 +0,0 @@-{-# LANGUAGE BangPatterns #-}-{-# LANGUAGE CPP #-}-{-# LANGUAGE EmptyDataDecls #-}-{-# LANGUAGE ForeignFunctionInterface #-}-#ifdef USE_EMPTY_CASE-{-# LANGUAGE EmptyCase #-}-#endif------------------------------------------------------------------------------------ |--- Module : Foreign.CUDA.Driver.Context--- Copyright : [2009..2014] Trevor L. McDonell--- License : BSD------ Context management for low-level driver interface--------------------------------------------------------------------------------------module Foreign.CUDA.Driver.Context (-- -- * Context Management- Context(..), ContextFlag(..),- create, attach, detach, destroy, device, pop, push, sync, get, set,-- -- * Peer Access- PeerFlag,- accessible, add, remove,-- -- * Cache Configuration- Cache(..), Limit(..),- getLimit, setLimit, setCacheConfig--) where--#include "cbits/stubs.h"-{# context lib="cuda" #}---- Friends-import Foreign.CUDA.Driver.Device ( Device(..) )-import Foreign.CUDA.Driver.Error-import Foreign.CUDA.Internal.C2HS---- System-import Foreign-import Foreign.C-import Control.Monad ( liftM )-------------------------------------------------------------------------------------- Data Types------------------------------------------------------------------------------------- |--- A device context----newtype Context = Context { useContext :: {# type CUcontext #}}- deriving (Eq, Show)----- |--- Context creation flags----{# enum CUctx_flags as ContextFlag- { underscoreToCase }- with prefix="CU_CTX" deriving (Eq, Show) #}---- |--- Device limits flags----#if CUDA_VERSION < 3010-data Limit-#else-{# enum CUlimit_enum as Limit- { underscoreToCase }- with prefix="CU_LIMIT" deriving (Eq, Show) #}-#endif---- |--- Device cache configuration preference----#if CUDA_VERSION < 3000-data Cache-#else-{# enum CUfunc_cache_enum as Cache- { underscoreToCase }- with prefix="CU_FUNC_CACHE" deriving (Eq, Show) #}-#endif---- |--- Possible option values for direct peer memory access----data PeerFlag-instance Enum PeerFlag where-#ifdef USE_EMPTY_CASE- toEnum x = case x of {}- fromEnum x = case x of {}-#endif---#if CUDA_VERSION >= 4000-{-# DEPRECATED attach, detach "deprecated as of CUDA-4.0" #-}-{-# DEPRECATED BlockingSync "use SchedBlockingSync instead" #-}-#endif-------------------------------------------------------------------------------------- Context management------------------------------------------------------------------------------------- |--- Create a new CUDA context and associate it with the calling thread----{-# INLINEABLE create #-}-create :: Device -> [ContextFlag] -> IO Context-create !dev !flags = resultIfOk =<< cuCtxCreate flags dev--{-# INLINE cuCtxCreate #-}-{# fun unsafe cuCtxCreate- { alloca- `Context' peekCtx*- , combineBitMasks `[ContextFlag]'- , useDevice `Device' } -> `Status' cToEnum #}- where peekCtx = liftM Context . peek----- |--- Increments the usage count of the context. API: no context flags are--- currently supported, so this parameter must be empty.----{-# INLINEABLE attach #-}-attach :: Context -> [ContextFlag] -> IO ()-attach !ctx !flags = nothingIfOk =<< cuCtxAttach ctx flags--{-# INLINE cuCtxAttach #-}-{# fun unsafe cuCtxAttach- { withCtx* `Context'- , combineBitMasks `[ContextFlag]' } -> `Status' cToEnum #}- where withCtx = with . useContext----- |--- Detach the context, and destroy if no longer used----{-# INLINEABLE detach #-}-detach :: Context -> IO ()-detach !ctx = nothingIfOk =<< cuCtxDetach ctx--{-# INLINE cuCtxDetach #-}-{# fun unsafe cuCtxDetach- { useContext `Context' } -> `Status' cToEnum #}----- |--- Destroy the specified context. This fails if the context is more than a--- single attachment (including that from initial creation).----{-# INLINEABLE destroy #-}-destroy :: Context -> IO ()-destroy !ctx = nothingIfOk =<< cuCtxDestroy ctx--{-# INLINE cuCtxDestroy #-}-{# fun unsafe cuCtxDestroy- { useContext `Context' } -> `Status' cToEnum #}----- |--- Return the context bound to the calling CPU thread. Requires cuda-4.0.----{-# INLINEABLE get #-}-get :: IO Context-#if CUDA_VERSION < 4000-get = requireSDK 4.0 "get"-#else-get = resultIfOk =<< cuCtxGetCurrent--{-# INLINE cuCtxGetCurrent #-}-{# fun unsafe cuCtxGetCurrent- { alloca- `Context' peekCtx* } -> `Status' cToEnum #}- where peekCtx = liftM Context . peek-#endif----- |--- Bind the specified context to the calling thread. Requires cuda-4.0.----{-# INLINEABLE set #-}-set :: Context -> IO ()-#if CUDA_VERSION < 4000-set _ = requireSDK 4.0 "set"-#else-set !ctx = nothingIfOk =<< cuCtxSetCurrent ctx--{-# INLINE cuCtxSetCurrent #-}-{# fun unsafe cuCtxSetCurrent- { useContext `Context' } -> `Status' cToEnum #}-#endif---- |--- Return the device of the currently active context----{-# INLINEABLE device #-}-device :: IO Device-device = resultIfOk =<< cuCtxGetDevice--{-# INLINE cuCtxGetDevice #-}-{# fun unsafe cuCtxGetDevice- { alloca- `Device' dev* } -> `Status' cToEnum #}- where dev = liftM Device . peekIntConv----- |--- Pop the current CUDA context from the CPU thread. The context must have a--- single usage count (matching calls to 'attach' and 'detach'). If successful,--- the new context is returned, and the old may be attached to a different CPU.----{-# INLINEABLE pop #-}-pop :: IO Context-pop = resultIfOk =<< cuCtxPopCurrent--{-# INLINE cuCtxPopCurrent #-}-{# fun unsafe cuCtxPopCurrent- { alloca- `Context' peekCtx* } -> `Status' cToEnum #}- where peekCtx = liftM Context . peek----- |--- Push the given context onto the CPU's thread stack of current contexts. The--- context must be floating (via 'pop'), i.e. not attached to any thread.----{-# INLINEABLE push #-}-push :: Context -> IO ()-push !ctx = nothingIfOk =<< cuCtxPushCurrent ctx--{-# INLINE cuCtxPushCurrent #-}-{# fun unsafe cuCtxPushCurrent- { useContext `Context' } -> `Status' cToEnum #}----- |--- Block until the device has completed all preceding requests----{-# INLINEABLE sync #-}-sync :: IO ()-sync = nothingIfOk =<< cuCtxSynchronize--{-# INLINE cuCtxSynchronize #-}-{# fun cuCtxSynchronize- { } -> `Status' cToEnum #}-------------------------------------------------------------------------------------- Peer access------------------------------------------------------------------------------------- |--- 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 CUDA_VERSION < 4000-accessible _ _ = requireSDK 4.0 "accessible"-#else-accessible !dev !peer = resultIfOk =<< cuDeviceCanAccessPeer dev peer--{-# INLINE cuDeviceCanAccessPeer #-}-{# fun unsafe cuDeviceCanAccessPeer- { alloca- `Bool' peekBool*- , useDevice `Device'- , useDevice `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 :: Context -> [PeerFlag] -> IO ()-#if CUDA_VERSION < 4000-add _ _ = requireSDK 4.0 "add"-#else-add !ctx !flags = nothingIfOk =<< cuCtxEnablePeerAccess ctx flags--{-# INLINE cuCtxEnablePeerAccess #-}-{# fun unsafe cuCtxEnablePeerAccess- { useContext `Context'- , combineBitMasks `[PeerFlag]' } -> `Status' cToEnum #}-#endif----- |--- Disable direct memory access from the current context to the supplied--- context. Requires cuda-4.0.----{-# INLINEABLE remove #-}-remove :: Context -> IO ()-#if CUDA_VERSION < 4000-remove _ = requireSDK 4.0 "remove"-#else-remove !ctx = nothingIfOk =<< cuCtxDisablePeerAccess ctx--{-# INLINE cuCtxDisablePeerAccess #-}-{# fun unsafe cuCtxDisablePeerAccess- { useContext `Context' } -> `Status' cToEnum #}-#endif-------------------------------------------------------------------------------------- Cache configuration------------------------------------------------------------------------------------- |--- Query compute 2.0 call stack limits. Requires cuda-3.1.----{-# INLINEABLE getLimit #-}-getLimit :: Limit -> IO Int-#if CUDA_VERSION < 3010-getLimit _ = requireSDK 3.1 "getLimit"-#else-getLimit !l = resultIfOk =<< cuCtxGetLimit l--{-# INLINE cuCtxGetLimit #-}-{# fun unsafe cuCtxGetLimit- { alloca- `Int' peekIntConv*- , cFromEnum `Limit' } -> `Status' cToEnum #}-#endif---- |--- Specify the size of the call stack, for compute 2.0 devices. Requires--- cuda-3.1.----{-# INLINEABLE setLimit #-}-setLimit :: Limit -> Int -> IO ()-#if CUDA_VERSION < 3010-setLimit _ _ = requireSDK 3.1 "setLimit"-#else-setLimit !l !n = nothingIfOk =<< cuCtxSetLimit l n--{-# INLINE cuCtxSetLimit #-}-{# fun unsafe cuCtxSetLimit- { cFromEnum `Limit'- , cIntConv `Int' } -> `Status' cToEnum #}-#endif---- |--- On devices where the L1 cache and shared memory use the same hardware--- resources, this sets the preferred cache configuration for the current--- context. This is only a preference. Requires cuda-3.2.----{-# INLINEABLE setCacheConfig #-}-setCacheConfig :: Cache -> IO ()-#if CUDA_VERSION < 3020-setCacheConfig _ = requireSDK 3.2 "setCacheConfig"-#else-setCacheConfig !c = nothingIfOk =<< cuCtxSetCacheConfig c--{-# INLINE cuCtxSetCacheConfig #-}-{# fun unsafe cuCtxSetCacheConfig- { cFromEnum `Cache' } -> `Status' cToEnum #}-#endif-
+ Foreign/CUDA/Driver/Context.hs view
@@ -0,0 +1,23 @@+--------------------------------------------------------------------------------+-- |+-- Module : Foreign.CUDA.Driver.Context+-- Copyright : [2009..2015] Trevor L. McDonell+-- License : BSD+--+-- Context management for low-level driver interface+--+--------------------------------------------------------------------------------+++module Foreign.CUDA.Driver.Context (++ module Foreign.CUDA.Driver.Context.Base,+ module Foreign.CUDA.Driver.Context.Config,+ module Foreign.CUDA.Driver.Context.Peer,++) where++import Foreign.CUDA.Driver.Context.Base+import Foreign.CUDA.Driver.Context.Config+import Foreign.CUDA.Driver.Context.Peer+
+ Foreign/CUDA/Driver/Context/Base.chs view
@@ -0,0 +1,241 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE EmptyDataDecls #-}+{-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-}+#ifdef USE_EMPTY_CASE+{-# LANGUAGE EmptyCase #-}+#endif+--------------------------------------------------------------------------------+-- |+-- Module : Foreign.CUDA.Driver.Context.Base+-- Copyright : [2009..2015] Trevor L. McDonell+-- License : BSD+--+-- Context management for the low-level driver interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Context.Base (++ -- * Context Management+ Context(..), ContextFlag(..),+ create, destroy, device, pop, push, sync, get, set,++ -- Deprecated in CUDA-4.0+ attach, detach,++) where++#include "cbits/stubs.h"+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Driver.Device ( Device(..) )+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Internal.C2HS++-- System+import Foreign+import Foreign.C+import Control.Monad ( liftM )+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- A device context+--+newtype Context = Context { useContext :: {# type CUcontext #}}+ deriving (Eq, Show)+++-- |+-- Context creation flags+--+{# enum CUctx_flags as ContextFlag+ { underscoreToCase }+ with prefix="CU_CTX" deriving (Eq, Show, Bounded) #}+++#if CUDA_VERSION >= 4000+{-# DEPRECATED attach, detach "as of CUDA-4.0" #-}+{-# DEPRECATED BlockingSync "use SchedBlockingSync instead" #-}+#endif+++--------------------------------------------------------------------------------+-- Context management+--------------------------------------------------------------------------------++-- |+-- Create a new CUDA context and associate it with the calling thread. The+-- context is created with a usage count of one, and the caller of 'create'+-- must call 'destroy' when done using the context. If a context is already+-- current to the thread, it is supplanted by the newly created context and+-- must be restored by a subsequent call to 'pop'.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g65dc0012348bc84810e2103a40d8e2cf>+--+{-# INLINEABLE create #-}+create :: Device -> [ContextFlag] -> IO Context+create !dev !flags = resultIfOk =<< cuCtxCreate flags dev++{-# INLINE cuCtxCreate #-}+{# fun unsafe cuCtxCreate+ { alloca- `Context' peekCtx*+ , combineBitMasks `[ContextFlag]'+ , useDevice `Device' } -> `Status' cToEnum #}+ where peekCtx = liftM Context . peek+++-- |+-- Increments the usage count of the context. API: no context flags are+-- currently supported, so this parameter must be empty.+--+{-# INLINEABLE attach #-}+attach :: Context -> [ContextFlag] -> IO ()+attach !ctx !flags = nothingIfOk =<< cuCtxAttach ctx flags++{-# INLINE cuCtxAttach #-}+{# fun unsafe cuCtxAttach+ { withCtx* `Context'+ , combineBitMasks `[ContextFlag]' } -> `Status' cToEnum #}+ where withCtx = with . useContext+++-- |+-- Detach the context, and destroy if no longer used+--+{-# INLINEABLE detach #-}+detach :: Context -> IO ()+detach !ctx = nothingIfOk =<< cuCtxDetach ctx++{-# INLINE cuCtxDetach #-}+{# fun unsafe cuCtxDetach+ { useContext `Context' } -> `Status' cToEnum #}+++-- |+-- Destroy the specified context, regardless of how many threads it is+-- current to. The context will be 'pop'ed from the current thread's+-- context stack, but if it is current on any other threads it will remain+-- current to those threads, and attempts to access it will result in an+-- error.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g27a365aebb0eb548166309f58a1e8b8e>+--+{-# INLINEABLE destroy #-}+destroy :: Context -> IO ()+destroy !ctx = nothingIfOk =<< cuCtxDestroy ctx++{-# INLINE cuCtxDestroy #-}+{# fun unsafe cuCtxDestroy+ { useContext `Context' } -> `Status' cToEnum #}+++-- |+-- Return the context bound to the calling CPU thread.+--+-- Requires CUDA-4.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g8f13165846b73750693640fb3e8380d0>+--+{-# INLINEABLE get #-}+get :: IO (Maybe Context)+#if CUDA_VERSION < 4000+get = requireSDK 'get 4.0+#else+get = resultIfOk =<< cuCtxGetCurrent++{-# INLINE cuCtxGetCurrent #-}+{# fun unsafe cuCtxGetCurrent+ { alloca- `Maybe Context' peekCtx* } -> `Status' cToEnum #}+ where peekCtx = liftM (nothingIfNull Context) . peek+#endif+++-- |+-- Bind the specified context to the calling thread.+--+-- Requires CUDA-4.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1gbe562ee6258b4fcc272ca6478ca2a2f7>+--+{-# INLINEABLE set #-}+set :: Context -> IO ()+#if CUDA_VERSION < 4000+set _ = requireSDK 'set 4.0+#else+set !ctx = nothingIfOk =<< cuCtxSetCurrent ctx++{-# INLINE cuCtxSetCurrent #-}+{# fun unsafe cuCtxSetCurrent+ { useContext `Context' } -> `Status' cToEnum #}+#endif+++-- |+-- Return the device of the currently active context+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g4e84b109eba36cdaaade167f34ae881e>+--+{-# INLINEABLE device #-}+device :: IO Device+device = resultIfOk =<< cuCtxGetDevice++{-# INLINE cuCtxGetDevice #-}+{# fun unsafe cuCtxGetDevice+ { alloca- `Device' dev* } -> `Status' cToEnum #}+ where dev = liftM Device . peekIntConv+++-- |+-- Pop the current CUDA context from the CPU thread. The context may then+-- be attached to a different CPU thread by calling 'push'.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g2fac188026a062d92e91a8687d0a7902>+--+{-# INLINEABLE pop #-}+pop :: IO Context+pop = resultIfOk =<< cuCtxPopCurrent++{-# INLINE cuCtxPopCurrent #-}+{# fun unsafe cuCtxPopCurrent+ { alloca- `Context' peekCtx* } -> `Status' cToEnum #}+ where peekCtx = liftM Context . peek+++-- |+-- Push the given context onto the CPU's thread stack of current contexts.+-- The specified context becomes the CPU thread's current context, so all+-- operations that operate on the current context are affected.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1gb02d4c850eb16f861fe5a29682cc90ba>+--+{-# INLINEABLE push #-}+push :: Context -> IO ()+push !ctx = nothingIfOk =<< cuCtxPushCurrent ctx++{-# INLINE cuCtxPushCurrent #-}+{# fun unsafe cuCtxPushCurrent+ { useContext `Context' } -> `Status' cToEnum #}+++-- |+-- Block until the device has completed all preceding requests. If the+-- context was created with the 'SchedBlockingSync' flag, the CPU thread+-- will block until the GPU has finished its work.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g7a54725f28d34b8c6299f0c6ca579616>+--+{-# INLINEABLE sync #-}+sync :: IO ()+sync = nothingIfOk =<< cuCtxSynchronize++{-# INLINE cuCtxSynchronize #-}+{# fun cuCtxSynchronize+ { } -> `Status' cToEnum #}+
+ Foreign/CUDA/Driver/Context/Config.chs view
@@ -0,0 +1,297 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE EmptyDataDecls #-}+{-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-}+#ifdef USE_EMPTY_CASE+{-# LANGUAGE EmptyCase #-}+#endif+--------------------------------------------------------------------------------+-- |+-- Module : Foreign.CUDA.Driver.Context.Config+-- Copyright : [2009..2015] Trevor L. McDonell+-- License : BSD+--+-- Context configuration for the low-level driver interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Context.Config (++ -- * Context configuration+ getFlags,++ -- ** Resource limits+ Limit(..),+ getLimit, setLimit,++ -- ** Cache+ Cache(..),+ getCache, setCache,++ -- ** Shared memory+ SharedMem(..),+ getSharedMem, setSharedMem,++ -- ** Streams+ StreamPriority,+ getStreamPriorityRange,++) where++#include "cbits/stubs.h"+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Driver.Context.Base+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Internal.C2HS+import Foreign.CUDA.Types++-- System+import Control.Monad+import Foreign+import Foreign.C+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- Device limits flags+--+#if CUDA_VERSION < 3010+data Limit+#else+{# enum CUlimit_enum as Limit+ { underscoreToCase }+ with prefix="CU_LIMIT" deriving (Eq, Show) #}+#endif+++-- |+-- Device cache configuration preference+--+#if CUDA_VERSION < 3020+data Cache+#else+{# enum CUfunc_cache_enum as Cache+ { underscoreToCase }+ with prefix="CU_FUNC_CACHE" deriving (Eq, Show) #}+#endif+++-- |+-- Device shared memory configuration preference+--+#if CUDA_VERSION < 3020+data SharedMem+#else+{# enum CUsharedconfig as SharedMem+ { underscoreToCase }+ with prefix="CU_SHARED_MEM_CONFIG" deriving (Eq, Show) #}+#endif+++--------------------------------------------------------------------------------+-- Context configuration+--------------------------------------------------------------------------------++-- |+-- Return the flags that were used to create the current context.+--+-- Requires CUDA-7.0+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1gf81eef983c1e3b2ef4f166d7a930c86d>+--+{-# INLINEABLE getFlags #-}+getFlags :: IO [ContextFlag]+#if CUDA_VERSION < 7000+getFlags = requireSDK 'getFlags 7.0+#else+getFlags = resultIfOk =<< cuCtxGetFlags++{-# INLINE cuCtxGetFlags #-}+{# fun unsafe cuCtxGetFlags+ { alloca- `[ContextFlag]' peekFlags* } -> `Status' cToEnum #}+ where+ peekFlags = liftM extractBitMasks . peek+#endif+++-- |+-- Query compute 2.0 call stack limits.+--+-- Requires CUDA-3.1.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g9f2d47d1745752aa16da7ed0d111b6a8>+--+{-# INLINEABLE getLimit #-}+getLimit :: Limit -> IO Int+#if CUDA_VERSION < 3010+getLimit _ = requireSDK 'getLimit 3.1+#else+getLimit !l = resultIfOk =<< cuCtxGetLimit l++{-# INLINE cuCtxGetLimit #-}+{# fun unsafe cuCtxGetLimit+ { alloca- `Int' peekIntConv*+ , cFromEnum `Limit' } -> `Status' cToEnum #}+#endif+++-- |+-- Specify the size of the call stack, for compute 2.0 devices.+--+-- Requires CUDA-3.1.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g0651954dfb9788173e60a9af7201e65a>+--+{-# INLINEABLE setLimit #-}+setLimit :: Limit -> Int -> IO ()+#if CUDA_VERSION < 3010+setLimit _ _ = requireSDK 'setLimit 3.1+#else+setLimit !l !n = nothingIfOk =<< cuCtxSetLimit l n++{-# INLINE cuCtxSetLimit #-}+{# fun unsafe cuCtxSetLimit+ { cFromEnum `Limit'+ , cIntConv `Int' } -> `Status' cToEnum #}+#endif+++-- |+-- On devices where the L1 cache and shared memory use the same hardware+-- resources, this function returns the preferred cache configuration for+-- the current context.+--+-- Requires CUDA-3.2.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g40b6b141698f76744dea6e39b9a25360>+--+{-# INLINEABLE getCache #-}+getCache :: IO Cache+#if CUDA_VERSION < 3020+getCache = requireSDK 'getCache 3.2+#else+getCache = resultIfOk =<< cuCtxGetCacheConfig++{-# INLINE cuCtxGetCacheConfig #-}+{# fun unsafe cuCtxGetCacheConfig+ { alloca- `Cache' peekEnum* } -> `Status' cToEnum #}+#endif+++-- |+-- On devices where the L1 cache and shared memory use the same hardware+-- resources, this sets the preferred cache configuration for the current+-- context. This is only a preference.+--+-- Any function configuration set via+-- 'Foreign.CUDA.Driver.Exec.setCacheConfigFun' will be preferred over this+-- context-wide setting.+--+-- Requires CUDA-3.2.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g54699acf7e2ef27279d013ca2095f4a3>+--+{-# INLINEABLE setCache #-}+setCache :: Cache -> IO ()+#if CUDA_VERSION < 3020+setCache _ = requireSDK 'setCache 3.2+#else+setCache !c = nothingIfOk =<< cuCtxSetCacheConfig c++{-# INLINE cuCtxSetCacheConfig #-}+{# fun unsafe cuCtxSetCacheConfig+ { cFromEnum `Cache' } -> `Status' cToEnum #}+#endif+++-- |+-- Return the current size of the shared memory banks in the current+-- context. On devices with configurable shared memory banks,+-- 'setSharedMem' can be used to change the configuration, so that+-- subsequent kernel launches will by default us the new bank size. On+-- devices without configurable shared memory, this function returns the+-- fixed bank size of the hardware.+--+-- Requires CUDA-3.2+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g17153a1b8b8c756f7ab8505686a4ad74>+--+{-# INLINEABLE getSharedMem #-}+getSharedMem :: IO SharedMem+#if CUDA_VERSION < 3020+getSharedMem = requireSDK 'getSharedMem 3.2+#else+getSharedMem = resultIfOk =<< cuCtxGetSharedMemConfig++{-# INLINE cuCtxGetSharedMemConfig #-}+{# fun unsafe cuCtxGetSharedMemConfig+ { alloca- `SharedMem' peekEnum*+ } -> `Status' cToEnum #}+#endif+++-- |+-- On devices with configurable shared memory banks, this function will set+-- the context's shared memory bank size that will be used by default for+-- subsequent kernel launches.+--+-- Changing the shared memory configuration between launches may insert+-- a device synchronisation.+--+-- Shared memory bank size does not affect shared memory usage or kernel+-- occupancy, but may have major effects on performance. Larger bank sizes+-- allow for greater potential bandwidth to shared memory, but change the+-- kinds of accesses which result in bank conflicts.+--+-- Requires CUDA-3.2+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g2574235fa643f8f251bf7bc28fac3692>+--+{-# INLINEABLE setSharedMem #-}+setSharedMem :: SharedMem -> IO ()+#if CUDA_VERSION < 3020+setSharedMem _ = requireSDK 'setSharedMem 3.2+#else+setSharedMem !c = nothingIfOk =<< cuCtxSetSharedMemConfig c++{-# INLINE cuCtxSetSharedMemConfig #-}+{# fun unsafe cuCtxSetSharedMemConfig+ { cFromEnum `SharedMem' } -> `Status' cToEnum #}+#endif++++-- |+-- Returns the numerical values that correspond to the greatest and least+-- priority execution streams in the current context respectively. Stream+-- priorities follow the convention that lower numerical numbers correspond+-- to higher priorities. The range of meaningful stream priorities is given+-- by the inclusive range [greatestPriority,leastPriority].+--+-- Requires CUDA-5.5.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g137920ab61a71be6ce67605b9f294091>+--+{-# INLINEABLE getStreamPriorityRange #-}+getStreamPriorityRange :: IO (StreamPriority, StreamPriority)+#if CUDA_VERSION < 5050+getStreamPriorityRange = requireSDK 'getStreamPriorityRange 5.5+#else+getStreamPriorityRange = do+ (r,l,h) <- cuCtxGetStreamPriorityRange+ resultIfOk (r, (h,l))++{-# INLINE cuCtxGetStreamPriorityRange #-}+{# fun unsafe cuCtxGetStreamPriorityRange+ { alloca- `Int' peekIntConv*+ , alloca- `Int' peekIntConv*+ }+ -> `Status' cToEnum #}+#endif+
+ Foreign/CUDA/Driver/Context/Peer.chs view
@@ -0,0 +1,131 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE EmptyDataDecls #-}+{-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-}+#ifdef USE_EMPTY_CASE+{-# LANGUAGE EmptyCase #-}+#endif+--------------------------------------------------------------------------------+-- |+-- Module : Foreign.CUDA.Driver.Context.Peer+-- Copyright : [2009..2015] Trevor L. McDonell+-- License : BSD+--+-- Direct peer context access functions for the low-level driver interface.+--+-- Since: CUDA-4.0+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Context.Peer (++ -- * Peer Access+ PeerFlag,+ accessible, add, remove,++) where++#include "cbits/stubs.h"+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Driver.Context.Base ( Context(..) )+import Foreign.CUDA.Driver.Device ( Device(..) )+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Internal.C2HS++-- System+import Foreign+import Foreign.C+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- Possible option values for direct peer memory access+--+data PeerFlag+instance Enum PeerFlag where+#ifdef USE_EMPTY_CASE+ toEnum x = case x of {}+ fromEnum x = case x of {}+#endif+++--------------------------------------------------------------------------------+-- Peer access+--------------------------------------------------------------------------------++-- |+-- 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.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__PEER__ACCESS.html#group__CUDA__PEER__ACCESS_1g496bdaae1f632ebfb695b99d2c40f19e>+--+{-# INLINEABLE accessible #-}+accessible :: Device -> Device -> IO Bool+#if CUDA_VERSION < 4000+accessible _ _ = requireSDK 'accessible 4.0+#else+accessible !dev !peer = resultIfOk =<< cuDeviceCanAccessPeer dev peer++{-# INLINE cuDeviceCanAccessPeer #-}+{# fun unsafe cuDeviceCanAccessPeer+ { alloca- `Bool' peekBool*+ , useDevice `Device'+ , useDevice `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.+--+-- Note that access is unidirectional, and in order to access memory in the+-- current context from the peer context, a separate symmetric call to+-- 'add' is required.+--+-- Requires CUDA-4.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__PEER__ACCESS.html#group__CUDA__PEER__ACCESS_1g0889ec6728e61c05ed359551d67b3f5a>+--+{-# INLINEABLE add #-}+add :: Context -> [PeerFlag] -> IO ()+#if CUDA_VERSION < 4000+add _ _ = requireSDK 'add 4.0+#else+add !ctx !flags = nothingIfOk =<< cuCtxEnablePeerAccess ctx flags++{-# INLINE cuCtxEnablePeerAccess #-}+{# fun unsafe cuCtxEnablePeerAccess+ { useContext `Context'+ , combineBitMasks `[PeerFlag]' } -> `Status' cToEnum #}+#endif+++-- |+-- Disable direct memory access from the current context to the supplied+-- peer context, and unregisters any registered allocations.+--+-- Requires CUDA-4.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__PEER__ACCESS.html#group__CUDA__PEER__ACCESS_1g5b4b6936ea868d4954ce4d841a3b4810>+--+{-# INLINEABLE remove #-}+remove :: Context -> IO ()+#if CUDA_VERSION < 4000+remove _ = requireSDK 'remave 4.0+#else+remove !ctx = nothingIfOk =<< cuCtxDisablePeerAccess ctx++{-# INLINE cuCtxDisablePeerAccess #-}+{# fun unsafe cuCtxDisablePeerAccess+ { useContext `Context' } -> `Status' cToEnum #}+#endif+
+ Foreign/CUDA/Driver/Context/Primary.chs view
@@ -0,0 +1,170 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-}+--------------------------------------------------------------------------------+-- |+-- Module : Foreign.CUDA.Driver.Context.Primary+-- Copyright : [2009..2014] Trevor L. McDonell+-- License : BSD+--+-- Primary context management for low-level driver interface. The primary+-- context is unique per device and shared with the Runtime API. This+-- allows integration with other libraries using CUDA.+--+-- Since: CUDA-7.0+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Context.Primary (++ status, setup, reset, retain, release,++) where++#include "cbits/stubs.h"+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Driver.Context.Base+import Foreign.CUDA.Driver.Device+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Internal.C2HS++-- System+import Control.Exception+import Control.Monad+import Foreign+import Foreign.C+++--------------------------------------------------------------------------------+-- Primary context management+--------------------------------------------------------------------------------+++-- |+-- Get the status of the primary context. Returns whether the current+-- context is active, and the flags it was (or will be) created with.+--+-- Requires CUDA-7.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__PRIMARY__CTX.html#group__CUDA__PRIMARY__CTX_1g65f3e018721b6d90aa05cfb56250f469>+--+{-# INLINEABLE status #-}+status :: Device -> IO (Bool, [ContextFlag])+#if CUDA_VERSION < 7000+status _ = requireSDK 'status 7.0+#else+status !dev =+ cuDevicePrimaryCtxGetState dev >>= \(rv, !flags, !active) ->+ case rv of+ Success -> return (active, flags)+ _ -> throwIO (ExitCode rv)++{# fun unsafe cuDevicePrimaryCtxGetState+ { useDevice `Device'+ , alloca- `[ContextFlag]' peekFlags*+ , alloca- `Bool' peekBool*+ } -> `Status' cToEnum #}+ where+ peekFlags = liftM extractBitMasks . peek+#endif+++-- |+-- Specify the flags that the primary context should be created with. Note+-- that this is an error if the primary context is already active.+--+-- Requires CUDA-7.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__PRIMARY__CTX.html#group__CUDA__PRIMARY__CTX_1gd779a84f17acdad0d9143d9fe719cfdf>+--+{-# INLINEABLE setup #-}+setup :: Device -> [ContextFlag] -> IO ()+#if CUDA_VERSION < 7000+setup _ _ = requireSDK 'setup 7.0+#else+setup !dev !flags = nothingIfOk =<< cuDevicePrimaryCtxSetFlags dev flags++{-# INLINE cuDevicePrimaryCtxSetFlags #-}+{# fun unsafe cuDevicePrimaryCtxSetFlags+ { useDevice `Device'+ , combineBitMasks `[ContextFlag]'+ } -> `Status' cToEnum #}+#endif+++-- |+-- Destroy all allocations and reset all state on the primary context of+-- the given device in the current process. Requires cuda-7.0+--+-- Requires CUDA-7.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__PRIMARY__CTX.html#group__CUDA__PRIMARY__CTX_1g5d38802e8600340283958a117466ce12>+--+{-# INLINEABLE reset #-}+reset :: Device -> IO ()+#if CUDA_VERSION < 7000+reset _ = requireSDK 'reset 7.0+#else+reset !dev = nothingIfOk =<< cuDevicePrimaryCtxReset dev++{-# INLINE cuDevicePrimaryCtxReset #-}+{# fun unsafe cuDevicePrimaryCtxReset+ { useDevice `Device' } -> `Status' cToEnum #}+#endif+++-- |+-- Release the primary context on the given device. If there are no more+-- references to the primary context it will be destroyed, regardless of+-- how many threads it is current to.+--+-- Unlike 'Foreign.CUDA.Driver.Context.Base.pop' this does not pop the+-- context from the stack in any circumstances.+--+-- Requires CUDA-7.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__PRIMARY__CTX.html#group__CUDA__PRIMARY__CTX_1gf2a8bc16f8df0c88031f6a1ba3d6e8ad>+--+{-# INLINEABLE release #-}+release :: Device -> IO ()+#if CUDA_VERSION < 7000+release _ = requireSDK 'release 7.0+#else+release !dev = nothingIfOk =<< cuDevicePrimaryCtxRelease dev++{-# INLINE cuDevicePrimaryCtxRelease #-}+{# fun unsafe cuDevicePrimaryCtxRelease+ { useDevice `Device' } -> `Status' cToEnum #}+#endif+++-- |+-- Retain the primary context for the given device, creating it if+-- necessary, and increasing its usage count. The caller must call+-- 'release' when done using the context. Unlike+-- 'Foreign.CUDA.Driver.Context.Base.create' the newly retained context is+-- not pushed onto the stack.+--+-- Requires CUDA-7.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__PRIMARY__CTX.html#group__CUDA__PRIMARY__CTX_1g9051f2d5c31501997a6cb0530290a300>+--+{-# INLINEABLE retain #-}+retain :: Device -> IO Context+#if CUDA_VERSION < 7000+retain _ = requireSDK 'retain 7.0+#else+retain !dev = resultIfOk =<< cuDevicePrimaryCtxRetain dev++{-# INLINE cuDevicePrimaryCtxRetain #-}+{# fun unsafe cuDevicePrimaryCtxRetain+ { alloca- `Context' peekCtx*+ , useDevice `Device'+ } -> `Status' cToEnum #}+ where+ peekCtx = liftM Context . peek+#endif+
Foreign/CUDA/Driver/Device.chs view
@@ -18,7 +18,7 @@ module Foreign.CUDA.Driver.Device ( -- * Device Management- Device(..), -- should be exported abstractly+ Device(..), DeviceProperties(..), DeviceAttribute(..), Compute(..), ComputeMode(..), InitFlag, initialise, capability, device, attribute, count, name, props, totalMem @@ -36,12 +36,17 @@ import Foreign import Foreign.C import Control.Monad ( liftM )+import Control.Applicative+import Prelude -------------------------------------------------------------------------------- -- Data Types -------------------------------------------------------------------------------- +-- |+-- A CUDA device+-- newtype Device = Device { useDevice :: {# type CUdevice #}} deriving (Eq, Show) @@ -54,11 +59,13 @@ , MAX as CU_DEVICE_ATTRIBUTE_MAX } -- ignore with prefix="CU_DEVICE_ATTRIBUTE" deriving (Eq, Show) #} -{# pointer *CUdevprop as ^ foreign -> CUDevProp nocode #} +#if CUDA_VERSION < 5000+{# pointer *CUdevprop as ^ foreign -> CUDevProp nocode #} ----- Properties of the compute device (internal helper)+-- Properties of the compute device (internal helper).+-- Replaced by cuDeviceGetAttribute in CUDA-5.0 and later. -- data CUDevProp = CUDevProp {@@ -107,6 +114,7 @@ cuClockRate = cl, cuTextureAlignment = ta }+#endif -- |@@ -126,9 +134,11 @@ -------------------------------------------------------------------------------- -- |--- Initialise the CUDA driver API. Must be called before any other driver--- function.+-- Initialise the CUDA driver API. This must be called before any other+-- driver function. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__INITIALIZE.html#group__CUDA__INITIALIZE_1g0a2f1517e1bd8502c7194c3a8c134bc3>+-- {-# INLINEABLE initialise #-} initialise :: [InitFlag] -> IO () initialise !flags = nothingIfOk =<< cuInit flags@@ -147,7 +157,14 @@ -- {-# INLINEABLE capability #-} capability :: Device -> IO Compute+#if CUDA_VERSION >= 5000 capability !dev =+ Compute <$> attribute dev ComputeCapabilityMajor+ <*> attribute dev ComputeCapabilityMinor+#else+-- Deprecated as of CUDA-5.0+--+capability !dev = (\(!s,!a,!b) -> resultIfOk (s,Compute a b)) =<< cuDeviceComputeCapability dev {-# INLINE cuDeviceComputeCapability #-}@@ -155,11 +172,14 @@ { alloca- `Int' peekIntConv* , alloca- `Int' peekIntConv* , useDevice `Device' } -> `Status' cToEnum #}+#endif -- |--- Return a device handle+-- Return a handle to the compute device at the given ordinal. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__DEVICE.html#group__CUDA__DEVICE_1g8bdd1cc7201304b01357b8034f6587cb>+-- {-# INLINEABLE device #-} device :: Int -> IO Device device !d = resultIfOk =<< cuDeviceGet d@@ -172,8 +192,10 @@ -- |--- Return the selected attribute for the given device+-- Return the selected attribute for the given device. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__DEVICE.html#group__CUDA__DEVICE_1g9c3e1414f0ad901d3278a4d6645fc266>+-- {-# INLINEABLE attribute #-} attribute :: Device -> DeviceAttribute -> IO Int attribute !d !a = resultIfOk =<< cuDeviceGetAttribute a d@@ -186,8 +208,10 @@ -- |--- Return the number of device with compute capability > 1.0+-- Return the number of device with compute capability > 1.0. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__DEVICE.html#group__CUDA__DEVICE_1g52b5ce05cb8c5fb6831b2c0ff2887c74>+-- {-# INLINEABLE count #-} count :: IO Int count = resultIfOk =<< cuDeviceGetCount@@ -198,8 +222,10 @@ -- |--- Name of the device+-- The identifying name of the device. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__DEVICE.html#group__CUDA__DEVICE_1gef75aa30df95446a845f2a7b9fffbb7f>+-- {-# INLINEABLE name #-} name :: Device -> IO String name !d = resultIfOk =<< cuDeviceGetName d@@ -217,16 +243,45 @@ -- | -- Return the properties of the selected device ------ Annoyingly, the driver API requires several different functions to extract--- all device properties that are part of a single structure in the runtime API--- {-# INLINEABLE props #-} props :: Device -> IO DeviceProperties props !d = do++#if CUDA_VERSION < 5000+ -- Old versions of the CUDA API used the separate cuDeviceGetProperties+ -- function to probe some properties, and cuDeviceGetAttribute for+ -- others. As of CUDA-5.0, the former was deprecated and its+ -- functionality subsumed by the latter, which we use below.+ -- p <- resultIfOk =<< cuDeviceGetProperties d+ let cm = cuTotalConstMem p+ sm = cuSharedMemPerBlock p+ rb = cuRegsPerBlock p+ ws = cuWarpSize p+ tb = cuMaxThreadsPerBlock p+ bs = cuMaxBlockSize p+ gs = cuMaxGridSize p+ cl = cuClockRate p+ mp = cuMemPitch p+ ta = cuTextureAlignment p+#else+ cm <- fromIntegral <$> attribute d TotalConstantMemory+ sm <- fromIntegral <$> attribute d SharedMemoryPerBlock+ mp <- fromIntegral <$> attribute d MaxPitch+ ta <- fromIntegral <$> attribute d TextureAlignment+ cl <- attribute d ClockRate+ ws <- attribute d WarpSize+ rb <- attribute d RegistersPerBlock+ tb <- attribute d MaxThreadsPerBlock+ bs <- (,,) <$> attribute d MaxBlockDimX+ <*> attribute d MaxBlockDimY+ <*> attribute d MaxBlockDimZ+ gs <- (,,) <$> attribute d MaxGridDimX+ <*> attribute d MaxGridDimY+ <*> attribute d MaxGridDimZ+#endif - -- And the remaining properties+ -- The rest of the properties. -- n <- name d cc <- capability d@@ -259,56 +314,80 @@ ua <- toBool `fmap` attribute d UnifiedAddressing tcc <- toBool `fmap` attribute d TccDriver #endif+#if CUDA_VERSION >= 5050+ sp <- toBool `fmap` attribute d StreamPrioritiesSupported+#endif+#if CUDA_VERSION >= 6000+ gl1 <- toBool `fmap` attribute d GlobalL1CacheSupported+ ll1 <- toBool `fmap` attribute d LocalL1CacheSupported+ mm <- toBool `fmap` attribute d ManagedMemory+ mg <- toBool `fmap` attribute d MultiGpuBoard+ mid <- attribute d MultiGpuBoardGroupId+#endif return DeviceProperties {- deviceName = n,- computeCapability = cc,- totalGlobalMem = gm,- totalConstMem = cuTotalConstMem p,- sharedMemPerBlock = cuSharedMemPerBlock p,- regsPerBlock = cuRegsPerBlock p,- warpSize = cuWarpSize p,- maxThreadsPerBlock = cuMaxThreadsPerBlock p,- maxBlockSize = cuMaxBlockSize p,- maxGridSize = cuMaxGridSize p,- clockRate = cuClockRate p,- multiProcessorCount = pc,- memPitch = cuMemPitch p,- textureAlignment = cuTextureAlignment p,- computeMode = md,- deviceOverlap = ov,+ deviceName = n+ , computeCapability = cc+ , totalGlobalMem = gm+ , totalConstMem = cm+ , sharedMemPerBlock = sm+ , regsPerBlock = rb+ , warpSize = ws+ , maxThreadsPerBlock = tb+ , maxBlockSize = bs+ , maxGridSize = gs+ , clockRate = cl+ , multiProcessorCount = pc+ , memPitch = mp+ , textureAlignment = ta+ , computeMode = md+ , deviceOverlap = ov+ , kernelExecTimeoutEnabled = ke+ , integrated = tg+ , canMapHostMemory = hm #if CUDA_VERSION >= 3000- concurrentKernels = ck,- eccEnabled = ee,- maxTextureDim1D = u1,- maxTextureDim2D = (u21,u22),- maxTextureDim3D = (u31,u32,u33),+ , concurrentKernels = ck+ , eccEnabled = ee+ , maxTextureDim1D = u1+ , maxTextureDim2D = (u21,u22)+ , maxTextureDim3D = (u31,u32,u33) #endif #if CUDA_VERSION >= 4000- asyncEngineCount = ae,- cacheMemL2 = l2,- maxThreadsPerMultiProcessor = tm,- memBusWidth = mw,- memClockRate = mc,- pciInfo = PCI pb pd pm,- tccDriverEnabled = tcc,- unifiedAddressing = ua,+ , asyncEngineCount = ae+ , cacheMemL2 = l2+ , maxThreadsPerMultiProcessor = tm+ , memBusWidth = mw+ , memClockRate = mc+ , pciInfo = PCI pb pd pm+ , tccDriverEnabled = tcc+ , unifiedAddressing = ua #endif- kernelExecTimeoutEnabled = ke,- integrated = tg,- canMapHostMemory = hm+#if CUDA_VERSION >= 5050+ , streamPriorities = sp+#endif+#if CUDA_VERSION >= 6000+ , globalL1Cache = gl1+ , localL1Cache = ll1+ , managedMemory = mm+ , multiGPUBoard = mg+ , multiGPUBoardGroupID = mid+#endif } -+#if CUDA_VERSION < 5000+-- Deprecated as of CUDA-5.0 {-# INLINE cuDeviceGetProperties #-} {# fun unsafe cuDeviceGetProperties { alloca- `CUDevProp' peek* , useDevice `Device' } -> `Status' cToEnum #}+#endif -- |--- Total memory available on the device (bytes)+-- The total memory available on the device (bytes).+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__DEVICE.html#group__CUDA__DEVICE_1gc6a0d6551335a3780f9f3c967a0fde5d> -- {-# INLINEABLE totalMem #-} totalMem :: Device -> IO Int64
Foreign/CUDA/Driver/Error.chs view
@@ -11,22 +11,32 @@ -- -------------------------------------------------------------------------------- -module Foreign.CUDA.Driver.Error- where+module Foreign.CUDA.Driver.Error ( + -- * CUDA Errors+ Status(..), CUDAException(..),+ describe,+ cudaError, requireSDK,+ resultIfOk, nothingIfOk,++) where+ -- Friends import Foreign.CUDA.Internal.C2HS -- System-import Data.Typeable import Control.Exception import Control.Monad+import Data.Typeable import Foreign.C-import Foreign.Ptr import Foreign.Marshal+import Foreign.Ptr import Foreign.Storable+import Language.Haskell.TH import System.IO.Unsafe+import Text.Printf + #include "cbits/stubs.h" {# context lib="cuda" #} @@ -157,8 +167,8 @@ -- | -- A specially formatted error message ---requireSDK :: Double -> String -> IO a-requireSDK v s = cudaError ("'" ++ s ++ "' requires at least cuda-" ++ show v)+requireSDK :: Name -> Double -> IO a+requireSDK n v = cudaError $ printf "'%s' requires at least cuda-%3.1f\n" (show n) v --------------------------------------------------------------------------------
Foreign/CUDA/Driver/Event.chs view
@@ -2,6 +2,7 @@ {-# LANGUAGE CPP #-} {-# LANGUAGE EmptyDataDecls #-} {-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-} -------------------------------------------------------------------------------- -- | -- Module : Foreign.CUDA.Driver.Event@@ -43,6 +44,8 @@ -- | -- Create a new event --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EVENT.html#group__CUDA__EVENT_1g450687e75f3ff992fe01662a43d9d3db>+-- {-# INLINEABLE create #-} create :: [EventFlag] -> IO Event create !flags = resultIfOk =<< cuEventCreate flags@@ -57,6 +60,8 @@ -- | -- Destroy an event --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EVENT.html#group__CUDA__EVENT_1g593ec73a8ec5a5fc031311d3e4dca1ef>+-- {-# INLINEABLE destroy #-} destroy :: Event -> IO () destroy !ev = nothingIfOk =<< cuEventDestroy ev@@ -69,6 +74,8 @@ -- | -- Determine the elapsed time (in milliseconds) between two events --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EVENT.html#group__CUDA__EVENT_1gdfb1178807353bbcaa9e245da497cf97>+-- {-# INLINEABLE elapsedTime #-} elapsedTime :: Event -> Event -> IO Float elapsedTime !ev1 !ev2 = resultIfOk =<< cuEventElapsedTime ev1 ev2@@ -83,6 +90,8 @@ -- | -- Determines if a event has actually been recorded --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EVENT.html#group__CUDA__EVENT_1g6f0704d755066b0ee705749ae911deef>+-- {-# INLINEABLE query #-} query :: Event -> IO Bool query !ev =@@ -101,6 +110,8 @@ -- Record an event once all operations in the current context (or optionally -- specified stream) have completed. This operation is asynchronous. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EVENT.html#group__CUDA__EVENT_1g95424d3be52c4eb95d83861b70fb89d1>+-- {-# INLINEABLE record #-} record :: Event -> Maybe Stream -> IO () record !ev !mst =@@ -116,12 +127,16 @@ -- 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.+-- different device contexts. --+-- Requires CUDA-3.2.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__STREAM.html#group__CUDA__STREAM_1g6a898b652dfc6aa1d5c8d97062618b2f>+-- {-# INLINEABLE wait #-} wait :: Event -> Maybe Stream -> [WaitFlag] -> IO () #if CUDA_VERSION < 3020-wait _ _ _ = requireSDK 3.2 "wait"+wait _ _ _ = requireSDK 'wait 3.2 #else wait !ev !mst !flags = nothingIfOk =<< cuStreamWaitEvent (fromMaybe defaultStream mst) ev flags@@ -135,6 +150,8 @@ -- | -- Wait until the event has been recorded+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EVENT.html#group__CUDA__EVENT_1g9e520d34e51af7f5375610bca4add99c> -- {-# INLINEABLE block #-} block :: Event -> IO ()
Foreign/CUDA/Driver/Exec.chs view
@@ -3,6 +3,7 @@ {-# LANGUAGE EmptyDataDecls #-} {-# LANGUAGE ForeignFunctionInterface #-} {-# LANGUAGE GADTs #-}+{-# LANGUAGE TemplateHaskell #-} -------------------------------------------------------------------------------- -- | -- Module : Foreign.CUDA.Driver.Exec@@ -16,10 +17,15 @@ module Foreign.CUDA.Driver.Exec ( -- * Kernel Execution- Fun(Fun), FunParam(..), FunAttribute(..),- requires, setBlockShape, setSharedSize, setParams, setCacheConfigFun,- launch, launchKernel, launchKernel'+ Fun(Fun), FunParam(..), FunAttribute(..), SharedMem(..),+ requires,+ setCacheConfigFun,+ setSharedMemConfigFun,+ launchKernel, launchKernel', + -- Deprecated since CUDA-4.0+ setBlockShape, setSharedSize, setParams, launch,+ ) where #include "cbits/stubs.h"@@ -28,7 +34,7 @@ -- Friends import Foreign.CUDA.Internal.C2HS import Foreign.CUDA.Driver.Error-import Foreign.CUDA.Driver.Context ( Cache(..) )+import Foreign.CUDA.Driver.Context ( Cache(..), SharedMem(..) ) import Foreign.CUDA.Driver.Stream ( Stream(..), defaultStream ) -- System@@ -49,7 +55,7 @@ -------------------------------------------------------------------------------- -- |--- A @__global__@ device function+-- A @\_\_global\_\_@ device function -- newtype Fun = Fun { useFun :: {# type CUfunction #}} @@ -92,8 +98,10 @@ -------------------------------------------------------------------------------- -- |--- Returns the value of the selected attribute requirement for the given kernel+-- Returns the value of the selected attribute requirement for the given kernel. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EXEC.html#group__CUDA__EXEC_1g5e92a1b0d8d1b82cb00dcfb2de15961b>+-- {-# INLINEABLE requires #-} requires :: Fun -> FunAttribute -> IO Int requires !fn !att = resultIfOk =<< cuFuncGetAttribute att fn@@ -106,36 +114,6 @@ -- |--- Specify the @(x,y,z)@ dimensions of the thread blocks that are created when--- the given kernel function is launched.----{-# INLINEABLE setBlockShape #-}-setBlockShape :: Fun -> (Int,Int,Int) -> IO ()-setBlockShape !fn (!x,!y,!z) = nothingIfOk =<< cuFuncSetBlockShape fn x y z--{-# INLINE cuFuncSetBlockShape #-}-{# fun unsafe cuFuncSetBlockShape- { useFun `Fun'- , `Int'- , `Int'- , `Int' } -> `Status' cToEnum #}----- |--- Set the number of bytes of dynamic shared memory to be available to each--- thread block when the function is launched----{-# INLINEABLE setSharedSize #-}-setSharedSize :: Fun -> Integer -> IO ()-setSharedSize !fn !bytes = nothingIfOk =<< cuFuncSetSharedSize fn bytes--{-# INLINE cuFuncSetSharedSize #-}-{# fun unsafe cuFuncSetSharedSize- { useFun `Fun'- , cIntConv `Integer' } -> `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@@ -144,10 +122,14 @@ -- Switching between configuration modes may insert a device-side -- synchronisation point for streamed kernel launches. --+-- Requires CUDA-3.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EXEC.html#group__CUDA__EXEC_1g40f8c11e81def95dc0072a375f965681>+-- {-# INLINEABLE setCacheConfigFun #-} setCacheConfigFun :: Fun -> Cache -> IO () #if CUDA_VERSION < 3000-setCacheConfigFun _ _ = requireSDK 3.0 "setCacheConfigFun"+setCacheConfigFun _ _ = requireSDK 'setCacheConfigFun 3.0 #else setCacheConfigFun !fn !pref = nothingIfOk =<< cuFuncSetCacheConfig fn pref @@ -157,22 +139,49 @@ , cFromEnum `Cache' } -> `Status' cToEnum #} #endif + -- |--- Invoke the kernel on a size @(w,h)@ grid of blocks. Each block contains the--- number of threads specified by a previous call to 'setBlockShape'. The launch--- may also be associated with a specific 'Stream'.+-- Set the shared memory configuration of a device function. ---{-# INLINEABLE launch #-}-launch :: Fun -> (Int,Int) -> Maybe Stream -> IO ()-launch !fn (!w,!h) mst =- nothingIfOk =<< cuLaunchGridAsync fn w h (fromMaybe defaultStream mst)+-- On devices with configurable shared memory banks, this will force all+-- subsequent launches of the given device function to use the specified+-- shared memory bank size configuration. On launch of the function, the+-- shared memory configuration of the device will be temporarily changed if+-- needed to suit the function configuration. Changes in shared memory+-- configuration may introduction a device side synchronisation between+-- kernel launches.+--+-- Any per-function configuration specified by 'setSharedMemConfig' will+-- override the context-wide configuration set with+-- 'Foreign.CUDA.Driver.Context.Config.setSharedMem'.+--+-- Changing the shared memory bank size will not increase shared memory+-- usage or affect occupancy of kernels, but may have major effects on+-- performance. Larger bank sizes will allow for greater potential+-- bandwidth to shared memory, but will change what kinds of accesses to+-- shared memory will result in bank conflicts.+--+-- This function will do nothing on devices with fixed shared memory bank+-- size.+--+-- Requires CUDA-5.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EXEC.html#group__CUDA__EXEC_1g430b913f24970e63869635395df6d9f5>+--+{-# INLINEABLE setSharedMemConfigFun #-}+setSharedMemConfigFun :: Fun -> SharedMem -> IO ()+#if CUDA_VERSION < 5000+setSharedMemConfigFun _ _ = requireSDK 'setSharedMemConfigFun 5.0+#else+setSharedMemConfigFun !fun !pref = nothingIfOk =<< cuFuncSetSharedMemConfig fun pref -{-# INLINE cuLaunchGridAsync #-}-{# fun unsafe cuLaunchGridAsync+{-# INLINE cuFuncSetSharedMemConfig #-}+{# fun unsafe cuFuncSetSharedMemConfig { useFun `Fun'- , `Int'- , `Int'- , useStream `Stream' } -> `Status' cToEnum #}+ , cFromEnum `SharedMem'+ }+ -> `Status' cToEnum #}+#endif -- |@@ -189,6 +198,8 @@ -- requiring the application to know the size and alignment/padding of each -- kernel parameter. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EXEC.html#group__CUDA__EXEC_1gb8f3dc3031b40da29d5f9a7139e52e15>+-- {-# INLINEABLE launchKernel #-} {-# INLINEABLE launchKernel' #-} launchKernel, launchKernel'@@ -270,10 +281,58 @@ ----------------------------------------------------------------------------------- Kernel function parameters+-- Deprecated -------------------------------------------------------------------------------- -- |+-- Invoke the kernel on a size @(w,h)@ grid of blocks. Each block contains the+-- number of threads specified by a previous call to 'setBlockShape'. The launch+-- may also be associated with a specific 'Stream'.+--+{-# INLINEABLE launch #-}+launch :: Fun -> (Int,Int) -> Maybe Stream -> IO ()+launch !fn (!w,!h) mst =+ nothingIfOk =<< cuLaunchGridAsync fn w h (fromMaybe defaultStream mst)++{-# INLINE cuLaunchGridAsync #-}+{# fun unsafe cuLaunchGridAsync+ { useFun `Fun'+ , `Int'+ , `Int'+ , useStream `Stream' } -> `Status' cToEnum #}+++-- |+-- Specify the @(x,y,z)@ dimensions of the thread blocks that are created when+-- the given kernel function is launched.+--+{-# INLINEABLE setBlockShape #-}+setBlockShape :: Fun -> (Int,Int,Int) -> IO ()+setBlockShape !fn (!x,!y,!z) = nothingIfOk =<< cuFuncSetBlockShape fn x y z++{-# INLINE cuFuncSetBlockShape #-}+{# fun unsafe cuFuncSetBlockShape+ { useFun `Fun'+ , `Int'+ , `Int'+ , `Int' } -> `Status' cToEnum #}+++-- |+-- Set the number of bytes of dynamic shared memory to be available to each+-- thread block when the function is launched+--+{-# INLINEABLE setSharedSize #-}+setSharedSize :: Fun -> Integer -> IO ()+setSharedSize !fn !bytes = nothingIfOk =<< cuFuncSetSharedSize fn bytes++{-# INLINE cuFuncSetSharedSize #-}+{# fun unsafe cuFuncSetSharedSize+ { useFun `Fun'+ , cIntConv `Integer' } -> `Status' cToEnum #}+++-- | -- Set the parameters that will specified next time the kernel is invoked -- {-# INLINEABLE setParams #-}@@ -317,4 +376,5 @@ , `Int' , castPtr `Ptr a' , `Int' } -> `Status' cToEnum #}+
+ Foreign/CUDA/Driver/IPC/Event.chs view
@@ -0,0 +1,136 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-}+--------------------------------------------------------------------------------+-- |+-- Module : Foreign.CUDA.Driver.IPC.Event+-- Copyright : [2009..2015] Trevor L. McDonell+-- License : BSD+--+-- IPC event management for low-level driver interface.+--+-- Restricted to devices which support unified addressing on Linux+-- operating systems.+--+-- Since CUDA-4.0.+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.IPC.Event (++ IPCEvent,+ export, open,++) where++#include "cbits/stubs.h"+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Driver.Event+import Foreign.CUDA.Internal.C2HS++-- System+import Control.Monad+import Prelude++import Foreign.C+import Foreign.Ptr+import Foreign.ForeignPtr+import Foreign.Marshal+import Foreign.Storable+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- A CUDA inter-process event handle.+--+newtype IPCEvent = IPCEvent { useIPCEvent :: IPCEventHandle }+ deriving (Eq, Show)+++--------------------------------------------------------------------------------+-- IPC event management+--------------------------------------------------------------------------------++-- |+-- Create an inter-process event handle for a previously allocated event.+-- The event must be created with the 'Interprocess' and 'DisableTiming'+-- event flags. The returned handle may then be sent to another process and+-- 'open'ed to allow efficient hardware synchronisation between GPU work in+-- other processes.+--+-- After the event has been opened in the importing process, 'record',+-- 'block', 'wait', 'query' may be used in either process.+--+-- Performing operations on the imported event after the event has been+-- 'destroy'ed in the exporting process is undefined.+--+-- Requires CUDA-4.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1gea02eadd12483de5305878b13288a86c>+--+{-# INLINEABLE export #-}+export :: Event -> IO IPCEvent+#if CUDA_VERSION < 4000+export _ = requireSDK 'create 4.0+#else+export !ev = do+ h <- newIPCEventHandle+ r <- cuIpcGetEventHandle h ev+ resultIfOk (r, IPCEvent h)++{-# INLINE cuIpcGetEventHandle #-}+{# fun unsafe cuIpcGetEventHandle+ { withForeignPtr* `IPCEventHandle'+ , useEvent `Event'+ }+ -> `Status' cToEnum #}+#endif+++-- |+-- Open an inter-process event handle for use in the current process,+-- returning an event that can be used in the current process and behaving+-- as a locally created event with the 'DisableTiming' flag specified.+--+-- The event must be freed with 'destroy'. Performing operations on the+-- imported event after the exported event has been 'destroy'ed in the+-- exporting process is undefined.+--+-- Requires CUDA-4.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1gf1d525918b6c643b99ca8c8e42e36c2e>+--+{-# INLINEABLE open #-}+open :: IPCEvent -> IO Event+#if CUDA_VERSION < 4000+open _ = requireSDK 'open 4.0+#else+open !ev = resultIfOk =<< cuIpcOpenEventHandle (useIPCEvent ev)++{-# INLINE cuIpcOpenEventHandle #-}+{# fun unsafe cuIpcOpenEventHandle+ { alloca- `Event' peekEvent*+ , withForeignPtr* `IPCEventHandle'+ }+ -> `Status' cToEnum #}+ where+ peekEvent = liftM Event . peek+#endif+++--------------------------------------------------------------------------------+-- Internal+--------------------------------------------------------------------------------++type IPCEventHandle = ForeignPtr ()++newIPCEventHandle :: IO IPCEventHandle+newIPCEventHandle = mallocForeignPtrBytes {#sizeof CUipcEventHandle#}+
+ Foreign/CUDA/Driver/IPC/Marshal.chs view
@@ -0,0 +1,168 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-}+--------------------------------------------------------------------------------+-- |+-- Module : Foreign.CUDA.Driver.IPC.Marshal+-- Copyright : [2009..2015] Trevor L. McDonell+-- License : BSD+--+-- IPC memory management for low-level driver interface.+--+-- Restricted to devices which support unified addressing on Linux+-- operating systems.+--+-- Since CUDA-4.0.+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.IPC.Marshal (++ -- ** IPC memory management+ IPCDevicePtr, IPCFlag(..),+ export, open, close,++) where++#include "cbits/stubs.h"+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Ptr+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Internal.C2HS+import Foreign.CUDA.Driver.Marshal++-- System+import Control.Monad+import Prelude++import Foreign.C+import Foreign.Ptr+import Foreign.ForeignPtr+import Foreign.Marshal+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- A CUDA memory handle used for inter-process communication.+--+newtype IPCDevicePtr a = IPCDevicePtr { useIPCDevicePtr :: IPCMemHandle }+ deriving (Eq, Show)+++-- |+-- Flags for controlling IPC memory access+--+{# enum CUipcMem_flags as IPCFlag+ { underscoreToCase }+ with prefix="CU_IPC_MEM" deriving (Eq, Show, Bounded) #}+++--------------------------------------------------------------------------------+-- IPC memory management+--------------------------------------------------------------------------------++-- |+-- Create an inter-process memory handle for an existing device memory+-- allocation. The handle can then be sent to another process and made+-- available to that process via 'open'.+--+-- Requires CUDA-4.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g6f1b5be767b275f016523b2ac49ebec1>+--+{-# INLINEABLE export #-}+export :: DevicePtr a -> IO (IPCDevicePtr a)+#if CUDA_VERSION < 4000+export _ = requireSDK 'create 4.0+#else+export !dptr = do+ h <- newIPCMemHandle+ r <- cuIpcGetMemHandle h dptr+ resultIfOk (r, IPCDevicePtr h)++{-# INLINE cuIpcGetMemHandle #-}+{# fun unsafe cuIpcGetMemHandle+ { withForeignPtr* `IPCMemHandle'+ , useDeviceHandle `DevicePtr a'+ }+ -> `Status' cToEnum #}+#endif+++-- |+-- Open an inter-process memory handle exported from another process,+-- returning a device pointer usable in the current process.+--+-- Maps memory exported by another process with 'create' into the current+-- device address space. For contexts on different devices, 'open' can+-- attempt to enable peer access if the user called+-- 'Foreign.CUDA.Driver.Context.Peer.add', and is controlled by the+-- 'LazyEnablePeerAccess' flag.+--+-- Each handle from a given device and context may only be 'open'ed by one+-- context per device per other process. Memory returned by 'open' must be+-- freed via 'close'.+--+-- Requires CUDA-4.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1ga8bd126fcff919a0c996b7640f197b79>+--+{-# INLINEABLE open #-}+open :: IPCDevicePtr a -> [IPCFlag]-> IO (DevicePtr a)+#if CUDA_VERSION < 4000+open _ _ = requireSDK 'open 4.0+#else+open !hdl !flags = resultIfOk =<< cuIpcOpenMemHandle (useIPCDevicePtr hdl) flags++{-# INLINE cuIpcOpenMemHandle #-}+{# fun unsafe cuIpcOpenMemHandle+ { alloca- `DevicePtr a' peekDeviceHandle*+ , withForeignPtr* `IPCMemHandle'+ , combineBitMasks `[IPCFlag]'+ }+ -> `Status' cToEnum #}+#endif+++-- |+-- Close and unmap memory returned by 'open'. The original allocation in+-- the exporting process as well as imported mappings in other processes+-- are unaffected.+--+-- Any resources used to enable peer access will be freed if this is the+-- last mapping using them.+--+-- Requires CUDA-4.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1gd6f5d5bcf6376c6853b64635b0157b9e>+--+{-# INLINEABLE close #-}+close :: DevicePtr a -> IO ()+#if CUDA_VERSION < 4000+close _ = requireSDK 'close 4.0+#else+close !dptr = nothingIfOk =<< cuIpcCloseMemHandle dptr++{-# INLINE cuIpcCloseMemHandle #-}+{# fun unsafe cuIpcCloseMemHandle+ { useDeviceHandle `DevicePtr a'+ }+ -> `Status' cToEnum #}+#endif+++--------------------------------------------------------------------------------+-- Internal+--------------------------------------------------------------------------------++type IPCMemHandle = ForeignPtr ()++newIPCMemHandle :: IO IPCMemHandle+newIPCMemHandle = mallocForeignPtrBytes {#sizeof CUipcMemHandle#}+
Foreign/CUDA/Driver/Marshal.chs view
@@ -2,6 +2,7 @@ {-# LANGUAGE CPP #-} {-# LANGUAGE EmptyDataDecls #-} {-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-} {-# OPTIONS_HADDOCK prune #-} -------------------------------------------------------------------------------- -- |@@ -52,7 +53,7 @@ import Foreign.CUDA.Ptr import Foreign.CUDA.Driver.Error import Foreign.CUDA.Driver.Stream ( Stream(..), defaultStream )-import Foreign.CUDA.Driver.Context ( Context(..) )+import Foreign.CUDA.Driver.Context.Base ( Context(..) ) import Foreign.CUDA.Internal.C2HS -- System@@ -86,7 +87,7 @@ -- {# enum CUmemhostalloc_option as AllocFlag { underscoreToCase }- with prefix="CU_MEMHOSTALLOC_OPTION" deriving (Eq, Show) #}+ with prefix="CU_MEMHOSTALLOC_OPTION" deriving (Eq, Show, Bounded) #} -- | -- Allocate a section of linear memory on the host which is page-locked and@@ -97,6 +98,12 @@ -- system performance may suffer. This is best used sparingly to allocate -- staging areas for data exchange. --+-- Host memory allocated in this way is automatically and immediately+-- accessible to all contexts on all devices which support unified+-- addressing.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1gdd8311286d2c2691605362c689bc64e0>+-- {-# INLINEABLE mallocHostArray #-} mallocHostArray :: Storable a => [AllocFlag] -> Int -> IO (HostPtr a) mallocHostArray !flags = doMalloc undefined@@ -115,8 +122,10 @@ -- |--- Free a section of page-locked host memory+-- Free a section of page-locked host memory. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g62e0fdbe181dab6b1c90fa1a51c7b92c>+-- {-# INLINEABLE freeHost #-} freeHost :: HostPtr a -> IO () freeHost !p = nothingIfOk =<< cuMemFreeHost p@@ -140,12 +149,17 @@ -- performance, since it reduces the amount of pageable memory available. This -- is best used sparingly to allocate staging areas for data exchange. ----- This function is not yet implemented on Mac OS X. Requires cuda-4.0.+-- This function has limited support on Mac OS X. OS 10.7 or later is+-- required. --+-- Requires CUDA-4.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1gf0a9fe11544326dabd743b7aa6b54223>+-- {-# INLINEABLE registerArray #-} registerArray :: Storable a => [AllocFlag] -> Int -> Ptr a -> IO (HostPtr a) #if CUDA_VERSION < 4000-registerArray _ _ _ = requireSDK 4.0 "registerArray"+registerArray _ _ _ = requireSDK 'registerArray 4.0 #else registerArray !flags !n = go undefined where@@ -165,12 +179,14 @@ -- | -- Unmaps the memory from the given pointer, and makes it pageable again. ----- This function is not yet implemented on Mac OS X. Requires cuda-4.0.+-- Requires CUDA-4.0. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g63f450c8125359be87b7623b1c0b2a14>+-- {-# INLINEABLE unregisterArray #-} unregisterArray :: HostPtr a -> IO (Ptr a) #if CUDA_VERSION < 4000-unregisterArray _ = requireSDK 4.0 "unregisterArray"+unregisterArray _ = requireSDK 'unregisterArray 4.0 #else unregisterArray (HostPtr !p) = do status <- cuMemHostUnregister p@@ -191,6 +207,8 @@ -- it. The memory is sufficient to hold the given number of elements of storable -- type. It is suitably aligned for any type, and is not cleared. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1gb82d2a09844a58dd9e744dc31e8aa467>+-- {-# INLINEABLE mallocArray #-} mallocArray :: Storable a => Int -> IO (DevicePtr a) mallocArray = doMalloc undefined@@ -213,7 +231,8 @@ -- or via an exception), so the pointer must not be used after this. -- -- Note that kernel launches can be asynchronous, so you may want to add a--- synchronisation point using 'sync' as part of the computation.+-- synchronisation point using 'Foreign.CUDA.Driver.Context.sync' as part+-- of the continuation. -- {-# INLINEABLE allocaArray #-} allocaArray :: Storable a => Int -> (DevicePtr a -> IO b) -> IO b@@ -221,8 +240,10 @@ -- |--- Release a section of device memory+-- Release a section of device memory. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g89b3f154e17cc89b6eea277dbdf5c93a>+-- {-# INLINEABLE free #-} free :: DevicePtr a -> IO () free !dp = nothingIfOk =<< cuMemFree dp@@ -244,17 +265,35 @@ #else {# enum CUmemAttach_flags as AttachFlag { underscoreToCase }- with prefix="CU_MEM_ATTACH_OPTION" deriving (Eq, Show) #}+ with prefix="CU_MEM_ATTACH_OPTION" deriving (Eq, Show, Bounded) #} #endif -- | -- Allocates memory that will be automatically managed by the Unified Memory--- system+-- system. The returned pointer is valid on the CPU and on all GPUs which+-- supported managed memory. All accesses to this pointer must obey the+-- Unified Memory programming model. --+-- On a multi-GPU system with peer-to-peer support, where multiple GPUs+-- support managed memory, the physical storage is created on the GPU which+-- is active at the time 'mallocManagedArray' is called. All other GPUs+-- will access the array at reduced bandwidth via peer mapping over the+-- PCIe bus. The Unified Memory system does not migrate memory between+-- GPUs.+--+-- On a multi-GPU system where multiple GPUs support managed memory, but+-- not all pairs of such GPUs have peer-to-peer support between them, the+-- physical storage is allocated in system memory (zero-copy memory) and+-- all GPUs will access the data at reduced bandwidth over the PCIe bus.+--+-- Requires CUDA-6.0+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1gb347ded34dc326af404aa02af5388a32>+-- {-# INLINEABLE mallocManagedArray #-} mallocManagedArray :: Storable a => [AttachFlag] -> Int -> IO (DevicePtr a) #if CUDA_VERSION < 6000-mallocManagedArray _ _ = requireSDK 6.0 "mallocManagedArray"+mallocManagedArray _ _ = requireSDK 'mallocManagedArray 6.0 #else mallocManagedArray !flags = doMalloc undefined where@@ -280,8 +319,10 @@ -- | -- Copy a number of elements from the device to host memory. This is a--- synchronous operation+-- synchronous operation. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g3480368ee0208a98f75019c9a8450893>+-- {-# INLINEABLE peekArray #-} peekArray :: Storable a => Int -> DevicePtr a -> Ptr a -> IO () peekArray !n !dptr !hptr = doPeek undefined dptr@@ -300,6 +341,8 @@ -- Copy memory from the device asynchronously, possibly associated with a -- particular stream. The destination host memory must be page-locked. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g56f30236c7c5247f8e061b59d3268362>+-- {-# INLINEABLE peekArrayAsync #-} peekArrayAsync :: Storable a => Int -> DevicePtr a -> HostPtr a -> Maybe Stream -> IO () peekArrayAsync !n !dptr !hptr !mst = doPeek undefined dptr@@ -320,6 +363,8 @@ -- | -- Copy a 2D array from the device to the host. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g27f885b30c34cc20a663a671dbf6fc27>+-- {-# INLINEABLE peekArray2D #-} peekArray2D :: Storable a@@ -368,6 +413,8 @@ -- associated with a particular execution stream. The destination host memory -- must be page-locked. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g4acf155faeb969d9d21f5433d3d0f274>+-- {-# INLINEABLE peekArray2DAsync #-} peekArray2DAsync :: Storable a@@ -433,8 +480,10 @@ -- -------------- -- |--- Copy a number of elements onto the device. This is a synchronous operation+-- Copy a number of elements onto the device. This is a synchronous operation. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g4d32266788c440b0220b1a9ba5795169>+-- {-# INLINEABLE pokeArray #-} pokeArray :: Storable a => Int -> Ptr a -> DevicePtr a -> IO () pokeArray !n !hptr !dptr = doPoke undefined dptr@@ -453,6 +502,8 @@ -- Copy memory onto the device asynchronously, possibly associated with a -- particular stream. The source host memory must be page-locked. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g1572263fe2597d7ba4f6964597a354a3>+-- {-# INLINEABLE pokeArrayAsync #-} pokeArrayAsync :: Storable a => Int -> HostPtr a -> DevicePtr a -> Maybe Stream -> IO () pokeArrayAsync !n !hptr !dptr !mst = dopoke undefined dptr@@ -473,6 +524,8 @@ -- | -- Copy a 2D array from the host to the device. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g27f885b30c34cc20a663a671dbf6fc27>+-- {-# INLINEABLE pokeArray2D #-} pokeArray2D :: Storable a@@ -521,6 +574,8 @@ -- associated with a particular execution stream. The source host memory must be -- page-locked. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g4acf155faeb969d9d21f5433d3d0f274>+-- {-# INLINEABLE pokeArray2DAsync #-} pokeArray2DAsync :: Storable a@@ -588,6 +643,8 @@ -- is asynchronous with respect to the host, but will never overlap with kernel -- execution. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g1725774abf8b51b91945f3336b778c8b>+-- {-# INLINEABLE copyArray #-} copyArray :: Storable a => Int -> DevicePtr a -> DevicePtr a -> IO () copyArray !n = docopy undefined@@ -608,6 +665,8 @@ -- operation is asynchronous with respect to the host, and can be asynchronous -- to other device operations by associating it with a particular stream. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g39ea09ba682b8eccc9c3e0c04319b5c8>+-- {-# INLINEABLE copyArrayAsync #-} copyArrayAsync :: Storable a => Int -> DevicePtr a -> DevicePtr a -> Maybe Stream -> IO () copyArrayAsync !n !src !dst !mst = docopy undefined src@@ -629,6 +688,8 @@ -- asynchronous with respect to the host, but will never overlap with kernel -- execution. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g27f885b30c34cc20a663a671dbf6fc27>+-- {-# INLINEABLE copyArray2D #-} copyArray2D :: Storable a@@ -678,6 +739,8 @@ -- asynchronous with respect to the host, and can be asynchronous to other -- device operations by associating it with a particular execution stream. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g4acf155faeb969d9d21f5433d3d0f274>+-- {-# INLINEABLE copyArray2DAsync #-} copyArray2DAsync :: Storable a@@ -735,6 +798,10 @@ -- the source and destination contexts. To avoid this synchronisation, use -- 'copyArrayPeerAsync' instead. --+-- Requires CUDA-4.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1ge1f5c7771544fee150ada8853c7cbf4a>+-- {-# INLINEABLE copyArrayPeer #-} copyArrayPeer :: Storable a => Int -- ^ number of array elements@@ -742,7 +809,7 @@ -> DevicePtr a -> Context -- ^ destination array and context -> IO () #if CUDA_VERSION < 4000-copyArrayPeer _ _ _ _ _ = requireSDK 4.0 "copyArrayPeer"+copyArrayPeer _ _ _ _ _ = requireSDK 'copyArrayPeer 4.0 #else copyArrayPeer !n !src !srcCtx !dst !dstCtx = go undefined src dst where@@ -764,6 +831,10 @@ -- Note that this function is asynchronous with respect to the host and all work -- in other streams and devices. --+-- Requires CUDA-4.0.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g82fcecb38018e64b98616a8ac30112f2>+-- {-# INLINEABLE copyArrayPeerAsync #-} copyArrayPeerAsync :: Storable a => Int -- ^ number of array elements@@ -772,7 +843,7 @@ -> Maybe Stream -- ^ stream to associate with -> IO () #if CUDA_VERSION < 4000-copyArrayPeerAsync _ _ _ _ _ _ = requireSDK 4.0 "copyArrayPeerAsync"+copyArrayPeerAsync _ _ _ _ _ _ = requireSDK 'copyArrayPeerAsync 4.0 #else copyArrayPeerAsync !n !src !srcCtx !dst !dstCtx !st = go undefined src dst where@@ -856,6 +927,12 @@ -- Set a number of data elements to the specified value, which may be either 8-, -- 16-, or 32-bits wide. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g6e582bf866e9e2fb014297bfaf354d7b>+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g7d805e610054392a4d11e8a8bf5eb35c>+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g983e8d8759acd1b64326317481fbf132>+-- {-# INLINEABLE memset #-} memset :: Storable a => DevicePtr a -> Int -> a -> IO () memset !dptr !n !val = case sizeOf val of@@ -890,12 +967,20 @@ -- | -- Set the number of data elements to the specified value, which may be either -- 8-, 16-, or 32-bits wide. The operation is asynchronous and may optionally be--- associated with a stream. Requires cuda-3.2.+-- associated with a stream. --+-- Requires CUDA-3.2.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1gaef08a7ccd61112f94e82f2b30d43627>+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1gf731438877dd8ec875e4c43d848c878c>+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g58229da5d30f1c0cdf667b320ec2c0f5>+-- {-# INLINEABLE memsetAsync #-} memsetAsync :: Storable a => DevicePtr a -> Int -> a -> Maybe Stream -> IO () #if CUDA_VERSION < 3020-memsetAsync _ _ _ _ = requireSDK 3.2 "memsetAsync"+memsetAsync _ _ _ _ = requireSDK 'memsetAsync 3.2 #else memsetAsync !dptr !n !val !mst = case sizeOf val of 1 -> nothingIfOk =<< cuMemsetD8Async dptr val n stream@@ -934,6 +1019,8 @@ -- -- Currently, no options are supported and this must be empty. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g57a39e5cba26af4d06be67fc77cc62f0>+-- {-# INLINEABLE getDevicePtr #-} getDevicePtr :: [AllocFlag] -> HostPtr a -> IO (DevicePtr a) getDevicePtr !flags !hp = resultIfOk =<< cuMemHostGetDevicePointer hp flags@@ -947,9 +1034,12 @@ alloca' = F.alloca useHP = castPtr . useHostPtr + -- |--- Return the base address and allocation size of the given device pointer+-- Return the base address and allocation size of the given device pointer. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g64fee5711274a2a0573a789c94d8299b>+-- {-# INLINEABLE getBasePtr #-} getBasePtr :: DevicePtr a -> IO (DevicePtr a, Int64) getBasePtr !dptr = do@@ -967,7 +1057,9 @@ -- | -- Return the amount of free and total memory respectively available to the--- current context (bytes)+-- current context (bytes).+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html#group__CUDA__MEM_1g808f555540d0143a331cc42aa98835c0> -- {-# INLINEABLE getMemInfo #-} getMemInfo :: IO (Int64, Int64)
− Foreign/CUDA/Driver/Module.chs
@@ -1,331 +0,0 @@-{-# LANGUAGE BangPatterns #-}-{-# LANGUAGE ForeignFunctionInterface #-}------------------------------------------------------------------------------------ |--- Module : Foreign.CUDA.Driver.Module--- Copyright : [2009..2014] Trevor L. McDonell--- License : BSD------ Module management for low-level driver interface--------------------------------------------------------------------------------------module Foreign.CUDA.Driver.Module (-- -- * Module Management- Module, JITOption(..), JITTarget(..), JITResult(..), JITFallback(..),-- -- ** Querying module inhabitants- getFun, getPtr, getTex,-- -- ** Loading and unloading modules- loadFile,- loadData, loadDataFromPtr,- loadDataEx, loadDataFromPtrEx,- unload--) where--#include "cbits/stubs.h"-{# context lib="cuda" #}---- Friends-import Foreign.CUDA.Analysis.Device-import Foreign.CUDA.Ptr-import Foreign.CUDA.Driver.Error-import Foreign.CUDA.Driver.Exec-import Foreign.CUDA.Driver.Marshal ( peekDeviceHandle )-import Foreign.CUDA.Driver.Texture-import Foreign.CUDA.Internal.C2HS---- System-import Foreign-import Foreign.C-import Unsafe.Coerce--import Control.Monad ( liftM )-import Control.Exception ( throwIO )-import Data.Maybe ( mapMaybe )-import Data.ByteString.Char8 ( ByteString )-import qualified Data.ByteString.Char8 as B-import qualified Data.ByteString.Internal as B--#if CUDA_VERSION < 5050-import Debug.Trace ( trace )-#endif------------------------------------------------------------------------------------- Data Types------------------------------------------------------------------------------------- |--- A reference to a Module object, containing collections of device functions----newtype Module = Module { useModule :: {# type CUmodule #}}- deriving (Eq, Show)----- |--- Just-in-time compilation options----data JITOption- = MaxRegisters !Int -- ^ maximum number of registers per thread- | ThreadsPerBlock !Int -- ^ number of threads per block to target for- | OptimisationLevel !Int -- ^ level of optimisation to apply (1-4, default 4)- | Target !Compute -- ^ compilation target, otherwise determined from context- | FallbackStrategy !JITFallback -- ^ fallback strategy if matching cubin not found- | GenerateDebugInfo -- ^ generate debug info (-g) (requires cuda >= 5.5)- | GenerateLineInfo -- ^ generate line number information (-lineinfo) (requires cuda >= 5.5)- | Verbose -- ^ verbose log messages (requires cuda >= 5.5)- deriving (Show)---- |--- Results of online compilation----data JITResult = JITResult- {- jitTime :: !Float, -- ^ milliseconds spent compiling PTX- jitInfoLog :: !ByteString, -- ^ information about PTX assembly- jitModule :: !Module -- ^ compilation error log or compiled module- }- deriving (Show)---{# enum CUjit_option as JITOptionInternal- { }- with prefix="CU" deriving (Eq, Show) #}--{# enum CUjit_target as JITTarget- { underscoreToCase }- with prefix="CU_TARGET" deriving (Eq, Show) #}--{# enum CUjit_fallback as JITFallback- { underscoreToCase- , CU_PREFER_PTX as PTX }- with prefix="CU_PREFER" deriving (Eq, Show) #}-------------------------------------------------------------------------------------- Module management------------------------------------------------------------------------------------- |--- Returns a function handle----{-# INLINEABLE getFun #-}-getFun :: Module -> String -> IO Fun-getFun !mdl !fn = resultIfFound "function" fn =<< cuModuleGetFunction mdl fn--{-# INLINE cuModuleGetFunction #-}-{# fun unsafe cuModuleGetFunction- { alloca- `Fun' peekFun*- , useModule `Module'- , withCString* `String' } -> `Status' cToEnum #}- where peekFun = liftM Fun . peek----- |--- Return a global pointer, and size of the global (in bytes)----{-# INLINEABLE getPtr #-}-getPtr :: Module -> String -> IO (DevicePtr a, Int)-getPtr !mdl !name = do- (!status,!dptr,!bytes) <- cuModuleGetGlobal mdl name- resultIfFound "global" name (status,(dptr,bytes))--{-# INLINE cuModuleGetGlobal #-}-{# fun unsafe cuModuleGetGlobal- { alloca- `DevicePtr a' peekDeviceHandle*- , alloca- `Int' peekIntConv*- , useModule `Module'- , withCString* `String' } -> `Status' cToEnum #}----- |--- Return a handle to a texture reference----{-# INLINEABLE getTex #-}-getTex :: Module -> String -> IO Texture-getTex !mdl !name = resultIfFound "texture" name =<< cuModuleGetTexRef mdl name--{-# INLINE cuModuleGetTexRef #-}-{# fun unsafe cuModuleGetTexRef- { alloca- `Texture' peekTex*- , useModule `Module'- , withCString* `String' } -> `Status' cToEnum #}----- |--- Load the contents of the specified file (either a ptx or cubin file) to--- create a new module, and load that module into the current context.----{-# INLINEABLE loadFile #-}-loadFile :: FilePath -> IO Module-loadFile !ptx = resultIfOk =<< cuModuleLoad ptx--{-# INLINE cuModuleLoad #-}-{# fun unsafe cuModuleLoad- { alloca- `Module' peekMod*- , withCString* `FilePath' } -> `Status' cToEnum #}----- |--- Load the contents of the given image into a new module, and load that module--- into the current context. The image is (typically) the contents of a cubin or--- PTX file.------ Note that the 'ByteString' will be copied into a temporary staging area so--- that it can be passed to C.----{-# INLINEABLE loadData #-}-loadData :: ByteString -> IO Module-loadData !img =- B.useAsCString img (\p -> loadDataFromPtr (castPtr p))---- |--- As 'loadData', but read the image data from the given pointer. The image is a--- NULL-terminated sequence of bytes.----{-# INLINEABLE loadDataFromPtr #-}-loadDataFromPtr :: Ptr Word8 -> IO Module-loadDataFromPtr !img = resultIfOk =<< cuModuleLoadData img--{-# INLINE cuModuleLoadData #-}-{# fun unsafe cuModuleLoadData- { alloca- `Module' peekMod*- , castPtr `Ptr Word8' } -> ` Status' cToEnum #}----- |--- Load the contents of the given image into a module with online compiler--- options, and load the module into the current context. The image is--- (typically) the contents of a cubin or PTX file. The actual attributes of the--- compiled kernel can be probed using 'requires'.------ Note that the 'ByteString' will be copied into a temporary staging area so--- that it can be passed to C.----{-# INLINEABLE loadDataEx #-}-loadDataEx :: ByteString -> [JITOption] -> IO JITResult-loadDataEx !img !options =- B.useAsCString img (\p -> loadDataFromPtrEx (castPtr p) options)---- |--- As 'loadDataEx', but read the image data from the given pointer. The image is--- a NULL-terminated sequence of bytes.----{-# INLINEABLE loadDataFromPtrEx #-}-loadDataFromPtrEx :: Ptr Word8 -> [JITOption] -> IO JITResult-loadDataFromPtrEx !img !options = do- fp_ilog <- B.mallocByteString logSize-- allocaArray logSize $ \p_elog -> do- withForeignPtr fp_ilog $ \p_ilog -> do-- let (opt,val) = unzip $- [ (JIT_WALL_TIME, 0) -- must be first, this is extracted below- , (JIT_INFO_LOG_BUFFER_SIZE_BYTES, logSize)- , (JIT_ERROR_LOG_BUFFER_SIZE_BYTES, logSize)- , (JIT_INFO_LOG_BUFFER, unsafeCoerce (p_ilog :: CString))- , (JIT_ERROR_LOG_BUFFER, unsafeCoerce (p_elog :: CString)) ] ++ mapMaybe unpack options-- withArray (map cFromEnum opt) $ \p_opts -> do- withArray (map unsafeCoerce val) $ \p_vals -> do-- (s,mdl) <- cuModuleLoadDataEx img (length opt) p_opts p_vals-- case s of- Success -> do- time <- peek (castPtr p_vals)- infoLog <- B.fromForeignPtr (castForeignPtr fp_ilog) 0 `fmap` c_strnlen p_ilog logSize- return $! JITResult time infoLog mdl-- _ -> do- errLog <- peekCString p_elog- cudaError (unlines [describe s, errLog])-- where- logSize = 2048-- unpack (MaxRegisters x) = Just (JIT_MAX_REGISTERS, x)- unpack (ThreadsPerBlock x) = Just (JIT_THREADS_PER_BLOCK, x)- unpack (OptimisationLevel x) = Just (JIT_OPTIMIZATION_LEVEL, x)- unpack (Target x) = Just (JIT_TARGET, jitTargetOfCompute x)- unpack (FallbackStrategy x) = Just (JIT_FALLBACK_STRATEGY, fromEnum x)-#if CUDA_VERSION >= 5050- unpack GenerateDebugInfo = Just (JIT_GENERATE_DEBUG_INFO, fromEnum True)- unpack GenerateLineInfo = Just (JIT_GENERATE_LINE_INFO, fromEnum True)- unpack Verbose = Just (JIT_LOG_VERBOSE, fromEnum True)-#else- unpack x = trace ("Warning: JITOption '" ++ show x ++ "' requires at least cuda-5.5") Nothing-#endif-- jitTargetOfCompute (Compute x y)- = fromEnum- $ case (x,y) of- (1,0) -> Compute10- (1,1) -> Compute11- (1,2) -> Compute12- (1,3) -> Compute13- (2,0) -> Compute20- (2,1) -> Compute21- (3,0) -> Compute30- (3,5) -> Compute35- _ -> error ("Unknown JIT Target for Compute " ++ show (Compute x y))---{-# INLINE cuModuleLoadDataEx #-}-{# fun unsafe cuModuleLoadDataEx- { alloca- `Module' peekMod*- , castPtr `Ptr Word8'- , `Int'- , id `Ptr CInt'- , id `Ptr (Ptr ())' } -> `Status' cToEnum #}----- |--- Unload a module from the current context----{-# INLINEABLE unload #-}-unload :: Module -> IO ()-unload !m = nothingIfOk =<< cuModuleUnload m--{-# INLINE cuModuleUnload #-}-{# fun unsafe cuModuleUnload- { useModule `Module' } -> `Status' cToEnum #}-------------------------------------------------------------------------------------- Internal-----------------------------------------------------------------------------------{-# INLINE resultIfFound #-}-resultIfFound :: String -> String -> (Status, a) -> IO a-resultIfFound kind name (!status,!result) =- case status of- Success -> return result- NotFound -> cudaError (kind ++ ' ' : describe status ++ ": " ++ name)- _ -> throwIO (ExitCode status)--{-# INLINE peekMod #-}-peekMod :: Ptr {# type CUmodule #} -> IO Module-peekMod = liftM Module . peek--{-# INLINE c_strnlen' #-}-#if defined(WIN32)-c_strnlen' :: CString -> CSize -> IO CSize-c_strnlen' str size = do- str' <- peekCStringLen (str, fromIntegral size)- return $ stringLen 0 str'- where stringLen acc [] = acc- stringLen acc ('\0':xs) = acc- stringLen acc (_:xs) = stringLen (acc+1) xs-#else-foreign import ccall unsafe "string.h strnlen" c_strnlen'- :: CString -> CSize -> IO CSize-#endif--{-# INLINE c_strnlen #-}-c_strnlen :: CString -> Int -> IO Int-c_strnlen str maxlen = cIntConv `fmap` c_strnlen' str (cIntConv maxlen)-
+ Foreign/CUDA/Driver/Module.hs view
@@ -0,0 +1,20 @@+--------------------------------------------------------------------------------+-- |+-- Module : Foreign.CUDA.Driver.Module+-- Copyright : [2009..2014] Trevor L. McDonell+-- License : BSD+--+-- Module management for low-level driver interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Module (++ module Foreign.CUDA.Driver.Module.Base,+ module Foreign.CUDA.Driver.Module.Query,++) where++import Foreign.CUDA.Driver.Module.Base hiding ( JITOptionInternal(..), useModule, jitOptionUnpack, jitTargetOfCompute )+import Foreign.CUDA.Driver.Module.Query+
+ Foreign/CUDA/Driver/Module/Base.chs view
@@ -0,0 +1,311 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE ForeignFunctionInterface #-}+{-# OPTIONS_HADDOCK prune #-}+--------------------------------------------------------------------------------+-- |+-- Module : Foreign.CUDA.Driver.Module.Base+-- Copyright : [2009..2014] Trevor L. McDonell+-- License : BSD+--+-- Module loading for low-level driver interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Module.Base (++ -- * Module Management+ Module(..),+ JITOption(..), JITTarget(..), JITResult(..), JITFallback(..), JITInputType(..),+ JITOptionInternal(..),++ -- ** Loading and unloading modules+ loadFile,+ loadData, loadDataFromPtr,+ loadDataEx, loadDataFromPtrEx,+ unload,++ -- Internal+ jitOptionUnpack, jitTargetOfCompute,++) where++#include "cbits/stubs.h"+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Analysis.Device+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Internal.C2HS++-- System+import Foreign+import Foreign.C+import Unsafe.Coerce++import Control.Monad ( liftM )+import Data.ByteString.Char8 ( ByteString )+import qualified Data.ByteString.Char8 as B+import qualified Data.ByteString.Internal as B+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- A reference to a Module object, containing collections of device functions+--+newtype Module = Module { useModule :: {# type CUmodule #}}+ deriving (Eq, Show)++-- |+-- Just-in-time compilation and linking options+--+data JITOption+ = MaxRegisters !Int -- ^ maximum number of registers per thread+ | ThreadsPerBlock !Int -- ^ number of threads per block to target for+ | OptimisationLevel !Int -- ^ level of optimisation to apply (1-4, default 4)+ | Target !Compute -- ^ compilation target, otherwise determined from context+ | FallbackStrategy !JITFallback -- ^ fallback strategy if matching cubin not found+ | GenerateDebugInfo -- ^ generate debug info (-g) (requires cuda >= 5.5)+ | GenerateLineInfo -- ^ generate line number information (-lineinfo) (requires cuda >= 5.5)+ | Verbose -- ^ verbose log messages (requires cuda >= 5.5)+ deriving (Show)++-- |+-- Results of online compilation+--+data JITResult = JITResult+ {+ jitTime :: !Float, -- ^ milliseconds spent compiling PTX+ jitInfoLog :: !ByteString, -- ^ information about PTX assembly+ jitModule :: !Module -- ^ the compiled module+ }+ deriving (Show)+++-- |+-- Online compilation target architecture+--+{# enum CUjit_target as JITTarget+ { underscoreToCase }+ with prefix="CU_TARGET" deriving (Eq, Show) #}++-- |+-- Online compilation fallback strategy+--+{# enum CUjit_fallback as JITFallback+ { underscoreToCase+ , CU_PREFER_PTX as PreferPTX }+ with prefix="CU" deriving (Eq, Show) #}++-- |+-- Device code formats that can be used for online linking+--+{# enum CUjitInputType as JITInputType+ { underscoreToCase+ , CU_JIT_INPUT_PTX as PTX }+ with prefix="CU_JIT_INPUT" deriving (Eq, Show) #}++{# enum CUjit_option as JITOptionInternal+ { }+ with prefix="CU" deriving (Eq, Show) #}+++--------------------------------------------------------------------------------+-- Module management+--------------------------------------------------------------------------------++-- |+-- Load the contents of the specified file (either a ptx or cubin file) to+-- create a new module, and load that module into the current context.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE_1g366093bd269dafd0af21f1c7d18115d3>+--+{-# INLINEABLE loadFile #-}+loadFile :: FilePath -> IO Module+loadFile !ptx = resultIfOk =<< cuModuleLoad ptx++{-# INLINE cuModuleLoad #-}+{# fun unsafe cuModuleLoad+ { alloca- `Module' peekMod*+ , withCString* `FilePath' } -> `Status' cToEnum #}+++-- |+-- Load the contents of the given image into a new module, and load that module+-- into the current context. The image is (typically) the contents of a cubin or+-- PTX file.+--+-- Note that the 'ByteString' will be copied into a temporary staging area so+-- that it can be passed to C.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE_1g04ce266ce03720f479eab76136b90c0b>+--+{-# INLINEABLE loadData #-}+loadData :: ByteString -> IO Module+loadData !img =+ B.useAsCString img (\p -> loadDataFromPtr (castPtr p))++-- |+-- As 'loadData', but read the image data from the given pointer. The image is a+-- NULL-terminated sequence of bytes.+--+{-# INLINEABLE loadDataFromPtr #-}+loadDataFromPtr :: Ptr Word8 -> IO Module+loadDataFromPtr !img = resultIfOk =<< cuModuleLoadData img++{-# INLINE cuModuleLoadData #-}+{# fun unsafe cuModuleLoadData+ { alloca- `Module' peekMod*+ , castPtr `Ptr Word8' } -> ` Status' cToEnum #}+++-- |+-- Load the contents of the given image into a module with online compiler+-- options, and load the module into the current context. The image is+-- (typically) the contents of a cubin or PTX file. The actual attributes of the+-- compiled kernel can be probed using 'requires'.+--+-- Note that the 'ByteString' will be copied into a temporary staging area so+-- that it can be passed to C.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE_1g9e8047e9dbf725f0cd7cafd18bfd4d12>+--+{-# INLINEABLE loadDataEx #-}+loadDataEx :: ByteString -> [JITOption] -> IO JITResult+loadDataEx !img !options =+ B.useAsCString img (\p -> loadDataFromPtrEx (castPtr p) options)++-- |+-- As 'loadDataEx', but read the image data from the given pointer. The image is+-- a NULL-terminated sequence of bytes.+--+{-# INLINEABLE loadDataFromPtrEx #-}+loadDataFromPtrEx :: Ptr Word8 -> [JITOption] -> IO JITResult+loadDataFromPtrEx !img !options = do+ let logSize = 2048++ fp_ilog <- B.mallocByteString logSize++ allocaArray logSize $ \p_elog -> do+ withForeignPtr fp_ilog $ \p_ilog -> do++ let (opt,val) = unzip $+ [ (JIT_WALL_TIME, 0) -- must be first, this is extracted below+ , (JIT_INFO_LOG_BUFFER_SIZE_BYTES, logSize)+ , (JIT_ERROR_LOG_BUFFER_SIZE_BYTES, logSize)+ , (JIT_INFO_LOG_BUFFER, unsafeCoerce (p_ilog :: CString))+ , (JIT_ERROR_LOG_BUFFER, unsafeCoerce (p_elog :: CString))+ ]+ +++ map jitOptionUnpack options++ withArrayLen (map cFromEnum opt) $ \i p_opts -> do+ withArray (map unsafeCoerce val) $ \ p_vals -> do++ (s,mdl) <- cuModuleLoadDataEx img i p_opts p_vals++ case s of+ Success -> do+ time <- peek (castPtr p_vals)+ infoLog <- B.fromForeignPtr (castForeignPtr fp_ilog) 0 `fmap` c_strnlen p_ilog logSize+ return $! JITResult time infoLog mdl++ _ -> do+ errLog <- peekCString p_elog+ cudaError (unlines [describe s, errLog])+++{-# INLINE cuModuleLoadDataEx #-}+{# fun unsafe cuModuleLoadDataEx+ { alloca- `Module' peekMod*+ , castPtr `Ptr Word8'+ , `Int'+ , id `Ptr CInt'+ , id `Ptr (Ptr ())' } -> `Status' cToEnum #}+++-- |+-- Unload a module from the current context.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE_1g8ea3d716524369de3763104ced4ea57b>+--+{-# INLINEABLE unload #-}+unload :: Module -> IO ()+unload !m = nothingIfOk =<< cuModuleUnload m++{-# INLINE cuModuleUnload #-}+{# fun unsafe cuModuleUnload+ { useModule `Module' } -> `Status' cToEnum #}+++--------------------------------------------------------------------------------+-- Internal+--------------------------------------------------------------------------------++{-# INLINE peekMod #-}+peekMod :: Ptr {# type CUmodule #} -> IO Module+peekMod = liftM Module . peek+++{-# INLINE jitOptionUnpack #-}+jitOptionUnpack :: JITOption -> (JITOptionInternal, Int)+jitOptionUnpack (MaxRegisters x) = (JIT_MAX_REGISTERS, x)+jitOptionUnpack (ThreadsPerBlock x) = (JIT_THREADS_PER_BLOCK, x)+jitOptionUnpack (OptimisationLevel x) = (JIT_OPTIMIZATION_LEVEL, x)+jitOptionUnpack (Target x) = (JIT_TARGET, fromEnum (jitTargetOfCompute x))+jitOptionUnpack (FallbackStrategy x) = (JIT_FALLBACK_STRATEGY, fromEnum x)+#if CUDA_VERSION >= 5050+jitOptionUnpack GenerateDebugInfo = (JIT_GENERATE_DEBUG_INFO, fromEnum True)+jitOptionUnpack GenerateLineInfo = (JIT_GENERATE_LINE_INFO, fromEnum True)+jitOptionUnpack Verbose = (JIT_LOG_VERBOSE, fromEnum True)+#else+jitOptionUnpack GenerateDebugInfo = requireSDK 'GenerateDebugInfo 5.5+jitOptionUnpack GenerateLineInfo = requireSDK 'GenerateLineInfo 5.5+jitOptionUnpack Verbose = requireSDK 'Verbose 5.5+#endif+++{-# INLINE jitTargetOfCompute #-}+jitTargetOfCompute :: Compute -> JITTarget+jitTargetOfCompute (Compute 1 0) = Compute10+jitTargetOfCompute (Compute 1 1) = Compute11+jitTargetOfCompute (Compute 1 2) = Compute12+jitTargetOfCompute (Compute 1 3) = Compute13+jitTargetOfCompute (Compute 2 0) = Compute20+jitTargetOfCompute (Compute 2 1) = Compute21+jitTargetOfCompute (Compute 3 0) = Compute30+jitTargetOfCompute (Compute 3 5) = Compute35+#if CUDA_VERSION >= 6000+jitTargetOfCompute (Compute 3 2) = Compute32+jitTargetOfCompute (Compute 5 0) = Compute50+#endif+#if CUDA_VERSION >= 6050+jitTargetOfCompute (Compute 3 7) = Compute37+#endif+#if CUDA_VERSION >= 7000+jitTargetOfCompute (Compute 5 2) = Compute52+#endif+jitTargetOfCompute compute = error ("Unknown JIT Target for Compute " ++ show compute)+++#if defined(WIN32)+{-# INLINE c_strnlen' #-}+c_strnlen' :: CString -> CSize -> IO CSize+c_strnlen' str size = do+ str' <- peekCStringLen (str, fromIntegral size)+ return $ stringLen 0 str'+ where+ stringLen acc [] = acc+ stringLen acc ('\0':_) = acc+ stringLen acc (_:xs) = stringLen (acc+1) xs+#else+foreign import ccall unsafe "string.h strnlen" c_strnlen'+ :: CString -> CSize -> IO CSize+#endif++{-# INLINE c_strnlen #-}+c_strnlen :: CString -> Int -> IO Int+c_strnlen str maxlen = cIntConv `fmap` c_strnlen' str (cIntConv maxlen)+
+ Foreign/CUDA/Driver/Module/Link.chs view
@@ -0,0 +1,219 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-}+--------------------------------------------------------------------------------+-- |+-- Module : Foreign.CUDA.Driver.Module.Link+-- Copyright : [2009..2014] Trevor L. McDonell+-- License : BSD+--+-- Module linking for low-level driver interface+--+-- Since CUDA-5.5+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Module.Link (++ -- ** JIT module linking+ LinkState, JITOption(..), JITInputType(..),++ create, destroy, complete,+ addFile,+ addData, addDataFromPtr,++) where++#include "cbits/stubs.h"+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Driver.Module.Base+import Foreign.CUDA.Internal.C2HS++-- System+import Control.Monad ( liftM )+import Foreign+import Foreign.C+import Unsafe.Coerce++import Data.ByteString.Char8 ( ByteString )+import qualified Data.ByteString.Char8 as B+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- A pending JIT linker state+--+newtype LinkState = LinkState { useLinkState :: {# type CUlinkState #} }+ deriving (Show)+++--------------------------------------------------------------------------------+-- JIT linking+--------------------------------------------------------------------------------++-- |+-- Create a pending JIT linker invocation. The returned 'LinkState' should+-- be 'destroy'ed once no longer needed. The device code machine size will+-- match the calling application.+--+-- Requires CUDA-5.5.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE_1g86ca4052a2fab369cb943523908aa80d>+--+{-# INLINEABLE create #-}+create :: [JITOption] -> IO LinkState+#if CUDA_VERSION < 5050+create _ = requireSDK 'create 5.5+#else+create !options =+ let (opt,val) = unzip $ map jitOptionUnpack options+ in+ withArray (map cFromEnum opt) $ \p_opts ->+ withArray (map unsafeCoerce val) $ \p_vals ->+ resultIfOk =<< cuLinkCreate (length opt) p_opts p_vals++{-# INLINE cuLinkCreate #-}+{# fun unsafe cuLinkCreate+ { `Int'+ , id `Ptr CInt'+ , id `Ptr (Ptr ())'+ , alloca- `LinkState' peekLS* } -> `Status' cToEnum #}+ where+ peekLS = liftM LinkState . peek+#endif+++-- |+-- Destroy the state of a JIT linker invocation.+--+-- Requires CUDA-5.5.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE_1g01b7ae2a34047b05716969af245ce2d9>+--+{-# INLINEABLE destroy #-}+destroy :: LinkState -> IO ()+#if CUDA_VERSION < 5050+destroy _ = requireSDK 'destroy 5.5+#else+destroy !s = nothingIfOk =<< cuLinkDestroy s++{-# INLINE cuLinkDestroy #-}+{# fun unsafe cuLinkDestroy+ { useLinkState `LinkState' } -> `Status' cToEnum #}+#endif+++-- |+-- Complete a pending linker invocation and load the current module. The+-- link state will be destroyed.+--+-- Requires CUDA-5.5.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE_1g818fcd84a4150a997c0bba76fef4e716>+--+{-# INLINEABLE complete #-}+complete :: LinkState -> IO Module+#if CUDA_VERSION < 5050+complete _ = requireSDK 'complete 5.5+#else+complete !ls = do+ cubin <- resultIfOk =<< cuLinkComplete ls nullPtr+ mdl <- loadDataFromPtr (castPtr cubin)+ destroy ls+ return mdl++{-# INLINE cuLinkComplete #-}+{# fun unsafe cuLinkComplete+ { useLinkState `LinkState'+ , alloca- `Ptr ()' peek*+ , castPtr `Ptr Int'+ }+ -> `Status' cToEnum #}+#endif+++-- |+-- Add an input file to a pending linker invocation.+--+-- Requires CUDA-5.5.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE_1g1224c0fd48d4a683f3ce19997f200a8c>+--+{-# INLINEABLE addFile #-}+addFile :: LinkState -> FilePath -> JITInputType -> [JITOption] -> IO ()+#if CUDA_VERSION < 5050+addFile _ _ _ _ = requireSDK 'addile 5.5+#else+addFile !ls !fp !t !options =+ let (opt,val) = unzip $ map jitOptionUnpack options+ in+ withArrayLen (map cFromEnum opt) $ \i p_opts ->+ withArray (map unsafeCoerce val) $ \ p_vals ->+ nothingIfOk =<< cuLinkAddFile ls t fp i p_opts p_vals++{-# INLINE cuLinkAddFile #-}+{# fun unsafe cuLinkAddFile+ { useLinkState `LinkState'+ , cFromEnum `JITInputType'+ , withCString* `FilePath'+ , `Int'+ , id `Ptr CInt'+ , id `Ptr (Ptr ())'+ }+ -> `Status' cToEnum #}+#endif+++-- |+-- Add an input to a pending linker invocation.+--+-- Requires CUDA-5.5.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE_1g3ebcd2ccb772ba9c120937a2d2831b77>+--+{-# INLINEABLE addData #-}+addData :: LinkState -> ByteString -> JITInputType -> [JITOption] -> IO ()+#if CUDA_VERSION < 5050+addData _ _ _ = requireSDK 'addData 5.5+#else+addData !ls !img !k !options =+ B.useAsCStringLen img (\(p, n) -> addDataFromPtr ls n (castPtr p) k options)+#endif++-- |+-- As 'addData', but read the specified number of bytes of image data from+-- the given pointer.+--+{-# INLINEABLE addDataFromPtr #-}+addDataFromPtr :: LinkState -> Int -> Ptr Word8 -> JITInputType -> [JITOption] -> IO ()+#if CUDA_VERSION < 5050+addDataFromPtr _ _ _ _ = requireSDK 'addDataFromPtr 5.5+#else+addDataFromPtr !ls !n !img !t !options =+ let (opt,val) = unzip $ map jitOptionUnpack options+ in+ withArrayLen (map cFromEnum opt) $ \i p_opts ->+ withArray (map unsafeCoerce val) $ \ p_vals ->+ nothingIfOk =<< cuLinkAddData ls t img n "<unknown>" i p_opts p_vals++{-# INLINE cuLinkAddData #-}+{# fun unsafe cuLinkAddData+ { useLinkState `LinkState'+ , cFromEnum `JITInputType'+ , castPtr `Ptr Word8'+ , `Int'+ , `String'+ , `Int'+ , id `Ptr CInt'+ , id `Ptr (Ptr ())'+ }+ -> `Status' cToEnum #}+#endif+
+ Foreign/CUDA/Driver/Module/Query.chs view
@@ -0,0 +1,108 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module : Foreign.CUDA.Driver.Module.Query+-- Copyright : [2009..2014] Trevor L. McDonell+-- License : BSD+--+-- Querying module attributes for low-level driver interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Module.Query (++ -- ** Querying module inhabitants+ getFun, getPtr, getTex,++) where++#include "cbits/stubs.h"+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Driver.Exec+import Foreign.CUDA.Driver.Marshal ( peekDeviceHandle )+import Foreign.CUDA.Driver.Module.Base+import Foreign.CUDA.Driver.Texture+import Foreign.CUDA.Internal.C2HS+import Foreign.CUDA.Ptr++-- System+import Foreign+import Foreign.C+import Control.Exception ( throwIO )+import Control.Monad ( liftM )+++--------------------------------------------------------------------------------+-- Querying module attributes+--------------------------------------------------------------------------------++-- |+-- Returns a function handle.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE_1ga52be009b0d4045811b30c965e1cb2cf>+--+{-# INLINEABLE getFun #-}+getFun :: Module -> String -> IO Fun+getFun !mdl !fn = resultIfFound "function" fn =<< cuModuleGetFunction mdl fn++{-# INLINE cuModuleGetFunction #-}+{# fun unsafe cuModuleGetFunction+ { alloca- `Fun' peekFun*+ , useModule `Module'+ , withCString* `String' } -> `Status' cToEnum #}+ where peekFun = liftM Fun . peek+++-- |+-- Return a global pointer, and size of the global (in bytes).+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE_1gf3e43672e26073b1081476dbf47a86ab>+--+{-# INLINEABLE getPtr #-}+getPtr :: Module -> String -> IO (DevicePtr a, Int)+getPtr !mdl !name = do+ (!status,!dptr,!bytes) <- cuModuleGetGlobal mdl name+ resultIfFound "global" name (status,(dptr,bytes))++{-# INLINE cuModuleGetGlobal #-}+{# fun unsafe cuModuleGetGlobal+ { alloca- `DevicePtr a' peekDeviceHandle*+ , alloca- `Int' peekIntConv*+ , useModule `Module'+ , withCString* `String' } -> `Status' cToEnum #}+++-- |+-- Return a handle to a texture reference. This texture reference handle+-- should not be destroyed, as the texture will be destroyed automatically+-- when the module is unloaded.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE_1g9607dcbf911c16420d5264273f2b5608>+--+{-# INLINEABLE getTex #-}+getTex :: Module -> String -> IO Texture+getTex !mdl !name = resultIfFound "texture" name =<< cuModuleGetTexRef mdl name++{-# INLINE cuModuleGetTexRef #-}+{# fun unsafe cuModuleGetTexRef+ { alloca- `Texture' peekTex*+ , useModule `Module'+ , withCString* `String' } -> `Status' cToEnum #}+++--------------------------------------------------------------------------------+-- Internal+--------------------------------------------------------------------------------++{-# INLINE resultIfFound #-}+resultIfFound :: String -> String -> (Status, a) -> IO a+resultIfFound kind name (!status,!result) =+ case status of+ Success -> return result+ NotFound -> cudaError (kind ++ ' ' : describe status ++ ": " ++ name)+ _ -> throwIO (ExitCode status)+
Foreign/CUDA/Driver/Stream.chs view
@@ -1,6 +1,7 @@ {-# LANGUAGE BangPatterns #-} {-# LANGUAGE EmptyDataDecls #-} {-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-} -------------------------------------------------------------------------------- -- | -- Module : Foreign.CUDA.Driver.Stream@@ -15,8 +16,10 @@ -- * Stream Management Stream(..), StreamFlag,- create, destroy, finished, block, defaultStream+ create, createWithPriority, destroy, finished, block, getPriority, + defaultStream,+ ) where #include "cbits/stubs.h"@@ -39,8 +42,10 @@ -------------------------------------------------------------------------------- -- |--- Create a new stream+-- Create a new stream. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__STREAM.html#group__CUDA__STREAM_1ga581f0c5833e21ded8b5a56594e243f4>+-- {-# INLINEABLE create #-} create :: [StreamFlag] -> IO Stream create !flags = resultIfOk =<< cuStreamCreate flags@@ -52,9 +57,50 @@ where peekStream = liftM Stream . peek + -- |--- Destroy a stream+-- Create a stream with the given priority. Work submitted to+-- a higher-priority stream may preempt work already executing in a lower+-- priority stream. --+-- The convention is that lower numbers represent higher priorities. The+-- default priority is zero. The range of meaningful numeric priorities can+-- be queried using 'Foreign.CUDA.Driver.Context.Config.getStreamPriorityRange'.+-- If the specified priority is outside the supported numerical range, it+-- will automatically be clamped to the highest or lowest number in the+-- range.+--+-- Requires CUDA-5.5.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__STREAM.html#group__CUDA__STREAM_1g95c1a8c7c3dacb13091692dd9c7f7471>+--+{-# INLINEABLE createWithPriority #-}+createWithPriority :: StreamPriority -> [StreamFlag] -> IO Stream+#if CUDA_VERSION < 5050+createWithPriority _ _ = requireSDK 'createWithPriority 5.5+#else+createWithPriority !priority !flags = resultIfOk =<< cuStreamCreateWithPriority flags priority++{-# INLINE cuStreamCreateWithPriority #-}+{# fun unsafe cuStreamCreateWithPriority+ { alloca- `Stream' peekStream*+ , combineBitMasks `[StreamFlag]'+ , cIntConv `StreamPriority'+ }+ -> `Status' cToEnum #}+ where+ peekStream = liftM Stream . peek+#endif+++-- |+-- Destroy a stream. If the device is still doing work in the stream when+-- 'destroy' is called, the function returns immediately and the resources+-- associated with the stream will be released automatically once the+-- device has completed all work.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__STREAM.html#group__CUDA__STREAM_1g244c8833de4596bcd31a06cdf21ee758>+-- {-# INLINEABLE destroy #-} destroy :: Stream -> IO () destroy !st = nothingIfOk =<< cuStreamDestroy st@@ -65,8 +111,10 @@ -- |--- Check if all operations in the stream have completed+-- Check if all operations in the stream have completed. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__STREAM.html#group__CUDA__STREAM_1g1b0d24bbe97fa68e4bc511fb6adfeb0b>+-- {-# INLINEABLE finished #-} finished :: Stream -> IO Bool finished !st =@@ -82,8 +130,10 @@ -- |--- Wait until the device has completed all operations in the Stream+-- Wait until the device has completed all operations in the Stream. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__STREAM.html#group__CUDA__STREAM_1g15e49dd91ec15991eb7c0a741beb7dad>+-- {-# INLINEABLE block #-} block :: Stream -> IO () block !st = nothingIfOk =<< cuStreamSynchronize st@@ -91,4 +141,27 @@ {-# INLINE cuStreamSynchronize #-} {# fun cuStreamSynchronize { useStream `Stream' } -> `Status' cToEnum #}+++-- |+-- Query the priority of a stream.+--+-- Requires CUDA-5.5.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__STREAM.html#group__CUDA__STREAM_1g5bd5cb26915a2ecf1921807339488484>+--+{-# INLINEABLE getPriority #-}+getPriority :: Stream -> IO StreamPriority+#if CUDA_VERSION < 5050+getPriority _ = requireSDK 'getPriority 5.5+#else+getPriority !st = resultIfOk =<< cuStreamGetPriority st++{-# INLINE cuStreamGetPriority #-}+{# fun unsafe cuStreamGetPriority+ { useStream `Stream'+ , alloca- `StreamPriority' peekIntConv*+ }+ -> `Status' cToEnum #}+#endif
Foreign/CUDA/Driver/Texture.chs view
@@ -15,11 +15,13 @@ -- * Texture Reference Management Texture(..), Format(..), AddressMode(..), FilterMode(..), ReadMode(..),- create, destroy, bind, bind2D, getAddressMode, getFilterMode, getFormat, setAddressMode, setFilterMode, setFormat, setReadMode, + -- Deprecated+ create, destroy,+ -- Internal peekTex @@ -114,6 +116,8 @@ -- reference functions are used to specify the format and interpretation to be -- used when the memory is read through this reference. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TEXREF__DEPRECATED.html#group__CUDA__TEXREF__DEPRECATED_1g0084fabe2c6d28ffcf9d9f5c7164f16c>+-- {-# INLINEABLE create #-} create :: IO Texture create = resultIfOk =<< cuTexRefCreate@@ -124,8 +128,10 @@ -- |--- Destroy a texture reference+-- Destroy a texture reference. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TEXREF__DEPRECATED.html#group__CUDA__TEXREF__DEPRECATED_1gea8edbd6cf9f97e6ab2b41fc6785519d>+-- {-# INLINEABLE destroy #-} destroy :: Texture -> IO () destroy !tex = nothingIfOk =<< cuTexRefDestroy tex@@ -139,6 +145,8 @@ -- Bind a linear array address of the given size (bytes) as a texture -- reference. Any previously bound references are unbound. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TEXREF.html#group__CUDA__TEXREF_1g44ef7e5055192d52b3d43456602b50a8>+-- {-# INLINEABLE bind #-} bind :: Texture -> DevicePtr a -> Int64 -> IO () bind !tex !dptr !bytes = nothingIfOk =<< cuTexRefSetAddress tex dptr bytes@@ -157,6 +165,8 @@ -- calls to 'setFormat' can not follow a call to 'bind2D' for the same texture -- reference. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TEXREF.html#group__CUDA__TEXREF_1g26f709bbe10516681913d1ffe8756ee2>+-- {-# INLINEABLE bind2D #-} bind2D :: Texture -> Format -> Int -> DevicePtr a -> (Int,Int) -> Int64 -> IO () bind2D !tex !fmt !chn !dptr (!width,!height) !pitch =@@ -177,6 +187,8 @@ -- Get the addressing mode used by a texture reference, corresponding to the -- given dimension (currently the only supported dimension values are 0 or 1). --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TEXREF.html#group__CUDA__TEXREF_1gfb367d93dc1d20aab0cf8ce70d543b33>+-- {-# INLINEABLE getAddressMode #-} getAddressMode :: Texture -> Int -> IO AddressMode getAddressMode !tex !dim = resultIfOk =<< cuTexRefGetAddressMode tex dim@@ -189,8 +201,10 @@ -- |--- Get the filtering mode used by a texture reference+-- Get the filtering mode used by a texture reference. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TEXREF.html#group__CUDA__TEXREF_1g2439e069746f69b940f2f4dbc78cdf87>+-- {-# INLINEABLE getFilterMode #-} getFilterMode :: Texture -> IO FilterMode getFilterMode !tex = resultIfOk =<< cuTexRefGetFilterMode tex@@ -202,8 +216,10 @@ -- |--- Get the data format and number of channel components of the bound texture+-- Get the data format and number of channel components of the bound texture. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TEXREF.html#group__CUDA__TEXREF_1g90936eb6c7c4434a609e1160c278ae53>+-- {-# INLINEABLE getFormat #-} getFormat :: Texture -> IO (Format, Int) getFormat !tex = do@@ -218,8 +234,10 @@ -- |--- Specify the addressing mode for the given dimension of a texture reference+-- Specify the addressing mode for the given dimension of a texture reference. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TEXREF.html#group__CUDA__TEXREF_1g85f4a13eeb94c8072f61091489349bcb>+-- {-# INLINEABLE setAddressMode #-} setAddressMode :: Texture -> Int -> AddressMode -> IO () setAddressMode !tex !dim !mode = nothingIfOk =<< cuTexRefSetAddressMode tex dim mode@@ -233,8 +251,10 @@ -- | -- Specify the filtering mode to be used when reading memory through a texture--- reference+-- reference. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TEXREF.html#group__CUDA__TEXREF_1g595d0af02c55576f8c835e4efd1f39c0>+-- {-# INLINEABLE setFilterMode #-} setFilterMode :: Texture -> FilterMode -> IO () setFilterMode !tex !mode = nothingIfOk =<< cuTexRefSetFilterMode tex mode@@ -247,8 +267,10 @@ -- | -- Specify additional characteristics for reading and indexing the texture--- reference+-- reference. --+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TEXREF.html#group__CUDA__TEXREF_1g554ffd896487533c36810f2e45bb7a28>+-- {-# INLINEABLE setReadMode #-} setReadMode :: Texture -> ReadMode -> IO () setReadMode !tex !mode = nothingIfOk =<< cuTexRefSetFlags tex mode@@ -261,7 +283,9 @@ -- | -- Specify the format of the data and number of packed components per element to--- be read by the texture reference+-- be read by the texture reference.+--+-- <http://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TEXREF.html#group__CUDA__TEXREF_1g05585ef8ea2fec728a03c6c8f87cf07a> -- {-# INLINEABLE setFormat #-} setFormat :: Texture -> Format -> Int -> IO ()
Foreign/CUDA/Driver/Utils.chs view
@@ -25,7 +25,7 @@ -- |--- Return the version number of the installed CUDA driver+-- Return the version number of the installed CUDA driver. -- {-# INLINEABLE driverVersion #-} driverVersion :: IO Int
Foreign/CUDA/Runtime/Device.chs view
@@ -2,6 +2,7 @@ {-# LANGUAGE CPP #-} {-# LANGUAGE EmptyDataDecls #-} {-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-} {-# OPTIONS_GHC -fno-warn-orphans #-} #ifdef USE_EMPTY_CASE {-# LANGUAGE EmptyCase #-}@@ -76,7 +77,7 @@ -- Device execution flags -- {# enum cudaDeviceFlags as DeviceFlag { }- with prefix="cudaDeviceFlag" deriving (Eq, Show) #}+ with prefix="cudaDeviceFlag" deriving (Eq, Show, Bounded) #} instance Storable DeviceProperties where@@ -122,58 +123,77 @@ tc <- cToBool `fmap` {#get cudaDeviceProp.tccDriver#} p ua <- cToBool `fmap` {#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 #if CUDART_VERSION >= 3000 [u21,u22] <- peekArrayWith cIntConv 2 =<< {#get cudaDeviceProp.maxTexture2D#} p [u31,u32,u33] <- peekArrayWith cIntConv 3 =<< {#get cudaDeviceProp.maxTexture3D#} p #endif+#if CUDART_VERSION >= 5050+ sp <- cToBool `fmap` {#get cudaDeviceProp.streamPrioritiesSupported#} p+#endif+#if CUDART_VERSION >= 6000+ gl1 <- cToBool `fmap` {#get cudaDeviceProp.globalL1CacheSupported#} p+ ll1 <- cToBool `fmap` {#get cudaDeviceProp.localL1CacheSupported#} p+ mm <- cToBool `fmap` {#get cudaDeviceProp.managedMemory#} p+ mg <- cToBool `fmap` {#get cudaDeviceProp.isMultiGpuBoard#} p+ mid <- cIntConv `fmap` {#get cudaDeviceProp.multiGpuBoardGroupID#} p+#endif return DeviceProperties {- deviceName = n,- computeCapability = Compute v1 v2,- totalGlobalMem = gm,- totalConstMem = cm,- sharedMemPerBlock = sm,- regsPerBlock = rb,- warpSize = ws,- maxThreadsPerBlock = tb,- maxBlockSize = (t1,t2,t3),- maxGridSize = (g1,g2,g3),- clockRate = cl,- multiProcessorCount = pc,- memPitch = mp,- textureAlignment = ta,- computeMode = md,- deviceOverlap = ov,+ deviceName = n+ , computeCapability = Compute v1 v2+ , totalGlobalMem = gm+ , totalConstMem = cm+ , sharedMemPerBlock = sm+ , regsPerBlock = rb+ , warpSize = ws+ , maxThreadsPerBlock = tb+ , maxBlockSize = (t1,t2,t3)+ , maxGridSize = (g1,g2,g3)+ , clockRate = cl+ , multiProcessorCount = pc+ , memPitch = mp+ , textureAlignment = ta+ , computeMode = md+ , deviceOverlap = ov+ , kernelExecTimeoutEnabled = ke+ , integrated = tg+ , canMapHostMemory = hm #if CUDART_VERSION >= 3000- concurrentKernels = ck,- maxTextureDim1D = u1,- maxTextureDim2D = (u21,u22),- maxTextureDim3D = (u31,u32,u33),+ , concurrentKernels = ck+ , maxTextureDim1D = u1+ , maxTextureDim2D = (u21,u22)+ , maxTextureDim3D = (u31,u32,u33) #endif #if CUDART_VERSION >= 3010- eccEnabled = ee,+ , eccEnabled = ee #endif #if CUDART_VERSION >= 3000 && CUDART_VERSION < 3010- -- not visible from runtime API < 3.1- eccEnabled = False,+ -- not visible from CUDA runtime API 3.0+ , eccEnabled = False #endif #if CUDART_VERSION >= 4000- asyncEngineCount = ae,- cacheMemL2 = l2,- maxThreadsPerMultiProcessor = tm,- memBusWidth = mw,- memClockRate = mc,- tccDriverEnabled = tc,- unifiedAddressing = ua,- pciInfo = PCI pb pd pm,+ , asyncEngineCount = ae+ , cacheMemL2 = l2+ , maxThreadsPerMultiProcessor = tm+ , memBusWidth = mw+ , memClockRate = mc+ , tccDriverEnabled = tc+ , unifiedAddressing = ua+ , pciInfo = PCI pb pd pm #endif- kernelExecTimeoutEnabled = ke,- integrated = tg,- canMapHostMemory = hm+#if CUDA_VERSION >= 5050+ , streamPriorities = sp+#endif+#if CUDA_VERSION >= 6000+ , globalL1Cache = gl1+ , localL1Cache = ll1+ , managedMemory = mm+ , multiGPUBoard = mg+ , multiGPUBoardGroupID = mid+#endif } @@ -332,7 +352,7 @@ {-# INLINEABLE accessible #-} accessible :: Device -> Device -> IO Bool #if CUDART_VERSION < 4000-accessible _ _ = requireSDK 4.0 "accessible"+accessible _ _ = requireSDK 'accessible 4.0 #else accessible !dev !peer = resultIfOk =<< cudaDeviceCanAccessPeer dev peer @@ -351,7 +371,7 @@ {-# INLINEABLE add #-} add :: Device -> [PeerFlag] -> IO () #if CUDART_VERSION < 4000-add _ _ = requireSDK 4.0 "add"+add _ _ = requireSDK 'add 4.0 #else add !dev !flags = nothingIfOk =<< cudaDeviceEnablePeerAccess dev flags @@ -369,7 +389,7 @@ {-# INLINEABLE remove #-} remove :: Device -> IO () #if CUDART_VERSION < 4000-remove _ = requireSDK 4.0 "remove"+remove _ = requireSDK 'remove 4.0 #else remove !dev = nothingIfOk =<< cudaDeviceDisablePeerAccess dev @@ -401,7 +421,7 @@ {-# INLINEABLE getLimit #-} getLimit :: Limit -> IO Int #if CUDART_VERSION < 3010-getLimit _ = requireSDK 3.1 "getLimit"+getLimit _ = requireSDK 'getLimit 3.1 #elif CUDART_VERSION < 4000 getLimit !l = resultIfOk =<< cudaThreadGetLimit l @@ -425,7 +445,7 @@ {-# INLINEABLE setLimit #-} setLimit :: Limit -> Int -> IO () #if CUDART_VERSION < 3010-setLimit _ _ = requireSDK 3.1 "setLimit"+setLimit _ _ = requireSDK 'setLimit 3.1 #elif CUDART_VERSION < 4000 setLimit !l !n = nothingIfOk =<< cudaThreadSetLimit l n
Foreign/CUDA/Runtime/Error.chs view
@@ -24,11 +24,13 @@ import Foreign.CUDA.Internal.C2HS -- System+import Control.Exception+import Data.Typeable import Foreign.C import Foreign.Ptr-import Data.Typeable-import Control.Exception+import Language.Haskell.TH import System.IO.Unsafe+import Text.Printf #include "cbits/stubs.h" {# context lib="cudart" #}@@ -70,8 +72,8 @@ -- | -- A specially formatted error message ---requireSDK :: Double -> String -> IO a-requireSDK v s = cudaError ("'" ++ s ++ "' requires at least cuda-" ++ show v)+requireSDK :: Name -> Double -> IO a+requireSDK n v = cudaError $ printf "'%s' requires at least cuda-%3.1f\n" (show n) v --------------------------------------------------------------------------------
Foreign/CUDA/Runtime/Event.chs view
@@ -2,6 +2,7 @@ {-# LANGUAGE CPP #-} {-# LANGUAGE EmptyDataDecls #-} {-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-} -------------------------------------------------------------------------------- -- | -- Module : Foreign.CUDA.Driver.Event@@ -121,7 +122,7 @@ {-# INLINEABLE wait #-} wait :: Event -> Maybe Stream -> [WaitFlag] -> IO () #if CUDART_VERSION < 3020-wait _ _ _ = requireSDK 3.2 "wait"+wait _ _ _ = requireSDK 'wait 3.2 #else wait !ev !mst !flags = let st = fromMaybe defaultStream mst
Foreign/CUDA/Runtime/Exec.chs view
@@ -2,6 +2,7 @@ {-# LANGUAGE CPP #-} {-# LANGUAGE ForeignFunctionInterface #-} {-# LANGUAGE GADTs #-}+{-# LANGUAGE TemplateHaskell #-} -------------------------------------------------------------------------------- -- | -- Module : Foreign.CUDA.Runtime.Exec@@ -216,7 +217,7 @@ {-# INLINEABLE setCacheConfig #-} setCacheConfig :: Fun -> CacheConfig -> IO () #if CUDART_VERSION < 3000-setCacheConfig _ _ = requireSDK 3.0 "setCacheConfig"+setCacheConfig _ _ = requireSDK 'setCacheConfig 3.0 #else setCacheConfig !fn !pref = nothingIfOk =<< cudaFuncSetCacheConfig fn pref @@ -265,10 +266,11 @@ -- CUDA 5.0 changed the type of a kernel function from char* to void* ---withFun :: Fun -> (Ptr a -> IO b) -> IO b #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
Foreign/CUDA/Runtime/Marshal.chs view
@@ -1,6 +1,7 @@ {-# LANGUAGE BangPatterns #-} {-# LANGUAGE EmptyDataDecls #-} {-# LANGUAGE ForeignFunctionInterface #-}+{-# LANGUAGE TemplateHaskell #-} -------------------------------------------------------------------------------- -- | -- Module : Foreign.CUDA.Runtime.Marshal@@ -86,7 +87,7 @@ -- {# enum cudaMemHostAlloc_option as AllocFlag { underscoreToCase }- with prefix="CUDA_MEMHOSTALLOC_OPTION" deriving (Eq, Show) #}+ with prefix="CUDA_MEMHOSTALLOC_OPTION" deriving (Eq, Show, Bounded) #} -- |@@ -194,7 +195,7 @@ #if CUDART_VERSION >= 6000 {# enum cudaMemAttachFlags_option as AttachFlag { underscoreToCase }- with prefix="CUDA_MEM_ATTACH_OPTION" deriving (Eq, Show) #}+ with prefix="CUDA_MEM_ATTACH_OPTION" deriving (Eq, Show, Bounded) #} #else data AttachFlag #endif@@ -206,7 +207,7 @@ {-# INLINEABLE mallocManagedArray #-} mallocManagedArray :: Storable a => [AttachFlag] -> Int -> IO (DevicePtr a) #if CUDART_VERSION < 6000-mallocManagedArray _ _ = requireSDK 6.0 "mallocManagedArray"+mallocManagedArray _ _ = requireSDK 'mallocManagedArray 6.0 #else mallocManagedArray !flags = doMalloc undefined where
Foreign/CUDA/Types.chs view
@@ -24,7 +24,7 @@ Event(..), EventFlag(..), WaitFlag, -- * Streams- Stream(..), StreamFlag,+ Stream(..), StreamPriority, StreamFlag, defaultStream, ) where@@ -99,7 +99,7 @@ -- {# enum CUevent_flags as EventFlag { underscoreToCase }- with prefix="CU_EVENT" deriving (Eq, Show) #}+ with prefix="CU_EVENT" deriving (Eq, Show, Bounded) #} -- | -- Possible option flags for waiting for events@@ -125,6 +125,14 @@ -- newtype Stream = Stream { useStream :: {# type CUstream #}} deriving (Eq, Show)+++-- |+-- Priority of an execution stream. Work submitted to a higher priority+-- stream may preempt execution of work already executing in a lower+-- priority stream. Lower numbers represent higher priorities.+--+type StreamPriority = Int -- | -- Possible option flags for stream initialisation. Dummy instance until the API
+ README.markdown view
@@ -0,0 +1,20 @@+Haskell FFI Bindings to CUDA+============================++[](https://travis-ci.org/tmcdonell/cuda)++The CUDA library provides a direct, general purpose C-like SPMD programming+model for NVIDIA graphics cards (G8x series onwards). This is a collection of+bindings to allow you to call and control, although not write, such functions+from Haskell-land. You will need to install the CUDA driver and developer+toolkit.++ <http://developer.nvidia.com/object/cuda.html>++The configure step will look for your CUDA installation in the standard places,+and if the `nvcc` compiler is found in your `PATH`, relative to that.++For important information on installing on Windows, see:++ <https://github.com/tmcdonell/cuda/blob/master/WINDOWS.markdown>+
Setup.hs view
@@ -1,51 +1,59 @@ import Distribution.PackageDescription import Distribution.PackageDescription.Parse-import Distribution.Verbosity-import Distribution.System import Distribution.Simple-import Distribution.Simple.Utils-import Distribution.Simple.Setup+import Distribution.Simple.BuildPaths import Distribution.Simple.Command-import Distribution.Simple.Program+import Distribution.Simple.Program.Db import Distribution.Simple.LocalBuildInfo import Distribution.Simple.PreProcess hiding (ppC2hs)--import Distribution.Simple.BuildPaths--import Data.List hiding (isInfixOf)-import Data.Maybe+import Distribution.Simple.Program+import Distribution.Simple.Setup+import Distribution.Simple.Utils+import Distribution.System+import Distribution.Verbosity -import Text.Printf import Control.Exception import Control.Monad-import System.Exit hiding (die)-import System.FilePath+import Data.Function+import Data.List hiding (isInfixOf)+import Data.Maybe import System.Directory import System.Environment+import System.Exit hiding (die)+import System.FilePath import System.IO.Error hiding (catch)+import Text.Printf import Prelude hiding (catch) -newtype CudaPath = CudaPath {- cudaPath :: String-} deriving (Eq, Ord, Show, Read) +newtype CudaPath = CudaPath { cudaPath :: String }+ deriving (Eq, Ord, Show, Read) + -- Windows compatibility function. ----- CUDA toolkit uses different names for import libraries and their respective DLLs.--- Eg. `cudart.lib` imports functions from `cudart32_70` (on 32-bit architecture and 7.0 version of toolkit).--- The ghci linker fails to resolve this. Therefore, it needs to be given the DLL filenames--- as `extra-ghci-libraries` option.+-- CUDA toolkit uses different names for import libraries and their+-- respective DLLs. For example, on 32-bit architecture and version 7.0 of+-- toolkit, `cudart.lib` imports functions from `cudart32_70`. ----- This function takes *a path to* import library and returns name of corresponding DLL.+-- The ghci linker fails to resolve this. Therefore, it needs to be given+-- the DLL filenames as `extra-ghci-libraries` option.+--+-- This function takes *a path to* import library and returns name of+-- corresponding DLL.+-- -- Eg: "C:/CUDA/Toolkit/Win32/cudart.lib" -> "cudart32_70.dll"--- Internally it assumes that nm tool is present in PATH. This should be always true, as nm is distributed along with GHC. ----- The function is meant to be used on Windows. Other platforms may or may not work.+-- Internally it assumes that 'nm' tool is present in PATH. This should be+-- always true, as 'nm' is distributed along with GHC. --+-- The function is meant to be used on Windows. Other platforms may or may+-- not work.+-- importLibraryToDllFileName :: FilePath -> IO (Maybe FilePath) importLibraryToDllFileName importLibPath = do -- Sample output nm generates on cudart.lib+ -- -- nvcuda.dll: -- 00000000 i .idata$2 -- 00000000 i .idata$4@@ -55,14 +63,16 @@ -- 00000000 I __IMPORT_DESCRIPTOR_nvcuda -- U __NULL_IMPORT_DESCRIPTOR -- U nvcuda_NULL_THUNK_DATA+ -- nmOutput <- getProgramInvocationOutput normal (simpleProgramInvocation "nm" [importLibPath]) return $ find (isInfixOf ("" <.> dllExtension)) (lines nmOutput) -- Windows compatibility function. ----- The function is used to populate the extraGHCiLibs list on Windows platform.--- It takes libraries directory and .lib filenames and returns their corresponding dll filename.--- (Both filenames are stripped from extensions)+-- The function is used to populate the extraGHCiLibs list on Windows+-- platform. It takes libraries directory and .lib filenames and returns+-- their corresponding dll filename. (Both filenames are stripped from+-- extensions) -- -- Eg: "C:\cuda\toolkit\lib\x64" -> ["cudart", "cuda"] -> ["cudart64_65", "ncuda"] --@@ -73,14 +83,16 @@ let dllNames = map (\(Just dllname) -> dropExtension dllname) (filter isJust candidateNames) return dllNames --- OSX compatibility function+-- Mac OS X compatibility function -- -- Returns [] or ["U__BLOCKS__"] -- getAppleBlocksOption :: IO [String] getAppleBlocksOption = do- let handler = (\_ -> return "") :: IOError -> IO String- fileContents <- catch (readFile "/usr/include/stdlib.h") handler -- If file does not exist, we'll end up wth an empty string.+ let handler :: IOError -> IO String+ handler = (\_ -> return "")+ -- If file does not exist, we'll end up wth an empty string.+ fileContents <- catch (readFile "/usr/include/stdlib.h") handler return ["-U__BLOCKS__" | "__BLOCKS__" `isInfixOf` fileContents] getCudaIncludePath :: CudaPath -> FilePath@@ -91,78 +103,219 @@ where libSubpath = case os of Windows -> "lib" </> case arch of- I386 -> "Win32"- X86_64 -> "x64"- _ -> error $ "Unexpected Windows architecture " ++ show arch ++ ". Please report this issue to https://github.com/tmcdonell/cuda/issues"+ I386 -> "Win32"+ X86_64 -> "x64"+ _ -> error $ printf "Unexpected Windows architecture %s.\nPlease report this issue to https://github.com/tmcdonell/cuda/issues\n" (show arch) - OSX -> "lib"+ OSX -> "lib" - -- For now just treat all non-Windows systems similarly+ -- For now just treat all other systems similarly _ -> case arch of- I386 -> "lib"- X86_64 -> "lib64"- _ -> "lib" -- TODO how should this be handled?+ X86_64 -> "lib64"+ I386 -> "lib"+ _ -> "lib" -- TODO: how should this be handled? -getCudaLibraries :: [String]-getCudaLibraries = ["cudart", "cuda"] +-- On OS X we don't link against the CUDA and CUDART libraries directly.+-- Instead, we only link against CUDA.framework. This means that we will+-- not need to set the DYLD_LIBRARY_PATH environment variable in order to+-- compile or execute programs.+--+getCudaLibraries :: Platform -> [String]+getCudaLibraries (Platform _ os) =+ case os of+ OSX -> []+ _ -> ["cudart", "cuda"] ++-- Slightly modified version of `words` from base - it takes predicate saying on which characters split.+splitOn :: (Char -> Bool) -> String -> [String]+splitOn p s = case dropWhile p s of+ "" -> []+ s' -> w : splitOn p s''+ where (w, s'') = break p s'++-- Tries to obtain the version `ld`.+-- Throws an exception if failed.+--+getLdVersion :: Verbosity -> FilePath -> IO (Maybe [Int])+getLdVersion verbosity ldPath = do+ -- Version string format is like `GNU ld (GNU Binutils) 2.25.1`+ -- or `GNU ld (GNU Binutils) 2.20.51.20100613`+ ldVersionString <- getProgramInvocationOutput normal (simpleProgramInvocation ldPath ["-v"])++ let versionText = last $ words ldVersionString -- takes e. g. "2.25.1"+ let versionParts = splitOn (== '.') versionText+ let versionParsed = Just $ map read versionParts++ -- last and read above may throw and message would be not understandable for user,+ -- so we'll intercept exception and rethrow it with more useful message.+ let handleError :: SomeException -> IO (Maybe [Int])+ handleError e = do+ warn verbosity $ printf "cannot parse ld version string: `%s`. Parsing exception: `%s`" ldVersionString (show e)+ return Nothing++ catch (evaluate versionParsed) handleError++++-- On Windows GHC package comes with two copies of ld.exe.+-- ProgramDb knows about the first one - ghcpath\mingw\bin\ld.exe+-- This function returns the other one - ghcpath\mingw\x86_64-w64-mingw32\bin\ld.exe+-- The second one is the one that does actual linking and code generation.+-- See: https://github.com/tmcdonell/cuda/issues/31#issuecomment-149181376+--+-- The function is meant to be used only on 64-bit GHC distributions.+--+getRealLdPath :: Verbosity -> ProgramDb -> IO (Maybe FilePath)+getRealLdPath verbosity programDb =+ -- This should ideally work `programFindVersion ldProgram` but for some reason it does not.+ -- The issue should be investigated at some time.+ case lookupProgram ghcProgram programDb of+ Nothing -> return Nothing+ Just configuredGhc -> do+ let ghcPath = locationPath $ programLocation configuredGhc+ presumedLdPath = (takeDirectory . takeDirectory) ghcPath </> "mingw" </> "x86_64-w64-mingw32" </> "bin" </> "ld.exe"+ info verbosity $ "Presuming ld location" ++ presumedLdPath+ presumedLdExists <- doesFileExist presumedLdPath+ return $ if presumedLdExists then Just presumedLdPath else Nothing++-- On Windows platform the binutils linker targeting x64 is bugged and cannot+-- properly link with import libraries generated by MS compiler (like the CUDA ones).+-- The programs would correctly compile and crash as soon as the first FFI call is made.+--+-- Therefore we fail configure process if the linker is too old and provide user+-- with guidelines on how to fix the problem.+--+validateLinker :: Verbosity -> Platform -> ProgramDb -> IO ()+validateLinker verbosity (Platform X86_64 Windows) db = do+ maybeLdPath <- getRealLdPath verbosity db+ case maybeLdPath of+ Nothing -> warn verbosity $ "Cannot find ld.exe to check if it is new enough. If generated executables crash when making calls to CUDA, please see " ++ helpfulPageLinkForWindows+ Just ldPath -> do+ debug verbosity $ "Checking if ld.exe at " ++ ldPath ++ " is new enough"+ maybeVersion <- getLdVersion verbosity ldPath+ case maybeVersion of+ Nothing -> warn verbosity $ "Unknown ld.exe version. If generated executables crash when making calls to CUDA, please see " ++ helpfulPageLinkForWindows+ Just ldVersion -> do+ debug verbosity $ "Found ld.exe version: " ++ show ldVersion+ when (ldVersion < [2,25,1]) $ die $ linkerBugOnWindowsMsg ldPath+validateLinker _ _ _ = return () -- The linker bug is present only on Win64 platform++helpfulPageLinkForWindows :: String+helpfulPageLinkForWindows = "https://github.com/tmcdonell/cuda/blob/master/WINDOWS.markdown"++linkerBugOnWindowsMsg :: FilePath -> String+linkerBugOnWindowsMsg ldPath = printf (unlines msg) ldPath+ where+ msg =+ [ "********************************************************************************"+ , ""+ , "The installed version of `ld.exe` has version < 2.25.1. This version has known bug on Windows x64 architecture, making it unable to correctly link programs using CUDA. The fix is available and MSys2 released fixed version of `ld.exe` as part of their binutils package (version 2.25.1)."+ , ""+ , "To fix this issue, replace the `ld.exe` in your GHC installation with the correct binary. See the following page for details:"+ , ""+ , " " ++ helpfulPageLinkForWindows+ , ""+ , "The full path to the outdated `ld.exe` detected in your installation:"+ , ""+ , "> %s"+ , ""+ , "Please download a recent version of binutils `ld.exe`, from, e.g.:"+ , ""+ , " http://repo.msys2.org/mingw/x86_64/mingw-w64-x86_64-binutils-2.25.1-1-any.pkg.tar.xz"+ , ""+ , "********************************************************************************"+ ]+ -- Generates build info with flags needed for CUDA Toolkit to be properly -- visible to underlying build tools. -- cudaLibraryBuildInfo :: CudaPath -> Platform -> Version -> IO HookedBuildInfo cudaLibraryBuildInfo cudaPath platform@(Platform arch os) ghcVersion = do- let cudaLibraryPath = getCudaLibraryPath cudaPath platform- -- Extra lib dirs are not needed on Windows somehow. On Linux their lack would cause an error: /usr/bin/ld: cannot find -lcudart- -- Still, they do not cause harm so let's have them regardless of OS.- let extraLibDirs_ = [cudaLibraryPath]- let includeDirs = [getCudaIncludePath cudaPath]- let ccOptions_ = map ("-I" ++) includeDirs- let ldOptions_ = map ("-L" ++) extraLibDirs_- let ghcOptions = map ("-optc" ++) ccOptions_ ++ map ("-optl" ++ ) ldOptions_- let extraLibs_ = getCudaLibraries+ let cudaLibraryPath = getCudaLibraryPath cudaPath platform + -- Extra lib dirs are not needed on Mac OS. On Windows or Linux their+ -- lack would cause an error: /usr/bin/ld: cannot find -lcudart+ let extraLibDirs_ = case os of+ OSX -> []+ _ -> [cudaLibraryPath]++ let includeDirs = [getCudaIncludePath cudaPath]+ let ccOptions_ = map ("-I" ++) includeDirs+ let ldOptions_ = map ("-L" ++) extraLibDirs_+ let ghcOptions = map ("-optc" ++) ccOptions_ ++ map ("-optl" ++ ) ldOptions_+ let extraLibs_ = getCudaLibraries platform+ -- Options for C2HS- let c2hsArchitectureFlag = case arch of I386 -> ["-m32"]- X86_64 -> ["-m64"]- _ -> []+ let c2hsArchitectureFlag = case arch of+ I386 -> ["-m32"]+ X86_64 -> ["-m64"]+ _ -> [] let c2hsEmptyCaseFlag = ["-DUSE_EMPTY_CASE" | versionBranch ghcVersion >= [7,8]]- let c2hsCppOptions = c2hsArchitectureFlag ++ c2hsEmptyCaseFlag ++ ["-E"]-- -- On OSX we might add one more options to c2hs cpp.- appleBlocksOption <- case os of OSX -> getAppleBlocksOption; _ -> return []+ let c2hsCppOptions = c2hsArchitectureFlag ++ c2hsEmptyCaseFlag ++ ["-E"] - let c2hsOptions = unwords $ map ("--cppopts=" ++) (c2hsCppOptions ++ appleBlocksOption)- let extraOptionsC2Hs = ("x-extra-c2hs-options", c2hsOptions)- let buildInfo = emptyBuildInfo+ let c2hsOptions = unwords . map ("--cppopts=" ++)+ let extraOptionsC2Hs = ("x-extra-c2hs-options", c2hsOptions c2hsCppOptions)+ let buildInfo = emptyBuildInfo { ccOptions = ccOptions_ , ldOptions = ldOptions_ , extraLibs = extraLibs_ , extraLibDirs = extraLibDirs_- , options = [(GHC, ghcOptions)] -- Is this needed for anything?+ -- Are ghc-options below needed for anything?+ -- On Windows they need to be disabled because Cabal does not escape+ -- them (quotes and backslashes) causing build fails on machines+ -- with CUDA_PATH containing spaces.+ , options = [(GHC, ghcOptions) | os /= Windows] , customFieldsBI = [extraOptionsC2Hs] } let addSystemSpecificOptions :: Platform -> IO BuildInfo addSystemSpecificOptions (Platform _ Windows) = do- -- Workaround issue with ghci linker not being able to find DLLs with names different from their import LIBs.+ -- Workaround issue with ghci linker not being able to find DLLs+ -- with names different from their import LIBs. extraGHCiLibs_ <- additionalGhciLibraries cudaLibraryPath extraLibs_ return buildInfo { extraGHCiLibs = extraGHCiLibs buildInfo ++ extraGHCiLibs_ }- addSystemSpecificOptions (Platform _ OSX) = return buildInfo- { customFieldsBI = customFieldsBI buildInfo ++ [("frameworks", "CUDA")]- , ldOptions = ldOptions buildInfo ++ ["-F/Library/Frameworks"]++ addSystemSpecificOptions (Platform _ OSX) = do+ -- On OS X tell the linker about the CUDA framework. It seems like+ -- this shouldn't be necessary, since we also specify this in the+ -- frameworks field. Possibly haskell/cabal#2724?+ --+ -- We also might need to add one or more options to c2hs cpp.+ appleBlocksOption <- getAppleBlocksOption+ return buildInfo+ { ldOptions = ldOptions buildInfo ++ ["-framework", "CUDA"]+ , customFieldsBI = unionWith (+++) []+ $ customFieldsBI buildInfo ++ [("frameworks", "CUDA")+ ,("x-extra-c2hs-options", c2hsOptions appleBlocksOption)] }+ addSystemSpecificOptions _ = return buildInfo adjustedBuildInfo <-addSystemSpecificOptions platform return (Just adjustedBuildInfo, []) ++unionWith :: Ord k => (a -> a -> a) -> a -> [(k,a)] -> [(k,a)]+unionWith f z+ = map (\kv -> let (k,v) = unzip kv in (head k, foldr f z v))+ . groupBy ((==) `on` fst)+ . sortBy (compare `on` fst)++(+++) :: String -> String -> String+[] +++ ys = ys+xs +++ [] = xs+xs +++ ys = xs ++ ' ':ys++ -- Checks whether given location looks like a valid CUDA toolkit directory -- validateLocation :: Verbosity -> FilePath -> IO Bool validateLocation verbosity path = do- -- TODO: Ideally this should check also for cudart.lib and whether cudart exports relevant symbols.- -- This should be achievable with some `nm` trickery+ -- TODO: Ideally this should check also for cudart.lib and whether cudart+ -- exports relevant symbols. This should be achievable with some `nm`+ -- trickery let testedPath = path </> "include" </> "cuda.h" exists <- doesFileExist testedPath info verbosity $@@ -199,14 +352,14 @@ nvccProgramName = "nvcc" -- NOTE: this function throws an exception when there is no `nvcc` in PATH.--- The exception contains meaningful message.+-- The exception contains a meaningful message. -- findProgramLocationThrowing :: String -> IO FilePath findProgramLocationThrowing execName = do location <- findProgramLocation normal execName case location of Just validLocation -> return validLocation- Nothing -> ioError $ mkIOError doesNotExistErrorType ("not found: " ++ execName) Nothing Nothing+ Nothing -> ioError $ mkIOError doesNotExistErrorType ("not found: " ++ execName) Nothing Nothing -- Returns pairs (action yielding candidate path, String description of that location) --@@ -264,9 +417,9 @@ -- Runs CUDA detection procedure and stores .buildinfo to a file. -- generateAndStoreBuildInfo :: Verbosity -> Platform -> CompilerId -> FilePath -> IO ()-generateAndStoreBuildInfo verbosity platform (CompilerId ghcFlavor ghcVersion) path = do+generateAndStoreBuildInfo verbosity platform (CompilerId _ghcFlavor ghcVersion) path = do cudalocation <- findCudaLocation verbosity- pbi <- cudaLibraryBuildInfo cudalocation platform ghcVersion+ pbi <- cudaLibraryBuildInfo cudalocation platform ghcVersion storeHookedBuildInfo verbosity path pbi customBuildinfoFilepath :: FilePath@@ -312,17 +465,19 @@ -- found, an error is raised. Otherwise the toolkit location is used to -- create a `cuda.buildinfo.generated` file with all the resulting flags. postConfHook :: Args -> ConfigFlags -> PackageDescription -> LocalBuildInfo -> IO ()- postConfHook args flags pkg_descr lbi- = let verbosity = fromFlag (configVerbosity flags)- currentPlatform = hostPlatform lbi- compilerId_ = (compilerId $ compiler lbi)- in do- noExtraFlags args- generateAndStoreBuildInfo verbosity currentPlatform compilerId_ generatedBuldinfoFilepath-- actualBuildInfoToUse <- getHookedBuildInfo verbosity- let pkg_descr' = updatePackageDescription actualBuildInfoToUse pkg_descr- postConf simpleUserHooks args flags pkg_descr' lbi+ postConfHook args flags pkg_descr lbi = do+ let+ verbosity = fromFlag (configVerbosity flags)+ currentPlatform = hostPlatform lbi+ compilerId_ = (compilerId $ compiler lbi)+ --+ noExtraFlags args+ generateAndStoreBuildInfo verbosity currentPlatform compilerId_ generatedBuldinfoFilepath+ validateLinker verbosity currentPlatform $ withPrograms lbi+ --+ actualBuildInfoToUse <- getHookedBuildInfo verbosity+ let pkg_descr' = updatePackageDescription actualBuildInfoToUse pkg_descr+ postConf simpleUserHooks args flags pkg_descr' lbi storeHookedBuildInfo :: Verbosity -> FilePath -> HookedBuildInfo -> IO ()@@ -338,16 +493,19 @@ getHookedBuildInfo :: Verbosity -> IO HookedBuildInfo getHookedBuildInfo verbosity = do doesCustomBuildInfoExists <- doesFileExist customBuildinfoFilepath- if doesCustomBuildInfoExists then do- notice verbosity $ "The user-provided buildinfo from file " ++ customBuildinfoFilepath ++ " will be used. To use default settings, delete this file."- readHookedBuildInfo verbosity customBuildinfoFilepath- else do- doesGeneratedBuildInfoExists <- doesFileExist generatedBuldinfoFilepath- if doesGeneratedBuildInfoExists then do- notice verbosity $ printf "Using build information from '%s'.\n" generatedBuldinfoFilepath- notice verbosity $ printf "Provide a '%s' file to override this behaviour.\n" customBuildinfoFilepath- readHookedBuildInfo verbosity generatedBuldinfoFilepath- else die $ "Unexpected failure. Neither the default " ++ generatedBuldinfoFilepath ++ " nor custom " ++ customBuildinfoFilepath ++ " do exist."+ 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+ else do+ doesGeneratedBuildInfoExists <- doesFileExist generatedBuldinfoFilepath+ if doesGeneratedBuildInfoExists+ then do+ notice verbosity $ printf "Using build information from '%s'.\n" generatedBuldinfoFilepath+ notice verbosity $ printf "Provide a '%s' file to override this behaviour.\n" customBuildinfoFilepath+ readHookedBuildInfo verbosity generatedBuldinfoFilepath+ else+ die $ printf "Unexpected failure. Neither the default %s nor custom %s exist.\n" generatedBuldinfoFilepath customBuildinfoFilepath -- Replicate the default C2HS preprocessor hook here, and inject a value for
+ WINDOWS.markdown view
@@ -0,0 +1,54 @@+Using the CUDA package on Windows+=================================++The CUDA package works on Windows and is actively maintained. If you encounter+any other issues, please report them.++Note that if you build your applications for the Windows 64-bit architecture,+you'll need to update your `ld.exe` as described below.+++Windows 64-bit+--------------++There is a known issue with the version of `ld.exe` that ships with the 64-bit+versions of (at least) GHC-7.8.4 and GHC-7.10.2. The version of `ld.exe` that+ships with these GHC distributions does not properly link against MS-style+dynamic libraries (such as those that ship with the CUDA toolkit), causing the+application to crash at runtime once those library routines are called. The+configure step will fail if it detects an old version of `ld.exe` (< 2.25.1),+which are known to be broken.++If you are using the 64-bit GHC distributions mentioned above, you will need to+apply the following steps. This bug does not affect 32-bit GHC distributions.+The bug has been fixed in MinGW binutils `ld.exe` >= 2.25.1, so it is expected+that newer releases of GHC will not have this issue.++The problem is fixed by replacing the linker binary `ld.exe` with the newer+(patched) version, available as part of the MSys2 binutils package here:++> <http://repo.msys2.org/mingw/x86_64/mingw-w64-x86_64-binutils-2.25.1-1-any.pkg.tar.xz>++The updated `ld.exe` binary must replace the version at the path:++> `GHC_PATH\mingw\x86_64-w64-mingw32\bin\`++Note that there is another copy of `ld.exe` located at `GHC_PATH\mingw\bin\`,+but this version does not seem to be used, so replacing it as well is not+necessary. It is not sufficient to replace whatever version of `ld.exe` appears+first in your `PATH`.++Please note that having another MinGW installation in `PATH` before the one+shipped with GHC may break things, particularly if you mix 32/64-bit+distributions of MinGW and GHC.++For further discussion of the bug, see:++ * [CUDA package issue][cuda31]+ * [GHC issue][ghc10885]+ * [binutils issue][binutils16598]++ [cuda31]: https://github.com/tmcdonell/cuda/issues/31+ [ghc10885]: https://ghc.haskell.org/trac/ghc/ticket/10885+ [binutils16598]: https://sourceware.org/bugzilla/show_bug.cgi?id=16598+
cbits/stubs.c view
@@ -382,5 +382,20 @@ { return cuMemHostRegister_v2(p, bytesize, Flags); }++CUresult CUDAAPI cuLinkCreate(unsigned int numOptions, CUjit_option *options, void **optionValues, CUlinkState *stateOut)+{+ return cuLinkCreate_v2(numOptions, options, optionValues, stateOut);+}++CUresult CUDAAPI cuLinkAddData(CUlinkState state, CUjitInputType type, void *data, size_t size, const char *name, unsigned int numOptions, CUjit_option *options, void **optionValues)+{+ return cuLinkAddData_v2(state, type, data, size, name, numOptions, options, optionValues);+}++CUresult CUDAAPI cuLinkAddFile(CUlinkState state, CUjitInputType type, const char *path, unsigned int numOptions, CUjit_option *options, void **optionValues)+{+ return cuLinkAddFile_v2(state, type, path, numOptions, options, optionValues);+} #endif
cbits/stubs.h view
@@ -11,19 +11,6 @@ #define CUDARTAPI __stdcall #endif -// #define __cdecl --/*- * We need to work around some shortcomings in the C parser of c2hs by disabling advanced attributes etc on Apple platforms.- */-#ifdef __APPLE__-#define _ANSI_SOURCE-#define __AVAILABILITY__-#define __OSX_AVAILABLE_STARTING(_mac, _iphone)-#define __OSX_AVAILABLE_BUT_DEPRECATED(_macIntro, _macDep, _iphoneIntro, _iphoneDep)-#define __OSX_AVAILABLE_BUT_DEPRECATED_MSG(_osxIntro, _osxDep, _iosIntro, _iosDep, _msg)-#endif- #include <cuda.h> #include <cuda_runtime_api.h> @@ -221,8 +208,14 @@ #if CUDA_VERSION >= 6050 #undef cuMemHostRegister+#undef cuLinkCreate+#undef cuLinkAddData+#undef cuLinkAddFile CUresult CUDAAPI cuMemHostRegister(void *p, size_t bytesize, unsigned int Flags);+CUresult CUDAAPI cuLinkCreate(unsigned int numOptions, CUjit_option *options, void **optionValues, CUlinkState *stateOut);+CUresult CUDAAPI cuLinkAddData(CUlinkState state, CUjitInputType type, void *data, size_t size, const char *name, unsigned int numOptions, CUjit_option *options, void **optionValues);+CUresult CUDAAPI cuLinkAddFile(CUlinkState state, CUjitInputType type, const char *path, unsigned int numOptions, CUjit_option *options, void **optionValues); #endif #ifdef __cplusplus
− changelog.md
@@ -1,34 +0,0 @@-0.6.7.0-- * Add support for building on Windows (thanks to mwu-tow)--0.6.6.2-- * Build fix--0.6.6.1-- * Build fixes for ghc-7.6 and ghc-7.10--0.6.6.0-- * Drop support for CUDA 3.0 and older.-- * Combine the definition of the 'Event' and 'Stream' data types. As of- CUDA-3.1 these data structures are equivalent, and can be safely shared- between runtime and driver API calls and libraries.-- * Mark FFI imports of potentially long-running API functions as safe. This- allows them to be safely called from Haskell threads without blocking the- entire HEC.-- * Add compute-capability data for 3.7, 5.2 devices.--0.6.5.1-- * Build fix for Mac OS X 10.10 (Yosemite)--0.6.5.0-- * Add support for the CUDA-6.5 release-
cuda.cabal view
@@ -1,5 +1,5 @@ Name: cuda-Version: 0.6.7.0+Version: 0.7.0.0 Synopsis: FFI binding to the CUDA interface for programming NVIDIA GPUs Description: The CUDA library provides a direct, general purpose C-like SPMD programming@@ -13,13 +13,19 @@ The configure script will look for your CUDA installation in the standard places, and if the nvcc compiler is found in your PATH, relative to that. .- This release is for version 6.5 of the CUDA toolkit.+ This library provides bindings to both the CUDA Driver and Runtime APIs. To+ get started, see one of: .- [/NOTE:/]+ * "Foreign.CUDA.Driver" .- Due to a bug in nvcc, this package is not compatible with c2hs-0.18.* or- c2hs-0.19.*. See tmcdonell/cuda#18.+ * "Foreign.CUDA.Runtime" .+ This release tested with versions 6.0, 6.5, and 7.0 of the CUDA toolkit.+ .+ For additional notes on installing on Windows, see:+ .+ <https://github.com/tmcdonell/cuda/blob/master/WINDOWS.markdown>+ . License: BSD3 License-file: LICENSE@@ -39,7 +45,9 @@ config.log Extra-source-files: cbits/stubs.h- changelog.md+ CHANGELOG.markdown+ README.markdown+ WINDOWS.markdown Library Exposed-Modules: Foreign.CUDA@@ -59,12 +67,21 @@ Foreign.CUDA.Runtime.Utils Foreign.CUDA.Driver Foreign.CUDA.Driver.Context+ Foreign.CUDA.Driver.Context.Base+ Foreign.CUDA.Driver.Context.Config+ Foreign.CUDA.Driver.Context.Peer+ Foreign.CUDA.Driver.Context.Primary Foreign.CUDA.Driver.Device Foreign.CUDA.Driver.Error Foreign.CUDA.Driver.Event Foreign.CUDA.Driver.Exec+ Foreign.CUDA.Driver.IPC.Event+ Foreign.CUDA.Driver.IPC.Marshal Foreign.CUDA.Driver.Marshal Foreign.CUDA.Driver.Module+ Foreign.CUDA.Driver.Module.Base+ Foreign.CUDA.Driver.Module.Link+ Foreign.CUDA.Driver.Module.Query Foreign.CUDA.Driver.Stream Foreign.CUDA.Driver.Texture Foreign.CUDA.Driver.Utils@@ -78,7 +95,7 @@ Build-depends: base >= 4 && < 5 , bytestring- , Cabal >= 1.22+ , template-haskell default-language: Haskell98 Extensions:@@ -90,9 +107,11 @@ Main-is: DeviceQuery.hs hs-source-dirs: examples/src/deviceQueryDrv - Build-depends: base >= 4 && < 5,- cuda,- pretty+ Build-depends:+ base >= 4 && < 5+ , cuda+ , pretty+ default-language: Haskell98