diff --git a/CHANGELOG.markdown b/CHANGELOG.markdown new file mode 100644 --- /dev/null +++ b/CHANGELOG.markdown @@ -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 + diff --git a/Foreign/CUDA/Analysis/Device.chs b/Foreign/CUDA/Analysis/Device.chs --- a/Foreign/CUDA/Analysis/Device.chs +++ b/Foreign/CUDA/Analysis/Device.chs @@ -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) diff --git a/Foreign/CUDA/Driver.hs b/Foreign/CUDA/Driver.hs --- a/Foreign/CUDA/Driver.hs +++ b/Foreign/CUDA/Driver.hs @@ -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 ) diff --git a/Foreign/CUDA/Driver/Context.chs b/Foreign/CUDA/Driver/Context.chs deleted file mode 100644 --- a/Foreign/CUDA/Driver/Context.chs +++ /dev/null @@ -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 - diff --git a/Foreign/CUDA/Driver/Context.hs b/Foreign/CUDA/Driver/Context.hs new file mode 100644 --- /dev/null +++ b/Foreign/CUDA/Driver/Context.hs @@ -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 + diff --git a/Foreign/CUDA/Driver/Context/Base.chs b/Foreign/CUDA/Driver/Context/Base.chs new file mode 100644 --- /dev/null +++ b/Foreign/CUDA/Driver/Context/Base.chs @@ -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'. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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 (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. +-- +-- +-- +{-# 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 +-- +-- +-- +{-# 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'. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# INLINEABLE sync #-} +sync :: IO () +sync = nothingIfOk =<< cuCtxSynchronize + +{-# INLINE cuCtxSynchronize #-} +{# fun cuCtxSynchronize + { } -> `Status' cToEnum #} + diff --git a/Foreign/CUDA/Driver/Context/Config.chs b/Foreign/CUDA/Driver/Context/Config.chs new file mode 100644 --- /dev/null +++ b/Foreign/CUDA/Driver/Context/Config.chs @@ -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 +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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 +-- +-- +-- +{-# 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 +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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 + diff --git a/Foreign/CUDA/Driver/Context/Peer.chs b/Foreign/CUDA/Driver/Context/Peer.chs new file mode 100644 --- /dev/null +++ b/Foreign/CUDA/Driver/Context/Peer.chs @@ -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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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 + diff --git a/Foreign/CUDA/Driver/Context/Primary.chs b/Foreign/CUDA/Driver/Context/Primary.chs new file mode 100644 --- /dev/null +++ b/Foreign/CUDA/Driver/Context/Primary.chs @@ -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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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 + diff --git a/Foreign/CUDA/Driver/Device.chs b/Foreign/CUDA/Driver/Device.chs --- a/Foreign/CUDA/Driver/Device.chs +++ b/Foreign/CUDA/Driver/Device.chs @@ -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. -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# INLINEABLE count #-} count :: IO Int count = resultIfOk =<< cuDeviceGetCount @@ -198,8 +222,10 @@ -- | --- Name of the device +-- The identifying name of the device. -- +-- +-- {-# 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). +-- +-- -- {-# INLINEABLE totalMem #-} totalMem :: Device -> IO Int64 diff --git a/Foreign/CUDA/Driver/Error.chs b/Foreign/CUDA/Driver/Error.chs --- a/Foreign/CUDA/Driver/Error.chs +++ b/Foreign/CUDA/Driver/Error.chs @@ -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 -------------------------------------------------------------------------------- diff --git a/Foreign/CUDA/Driver/Event.chs b/Foreign/CUDA/Driver/Event.chs --- a/Foreign/CUDA/Driver/Event.chs +++ b/Foreign/CUDA/Driver/Event.chs @@ -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 -- +-- +-- {-# INLINEABLE create #-} create :: [EventFlag] -> IO Event create !flags = resultIfOk =<< cuEventCreate flags @@ -57,6 +60,8 @@ -- | -- Destroy an event -- +-- +-- {-# INLINEABLE destroy #-} destroy :: Event -> IO () destroy !ev = nothingIfOk =<< cuEventDestroy ev @@ -69,6 +74,8 @@ -- | -- Determine the elapsed time (in milliseconds) between two events -- +-- +-- {-# 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 -- +-- +-- {-# 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. -- +-- +-- {-# 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. +-- +-- +-- {-# 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 +-- +-- -- {-# INLINEABLE block #-} block :: Event -> IO () diff --git a/Foreign/CUDA/Driver/Exec.chs b/Foreign/CUDA/Driver/Exec.chs --- a/Foreign/CUDA/Driver/Exec.chs +++ b/Foreign/CUDA/Driver/Exec.chs @@ -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. -- +-- +-- {-# 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. +-- +-- +-- {-# 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. +-- +-- +-- +{-# 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. -- +-- +-- {-# 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 #} + diff --git a/Foreign/CUDA/Driver/IPC/Event.chs b/Foreign/CUDA/Driver/IPC/Event.chs new file mode 100644 --- /dev/null +++ b/Foreign/CUDA/Driver/IPC/Event.chs @@ -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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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#} + diff --git a/Foreign/CUDA/Driver/IPC/Marshal.chs b/Foreign/CUDA/Driver/IPC/Marshal.chs new file mode 100644 --- /dev/null +++ b/Foreign/CUDA/Driver/IPC/Marshal.chs @@ -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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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#} + diff --git a/Foreign/CUDA/Driver/Marshal.chs b/Foreign/CUDA/Driver/Marshal.chs --- a/Foreign/CUDA/Driver/Marshal.chs +++ b/Foreign/CUDA/Driver/Marshal.chs @@ -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. +-- +-- +-- {-# 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. -- +-- +-- {-# 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. +-- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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 +-- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# INLINEABLE peekArray2D #-} peekArray2D :: Storable a @@ -368,6 +413,8 @@ -- associated with a particular execution stream. The destination host memory -- must be page-locked. -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# INLINEABLE pokeArray2D #-} pokeArray2D :: Storable a @@ -521,6 +574,8 @@ -- associated with a particular execution stream. The source host memory must be -- page-locked. -- +-- +-- {-# INLINEABLE pokeArray2DAsync #-} pokeArray2DAsync :: Storable a @@ -588,6 +643,8 @@ -- is asynchronous with respect to the host, but will never overlap with kernel -- execution. -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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. +-- +-- +-- {-# 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. +-- +-- +-- {-# 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. -- +-- +-- +-- +-- +-- +-- {-# 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. +-- +-- +-- +-- +-- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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). +-- +-- -- {-# INLINEABLE getMemInfo #-} getMemInfo :: IO (Int64, Int64) diff --git a/Foreign/CUDA/Driver/Module.chs b/Foreign/CUDA/Driver/Module.chs deleted file mode 100644 --- a/Foreign/CUDA/Driver/Module.chs +++ /dev/null @@ -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) - diff --git a/Foreign/CUDA/Driver/Module.hs b/Foreign/CUDA/Driver/Module.hs new file mode 100644 --- /dev/null +++ b/Foreign/CUDA/Driver/Module.hs @@ -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 + diff --git a/Foreign/CUDA/Driver/Module/Base.chs b/Foreign/CUDA/Driver/Module/Base.chs new file mode 100644 --- /dev/null +++ b/Foreign/CUDA/Driver/Module/Base.chs @@ -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. +-- +-- +-- +{-# 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 + 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. +-- +-- +-- +{-# 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) + diff --git a/Foreign/CUDA/Driver/Module/Link.chs b/Foreign/CUDA/Driver/Module/Link.chs new file mode 100644 --- /dev/null +++ b/Foreign/CUDA/Driver/Module/Link.chs @@ -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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- +{-# 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 "" 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 + diff --git a/Foreign/CUDA/Driver/Module/Query.chs b/Foreign/CUDA/Driver/Module/Query.chs new file mode 100644 --- /dev/null +++ b/Foreign/CUDA/Driver/Module/Query.chs @@ -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. +-- +-- +-- +{-# 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. This texture reference handle +-- should not be destroyed, as the texture will be destroyed automatically +-- when the module is unloaded. +-- +-- +-- +{-# 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) + diff --git a/Foreign/CUDA/Driver/Stream.chs b/Foreign/CUDA/Driver/Stream.chs --- a/Foreign/CUDA/Driver/Stream.chs +++ b/Foreign/CUDA/Driver/Stream.chs @@ -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. -- +-- +-- {-# 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. +-- +-- +-- +{-# 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. +-- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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. +-- +-- +-- +{-# 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 diff --git a/Foreign/CUDA/Driver/Texture.chs b/Foreign/CUDA/Driver/Texture.chs --- a/Foreign/CUDA/Driver/Texture.chs +++ b/Foreign/CUDA/Driver/Texture.chs @@ -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. -- +-- +-- {-# INLINEABLE create #-} create :: IO Texture create = resultIfOk =<< cuTexRefCreate @@ -124,8 +128,10 @@ -- | --- Destroy a texture reference +-- Destroy a texture reference. -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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). -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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. -- +-- +-- {-# 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. +-- +-- -- {-# INLINEABLE setFormat #-} setFormat :: Texture -> Format -> Int -> IO () diff --git a/Foreign/CUDA/Driver/Utils.chs b/Foreign/CUDA/Driver/Utils.chs --- a/Foreign/CUDA/Driver/Utils.chs +++ b/Foreign/CUDA/Driver/Utils.chs @@ -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 diff --git a/Foreign/CUDA/Runtime/Device.chs b/Foreign/CUDA/Runtime/Device.chs --- a/Foreign/CUDA/Runtime/Device.chs +++ b/Foreign/CUDA/Runtime/Device.chs @@ -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 diff --git a/Foreign/CUDA/Runtime/Error.chs b/Foreign/CUDA/Runtime/Error.chs --- a/Foreign/CUDA/Runtime/Error.chs +++ b/Foreign/CUDA/Runtime/Error.chs @@ -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 -------------------------------------------------------------------------------- diff --git a/Foreign/CUDA/Runtime/Event.chs b/Foreign/CUDA/Runtime/Event.chs --- a/Foreign/CUDA/Runtime/Event.chs +++ b/Foreign/CUDA/Runtime/Event.chs @@ -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 diff --git a/Foreign/CUDA/Runtime/Exec.chs b/Foreign/CUDA/Runtime/Exec.chs --- a/Foreign/CUDA/Runtime/Exec.chs +++ b/Foreign/CUDA/Runtime/Exec.chs @@ -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 diff --git a/Foreign/CUDA/Runtime/Marshal.chs b/Foreign/CUDA/Runtime/Marshal.chs --- a/Foreign/CUDA/Runtime/Marshal.chs +++ b/Foreign/CUDA/Runtime/Marshal.chs @@ -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 diff --git a/Foreign/CUDA/Types.chs b/Foreign/CUDA/Types.chs --- a/Foreign/CUDA/Types.chs +++ b/Foreign/CUDA/Types.chs @@ -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 diff --git a/README.markdown b/README.markdown new file mode 100644 --- /dev/null +++ b/README.markdown @@ -0,0 +1,20 @@ +Haskell FFI Bindings to CUDA +============================ + +[![Build status](https://travis-ci.org/tmcdonell/cuda.svg?branch=master)](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. + + + +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: + + + diff --git a/Setup.hs b/Setup.hs --- a/Setup.hs +++ b/Setup.hs @@ -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 diff --git a/WINDOWS.markdown b/WINDOWS.markdown new file mode 100644 --- /dev/null +++ b/WINDOWS.markdown @@ -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: + +> + +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 + diff --git a/cbits/stubs.c b/cbits/stubs.c --- a/cbits/stubs.c +++ b/cbits/stubs.c @@ -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 diff --git a/cbits/stubs.h b/cbits/stubs.h --- a/cbits/stubs.h +++ b/cbits/stubs.h @@ -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 #include @@ -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 diff --git a/changelog.md b/changelog.md deleted file mode 100644 --- a/changelog.md +++ /dev/null @@ -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 - diff --git a/cuda.cabal b/cuda.cabal --- a/cuda.cabal +++ b/cuda.cabal @@ -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: + . + + . 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