packages feed

cuda (empty) → 0.1

raw patch · 81 files changed

+8861/−0 lines, 81 filesdep +basedep +bytestringdep +extensible-exceptionssetup-changedbinary-added

Dependencies added: base, bytestring, extensible-exceptions, haskell98

Files

+ Foreign/CUDA.hs view
@@ -0,0 +1,19 @@+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Top level bindings. By default, expose the C-for-CUDA runtime API bindings,+-- as they are slightly more user friendly.+--+--------------------------------------------------------------------------------++module Foreign.CUDA+  (+    module Foreign.CUDA.Runtime+  )+  where++import Foreign.CUDA.Runtime+
+ Foreign/CUDA/Driver.hs view
@@ -0,0 +1,30 @@+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Driver+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Top level bindings to CUDA driver API+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver+  (+    module Foreign.CUDA.Driver.Context,+    module Foreign.CUDA.Driver.Device,+    module Foreign.CUDA.Driver.Error,+    module Foreign.CUDA.Driver.Exec,+    module Foreign.CUDA.Driver.Marshal,+    module Foreign.CUDA.Driver.Module,+    module Foreign.CUDA.Driver.Utils+  )+  where++import Foreign.CUDA.Driver.Context+import Foreign.CUDA.Driver.Device+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Driver.Exec+import Foreign.CUDA.Driver.Marshal+import Foreign.CUDA.Driver.Module+import Foreign.CUDA.Driver.Utils+
+ Foreign/CUDA/Driver/Context.chs view
@@ -0,0 +1,145 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Driver.Context+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Context management for low-level driver interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Context+  (+    Context, ContextFlag(..),+    create, attach, detach, destroy, current, pop, push, sync+  )+  where++#include <cuda.h>+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Driver.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 #}}+++-- |+-- Context creation flags+--+{# enum CUctx_flags as ContextFlag+    { underscoreToCase }+    with prefix="CU_CTX" deriving (Eq, Show) #}+++--------------------------------------------------------------------------------+-- Context management+--------------------------------------------------------------------------------++-- |+-- Create a new CUDA context and associate it with the calling thread+--+create :: Device -> [ContextFlag] -> IO Context+create dev flags = resultIfOk =<< cuCtxCreate flags dev++{# 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.+--+attach :: Context -> [ContextFlag] -> IO ()+attach ctx flags = nothingIfOk =<< cuCtxAttach ctx flags++{# fun unsafe cuCtxAttach+  { withCtx*        `Context'+  , combineBitMasks `[ContextFlag]' } -> `Status' cToEnum #}+  where withCtx = with . useContext+++-- |+-- Detach the context, and destroy if no longer used+--+detach :: Context -> IO ()+detach ctx = nothingIfOk =<< cuCtxDetach ctx++{# 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).+--+destroy :: Context -> IO ()+destroy ctx = nothingIfOk =<< cuCtxDestroy ctx++{# fun unsafe cuCtxDestroy+  { useContext `Context' } -> `Status' cToEnum #}+++-- |+-- Return the device of the currently active context+--+current :: IO Device+current = resultIfOk =<< 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/detach). If successful, the new+-- context is returned, and the old may be attached to a different CPU.+--+pop :: IO Context+pop = resultIfOk =<< 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.+--+push :: Context -> IO ()+push ctx = nothingIfOk =<< cuCtxPushCurrent ctx++{# fun unsafe cuCtxPushCurrent+  { useContext `Context' } -> `Status' cToEnum #}+++-- |+-- Block until the device has completed all preceding requests+--+sync :: IO ()+sync = nothingIfOk =<< cuCtxSynchronize++{# fun unsafe cuCtxSynchronize+  { } -> `Status' cToEnum #}+
+ Foreign/CUDA/Driver/Device.chs view
@@ -0,0 +1,215 @@+{-# LANGUAGE ForeignFunctionInterface, EmptyDataDecls #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Driver.Device+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Device management for low-level driver interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Device+  (+    Device(..), -- should be exported abstractly+    DeviceProperties(..), DeviceAttribute(..), InitFlag,++    initialise, capability, device, attribute, count, name, props, totalMem+  )+  where++#include <cuda.h>+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Internal.C2HS+import Foreign.CUDA.Internal.Offsets++-- System+import Foreign+import Foreign.C+import Control.Monad            (liftM)+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++newtype Device = Device { useDevice :: {# type CUdevice #}}+++-- |+-- Device attributes+--+{# enum CUdevice_attribute as DeviceAttribute+    { underscoreToCase }+    with prefix="CU_DEVICE_ATTRIBUTE" deriving (Eq, Show) #}++{# pointer *CUdevprop as ^ foreign -> DeviceProperties nocode #}+++-- |+-- Properties of the compute device+--+data DeviceProperties = DeviceProperties+  {+    maxThreadsPerBlock  :: Int,           -- ^ Maximum number of threads per block+    maxThreadsDim       :: (Int,Int,Int), -- ^ Maximum size of each dimension of a block+    maxGridSize         :: (Int,Int,Int), -- ^ Maximum size of each dimension of a grid+    sharedMemPerBlock   :: Int,           -- ^ Shared memory available per block in bytes+    totalConstantMemory :: Int,           -- ^ Constant memory available on device in bytes+    warpSize            :: Int,           -- ^ Warp size in threads (SIMD width)+    memPitch            :: Int,           -- ^ Maximum pitch in bytes allowed by memory copies+    regsPerBlock        :: Int,           -- ^ 32-bit registers available per block+    clockRate           :: Int,           -- ^ Clock frequency in kilohertz+    textureAlign        :: Int            -- ^ Alignment requirement for textures+  }+  deriving (Show)++instance Storable DeviceProperties where+  sizeOf _    = {#sizeof CUdevprop#}+  alignment _ = alignment (undefined :: Ptr ())++  peek p      = do+    tb <- cIntConv `fmap` {#get CUdevprop.maxThreadsPerBlock#} p+    sm <- cIntConv `fmap` {#get CUdevprop.sharedMemPerBlock#} p+    cm <- cIntConv `fmap` {#get CUdevprop.totalConstantMemory#} p+    ws <- cIntConv `fmap` {#get CUdevprop.SIMDWidth#} p+    mp <- cIntConv `fmap` {#get CUdevprop.memPitch#} p+    rb <- cIntConv `fmap` {#get CUdevprop.regsPerBlock#} p+    cl <- cIntConv `fmap` {#get CUdevprop.clockRate#} p+    ta <- cIntConv `fmap` {#get CUdevprop.textureAlign#} p++    (t1:t2:t3:_) <- map cIntConv `fmap` peekArray 3 (p `plusPtr` devMaxThreadDimOffset' :: Ptr CInt)+    (g1:g2:g3:_) <- map cIntConv `fmap` peekArray 3 (p `plusPtr` devMaxGridSizeOffset'  :: Ptr CInt)++    return DeviceProperties+      {+        maxThreadsPerBlock  = tb,+        maxThreadsDim       = (t1,t2,t3),+        maxGridSize         = (g1,g2,g3),+        sharedMemPerBlock   = sm,+        totalConstantMemory = cm,+        warpSize            = ws,+        memPitch            = mp,+        regsPerBlock        = rb,+        clockRate           = cl,+        textureAlign        = ta+      }++-- |+-- Possible option flags for CUDA initialisation. Dummy instance until the API+-- exports actual option values.+--+data InitFlag++instance Enum InitFlag where+++--------------------------------------------------------------------------------+-- Initialisation+--------------------------------------------------------------------------------++-- |+-- Initialise the CUDA driver API. Must be called before any other driver+-- function.+--+initialise :: [InitFlag] -> IO ()+initialise flags = nothingIfOk =<< cuInit flags++{# fun unsafe cuInit+  { combineBitMasks `[InitFlag]' } -> `Status' cToEnum #}+++--------------------------------------------------------------------------------+-- Device Management+--------------------------------------------------------------------------------++-- |+-- Return the compute compatibility revision supported by the device+--+capability :: Device -> IO Double+capability dev =+  (\(s,a,b) -> resultIfOk (s,cap a b)) =<< cuDeviceComputeCapability dev+  where+    cap a b = let a' = fromIntegral a in+              let b' = fromIntegral b in+              a' + b' / max 10 (10^ ((ceiling . logBase 10) b' :: Int))++{# fun unsafe cuDeviceComputeCapability+  { alloca-   `Int'    peekIntConv*+  , alloca-   `Int'    peekIntConv*+  , useDevice `Device'              } -> `Status' cToEnum #}+++-- |+-- Return a device handle+--+device :: Int -> IO Device+device d = resultIfOk =<< cuDeviceGet d++{# fun unsafe cuDeviceGet+  { alloca-  `Device' dev*+  , cIntConv `Int'           } -> `Status' cToEnum #}+  where dev = liftM Device . peek+++-- |+-- Return the selected attribute for the given device+--+attribute :: Device -> DeviceAttribute -> IO Int+attribute d a = resultIfOk =<< cuDeviceGetAttribute a d++{# fun unsafe cuDeviceGetAttribute+  { alloca-   `Int'             peekIntConv*+  , cFromEnum `DeviceAttribute'+  , useDevice `Device'                       } -> `Status' cToEnum #}+++-- |+-- Return the number of device with compute capability > 1.0+--+count :: IO Int+count = resultIfOk =<< cuDeviceGetCount++{# fun unsafe cuDeviceGetCount+  { alloca- `Int' peekIntConv* } -> `Status' cToEnum #}+++-- |+-- Name of the device+--+name :: Device -> IO String+name d = resultIfOk =<< cuDeviceGetName d++{# fun unsafe cuDeviceGetName+  { allocaS-  `String'& peekS*+  , useDevice `Device'         } -> `Status' cToEnum #}+  where+    len            = 512+    allocaS a      = allocaBytes len $ \p -> a (p, cIntConv len)+    peekS s _      = peekCString s+++-- |+-- Return the properties of the selected device+--+props :: Device -> IO DeviceProperties+props d = resultIfOk =<< cuDeviceGetProperties d++{# fun unsafe cuDeviceGetProperties+  { alloca-   `DeviceProperties' peek*+  , useDevice `Device'                 } -> `Status' cToEnum #}+++-- |+-- Total memory available on the device (bytes)+--+totalMem :: Device -> IO Int+totalMem d = resultIfOk =<< cuDeviceTotalMem d++{# fun unsafe cuDeviceTotalMem+  { alloca-   `Int' peekIntConv*+  , useDevice `Device'           } -> `Status' cToEnum #}+
+ Foreign/CUDA/Driver/Error.chs view
@@ -0,0 +1,129 @@+{-# LANGUAGE DeriveDataTypeable #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Driver.Error+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Error handling+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Error+  where+++-- System+import Data.Typeable+import Control.Exception.Extensible++#include <cuda.h>+{# context lib="cuda" #}+++--------------------------------------------------------------------------------+-- Return Status+--------------------------------------------------------------------------------++--+-- Error Codes+--+{# enum CUresult as Status+    { underscoreToCase+    , CUDA_SUCCESS as Success+    , CUDA_ERROR_NO_BINARY_FOR_GPU as NoBinaryForGPU }+    with prefix="CUDA_ERROR" deriving (Eq, Show) #}+++-- |+-- Return a descriptive error string associated with a particular error code+--+describe :: Status -> String+describe Success                     = "no error"+describe InvalidValue                = "invalid argument"+describe OutOfMemory                 = "out of memory"+describe NotInitialized              = "driver not initialised"+describe Deinitialized               = "driver deinitialised"+describe NoDevice                    = "no CUDA-capable device is available"+describe InvalidDevice               = "invalid device ordinal"+describe InvalidImage                = "invalid kernel image"+describe InvalidContext              = "invalid context handle"+describe ContextAlreadyCurrent       = "context already current"+describe MapFailed                   = "map failed"+describe UnmapFailed                 = "unmap failed"+describe ArrayIsMapped               = "array is mapped"+describe AlreadyMapped               = "already mapped"+describe NoBinaryForGPU              = "no binary available for this GPU"+describe AlreadyAcquired             = "resource already acquired"+describe NotMapped                   = "not mapped"+describe InvalidSource               = "invalid source"+describe FileNotFound                = "file not found"+describe InvalidHandle               = "invalid handle"+describe NotFound                    = "not found"+describe NotReady                    = "device not ready"+describe LaunchFailed                = "unspecified launch failure"+describe LaunchOutOfResources        = "too many resources requested for launch"+describe LaunchTimeout               = "the launch timed out and was terminated"+describe LaunchIncompatibleTexturing = "launch with incompatible texturing"+describe Unknown                     = "unknown error"+++--------------------------------------------------------------------------------+-- Exceptions+--------------------------------------------------------------------------------++data CUDAException+  = ExitCode Status+  | UserError String+  deriving Typeable++instance Exception CUDAException++instance Show CUDAException where+  showsPrec _ (ExitCode  s) = showString ("CUDA Exception: " ++ describe s)+  showsPrec _ (UserError s) = showString ("CUDA Exception: " ++ s)+++-- |+-- Raise a CUDAException in the IO Monad+--+cudaError :: String -> IO a+cudaError s = throwIO (UserError s)+++-- |+-- Run a CUDA computation+--+{-+runCUDA f = runEMT $ do+  f `catchWithSrcLoc` \l e -> lift (handle l e)+  where+    handle :: CallTrace -> CUDAException -> IO ()+    handle l e = putStrLn $ showExceptionWithTrace l e+-}++--------------------------------------------------------------------------------+-- Helper Functions+--------------------------------------------------------------------------------++-- |+-- Return the results of a function on successful execution, otherwise throw an+-- exception with an error string associated with the return code+--+resultIfOk :: (Status, a) -> IO a+resultIfOk (status,result) =+    case status of+        Success -> return  result+        _       -> throwIO (ExitCode status)+++-- |+-- Throw an exception with an error string associated with an unsuccessful+-- return code, otherwise return unit.+--+nothingIfOk :: Status -> IO ()+nothingIfOk status =+    case status of+        Success -> return  ()+        _       -> throwIO (ExitCode status)+
+ Foreign/CUDA/Driver/Event.chs view
@@ -0,0 +1,127 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Driver.Event+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Event management for low-level driver interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Event+  (+    Event, EventFlag(..),+    create, destroy, elapsedTime, query, record, block+  )+  where++#include <cuda.h>+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Internal.C2HS+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Driver.Stream               (Stream(..))++-- System+import Foreign+import Foreign.C+import Control.Monad                            (liftM)+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- Events+--+newtype Event = Event { useEvent :: {# type CUevent #}}+++-- |+-- Event creation flags+--+{# enum CUevent_flags as EventFlag+    { underscoreToCase }+    with prefix="CU_EVENT" deriving (Eq, Show) #}+++--------------------------------------------------------------------------------+-- Event management+--------------------------------------------------------------------------------++-- |+-- Create a new event+--+create :: [EventFlag] -> IO Event+create flags = resultIfOk =<< cuEventCreate flags++{# fun unsafe cuEventCreate+  { alloca-         `Event'       peekEvt*+  , combineBitMasks `[EventFlag]'          } -> `Status' cToEnum #}+  where peekEvt = liftM Event . peek+++-- |+-- Destroy an event+--+destroy :: Event -> IO ()+destroy ev = nothingIfOk =<< cuEventDestroy ev++{# fun unsafe cuEventDestroy+  { useEvent `Event' } -> `Status' cToEnum #}+++-- |+-- Determine the elapsed time (in milliseconds) between two events+--+elapsedTime :: Event -> Event -> IO Float+elapsedTime ev1 ev2 = resultIfOk =<< cuEventElapsedTime ev1 ev2++{# fun unsafe cuEventElapsedTime+  { alloca-  `Float' peekFloatConv*+  , useEvent `Event'+  , useEvent `Event'                } -> `Status' cToEnum #}+++-- |+-- Determines if a event has actually been recorded+--+query :: Event -> IO Bool+query ev =+  cuEventQuery ev >>= \rv ->+  case rv of+    Success  -> return True+    NotReady -> return False+    _        -> resultIfOk (rv,undefined)++{# fun unsafe cuEventQuery+  { useEvent `Event' } -> `Status' cToEnum #}+++-- |+-- Record an event once all operations in the current context (or optionally+-- specified stream) have completed. This operation is asynchronous.+--+record :: Event -> Maybe Stream -> IO ()+record ev mst =+  nothingIfOk =<< case mst of+    Just st -> cuEventRecord ev st+    Nothing -> cuEventRecord ev (Stream nullPtr)++{# fun unsafe cuEventRecord+  { useEvent  `Event'+  , useStream `Stream' } -> `Status' cToEnum #}+++-- |+-- Wait until the event has been recorded+--+block :: Event -> IO ()+block ev = nothingIfOk =<< cuEventSynchronize ev++{# fun unsafe cuEventSynchronize+  { useEvent `Event' } -> `Status' cToEnum #}+
+ Foreign/CUDA/Driver/Exec.chs view
@@ -0,0 +1,165 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Driver.Exec+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Kernel execution control for low-level driver interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Exec+  (+    Fun(Fun),  -- need to export the data constructor for use by Module )=+    FunParam(..), FunAttribute(..),+    requires, setBlockShape, setSharedSize, setParams, launch+  )+  where++#include <cuda.h>+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Internal.C2HS+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Driver.Stream               (Stream(..))++-- System+import Foreign+import Foreign.C+import Control.Monad                            (zipWithM_)+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- A @__global__@ device function+--+newtype Fun = Fun { useFun :: {# type CUfunction #}}+++-- |+-- Function attributes+--+{# enum CUfunction_attribute as FunAttribute+    { underscoreToCase+    , MAX_THREADS_PER_BLOCK as MaxKernelThreadsPerBlock }+    with prefix="CU_FUNC_ATTRIBUTE" deriving (Eq, Show) #}+++-- |+-- Kernel function parameters+--+data Storable a => FunParam a+    = IArg Int+    | FArg Float+    | VArg a+--  | TArg Texture++--------------------------------------------------------------------------------+-- Execution Control+--------------------------------------------------------------------------------++-- |+-- Returns the value of the selected attribute requirement for the given kernel+--+requires :: Fun -> FunAttribute -> IO Int+requires fn att = resultIfOk =<< cuFuncGetAttribute att fn++{# fun unsafe cuFuncGetAttribute+  { alloca-   `Int'          peekIntConv*+  , cFromEnum `FunAttribute'+  , useFun    `Fun'                       } -> `Status' cToEnum #}+++-- |+-- Specify the (x,y,z) dimensions of the thread blocks that are created when the+-- given kernel function is lanched+--+setBlockShape :: Fun -> (Int,Int,Int) -> IO ()+setBlockShape fn (x,y,z) = nothingIfOk =<< cuFuncSetBlockShape fn x y z++{# 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+--+setSharedSize :: Fun -> Integer -> IO ()+setSharedSize fn bytes = nothingIfOk =<< cuFuncSetSharedSize fn bytes++{# fun unsafe cuFuncSetSharedSize+  { useFun   `Fun'+  , cIntConv `Integer' } -> `Status' cToEnum #}+++-- |+-- 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'.+--+launch :: Fun -> (Int,Int) -> Maybe Stream -> IO ()+launch fn (w,h) mst =+  nothingIfOk =<< case mst of+    Nothing -> cuLaunchGridAsync fn w h (Stream nullPtr)+    Just st -> cuLaunchGridAsync fn w h st++{# fun unsafe cuLaunchGridAsync+  { useFun    `Fun'+  ,           `Int'+  ,           `Int'+  , useStream `Stream' } -> `Status' cToEnum #}+++--------------------------------------------------------------------------------+-- Kernel function parameters+--------------------------------------------------------------------------------++-- |+-- Set the parameters that will specified next time the kernel is invoked+--+setParams :: Storable a => Fun -> [FunParam a] -> IO ()+setParams fn prs = do+  zipWithM_ (set fn) offsets prs+  nothingIfOk =<< cuParamSetSize fn (last offsets)+  where+    offsets = scanl (\a b -> a + size b) 0 prs++    size (IArg _)    = sizeOf (undefined::CUInt)+    size (FArg _)    = sizeOf (undefined::CFloat)+    size (VArg v)    = sizeOf v++    set f o (IArg v) = nothingIfOk =<< cuParamSeti f o v+    set f o (FArg v) = nothingIfOk =<< cuParamSetf f o v+    set f o (VArg v) = with v $ \p -> (nothingIfOk =<< cuParamSetv f o p (sizeOf v))+++{# fun unsafe cuParamSetSize+  { useFun `Fun'+  ,        `Int' } -> `Status' cToEnum #}++{# fun unsafe cuParamSeti+  { useFun `Fun'+  ,        `Int'+  ,        `Int' } -> `Status' cToEnum #}++{# fun unsafe cuParamSetf+  { useFun `Fun'+  ,        `Int'+  ,        `Float' } -> `Status' cToEnum #}++{# fun unsafe cuParamSetv+  `Storable a' =>+  { useFun  `Fun'+  ,         `Int'+  , castPtr `Ptr a'+  ,         `Int'   } -> `Status' cToEnum #}+
+ Foreign/CUDA/Driver/Marshal.chs view
@@ -0,0 +1,419 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Driver.Marshal+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Memory management for low-level driver interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Marshal+  (+    -- * Host Allocation+    HostPtr(..), AllocFlag(..),+    withHostPtr, mallocHostArray, freeHost, nullHostPtr,++    -- * Device Allocation+    DevicePtr,+    mallocArray, allocaArray, free, nullDevPtr,++    -- * Marshalling+    peekArray, peekArrayAsync, peekListArray,+    pokeArray, pokeArrayAsync, pokeListArray,+               copyArrayAsync,++    -- * Combined Allocation and Marshalling+    newListArray, withListArray, withListArrayLen,++    -- * Utility+    memset, getDevicePtr+  )+  where++#include <cuda.h>+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Internal.C2HS+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Driver.Stream               (Stream(..))++-- System+import Unsafe.Coerce+import Control.Monad                            (liftM)+import Control.Exception.Extensible++import Foreign.C+import Foreign.Ptr+import Foreign.Storable+import qualified Foreign.Marshal as F++#c+typedef enum CUmemhostalloc_option_enum {+    CU_MEMHOSTALLOC_OPTION_PORTABLE       = CU_MEMHOSTALLOC_PORTABLE,+    CU_MEMHOSTALLOC_OPTION_DEVICE_MAPPED  = CU_MEMHOSTALLOC_DEVICEMAP,+    CU_MEMHOSTALLOC_OPTION_WRITE_COMBINED = CU_MEMHOSTALLOC_WRITECOMBINED+} CUmemhostalloc_option;+#endc++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- A reference to memory allocated on the device+--+newtype DevicePtr a = DevicePtr { useDevicePtr :: {# type CUdeviceptr #}}++instance Storable (DevicePtr a) where+  sizeOf _      = sizeOf    (undefined :: {# type CUdeviceptr #})+  alignment _   = alignment (undefined :: {# type CUdeviceptr #})+  peek p        = DevicePtr `fmap` peek (castPtr p)+  poke p v      = poke (castPtr p) (useDevicePtr v)+++-- |+-- A reference to memory on the host that is page-locked and directly-accessible+-- from the device. Since the memory can be accessed directly, it can be read or+-- written at a much higher bandwidth than pageable memory from the traditional+-- malloc.+--+-- The driver automatically accelerates calls to functions such as `memcpy'+-- which reference page-locked memory.+--+newtype HostPtr a = HostPtr { useHostPtr :: Ptr a }++-- |+-- Unwrap a host pointer and execute a computation using the base pointer object+--+withHostPtr :: HostPtr a -> (Ptr a -> IO b) -> IO b+withHostPtr p f = f (useHostPtr p)++-- |+-- Options for host allocation+--+{# enum CUmemhostalloc_option as AllocFlag+    { underscoreToCase }+    with prefix="CU_MEMHOSTALLOC_OPTION" deriving (Eq, Show) #}++--------------------------------------------------------------------------------+-- Host Allocation+--------------------------------------------------------------------------------++-- |+-- Allocate a section of linear memory on the host which is page-locked and+-- directly accessible from the device. The storage is sufficient to hold the+-- given number of elements of a storable type.+--+-- Note that since the amount of pageable memory is thusly reduced, overall+-- system performance may suffer. This is best used sparingly to allocate+-- staging areas for data exchange.+--+mallocHostArray :: Storable a => [AllocFlag] -> Int -> IO (HostPtr a)+mallocHostArray flags = doMalloc undefined+  where+    doMalloc :: Storable a' => a' -> Int -> IO (HostPtr a')+    doMalloc x n = resultIfOk =<< cuMemHostAlloc (n * sizeOf x) flags++{# fun unsafe cuMemHostAlloc+  { alloca'-        `HostPtr a'   peekHP*+  ,                 `Int'+  , combineBitMasks `[AllocFlag]'         } -> `Status' cToEnum #}+  where+    alloca'  = F.alloca+    peekHP p = (HostPtr . castPtr) `fmap` peek p+++-- |+-- Free a section of page-locked host memory+--+freeHost :: HostPtr a -> IO ()+freeHost p = nothingIfOk =<< cuMemFreeHost p++{# fun unsafe cuMemFreeHost+  { useHP `HostPtr a' } -> `Status' cToEnum #}+  where+    useHP = castPtr . useHostPtr+++-- |+-- The constant 'nullHostPtr' contains the distinguished memory location that is+-- not associated with a valid memory location+--+nullHostPtr :: HostPtr a+nullHostPtr =  HostPtr nullPtr+++--------------------------------------------------------------------------------+-- Device Allocation+--------------------------------------------------------------------------------++-- |+-- Allocate a section of linear memory on the device, and return a reference to+-- it. The memory is sufficient to hold the given number of elements of storable+-- type. It is suitably aligned for any type, and is not cleared.+--+mallocArray :: Storable a => Int -> IO (DevicePtr a)+mallocArray = doMalloc undefined+  where+    doMalloc :: Storable a' => a' -> Int -> IO (DevicePtr a')+    doMalloc x n = resultIfOk =<< cuMemAlloc (n * sizeOf x)++{# fun unsafe cuMemAlloc+  { alloca'- `DevicePtr a' peekDP*+  ,          `Int'               } -> `Status' cToEnum #}+  where+    alloca' = F.alloca+    peekDP  = liftM DevicePtr . peek+++-- |+-- Execute a computation on the device, passing a pointer to a temporarily+-- allocated block of memory sufficient to hold the given number of elements of+-- storable type. The memory is freed when the computation terminates (normally+-- or via an exception), so the pointer must not be used after this.+--+-- Note that kernel launches can be asynchronous, so you may want to add a+-- synchronisation point using 'sync' as part of the computation.+--+allocaArray :: Storable a => Int -> (DevicePtr a -> IO b) -> IO b+allocaArray n = bracket (mallocArray n) free+++-- |+-- Release a section of device memory+--+free :: DevicePtr a -> IO ()+free dp = nothingIfOk =<< cuMemFree dp++{# fun unsafe cuMemFree+  { useDevicePtr `DevicePtr a' } -> `Status' cToEnum #}+++-- |+-- The constant 'nullDevPtr' contains the distinguished memory location that is+-- not associated with a valid memory location+--+nullDevPtr :: DevicePtr a+nullDevPtr =  DevicePtr 0+++--------------------------------------------------------------------------------+-- Marshalling+--------------------------------------------------------------------------------++-- |+-- Copy a number of elements from the device to host memory. This is a+-- synchronous operation+--+peekArray :: Storable a => Int -> DevicePtr a -> Ptr a -> IO ()+peekArray n dptr hptr = doPeek undefined dptr+  where+    doPeek :: Storable a' => a' -> DevicePtr a' -> IO ()+    doPeek x _ = nothingIfOk =<< cuMemcpyDtoH hptr dptr (n * sizeOf x)++{# fun unsafe cuMemcpyDtoH+  { castPtr      `Ptr a'+  , useDevicePtr `DevicePtr a'+  ,              `Int'         } -> `Status' cToEnum #}+++-- |+-- Copy memory from the device asynchronously, possibly associated with a+-- particular stream. The destination host memory must be page-locked.+--+peekArrayAsync :: Storable a => Int -> DevicePtr a -> HostPtr a -> Maybe Stream -> IO ()+peekArrayAsync n dptr hptr mst = doPeek undefined dptr+  where+    doPeek :: Storable a' => a' -> DevicePtr a' -> IO ()+    doPeek x _ =+      nothingIfOk =<< case mst of+        Nothing -> cuMemcpyDtoHAsync hptr dptr (n * sizeOf x) (Stream nullPtr)+        Just st -> cuMemcpyDtoHAsync hptr dptr (n * sizeOf x) st++{# fun unsafe cuMemcpyDtoHAsync+  { useHP        `HostPtr a'+  , useDevicePtr `DevicePtr a'+  ,              `Int'+  , useStream    `Stream'  } -> `Status' cToEnum #}+  where+    useHP = castPtr . useHostPtr+++-- |+-- Copy a number of elements from the device into a new Haskell list. Note that+-- this requires two memory copies: firstly from the device into a heap+-- allocated array, and from there marshalled into a list.+--+peekListArray :: Storable a => Int -> DevicePtr a -> IO [a]+peekListArray n dptr =+  F.allocaArray n $ \p -> do+    peekArray   n dptr p+    F.peekArray n p+++-- |+-- Copy a number of elements onto the device. This is a synchronous operation+--+pokeArray :: Storable a => Int -> Ptr a -> DevicePtr a -> IO ()+pokeArray n hptr dptr = doPoke undefined dptr+  where+    doPoke :: Storable a' => a' -> DevicePtr a' -> IO ()+    doPoke x _ = nothingIfOk =<< cuMemcpyHtoD dptr hptr (n * sizeOf x)++{# fun unsafe cuMemcpyHtoD+  { useDevicePtr `DevicePtr a'+  , castPtr      `Ptr a'+  ,              `Int'         } -> `Status' cToEnum #}+++-- |+-- Copy memory onto the device asynchronously, possibly associated with a+-- particular stream. The source host memory must be page-locked.+--+pokeArrayAsync :: Storable a => Int -> HostPtr a -> DevicePtr a -> Maybe Stream -> IO ()+pokeArrayAsync n hptr dptr mst = dopoke undefined dptr+  where+    dopoke :: Storable a' => a' -> DevicePtr a' -> IO ()+    dopoke x _ =+      nothingIfOk =<< case mst of+        Nothing -> cuMemcpyHtoDAsync dptr hptr (n * sizeOf x) (Stream nullPtr)+        Just st -> cuMemcpyHtoDAsync dptr hptr (n * sizeOf x) st++{# fun unsafe cuMemcpyHtoDAsync+  { useDevicePtr `DevicePtr a'+  , useHP        `HostPtr a'+  ,              `Int'+  , useStream    `Stream'      } -> `Status' cToEnum #}+  where+    useHP = castPtr . useHostPtr+++-- |+-- Write a list of storable elements into a device array. The device array must+-- be sufficiently large to hold the entire list. This requires two marshalling+-- operations.+--+pokeListArray :: Storable a => [a] -> DevicePtr a -> IO ()+pokeListArray xs dptr = F.withArrayLen xs $ \len p -> pokeArray len p dptr+++-- |+-- Copy the given number of elements from the first device array (source) to the+-- second (destination). The copied areas may not overlap. This operation is+-- asynchronous with respect to the host, but will never overlap with kernel+-- execution.+--+copyArrayAsync :: Storable a => Int -> DevicePtr a -> DevicePtr a -> IO ()+copyArrayAsync n = docopy undefined+  where+    docopy :: Storable a' => a' -> DevicePtr a' -> DevicePtr a' -> IO ()+    docopy x src dst = nothingIfOk =<< cuMemcpyDtoD dst src (n * sizeOf x)++{# fun unsafe cuMemcpyDtoD+  { useDevicePtr `DevicePtr a'+  , useDevicePtr `DevicePtr a'+  ,              `Int'         } -> `Status' cToEnum #}+++--------------------------------------------------------------------------------+-- Combined Allocation and Marshalling+--------------------------------------------------------------------------------++-- |+-- Write a list of storable elements into a newly allocated device array. Note+-- that this requires two memory copies: firstly from a Haskell list to a heap+-- allocated array, and from there onto the graphics device. The memory should+-- be 'free'd when no longer required.+--+newListArray :: Storable a => [a] -> IO (DevicePtr a)+newListArray xs =+  F.withArrayLen xs                     $ \len p ->+  bracketOnError (mallocArray len) free $ \d_xs  -> do+    pokeArray len p d_xs+    return d_xs+++-- |+-- Temporarily store a list of elements into a newly allocated device array. An+-- IO action is applied to to the array, the result of which is returned.+-- Similar to 'newListArray', this requires copying the data twice.+--+-- As with 'allocaArray', the memory is freed once the action completes, so you+-- should not return the pointer from the action, and be wary of asynchronous+-- kernel execution.+--+withListArray :: Storable a => [a] -> (DevicePtr a -> IO b) -> IO b+withListArray xs = withListArrayLen xs . const+++-- |+-- A variant of 'withListArray' which also supplies the number of elements in+-- the array to the applied function+--+withListArrayLen :: Storable a => [a] -> (Int -> DevicePtr a -> IO b) -> IO b+withListArrayLen xs f =+  allocaArray   len $ \d_xs -> do+  F.allocaArray len $ \h_xs -> do+    F.pokeArray h_xs xs+    pokeArray len h_xs d_xs+  f len d_xs+  where+    len = length xs+++--------------------------------------------------------------------------------+-- Utility+--------------------------------------------------------------------------------++-- |+-- Set a number of data elements to the specified value, which may be either 8-,+-- 16-, or 32-bits wide.+--+memset :: Storable a => DevicePtr a -> Int -> a -> IO ()+memset dptr n val = case sizeOf val of+    1 -> nothingIfOk =<< cuMemsetD8  dptr val n+    2 -> nothingIfOk =<< cuMemsetD16 dptr val n+    4 -> nothingIfOk =<< cuMemsetD32 dptr val n+    _ -> cudaError "can only memset 8-, 16-, and 32-bit values"++--+-- We use unsafe coerce below to reinterpret the bits of the value to memset as,+-- into the integer type required by the setting functions.+--+{# fun unsafe cuMemsetD8+  { useDevicePtr `DevicePtr a'+  , unsafeCoerce `a'+  ,              `Int'         } -> `Status' cToEnum #}++{# fun unsafe cuMemsetD16+  { useDevicePtr `DevicePtr a'+  , unsafeCoerce `a'+  ,              `Int'         } -> `Status' cToEnum #}++{# fun unsafe cuMemsetD32+  { useDevicePtr `DevicePtr a'+  , unsafeCoerce `a'+  ,              `Int'         } -> `Status' cToEnum #}+++-- |+-- Return the device pointer associated with a mapped, pinned host buffer, which+-- was allocated with the 'DeviceMapped' option by 'mallocHostArray'.+--+-- Currently, no options are supported and this must be empty.+--+getDevicePtr :: [AllocFlag] -> HostPtr a -> IO (DevicePtr a)+getDevicePtr flags hp = resultIfOk =<< cuMemHostGetDevicePointer hp flags++{# fun unsafe cuMemHostGetDevicePointer+  { alloca'-        `DevicePtr a' peekDP*+  , useHP           `HostPtr a'+  , combineBitMasks `[AllocFlag]'         } -> `Status' cToEnum #}+  where+    alloca' = F.alloca+    useHP   = castPtr . useHostPtr+    peekDP  = liftM DevicePtr . peek+
+ Foreign/CUDA/Driver/Module.chs view
@@ -0,0 +1,182 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Driver.Module+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Module management for low-level driver interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Module+  (+    Module,+    JITOption(..), JITTarget(..), JITResult(..),+    getFun, loadFile, loadData, loadDataEx, unload+  )+  where++#include <cuda.h>+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Internal.C2HS+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Driver.Exec++-- 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+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- A reference to a Module object, containing collections of device functions+--+newtype Module = Module { useModule :: {# type CUmodule #}}+++-- |+-- 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             JITTarget -- ^ compilation target, otherwise determined from context+--  | FallbackStrategy   JITFallback+  deriving (Show)++-- |+-- Results of online compilation+--+data JITResult = JITResult+  {+    jitTime     :: Float,       -- ^ milliseconds spent compiling PTX+    jitInfoLog  :: ByteString,  -- ^ information about PTX asembly+    jitErrorLog :: ByteString   -- ^ compilation errors+  }+  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 }+    with prefix="CU_PREFER" deriving (Eq, Show) #}+++--------------------------------------------------------------------------------+-- Module management+--------------------------------------------------------------------------------++-- |+-- Returns a function handle+--+getFun :: Module -> String -> IO Fun+getFun mdl fn = resultIfOk =<< cuModuleGetFunction mdl fn++{# fun unsafe cuModuleGetFunction+  { alloca-      `Fun'    peekFun*+  , useModule    `Module'+  , withCString* `String'          } -> `Status' cToEnum #}+  where peekFun = liftM Fun . peek+++-- |+-- 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+--+loadFile :: String -> IO Module+loadFile ptx = resultIfOk =<< cuModuleLoad ptx++{# fun unsafe cuModuleLoad+  { alloca-      `Module' peekMod*+  , withCString* `String'          } -> `Status' cToEnum #}+  where peekMod = liftM Module . peek+++-- |+-- Load the contents of the given image into a new module, and load that module+-- into the current context. The image (typically) is the contents of a cubin or+-- ptx file as a NULL-terminated string.+--+loadData :: ByteString -> IO Module+loadData img = resultIfOk =<< cuModuleLoadData img++{# fun unsafe cuModuleLoadData+  { alloca- `Module'     peekMod*+  , useBS*  `ByteString'          } -> ` Status' cToEnum #}+  where+    peekMod      = liftM Module . peek+    useBS bs act = B.useAsCString bs $ \p -> act (castPtr p)+++-- |+-- Load a module with online compiler options. The actual attributes of the+-- compiled kernel can be probed using `requirements'.+--+loadDataEx :: ByteString -> [JITOption] -> IO (Module, JITResult)+loadDataEx img options =+  allocaArray logSize $ \p_ilog ->+  allocaArray logSize $ \p_elog ->+  let (opt,val) = unzip $+        [ (JIT_WALL_TIME, 0) -- must be first+        , (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 unpack options in++  withArray (map cFromEnum opt)    $ \p_opts ->+  withArray (map unsafeCoerce val) $ \p_vals -> do++  (s,mdl) <- cuModuleLoadDataEx img (length opt) p_opts p_vals+  infoLog <- B.packCString p_ilog+  errLog  <- B.packCString p_elog+  time    <- peek (castPtr p_vals)+  resultIfOk (s, (mdl, JITResult time infoLog errLog))++  where+    logSize = 2048++    unpack (MaxRegisters x)      = (JIT_MAX_REGISTERS, x)+    unpack (ThreadsPerBlock x)   = (JIT_THREADS_PER_BLOCK, x)+    unpack (OptimisationLevel x) = (JIT_OPTIMIZATION_LEVEL, x)+    unpack (Target x)            = (JIT_TARGET, fromEnum x)+++{# fun unsafe cuModuleLoadDataEx+  { alloca- `Module'       peekMod*+  , useBS*  `ByteString'+  ,         `Int'+  , id      `Ptr CInt'+  , id      `Ptr (Ptr ())'          } -> `Status' cToEnum #}+  where+    peekMod      = liftM Module . peek+    useBS bs act = B.useAsCString bs $ \p -> act (castPtr p)+++-- |+-- Unload a module from the current context+--+unload :: Module -> IO ()+unload m = nothingIfOk =<< cuModuleUnload m++{# fun unsafe cuModuleUnload+  { useModule `Module' } -> `Status' cToEnum #}+
+ Foreign/CUDA/Driver/Stream.chs view
@@ -0,0 +1,98 @@+{-# LANGUAGE ForeignFunctionInterface, EmptyDataDecls #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Driver.Stream+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Stream management for low-level driver interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Stream+  (+    Stream(..), StreamFlag,+    create, destroy, finished, block+  )+  where++#include <cuda.h>+{# context lib="cuda" #}++-- Friends+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Internal.C2HS++-- System+import Foreign+import Foreign.C+import Control.Monad                            (liftM)+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- A processing stream+--+newtype Stream = Stream { useStream :: {# type CUstream #}}+++-- |+-- Possible option flags for stream initialisation. Dummy instance until the API+-- exports actual option values.+--+data StreamFlag++instance Enum StreamFlag where++--------------------------------------------------------------------------------+-- Stream management+--------------------------------------------------------------------------------++-- |+-- Create a new stream+--+create :: [StreamFlag] -> IO Stream+create flags = resultIfOk =<< cuStreamCreate flags++{# fun unsafe cuStreamCreate+  { alloca-         `Stream'       peekStream*+  , combineBitMasks `[StreamFlag]'             } -> `Status' cToEnum #}+  where peekStream = liftM Stream . peek++-- |+-- Destroy a stream+--+destroy :: Stream -> IO ()+destroy st = nothingIfOk =<< cuStreamDestroy st++{# fun unsafe cuStreamDestroy+  { useStream `Stream' } -> `Status' cToEnum #}+++-- |+-- Check if all operations in the stream have completed+--+finished :: Stream -> IO Bool+finished st =+  cuStreamQuery st >>= \rv ->+  case rv of+    Success  -> return True+    NotReady -> return False+    _        -> resultIfOk (rv,undefined)++{# fun unsafe cuStreamQuery+  { useStream `Stream' } -> `Status' cToEnum #}+++-- |+-- Wait until the device has completed all operations in the Stream+--+block :: Stream -> IO ()+block st = nothingIfOk =<< cuStreamSynchronize st++{# fun unsafe cuStreamSynchronize+  { useStream `Stream' } -> `Status' cToEnum #}+
+ Foreign/CUDA/Driver/Utils.chs view
@@ -0,0 +1,39 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Driver.Utils+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Utility functions+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Driver.Utils+  (+    driverVersion+  )+  where++#include <cuda.h>+{# context lib="cuda" #}+++-- Friends+import Foreign.CUDA.Driver.Error+import Foreign.CUDA.Internal.C2HS++-- System+import Foreign+import Foreign.C+++-- |+-- Return the version number of the installed CUDA driver+--+driverVersion :: IO Int+driverVersion =  resultIfOk =<< cuDriverGetVersion++{# fun unsafe cuDriverGetVersion+  { alloca- `Int' peekIntConv* } -> `Status' cToEnum #}+
+ Foreign/CUDA/Internal/C2HS.hs view
@@ -0,0 +1,225 @@+--  C->Haskell Compiler: Marshalling library+--+--  Copyright (c) [1999...2005] Manuel M T Chakravarty+--+--  Redistribution and use in source and binary forms, with or without+--  modification, are permitted provided that the following conditions are met:+-- +--  1. Redistributions of source code must retain the above copyright notice,+--     this list of conditions and the following disclaimer. +--  2. Redistributions in binary form must reproduce the above copyright+--     notice, this list of conditions and the following disclaimer in the+--     documentation and/or other materials provided with the distribution. +--  3. The name of the author may not be used to endorse or promote products+--     derived from this software without specific prior written permission. +--+--  THIS SOFTWARE IS PROVIDED BY THE AUTHOR ``AS IS'' AND ANY EXPRESS OR+--  IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES+--  OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN+--  NO EVENT SHALL THE AUTHOR BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +--  SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED+--  TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR+--  PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF+--  LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING+--  NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS+--  SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.+--+--- Description ---------------------------------------------------------------+--+--  Language: Haskell 98+--+--  This module provides the marshaling routines for Haskell files produced by +--  C->Haskell for binding to C library interfaces.  It exports all of the+--  low-level FFI (language-independent plus the C-specific parts) together+--  with the C->HS-specific higher-level marshalling routines.+--++module Foreign.CUDA.Internal.C2HS (++--  -- * Re-export the language-independent component of the FFI +--  module Foreign,+--+--  -- * Re-export the C language component of the FFI+--  module CForeign,++  -- * Composite marshalling functions+  withCStringLenIntConv, peekCStringLenIntConv, withIntConv, withFloatConv,+  peekIntConv, peekFloatConv, withBool, peekBool, withEnum, peekEnum,++  -- * Conditional results using 'Maybe'+  nothingIf, nothingIfNull,++  -- * Bit masks+  combineBitMasks, containsBitMask, extractBitMasks,++  -- * Conversion between C and Haskell types+  cIntConv, cFloatConv, cToBool, cFromBool, cToEnum, cFromEnum+) where +++import Foreign+       hiding       (Word)+		    -- Should also hide the Foreign.Marshal.Pool exports in+		    -- compilers that export them+import CForeign++import Monad        (liftM)+++-- Composite marshalling functions+-- -------------------------------++-- Strings with explicit length+--+withCStringLenIntConv     :: String -> (CStringLen -> IO a) -> IO a+withCStringLenIntConv s f  = withCStringLen s $ \(p, n) -> f (p, cIntConv n)++peekCStringLenIntConv        :: CStringLen -> IO String+peekCStringLenIntConv (s, n)  = peekCStringLen (s, cIntConv n)++-- Marshalling of numerals+--++withIntConv   :: (Storable b, Integral a, Integral b) +	      => a -> (Ptr b -> IO c) -> IO c+withIntConv    = with . cIntConv++withFloatConv :: (Storable b, RealFloat a, RealFloat b) +	      => a -> (Ptr b -> IO c) -> IO c+withFloatConv  = with . cFloatConv++peekIntConv   :: (Storable a, Integral a, Integral b) +	      => Ptr a -> IO b+peekIntConv    = liftM cIntConv . peek++peekFloatConv :: (Storable a, RealFloat a, RealFloat b) +	      => Ptr a -> IO b+peekFloatConv  = liftM cFloatConv . peek++-- Passing Booleans by reference+--++withBool :: (Integral a, Storable a) => Bool -> (Ptr a -> IO b) -> IO b+withBool  = with . fromBool++peekBool :: (Integral a, Storable a) => Ptr a -> IO Bool+peekBool  = liftM toBool . peek+++-- Passing enums by reference+--++withEnum :: (Enum a, Integral b, Storable b) => a -> (Ptr b -> IO c) -> IO c+withEnum  = with . cFromEnum++peekEnum :: (Enum a, Integral b, Storable b) => Ptr b -> IO a+peekEnum  = liftM cToEnum . peek+++{-+-- Storing of 'Maybe' values+-- -------------------------++instance Storable a => Storable (Maybe a) where+  sizeOf    _ = sizeOf    (undefined :: Ptr ())+  alignment _ = alignment (undefined :: Ptr ())++  peek p = do+	     ptr <- peek (castPtr p)+	     if ptr == nullPtr+	       then return Nothing+	       else liftM Just $ peek ptr++  poke p v = do+	       ptr <- case v of+		        Nothing -> return nullPtr+			Just v' -> new v'+               poke (castPtr p) ptr+-}+++-- Conditional results using 'Maybe'+-- ---------------------------------++-- Wrap the result into a 'Maybe' type.+--+-- * the predicate determines when the result is considered to be non-existing,+--   ie, it is represented by `Nothing'+--+-- * the second argument allows to map a result wrapped into `Just' to some+--   other domain+--+nothingIf       :: (a -> Bool) -> (a -> b) -> a -> Maybe b+nothingIf p f x  = if p x then Nothing else Just $ f x++-- |Instance for special casing null pointers.+--+nothingIfNull :: (Ptr a -> b) -> Ptr a -> Maybe b+nothingIfNull  = nothingIf (== nullPtr)+++-- Support for bit masks+-- ---------------------++-- Given a list of enumeration values that represent bit masks, combine these+-- masks using bitwise disjunction.+--+combineBitMasks :: (Enum a, Bits b) => [a] -> b+combineBitMasks = foldl (.|.) 0 . map (fromIntegral . fromEnum)++-- Tests whether the given bit mask is contained in the given bit pattern+-- (i.e., all bits set in the mask are also set in the pattern).+--+containsBitMask :: (Bits a, Enum b) => a -> b -> Bool+bits `containsBitMask` bm = let bm' = fromIntegral . fromEnum $ bm+			    in+			    bm' .&. bits == bm'++-- |Given a bit pattern, yield all bit masks that it contains.+--+-- * This does *not* attempt to compute a minimal set of bit masks that when+--   combined yield the bit pattern, instead all contained bit masks are+--   produced.+--+extractBitMasks :: (Bits a, Enum b, Bounded b) => a -> [b]+extractBitMasks bits = +  [bm | bm <- [minBound..maxBound], bits `containsBitMask` bm]+++-- Conversion routines+-- -------------------++-- |Integral conversion+--+cIntConv :: (Integral a, Integral b) => a -> b+cIntConv  = fromIntegral++-- |Floating conversion+--+cFloatConv :: (RealFloat a, RealFloat b) => a -> b+cFloatConv  = realToFrac+-- As this conversion by default goes via `Rational', it can be very slow...+{-# RULES +  "cFloatConv/Float->Float"   forall (x::Float).  cFloatConv x = x;+  "cFloatConv/Double->Double" forall (x::Double). cFloatConv x = x+ #-}++-- |Obtain C value from Haskell 'Bool'.+--+cFromBool :: Num a => Bool -> a+cFromBool  = fromBool++-- |Obtain Haskell 'Bool' from C value.+--+cToBool :: Num a => a -> Bool+cToBool  = toBool++-- |Convert a C enumeration to Haskell.+--+cToEnum :: (Integral i, Enum e) => i -> e+cToEnum  = toEnum . cIntConv++-- |Convert a Haskell enumeration to C.+--+cFromEnum :: (Enum e, Integral i) => e -> i+cFromEnum  = cIntConv . fromEnum
+ Foreign/CUDA/Internal/Offsets.hsc view
@@ -0,0 +1,33 @@+{-# LANGUAGE ForeignFunctionInterface #-}+{-+ - Structure field offset constants.+ - Too difficult to extract using C->Haskell )=+ -}++module Foreign.CUDA.Internal.Offsets where+++--------------------------------------------------------------------------------+-- Runtime API+--------------------------------------------------------------------------------++#include <cuda_runtime_api.h>++devNameOffset, devMaxThreadDimOffset, devMaxGridSizeOffset :: Int++devNameOffset         = #{offset struct cudaDeviceProp, name}+devMaxThreadDimOffset = #{offset struct cudaDeviceProp, maxThreadsDim}+devMaxGridSizeOffset  = #{offset struct cudaDeviceProp, maxGridSize}+++--------------------------------------------------------------------------------+-- Driver API+--------------------------------------------------------------------------------++#include <cuda.h>++devMaxThreadDimOffset', devMaxGridSizeOffset' :: Int++devMaxThreadDimOffset' = #{offset struct CUdevprop_st, maxThreadsDim}+devMaxGridSizeOffset'  = #{offset struct CUdevprop_st, maxGridSize}+
+ Foreign/CUDA/Runtime.hs view
@@ -0,0 +1,30 @@+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Runtime+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Top level bindings to the C-for-CUDA runtime API+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Runtime+  (+    module Foreign.CUDA.Runtime.Device,+    module Foreign.CUDA.Runtime.Error,+    module Foreign.CUDA.Runtime.Exec,+    module Foreign.CUDA.Runtime.Marshal,+    module Foreign.CUDA.Runtime.Ptr,+    module Foreign.CUDA.Runtime.Thread,+    module Foreign.CUDA.Runtime.Utils+  )+  where++import Foreign.CUDA.Runtime.Device+import Foreign.CUDA.Runtime.Error+import Foreign.CUDA.Runtime.Exec+import Foreign.CUDA.Runtime.Marshal+import Foreign.CUDA.Runtime.Ptr+import Foreign.CUDA.Runtime.Thread+import Foreign.CUDA.Runtime.Utils+
+ Foreign/CUDA/Runtime/Device.chs view
@@ -0,0 +1,229 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Runtime.Device+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Device management routines+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Runtime.Device+  (+    ComputeMode(..), DeviceFlag(..), DeviceProperties(..),++    -- ** Device management+    choose, get, count, props, set, setFlags, setOrder+  )+  where++#include <cuda_runtime_api.h>+{# context lib="cudart" #}++-- Friends+import Foreign.CUDA.Runtime.Error+import Foreign.CUDA.Internal.C2HS+import Foreign.CUDA.Internal.Offsets++-- System+import Foreign+import Foreign.C++#c+typedef struct cudaDeviceProp   cudaDeviceProp;++typedef enum+{+    cudaDeviceFlagScheduleAuto  = cudaDeviceScheduleAuto,+    cudaDeviceFlagScheduleSpin  = cudaDeviceScheduleSpin,+    cudaDeviceFlagScheduleYield = cudaDeviceScheduleYield,+    cudaDeviceFlagBlockingSync  = cudaDeviceBlockingSync,+    cudaDeviceFlagMapHost       = cudaDeviceMapHost+} cudaDeviceFlags;+#endc+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++{# pointer *cudaDeviceProp as ^ foreign -> DeviceProperties nocode #}++-- |+-- The compute mode the device is currently in+--+{# enum cudaComputeMode as ComputeMode { }+    with prefix="cudaComputeMode" deriving (Eq, Show) #}++-- |+-- Device execution flags+--+{# enum cudaDeviceFlags as DeviceFlag { }+    with prefix="cudaDeviceFlag" deriving (Eq, Show) #}++-- |+-- The properties of a compute device+--+data DeviceProperties = DeviceProperties+  {+    deviceName               :: String,         -- ^ Identifier+    computeCapability        :: Double,         -- ^ 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+    maxThreadsPerBlock       :: Int,            -- ^ Max number of threads per block+    maxThreadsDim            :: (Int,Int,Int),  -- ^ Max size of each dimension of a block+    maxGridSize              :: (Int,Int,Int),  -- ^ Max size of each dimension of a grid+    clockRate                :: Int,            -- ^ Clock frequency in kilohertz+    multiProcessorCount      :: Int,            -- ^ Number of multiprocessors on the device+    memPitch                 :: Int64,          -- ^ Max pitch in bytes allowed by memory copies+    textureAlignment         :: Int64,          -- ^ Alignment requirement for textures+    computeMode              :: ComputeMode,+    deviceOverlap            :: Bool,           -- ^ Device can concurrently copy memory and execute a kernel+    kernelExecTimeoutEnabled :: Bool,           -- ^ Whether there is a runtime limit on kernels+    integrated               :: Bool,           -- ^ As opposed to discrete+    canMapHostMemory         :: Bool            -- ^ Device can use pinned memory+  }+  deriving (Show)+++instance Storable DeviceProperties where+  sizeOf _    = {#sizeof cudaDeviceProp#}+  alignment _ = alignment (undefined :: Ptr ())++  peek p      = do+    gm <- cIntConv     `fmap` {#get cudaDeviceProp.totalGlobalMem#} p+    sm <- cIntConv     `fmap` {#get cudaDeviceProp.sharedMemPerBlock#} p+    rb <- cIntConv     `fmap` {#get cudaDeviceProp.regsPerBlock#} p+    ws <- cIntConv     `fmap` {#get cudaDeviceProp.warpSize#} p+    mp <- cIntConv     `fmap` {#get cudaDeviceProp.memPitch#} p+    tb <- cIntConv     `fmap` {#get cudaDeviceProp.maxThreadsPerBlock#} p+    cl <- cIntConv     `fmap` {#get cudaDeviceProp.clockRate#} p+    cm <- cIntConv     `fmap` {#get cudaDeviceProp.totalConstMem#} p+    v1 <- fromIntegral `fmap` {#get cudaDeviceProp.major#} p+    v2 <- fromIntegral `fmap` {#get cudaDeviceProp.minor#} p+    ta <- cIntConv     `fmap` {#get cudaDeviceProp.textureAlignment#} p+    ov <- cToBool      `fmap` {#get cudaDeviceProp.deviceOverlap#} p+    pc <- cIntConv     `fmap` {#get cudaDeviceProp.multiProcessorCount#} p+    ke <- cToBool      `fmap` {#get cudaDeviceProp.kernelExecTimeoutEnabled#} p+    tg <- cToBool      `fmap` {#get cudaDeviceProp.integrated#} p+    hm <- cToBool      `fmap` {#get cudaDeviceProp.canMapHostMemory#} p+    md <- cToEnum      `fmap` {#get cudaDeviceProp.computeMode#} p++    --+    -- C->Haskell returns the wrong type when accessing static arrays in+    -- structs, returning the dereferenced element but with a Ptr type. Work+    -- around this with manual pointer arithmetic...+    --+    n            <- peekCString (p `plusPtr` devNameOffset)+    (t1:t2:t3:_) <- map cIntConv `fmap` peekArray 3 (p `plusPtr` devMaxThreadDimOffset :: Ptr CInt)+    (g1:g2:g3:_) <- map cIntConv `fmap` peekArray 3 (p `plusPtr` devMaxGridSizeOffset  :: Ptr CInt)++    return DeviceProperties+      {+        deviceName               = n,+        computeCapability        = v1 + v2 / max 10 (10^ ((ceiling . logBase 10) v2 :: Int)),+        totalGlobalMem           = gm,+        totalConstMem            = cm,+        sharedMemPerBlock        = sm,+        regsPerBlock             = rb,+        warpSize                 = ws,+        maxThreadsPerBlock       = tb,+        maxThreadsDim            = (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+      }+++--------------------------------------------------------------------------------+-- Functions+--------------------------------------------------------------------------------++-- |+-- Select the compute device which best matches the given criteria+--+choose     :: DeviceProperties -> IO Int+choose dev =  resultIfOk =<< cudaChooseDevice dev++{# fun unsafe cudaChooseDevice+  { alloca-      `Int'              peekIntConv*+  , withDevProp* `DeviceProperties'              } -> `Status' cToEnum #}+  where+      withDevProp = with+++-- |+-- Returns which device is currently being used+--+get :: IO Int+get =  resultIfOk =<< cudaGetDevice++{# fun unsafe cudaGetDevice+  { alloca- `Int' peekIntConv* } -> `Status' cToEnum #}+++-- |+-- Returns the number of devices available for execution, with compute+-- capability >= 1.0+--+count :: IO Int+count =  resultIfOk =<< cudaGetDeviceCount++{# fun unsafe cudaGetDeviceCount+  { alloca- `Int' peekIntConv* } -> `Status' cToEnum #}+++-- |+-- Return information about the selected compute device+--+props   :: Int -> IO DeviceProperties+props n =  resultIfOk =<< cudaGetDeviceProperties n++{# fun unsafe cudaGetDeviceProperties+  { alloca- `DeviceProperties' peek*+  ,         `Int'                    } -> `Status' cToEnum #}+++-- |+-- Set device to be used for GPU execution+--+set   :: Int -> IO ()+set n =  nothingIfOk =<< cudaSetDevice n++{# fun unsafe cudaSetDevice+  { `Int' } -> `Status' cToEnum #}+++-- |+-- Set flags to be used for device executions+--+setFlags   :: [DeviceFlag] -> IO ()+setFlags f =  nothingIfOk =<< cudaSetDeviceFlags (combineBitMasks f)++{# fun unsafe cudaSetDeviceFlags+  { `Int' } -> `Status' cToEnum #}+++-- |+-- Set list of devices for CUDA execution in priority order+--+setOrder   :: [Int] -> IO ()+setOrder l =  nothingIfOk =<< cudaSetValidDevices l (length l)++{# fun unsafe cudaSetValidDevices+  { withArrayIntConv* `[Int]'+  ,                   `Int'   } -> `Status' cToEnum #}+  where+      withArrayIntConv = withArray . map cIntConv+
+ Foreign/CUDA/Runtime/Error.chs view
@@ -0,0 +1,103 @@+{-# LANGUAGE ForeignFunctionInterface, DeriveDataTypeable #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Runtime.Error+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Error handling functions+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Runtime.Error+  (+    Status(..), CUDAException(..),+    cudaError, describe,+    resultIfOk, nothingIfOk+  )+  where+++-- Friends+import Foreign.CUDA.Internal.C2HS++-- System+import Foreign+import Foreign.C+import Data.Typeable+import Control.Exception.Extensible++#include <cuda_runtime_api.h>+{# context lib="cudart" #}+++--------------------------------------------------------------------------------+-- Return Status+--------------------------------------------------------------------------------++-- |+-- Return codes from API functions+--+{# enum cudaError as Status+    { cudaSuccess as Success }+    with prefix="cudaError" deriving (Eq, Show) #}++--------------------------------------------------------------------------------+-- Exceptions+--------------------------------------------------------------------------------++data CUDAException+  = ExitCode Status+  | UserError String+  deriving Typeable++instance Exception CUDAException++instance Show CUDAException where+  showsPrec _ (ExitCode  s) = showString ("CUDA Exception: " ++ describe s)+  showsPrec _ (UserError s) = showString ("CUDA Exception: " ++ s)+++-- |+-- Raise a 'CUDAException' in the IO Monad+--+cudaError :: String -> IO a+cudaError s = throwIO (UserError s)+++--------------------------------------------------------------------------------+-- Helper Functions+--------------------------------------------------------------------------------++-- |+-- Return the descriptive string associated with a particular error code+--+{# fun pure unsafe cudaGetErrorString as describe+    { cFromEnum `Status' } -> `String' #}+--+-- Logically, this must be a pure function, returning a pointer to a statically+-- defined string constant.+--+++-- |+-- Return the results of a function on successful execution, otherwise return+-- the error string associated with the return code+--+resultIfOk :: (Status, a) -> IO a+resultIfOk (status,result) =+    case status of+        Success -> return  result+        _       -> throwIO (ExitCode status)+++-- |+-- Return the error string associated with an unsuccessful return code,+-- otherwise Nothing+--+nothingIfOk :: Status -> IO ()+nothingIfOk status =+    case status of+        Success -> return  ()+        _       -> throwIO (ExitCode status)+
+ Foreign/CUDA/Runtime/Event.chs view
@@ -0,0 +1,133 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Driver.Event+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Event management for C-for-CUDA runtime environment+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Runtime.Event+  (+    Event, EventFlag(..),+    create, destroy, elapsedTime, query, record, block+  )+  where++#include <cuda_runtime_api.h>+{# context lib="cudart" #}++-- Friends+import Foreign.CUDA.Runtime.Error+import Foreign.CUDA.Runtime.Stream                      (Stream(..))+import Foreign.CUDA.Internal.C2HS++-- System+import Foreign+import Foreign.C+import Control.Monad                                    (liftM)+++#c+typedef enum cudaEvent_option_enum {+//  CUDA_EVENT_OPTION_DEFAULT       = cuduEventDefault,+    CUDA_EVENT_OPTION_BLOCKING_SYNC = cudaEventBlockingSync+} cudaEvent_option;+#endc++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- Events+--+newtype Event = Event { useEvent :: {# type cudaEvent_t #}}++-- |+-- Event creation flags+--+{# enum cudaEvent_option_enum as EventFlag+    { underscoreToCase }+    with prefix="CUDA_EVENT_OPTION" deriving (Eq,Show) #}+++--------------------------------------------------------------------------------+-- Event management+--------------------------------------------------------------------------------++-- |+-- Create a new event+--+create :: [EventFlag] -> IO Event+create flags = resultIfOk =<< cudaEventCreateWithFlags flags++{# fun unsafe cudaEventCreateWithFlags+  { alloca-         `Event'       peekEvt*+  , combineBitMasks `[EventFlag]'          } -> `Status' cToEnum #}+  where peekEvt = liftM Event . peek+++-- |+-- Destroy an event+--+destroy :: Event -> IO ()+destroy ev = nothingIfOk =<< cudaEventDestroy ev++{# fun unsafe cudaEventDestroy+  { useEvent `Event' } -> `Status' cToEnum #}+++-- |+-- Determine the elapsed time (in milliseconds) between two events+--+elapsedTime :: Event -> Event -> IO Float+elapsedTime ev1 ev2 = resultIfOk =<< cudaEventElapsedTime ev1 ev2++{# fun unsafe cudaEventElapsedTime+  { alloca-  `Float' peekFloatConv*+  , useEvent `Event'+  , useEvent `Event'                } -> `Status' cToEnum #}+++-- |+-- Determines if a event has actually been recorded+--+query :: Event -> IO Bool+query ev =+  cudaEventQuery ev >>= \rv ->+  case rv of+    Success  -> return True+    NotReady -> return False+    _        -> resultIfOk (rv,undefined)++{# fun unsafe cudaEventQuery+  { useEvent `Event' } -> `Status' cToEnum #}+++-- |+-- Record an event once all operations in the current context (or optionally+-- specified stream) have completed. This operation is asynchronous.+--+record :: Event -> Maybe Stream -> IO ()+record ev mst =+  nothingIfOk =<< case mst of+    Just st -> cudaEventRecord ev st+    Nothing -> cudaEventRecord ev (Stream 0)++{# fun unsafe cudaEventRecord+  { useEvent  `Event'+  , useStream `Stream' } -> `Status' cToEnum #}+++-- |+-- Wait until the event has been recorded+--+block :: Event -> IO ()+block ev = nothingIfOk =<< cudaEventSynchronize ev++{# fun unsafe cudaEventSynchronize+  { useEvent `Event' } -> `Status' cToEnum #}+
+ Foreign/CUDA/Runtime/Exec.chs view
@@ -0,0 +1,180 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Runtime.Exec+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Kernel execution control for C-for-CUDA runtime interface+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Runtime.Exec+  (+    FunAttributes(..), FunParam(..),+    attributes, setConfig, setParams, launch+  )+  where++#include "cbits/stubs.h"+#include <cuda_runtime_api.h>+{# context lib="cudart" #}++-- Friends+import Foreign.CUDA.Runtime.Stream+import Foreign.CUDA.Runtime.Error+import Foreign.CUDA.Internal.C2HS++-- System+import Foreign+import Foreign.C+import Control.Monad++#c+typedef struct cudaFuncAttributes cudaFuncAttributes;+#endc+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++--+-- Function Attributes+--+{# pointer *cudaFuncAttributes as ^ foreign -> FunAttributes nocode #}++data FunAttributes = FunAttributes+  {+    constSizeBytes           :: Int64,+    localSizeBytes           :: Int64,+    sharedSizeBytes          :: Int64,+    maxKernelThreadsPerBlock :: Int,	-- ^ maximum block size that can be successively launched (based on register usage)+    numRegs                  :: Int     -- ^ number of registers required for each thread+  }+  deriving (Show)++instance Storable FunAttributes where+  sizeOf _    = {# sizeof cudaFuncAttributes #}+  alignment _ = alignment (undefined :: Ptr ())++  peek p      = do+    cs <- cIntConv `fmap` {#get cudaFuncAttributes.constSizeBytes#} p+    ls <- cIntConv `fmap` {#get cudaFuncAttributes.localSizeBytes#} p+    ss <- cIntConv `fmap` {#get cudaFuncAttributes.sharedSizeBytes#} p+    tb <- cIntConv `fmap` {#get cudaFuncAttributes.maxThreadsPerBlock#} p+    nr <- cIntConv `fmap` {#get cudaFuncAttributes.numRegs#} p++    return FunAttributes+      {+        constSizeBytes           = cs,+        localSizeBytes           = ls,+        sharedSizeBytes          = ss,+        maxKernelThreadsPerBlock = tb,+        numRegs                  = nr+      }++-- |+-- Kernel function parameters. Doubles will be converted to an internal float+-- representation on devices that do not support doubles natively.+--+data Storable a => FunParam a+    = IArg Int+    | FArg Float+    | DArg Double+    | VArg a++--------------------------------------------------------------------------------+-- Execution Control+--------------------------------------------------------------------------------++-- |+-- Obtain the attributes of the named @__global__@ device function. This+-- itemises the requirements to successfully launch the given kernel.+--+attributes :: String -> IO FunAttributes+attributes fn = resultIfOk =<< cudaFuncGetAttributes fn++{# fun unsafe cudaFuncGetAttributes+  { alloca-      `FunAttributes' peek*+  , withCString* `String'              } -> `Status' cToEnum #}+++-- |+-- Specify the grid and block dimensions for a device call. Used in conjunction+-- with 'setParams', this pushes data onto the execution stack that will be+-- popped when a function is 'launch'ed.+--+setConfig :: (Int,Int)		-- ^ grid dimensions+	  -> (Int,Int,Int)	-- ^ block dimensions+	  -> Int64		-- ^ shared memory per block (bytes)+	  -> Maybe Stream	-- ^ associated processing stream+	  -> IO ()+setConfig (gx,gy) (bx,by,bz) sharedMem mst =+  nothingIfOk =<< case mst of+    Nothing -> cudaConfigureCallSimple gx gy bx by bz sharedMem (Stream 0)+    Just st -> cudaConfigureCallSimple gx gy bx by bz sharedMem st++--+-- The FFI does not support passing deferenced structures to C functions, as+-- this is highly platform/compiler dependent. Wrap our own function stub+-- accepting plain integers.+--+{# fun unsafe cudaConfigureCallSimple+  {	      `Int', `Int'+  ,           `Int', `Int', `Int'+  , cIntConv  `Int64'+  , useStream `Stream'            } -> `Status' cToEnum #}+++-- |+-- Set the argument parameters that will be passed to the next kernel+-- invocation. This is used in conjunction with 'setConfig' to control kernel+-- execution.+setParams :: Storable a => [FunParam a] -> IO ()+setParams = foldM_ k 0+  where+    k offset arg = do+      let s = size arg+      set arg s offset >>= nothingIfOk+      return (offset + s)++    size (IArg _) = sizeOf (undefined :: Int)+    size (FArg _) = sizeOf (undefined :: Float)+    size (DArg _) = sizeOf (undefined :: Double)+    size (VArg a) = sizeOf a++    set (IArg v) s o = cudaSetupArgument v s o+    set (FArg v) s o = cudaSetupArgument v s o+    set (VArg v) s o = cudaSetupArgument v s o+    set (DArg v) s o =+      cudaSetDoubleForDevice v >>= resultIfOk >>= \d ->+      cudaSetupArgument d s o+++{# fun unsafe cudaSetupArgument+  `Storable a' =>+  { with'* `a'+  ,        `Int'+  ,        `Int'   } -> `Status' cToEnum #}+  where+    with' v a = with v $ \p -> a (castPtr p)++{# fun unsafe cudaSetDoubleForDevice+  { with'* `Double' peek'* } -> `Status' cToEnum #}+  where+    with' v a = with v $ \p -> a (castPtr p)+    peek'     = peek . castPtr+++-- |+-- Invoke the named kernel on the device, which must have been declared+-- @__global__@. This must be preceded by a call to 'setConfig' and (if+-- appropriate) 'setParams'.+--+launch :: String -> IO ()+launch fn = nothingIfOk =<< cudaLaunch fn++{# fun unsafe cudaLaunch+  { withCString* `String' } -> `Status' cToEnum #}+
+ Foreign/CUDA/Runtime/Marshal.chs view
@@ -0,0 +1,359 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Runtime.Marshal+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Memory management for CUDA devices+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Runtime.Marshal+  (+    -- * Host Allocation+    AllocFlag(..),+    mallocHostArray, freeHost,++    -- * Device Allocation+    mallocArray, allocaArray, free,++    -- * Marshalling+    peekArray, peekArrayAsync, peekListArray,+    pokeArray, pokeArrayAsync, pokeListArray,+    copyArray, copyArrayAsync,++    -- * Combined Allocation and Marshalling+    newListArray, withListArray, withListArrayLen,++    -- * Utility+    memset+  )+  where++#include <cuda_runtime_api.h>+{# context lib="cudart" #}++-- Friends+import Foreign.CUDA.Runtime.Ptr+import Foreign.CUDA.Runtime.Error+import Foreign.CUDA.Runtime.Stream+import Foreign.CUDA.Internal.C2HS++-- System+import Data.Int+import Control.Exception.Extensible++import Foreign.C+import Foreign.Ptr+import Foreign.Storable+import qualified Foreign.Marshal as F++#c+typedef enum cudaMemHostAlloc_option_enum {+//  CUDA_MEMHOSTALLOC_OPTION_DEFAULT        = cudaHostAllocDefault,+    CUDA_MEMHOSTALLOC_OPTION_DEVICE_MAPPED  = cudaHostAllocMapped,+    CUDA_MEMHOSTALLOC_OPTION_PORTABLE       = cudaHostAllocPortable,+    CUDA_MEMHOSTALLOC_OPTION_WRITE_COMBINED = cudaHostAllocWriteCombined+} cudaMemHostAlloc_option;+#endc+++--------------------------------------------------------------------------------+-- Host Allocation+--------------------------------------------------------------------------------++-- |+-- Options for host allocation+--+{# enum cudaMemHostAlloc_option as AllocFlag+    { underscoreToCase }+    with prefix="CUDA_MEMHOSTALLOC_OPTION" deriving (Eq, Show) #}+++-- |+-- Allocate a section of linear memory on the host which is page-locked and+-- directly accessible from the device. The storage is sufficient to hold the+-- given number of elements of a storable type. The runtime system automatically+-- accelerates calls to functions such as 'memcpy' to page-locked memory.+--+-- Note that since the amount of pageable memory is thusly reduced, overall+-- system performance may suffer. This is best used sparingly to allocate+-- staging areas for data exchange+--+mallocHostArray :: Storable a => [AllocFlag] -> Int -> IO (HostPtr a)+mallocHostArray flags = doMalloc undefined+  where+    doMalloc :: Storable a' => a' -> Int -> IO (HostPtr a')+    doMalloc x n = resultIfOk =<< cudaHostAlloc (fromIntegral n * fromIntegral (sizeOf x)) flags++{# fun unsafe cudaHostAlloc+  { alloca'-        `HostPtr a' hptr*+  , cIntConv        `Int64'+  , combineBitMasks `[AllocFlag]'     } -> `Status' cToEnum #}+  where+    alloca' = F.alloca+    hptr p  = (HostPtr . castPtr) `fmap` peek p+++-- |+-- Free page-locked host memory previously allocated with 'mallecHost'+--+freeHost :: HostPtr a -> IO ()+freeHost p = nothingIfOk =<< cudaFreeHost p++{# fun unsafe cudaFreeHost+  { hptr `HostPtr a' } -> `Status' cToEnum #}+  where hptr = castPtr . useHostPtr+++--------------------------------------------------------------------------------+-- Device Allocation+--------------------------------------------------------------------------------++-- |+-- Allocate a section of linear memory on the device, and return a reference to+-- it. The memory is sufficient to hold the given number of elements of storable+-- type. It is suitable aligned, and not cleared.+--+mallocArray :: Storable a => Int -> IO (DevicePtr a)+mallocArray = doMalloc undefined+  where+    doMalloc :: Storable a' => a' -> Int -> IO (DevicePtr a')+    doMalloc x n = resultIfOk =<< cudaMalloc (fromIntegral n * fromIntegral (sizeOf x))++{# fun unsafe cudaMalloc+  { alloca'- `DevicePtr a' dptr*+  , cIntConv `Int64'             } -> `Status' cToEnum #}+  where+    -- C-> Haskell doesn't like qualified imports in marshaller specifications+    alloca' = F.alloca+    dptr p  = (castDevPtr . DevicePtr) `fmap` peek p+++-- |+-- Execute a computation, passing a pointer to a temporarily allocated block of+-- memory sufficient to hold the given number of elements of storable type. The+-- memory is freed when the computation terminates (normally or via an+-- exception), so the pointer must not be used after this.+--+-- Note that kernel launches can be asynchronous, so you may need to add a+-- synchronisation point at the end of the computation.+--+allocaArray :: Storable a => Int -> (DevicePtr a -> IO b) -> IO b+allocaArray n = bracket (mallocArray n) free+++-- |+-- Free previously allocated memory on the device+--+free :: DevicePtr a -> IO ()+free p = nothingIfOk =<< cudaFree p++{# fun unsafe cudaFree+  { dptr `DevicePtr a' } -> `Status' cToEnum #}+  where+    dptr = useDevicePtr . castDevPtr+++--------------------------------------------------------------------------------+-- Marshalling+--------------------------------------------------------------------------------++-- |+-- Copy a number of elements from the device to host memory. This is a+-- synchronous operation.+--+peekArray :: Storable a => Int -> DevicePtr a -> Ptr a -> IO ()+peekArray n dptr hptr = memcpy hptr (useDevicePtr dptr) n DeviceToHost+++-- |+-- Copy memory from the device asynchronously, possibly associated with a+-- particular stream. The destination memory must be page locked.+--+peekArrayAsync :: Storable a => Int -> DevicePtr a -> HostPtr a -> Maybe Stream -> IO ()+peekArrayAsync n dptr hptr mst =+  memcpyAsync (useHostPtr hptr) (useDevicePtr dptr) n DeviceToHost mst+++-- |+-- Copy a number of elements from the device into a new Haskell list. Note that+-- this requires two memory copies: firstly from the device into a heap+-- allocated array, and from there marshalled into a list+--+peekListArray :: Storable a => Int -> DevicePtr a -> IO [a]+peekListArray n dptr =+  F.allocaArray n $ \p -> do+    peekArray   n dptr p+    F.peekArray n p+++-- |+-- Copy a number of elements onto the device. This is a synchronous operation.+--+pokeArray :: Storable a => Int -> Ptr a -> DevicePtr a -> IO ()+pokeArray n hptr dptr = memcpy (useDevicePtr dptr) hptr n HostToDevice+++-- |+-- Copy memory onto the device asynchronously, possibly associated with a+-- particular stream. The source memory must be page-locked.+--+pokeArrayAsync :: Storable a => Int -> HostPtr a -> DevicePtr a -> Maybe Stream -> IO ()+pokeArrayAsync n hptr dptr mst =+  memcpyAsync (useDevicePtr dptr) (useHostPtr hptr) n HostToDevice mst+++-- |+-- Write a list of storable elements into a device array. The array must be+-- sufficiently large to hold the entire list. This requires two marshalling+-- operations+--+pokeListArray :: Storable a => [a] -> DevicePtr a -> IO ()+pokeListArray xs dptr = F.withArrayLen xs $ \len p -> pokeArray len p dptr+++-- |+-- Copy the given number of elements from the first device array (source) to the+-- second (destination). The copied areas may not overlap. This is a synchronous+-- operation.+--+copyArray :: Storable a => Int -> DevicePtr a -> DevicePtr a -> IO ()+copyArray n src dst = memcpy (useDevicePtr dst) (useDevicePtr src) n DeviceToDevice+++-- |+-- Copy the given number of elements from the first device array (source) to the+-- second (destination). The copied areas may not overlap. This operation is+-- asynchronous with respect to host, but will never overlap with kernel+-- execution.+--+copyArrayAsync :: Storable a => Int -> DevicePtr a -> DevicePtr a -> Maybe Stream -> IO ()+copyArrayAsync n src dst mst =+  memcpyAsync (useDevicePtr dst) (useDevicePtr src) n DeviceToDevice mst+++--+-- Memory copy kind+--+{# enum cudaMemcpyKind as CopyDirection {}+    with prefix="cudaMemcpy" deriving (Eq, Show) #}++-- |+-- Copy data between host and device. This is a synchronous operation.+--+memcpy :: Storable a+       => Ptr a                 -- ^ destination+       -> Ptr a                 -- ^ source+       -> Int                   -- ^ number of elements+       -> CopyDirection+       -> IO ()+memcpy dst src n dir = doMemcpy undefined dst+  where+    doMemcpy :: Storable a' => a' -> Ptr a' -> IO ()+    doMemcpy x _ =+      nothingIfOk =<< cudaMemcpy dst src (fromIntegral n * fromIntegral (sizeOf x)) dir++{# fun unsafe cudaMemcpy+  { castPtr   `Ptr a'+  , castPtr   `Ptr a'+  , cIntConv  `Int64'+  , cFromEnum `CopyDirection' } -> `Status' cToEnum #}+++-- |+-- Copy data between the host and device asynchronously, possibly associated+-- with a particular stream. The host-side memory must be page-locked (allocated+-- with 'mallocHostArray').+--+memcpyAsync :: Storable a+            => Ptr a            -- ^ destination+            -> Ptr a            -- ^ source+            -> Int              -- ^ number of elements+            -> CopyDirection+            -> Maybe Stream+            -> IO ()+memcpyAsync dst src n kind mst = doMemcpy undefined dst+  where+    doMemcpy :: Storable a' => a' -> Ptr a' -> IO ()+    doMemcpy x _ =+      let bytes = fromIntegral n * fromIntegral (sizeOf x) in+      nothingIfOk =<< case mst of+        Nothing -> cudaMemcpyAsync dst src bytes kind (Stream 0)+        Just st -> cudaMemcpyAsync dst src bytes kind st++{# fun unsafe cudaMemcpyAsync+  { castPtr   `Ptr a'+  , castPtr   `Ptr a'+  , cIntConv  `Int64'+  , cFromEnum `CopyDirection'+  , useStream `Stream'        } -> `Status' cToEnum #}+++--------------------------------------------------------------------------------+-- Combined Allocation and Marshalling+--------------------------------------------------------------------------------++-- |+-- Write a list of storable elements into a newly allocated device array. Note+-- that this requires two copy operations: firstly from a Haskell list into a+-- heap-allocated array, and from there into device memory. The array should be+-- 'free'd when no longer required.+--+newListArray :: Storable a => [a] -> IO (DevicePtr a)+newListArray xs =+  F.withArrayLen xs                     $ \len p ->+  bracketOnError (mallocArray len) free $ \d_xs  -> do+    pokeArray len p d_xs+    return d_xs+++-- |+-- Temporarily store a list of elements into a newly allocated device array. An+-- IO action is applied to the array, the result of which is returned. Similar+-- to 'newListArray', this requires two marshalling operations of the data.+--+-- As with 'allocaArray', the memory is freed once the action completes, so you+-- should not return the pointer from the action, and be sure that any+-- asynchronous operations (such as kernel execution) have completed.+--+withListArray :: Storable a => [a] -> (DevicePtr a -> IO b) -> IO b+withListArray xs = withListArrayLen xs . const+++-- |+-- A variant of 'withListArray' which also supplies the number of elements in+-- the array to the applied function+--+withListArrayLen :: Storable a => [a] -> (Int -> DevicePtr a -> IO b) -> IO b+withListArrayLen xs f =+  allocaArray   len $ \d_xs -> do+  F.allocaArray len $ \h_xs -> do+    F.pokeArray h_xs xs+    pokeArray len h_xs d_xs+  f len d_xs+  where+    len = length xs+++--------------------------------------------------------------------------------+-- Utility+--------------------------------------------------------------------------------++-- |+-- Initialise device memory to a given 8-bit value+--+memset :: DevicePtr a                   -- ^ The device memory+       -> Int64                         -- ^ Number of bytes+       -> Int8                          -- ^ Value to set for each byte+       -> IO ()+memset dptr bytes symbol = nothingIfOk =<< cudaMemset dptr symbol bytes++{# fun unsafe cudaMemset+  { dptr     `DevicePtr a'+  , cIntConv `Int8'+  , cIntConv `Int64'       } -> `Status' cToEnum #}+  where+    dptr = useDevicePtr . castDevPtr+
+ Foreign/CUDA/Runtime/Ptr.hs view
@@ -0,0 +1,161 @@+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Runtime.Ptr+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- References to objects stored on the CUDA devices+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Runtime.Ptr+  where++-- System+import Foreign.Ptr+import Foreign.Storable+++--------------------------------------------------------------------------------+-- Device Pointer+--------------------------------------------------------------------------------++-- |+-- A reference to data stored on the device+--+data DevicePtr a = DevicePtr { useDevicePtr :: Ptr a }+  deriving (Eq,Ord)++instance Show (DevicePtr a) where+  showsPrec n (DevicePtr p) = showsPrec n p++instance Storable (DevicePtr a) where+  sizeOf _    = sizeOf    (undefined :: Ptr a)+  alignment _ = alignment (undefined :: Ptr a)+  peek p      = DevicePtr `fmap` peek (castPtr p)+  poke p v    = poke (castPtr p) (useDevicePtr v)+++-- |+-- Look at the contents of device memory. This takes an IO action that will be+-- applied to that pointer, the result of which is returned. It would be silly+-- to return the pointer from the action.+--+withDevicePtr :: DevicePtr a -> (Ptr a -> IO b) -> IO b+withDevicePtr p f = f (useDevicePtr p)+++-- |+-- The constant 'nullDevPtr' contains the distinguished memory location that is+-- not associated with a valid memory location+--+nullDevPtr :: DevicePtr a+nullDevPtr =  DevicePtr nullPtr++-- |+-- Cast a device pointer from one type to another+--+castDevPtr :: DevicePtr a -> DevicePtr b+castDevPtr (DevicePtr p) = DevicePtr (castPtr p)++-- |+-- Advance the pointer address by the given offset in bytes.+--+plusDevPtr :: DevicePtr a -> Int -> DevicePtr a+plusDevPtr (DevicePtr p) d = DevicePtr (p `plusPtr` d)++-- |+-- Given an alignment constraint, align the device pointer to the next highest+-- address satisfying the constraint+--+alignDevPtr :: DevicePtr a -> Int -> DevicePtr a+alignDevPtr (DevicePtr p) i = DevicePtr (p `alignPtr` i)++-- |+-- Compute the difference between the second and first argument. This fulfils+-- the relation+--+-- > p2 == p1 `plusDevPtr` (p2 `minusDevPtr` p1)+--+minusDevPtr :: DevicePtr a -> DevicePtr a -> Int+minusDevPtr (DevicePtr a) (DevicePtr b) = a `minusPtr` b++-- |+-- Advance a pointer into a device array by the given number of elements+--+advanceDevPtr :: Storable a => DevicePtr a -> Int -> DevicePtr a+advanceDevPtr  = doAdvance undefined+  where+    doAdvance :: Storable a' => a' -> DevicePtr a' -> Int -> DevicePtr a'+    doAdvance x p i = p `plusDevPtr` (i * sizeOf x)+++--------------------------------------------------------------------------------+-- Host Pointer+--------------------------------------------------------------------------------++-- |+-- A reference to page-locked host memory+--+data HostPtr a = HostPtr { useHostPtr :: Ptr a }+  deriving (Eq,Ord)++instance Show (HostPtr a) where+  showsPrec n (HostPtr p) = showsPrec n p++instance Storable (HostPtr a) where+  sizeOf _    = sizeOf    (undefined :: Ptr a)+  alignment _ = alignment (undefined :: Ptr a)+  peek p      = HostPtr `fmap` peek (castPtr p)+  poke p v    = poke (castPtr p) (useHostPtr v)+++-- |+-- Apply an IO action to the memory reference living inside the host pointer+-- object. All uses of the pointer should be inside the 'withHostPtr' bracket.+--+withHostPtr :: HostPtr a -> (Ptr a -> IO b) -> IO b+withHostPtr p f = f (useHostPtr p)+++-- |+-- The constant 'nullHostPtr' contains the distinguished memory location that is+-- not associated with a valid memory location+--+nullHostPtr :: HostPtr a+nullHostPtr =  HostPtr nullPtr++-- |+-- Cast a host pointer from one type to another+--+castHostPtr :: HostPtr a -> HostPtr b+castHostPtr (HostPtr p) = HostPtr (castPtr p)++-- |+-- Advance the pointer address by the given offset in bytes+--+plusHostPtr :: HostPtr a -> Int -> HostPtr a+plusHostPtr (HostPtr p) d = HostPtr (p `plusPtr` d)++-- |+-- Given an alignment constraint, align the host pointer to the next highest+-- address satisfying the constraint+--+alignHostPtr :: HostPtr a -> Int -> HostPtr a+alignHostPtr (HostPtr p) i = HostPtr (p `alignPtr` i)++-- |+-- Compute the difference between the second and first argument+--+minusHostPtr :: HostPtr a -> HostPtr a -> Int+minusHostPtr (HostPtr a) (HostPtr b) = a `minusPtr` b++-- |+-- Advance a pointer into a host array by a given number of elements+--+advanceHostPtr :: Storable a => HostPtr a -> Int -> HostPtr a+advanceHostPtr  = doAdvance undefined+  where+    doAdvance :: Storable a' => a' -> HostPtr a' -> Int -> HostPtr a'+    doAdvance x p i = p `plusHostPtr` (i * sizeOf x)+
+ Foreign/CUDA/Runtime/Stream.chs view
@@ -0,0 +1,94 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Runtime.Stream+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Stream management routines+--+--------------------------------------------------------------------------------+++module Foreign.CUDA.Runtime.Stream+  (+    Stream(..),++    -- ** Stream management+    create, destroy, finished, block+  )+  where++#include <cuda_runtime_api.h>+{# context lib="cudart" #}++-- Friends+import Foreign.CUDA.Runtime.Error+import Foreign.CUDA.Internal.C2HS++-- System+import Foreign+import Foreign.C+import Control.Monad                                    (liftM)+++--------------------------------------------------------------------------------+-- Data Types+--------------------------------------------------------------------------------++-- |+-- A processing stream+--+newtype Stream = Stream { useStream :: {# type cudaStream_t #}}+  deriving (Show)+++--------------------------------------------------------------------------------+-- Functions+--------------------------------------------------------------------------------++-- |+-- Create a new asynchronous stream+--+create :: IO Stream+create = resultIfOk =<< cudaStreamCreate++{# fun unsafe cudaStreamCreate+  { alloca- `Stream' peekStream* } -> `Status' cToEnum #}+  where peekStream = liftM Stream . peekIntConv+++-- |+-- Destroy and clean up an asynchronous stream+--+destroy :: Stream -> IO ()+destroy s = nothingIfOk =<< cudaStreamDestroy s++{# fun unsafe cudaStreamDestroy+  { useStream `Stream' } -> `Status' cToEnum #}+++-- |+-- Determine if all operations in a stream have completed+--+finished   :: Stream -> IO Bool+finished s =+  cudaStreamQuery s >>= \rv -> do+  case rv of+      Success  -> return True+      NotReady -> return False+      _        -> resultIfOk (rv,undefined)++{# fun unsafe cudaStreamQuery+  { useStream `Stream' } -> `Status' cToEnum #}+++-- |+-- Block until all operations in a Stream have been completed+--+block :: Stream -> IO ()+block s = nothingIfOk =<< cudaStreamSynchronize s++{# fun unsafe cudaStreamSynchronize+  { useStream `Stream' } -> `Status' cToEnum #}+
+ Foreign/CUDA/Runtime/Thread.chs view
@@ -0,0 +1,55 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Runtime.Thread+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Thread management routines+--+--------------------------------------------------------------------------------+++module Foreign.CUDA.Runtime.Thread+  (+    sync, exit+  )+  where++#include <cuda_runtime_api.h>+{# context lib="cudart" #}++-- Friends+import Foreign.CUDA.Runtime.Error+import Foreign.CUDA.Internal.C2HS++-- System+import Foreign.C+++--------------------------------------------------------------------------------+-- Thread management+--------------------------------------------------------------------------------++-- |+-- Block until the device has completed all preceding requested tasks. Returns+-- an error if one of the tasks fails.+--+sync :: IO ()+sync =  nothingIfOk =<< cudaThreadSynchronize++{# fun unsafe cudaThreadSynchronize+  { } -> `Status' cToEnum #}+++-- |+-- Explicitly clean up all runtime related resources associated with the calling+-- host thread. Any subsequent API call re-initialised the environment.+-- Implicitly called on thread exit.+--+exit :: IO ()+exit =  nothingIfOk =<< cudaThreadExit++{# fun unsafe cudaThreadExit+  { } -> `Status' cToEnum #}+
+ Foreign/CUDA/Runtime/Utils.chs view
@@ -0,0 +1,48 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+-- |+-- Module    : Foreign.CUDA.Runtime.Utils+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Utility functions+--+--------------------------------------------------------------------------------++module Foreign.CUDA.Runtime.Utils+  (+    runtimeVersion, driverVersion+  )+  where++#include <cuda_runtime_api.h>+{# context lib="cudart" #}++-- Friends+import Foreign.CUDA.Runtime.Error+import Foreign.CUDA.Internal.C2HS++-- System+import Foreign+import Foreign.C+++-- |+-- Return the version number of the installed CUDA driver+--+runtimeVersion :: IO Int+runtimeVersion =  resultIfOk =<< cudaRuntimeGetVersion++{# fun unsafe cudaRuntimeGetVersion+  { alloca- `Int' peekIntConv* } -> `Status' cToEnum #}+++-- |+-- Return the version number of the installed CUDA runtime+--+driverVersion :: IO Int+driverVersion =  resultIfOk =<< cudaDriverGetVersion++{# fun unsafe cudaDriverGetVersion+  { alloca- `Int' peekIntConv* } -> `Status' cToEnum #}+
+ LICENSE view
@@ -0,0 +1,24 @@+Copyright (c) 2009 Trevor L. McDonell, University of New South Wales.+All rights reserved.++Redistribution and use in source and binary forms, with or without+modification, are permitted provided that the following conditions are met:+    * Redistributions of source code must retain the above copyright+      notice, this list of conditions and the following disclaimer.+    * Redistributions in binary form must reproduce the above copyright+      notice, this list of conditions and the following disclaimer in the+      documentation and/or other materials provided with the distribution.+    * Neither the name of the University of New South Wales nor the+      names of its contributors may be used to endorse or promote products+      derived from this software without specific prior written permission.++THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ''AS IS'' AND ANY+EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED+WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE+DISCLAIMED. IN NO EVENT SHALL COPYRIGHT HOLDERS BE LIABLE FOR ANY+DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES+(INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;+LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND+ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS+SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ Setup.hs view
@@ -0,0 +1,4 @@+import Distribution.Simple++main :: IO ()+main =  defaultMain
+ cbits/stubs.c view
@@ -0,0 +1,16 @@+/*+ * Extra bits for CUDA binding+ */++#include "cbits/stubs.h"+++cudaError_t+cudaConfigureCallSimple(int gx, int gy, int bx, int by, int bz, size_t sharedMem, cudaStream_t stream)+{+    dim3 gridDim  = {gx,gy,1};+    dim3 blockDim = {bx,by,bz};++    return cudaConfigureCall(gridDim, blockDim, sharedMem, stream);+}+
+ cbits/stubs.h view
@@ -0,0 +1,21 @@+/*+ * Extra bits for CUDA bindings+ */++#ifndef C_STUBS_H+#define C_STUBS_H++#include <cuda_runtime_api.h>++#ifdef __cplusplus+extern "C" {+#endif++cudaError_t+cudaConfigureCallSimple(int gx, int gy, int bx, int by, int bz, size_t sharedMem, cudaStream_t stream);+++#ifdef __cplusplus+}+#endif+#endif
+ configure view
@@ -0,0 +1,9 @@+#!/bin/sh++# substitute standard header path variables+if test -n "$CPPFLAGS" ; then+    echo "Found CPPFLAGS in environment: '$CPPFLAGS'"+    sed 's,@CPPFLAGS@,'"$CPPFLAGS"',g;s,@LDFLAGS@,'"$LDFLAGS"',g'  \+        < cuda.buildinfo.in > cuda.buildinfo+fi+
+ cuda.buildinfo.in view
@@ -0,0 +1,3 @@+ghc-options: -optc@CPPFLAGS@+cc-options:  @CPPFLAGS@+ld-options:  @LDFLAGS@
+ cuda.cabal view
@@ -0,0 +1,67 @@+Name:                   cuda+Version:                0.1+Synopsis:               A binding to the CUDA interface for programming NVIDIA GPUs+Description:    +    The CUDA library provides a direct, general purpose C-like SPMD programming+    model for NVIDIA graphics cards (G8x series onwards). This is a collection+    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 (tested with v2.3).+    .+    <http://developer.nvidia.com/object/cuda.html>+    .+    Note that on Snow Leopard, the c2hs preprocessor is confused by the notation+    for Apple's Blocks extension, so to work around this:+    .+    > cabal install --c2hs-option=-ccpp-4.0+    .++License:                BSD3+License-file:           LICENSE+Copyright:              Copyright (c) 2009-10. Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+Author:                 Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+Maintainer:             Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+Category:               Foreign+Cabal-version:          >=1.2++Build-type:             Configure+Extra-tmp-files:        cuda.buildinfo+Extra-source-files:     configure+                        cuda.buildinfo.in++Library+  Exposed-Modules:      Foreign.CUDA+                        Foreign.CUDA.Runtime+                        Foreign.CUDA.Runtime.Device+                        Foreign.CUDA.Runtime.Error+                        Foreign.CUDA.Runtime.Event+                        Foreign.CUDA.Runtime.Exec+                        Foreign.CUDA.Runtime.Marshal+                        Foreign.CUDA.Runtime.Ptr+                        Foreign.CUDA.Runtime.Stream+                        Foreign.CUDA.Runtime.Thread+                        Foreign.CUDA.Runtime.Utils++                        Foreign.CUDA.Driver+                        Foreign.CUDA.Driver.Context+                        Foreign.CUDA.Driver.Device+                        Foreign.CUDA.Driver.Error+                        Foreign.CUDA.Driver.Event+                        Foreign.CUDA.Driver.Exec+                        Foreign.CUDA.Driver.Marshal+                        Foreign.CUDA.Driver.Module+                        Foreign.CUDA.Driver.Stream+                        Foreign.CUDA.Driver.Utils++  Other-modules:        Foreign.CUDA.Internal.C2HS+                        Foreign.CUDA.Internal.Offsets++  Include-dirs:         .+  C-sources:            cbits/stubs.c++  Build-tools:          c2hs, hsc2hs+  Build-depends:        base >= 3 && < 5, haskell98, bytestring, extensible-exceptions+  Extensions:+  ghc-options:          -Wall+  extra-libraries:      cuda cudart+
+ examples/Makefile view
@@ -0,0 +1,18 @@+ifneq ($(emu),1)+    PROJECTS := $(shell find src -name Makefile)+else+    PROJECTS := $(shell find src -name Makefile | xargs grep -L 'USEDRVAPI')+endif+++%.do :+	$(MAKE) -C $(dir $*) $(MAKECMDGOALS)++all : $(addsuffix .do,$(PROJECTS))+	@echo "Finished building all"++clean : $(addsuffix .do,$(PROJECTS))+	@echo "Finished cleaning all"++clobber : $(addsuffix .do,$(PROJECTS))+	@echo "Finished cleaning all"
+ examples/common/common.mk view
@@ -0,0 +1,413 @@+# ------------------------------------------------------------------------------+#+# Copyright 1993-2009 NVIDIA Corporation.  All rights reserved.+#+# NVIDIA Corporation and its licensors retain all intellectual property and +# proprietary rights in and to this software and related documentation. +# Any use, reproduction, disclosure, or distribution of this software +# and related documentation without an express license agreement from+# NVIDIA Corporation is strictly prohibited.+#+# Please refer to the applicable NVIDIA end user license agreement (EULA) +# associated with this source code for terms and conditions that govern +# your use of this NVIDIA software.+#+# ------------------------------------------------------------------------------++# ------------------------------------------------------------------------------+# Common Haskell/CUDA build system+# ------------------------------------------------------------------------------++.SUFFIXES : .cu .cu_dbg.o .c_dbg.o .cpp_dbg.o .cu_rel.o .c_rel.o .cpp_rel.o .cubin .ptx++# No CUDA compatible device present+# MAIN_DEVICE     := $(shell ghc -e "Control.Monad.liftM (Data.Either.either id Foreign.CUDA.deviceName) (Foreign.CUDA.props 0)")+# ifeq ($(MAIN_DEVICE),"Device Emulation (CPU)")+#     emu		:= 1+# endif++# Add new SM Versions here as devices with new Compute Capability are released+SM_VERSIONS     := sm_10 sm_11 sm_12 sm_13++# detect OS+OSUPPER         := $(shell uname -s 2>/dev/null | tr [:lower:] [:upper:])+OSLOWER         := $(shell uname -s 2>/dev/null | tr [:upper:] [:lower:])+DARWIN          := $(strip $(findstring DARWIN, $(OSUPPER)))+ifneq ($(DARWIN),)+    SNOWLEOPARD := $(strip $(findstring 10.6, $(shell egrep "<string>10\.6.*</string>" /System/Library/CoreServices/SystemVersion.plist)))+endif++# detect if 32 bit or 64 bit system+HP_64           := $(strip $(shell uname -m | grep 64))++# Basic directory setup+CUDA_INSTALL_PATH ?= /usr/local/cuda+CUDA_SDK_PATH     ?= /Developer/GPU\ Computing/C++SRCDIR          ?= .+ROOTDIR         ?= ..+ROOTBINDIR      := $(ROOTDIR)/../bin+BINDIR		:= $(ROOTBINDIR)/$(OSLOWER)+ROOTOBJDIR      := obj+LIBDIR          := $(ROOTDIR)/../lib+COMMONDIR	:= $(ROOTDIR)/../common++# Compilers+NVCC            := $(CUDA_INSTALL_PATH)/bin/nvcc+GHC             := ghc+C2HS            := c2hs+HSC2HS          := hsc2hs+CXX             := g+++CC              := gcc+LINK            := g++ -fPIC++# Includes+INCLUDES        += -I. -I$(CUDA_INSTALL_PATH)/include -I$(COMMONDIR)/include -I$(COMMONDIR)/src++# architecture flag for cubin build+CUBIN_ARCH_FLAG :=++# Warning flags+CXXWARN_FLAGS := \+        -W -Wall \+        -Wimplicit \+        -Wswitch \+        -Wformat \+        -Wchar-subscripts \+        -Wparentheses \+        -Wmultichar \+        -Wtrigraphs \+        -Wpointer-arith \+        -Wcast-align \+        -Wreturn-type \+        -Wno-unused-function++CWARN_FLAGS := $(CXXWARN_FLAGS) \+        -Wstrict-prototypes \+        -Wmissing-prototypes \+        -Wmissing-declarations \+        -Wnested-externs \+        -Wmain++GHCWARN_FLAGS := \+        -Wall++# cross-compilation flags+ifeq ($(x86_64),1)+    NVCCFLAGS	+= -m64+    ifneq ($(DARWIN),)+        CXX_ARCH_FLAGS 	+= -arch x86_64+    else+        CXX_ARCH_FLAGS	+= -m64+    endif+else+ifeq ($(i386),1)+    NVCCFLAGS	+= -m32+    ifneq ($(DARWIN),)+        CXX_ARCH_FLAGS	+= -arch i386+    else+        CXX_ARCH_FLAGS	+= -m32+    endif+else+    ifneq ($(SNOWLEOPARD),)+        NVCCFLAGS	+= -m32+        CXX_ARCH_FLAGS	+= -arch i386 -m32+    endif+endif+endif++# Compiler-specific flags+NVCCFLAGS       +=+GHCFLAGS        += $(GHCWARN_FLAGS) -i$(SRCDIR) -i$(COMMONDIR)/src -i$(OBJDIR) -odir $(OBJDIR) -hidir $(OBJDIR) --make+CXXFLAGS        += $(CXXWARN_FLAGS) $(CXX_ARCH_FLAGS)+CFLAGS          += $(CWARN_FLAGS) $(CXX_ARCH_FLAGS)+LINK		+= $(CXX_ARCH_FLAGS)++# Common flags+COMMONFLAGS     += $(INCLUDES) -DUNIX++# Debug/release configuration+ifeq ($(dbg),1)+    COMMONFLAGS += -g+    GHCFLAGS	+= -prof -auto-all -fhpc+    NVCCFLAGS   += -D_DEBUG+    CXXFLAGS    += -D_DEBUG+    CFLAGS      += -D_DEBUG+    BINSUBDIR   := debug+    LIBSUFFIX   := D+else+    COMMONFLAGS += -O2+    GHCFLAGS	+= -O2+    BINSUBDIR   := release+    LIBSUFFIX   :=+    NVCCFLAGS   += --compiler-options -fno-strict-aliasing+    CXXFLAGS    += -fno-strict-aliasing+    CFLAGS      += -fno-strict-aliasing+endif++# Device emulation configuration+ifeq ($(emu),1)+    NVCCFLAGS   += -deviceemu+    CUDACCFLAGS +=+    BINSUBDIR   := emu$(BINSUBDIR)+    LIBSUFFIX   := $(LIBSUFFIX)_emu+    # consistency, makes developing easier+    C2HSFLAGS   += --cppopts=-D__DEVICE_EMULATION__+    GHCFLAGS    += -D__DEVICE_EMULATION__+    CXXFLAGS    += -D__DEVICE_EMULATION__+    CFLAGS      += -D__DEVICE_EMULATION__+endif++# architecture flag for cubin build+CUBIN_ARCH_FLAG :=++# Libraries+LIB             += -L$(LIBDIR) $(addprefix -l,$(EXTRALIBS))+ifeq ($(HP_64),)+   LIB          += -L$(CUDA_INSTALL_PATH)/lib+else+   LIB          += -L$(CUDA_INSTALL_PATH)/lib64+endif++# If dynamically linking to CUDA and CUDART, we exclude the libraries from the LIB+ifeq ($(USECUDADYNLIB),1)+    LIB         += -ldl -rdynamic+else+    # static linking, we will statically link against CUDA and CUDART+    ifeq ($(USEDRVAPI),1)+        LIB     += -lcuda+    else+        LIB     += -lcudart+    endif+endif++ifeq ($(USECUFFT),1)+    ifeq ($(emu),1)+        LIB     += -lcufftemu+    else+        LIB     += -lcufft+    endif+endif++ifeq ($(USECUBLAS),1)+    ifeq ($(emu),1)+        LIB     += -lcublasemu+    else+        LIB     += -lcublas+    endif+endif++ifeq ($(USECUDPP),1)+    CUDPPLIB := cudpp++    ifneq ($(HP_64),)+        CUDPPLIB := $(CUDPPLIB)64+    endif++    ifeq ($(emu),1)+        CUDPPLIB := $(CUDPPLIB)_emu+    endif++    LIB     += -l$(CUDPPLIB)+endif++# Library/executable configuration+ifneq ($(STATIC_LIB),)+    TARGETDIR   := $(LIBDIR)+    TARGET      := $(LIBDIR)/$(basename $(STATIC_LIB))$(LIBSUFFIX)$(suffix $(STATIC_LIB))+    LINKLINE     = ar rucv $(TARGET) $(OBJS); ranlib $(TARGET)+else+ifneq ($(DYNAMIC_LIB),)+    TARGETDIR   := $(LIBDIR)+    TARGET      := $(LIBDIR)/$(basename $(DYNAMIC_LIB))$(LIBSUFFIX)$(suffix $(DYNAMIC_LIB))+    CFLAGS      += -fPIC+    CXXFLAGS    += -fPIC+    NVCCFLAGS   += -Xcompiler -fPIC+    ifneq ($(DARWIN),)+    LINKLINE     = $(LINK) -dynamiclib -o $(TARGET) -install_name "@rpath/$(notdir $(TARGET))" $(OBJS) $(LIB)+    else+    LINKLINE     = $(LINK) -shared -o $(TARGET) -Wl,-rpath,$(notdir $(TARGET)) $(OBJS) $(LIB)+    endif+else+    TARGETDIR   := $(BINDIR)/$(BINSUBDIR)+    TARGET      := $(TARGETDIR)/$(EXECUTABLE)+    ifneq ($(HSMAIN),)+        ifeq ($(suffix $(HSMAIN)),.chs)+            FOO      := $(HSMAIN)+            CHSFILES += $(FOO)+            HSMAIN    = $(OBJDIR)/$(basename $(FOO)).hs+        endif++        ifeq ($(dbg),1)+#            OBJS     += $(OBJDIR)/ptxvars.cu.o+            LINKLINE  = $(GHC) -o $(TARGET) $(LIB) $(OBJS) $(GHCFLAGS) $(HSMAIN)+        else+            LINKLINE  = $(GHC) -o $(TARGET) $(LIB) $(OBJS) $(GHCFLAGS) $(HSMAIN)+        endif+    else+        LINKLINE = $(LINK) -o $(TARGET) $(OBJS) $(LIB)+    endif+endif+endif++# check if verbose+ifeq ($(verbose),1)+    VERBOSE     :=+else+    VERBOSE     := @+endif+++# ------------------------------------------------------------------------------+# Check for input flags and set compiler flags appropriately+# ------------------------------------------------------------------------------+ifeq ($(fastmath),1)+    NVCCFLAGS   += -use_fast_math+endif++ifeq ($(keep),1)+    NVCCFLAGS       += -keep+    NVCC_KEEP_CLEAN := *.i* *.cubin *.cu.c *.cudafe* *.fatbin.c *.ptx+endif++ifdef maxregisters+    NVCCFLAGS   += -maxrregcount $(maxregisters)+endif++# Add cudacc flags+NVCCFLAGS += $(CUDACCFLAGS)++# Add common flags+NVCCFLAGS += $(COMMONFLAGS)+CXXFLAGS  += $(COMMONFLAGS)+CFLAGS    += $(COMMONFLAGS)++ifeq ($(nvcc_warn_verbose),1)+    NVCCFLAGS += $(addprefix --compiler-options ,$(CXXWARN_FLAGS))+    NVCCFLAGS += --compiler-options -fno-strict-aliasing+endif+++# ------------------------------------------------------------------------------+# Set up object files+# ------------------------------------------------------------------------------+OBJDIR  := $(ROOTOBJDIR)/$(BINSUBDIR)+OBJS    += $(patsubst %.cpp,$(OBJDIR)/%.cpp.o,$(notdir $(CCFILES)))+OBJS    += $(patsubst %.c,$(OBJDIR)/%.c.o,$(notdir $(CFILES)))+OBJS    += $(patsubst %.cu,$(OBJDIR)/%.cu.o,$(notdir $(CUFILES)))++# ------------------------------------------------------------------------------+# Set up preprocessed Haskell files+# ------------------------------------------------------------------------------+DEPS	+= $(patsubst %.chs,$(OBJDIR)/%.hs,$(notdir $(CHSFILES)))+DEPS	+= $(patsubst %.hsc,$(OBJDIR)/%.hs,$(notdir $(HSCFILES)))++# ------------------------------------------------------------------------------+# Set up cubin output files+# ------------------------------------------------------------------------------+CUBINDIR := $(SRCDIR)/data+CUBINS   += $(patsubst %.cu,$(CUBINDIR)/%.cubin,$(notdir $(CUBINFILES)))++# ------------------------------------------------------------------------------+# Set up PTX output files+# ------------------------------------------------------------------------------+PTXDIR  := $(SRCDIR)/data+PTXBINS += $(patsubst %.cu,$(PTXDIR)/%.ptx,$(notdir $(PTXFILES)))+++# ------------------------------------------------------------------------------+# Rules+# ------------------------------------------------------------------------------+default: $(TARGET)++%.subdir :+	$(VERBOSE)$(MAKE) -C $* $(MAKECMDGOALS)++$(OBJDIR)/%.c.o : $(SRCDIR)/%.c $(C_DEPS)+	$(VERBOSE)$(CC) $(CFLAGS) -o $@ -c $<++$(OBJDIR)/%.cpp.o : $(SRCDIR)/%.cpp $(C_DEPS)+	$(VERBOSE)$(CXX) $(CXXFLAGS) -o $@ -c $<++$(OBJDIR)/ptxvars.cu.o: makedirectories+	$(VERBOSE)$(NVCC) -g -G --host-compilation=C -D__DEVICE_LAUNCH_PARAMETERS_H__ -Xptxas -fext -o $@ -c $(CUDA_INSTALL_PATH)/bin/ptxvars.cu++$(OBJDIR)/%.cu.o : $(SRCDIR)/%.cu $(CU_DEPS)+	$(VERBOSE)$(NVCC) $(NVCCFLAGS) $(SMVERSIONFLAGS) -o $@ -c $<++$(OBJDIR)/%.hs : $(SRCDIR)/%.chs+	$(VERBOSE)$(C2HS) $(C2HSFLAGS) --include=$(OBJDIR) $(addprefix --cppopts=,$(INCLUDES)) --output-dir=$(OBJDIR) --output=$(notdir $@) $<++$(OBJDIR)/%.hs : $(SRCDIR)/%.hsc+	$(VERBOSE)$(HSC2HS) $(INCLUDES) -o $@ $<++$(CUBINDIR)/%.cubin : $(SRCDIR)/%.cu cubindirectory+	$(VERBOSE)$(NVCC) $(CUBIN_ARCH_FLAG) $(NVCCFLAGS) $(SMVERSIONFLAGS) -o $@ -cubin $<++$(PTXDIR)/%.ptx : $(SRCDIR)/%.cu ptxdirectory+	$(VERBOSE)$(NVCC) $(CUBIN_ARCH_FLAG) $(NVCCFLAGS) $(SMVERSIONFLAGS) -o $@ -ptx $<++#+# The following definition is a template that gets instantiated for each SM+# version (sm_10, sm_13, etc.) stored in SMVERSIONS.  It does 2 things:+# 1. It adds to OBJS a .cu_sm_XX.o for each .cu file it finds in CUFILES_sm_XX.+# 2. It generates a rule for building .cu_sm_XX.o files from the corresponding+#    .cu file.+#+# The intended use for this is to allow Makefiles that use common.mk to compile+# files to different Compute Capability targets (aka SM arch version).  To do+# so, in the Makefile, list files for each SM arch separately, like so:+#+# CUFILES_sm_10 := mycudakernel_sm10.cu app.cu+# CUFILES_sm_12 := anothercudakernel_sm12.cu+#+define SMVERSION_template+OBJS += $(patsubst %.cu,$(OBJDIR)/%.cu_$(1).o,$(notdir $(CUFILES_$(1))))+$(OBJDIR)/%.cu_$(1).o : $(SRCDIR)/%.cu $(CU_DEPS)+	$(VERBOSE)$(NVCC) -o $$@ -c $$< $(NVCCFLAGS) -arch $(1)+endef++# This line invokes the above template for each arch version stored in+# SM_VERSIONS.  The call funtion invokes the template, and the eval+# function interprets it as make commands.+$(foreach smver,$(SM_VERSIONS),$(eval $(call SMVERSION_template,$(smver))))++$(TARGET): makedirectories $(DEPS) $(OBJS) $(CUBINS) $(PTXBINS) Makefile $(addsuffix .subdir,$(SUBDIRS))+	$(VERBOSE)$(LINKLINE)++cubindirectory:+	$(VERBOSE)mkdir -p $(CUBINDIR)++ptxdirectory:+	$(VERBOSE)mkdir -p $(PTXDIR)++makedirectories:+	$(VERBOSE)mkdir -p $(LIBDIR)+	$(VERBOSE)mkdir -p $(OBJDIR)+	$(VERBOSE)mkdir -p $(TARGETDIR)+++tidy : $(addsuffix .subdir,$(SUBDIRS))+	$(VERBOSE)find . | egrep "#"  | xargs rm -f+	$(VERBOSE)find . | egrep "\~" | xargs rm -f++clean : tidy+	$(VERBOSE)rm -f $(OBJS)+	$(VERBOSE)rm -f $(CUBINS)+	$(VERBOSE)rm -f $(PTXBINS)+	$(VERBOSE)rm -f $(TARGET)+	$(VERBOSE)rm -f $(NVCC_KEEP_CLEAN)+	$(VERBOSE)rm -f $(ROOTBINDIR)/$(OSLOWER)/$(BINSUBDIR)/*.ppm+	$(VERBOSE)rm -f $(ROOTBINDIR)/$(OSLOWER)/$(BINSUBDIR)/*.pgm+	$(VERBOSE)rm -f $(ROOTBINDIR)/$(OSLOWER)/$(BINSUBDIR)/*.bin+	$(VERBOSE)rm -f $(ROOTBINDIR)/$(OSLOWER)/$(BINSUBDIR)/*.bmp++clobber : clean+	$(VERBOSE)rm -rf $(ROOTOBJDIR)++spotless:+	$(VERBOSE)rm -rf .hpc+	$(VERBOSE)rm -f $(EXECUTABLE).{aux,hp,prof,ps}+	$(VERBOSE)rm -f *.html+	$(VERBOSE)find . -name "*.tix" -print0 | xargs -0 rm -f+
+ examples/common/include/cudpp/LICENSE view
@@ -0,0 +1,25 @@+Copyright (c) 2007-2009 The Regents of the University of California, Davis+campus ("The Regents") and NVIDIA Corporation ("NVIDIA"). All rights reserved.++Redistribution and use in source and binary forms, with or without modification, +are permitted provided that the following conditions are met:++    * Redistributions of source code must retain the above copyright notice, +      this list of conditions and the following disclaimer.+    * Redistributions in binary form must reproduce the above copyright notice, +      this list of conditions and the following disclaimer in the documentation +      and/or other materials provided with the distribution.+    * Neither the name of the The Regents, nor NVIDIA, nor the names of its +      contributors may be used to endorse or promote products derived from this +      software without specific prior written permission.++THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND +ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED +WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. +IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, +INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, +BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF +LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE +OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF +ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ examples/common/include/cudpp/cudpp_globals.h view
@@ -0,0 +1,56 @@+// -------------------------------------------------------------+// cuDPP -- CUDA Data Parallel Primitives library+// -------------------------------------------------------------+// $Revision: 5632 $+// $Date: 2009-07-01 14:36:01 +1000 (Wed, 01 Jul 2009) $+// -------------------------------------------------------------+// This source code is distributed under the terms of license.txt in+// the root directory of this source distribution.+// -------------------------------------------------------------++/**+ * @file+ * cudpp_globals.h+ *+ * @brief Global declarations defining machine characteristics of GPU target+ * These are currently set for best performance on G8X GPUs.  The optimal+ * parameters may change on future GPUs. In the future, we hope to make+ * CUDPP a self-tuning library.+ */++#ifndef __CUDPP_GLOBALS_H__+#define __CUDPP_GLOBALS_H__++const int NUM_BANKS = 16;                        /**< Number of shared memory banks */+const int LOG_NUM_BANKS = 4;                     /**< log_2(NUM_BANKS) */+const int CTA_SIZE = 128;                        /**< Number of threads in a CTA */+const int WARP_SIZE = 32;                        /**< Number of threads in a warp */+const int LOG_CTA_SIZE = 7;                      /**< log_2(CTA_SIZE) */+const int LOG_WARP_SIZE = 5;                     /**< log_2(WARP_SIZE) */+const int LOG_SIZEOF_FLOAT = 2;                  /**< log_2(sizeof(float)) */+const int SCAN_ELTS_PER_THREAD = 8;              /**< Number of elements per scan thread */+const int SEGSCAN_ELTS_PER_THREAD = 8;     /**< Number of elements per segmented scan thread */++const int maxSharedMemoryPerBlock = 16384; /**< Number of bytes of shared+                                              memory in each block */+const int maxThreadsPerBlock = CTA_SIZE;   /**< Maximum number of+                                             * threads in a CTA */++#define AVOID_BANK_CONFLICTS /**< Set if by default, we want our+                              * shared memory allocation to perform+                              * additional computation to avoid bank+                              * conflicts */++#ifdef AVOID_BANK_CONFLICTS+#define CONFLICT_FREE_OFFSET(index) ((index) >> LOG_NUM_BANKS)+#else+#define CONFLICT_FREE_OFFSET(index) (0)+#endif++#endif // __CUDPP_GLOBALS_H__++// Leave this at the end of the file+// Local Variables:+// mode:c+++// c-file-style: "NVIDIA"+// End:
+ examples/common/include/cudpp/shared_mem.h view
@@ -0,0 +1,115 @@+// -------------------------------------------------------------+// cuDPP -- CUDA Data Parallel Primitives library+// -------------------------------------------------------------+// $Revision: 5633 $+// $Date: 2009-07-01 15:02:51 +1000 (Wed, 01 Jul 2009) $+// -------------------------------------------------------------+// This source code is distributed under the terms of license.txt+// in the root directory of this source distribution.+// -------------------------------------------------------------++/**+ * @file+ * sharedmem.h+ *+ * @brief Shared memory declaration struct for templatized types.+ *+ * Because dynamically sized shared memory arrays are declared "extern" in CUDA,+ * we can't templatize their types directly.  To get around this, we declare a+ * simple wrapper struct that will declare the extern array with a different+ * name depending on the type.  This avoids linker errors about multiple+ * definitions.+ *+ * To use dynamically allocated shared memory in a templatized __global__ or+ * __device__ function, just replace code like this:+ *+ * <pre>+ *  template<class T>+ *  __global__ void+ *  foo( T* d_out, T* d_in)+ *  {+ *      // Shared mem size is determined by the host app at run time+ *      extern __shared__  T sdata[];+ *      ...+ *      doStuff(sdata);+ *      ...+ *  }+ * </pre>+ *+ *  With this+ * <pre>+ *  template<class T>+ *  __global__ void+ *  foo( T* d_out, T* d_in)+ *  {+ *      // Shared mem size is determined by the host app at run time+ *      SharedMemory<T> smem;+ *      T* sdata = smem.getPointer();+ *      ...+ *      doStuff(sdata);+ *      ...+ *  }+ * </pre>+ */++#ifndef __SHARED_MEM_H__+#define __SHARED_MEM_H__+++/** @brief Wrapper class for templatized dynamic shared memory arrays.+  *+  * This struct uses template specialization on the type \a T to declare+  * a differently named dynamic shared memory array for each type+  * (\code extern __shared__ T s_type[] \endcode).+  *+  * Currently there are specializations for the following types:+  * \c int, \c uint, \c char, \c uchar, \c short, \c ushort, \c long,+  * \c unsigned long, \c bool, \c float, and \c double. One can also specialize it+  * for user defined types.+  */+template <typename T>+struct SharedMemory+{+    /** Return a pointer to the runtime-sized shared memory array. **/+    __device__ T* getPointer()+    {+        extern __device__ void Error_UnsupportedType(); // Ensure that we won't compile any un-specialized types+        Error_UnsupportedType();+        return (T*)0;+    }+    // TODO: Use operator overloading to make this class look like a regular array+};++// Following are the specializations for the following types.+// int, uint, char, uchar, short, ushort, long, ulong, bool, float, and double+// One could also specialize it for user-defined types.++#define SPEC_SHAREDMEM(T, name)                                                         \+    template <> struct SharedMemory <T>                                                 \+    {                                                                                   \+        __device__ T* getPointer() { extern __shared__ T s_##name[]; return s_##name; } \+    }++SPEC_SHAREDMEM(int,    int);+SPEC_SHAREDMEM(char,   char);+SPEC_SHAREDMEM(long,   long);+SPEC_SHAREDMEM(short,  short);+SPEC_SHAREDMEM(bool,   bool);+SPEC_SHAREDMEM(float,  float);+SPEC_SHAREDMEM(double, double);++SPEC_SHAREDMEM(unsigned int,   uint);+SPEC_SHAREDMEM(unsigned char,  uchar);+SPEC_SHAREDMEM(unsigned long,  ulong);+SPEC_SHAREDMEM(unsigned short, ushort);++SPEC_SHAREDMEM(uchar4, uchar4);++#undef SPEC_SHAREDMEM+#endif // __SHARED_MEM_H__++// Leave this at the end of the file+// Local Variables:+// mode:c+++// c-file-style: "NVIDIA"+// End:
+ examples/common/include/cudpp/type_vector.h view
@@ -0,0 +1,51 @@+// -------------------------------------------------------------+// cuDPP -- CUDA Data Parallel Primitives library+// -------------------------------------------------------------+// $Revision: 5632 $+// $Date: 2009-07-01 14:36:01 +1000 (Wed, 01 Jul 2009) $+// -------------------------------------------------------------+// This source code is distributed under the terms of license.txt in+// the root directory of this source distribution.+// -------------------------------------------------------------++#ifndef __TYPE_VECTOR_H__+#define __TYPE_VECTOR_H__++/** @brief Utility template struct for generating small vector types from scalar types+  *+  * Given a base scalar type (\c int, \c float, etc.) and a vector length (1 through 4) as+  * template parameters, this struct defines a vector type (\c float3, \c int4, etc.) of the+  * specified length and base type.  For example:+  * \code+  * template <class T>+  * __device__ void myKernel(T *data)+  * {+  *     typeToVector<T,4>::Result myVec4;             // create a vec4 of type T+  *     myVec4 = (typeToVector<T,4>::Result*)data[0]; // load first element of data as a vec4+  * }+  * \endcode+  *+  * This functionality is implemented using template specialization.  Currently specializations+  * for int, float, and unsigned int vectors of lengths 2-4 are defined.  Note that this results+  * in types being generated at compile time -- there is no runtime cost.  typeToVector is used by+  * the optimized scan \c __device__ functions in scan_cta.cu.+  */+template <typename T, int N>+struct typeToVector+{+    typedef T Result;+};++#define TYPE_VECTOR(type, name)                                                 \+    template <> struct typeToVector<type, 2> { typedef name##2 Result; };       \+    template <> struct typeToVector<type, 3> { typedef name##3 Result; };       \+    template <> struct typeToVector<type, 4> { typedef name##4 Result; }++TYPE_VECTOR(int,         int);+TYPE_VECTOR(float,       float);+TYPE_VECTOR(unsigned int, uint);+++#undef TYPE_VECTOR+#endif+
+ examples/common/include/operator.h view
@@ -0,0 +1,114 @@+/* -----------------------------------------------------------------------------+ *+ * Module    : Operator+ * Copyright : (c) 2009 Trevor L. McDonell+ * License   : BSD+ *+ * A template class for unary/binary kernel operations+ *+ * ---------------------------------------------------------------------------*/+++#ifndef __OPERATOR_H__+#define __OPERATOR_H__++#include <float.h>+#include <limits.h>+++/*+ * Template class for an operation that can be mapped over an array.+ */+template <typename Ta, typename Tb=Ta>+class Functor+{+public:+    /*+     * Apply the operation to the given operand.+     */+    static __device__ Tb apply(const Ta &x);+};++template <typename Ta, typename Tb>+class fromIntegral : Functor<Ta, Tb>+{+public:+    static __device__ Tb apply(const Ta &x) { return (Tb) x; }+};+++/*+ * Template class for binary operators. Certain algorithms may require the+ * operator to be associative (that is, Ta == Tb), such as parallel scan and+ * reduction.+ *+ * As this is template code, it should compile down to something efficient...+ */+template <typename Ta, typename Tb=Ta, typename Tc=Ta>+class BinaryOp+{+public:+    /*+     * Apply the operation to the given operands.+     */+    static __device__ Tc apply(const Ta &a, const Tb &b);++    /*+     * Return an identity element for the type Tc.+     *+     * This may have special meaning for a given implementation, for example a+     * `max' operation over integers may want to return INT_MIN.+     */+    static __device__ Tc identity();+};++/*+ * Return the minimum or maximum value of a type+ */+template <typename T> inline __device__ T getMin();+template <typename T> inline __device__ T getMax();++#define SPEC_MINMAX(type,vmin,vmax)                                            \+    template <> inline __device__ type getMin() { return vmin; };              \+    template <> inline __device__ type getMax() { return vmax; }               \++SPEC_MINMAX(float,        -FLT_MAX,  FLT_MAX);+SPEC_MINMAX(int,           INT_MIN,  INT_MAX);+SPEC_MINMAX(char,          CHAR_MIN, CHAR_MAX);+SPEC_MINMAX(unsigned int,  0,        UINT_MAX);+SPEC_MINMAX(unsigned char, 0,        UCHAR_MAX);+++/*+ * Basic binary arithmetic operations. We take advantage of automatic type+ * promotion to keep the parameters general.+ */+#define BASIC_OP(name,expr,id)                                                 \+    template <typename Ta, typename Tb=Ta, typename Tc=Ta>                     \+    class name : BinaryOp<Ta, Tb, Tc>                                          \+    {                                                                          \+    public:                                                                    \+        static __device__ Tc apply(const Ta &a, const Tb &b) { return expr; }  \+        static __device__ Tc identity() { return id; }                         \+    }++#define LOGICAL_OP(name,expr,id)                                               \+    template <typename Ta, typename Tb=Ta>                                     \+    class name : BinaryOp<Ta, Tb, bool>                                        \+    {                                                                          \+    public:                                                                    \+        static __device__ bool apply(const Ta &a, const Tb &b) { return expr; }\+        static __device__ bool identity() { return id; }                       \+    }++BASIC_OP(Plus,  a + b,    0);+BASIC_OP(Times, a * b,    1);+BASIC_OP(Min,   min(a,b), getMax<Ta>());+BASIC_OP(Max,   max(a,b), getMin<Ta>());++LOGICAL_OP(Eq,  a == b,   false);++#undef SPEC_MINMAX+#undef BASIC_OP+#endif+
+ examples/common/include/utils.h view
@@ -0,0 +1,146 @@+/* -----------------------------------------------------------------------------+ *+ * Module    : Utils+ * Copyright : (c) 2009 Trevor L. McDonell+ * License   : BSD+ *+ * ---------------------------------------------------------------------------*/+++#ifndef __UTILS_H__+#define __UTILS_H__++#include <math.h>+#include <stdio.h>+++/*+ * Core assert function. Don't let this escape...+ */+#if defined(__CUDACC__) || !defined(__DEVICE_EMULATION__)+#define __assert(e, file, line) ((void)0)+#else+#define __assert(e, file, line) \+    ((void) fprintf (stderr, "%s:%u: failed assertion `%s'\n", file, line, e), abort())+#endif++/*+ * Test the given expression, and abort the program if it evaluates to false.+ * Only available in debug mode.+ */+#ifndef _DEBUG+#define assert(e)               ((void)0)+#else+#define assert(e)  \+    ((void) ((e) ? (void(0)) : __assert (#e, __FILE__, __LINE__)))+#endif+++/*+ * Macro to insert __syncthreads() in device emulation mode+ */+#ifdef __DEVICE_EMULATION__+#define __EMUSYNC               __syncthreads()+#else+#define __EMUSYNC+#endif+++/*+ * Check the return status of CUDA API calls, and abort with an appropriate+ * error string on failure.+ */+#define CUDA_SAFE_CALL_NO_SYNC(call)                                           \+    do {                                                                       \+        cudaError err = call;                                                  \+        if(cudaSuccess != err) {                                               \+            const char *str = cudaGetErrorString(err);                         \+            __assert(str, __FILE__, __LINE__);                                 \+        }                                                                      \+    } while (0)++#define CUDA_SAFE_CALL(call)                                                   \+    do {                                                                       \+        CUDA_SAFE_CALL_NO_SYNC(call);                                          \+        CUDA_SAFE_CALL_NO_SYNC(cudaThreadSynchronize());                       \+    } while (0)+++#ifdef __cplusplus+extern "C" {+#endif++/*+ * Determine if the input is a power of two+ */+inline bool+isPow2(unsigned int x)+{+    return ((x&(x-1)) == 0);+}+++/*+ * Compute the next highest power of two+ */+inline unsigned int+ceilPow2(unsigned int x)+{+#if 0+    --x;+    x |= x >> 1;+    x |= x >> 2;+    x |= x >> 4;+    x |= x >> 8;+    x |= x >> 16;+    return ++x;+#endif++    return (isPow2(x)) ? x : 1u << (int) ceil(log2((double)x));+}+++/*+ * Compute the next lowest power of two+ */+inline unsigned int+floorPow2(unsigned int x)+{+#if 0+    float nf = (float) n;+    return 1 << (((*(int*)&nf) >> 23) - 127);+#endif++    int exp;+    frexp(x, &exp);+    return 1 << (exp - 1);+}+++/*+ * computes next highest multiple of f from x+ */+inline unsigned int+multiple(unsigned int x, unsigned int f)+{+    return ((x + (f-1)) / f);+}+++/*+ * MS Excel-style CEIL() function. Rounds x up to nearest multiple of f+ */+inline unsigned int+ceiling(unsigned int x, unsigned int f)+{+    return multiple(x, f) * f;+}+++#undef __asert++#ifdef __cplusplus+}+#endif+#endif+
+ examples/common/src/C2HS.hs view
@@ -0,0 +1,224 @@+--  C->Haskell Compiler: Marshalling library+--+--  Copyright (c) [1999...2005] Manuel M T Chakravarty+--+--  Redistribution and use in source and binary forms, with or without+--  modification, are permitted provided that the following conditions are met:+-- +--  1. Redistributions of source code must retain the above copyright notice,+--     this list of conditions and the following disclaimer. +--  2. Redistributions in binary form must reproduce the above copyright+--     notice, this list of conditions and the following disclaimer in the+--     documentation and/or other materials provided with the distribution. +--  3. The name of the author may not be used to endorse or promote products+--     derived from this software without specific prior written permission. +--+--  THIS SOFTWARE IS PROVIDED BY THE AUTHOR ``AS IS'' AND ANY EXPRESS OR+--  IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES+--  OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN+--  NO EVENT SHALL THE AUTHOR BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +--  SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED+--  TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR+--  PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF+--  LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING+--  NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS+--  SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.+--+--- Description ---------------------------------------------------------------+--+--  Language: Haskell 98+--+--  This module provides the marshaling routines for Haskell files produced by +--  C->Haskell for binding to C library interfaces.  It exports all of the+--  low-level FFI (language-independent plus the C-specific parts) together+--  with the C->HS-specific higher-level marshalling routines.+--++module C2HS (++  -- * Re-export the language-independent component of the FFI +  module Foreign,++  -- * Re-export the C language component of the FFI+  module Foreign.C,++  -- * Composite marshalling functions+  withCStringLenIntConv, peekCStringLenIntConv, withIntConv, withFloatConv,+  peekIntConv, peekFloatConv, withBool, peekBool, withEnum, peekEnum,++  -- * Conditional results using 'Maybe'+  nothingIf, nothingIfNull,++  -- * Bit masks+  combineBitMasks, containsBitMask, extractBitMasks,++  -- * Conversion between C and Haskell types+  cIntConv, cFloatConv, cToBool, cFromBool, cToEnum, cFromEnum+) where +++import Foreign+       hiding       (Word)+		    -- Should also hide the Foreign.Marshal.Pool exports in+		    -- compilers that export them+import Foreign.C++import Monad        (liftM)+++-- Composite marshalling functions+-- -------------------------------++-- Strings with explicit length+--+withCStringLenIntConv    :: String -> (CStringLen -> IO a) -> IO a+withCStringLenIntConv s f = withCStringLen s $ \(p, n) -> f (p, cIntConv n)++peekCStringLenIntConv      :: CStringLen -> IO String+peekCStringLenIntConv (s,n) = peekCStringLen (s, cIntConv n)++-- Marshalling of numerals+--++withIntConv   :: (Storable b, Integral a, Integral b) +	      => a -> (Ptr b -> IO c) -> IO c+withIntConv    = with . cIntConv++withFloatConv :: (Storable b, RealFloat a, RealFloat b) +	      => a -> (Ptr b -> IO c) -> IO c+withFloatConv  = with . cFloatConv++peekIntConv   :: (Storable a, Integral a, Integral b) +	      => Ptr a -> IO b+peekIntConv    = liftM cIntConv . peek++peekFloatConv :: (Storable a, RealFloat a, RealFloat b) +	      => Ptr a -> IO b+peekFloatConv  = liftM cFloatConv . peek++-- Passing Booleans by reference+--++withBool :: (Integral a, Storable a) => Bool -> (Ptr a -> IO b) -> IO b+withBool  = with . fromBool++peekBool :: (Integral a, Storable a) => Ptr a -> IO Bool+peekBool  = liftM toBool . peek+++-- Passing enums by reference+--++withEnum :: (Enum a, Integral b, Storable b) => a -> (Ptr b -> IO c) -> IO c+withEnum  = with . cFromEnum++peekEnum :: (Enum a, Integral b, Storable b) => Ptr b -> IO a+peekEnum  = liftM cToEnum . peek+++{-+-- Storing of 'Maybe' values+-- -------------------------++instance Storable a => Storable (Maybe a) where+  sizeOf    _ = sizeOf    (undefined :: Ptr ())+  alignment _ = alignment (undefined :: Ptr ())++  peek p = do+	     ptr <- peek (castPtr p)+	     if ptr == nullPtr+	       then return Nothing+	       else liftM Just $ peek ptr++  poke p v = do+	       ptr <- case v of+		        Nothing -> return nullPtr+			Just v' -> new v'+               poke (castPtr p) ptr+-}++-- Conditional results using 'Maybe'+-- ---------------------------------++-- Wrap the result into a 'Maybe' type.+--+-- * the predicate determines when the result is considered to be non-existing,+--   ie, it is represented by `Nothing'+--+-- * the second argument allows to map a result wrapped into `Just' to some+--   other domain+--+nothingIf       :: (a -> Bool) -> (a -> b) -> a -> Maybe b+nothingIf p f x  = if p x then Nothing else Just $ f x++-- |Instance for special casing null pointers.+--+nothingIfNull :: (Ptr a -> b) -> Ptr a -> Maybe b+nothingIfNull  = nothingIf (== nullPtr)+++-- Support for bit masks+-- ---------------------++-- Given a list of enumeration values that represent bit masks, combine these+-- masks using bitwise disjunction.+--+combineBitMasks :: (Enum a, Bits b) => [a] -> b+combineBitMasks = foldl (.|.) 0 . map (fromIntegral . fromEnum)++-- Tests whether the given bit mask is contained in the given bit pattern+-- (i.e., all bits set in the mask are also set in the pattern).+--+containsBitMask :: (Bits a, Enum b) => a -> b -> Bool+bits `containsBitMask` bm = let bm' = fromIntegral . fromEnum $ bm+			    in+			    bm' .&. bits == bm'++-- |Given a bit pattern, yield all bit masks that it contains.+--+-- * This does *not* attempt to compute a minimal set of bit masks that when+--   combined yield the bit pattern, instead all contained bit masks are+--   produced.+--+extractBitMasks :: (Bits a, Enum b, Bounded b) => a -> [b]+extractBitMasks bits = +  [bm | bm <- [minBound..maxBound], bits `containsBitMask` bm]+++-- Conversion routines+-- -------------------++-- |Integral conversion+--+cIntConv :: (Integral a, Integral b) => a -> b+cIntConv  = fromIntegral++-- |Floating conversion+--+cFloatConv :: (RealFloat a, RealFloat b) => a -> b+cFloatConv  = realToFrac+-- As this conversion by default goes via `Rational', it can be very slow...+{-# RULES +  "cFloatConv/Float->Float"   forall (x::Float).  cFloatConv x = x;+  "cFloatConv/Double->Double" forall (x::Double). cFloatConv x = x+ #-}++-- |Obtain C value from Haskell 'Bool'.+--+cFromBool :: Num a => Bool -> a+cFromBool  = fromBool++-- |Obtain Haskell 'Bool' from C value.+--+cToBool :: Num a => a -> Bool+cToBool  = toBool++-- |Convert a C enumeration to Haskell.+--+cToEnum :: (Integral i, Enum e) => i -> e+cToEnum  = toEnum . cIntConv++-- |Convert a Haskell enumeration to C.+--+cFromEnum :: (Enum e, Integral i) => e -> i+cFromEnum  = cIntConv . fromEnum
+ examples/common/src/PrettyPrint.hs view
@@ -0,0 +1,61 @@+--------------------------------------------------------------------------------+-- |+-- Module    : PrettyPrint+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Simple layout and pretty printing+--+--------------------------------------------------------------------------------++module PrettyPrint where++import Data.List+import Text.PrettyPrint+import System.IO+++--------------------------------------------------------------------------------+-- Printing+--------------------------------------------------------------------------------++printDoc :: Doc -> IO ()+printDoc =  putStrLn . flip (++) "\n" . render++--+-- stolen from $fptools/ghc/compiler/utils/Pretty.lhs+--+-- This code has a BSD-style license+--+printDocFull :: Mode -> Handle -> Doc -> IO ()+printDocFull m hdl doc = do+  fullRender m cols 1.5 put done doc+  hFlush hdl+  where+    put (Chr c) next  = hPutChar hdl c >> next+    put (Str s) next  = hPutStr  hdl s >> next+    put (PStr s) next = hPutStr  hdl s >> next++    done = hPutChar hdl '\n'+    cols = 80+++--------------------------------------------------------------------------------+-- Layout+--------------------------------------------------------------------------------++--+-- Display the given grid of renderable data, given as either a list of rows or+-- columns, using the minimum size required for each column. An additional+-- parameter specifies extra space to be inserted between each column.+--+ppAsRows      :: Int -> [[Doc]] -> Doc+ppAsRows q    =  ppAsColumns q . transpose++ppAsColumns   :: Int -> [[Doc]] -> Doc+ppAsColumns q =  vcat . map hsep . transpose . map (\col -> pad (width col) col)+  where+    len   = length . render+    width = maximum . map len+    pad w = map (\x -> x <> (hcat $ replicate (w - (len x) + q) space))+
+ examples/common/src/RandomVector.hs view
@@ -0,0 +1,101 @@+{-# LANGUAGE FlexibleContexts, ParallelListComp #-}+--------------------------------------------------------------------------------+--+-- Module    : RandomVector+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Storable multi-dimensional arrays and lists of random numbers+--+--------------------------------------------------------------------------------++module RandomVector+  (+    Storable,+    module RandomVector,+    module Data.Array.Storable+  )+  where++import Foreign                                          (Ptr, Storable)+import Control.Monad                                    (join)+import Control.Exception                                (evaluate)+import Data.Array.Storable+import System.Random+++--------------------------------------------------------------------------------+-- Arrays+--------------------------------------------------------------------------------++type Vector e = StorableArray Int e+type Matrix e = StorableArray (Int,Int) e++withVector :: Vector e -> (Ptr e -> IO a) -> IO a+withVector = withStorableArray++withMatrix :: Matrix e -> (Ptr e -> IO a) -> IO a+withMatrix = withStorableArray+++--+-- To ensure the array is fully evaluated, force one element+--+evaluateArr :: (Ix i, MArray StorableArray e IO)+            => i -> StorableArray i e -> IO (StorableArray i e)+evaluateArr l arr = (join $ evaluate (arr `readArray` l)) >> return arr+++--+-- Generate a new random array+--+randomArrR :: (Ix i, Num e, Storable e, Random e, MArray StorableArray e IO)+           => (i,i) -> (e,e) -> IO (StorableArray i e)+randomArrR (l,u) bnds = do+  rg <- newStdGen+  let -- The standard random number generator is too slow to generate really+      -- large vectors. Instead, we generate a short vector and repeat that.+      k     = 1000+      rands = take k (randomRs bnds rg)++  newListArray (l,u) [rands !! (index (l,u) i`mod`k) | i <- range (l,u)] >>= evaluateArr l+++randomArr :: (Ix i, Num e, Storable e, Random e, MArray StorableArray e IO)+          => (i,i) -> IO (StorableArray i e)+randomArr (l,u) = randomArrR (l,u) (-1,1)+++--+-- Verify similarity of two arrays+--+verify :: (Ix i, Ord e, Fractional e, Storable e)+       => StorableArray i e -> StorableArray i e -> IO (Bool)+verify ref arr = do+  as <- getElems arr+  bs <- getElems ref+  return (verifyList as bs)+++--------------------------------------------------------------------------------+-- Lists+--------------------------------------------------------------------------------++randomListR :: (Num e, Random e, Storable e) => Int -> (e,e) -> IO [e]+randomListR len bnds = do+  rg <- newStdGen+  let -- The standard random number generator is too slow to generate really+      -- large vectors. Instead, we generate a short vector and repeat that.+      k     = 1000+      rands = take k (randomRs bnds rg)++  evaluate [rands !! (i`mod`k) | i <- [0..len-1]]++randomList :: (Num e, Random e, Storable e) => Int -> IO [e]+randomList len = randomListR len (-1,1)+++verifyList :: (Ord e, Fractional e) => [e] -> [e] -> Bool+verifyList xs ys = all (< epsilon) [abs ((x-y)/(x+y+epsilon)) | x <- xs | y <- ys]+  where epsilon = 0.0005+
+ examples/common/src/Time.hs view
@@ -0,0 +1,52 @@+--------------------------------------------------------------------------------+--+-- Module    : Time+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Simple timing benchmarks+--+--------------------------------------------------------------------------------++module Time where++import System.CPUTime+import Control.Monad+++-- Timing+--+data Time = Time { cpu_time :: Integer }++type TimeUnit = Integer -> Integer++picosecond, millisecond, second :: TimeUnit+picosecond  n = n+millisecond n = n `div` 1000000000+second      n = n `div` 1000000000000++getTime :: IO Time+getTime = Time `fmap` getCPUTime++timeIn :: TimeUnit -> Time -> Integer+timeIn u (Time t) = u t++elapsedTime :: Time -> Time -> Time+elapsedTime (Time t1) (Time t2) = Time (t2 - t1)+++-- Simple benchmarking+--+{-# NOINLINE benchmark #-}+benchmark+  :: Int                -- Number of times to repeat test+  -> IO a               -- Test to run+  -> IO b               -- Finaliser to before measuring elapsed time+  -> IO (Time,a)+benchmark n testee finaliser = do+  t1    <- getTime+  (r:_) <- replicateM n testee+  _     <- finaliser+  t2    <- getTime+  return (elapsedTime t1 t2, r)+
+ examples/common/src/cudpp/LICENSE view
@@ -0,0 +1,25 @@+Copyright (c) 2007-2009 The Regents of the University of California, Davis
+campus ("The Regents") and NVIDIA Corporation ("NVIDIA"). All rights reserved.
+
+Redistribution and use in source and binary forms, with or without modification, 
+are permitted provided that the following conditions are met:
+
+    * Redistributions of source code must retain the above copyright notice, 
+      this list of conditions and the following disclaimer.
+    * Redistributions in binary form must reproduce the above copyright notice, 
+      this list of conditions and the following disclaimer in the documentation 
+      and/or other materials provided with the distribution.
+    * Neither the name of the The Regents, nor NVIDIA, nor the names of its 
+      contributors may be used to endorse or promote products derived from this 
+      software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND 
+ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED 
+WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. 
+IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, 
+INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, 
+BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, 
+DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF 
+LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE 
+OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF 
+ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ examples/common/src/cudpp/scan_cta.cu view
@@ -0,0 +1,617 @@+// -------------------------------------------------------------+//  cuDPP -- CUDA Data Parallel Primitives library+//  -------------------------------------------------------------+//  $Revision: 5633 $+//  $Date: 2009-07-01 15:02:51 +1000 (Wed, 01 Jul 2009) $+// -------------------------------------------------------------+// This source code is distributed under the terms of license.txt+// in the root directory of this source distribution.+// -------------------------------------------------------------++/**+ * @file+ * scan_cta.cu+ *+ * @brief CUDPP CTA-level scan routines+ */++/** \defgroup cudpp_cta CUDPP CTA-Level API+  * The CUDPP CTA-Level API contains functions that run on the GPU+  * device.  These are CUDA \c __device__ functions that are called+  * from within other CUDA device functions (typically+  * \link cudpp_kernel CUDPP Kernel-Level API\endlink functions).+  * They are called CTA-level functions because they typically process+  * s_data "owned" by each CTA within shared memory, and are agnostic of+  * any other CTAs that may be running (or how many CTAs are running),+  * other than to compute appropriate global memory addresses.+  * @{+  */++/** @name Scan Functions+* @{+*/++#include "cudpp/cudpp_globals.h"+#include "cudpp/type_vector.h"+// #include <cudpp_util.h>+#include <math.h>+// #include <cudpp.h>++/**+ * @brief Macro to insert necessary __syncthreads() in device emulation mode+ */+#ifdef __DEVICE_EMULATION__+#define __EMUSYNC __syncthreads()+#else+#define __EMUSYNC+#endif++/**+  * @brief Template class containing compile-time parameters to the scan functions+  *+  * ScanTraits is passed as a template parameter to all scan functions.  By+  * using these compile-time functions we can enable generic code while+  * maintaining the highest performance.  This is crucial for the performance+  * of low-level workhorse algorithms like scan.+  *+  * @param T The datatype of the scan+  * @param oper The ::CUDPPOperator to use for the scan (add, max, etc.)+  * @param multiRow True if this is a multi-row scan+  * @param unroll True if scan inner loops should be unrolled+  * @param sums True if each block should write it's sum to the d_blockSums array (false for single-block scans)+  * @param backward True if this is a backward scan+  * @param fullBlock True if all blocks in this scan are full (CTA_SIZE * SCAN_ELEMENTS_PER_THREAD elements)+  * @param exclusive True for exclusive scans, false for inclusive scans+  */+template <class T, class oper, bool backward, bool exclusive,+          bool multiRow, bool sums, bool fullBlock>+class ScanTraits+{+public:++    //! Returns true if this is a backward scan+    static __device__ bool isBackward()    { return backward; };+    //! Returns true if this is an exclusive scan+    static __device__ bool isExclusive()  { return exclusive; };+    //! Returns true if this a multi-row scan.+    static __device__ bool isMultiRow()    { return multiRow; };+    //! Returns true if this scan writes the sum of each block to the d_blockSums array (multi-block scans)+    static __device__ bool writeSums()     { return sums; };+    //! Returns true if this is a full scan -- all blocks process CTA_SIZE * SCAN_ELEMENTS_PER_THREAD elements+    static __device__ bool isFullBlock()   { return fullBlock; };+++    //! The operator function used for the scan+    static __device__ T op(const T &a, const T &b) { return oper::apply(a, b); }++    //! The identity value used by the scan+    static __device__ T identity() { return oper::identity(); }+};++//! This is used to insert syncthreads to avoid perf loss caused by 128-bit+//! load overlap that happens on G80.  This gives about a 15% boost on scans on+//! G80.+//! @todo Parameterize this in case this perf detail changes on future GPUs.+#define DISALLOW_LOADSTORE_OVERLAP 1++/**+* @brief Handles loading input s_data from global memory to shared memory+* (vec4 version)+*+* Load a chunk of 8*blockDim.x elements from global memory into a+* shared memory array.  Each thread loads two T4 elements (where+* T4 is, e.g. int4 or float4), computes the scan of those two vec4s in+* thread local arrays (in registers), and writes the two total sums of the+* vec4s into shared memory, where they will be cooperatively scanned with+* the other partial sums by all threads in the CTA.+*+* @param[out] s_out The output (shared) memory array+* @param[out] threadScan0 Intermediate per-thread partial sums array 1+* @param[out] threadScan1 Intermediate per-thread partial sums array 2+* @param[in] d_in The input (device) memory array+* @param[in] numElements The number of elements in the array being scanned+* @param[in] iDataOffset the offset of the input array in global memory for this+* thread block+* @param[out] ai The shared memory address for the thread's first element+* (returned for reuse)+* @param[out] bi The shared memory address for the thread's second element+* (returned for reuse)+* @param[out] aiDev The device memory address for this thread's first element+* (returned for reuse)+* @param[out] biDev The device memory address for this thread's second element+* (returned for reuse)+*/+template <class T, class traits>+__device__ void loadSharedChunkFromMem4(T        *s_out,+                                        T        threadScan0[4],+                                        T        threadScan1[4],+                                        const T  *d_in,+                                        int      numElements,+                                        int      iDataOffset,+                                        int      &ai,+                                        int      &bi,+                                        int      &aiDev,+                                        int      &biDev)+{+    int thid = threadIdx.x;+    aiDev = iDataOffset + thid;+    biDev = aiDev + blockDim.x;++    // convert to 4-vector+    typename typeToVector<T,4>::Result  tempData;+    typename typeToVector<T,4>::Result* inData = (typename typeToVector<T,4>::Result*)d_in;++    ai = thid;+    bi = thid + blockDim.x;++    // read into tempData;+    if (traits::isBackward())+    {+        int i = aiDev * 4;+        if (traits::isFullBlock() || i + 3 < numElements)+        {+            tempData       = inData[aiDev];+            threadScan0[3] = tempData.w;+            threadScan0[2] = traits::op(tempData.z, threadScan0[3]);+            threadScan0[1] = traits::op(tempData.y, threadScan0[2]);+            threadScan0[0] = s_out[ai]+                           = traits::op(tempData.x, threadScan0[1]);+        }+        else+        {+            threadScan0[3] = traits::identity();+            threadScan0[2] = traits::op(((i+2) < numElements) ? d_in[i+2] : traits::identity(), threadScan0[3]);+            threadScan0[1] = traits::op(((i+1) < numElements) ? d_in[i+1] : traits::identity(), threadScan0[2]);+            threadScan0[0] = s_out[ai]+                           = traits::op((i     < numElements) ? d_in[i]   : traits::identity(), threadScan0[1]);+        }++#ifdef DISALLOW_LOADSTORE_OVERLAP+        __syncthreads();+#endif++        i = biDev * 4;+        if (traits::isFullBlock() || i + 3 < numElements)+        {+            tempData       = inData[biDev];+            threadScan1[3] = tempData.w;+            threadScan1[2] = traits::op(tempData.z, threadScan1[3]);+            threadScan1[1] = traits::op(tempData.y, threadScan1[2]);+            threadScan1[0] = s_out[bi]+                           = traits::op(tempData.x, threadScan1[1]);+        }+        else+        {+            threadScan1[3] = traits::identity();+            threadScan1[2] = traits::op(((i+2) < numElements) ? d_in[i+2] : traits::identity(), threadScan1[3]);+            threadScan1[1] = traits::op(((i+1) < numElements) ? d_in[i+1] : traits::identity(), threadScan1[2]);+            threadScan1[0] = s_out[bi]+                           = traits::op((i     < numElements) ? d_in[i]   : traits::identity(), threadScan1[1]);+        }+        __syncthreads();++        // reverse s_data in shared memory+        if (ai < CTA_SIZE)+        {+            unsigned int leftIdx = ai;+            unsigned int rightIdx = (2 * CTA_SIZE - 1) - ai;++            if (leftIdx < rightIdx)+            {+                T tmp           = s_out[leftIdx];+                s_out[leftIdx]  = s_out[rightIdx];+                s_out[rightIdx] = tmp;+            }+        }+        __syncthreads();+    }+    else+    {+        int i = aiDev * 4;+        if (traits::isFullBlock() || i + 3 < numElements)+        {+            tempData       = inData[aiDev];+            threadScan0[0] = tempData.x;+            threadScan0[1] = traits::op(tempData.y, threadScan0[0]);+            threadScan0[2] = traits::op(tempData.z, threadScan0[1]);+            threadScan0[3] = s_out[ai]+                           = traits::op(tempData.w, threadScan0[2]);+        }+        else+        {+            threadScan0[0] = (i < numElements) ? d_in[i] : traits::identity();+            threadScan0[1] = traits::op(((i+1) < numElements) ? d_in[i+1] : traits::identity(), threadScan0[0]);+            threadScan0[2] = traits::op(((i+2) < numElements) ? d_in[i+2] : traits::identity(), threadScan0[1]);+            threadScan0[3] = s_out[ai]+                           = traits::op(((i+3) < numElements) ? d_in[i+3] : traits::identity(), threadScan0[2]);+        }+++#ifdef DISALLOW_LOADSTORE_OVERLAP+        __syncthreads();+#endif++        i = biDev * 4;+        if (traits::isFullBlock() || i + 3 < numElements)+        {+            tempData       = inData[biDev];+            threadScan1[0] = tempData.x;+            threadScan1[1] = traits::op(tempData.y, threadScan1[0]);+            threadScan1[2] = traits::op(tempData.z, threadScan1[1]);+            threadScan1[3] = s_out[bi]+                           = traits::op(tempData.w, threadScan1[2]);+        }+        else+        {+            threadScan1[0] = (i < numElements) ? d_in[i] : traits::identity();+            threadScan1[1] = traits::op(((i+1) < numElements) ? d_in[i+1] : traits::identity(), threadScan1[0]);+            threadScan1[2] = traits::op(((i+2) < numElements) ? d_in[i+2] : traits::identity(), threadScan1[1]);+            threadScan1[3] = s_out[bi]+                           = traits::op(((i+3) < numElements) ? d_in[i+3] : traits::identity(), threadScan1[2]);+        }+        __syncthreads();+    }+}+++/**+* @brief Handles storing result s_data from shared memory to global memory+* (vec4 version)+*+* Store a chunk of SCAN_ELTS_PER_THREAD*blockDim.x elements from shared memory+* into a device memory array.  Each thread stores reads two elements from shared+* memory, adds them to the intermediate sums computed in+* loadSharedChunkFromMem4(), and writes two T4 elements (where+* T4 is, e.g. int4 or float4) to global memory.+*+* @param[out] d_out The output (device) memory array+* @param[in] threadScan0 Intermediate per-thread partial sums array 1+* (contents computed in loadSharedChunkFromMem4())+* @param[in] threadScan1 Intermediate per-thread partial sums array 2+* (contents computed in loadSharedChunkFromMem4())+* @param[in] s_in The input (shared) memory array+* @param[in] numElements The number of elements in the array being scanned+* @param[in] oDataOffset the offset of the output array in global memory+* for this thread block+* @param[in] ai The shared memory address for the thread's first element+* (computed in loadSharedChunkFromMem4())+* @param[in] bi The shared memory address for the thread's second element+* (computed in loadSharedChunkFromMem4())+* @param[in] aiDev The device memory address for this thread's first element+* (computed in loadSharedChunkFromMem4())+* @param[in] biDev The device memory address for this thread's second element+* (computed in loadSharedChunkFromMem4())+*/+template <class T, class traits>+__device__ void storeSharedChunkToMem4(T   *d_out,+                                       T   threadScan0[4],+                                       T   threadScan1[4],+                                       T   *s_in,+                                       int numElements,+                                       int oDataOffset,+                                       int ai,+                                       int bi,+                                       int aiDev,+                                       int biDev)+{+    // Convert to 4-vector+    typename typeToVector<T,4>::Result tempData;+    typename typeToVector<T,4>::Result* outData = (typename typeToVector<T,4>::Result*)d_out;++    // write results to global memory+    if (traits::isBackward())+    {+        if (ai < CTA_SIZE)+        {++            unsigned int leftIdx = ai;+            unsigned int rightIdx = (2 * CTA_SIZE - 1) - ai;++            if (leftIdx < rightIdx)+            {+                T tmp = s_in[leftIdx];+                s_in[leftIdx] = s_in[rightIdx];+                s_in[rightIdx] = tmp;+            }+        }+        __syncthreads();++        T temp = s_in[ai];++        if (traits::isExclusive())+        {+            tempData.w = temp;+            tempData.z = traits::op(temp, threadScan0[3]);+            tempData.y = traits::op(temp, threadScan0[2]);+            tempData.x = traits::op(temp, threadScan0[1]);+        }+        else+        {+            tempData.w = traits::op(temp, threadScan0[3]);+            tempData.z = traits::op(temp, threadScan0[2]);+            tempData.y = traits::op(temp, threadScan0[1]);+            tempData.x = traits::op(temp, threadScan0[0]);+        }++        int i = aiDev * 4;+        if (traits::isFullBlock() || i + 3 < numElements)+        {+            outData[aiDev] = tempData;+        }+        else+        {+            if (i   < numElements) { d_out[i]   = tempData.x;+            if (i+1 < numElements) { d_out[i+1] = tempData.y;+            if (i+2 < numElements) { d_out[i+2] = tempData.z; }}}+        }++#ifdef DISALLOW_LOADSTORE_OVERLAP+        __syncthreads();+#endif++        temp = s_in[bi];++        if (traits::isExclusive())+        {+            tempData.w = temp;+            tempData.z = traits::op(temp, threadScan1[3]);+            tempData.y = traits::op(temp, threadScan1[2]);+            tempData.x = traits::op(temp, threadScan1[1]);+        }+        else+        {+            tempData.w = traits::op(temp, threadScan1[3]);+            tempData.z = traits::op(temp, threadScan1[2]);+            tempData.y = traits::op(temp, threadScan1[1]);+            tempData.x = traits::op(temp, threadScan1[0]);+        }++        i = biDev * 4;+        if (traits::isFullBlock() || i + 3 < numElements)+        {+            outData[biDev] = tempData;+        }+        else+        {+            if (i   < numElements) { d_out[i]   = tempData.x;+            if (i+1 < numElements) { d_out[i+1] = tempData.y;+            if (i+2 < numElements) { d_out[i+2] = tempData.z; }}}+        }+    }+    else+    {+        T temp;+        temp = s_in[ai];++        if (traits::isExclusive())+        {+            tempData.x = temp;+            tempData.y = traits::op(temp, threadScan0[0]);+            tempData.z = traits::op(temp, threadScan0[1]);+            tempData.w = traits::op(temp, threadScan0[2]);+        }+        else+        {+            tempData.x = traits::op(temp, threadScan0[0]);+            tempData.y = traits::op(temp, threadScan0[1]);+            tempData.z = traits::op(temp, threadScan0[2]);+            tempData.w = traits::op(temp, threadScan0[3]);+        }++        int i = aiDev * 4;+        if (traits::isFullBlock() || i + 3 < numElements)+        {+            outData[aiDev] = tempData;+        }+        else+        {+            // we can't use vec4 because the original array isn't a multiple of+            // 4 elements+            if ( i    < numElements) { d_out[i]   = tempData.x;+            if ((i+1) < numElements) { d_out[i+1] = tempData.y;+            if ((i+2) < numElements) { d_out[i+2] = tempData.z; } } }+        }++#ifdef DISALLOW_LOADSTORE_OVERLAP+        __syncthreads();+#endif++        temp       = s_in[bi];++        if (traits::isExclusive())+        {+            tempData.x = temp;+            tempData.y = traits::op(temp, threadScan1[0]);+            tempData.z = traits::op(temp, threadScan1[1]);+            tempData.w = traits::op(temp, threadScan1[2]);+        }+        else+        {+            tempData.x = traits::op(temp, threadScan1[0]);+            tempData.y = traits::op(temp, threadScan1[1]);+            tempData.z = traits::op(temp, threadScan1[2]);+            tempData.w = traits::op(temp, threadScan1[3]);+        }++        i = biDev * 4;+        if (traits::isFullBlock() || i + 3 < numElements)+        {+            outData[biDev] = tempData;+        }+        else+        {+            // we can't use vec4 because the original array isn't a multiple of+            // 4 elements+            if ( i    < numElements) { d_out[i]   = tempData.x;+            if ((i+1) < numElements) { d_out[i+1] = tempData.y;+            if ((i+2) < numElements) { d_out[i+2] = tempData.z; } } }+        }+    }+}++/** @brief Scan all warps of a CTA without synchronization+  *+  * The warp-scan algorithm breaks a block of data into warp-sized chunks, and+  * scans the chunks independently with a warp of threads each.  Because warps+  * execute instructions in SIMD fashion, there is no need to synchronize in+  * order to share data within a warp (only across warps).  Also, in SIMD the+  * most efficient algorithm is a step-efficient algorithm.  Therefore, within+  * each warp we use a Hillis-and-Steele-style scan that takes log2(N) steps+  * to scan the warp [Daniel Hillis and Guy Steele 1986], rather than the+  * work-efficient tree-based algorithm described by Guy Blelloch [1990] that+  * takes 2 * log(N) steps and is in general more complex to implement.+  * Previous versions of CUDPP used the Blelloch algorithm.  For current GPUs,+  * the warp size is 32, so this takes five steps per warp.+  *+  * Each thread is responsible for a single element of the array to be scanned.+  * Each thread inputs a single value to the scan via \a val and returns+  * its own scanned result element.  The threads of each warp cooperate+  * via the shared memory array \a s_data to scan WARP_SIZE elements.+  *+  * Template parameter \a maxlevel allows this warpscan to be performed on+  * partial warps.  For example, if only the first 8 elements of each warp need+  * to be scanned, then warpscan only performs log2(8)=3 steps rather than 5.+  *+  * The computation uses 2 * WARP_SIZE elements of shared memory per warp to+  * enable warps to offset beyond their input data and receive the identity+  * element without using any branch instructions.+  *+  * \note s_data is declared volatile here to prevent the compiler from+  * optimizing away writes to shared memory, and ensure correct intrawarp+  * communication in the absence of __syncthreads.+  *+  * @return The result of the warp scan for the current thread+  * @param[in] val The current threads's input to the scan+  * @param[in,out] s_data A pointer to a temporary shared array of 2*CTA_SIZE+  * elements used to compute the warp scans+  */+template<class T, class traits,int maxlevel>+__device__ T warpscan(T val, volatile T* s_data)+{+    // The following is the same as 2 * 32 * warpId + threadInWarp =+    // 64*(threadIdx.x >> 5) + (threadIdx.x & (WARP_SIZE-1))+    int idx = 2 * threadIdx.x - (threadIdx.x & (WARP_SIZE-1));+    s_data[idx] = traits::identity();+    idx += WARP_SIZE;+    s_data[idx] = val;                                 __EMUSYNC;++        // This code is needed because the warp size of device emulation+        // is only 1 thread, so sync-less cooperation within a warp doesn't+        // work.+#ifdef __DEVICE_EMULATION__+    T t = s_data[idx -  1]; __EMUSYNC;+    s_data[idx] = traits::op((const T&)s_data[idx],t); __EMUSYNC;+    t = s_data[idx -  2]; __EMUSYNC;+    s_data[idx] = traits::op((const T&)s_data[idx],t); __EMUSYNC;+    t = s_data[idx -  4]; __EMUSYNC;+    s_data[idx] = traits::op((const T&)s_data[idx],t); __EMUSYNC;+    t = s_data[idx -  8]; __EMUSYNC;+    s_data[idx] = traits::op((const T&)s_data[idx],t); __EMUSYNC;+    t = s_data[idx - 16]; __EMUSYNC;+    s_data[idx] = traits::op((const T&)s_data[idx],t); __EMUSYNC;+#else+    if (0 <= maxlevel) { s_data[idx] = traits::op((const T&)s_data[idx], (const T&)s_data[idx - 1]); }+    if (1 <= maxlevel) { s_data[idx] = traits::op((const T&)s_data[idx], (const T&)s_data[idx - 2]); }+    if (2 <= maxlevel) { s_data[idx] = traits::op((const T&)s_data[idx], (const T&)s_data[idx - 4]); }+    if (3 <= maxlevel) { s_data[idx] = traits::op((const T&)s_data[idx], (const T&)s_data[idx - 8]); }+    if (4 <= maxlevel) { s_data[idx] = traits::op((const T&)s_data[idx], (const T&)s_data[idx -16]); }+#endif++    return s_data[idx-1];      // convert inclusive -> exclusive+}++/** @brief Perform a full CTA scan using the warp-scan algorithm+  *+  * As described in the comment for warpscan(), the warp-scan algorithm breaks+  * a block of data into warp-sized chunks, and scans the chunks independently+  * with a warp of threads each.  To complete the scan, each warp <i>j</i> then+  * writes its last element to element <i>j</i> of a temporary shared array.+  * Then a single warp exclusive-scans these "warp sums".  Finally, each thread+  * adds the result of the warp sum scan to the result of the scan from the+  * first pass.+  *+  * Because we scan 2*CTA_SIZE elements per thread, we have to call warpscan+  * twice.+  *+  * @param x The first input value for the current thread+  * @param y The second input value for the current thread+  * @param s_data Temporary shared memory space of 2*CTA_SIZE elements for+  * performing the scan+  */+template <class T, class traits>+__device__ void scanWarps(T x, T y,+                          T *s_data)+{+    T val  = warpscan<T, traits, 4>(x, s_data);+    __syncthreads();+    T val2 = warpscan<T, traits, 4>(y, s_data);++    int idx = threadIdx.x;++    if ((idx & 31)==31)+    {+        s_data[idx >> 5]                = traits::op(val, x);+        s_data[(idx + blockDim.x) >> 5] = traits::op(val2, y);+    }+    __syncthreads();++#ifndef __DEVICE_EMULATION__+    if (idx < 32)+#endif+    {+        s_data[idx] = warpscan<T,traits,(LOG_CTA_SIZE-LOG_WARP_SIZE+1)>(s_data[idx], s_data);+    }+    __syncthreads();++    val  = traits::op(val, s_data[idx >> 5]);++    val2 = traits::op(val2, s_data[(idx + blockDim.x) >> 5]);++    __syncthreads();++    s_data[idx] = val;+    s_data[idx+blockDim.x] = val2;+}++/**+* @brief CTA-level scan routine; scans s_data in shared memory in each thread block+*+* This function is the main CTA-level scan function.  It may be called by other+* CUDA __global__ or __device__ functions. This function scans 2 * CTA_SIZE elements.+* Each thread is responsible for one element in each half of the input array.+* \note This code is intended to be run on a CTA of 128 threads.  Other sizes are+* untested.+*+* @param[in] s_data The array to be scanned in shared memory+* @param[out] d_blockSums Array of per-block sums+* @param[in] blockSumIndex Location in \a d_blockSums to which to write this block's sum+*/+template <class T, class traits>+__device__ void scanCTA(T            *s_data,+                        T            *d_blockSums,+                        unsigned int blockSumIndex)+{+    T val  = s_data[threadIdx.x];+    T val2 = s_data[threadIdx.x + blockDim.x];+    __syncthreads();++    scanWarps<T,traits>(val, val2, s_data);+    __syncthreads();++    if (traits::writeSums() && threadIdx.x == blockDim.x - 1)+    {+        d_blockSums[blockSumIndex] = traits::op(val2, s_data[threadIdx.x + blockDim.x]);+    }+++#ifdef __DEVICE_EMULATION__+    // must sync in emulation mode when doing backward scans, because otherwise the+    // shared memory array will get reversed before the block sums are read!+    if (traits::isBackward())+        __syncthreads();+#endif+}+++/** @} */ // end scan functions+/** @} */ // end cudpp_cta
+ examples/common/src/cudpp/scan_kernel.cu view
@@ -0,0 +1,113 @@+// -------------------------------------------------------------+// cuDPP -- CUDA Data Parallel Primitives library+// -------------------------------------------------------------+//  $Revision: 5633 $+//  $Date: 2009-07-01 15:02:51 +1000 (Wed, 01 Jul 2009) $+// -------------------------------------------------------------+// This source code is distributed under the terms of license.txt+// in the root directory of this source distribution.+// -------------------------------------------------------------++/**+ * @file+ * scan_kernel.cu+ *+ * @brief CUDPP kernel-level scan routines+ */++/** \defgroup cudpp_kernel CUDPP Kernel-Level API+  * The CUDPP Kernel-Level API contains functions that run on the GPU+  * device across a grid of Cooperative Thread Array (CTA, aka Thread+  * Block).  These kernels are declared \c __global__ so that they+  * must be invoked from host (CPU) code.  They generally invoke GPU+  * \c __device__ routines in the CUDPP \link cudpp_cta CTA-Level API\endlink.+  * Kernel-Level API functions are used by CUDPP+  * \link cudpp_app Application-Level\endlink functions to implement their+  * functionality.+  * @{+  */++/** @name Scan Functions+* @{+*/++#include "cudpp/cudpp_globals.h"+#include "cudpp/scan_cta.cu"+#include "cudpp/shared_mem.h"++/**+  * @brief Main scan kernel+  *+  * This __global__ device function performs one level of a multiblock scan on+  * an arbitrary-dimensioned array in \a d_in, returning the result in \a d_out+  * (which may point to the same array).  The same function may be used for+  * single or multi-row scans.  To perform a multirow scan, pass the width of+  * each row of the input row (in elements) in \a dataRowPitch, and the width of+  * the rows of \a d_blockSums (in elements) in \a blockSumRowPitch, and invoke+  * with a thread block grid with height greater than 1.+  *+  * This function peforms one level of a recursive, multiblock scan.  At the+  * app level, this function is called by cudppScan and cudppMultiScan and used+  * in combination with vectorAddUniform4() to produce a complete scan.+  *+  * Template parameter \a T is the datatype of the array to be scanned.+  * Template parameter \a traits is the ScanTraits struct containing+  * compile-time options for the scan, such as whether it is forward or+  * backward, exclusive or inclusive, multi- or single-row, etc.+  *+  * @param[out] d_out The output (scanned) array+  * @param[in]  d_in The input array to be scanned+  * @param[out] d_blockSums The array of per-block sums+  * @param[in]  numElements The number of elements to scan+  * @param[in]  dataRowPitch The width of each row of \a d_in in elements+  * (for multi-row scans)+  * @param[in]  blockSumRowPitch The with of each row of \a d_blockSums in elements+  * (for multi-row scans)+  */+template<class T, class traits>+__global__ void scan4(T            *d_out,+                      const T      *d_in,+                      T            *d_blockSums,+                      int          numElements,+                      unsigned int dataRowPitch,+                      unsigned int blockSumRowPitch)+{+    SharedMemory<T> smem;+    T* temp = smem.getPointer();++    int devOffset, ai, bi, aiDev, biDev;+    T threadScan0[4], threadScan1[4];++    unsigned int blockN = numElements;+    unsigned int blockSumIndex = blockIdx.x;++    if (traits::isMultiRow())+    {+        //int width = __mul24(gridDim.x, blockDim.x) << 1;+        int yIndex     = __umul24(blockDim.y, blockIdx.y) + threadIdx.y;+        devOffset      = __umul24(dataRowPitch, yIndex);+        blockN        += (devOffset << 2);+        devOffset     += __umul24(blockIdx.x, blockDim.x << 1);+        blockSumIndex += __umul24(blockSumRowPitch << 2, yIndex) ;+    }+    else+    {+        devOffset = __umul24(blockIdx.x, (blockDim.x << 1));+    }++    // load data into shared memory+    loadSharedChunkFromMem4<T, traits>+        (temp, threadScan0, threadScan1, d_in,+         blockN, devOffset, ai, bi, aiDev, biDev);++    scanCTA<T, traits>(temp, d_blockSums, blockSumIndex);++    // write results to device memory+    storeSharedChunkToMem4<T, traits>+        (d_out, threadScan0, threadScan1, temp,+         blockN, devOffset, ai, bi, aiDev, biDev);++}++/** @} */ // end scan functions+/** @} */ // end cudpp_kernel
+ examples/common/src/cudpp/vector_kernel.cu view
@@ -0,0 +1,444 @@+// -------------------------------------------------------------+// CUDPP -- CUDA Data Parallel Primitives library+// -------------------------------------------------------------+//  $Revision: 5632 $+//  $Date: 2009-07-01 14:36:01 +1000 (Wed, 01 Jul 2009) $+// -------------------------------------------------------------+// This source code is distributed under the terms of license.txt in+// the root directory of this source distribution.+// -------------------------------------------------------------++/**+ * @file+ * vector_kernel.cu+ *+ * @brief CUDA kernel methods for basic operations on vectors.+ *+ * CUDA kernel methods for basic operations on vectors.+ *+ * Examples:+ * - vectorAddConstant(): d_vector + constant+ * - vectorAddUniform():  d_vector + uniform (per-block constants)+ * - vectorAddVectorVector(): d_vector + d_vector+ */++// MJH: these functions assume there are 2N elements for N threads.+// Is this always going to be a good idea?  There may be cases where+// we have as many threads as elements, but for large problems+// we are probably limited by max CTA size for simple kernels like+// this so we should process multiple elements per thread.+// we may want to extend these with looping versions that process+// many elements per thread.++// #include "cudpp_util.h"+// #include "sharedmem.h"+// #include "cudpp.h"++/** \addtogroup cudpp_kernel+  * @{+  */++/** @name Vector Functions+ * CUDA kernel methods for basic operations on vectors.+ * @{+ */++#if 0+/** @brief Adds a constant value to all values in the input d_vector+ *+ * Each thread adds two pairs of elements.+ * @todo Test this function -- it is currently not yet used.+ *+ * @param[in,out] d_vector The array of elements to be modified+ * @param[in] constant The constant value to be added to elements of+ * \a d_vector+ * @param[in] n The number of elements in the d_vector to be modified+ * @param[in] baseIndex An optional offset to the beginning of the+ * elements in the input array to be processed+ */+template <class T>+__global__  void vectorAddConstant(T   *d_vector,+                                   T   constant,+                                   int n,+                                   int baseIndex)+{+    // Compute this thread's output address+    unsigned int address = baseIndex + threadIdx.x ++        __mul24(blockIdx.x, (blockDim.x << 1));++    // note two adds per thread: one in first half of the block, one in last+    d_vector[address]              += constant;+    d_vector[address + blockDim.x] += (threadIdx.x + blockDim.x < n) * constant;+}+#endif+#if 0+ /** @brief Add a uniform value to each data element of an array+  *+  * This function reads one value per CTA from \a d_uniforms into shared+  * memory and adds that value to all values "owned" by the CTA in \a+  * d_vector.  Each thread adds two pairs of values.+  *+  * @param[out] d_vector The d_vector whose values will have the uniform added+  * @param[in] d_uniforms The array of uniform values (one per CTA)+  * @param[in] numElements The number of elements in \a d_vector to process+  * @param[in] blockOffset an optional offset to the beginning of this block's+  * data.+  * @param[in] baseIndex an optional offset to the beginning of the array+  * within \a d_vector.+  */+template <class T>+__global__ void vectorAddUniform(T       *d_vector,+                                 const T *d_uniforms,+                                 int     numElements,+                                 int     blockOffset,+                                 int     baseIndex)+{+    __shared__ T uni;+    // Get this block's uniform value from the uniform array in device memory+    // We store it in shared memory so that the hardware's shared memory+    // broadcast capability can be used to share among all threads in each warp+    // in a single cycle+    if (threadIdx.x == 0)+    {+        uni = d_uniforms[blockIdx.x + __mul24(gridDim.x, blockIdx.y) + blockOffset];+    }++    // Compute this thread's output address+    int width = __mul24(gridDim.x,(blockDim.x << 1));++    unsigned int address = baseIndex + __mul24(width, blockIdx.y)+        + threadIdx.x + __mul24(blockIdx.x, (blockDim.x << 1));++    __syncthreads();++    // note two adds per thread: one in first half of the block, one in last+    d_vector[address]              += uni;+    if (threadIdx.x + blockDim.x < numElements) d_vector[address + blockDim.x] += uni;+}+#endif++/** @brief Add a uniform value to each data element of an array (vec4 version)+  *+  * This function reads one value per CTA from \a d_uniforms into shared+  * memory and adds that value to all values "owned" by the CTA in \a d_vector.+  * Each thread adds the uniform value to eight values in \a d_vector.+  *+  * @param[out] d_vector The d_vector whose values will have the uniform added+  * @param[in] d_uniforms The array of uniform values (one per CTA)+  * @param[in] numElements The number of elements in \a d_vector to process+  * @param[in] vectorRowPitch For 2D arrays, the pitch (in elements) of the+  * rows of \a d_vector.+  * @param[in] uniformRowPitch For 2D arrays, the pitch (in elements) of the+  * rows of \a d_uniforms.+  * @param[in] blockOffset an optional offset to the beginning of this block's+  * data.+  * @param[in] baseIndex an optional offset to the beginning of the array+  * within \a d_vector.+  */+template <class T, class op, int elementsPerThread>+__global__ void vectorAddUniform4(T       *d_vector,+                                  const T *d_uniforms,+                                  int      numElements,+                                  int      vectorRowPitch,     // width of input array in elements+                                  int      uniformRowPitch,    // width of uniform array in elements+                                  int      blockOffset,+                                  int      baseIndex)+{+    __shared__ T uni;+    // Get this block's uniform value from the uniform array in device memory+    // We store it in shared memory so that the hardware's shared memory+    // broadcast capability can be used to share among all threads in each warp+    // in a single cycle+    if (threadIdx.x == 0)+    {+        uni = d_uniforms[blockIdx.x + __umul24(uniformRowPitch, blockIdx.y) + blockOffset];+    }++    // Compute this thread's output address+    //int width = __mul24(gridDim.x,(blockDim.x << 1));++    unsigned int address = baseIndex + __umul24(vectorRowPitch, blockIdx.y)+        + threadIdx.x + __umul24(blockIdx.x, (blockDim.x * elementsPerThread));+    numElements += __umul24(vectorRowPitch, blockIdx.y);++    __syncthreads();++    for (int i = 0; i < elementsPerThread && address < numElements; i++)+    {+        d_vector[address] = op::apply(d_vector[address], uni);+        address += blockDim.x;+    }+}++#if 0+/** @brief Adds together two vectors+ *+ * Each thread adds two pairs of elements.+ * @todo Test this function -- it is currently not yet used.+ *+ * @param[out] d_vectorA The left operand array and the result+ * @param[in] d_vectorB The right operand array+ * @param[in] numElements The number of elements in the vectors to be added.+ * @param[in] baseIndex An optional offset to the beginning of the+ * elements in the input arrays to be processed+ */+template <class T>+__global__ void vectorAddVector(T       *d_vectorA,        // A += B+                                const T *d_vectorB,+                                int     numElements,+                                int     baseIndex)+{+    // Compute this thread's output address+    unsigned int address = baseIndex + threadIdx.x ++        __mul24(blockIdx.x, (blockDim.x << 1));++    // note two adds per thread: one in first half of the block, one in last+    d_vectorA[address]              += d_vectorB[address];+    d_vectorA[address + blockDim.x] +=+        (threadIdx.x + blockDim.x < numElements) * d_vectorB[address];+}+#endif++/** @brief Add a uniform value to data elements of an array (vec4 version)+  *+  * This function reads one value per CTA from \a d_uniforms into shared+  * memory and adds that value to values "owned" by the CTA in \a d_vector.+  * The uniform value is added to only those values "owned" by the CTA which+  * have an index less than d_maxIndex. If d_maxIndex for that CTA is UINT_MAX+  * it adds the uniform to all values "owned" by the CTA.+  * Each thread adds the uniform value to eight values in \a d_vector.+  *+  * @param[out] d_vector The d_vector whose values will have the uniform added+  * @param[in] d_uniforms The array of uniform values (one per CTA)+  * @param[in] d_maxIndices The array of maximum indices (one per CTA). This is+  *            index upto which the uniform would be added. If this is UINT_MAX+  *            the uniform is added to all elements of the CTA. This index is+  *            1-based.+  * @param[in] numElements The number of elements in \a d_vector to process+  * @param[in] blockOffset an optional offset to the beginning of this block's+  * data.+  * @param[in] baseIndex an optional offset to the beginning of the array+  * within \a d_vector.+  */+template <class T, class op, bool isLastBlockFull>+__global__ void vectorSegmentedAddUniform4(T                  *d_vector,+                                           const T            *d_uniforms,+                                           const unsigned int *d_maxIndices,+                                           unsigned int       numElements,+                                           int                blockOffset,+                                           int                baseIndex)+{+    __shared__ T uni[2];++    unsigned int blockAddress =+        blockIdx.x + __mul24(gridDim.x, blockIdx.y) + blockOffset;++    // Get this block's uniform value from the uniform array in device memory+    // We store it in shared memory so that the hardware's shared memory+    // broadcast capability can be used to share among all threads in each warp+    // in a single cycle++    if (threadIdx.x == 0)+    {+        if (blockAddress > 0)+            uni[0] = d_uniforms[blockAddress-1];+        else+            uni[0] = op::identity();++        // Tacit assumption that T is four-byte wide+        uni[1] = (T)(d_maxIndices[blockAddress]);+    }++    // Compute this thread's output address+    int width = __mul24(gridDim.x,(blockDim.x << 1));++    unsigned int address = baseIndex + __mul24(width, blockIdx.y)+                           + threadIdx.x + __mul24(blockIdx.x, (blockDim.x << 3));++    __syncthreads();++    unsigned int maxIndex = (unsigned int)(uni[1]);++    bool isLastBlock = (blockIdx.x == (gridDim.x-1));++    if (maxIndex < UINT_MAX)+    {+        // Since maxIndex is a 1 based index+        --maxIndex;+        bool leftLess = address < maxIndex;+        bool rightLess = (address + 7 * blockDim.x) < maxIndex;++        if (leftLess)+        {+            if (rightLess)+            {+                for (unsigned int i = 0; i < 8; ++i)+                    d_vector[address + i * blockDim.x] =+                        op::apply(d_vector[address + i * blockDim.x], uni[0]);+            }+            else+            {+                for (unsigned int i=0; i < 8; ++i)+                {+                    if (address < maxIndex)+                        d_vector[address] =+                            op::apply(d_vector[address], uni[0]);++                    address += blockDim.x;+                }+            }+        }+    }+    else+    {+        if (!isLastBlockFull && isLastBlock)+        {+            for (unsigned int i = 0; i < 8; ++i)+            {+                if (address < numElements)+                    d_vector[address] =+                        op::apply(d_vector[address], uni[0]);++                address += blockDim.x;+            }+        }+        else+        {+            for (unsigned int i=0; i<8; ++i)+            {+                d_vector[address] =+                    op::apply(d_vector[address], uni[0]);++                address += blockDim.x;+            }+        }+    }+}+++/** @brief Add a uniform value to data elements of an array (vec4 version)+  *+  * This function reads one value per CTA from \a d_uniforms into shared+  * memory and adds that value to values "owned" by the CTA in \a d_vector.+  * The uniform value is added to only those values "owned" by the CTA which+  * have an index greater than d_minIndex. If d_minIndex for that CTA is 0+  * it adds the uniform to all values "owned" by the CTA.+  * Each thread adds the uniform value to eight values in \a d_vector.+  *+  * @param[out] d_vector The d_vector whose values will have the uniform added+  * @param[in] d_uniforms The array of uniform values (one per CTA)+  * @param[in] d_minIndices The array of minimum indices (one per CTA). The+  *            uniform is added to the right of this index (that is, to every index+  *            that is greater than this index). If this is 0, the uniform is+  *            added to all elements of the CTA. This index is 1-based to+  *            prevent overloading of what 0 means. In our case it means+  *            absence of a flag. But if the first element of a CTA has+  *            flag the index will also be 0. Hence we use 1-based indices+  *            so the index is 1 in the latter case.+  * @param[in] numElements The number of elements in \a d_vector to process+  * @param[in] blockOffset an optional offset to the beginning of this block's+  * data.+  * @param[in] baseIndex an optional offset to the beginning of the array+  * within \a d_vector.+  *+  */+template <class T, class op, bool isLastBlockFull>+__global__ void vectorSegmentedAddUniformToRight4(T                  *d_vector,+                                                  const T            *d_uniforms,+                                                  const unsigned int *d_minIndices,+                                                  unsigned int       numElements,+                                                  int                blockOffset,+                                                  int                baseIndex)+{+    __shared__ T uni[2];++    unsigned int blockAddress =+        blockIdx.x + __mul24(gridDim.x, blockIdx.y) + blockOffset;++    // Get this block's uniform value from the uniform array in device memory+    // We store it in shared memory so that the hardware's shared memory+    // broadcast capability can be used to share among all threads in each warp+    // in a single cycle++    if (threadIdx.x == 0)+    {+        // FIXME - blockAddress test here is incompatible with how it is calculated+        // above+        if (blockAddress < (gridDim.x-1))+            uni[0] = d_uniforms[blockAddress+1];+        else+            uni[0] = op::identity();++        // Tacit assumption that T is four-byte wide+        uni[1] = (T)(d_minIndices[blockAddress]);+    }++    // Compute this thread's output address+    int width = __mul24(gridDim.x,(blockDim.x << 1));++    unsigned int address = baseIndex + __mul24(width, blockIdx.y)+                           + threadIdx.x + __mul24(blockIdx.x, (blockDim.x << 3));++    __syncthreads();++    unsigned int minIndex = (unsigned int)(uni[1]);++    bool isLastBlock = (blockIdx.x == (gridDim.x-1));++    if (minIndex > 0)+    {+        // Since minIndex is a 1 based index+        --minIndex;+        bool leftInRange = address > minIndex;+        bool rightInRange = (address + 7 * blockDim.x) > minIndex;++        if (rightInRange)+        {+            if (leftInRange)+            {+                for (unsigned int i = 0; i < 8; ++i)+                    d_vector[address + i * blockDim.x] =+                        op::apply(d_vector[address + i * blockDim.x], uni[0]);+            }+            else+            {+                for (unsigned int i=0; i < 8; ++i)+                {+                    if (address > minIndex)+                        d_vector[address] =+                            op::apply(d_vector[address], uni[0]);++                    address += blockDim.x;+                }+            }+        }+    }+    else+    {+        if (!isLastBlockFull && isLastBlock)+        {+            for (unsigned int i = 0; i < 8; ++i)+            {+                if (address < numElements)+                    d_vector[address] =+                        op::apply(d_vector[address], uni[0]);++                address += blockDim.x;+            }+        }+        else+        {+            for (unsigned int i=0; i<8; ++i)+            {+                d_vector[address] =+                    op::apply(d_vector[address], uni[0]);++                address += blockDim.x;+            }+        }+    }+}+++/** @} */ // end d_vector functions+/** @} */ // end cudpp_kernel
+ examples/src/bandwidthTest/BandwidthTest.hs view
@@ -0,0 +1,221 @@+--------------------------------------------------------------------------------+--+-- Module    : Fold+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Reduce a vector to a single value+--+--------------------------------------------------------------------------------++module Main where++-- Friends+import PrettyPrint++-- System+import Numeric+import Data.List+import Control.Monad+import Control.Exception+import System.Exit+import System.Environment+import System.Console.GetOpt+import Text.PrettyPrint++import Foreign+import qualified Foreign.CUDA               as CUDA+import qualified Foreign.CUDA.Runtime.Event as CUDA+++--------------------------------------------------------------------------------+-- Options+--------------------------------------------------------------------------------++data TestMode = Quick | Range | Shmoo+  deriving (Eq,Show,Read)++data MemoryMode = List | Pageable | Pinned | WriteCombined+  deriving (Eq,Show,Read,Ord,Enum)++data CopyMode = HostToDevice | DeviceToHost | DeviceToDevice+  deriving (Eq,Show,Read,Ord,Enum)++data Options = Options+  { testMode   :: TestMode+--  , memoryMode :: [MemoryMode]+--  , copyMode   :: [CopyMode]+  , device     :: Int+  , range      :: (Int,Int,Int)+  }+  deriving (Show)++defaultOptions :: Options+defaultOptions = Options+  { device     = 0+  , testMode   = Quick+--  , memoryMode = []+--  , copyMode   = []+  , range      = (kilobyte, kilobyte, 10*kilobyte)+  }++kilobyte, megabyte :: Int+kilobyte = 1 `shift` 10+megabyte = 1 `shift` 20++options :: [OptDescr (Options -> Options)]+options =+  [ Option ['d'] ["device"]    (ReqArg (\d opts -> opts { device = read d }) "ID")                           "numeral of device to test"+  , Option ['t'] ["test"]      (ReqArg (\t opts -> opts { testMode = read t}) "MODE")                        "testing mode: Quick | Range | Shmoo"+--  , Option ['m'] ["memory"]    (ReqArg (\m opts -> opts { memoryMode = read m : memoryMode opts}) "MODE")    "memory copy mode: List | Pageable | Pinned | WriteCombined"+--  , Option ['c'] ["copy"]      (ReqArg (\c opts -> opts { copyMode = read c : copyMode opts}) "DIRECTION")   "memory copy direction: DeviceToHost | HostToDevice | DeviceToDevice"+  , Option ['s'] ["start"]     (ReqArg (\s opts -> opts { range = let (_,i,e) = range opts in (read s,i,e)}) "BYTES") "starting transfer size"+  , Option ['i'] ["increment"] (ReqArg (\i opts -> opts { range = let (s,_,e) = range opts in (s,read i,e)}) "BYTES") "transfer test size increment"+  , Option ['e'] ["end"]       (ReqArg (\e opts -> opts { range = let (s,i,_) = range opts in (s,i,read e)}) "BYTES") "ending transfer size"+  ]+++--------------------------------------------------------------------------------+-- Testing+--------------------------------------------------------------------------------++-- Benchmarking+--+bench :: Int -> IO a -> IO Double+bench n testee =+  let iter = 10+      size = fromIntegral (iter * n) * fromIntegral (sizeOf (undefined::Int))+  in+  bracket (CUDA.create []) CUDA.destroy $ \start ->+  bracket (CUDA.create []) CUDA.destroy $ \stop  -> do+    CUDA.record start Nothing+    replicateM_ iter testee+    CUDA.record stop  Nothing+    CUDA.sync+    ms <- realToFrac `fmap` CUDA.elapsedTime start stop+    return $ 1E3 * size / (fromIntegral megabyte * ms)+++-- Bandwidth testing for the various copy modes+--+bandwidth :: CopyMode -> MemoryMode -> Int -> IO Double++bandwidth HostToDevice List     n =+  bench n (CUDA.withListArray [1..n] (\_ -> return ()))++bandwidth HostToDevice Pageable n =+  CUDA.allocaArray n $ \d_ptr ->+  withArray [1..n]   $ \h_ptr ->+  bench n (CUDA.pokeArray n h_ptr d_ptr)++bandwidth HostToDevice x n        =+  CUDA.allocaArray n $ \d_ptr ->+  let f = if x == WriteCombined then [CUDA.WriteCombined] else [] in+  bracket (CUDA.mallocHostArray f n) (CUDA.freeHost) $ \h_ptr -> do+  pokeArray (CUDA.useHostPtr h_ptr) [1..n]+  bench n (CUDA.pokeArrayAsync n h_ptr d_ptr Nothing)++bandwidth DeviceToHost List     n =+  CUDA.withListArray [1..n] $ \d_ptr ->+  bench n (CUDA.peekListArray n d_ptr)++bandwidth DeviceToHost Pageable n =+  allocaArray n             $ \h_ptr ->+  CUDA.withListArray [1..n] $ \d_ptr ->+  bench n (CUDA.peekArray n d_ptr h_ptr)++bandwidth DeviceToHost x n        =+  let f = if x == WriteCombined then [CUDA.WriteCombined] else [] in+  bracket (CUDA.mallocHostArray f n) (CUDA.freeHost) $ \h_ptr ->+  CUDA.withListArray [1..n]                          $ \d_ptr ->+  bench n (CUDA.peekArrayAsync n d_ptr h_ptr Nothing)++bandwidth DeviceToDevice _ n      =+  CUDA.withListArray [1..n] $ \d_src ->+  CUDA.allocaArray n        $ \d_dst ->+  (\x->2*x) `fmap` bench n (CUDA.copyArray n d_src d_dst)+++-- Testing modes+--+runTests :: [Int] -> IO [(String,[Double])]+runTests bytes = sequence $+  [ run m c | m <- [List ..], c <- [HostToDevice,DeviceToHost]] ++ [ run Pageable DeviceToDevice ]+  where+    run m c =+      mapM (\b -> bandwidth c m (b `div` sizeOf (undefined::Int))) bytes >>= \t ->+      return (sc c ++ sm m, t)++    sc DeviceToHost   = "D->H"+    sc HostToDevice   = "H->D"+    sc DeviceToDevice = "D->D"+    sm List           = " (List)"+    sm Pinned         = " (Pinned)"+    sm WriteCombined  = " (WC)"+    sm Pageable       = ""+++--------------------------------------------------------------------------------+-- Main+--------------------------------------------------------------------------------++printQuick :: [(String,[Double])] -> IO ()+printQuick = printDoc . ppAsRows 1 . (++) header . map k+  where+    k (x,y) = text x : map float' y+    float'  = text . flip (showFFloat (Just 2)) ""+    header  = map (map text) [["Mode", "Bandwidth (MB/s)"]+                             ,["----", "----------------"]]++printMany :: [Int] -> [(String,[Double])] -> IO ()+printMany xs tests =+  printDoc . ppAsRows 1 . (++) header . zipWith k xs . transpose . snd . unzip $ tests+  where+    k b ys  = int b : map float' ys+    float'  = text . flip (showFFloat (Just 2)) ""++    titles    = "Size (bytes)" : fst (unzip tests)+    seperator = map (flip replicate '-' . length) $ titles+    header    = map (map text) [titles,seperator]+++parseOptions :: [String] -> IO (Options, [String])+parseOptions argv =+  case getOpt Permute options argv of+    (o,n,[])   -> return (foldl (flip id) defaultOptions o, n)+    (_,_,errs) -> do putStrLn $ concat errs ++ usageInfo header options+                     exitFailure+  where+    header = "Usage: bandwidthTest [OPTION...]"++shmooBytes :: [Int]+shmooBytes =+  [    kilobyte,   2*kilobyte  ..  20*kilobyte ] +++  [ 22*kilobyte,  24*kilobyte  ..  50*kilobyte ] +++  [ 60*kilobyte,  70*kilobyte  .. 100*kilobyte ] +++  [ 200*kilobyte, 300*kilobyte ..   1*megabyte ] +++  [   2*megabyte,   3*megabyte ..  16*megabyte ] +++  [  18*megabyte,  20*megabyte ..  32*megabyte ] +++  [  36*megabyte,  40*megabyte ..  64*megabyte ]++-- Main+--+main :: IO ()+main = do+  (opts,_) <- parseOptions =<< getArgs+  props    <- CUDA.props (device opts)+  putStrLn $  "Device " ++ show (device opts) ++ ": " ++ (CUDA.deviceName props)++  let (s,i,e) = range opts+      bytes   = [s,(s+i)..e]++  case testMode opts of+    Quick -> putStrLn "Quick mode: Bandwidth of 32MB transfer"+    _     -> putStrLn "Bandwidth measured in MB/s"++  putStrLn "Please wait...\n"+  case testMode opts of+    Range -> runTests bytes         >>= printMany bytes+    Shmoo -> runTests shmooBytes    >>= printMany shmooBytes+    Quick -> runTests [32*megabyte] >>= printQuick+
+ examples/src/bandwidthTest/Makefile view
@@ -0,0 +1,18 @@+#+# Baking!+#++# ------------------------------------------------------------------------------+# Input files+# ------------------------------------------------------------------------------+EXECUTABLE      := bandwidthTest++HSMAIN          := BandwidthTest.hs+CUFILES         :=++EXTRALIBS       :=++# ------------------------------------------------------------------------------+# Haskell/CUDA build system+# ------------------------------------------------------------------------------+include ../../common/common.mk
+ examples/src/bandwidthTest/results/GT120.pdf view

binary file changed (absent → 38879 bytes)

+ examples/src/bandwidthTest/results/Tesla.pdf view

binary file changed (absent → 38483 bytes)

+ examples/src/fold/Fold.chs view
@@ -0,0 +1,79 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+--+-- Module    : Fold+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Reduce a vector to a single value+--+--------------------------------------------------------------------------------++module Main where++#include "fold.h"++-- Friends+import C2HS+import Time+import RandomVector++-- System+import Control.Exception+import qualified Foreign.CUDA.Runtime as CUDA+++--------------------------------------------------------------------------------+-- Reference+--------------------------------------------------------------------------------++foldRef :: Num e => [e] -> IO e+foldRef xs = do+  (t,r) <- benchmark 100 (evaluate (foldl (+) 0 xs)) (return ())+  putStrLn $ "== Reference: " ++ shows (fromInteger (timeIn millisecond t)/100::Float) " ms"+  return r++--------------------------------------------------------------------------------+-- CUDA+--------------------------------------------------------------------------------++--+-- Note that this requires two memory copies: once from a Haskell list to the C+-- heap, and from there into the graphics card memory. See the `bandwidthTest'+-- example for the atrocious performance of this operation.+--+-- For this test, cheat a little and just time the pure computation.+--+foldCUDA :: [Float] -> IO Float+foldCUDA xs = do+  let len = length xs+  CUDA.withListArray xs $ \d_xs -> do+    (t,r) <- benchmark 100 (fold_plusf d_xs len) CUDA.sync+    putStrLn $ "== CUDA: " ++ shows (fromInteger (timeIn millisecond t)/100::Float) " ms"+    return r++{# fun unsafe fold_plusf+  { withDP* `CUDA.DevicePtr Float'+  ,         `Int'+  }+  -> `Float' cFloatConv #}+  where+    withDP p a = CUDA.withDevicePtr p $ \p' -> a (castPtr p')+++--------------------------------------------------------------------------------+-- Main+--------------------------------------------------------------------------------++main :: IO ()+main = do+  dev   <- CUDA.get+  props <- CUDA.props dev+  putStrLn $ "Using device " ++ show dev ++ ": " ++ CUDA.deviceName props++  xs   <- randomList 30000+  ref  <- foldRef xs+  cuda <- foldCUDA xs++  putStrLn $ "== Validating: " ++ if ((ref-cuda)/ref) < 0.0001 then "Ok!" else "INVALID!"+
+ examples/src/fold/Makefile view
@@ -0,0 +1,18 @@+#+# Baking!+#++# ------------------------------------------------------------------------------+# Input files+# ------------------------------------------------------------------------------+EXECUTABLE      := fold++HSMAIN          := Fold.chs+CUFILES         := fold.cu++EXTRALIBS       := stdc++++# ------------------------------------------------------------------------------+# Haskell/CUDA build system+# ------------------------------------------------------------------------------+include ../../common/common.mk
+ examples/src/fold/fold.cu view
@@ -0,0 +1,223 @@+/* -----------------------------------------------------------------------------+ *+ * Module    : Fold+ * Copyright : (c) 2009 Trevor L. McDonell+ * License   : BSD+ *+ * ---------------------------------------------------------------------------*/++#include "fold.h"++#include "utils.h"+#include "operator.h"+#include "cudpp/shared_mem.h"+++/*+ * Compute multiple elements per thread sequentially. This reduces the overall+ * cost of the algorithm while keeping the work complexity O(n) and the step+ * complexity O(log n). c.f. Brent's Theorem optimisation.+ *+ * Stolen from the CUDA SDK examples+ */+template <unsigned int blockSize, bool lengthIsPow2, class op, typename T>+__global__ static void+fold_recursive+(+    const T     *d_xs,+    T           *d_ys,+    int         length+)+{+    SharedMemory<T> smem;+    T *scratch = smem.getPointer();++    /*+     * Calculate first level of reduction reading into shared memory+     */+    unsigned int i;+    unsigned int tid      = threadIdx.x;+    unsigned int gridSize = blockSize * 2 * gridDim.x;++    scratch[tid] = op::identity();++    /*+     * Reduce multiple elements per thread. The number is determined by the+     * number of active thread blocks (via gridDim). More blocks will result in+     * a larger `gridSize', and hence fewer elements per thread+     *+     * The loop stride of `gridSize' is used to maintain coalescing.+     */+    for (i =  blockIdx.x * blockSize * 2 + tid; i <  length; i += gridSize)+    {+        scratch[tid] = op::apply(scratch[tid], d_xs[i]);++        /*+         * Ensure we don't read out of bounds. This is optimised away if the+         * input length is a power of two+         */+        if (lengthIsPow2 || i + blockSize < length)+            scratch[tid] = op::apply(scratch[tid], d_xs[i+blockSize]);+    }+    __syncthreads();++    /*+     * Now, calculate the reduction in shared memory+     */+    if (blockSize >= 512) { if (tid < 256) { scratch[tid] = op::apply(scratch[tid], scratch[tid+256]); } __syncthreads(); }+    if (blockSize >= 256) { if (tid < 128) { scratch[tid] = op::apply(scratch[tid], scratch[tid+128]); } __syncthreads(); }+    if (blockSize >= 128) { if (tid <  64) { scratch[tid] = op::apply(scratch[tid], scratch[tid+ 64]); } __syncthreads(); }++#ifndef __DEVICE_EMULATION__+    if (tid < 32)+#endif+    {+        if (blockSize >= 64) { scratch[tid] = op::apply(scratch[tid], scratch[tid+32]);  __EMUSYNC; }+        if (blockSize >= 32) { scratch[tid] = op::apply(scratch[tid], scratch[tid+16]);  __EMUSYNC; }+        if (blockSize >= 16) { scratch[tid] = op::apply(scratch[tid], scratch[tid+ 8]);  __EMUSYNC; }+        if (blockSize >=  8) { scratch[tid] = op::apply(scratch[tid], scratch[tid+ 4]);  __EMUSYNC; }+        if (blockSize >=  4) { scratch[tid] = op::apply(scratch[tid], scratch[tid+ 2]);  __EMUSYNC; }+        if (blockSize >=  2) { scratch[tid] = op::apply(scratch[tid], scratch[tid+ 1]);  __EMUSYNC; }+    }++    /*+     * Write the results of this block back to global memory+     */+    if (tid == 0)+        d_ys[blockIdx.x] = scratch[0];+}+++/*+ * Wrapper function for kernel launch+ */+template <class op, typename T>+static void+fold_dispatch+(+    const T     *d_xs,+    T           *d_ys,+    int         length,+    int         blocks,+    int         threads+)+{+    unsigned int smem = threads * sizeof(T);++    if (isPow2(length))+    {+        switch (threads)+        {+        case 512: fold_recursive<512,true,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case 256: fold_recursive<256,true,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case 128: fold_recursive<128,true,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case  64: fold_recursive< 64,true,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case  32: fold_recursive< 32,true,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case  16: fold_recursive< 16,true,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case   8: fold_recursive<  8,true,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case   4: fold_recursive<  4,true,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case   2: fold_recursive<  2,true,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case   1: fold_recursive<  1,true,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        default:+            assert(!"Non-exhaustive patterns in match");+        }+    }+    else+    {+        switch (threads)+        {+        case 512: fold_recursive<512,false,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case 256: fold_recursive<256,false,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case 128: fold_recursive<128,false,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case  64: fold_recursive< 64,false,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case  32: fold_recursive< 32,false,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case  16: fold_recursive< 16,false,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case   8: fold_recursive<  8,false,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case   4: fold_recursive<  4,false,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case   2: fold_recursive<  2,false,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        case   1: fold_recursive<  1,false,op,T><<<blocks,threads,smem>>>(d_xs, d_ys, length); break;+        default:+            assert(!"Non-exhaustive patterns in match");+        }+    }+}+++/*+ * Compute the number of blocks and threads to use for the reduction kernel+ */+static void+fold_control+(+    int         n,+    int         &blocks,+    int         &threads,+    int         maxThreads = MAX_THREADS,+    int         maxBlocks  = MAX_BLOCKS+)+{+    threads = (n < maxThreads*2) ? ceilPow2((n+1)/2) : maxThreads;+    blocks  = (n + threads * 2 - 1) / (threads * 2);+    blocks  = min(blocks, maxBlocks);+}+++/*+ * Apply a binary operator to an array, reducing the array to a single value.+ * The reduction will take place in parallel, so the operator must be+ * associative.+ */+template <class op, typename T>+T+fold+(+    const T     *d_xs,+    int         n+)+{+    int blocks;+    int threads;+    T   gpu_result;+    T*  d_data          = NULL;++    /*+     * Allocate temporary storage for the block-level reduction+     */+    fold_control(n, blocks, threads);+    cudaMalloc((void **) &d_data, sizeof(T) * blocks);++    /*+     * Recursively fold the partial block sums to a single value+     */+    fold_dispatch<op,T>(d_xs, d_data, n, blocks, threads);++    n = blocks;+    while (n > 1)+    {+        fold_control(n, blocks, threads);+        fold_dispatch<op,T>(d_data, d_data, n, blocks, threads);++        n = (n + threads * 2 - 1) / (threads * 2);+    }+    assert(n == 1);++    /*+     * Read back the final result+     */+    cudaMemcpy(&gpu_result, d_data, sizeof(T), cudaMemcpyDeviceToHost);+    cudaFree(d_data);++    return gpu_result;+}+++// -----------------------------------------------------------------------------+// Instances+// -----------------------------------------------------------------------------++float fold_plusf(float *xs, int N)+{+    float  result = fold< Plus<float> >(xs, N);+    return result;+}+
+ examples/src/fold/fold.h view
@@ -0,0 +1,32 @@+/* -----------------------------------------------------------------------------+ *+ * Module    : Fold+ * Copyright : (c) 2009 Trevor L. McDonell+ * License   : BSD+ *+ * ---------------------------------------------------------------------------*/++#ifndef __FOLD_H__+#define __FOLD_H__++/*+ * Optimised for Tesla.+ * Maximum thread occupancy for your card may be achieved with different values.+ */+#define MAX_THREADS     128+#define MAX_BLOCKS      64++#ifdef __cplusplus+extern "C" {+#endif++/*+ * Instances+ */+float fold_plusf(float *xs, int N);+++#ifdef __cplusplus+}+#endif+#endif
+ examples/src/matrixMul/LICENSE view
@@ -0,0 +1,187 @@+NVIDIA Corporation GPU COMPUTING SDK END USER LICENSE AGREEMENT ("Agreement")++BY DOWNLOADING THE SOFTWARE AND OTHER AVAILABLE MATERIALS, YOU  ("DEVELOPER")+AGREE TO BE BOUND BY THE FOLLOWING TERMS AND CONDITIONS OF THIS AGREEMENT.  IF+DEVELOPER DOES NOT AGREE TO THE TERMS AND CONDITION OF THIS AGREEMENT, THEN DO+NOT DOWNLOAD THE SOFTWARE AND MATERIALS.++The materials available for download to Developers may include software in both+sample source ("Source Code") and object code ("Object Code") versions,+documentation ("Documentation"), certain art work ("Art Assets") and other+materials (collectively, these materials referred to herein as "Materials").+Except as expressly indicated herein, all terms and conditions of this Agreement+apply to all of the Materials.++Except as expressly set forth herein, NVIDIA owns all of the Materials and makes+them available to Developer only under the terms and conditions set forth in+this Agreement.++License:  Subject to the terms of this Agreement, NVIDIA hereby grants to+Developer a royalty-free, non-exclusive license to possess and to use the+Materials.  Developer may install and use multiple copies of the Materials on a+shared computer or concurrently on different computers, and make multiple+back-up copies of the Materials, solely for Licensee's use within Licensee's+Enterprise. "Enterprise" shall mean individual use by Licensee or any legal+entity (such as a corporation or university) and the subsidiaries it owns by+more than 50 percent.  The following terms apply to the specified type of+Material:++Source Code:  Developer shall have the right to modify and create derivative+works with the Source Code.  Developer shall own any derivative works+("Derivatives") it creates to the Source Code, provided that Developer uses the+Materials in accordance with the terms and conditions of this Agreement.+Developer may distribute the Derivatives, provided that all NVIDIA copyright+notices and trademarks are used properly and the Derivatives include the+following statement: "This software contains source code provided by NVIDIA+Corporation."  ++Object Code:  Developer agrees not to disassemble, decompile or reverse engineer+the Object Code versions of any of the Materials.  Developer acknowledges that+certain of the Materials provided in Object Code version may contain third party+components that may be subject to restrictions, and expressly agrees not to+attempt to modify or distribute such Materials without first receiving consent+from NVIDIA.++Art Assets:  Developer shall have the right to modify and create Derivatives of+the Art Assets, but may not distribute any of the Art Assets or Derivatives+created therefrom without NVIDIA's prior written consent.++No Other License: No rights or licenses are granted by NVIDIA to Developer under+this Agreement, expressly or by implication, with respect to any proprietary+information or patent, copyright, trade secret or other intellectual property+right owned or controlled by NVIDIA, except as expressly provided in this+Agreement.++Intellectual Property Ownership: All rights, title, interest and copyrights in+and to the Materials (including but not limited to all images, photographs,+animations, video, audio, music, text, and other information incorporated into+the Materials), are owned by NVIDIA, or its suppliers. The Materials are+protected by copyright laws and international treaty provisions. Accordingly,+Developer is required to treat the Materials like any other copyrighted+material, except as otherwise allowed pursuant to this Agreement. +++Term of Agreement:  This Agreement is effective until (i) automatically+terminated if Developer fails to comply with any of the terms and conditions of+this Agreement; or (ii) terminated by NVIDIA.  NVIDIA may terminate this+Agreement (and with it, all of Developer's right to the Materials) immediately+upon written notice (which may include email) to Developer, with or without+cause.++Defensive Suspension: If Developer commences or participates in any legal+proceeding against NVIDIA, then NVIDIA may, in its sole discretion, suspend or+terminate all license grants and any other rights provided under this Agreement+during the pendency of such legal proceedings. +++No Support:  NVIDIA has no obligation to support or to continue providing or+updating any of the Materials.++No Warranty:  THE SOFTWARE AND ANY OTHER MATERIALS PROVIDED BY NVIDIA TO+DEVELOPER HEREUNDER ARE PROVIDED "AS IS."  NVIDIA DISCLAIMS ALL WARRANTIES,+EXPRESS, IMPLIED OR STATUTORY, INCLUDING, WITHOUT LIMITATION, THE IMPLIED+WARRANTIES OF TITLE, MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND+NONINFRINGEMENT.++Limitation of Liability: NVIDIA SHALL NOT BE LIABLE TO DEVELOPER, DEVELOPER'S+CUSTOMERS, OR ANY OTHER PERSON OR ENTITY CLAIMING THROUGH OR UNDER DEVELOPER FOR+ANY LOSS OF PROFITS, INCOME, SAVINGS, OR ANY OTHER CONSEQUENTIAL, INCIDENTAL,+SPECIAL, PUNITIVE, DIRECT OR INDIRECT DAMAGES (WHETHER IN AN ACTION IN CONTRACT,+TORT OR BASED ON A WARRANTY), EVEN IF NVIDIA HAS BEEN ADVISED OF THE POSSIBILITY+OF SUCH DAMAGES.  THESE LIMITATIONS SHALL APPLY NOTWITHSTANDING ANY FAILURE OF+THE ESSENTIAL PURPOSE OF ANY LIMITED REMEDY.  IN NO EVENT SHALL NVIDIA'S+AGGREGATE LIABILITY TO DEVELOPER OR ANY OTHER PERSON OR ENTITY CLAIMING THROUGH+OR UNDER DEVELOPER EXCEED THE AMOUNT OF MONEY ACTUALLY PAID BY DEVELOPER TO+NVIDIA FOR THE SOFTWARE OR ANY OTHER MATERIALS.++Applicable Law: This Agreement shall be deemed to have been made in, and shall+be construed pursuant to, the laws of the State of Delaware. The United Nations+Convention on Contracts for the International Sale of Goods is specifically+disclaimed.++Feedback: In the event Developer contacts NVIDIA to request Feedback (as defined+below) on how to optimize Developer's product for use with the Materials, the+following terms and conditions apply the Feedback:++1.	Exchange of Feedback. Both parties agree that neither party has an+obligation to give the other party any suggestions, comments or other feedback,+whether verbally or in code form ("Feedback"), relating to (i) the Materials;+(ii) Developer's products; (iii) Developer's use of the Materials; or (iv)+optimization of Developer's product with NVIDIA product(s).  In the event either+party provides Feedback to the other party, the party receiving the Feedback may+use and include any Feedback that the other party voluntarily provides to+improve the (i) Materials or other related NVIDIA technologies, respectively for+the benefit of NVIDIA; or (ii) Developer's product or other related Developer+technologies, respectively for the benefit of Developer.  Accordingly, if either+party provides Feedback to the other party, both parties agree that the other+party and its respective Developers may freely use, reproduce, license,+distribute, and otherwise commercialize the Feedback in the (i) Materials or+other related technologies; or (ii) Developer's products or other related+technologies, respectively, without the payment of any royalties or fees.+  +2.	Residual Rights. Developer agrees that NVIDIA shall be free to use any+general knowledge, skills and experience, (including, but not limited to, ideas,+concepts, know-how, or techniques) ("Residuals"), contained in the (i) Feedback+provided by Developer to NVIDIA; (ii) Developer's products, in source or object+code form, shared or disclosed to NVIDIA in connection with the Feedback; or (c)+Developer's confidential information voluntarily provided to NVIDIA in+connection with the Feedback, which are retained in the memories of NVIDIA's+employees, agents, or contractors who have had access to such (i) Feedback+provided by Developer to NVIDIA; (ii) Developer's products; or (c) Developer's+confidential information voluntarily provided to NVIDIA, in connection with the+Feedback.  Subject to the terms and conditions of this Agreement, NVIDIA's+employees, agents, or contractors shall not be prevented from using Residuals as+part of such employee's, agent's or contractor's general knowledge, skills,+experience, talent, and/or expertise. NVIDIA shall not have any obligation to+limit or restrict the assignment of such employees, agents or contractors or to+pay royalties for any work resulting from the use of Residuals.++3.	Disclaimer of Warranty. FEEDBACK FROM EITHER PARTY IS PROVIDED FOR THE+OTHER PARTY'S USE "AS IS" AND BOTH PARTIES DISCLAIM ALL WARRANTIES, EXPRESS,+IMPLIED AND STATUTORY INCLUDING, WITHOUT LIMITATION, THE IMPLIED WARRANTIES OF+MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE.  BOTH PARTIES DO NOT+REPRESENT OR WARRANT THAT THE FEEDBACK WILL MEET THE OTHER PARTY'S REQUIREMENTS+OR THAT THE OPERATION OR IMPLEMENTATION OF THE FEEDBACK WILL BE UNINTERRUPTED OR+ERROR-FREE.++4.	No Liability for Consequential Damages. TO THE MAXIMUM EXTENT PERMITTED+BY APPLICABLE LAW, IN NO EVENT SHALL EITHER PARTY OR ITS SUPPLIERS BE LIABLE FOR+ANY SPECIAL, INCIDENTAL, INDIRECT, OR CONSEQUENTIAL DAMAGES WHATSOEVER+(INCLUDING, WITHOUT LIMITATION, DAMAGES FOR LOSS OF BUSINESS PROFITS, BUSINESS+INTERRUPTION, LOSS OF BUSINESS INFORMATION, OR ANY OTHER PECUNIARY LOSS) ARISING+OUT OF THE USE OF OR INABILITY TO USE THE FEEDBACK, EVEN IF THE OTHER PARTY HAS+BEEN ADVISED OF THE POSSIBILITY OF SUCH DAMAGES.++5.	Freedom of Action.  Developer agrees that this Agreement is nonexclusive+and NVIDIA may currently or in the future be developing software, other+technology or confidential information internally, or receiving confidential+information from other parties that maybe similar to the Feedback and+Developer's confidential information (as provided in subsection 2 above), which+may be provided to NVIDIA in connection with Feedback by Developer. Accordingly,+Developer agrees that nothing in this Agreement will be construed as a+representation or inference that NVIDIA will not develop, design, manufacture,+acquire, market products, or have products developed, designed, manufactured,+acquired, or marketed for NVIDIA, that compete with the Developer's products or+confidential information.++RESTRICTED RIGHTS NOTICE: Materials has been developed entirely at private+expense and is commercial computer software provided with RESTRICTED RIGHTS.+Use, duplication or disclosure by the U.S. Government or a U.S. Government+subcontractor is subject to the restrictions set forth in the license agreement+under which Materials was obtained pursuant to DFARS 227.7202-3(a) or as set+forth in subparagraphs (c)(1) and (2) of the Commercial Computer Software -+Restricted Rights clause at FAR 52.227-19, as applicable.+Contractor/manufacturer is NVIDIA, 2701 San Tomas Expressway, Santa Clara, CA+95050.++Miscellaneous: If any provision of this Agreement is inconsistent with, or+cannot be fully enforced under, the law, such provision will be construed as+limited to the extent necessary to be consistent with and fully enforceable+under the law. This Agreement is the final, complete and exclusive agreement+between the parties relating to the subject matter hereof, and supersedes all+prior or contemporaneous understandings and agreements relating to such subject+matter, whether oral or written. This Agreement may only be modified in writing+signed by an authorized officer of NVIDIA. Developer agrees that it will not+ship, transfer or export the Materials into any country, or use the Materials in+any manner, prohibited by the United States Bureau of Industry and Security or+any export laws, restrictions or regulations.
+ examples/src/matrixMul/Makefile view
@@ -0,0 +1,18 @@+#+# Baking!+#++# ------------------------------------------------------------------------------+# Input files+# ------------------------------------------------------------------------------+EXECUTABLE      := matrixMul++HSMAIN          := MatrixMul.hs+CUFILES         := matrix_mul.cu++EXTRALIBS       := stdc++++# ------------------------------------------------------------------------------+# Haskell/CUDA build system+# ------------------------------------------------------------------------------+include ../../common/common.mk
+ examples/src/matrixMul/MatrixMul.hs view
@@ -0,0 +1,114 @@+{-# LANGUAGE CPP #-}+--------------------------------------------------------------------------------+--+-- Module    : MatrixMul+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Matrix multiplication using runtime interface and execution control instead+-- of calling C functions via the FFI.+--+--------------------------------------------------------------------------------++module Main where++#include "matrix_mul.h"++-- Friends+import Time+import RandomVector++-- System+import Data.Array+import System.IO+import Foreign+import qualified Foreign.CUDA as CUDA+++--------------------------------------------------------------------------------+-- Reference+--------------------------------------------------------------------------------++matMult :: (Num e, Storable e) => Matrix e -> Matrix e -> IO (Matrix e)+matMult mx my = do+  x <- unsafeFreeze mx+  y <- unsafeFreeze my+  let ((li, lj), (ui, uj))  = bounds x+      ((li',lj'),(ui',uj')) = bounds y+      resBnds | (lj,uj) == (li',ui') = ((li,lj'),(ui,uj'))+              | otherwise            = error "matrix dimensions must agree"++  newListArray resBnds [sum [x!(i,k) * y!(k,j) | k <- range (lj,uj)]+                         | i <- range (li,ui)+                         , j <- range (lj',uj') ]+++--------------------------------------------------------------------------------+-- CUDA+--------------------------------------------------------------------------------++matMultCUDA :: (Num e, Storable e) => Matrix e -> Matrix e -> IO (Matrix e)+matMultCUDA xs' ys' = doMult undefined xs' ys'+  where+    doMult :: (Num e', Storable e') => e' -> Matrix e' -> Matrix e' -> IO (Matrix e')+    doMult dummy xs ys = do++      -- Setup matrix parameters+      --+      ((li, lj), (ui, uj))  <- getBounds xs+      ((li',lj'),(ui',uj')) <- getBounds ys+      let wx = rangeSize (lj,uj)+          hx = rangeSize (li,ui)+          wy = rangeSize (lj',uj')+          hy = rangeSize (li',ui')+          resBnds | wx == hy  = ((li,lj'),(ui,uj'))+                  | otherwise = error "matrix dimensions must agree"++      -- Allocate memory and copy test data+      --+      CUDA.allocaArray (wx*hx) $ \d_xs -> do+      CUDA.allocaArray (wy*hy) $ \d_ys -> do+      CUDA.allocaArray (wy*hx) $ \d_zs -> do+      withMatrix xs $ \p -> CUDA.pokeArray (wx*hx) p d_xs+      withMatrix ys $ \p -> CUDA.pokeArray (wy*hy) p d_ys++      -- Launch the kernel+      --+      let gridDim   = (wy`div`BLOCK_SIZE, hx`div`BLOCK_SIZE)+          blockDim  = (BLOCK_SIZE,BLOCK_SIZE,1)+          sharedMem = 2 * BLOCK_SIZE * BLOCK_SIZE * fromIntegral (sizeOf dummy)++      CUDA.setConfig gridDim blockDim sharedMem Nothing+      CUDA.setParams [CUDA.VArg d_xs, CUDA.VArg d_ys, CUDA.VArg d_zs, CUDA.IArg wx, CUDA.IArg wy]+      CUDA.launch "matrixMul"++      -- Copy back result+      zs <- newArray_ resBnds+      withMatrix zs $ \p -> CUDA.peekArray (wy*hx) d_zs p+      return zs+++--------------------------------------------------------------------------------+-- Main+--------------------------------------------------------------------------------++main :: IO ()+main = do+  dev   <- CUDA.get+  props <- CUDA.props dev+  putStrLn $ "Using device " ++ show dev ++ ": " ++ CUDA.deviceName props++  xs <- randomArr ((1,1),(8*BLOCK_SIZE, 4*BLOCK_SIZE)) :: IO (Matrix Float)+  ys <- randomArr ((1,1),(4*BLOCK_SIZE,12*BLOCK_SIZE)) :: IO (Matrix Float)++  putStr   "== Reference: " >> hFlush stdout+  (tr,ref) <- benchmark 100 (matMult xs ys) (return ())+  putStrLn $  shows (fromInteger (timeIn millisecond tr) / 100::Float) " ms"++  putStr   "== CUDA: " >> hFlush stdout+  (tc,mat) <- benchmark 100 (matMultCUDA xs ys) (CUDA.sync)+  putStrLn $  shows (fromInteger (timeIn millisecond tc) / 100::Float) " ms"++  putStr "== Validating: "+  verify ref mat >>= \rv -> putStrLn $ if rv then "Ok!" else "INVALID!"+
+ examples/src/matrixMul/matrix_mul.cu view
@@ -0,0 +1,110 @@+/*+ * Copyright 1993-2009 NVIDIA Corporation.  All rights reserved.+ *+ * NVIDIA Corporation and its licensors retain all intellectual property and + * proprietary rights in and to this software and related documentation. + * Any use, reproduction, disclosure, or distribution of this software + * and related documentation without an express license agreement from+ * NVIDIA Corporation is strictly prohibited.+ *+ * Please refer to the applicable NVIDIA end user license agreement (EULA) + * associated with this source code for terms and conditions that govern + * your use of this NVIDIA software.+ * + */++/* Matrix multiplication: C = A * B.+ * Device code.+ */++#ifndef _MATRIXMUL_KERNEL_H_+#define _MATRIXMUL_KERNEL_H_++#include <stdio.h>+#include "matrix_mul.h"++#define CHECK_BANK_CONFLICTS 0+#if CHECK_BANK_CONFLICTS+#define AS(i, j) cutilBankChecker(((float*)&As[0][0]), (BLOCK_SIZE * i + j))+#define BS(i, j) cutilBankChecker(((float*)&Bs[0][0]), (BLOCK_SIZE * i + j))+#else+#define AS(i, j) As[i][j]+#define BS(i, j) Bs[i][j]+#endif++////////////////////////////////////////////////////////////////////////////////+//! Matrix multiplication on the device: C = A * B+//! wA is A's width and wB is B's width+////////////////////////////////////////////////////////////////////////////////+extern "C" __global__ void+matrixMul(float* A, float* B, float* C, int wA, int wB)+{+    // Block index+    int bx = blockIdx.x;+    int by = blockIdx.y;++    // Thread index+    int tx = threadIdx.x;+    int ty = threadIdx.y;++    // Index of the first sub-matrix of A processed by the block+    int aBegin = wA * BLOCK_SIZE * by;++    // Index of the last sub-matrix of A processed by the block+    int aEnd   = aBegin + wA - 1;++    // Step size used to iterate through the sub-matrices of A+    int aStep  = BLOCK_SIZE;++    // Index of the first sub-matrix of B processed by the block+    int bBegin = BLOCK_SIZE * bx;++    // Step size used to iterate through the sub-matrices of B+    int bStep  = BLOCK_SIZE * wB;++    // Csub is used to store the element of the block sub-matrix+    // that is computed by the thread+    float Csub = 0;++    // Loop over all the sub-matrices of A and B+    // required to compute the block sub-matrix+    for (int a = aBegin, b = bBegin;+             a <= aEnd;+             a += aStep, b += bStep) {++        // Declaration of the shared memory array As used to+        // store the sub-matrix of A+        __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];++        // Declaration of the shared memory array Bs used to+        // store the sub-matrix of B+        __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];++        // Load the matrices from device memory+        // to shared memory; each thread loads+        // one element of each matrix+        AS(ty, tx) = A[a + wA * ty + tx];+        BS(ty, tx) = B[b + wB * ty + tx];++        // Synchronize to make sure the matrices are loaded+        __syncthreads();++        // Multiply the two matrices together;+        // each thread computes one element+        // of the block sub-matrix+        for (int k = 0; k < BLOCK_SIZE; ++k)+            Csub += AS(ty, k) * BS(k, tx);++        // Synchronize to make sure that the preceding+        // computation is done before loading two new+        // sub-matrices of A and B in the next iteration+        __syncthreads();+    }++    // Write the block sub-matrix to device memory;+    // each thread writes one element+    int c = wB * BLOCK_SIZE * by + BLOCK_SIZE * bx;+    C[c + wB * ty + tx] = Csub;+}++#endif // #ifndef _MATRIXMUL_KERNEL_H_
+ examples/src/matrixMul/matrix_mul.h view
@@ -0,0 +1,33 @@+/*+ * Copyright 1993-2009 NVIDIA Corporation.  All rights reserved.+ *+ * NVIDIA Corporation and its licensors retain all intellectual property and + * proprietary rights in and to this software and related documentation. + * Any use, reproduction, disclosure, or distribution of this software + * and related documentation without an express license agreement from+ * NVIDIA Corporation is strictly prohibited.+ *+ * Please refer to the applicable NVIDIA end user license agreement (EULA) + * associated with this source code for terms and conditions that govern + * your use of this NVIDIA software.+ * + */++#ifndef _MATRIXMUL_H_+#define _MATRIXMUL_H_++/* Thread block size */+#define BLOCK_SIZE 16++/* Matrix dimensions+ * (chosen as multiples of the thread block size for simplicity)+ */+#define WA (3 * BLOCK_SIZE) /* Matrix A width  */+#define HA (5 * BLOCK_SIZE) /* Matrix A height */+#define WB (8 * BLOCK_SIZE) /* Matrix B width  */+#define HB WA  /* Matrix B height */+#define WC WB  /* Matrix C width  */+#define HC HA  /* Matrix C height */++#endif /* _MATRIXMUL_H_ */+
+ examples/src/matrixMulDrv/LICENSE view
@@ -0,0 +1,187 @@+NVIDIA Corporation GPU COMPUTING SDK END USER LICENSE AGREEMENT ("Agreement")++BY DOWNLOADING THE SOFTWARE AND OTHER AVAILABLE MATERIALS, YOU  ("DEVELOPER")+AGREE TO BE BOUND BY THE FOLLOWING TERMS AND CONDITIONS OF THIS AGREEMENT.  IF+DEVELOPER DOES NOT AGREE TO THE TERMS AND CONDITION OF THIS AGREEMENT, THEN DO+NOT DOWNLOAD THE SOFTWARE AND MATERIALS.++The materials available for download to Developers may include software in both+sample source ("Source Code") and object code ("Object Code") versions,+documentation ("Documentation"), certain art work ("Art Assets") and other+materials (collectively, these materials referred to herein as "Materials").+Except as expressly indicated herein, all terms and conditions of this Agreement+apply to all of the Materials.++Except as expressly set forth herein, NVIDIA owns all of the Materials and makes+them available to Developer only under the terms and conditions set forth in+this Agreement.++License:  Subject to the terms of this Agreement, NVIDIA hereby grants to+Developer a royalty-free, non-exclusive license to possess and to use the+Materials.  Developer may install and use multiple copies of the Materials on a+shared computer or concurrently on different computers, and make multiple+back-up copies of the Materials, solely for Licensee's use within Licensee's+Enterprise. "Enterprise" shall mean individual use by Licensee or any legal+entity (such as a corporation or university) and the subsidiaries it owns by+more than 50 percent.  The following terms apply to the specified type of+Material:++Source Code:  Developer shall have the right to modify and create derivative+works with the Source Code.  Developer shall own any derivative works+("Derivatives") it creates to the Source Code, provided that Developer uses the+Materials in accordance with the terms and conditions of this Agreement.+Developer may distribute the Derivatives, provided that all NVIDIA copyright+notices and trademarks are used properly and the Derivatives include the+following statement: "This software contains source code provided by NVIDIA+Corporation."  ++Object Code:  Developer agrees not to disassemble, decompile or reverse engineer+the Object Code versions of any of the Materials.  Developer acknowledges that+certain of the Materials provided in Object Code version may contain third party+components that may be subject to restrictions, and expressly agrees not to+attempt to modify or distribute such Materials without first receiving consent+from NVIDIA.++Art Assets:  Developer shall have the right to modify and create Derivatives of+the Art Assets, but may not distribute any of the Art Assets or Derivatives+created therefrom without NVIDIA's prior written consent.++No Other License: No rights or licenses are granted by NVIDIA to Developer under+this Agreement, expressly or by implication, with respect to any proprietary+information or patent, copyright, trade secret or other intellectual property+right owned or controlled by NVIDIA, except as expressly provided in this+Agreement.++Intellectual Property Ownership: All rights, title, interest and copyrights in+and to the Materials (including but not limited to all images, photographs,+animations, video, audio, music, text, and other information incorporated into+the Materials), are owned by NVIDIA, or its suppliers. The Materials are+protected by copyright laws and international treaty provisions. Accordingly,+Developer is required to treat the Materials like any other copyrighted+material, except as otherwise allowed pursuant to this Agreement. +++Term of Agreement:  This Agreement is effective until (i) automatically+terminated if Developer fails to comply with any of the terms and conditions of+this Agreement; or (ii) terminated by NVIDIA.  NVIDIA may terminate this+Agreement (and with it, all of Developer's right to the Materials) immediately+upon written notice (which may include email) to Developer, with or without+cause.++Defensive Suspension: If Developer commences or participates in any legal+proceeding against NVIDIA, then NVIDIA may, in its sole discretion, suspend or+terminate all license grants and any other rights provided under this Agreement+during the pendency of such legal proceedings. +++No Support:  NVIDIA has no obligation to support or to continue providing or+updating any of the Materials.++No Warranty:  THE SOFTWARE AND ANY OTHER MATERIALS PROVIDED BY NVIDIA TO+DEVELOPER HEREUNDER ARE PROVIDED "AS IS."  NVIDIA DISCLAIMS ALL WARRANTIES,+EXPRESS, IMPLIED OR STATUTORY, INCLUDING, WITHOUT LIMITATION, THE IMPLIED+WARRANTIES OF TITLE, MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND+NONINFRINGEMENT.++Limitation of Liability: NVIDIA SHALL NOT BE LIABLE TO DEVELOPER, DEVELOPER'S+CUSTOMERS, OR ANY OTHER PERSON OR ENTITY CLAIMING THROUGH OR UNDER DEVELOPER FOR+ANY LOSS OF PROFITS, INCOME, SAVINGS, OR ANY OTHER CONSEQUENTIAL, INCIDENTAL,+SPECIAL, PUNITIVE, DIRECT OR INDIRECT DAMAGES (WHETHER IN AN ACTION IN CONTRACT,+TORT OR BASED ON A WARRANTY), EVEN IF NVIDIA HAS BEEN ADVISED OF THE POSSIBILITY+OF SUCH DAMAGES.  THESE LIMITATIONS SHALL APPLY NOTWITHSTANDING ANY FAILURE OF+THE ESSENTIAL PURPOSE OF ANY LIMITED REMEDY.  IN NO EVENT SHALL NVIDIA'S+AGGREGATE LIABILITY TO DEVELOPER OR ANY OTHER PERSON OR ENTITY CLAIMING THROUGH+OR UNDER DEVELOPER EXCEED THE AMOUNT OF MONEY ACTUALLY PAID BY DEVELOPER TO+NVIDIA FOR THE SOFTWARE OR ANY OTHER MATERIALS.++Applicable Law: This Agreement shall be deemed to have been made in, and shall+be construed pursuant to, the laws of the State of Delaware. The United Nations+Convention on Contracts for the International Sale of Goods is specifically+disclaimed.++Feedback: In the event Developer contacts NVIDIA to request Feedback (as defined+below) on how to optimize Developer's product for use with the Materials, the+following terms and conditions apply the Feedback:++1.	Exchange of Feedback. Both parties agree that neither party has an+obligation to give the other party any suggestions, comments or other feedback,+whether verbally or in code form ("Feedback"), relating to (i) the Materials;+(ii) Developer's products; (iii) Developer's use of the Materials; or (iv)+optimization of Developer's product with NVIDIA product(s).  In the event either+party provides Feedback to the other party, the party receiving the Feedback may+use and include any Feedback that the other party voluntarily provides to+improve the (i) Materials or other related NVIDIA technologies, respectively for+the benefit of NVIDIA; or (ii) Developer's product or other related Developer+technologies, respectively for the benefit of Developer.  Accordingly, if either+party provides Feedback to the other party, both parties agree that the other+party and its respective Developers may freely use, reproduce, license,+distribute, and otherwise commercialize the Feedback in the (i) Materials or+other related technologies; or (ii) Developer's products or other related+technologies, respectively, without the payment of any royalties or fees.+  +2.	Residual Rights. Developer agrees that NVIDIA shall be free to use any+general knowledge, skills and experience, (including, but not limited to, ideas,+concepts, know-how, or techniques) ("Residuals"), contained in the (i) Feedback+provided by Developer to NVIDIA; (ii) Developer's products, in source or object+code form, shared or disclosed to NVIDIA in connection with the Feedback; or (c)+Developer's confidential information voluntarily provided to NVIDIA in+connection with the Feedback, which are retained in the memories of NVIDIA's+employees, agents, or contractors who have had access to such (i) Feedback+provided by Developer to NVIDIA; (ii) Developer's products; or (c) Developer's+confidential information voluntarily provided to NVIDIA, in connection with the+Feedback.  Subject to the terms and conditions of this Agreement, NVIDIA's+employees, agents, or contractors shall not be prevented from using Residuals as+part of such employee's, agent's or contractor's general knowledge, skills,+experience, talent, and/or expertise. NVIDIA shall not have any obligation to+limit or restrict the assignment of such employees, agents or contractors or to+pay royalties for any work resulting from the use of Residuals.++3.	Disclaimer of Warranty. FEEDBACK FROM EITHER PARTY IS PROVIDED FOR THE+OTHER PARTY'S USE "AS IS" AND BOTH PARTIES DISCLAIM ALL WARRANTIES, EXPRESS,+IMPLIED AND STATUTORY INCLUDING, WITHOUT LIMITATION, THE IMPLIED WARRANTIES OF+MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE.  BOTH PARTIES DO NOT+REPRESENT OR WARRANT THAT THE FEEDBACK WILL MEET THE OTHER PARTY'S REQUIREMENTS+OR THAT THE OPERATION OR IMPLEMENTATION OF THE FEEDBACK WILL BE UNINTERRUPTED OR+ERROR-FREE.++4.	No Liability for Consequential Damages. TO THE MAXIMUM EXTENT PERMITTED+BY APPLICABLE LAW, IN NO EVENT SHALL EITHER PARTY OR ITS SUPPLIERS BE LIABLE FOR+ANY SPECIAL, INCIDENTAL, INDIRECT, OR CONSEQUENTIAL DAMAGES WHATSOEVER+(INCLUDING, WITHOUT LIMITATION, DAMAGES FOR LOSS OF BUSINESS PROFITS, BUSINESS+INTERRUPTION, LOSS OF BUSINESS INFORMATION, OR ANY OTHER PECUNIARY LOSS) ARISING+OUT OF THE USE OF OR INABILITY TO USE THE FEEDBACK, EVEN IF THE OTHER PARTY HAS+BEEN ADVISED OF THE POSSIBILITY OF SUCH DAMAGES.++5.	Freedom of Action.  Developer agrees that this Agreement is nonexclusive+and NVIDIA may currently or in the future be developing software, other+technology or confidential information internally, or receiving confidential+information from other parties that maybe similar to the Feedback and+Developer's confidential information (as provided in subsection 2 above), which+may be provided to NVIDIA in connection with Feedback by Developer. Accordingly,+Developer agrees that nothing in this Agreement will be construed as a+representation or inference that NVIDIA will not develop, design, manufacture,+acquire, market products, or have products developed, designed, manufactured,+acquired, or marketed for NVIDIA, that compete with the Developer's products or+confidential information.++RESTRICTED RIGHTS NOTICE: Materials has been developed entirely at private+expense and is commercial computer software provided with RESTRICTED RIGHTS.+Use, duplication or disclosure by the U.S. Government or a U.S. Government+subcontractor is subject to the restrictions set forth in the license agreement+under which Materials was obtained pursuant to DFARS 227.7202-3(a) or as set+forth in subparagraphs (c)(1) and (2) of the Commercial Computer Software -+Restricted Rights clause at FAR 52.227-19, as applicable.+Contractor/manufacturer is NVIDIA, 2701 San Tomas Expressway, Santa Clara, CA+95050.++Miscellaneous: If any provision of this Agreement is inconsistent with, or+cannot be fully enforced under, the law, such provision will be construed as+limited to the extent necessary to be consistent with and fully enforceable+under the law. This Agreement is the final, complete and exclusive agreement+between the parties relating to the subject matter hereof, and supersedes all+prior or contemporaneous understandings and agreements relating to such subject+matter, whether oral or written. This Agreement may only be modified in writing+signed by an authorized officer of NVIDIA. Developer agrees that it will not+ship, transfer or export the Materials into any country, or use the Materials in+any manner, prohibited by the United States Bureau of Industry and Security or+any export laws, restrictions or regulations.
+ examples/src/matrixMulDrv/Makefile view
@@ -0,0 +1,18 @@+#+# Baking!+#++# ------------------------------------------------------------------------------+# Input files+# ------------------------------------------------------------------------------+EXECUTABLE	:= matrixMulDrv++HSMAIN		:= MatrixMul.hs+PTXFILES	:= matrix_mul.cu++USEDRVAPI	:= 1++# ------------------------------------------------------------------------------+# Haskell/CUDA build system+# ------------------------------------------------------------------------------+include ../../common/common.mk
+ examples/src/matrixMulDrv/MatrixMul.hs view
@@ -0,0 +1,155 @@+{-# LANGUAGE CPP #-}+--------------------------------------------------------------------------------+--+-- Module    : MatrixMul+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Matrix multiplication using driver interface+--+--------------------------------------------------------------------------------++module Main where++#include "matrix_mul.h"++-- Friends+import RandomVector++-- System+import Numeric+import Data.Array+import Control.Exception+import Data.Array.Storable+import Foreign.Storable+import qualified Data.ByteString.Char8 as B+import qualified Foreign.CUDA.Driver   as CUDA+++-- Return the (width,height) of a matrix+--+getSize :: Storable e => Matrix e -> IO (Int,Int)+getSize mat = do+  ((li,lj),(ui,uj)) <- getBounds mat+  return (rangeSize (lj,uj), rangeSize (li,ui))++--------------------------------------------------------------------------------+-- Reference implementation+--------------------------------------------------------------------------------++matMult :: (Num e, Storable e) => Matrix e -> Matrix e -> IO (Matrix e)+matMult mx my = do+  x <- unsafeFreeze mx+  y <- unsafeFreeze my+  let ((li, lj), (ui, uj))  = bounds x+      ((li',lj'),(ui',uj')) = bounds y+      resBnds | (lj,uj) == (li',ui') = ((li,lj'),(ui,uj'))+              | otherwise            = error "matrix dimensions must agree"++  newListArray resBnds [sum [x!(i,k) * y!(k,j) | k <- range (lj,uj)]+                         | i <- range (li,ui)+                         , j <- range (lj',uj') ]+++--------------------------------------------------------------------------------+-- CUDA+--------------------------------------------------------------------------------++--+-- Initialise the device and context. Load the PTX source code, and return a+-- reference to the kernel function.+--+initCUDA :: IO (CUDA.Context, CUDA.Fun)+initCUDA = do+  CUDA.initialise []+  dev     <- CUDA.device 0+  ctx     <- CUDA.create dev []+  ptx     <- B.readFile "data/matrix_mul.ptx"+  (mdl,r) <- CUDA.loadDataEx ptx [CUDA.ThreadsPerBlock (BLOCK_SIZE*BLOCK_SIZE)]+  fun     <- CUDA.getFun mdl "matrixMul"++  putStrLn $ ">> PTX JIT compilation (" ++ showFFloat (Just 2) (CUDA.jitTime r) " ms)"+  B.putStrLn (CUDA.jitInfoLog r)+  return (ctx,fun)+++--+-- Allocate some memory, and copy over the input data to the device. Should+-- probably catch allocation exceptions individually...+--+initData :: (Num e, Storable e)+         => Matrix e -> Matrix e -> IO (CUDA.DevicePtr e, CUDA.DevicePtr e, CUDA.DevicePtr e)+initData xs ys = do+  (wx,hx) <- getSize xs+  (wy,hy) <- getSize ys+  dxs     <- CUDA.mallocArray (wx*hx)+  dys     <- CUDA.mallocArray (wy*hy)+  res     <- CUDA.mallocArray (wy*hx)++  flip onException (mapM_ CUDA.free [dxs,dys,res]) $ do+  withMatrix xs $ \p -> CUDA.pokeArray (wx*hx) p dxs+  withMatrix ys $ \p -> CUDA.pokeArray (wy*hy) p dys+  return (dxs, dys, res)+++--+-- Run the test+--+testCUDA :: (Num e, Storable e) => Matrix e -> Matrix e -> IO (Matrix e)+testCUDA xs' ys' = doTest undefined xs' ys'+  where+    doTest :: (Num e', Storable e') => e' -> Matrix e' -> Matrix e' -> IO (Matrix e')+    doTest dummy xs ys = do+      (widthX,heightX)      <- getSize xs+      (widthY,_)            <- getSize ys+      ((li, lj), (ui, uj))  <- getBounds xs+      ((li',lj'),(ui',uj')) <- getBounds ys+      let resBnds | (lj,uj) == (li',ui') = ((li,lj'),(ui,uj'))+                  | otherwise            = error "matrix dimensions must agree"++      -- Initialise environment and copy over test data+      --+      putStrLn ">> Initialising"+      bracket initCUDA (\(ctx,_) -> CUDA.destroy ctx) $ \(_,matMul) -> do++      -- Ensure we release the memory, even if there was an error+      --+      putStrLn ">> Executing"+      bracket+        (initData xs ys)+        (\(dx,dy,dz) -> mapM_ CUDA.free [dx,dy,dz]) $+         \(dx,dy,dz) -> do+          -- Repeat test many times...+          --+          CUDA.setParams     matMul [CUDA.VArg dx, CUDA.VArg dy, CUDA.VArg dz, CUDA.IArg widthX, CUDA.IArg widthY]+          CUDA.setBlockShape matMul (BLOCK_SIZE,BLOCK_SIZE,1)+          CUDA.setSharedSize matMul (fromIntegral (2 * BLOCK_SIZE * BLOCK_SIZE * sizeOf dummy))+          CUDA.launch        matMul (widthY `div` BLOCK_SIZE, heightX `div` BLOCK_SIZE) Nothing+          CUDA.sync++          -- Copy back result+          --+          zs <- newArray_ resBnds+          withMatrix zs $ \p -> CUDA.peekArray (widthY*heightX) dz p+          return zs+++--------------------------------------------------------------------------------+-- Test & Verify+--------------------------------------------------------------------------------++main :: IO ()+main = do+  putStrLn "== Generating random matrices"+  xs <- randomArr ((1,1),(8*BLOCK_SIZE, 4*BLOCK_SIZE)) :: IO (Matrix Float)+  ys <- randomArr ((1,1),(4*BLOCK_SIZE,12*BLOCK_SIZE)) :: IO (Matrix Float)++  putStrLn "== Generating reference solution"+  ref <- matMult xs ys++  putStrLn "== Testing CUDA"+  mat <- testCUDA xs ys++  putStr "== Validating: "+  verify ref mat >>= \rv -> putStrLn $ if rv then "Ok!" else "INVALID!"+
+ examples/src/matrixMulDrv/matrix_mul.cu view
@@ -0,0 +1,110 @@+/*+ * Copyright 1993-2009 NVIDIA Corporation.  All rights reserved.+ *+ * NVIDIA Corporation and its licensors retain all intellectual property and + * proprietary rights in and to this software and related documentation. + * Any use, reproduction, disclosure, or distribution of this software + * and related documentation without an express license agreement from+ * NVIDIA Corporation is strictly prohibited.+ *+ * Please refer to the applicable NVIDIA end user license agreement (EULA) + * associated with this source code for terms and conditions that govern + * your use of this NVIDIA software.+ * + */++/* Matrix multiplication: C = A * B.+ * Device code.+ */++#ifndef _MATRIXMUL_KERNEL_H_+#define _MATRIXMUL_KERNEL_H_++#include <stdio.h>+#include "matrix_mul.h"++#define CHECK_BANK_CONFLICTS 0+#if CHECK_BANK_CONFLICTS+#define AS(i, j) cutilBankChecker(((float*)&As[0][0]), (BLOCK_SIZE * i + j))+#define BS(i, j) cutilBankChecker(((float*)&Bs[0][0]), (BLOCK_SIZE * i + j))+#else+#define AS(i, j) As[i][j]+#define BS(i, j) Bs[i][j]+#endif++////////////////////////////////////////////////////////////////////////////////+//! Matrix multiplication on the device: C = A * B+//! wA is A's width and wB is B's width+////////////////////////////////////////////////////////////////////////////////+extern "C" __global__ void+matrixMul(float* A, float* B, float* C, int wA, int wB)+{+    // Block index+    int bx = blockIdx.x;+    int by = blockIdx.y;++    // Thread index+    int tx = threadIdx.x;+    int ty = threadIdx.y;++    // Index of the first sub-matrix of A processed by the block+    int aBegin = wA * BLOCK_SIZE * by;++    // Index of the last sub-matrix of A processed by the block+    int aEnd   = aBegin + wA - 1;++    // Step size used to iterate through the sub-matrices of A+    int aStep  = BLOCK_SIZE;++    // Index of the first sub-matrix of B processed by the block+    int bBegin = BLOCK_SIZE * bx;++    // Step size used to iterate through the sub-matrices of B+    int bStep  = BLOCK_SIZE * wB;++    // Csub is used to store the element of the block sub-matrix+    // that is computed by the thread+    float Csub = 0;++    // Loop over all the sub-matrices of A and B+    // required to compute the block sub-matrix+    for (int a = aBegin, b = bBegin;+             a <= aEnd;+             a += aStep, b += bStep) {++        // Declaration of the shared memory array As used to+        // store the sub-matrix of A+        __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];++        // Declaration of the shared memory array Bs used to+        // store the sub-matrix of B+        __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];++        // Load the matrices from device memory+        // to shared memory; each thread loads+        // one element of each matrix+        AS(ty, tx) = A[a + wA * ty + tx];+        BS(ty, tx) = B[b + wB * ty + tx];++        // Synchronize to make sure the matrices are loaded+        __syncthreads();++        // Multiply the two matrices together;+        // each thread computes one element+        // of the block sub-matrix+        for (int k = 0; k < BLOCK_SIZE; ++k)+            Csub += AS(ty, k) * BS(k, tx);++        // Synchronize to make sure that the preceding+        // computation is done before loading two new+        // sub-matrices of A and B in the next iteration+        __syncthreads();+    }++    // Write the block sub-matrix to device memory;+    // each thread writes one element+    int c = wB * BLOCK_SIZE * by + BLOCK_SIZE * bx;+    C[c + wB * ty + tx] = Csub;+}++#endif // #ifndef _MATRIXMUL_KERNEL_H_
+ examples/src/matrixMulDrv/matrix_mul.h view
@@ -0,0 +1,33 @@+/*+ * Copyright 1993-2009 NVIDIA Corporation.  All rights reserved.+ *+ * NVIDIA Corporation and its licensors retain all intellectual property and + * proprietary rights in and to this software and related documentation. + * Any use, reproduction, disclosure, or distribution of this software + * and related documentation without an express license agreement from+ * NVIDIA Corporation is strictly prohibited.+ *+ * Please refer to the applicable NVIDIA end user license agreement (EULA) + * associated with this source code for terms and conditions that govern + * your use of this NVIDIA software.+ * + */++#ifndef _MATRIXMUL_H_+#define _MATRIXMUL_H_++/* Thread block size */+#define BLOCK_SIZE 16++/* Matrix dimensions+ * (chosen as multiples of the thread block size for simplicity)+ */+#define WA (3 * BLOCK_SIZE) /* Matrix A width  */+#define HA (5 * BLOCK_SIZE) /* Matrix A height */+#define WB (8 * BLOCK_SIZE) /* Matrix B width  */+#define HB WA  /* Matrix B height */+#define WC WB  /* Matrix C width  */+#define HC HA  /* Matrix C height */++#endif /* _MATRIXMUL_H_ */+
+ examples/src/scan/Makefile view
@@ -0,0 +1,18 @@+#+# Baking!+#++# ------------------------------------------------------------------------------+# Input files+# ------------------------------------------------------------------------------+EXECUTABLE      := scan++HSMAIN          := Scan.chs+CUFILES         := scan.cu++EXTRALIBS       := stdc++++# ------------------------------------------------------------------------------+# Haskell/CUDA build system+# ------------------------------------------------------------------------------+include ../../common/common.mk
+ examples/src/scan/Scan.chs view
@@ -0,0 +1,112 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+--+-- Module    : Scan+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Apply a binary operator to an array similar to 'fold', but return a+-- successive list of values reduced from the left (or right).+--+--------------------------------------------------------------------------------++module Main where++#include "scan.h"++-- Friends+import C2HS                                     hiding (newArray)+import Time+import RandomVector++-- System+import Control.Monad+import Control.Exception+import qualified Foreign.CUDA as CUDA+++--------------------------------------------------------------------------------+-- Reference+--------------------------------------------------------------------------------++scanList :: (Num e, Storable e) => Vector e -> IO (Vector e)+scanList xs = do+  bnds    <- getBounds xs+  xs'     <- getElems  xs+  (t,zs') <- benchmark 100 (return (scanl1 (+) xs')) (return ())+  putStrLn $ "List: " ++ shows (fromInteger (timeIn millisecond t`div`100)::Float) " ms"+  newListArray bnds zs'+++scanArr :: (Num e, Storable e) => Vector e -> IO (Vector e)+scanArr xs = do+  bnds  <- getBounds xs+  zs    <- newArray_ bnds+  let idx = range bnds+  (t,_) <- benchmark 100 (foldM_ (k zs) 0 idx) (return ())+  putStrLn $ "Array: " ++ shows (fromInteger (timeIn millisecond t)/100::Float) " ms"+  return zs+  where+    k zs a i = do+      x <- readArray xs i+      let z = x+a+      writeArray zs i z+      return z+++--------------------------------------------------------------------------------+-- CUDA+--------------------------------------------------------------------------------++--+-- Include the time to copy the data to/from the storable array (significantly+-- faster than from a Haskell list)+--+scanCUDA :: Vector Float -> IO (Vector Float)+scanCUDA xs = do+  bnds  <- getBounds xs+  zs    <- newArray_ bnds+  let len = rangeSize bnds+  CUDA.allocaArray len $ \d_xs -> do+  CUDA.allocaArray len $ \d_zs -> do+  (t,_) <- flip (benchmark 100) CUDA.sync $ do+    withVector xs $ \p -> CUDA.pokeArray len p d_xs+    scanl1_plusf d_xs d_zs len+    withVector zs $ \p -> CUDA.peekArray len d_zs p+  putStrLn $ "CUDA: " ++ shows (fromInteger (timeIn millisecond t)/100::Float) " ms (with copy)"++  (t',_) <- benchmark 100 (scanl1_plusf d_xs d_zs len) CUDA.sync+  putStrLn $ "CUDA: " ++ shows (fromInteger (timeIn millisecond t')/100::Float) " ms (compute only)"++  return zs++{# fun unsafe scanl1_plusf+  { withDP* `CUDA.DevicePtr Float'+  , withDP* `CUDA.DevicePtr Float'+  ,         `Int'+  } -> `()' #}+  where+    withDP p a = CUDA.withDevicePtr p $ \p' -> a (castPtr p')+++--------------------------------------------------------------------------------+-- Main+--------------------------------------------------------------------------------++main :: IO ()+main = do+  dev   <- CUDA.get+  props <- CUDA.props dev+  putStrLn $ "Using device " ++ show dev ++ ": " ++ CUDA.deviceName props++  arr  <- randomArr (1,100000) :: IO (Vector Float)+  ref  <- scanList arr+  ref' <- scanArr  arr+  cuda <- scanCUDA arr++  return ()++  putStr   "== Validating: "+  verify ref ref' >>= \rv -> assert rv (return ())+  verify ref cuda >>= \rv -> putStrLn $ if rv then "Ok!" else "INVALID!"+
+ examples/src/scan/scan.cu view
@@ -0,0 +1,204 @@+/* -----------------------------------------------------------------------------+ *+ * Module    : Scan+ * Copyright : (c) 2009 Trevor L. McDonell+ * License   : BSD+ *+ * ---------------------------------------------------------------------------*/++#include "scan.h"++#include "utils.h"+#include "operator.h"+#include "cudpp/cudpp_globals.h"+#include "cudpp/scan_kernel.cu"+#include "cudpp/vector_kernel.cu"++template <typename T>+struct scan_plan+{+    T           **block_sums;+    size_t      num_levels;+};++static inline unsigned int+calc_num_blocks(unsigned int N)+{+    return max(1u, (unsigned int)ceil((double)N / (SCAN_ELTS_PER_THREAD * CTA_SIZE)));+}+++/*+ * This is the CPU-side workhorse of the scan operation, invoking the kernel on+ * each of the reduction blocks.+ */+template <class op, typename T, bool backward, bool exclusive>+static void+scan_recursive+(+    const T             *in,+    T                   *out,+    scan_plan<T>        *plan,+    int                 N,+    int                 level+)+{+    size_t      num_blocks = calc_num_blocks(N);+    bool        is_full    = N == num_blocks * SCAN_ELTS_PER_THREAD * CTA_SIZE;++    dim3        grid(num_blocks, 1, 1);+    dim3        block(CTA_SIZE, 1, 1);+    size_t      smem = sizeof(T) * CTA_SIZE * 2;++#define MULTIBLOCK      0x01+#define FULLBLOCK       0x04+    int traits = 0;+    if (num_blocks > 1) traits |= MULTIBLOCK;+    if (is_full)        traits |= FULLBLOCK;++    /*+     * Set up execution parameters, and execute the scan+     */+    switch (traits)+    {+    case 0:+        scan4+            < T, ScanTraits<T, op, backward, exclusive, false, false, false> >+            <<<grid, block, smem>>>(out, in, NULL, N, 1, 1);+        break;++    case MULTIBLOCK:+        scan4+            < T, ScanTraits<T, op, backward, exclusive, false, true, false> >+            <<<grid, block, smem>>>(out, in, plan->block_sums[level], N, 1, 1);+        break;++    case FULLBLOCK:+        scan4+            < T, ScanTraits<T, op, backward, exclusive, false, false, true> >+            <<<grid, block, smem>>>(out, in, NULL, N, 1, 1);+        break;++    case MULTIBLOCK | FULLBLOCK:+        scan4+            < T, ScanTraits<T, op, backward, exclusive, false, true, true> >+            <<<grid, block, smem>>>(out, in, plan->block_sums[level], N, 1, 1);+        break;++    default:+        assert(!"Non-exhaustive patterns in match");+    }++    /*+     * After scanning the sub-blocks, we now need to combine those results by+     * taking the last value from each sub-block, and adding that to each of the+     * successive blocks (i.e. scan across the sub-computations)+     */+    if (num_blocks > 1)+    {+        T *sums = plan->block_sums[level];++        scan_recursive+            <op, T, backward, true>+            (sums, sums, plan, num_blocks, level+1);++        vectorAddUniform4+            <T, op, SCAN_ELTS_PER_THREAD>+            <<<grid,block>>>+            (out, sums, N, 4, 4, 0, 0);+    }++#undef MULTIBLOCK+#undef FULLBLOCK+}+++/*+ * Allocate temporary memory used by the scan.+ */+template <typename T>+static void+scan_init(int N, scan_plan<T> *plan)+{+    size_t level        = 0;+    size_t elements     = N;+    size_t num_blocks;++    /*+     * Determine how many intermediate block-level summations will be required+     */+    for (elements = N; elements > 1; elements = num_blocks)+    {+        num_blocks = calc_num_blocks(elements);++        if (num_blocks > 1)+            ++level;+    }++    plan->block_sums = (T**) malloc(level * sizeof(T*));+    plan->num_levels = level;++    /*+     * Now, allocate the necessary storage at each level+     */+    for (elements = N, level = 0; elements > 1; elements = num_blocks, level++)+    {+        num_blocks = calc_num_blocks(elements);++        if (num_blocks > 1)+            cudaMalloc((void**) &plan->block_sums[level], num_blocks * sizeof(T));+    }+}+++/*+ * Clean up temporary memory used by the scan+ */+template <typename T>+static void+scan_finalise(scan_plan<T> *p)+{+    for (size_t l = 0; l < p->num_levels; ++l)+        cudaFree(p->block_sums[l]);++    free(p->block_sums);+}+++/*+ * Apply a binary operator to an array similar to `fold', but return a+ * successive list of values reduced from the left. The reduction will take+ * place in parallel, so the operator must be associative.+ */+template <class op, typename T, bool backward, bool exclusive>+void+scan+(+    const T     *in,+    T           *out,+    int         length+)+{+    scan_plan<T> plan;+    scan_init<T>(length, &plan);++    scan_recursive<op, T, backward, exclusive>(in, out, &plan, length, 0);++    scan_finalise<T>(&plan);+}+++// -----------------------------------------------------------------------------+// Instances+// -----------------------------------------------------------------------------++void scanl_plusf(float *in, float *out, int N)+{+    scan< Plus<float>, float, false, true >(in, out, N);+}++void scanl1_plusf(float *in, float *out, int N)+{+    scan< Plus<float>, float, false, false >(in, out, N);+}+
+ examples/src/scan/scan.h view
@@ -0,0 +1,26 @@+/* -----------------------------------------------------------------------------+ *+ * Module    : Scan+ * Copyright : (c) 2009 Trevor L. McDonell+ * License   : BSD+ *+ * ---------------------------------------------------------------------------*/++#ifndef __SCAN_H__+#define __SCAN_H__++#ifdef __cplusplus+extern "C" {+#endif++/*+ * Instances+ */+void scanl_plusf(float *in, float *out, int N);+void scanl1_plusf(float *in, float *out, int N);+++#ifdef __cplusplus+}+#endif+#endif
+ examples/src/smvm/Makefile view
@@ -0,0 +1,20 @@+#+# Baking!+#++# ------------------------------------------------------------------------------+# Input files+# ------------------------------------------------------------------------------+EXECUTABLE      := smvm++HSMAIN          := SMVM.chs+CUFILES         := smvm-csr.cu \+                   smvm-cudpp.cu++USECUDPP        := 1+EXTRALIBS       := stdc++++# ------------------------------------------------------------------------------+# Haskell/CUDA build system+# ------------------------------------------------------------------------------+include ../../common/common.mk
+ examples/src/smvm/SMVM.chs view
@@ -0,0 +1,201 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+--+-- Module    : SMVM+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Sparse-matrix dense-vector multiplication+--+--------------------------------------------------------------------------------++module Main where++#include "smvm.h"++-- Friends+import Time+import C2HS+import RandomVector                             (randomList,randomListR,verifyList)++-- System+import Numeric+import Data.List+import Control.Monad+import Control.Applicative+import System.Random+import Foreign.CUDA                             (withDevicePtr)+import qualified Foreign.CUDA as CUDA++--+-- A very simple sparse-matrix / vector representation+-- (confusingly, different from that in RandomVector and used elsewhere)+--+type Vector e       = [e]+type SparseVector e = [(Int,e)]+type SparseMatrix e = [SparseVector e]++--------------------------------------------------------------------------------+-- Reference+--------------------------------------------------------------------------------++smvm :: Num e => SparseMatrix e -> Vector e -> Vector e+smvm sm v = [ sum [ x * (v!!col) | (col,x) <- sv ]  | sv <- sm ]++--------------------------------------------------------------------------------+-- CUDA+--------------------------------------------------------------------------------++--+-- Sparse-matrix vector multiplication, using compressed-sparse row format.+--+-- Lots of boilerplate to copy data to the device. Our simple list+-- representation has atrocious copy performance (see the `bandwidthTest'+-- example), so don't include that in the benchmarking+--+smvm_csr :: SparseMatrix Float -> Vector Float -> IO (Float, Vector Float)+smvm_csr sm v =+  let matData = concatMap (map cFloatConv . snd . unzip) sm+      colIdx  = concatMap (map cIntConv   . fst . unzip) sm+      rowPtr  = scanl (+) 0 (map (cIntConv . length) sm)+      v'      = map cFloatConv v+#ifdef __DEVICE_EMULATION__+      iters   = 1+#else+      iters   = 100+#endif+  in+  CUDA.withListArray    matData  $ \d_data       ->+  CUDA.withListArray    rowPtr   $ \d_ptr        ->+  CUDA.withListArray    colIdx   $ \d_indices    ->+  CUDA.withListArrayLen v'       $ \num_rows d_x ->+  CUDA.allocaArray      num_rows $ \d_y          -> do+    (t,_) <- benchmark iters (smvm_csr_f d_y d_x d_data d_ptr d_indices num_rows) CUDA.sync+    y     <- map cFloatConv <$> CUDA.peekListArray num_rows d_y+    return (fromInteger (timeIn millisecond t) / fromIntegral iters, y)+++{# fun unsafe smvm_csr_f+  { withDevicePtr* `CUDA.DevicePtr CFloat'+  , withDevicePtr* `CUDA.DevicePtr CFloat'+  , withDevicePtr* `CUDA.DevicePtr CFloat'+  , withDevicePtr* `CUDA.DevicePtr CUInt'+  , withDevicePtr* `CUDA.DevicePtr CUInt'+  ,                `Int'                   } -> `()' #}+++--+-- Sparse-matrix vector multiplication from CUDPP+--+smvm_cudpp :: SparseMatrix Float -> Vector Float -> IO (Float, Vector Float)+smvm_cudpp sm v =+  let matData = concatMap (map cFloatConv . snd . unzip) sm+      colIdx  = concatMap (map cIntConv   . fst . unzip) sm+      rowPtr  = scanl (+) 0 (map (cIntConv . length) sm)+      v'      = map cFloatConv v+#ifdef __DEVICE_EMULATION__+      iters   = 1+#else+      iters   = 100+#endif+  in+  CUDA.withListArrayLen v'       $ \num_rows     d_x    ->+  CUDA.allocaArray      num_rows $ \d_y                 ->+  withArrayLen          matData  $ \num_nonzeros h_data ->+  withArray             rowPtr   $ \h_rowPtr            ->+  withArray             colIdx   $ \h_colIdx            -> do+    (t,_) <- benchmark iters (smvm_cudpp_f d_y d_x h_data h_rowPtr h_colIdx num_rows num_nonzeros) CUDA.sync+    y     <- map cFloatConv <$> CUDA.peekListArray num_rows d_y+    return (fromInteger (timeIn millisecond t) / fromIntegral iters, y)++{# fun unsafe smvm_cudpp_f+  { withDevicePtr* `CUDA.DevicePtr CFloat'+  , withDevicePtr* `CUDA.DevicePtr CFloat'+  , id             `Ptr CFloat'+  , id             `Ptr CUInt'+  , id             `Ptr CUInt'+  ,                `Int'+  ,                `Int'                   } -> `()' #}+++--------------------------------------------------------------------------------+-- Main+--------------------------------------------------------------------------------++--+-- Generate random matrices+--+sparseMat :: (Num e, Random e, Storable e) => (Int,Int) -> (Int,Int) -> IO (SparseMatrix e)+sparseMat bnds (h,w) = replicateM h sparseVec+  where+    sparseVec = do+      nz  <- randomRIO bnds                         -- number of non-zero elements+      idx <- nub . sort <$> randomListR nz (0,w-1)  -- remove duplicate column indices+      zip idx <$> randomList (length idx)           -- (column indices don't actually need to be sorted)++denseMat :: (Num e, Random e, Storable e) => (Int,Int) -> IO (SparseMatrix e)+denseMat (h,w) = replicateM h (zip [0..] <$> randomList w)++--+-- Some test-harness utilities+--+stats :: (Floating a, Ord a) => [a] -> (a,a,a,a,a,a)+stats (x:xs) = finish . foldl' stats' (x,x,x,x*x,1) $ xs+  where+    stats' (mn,mx,s,ss,n) v = (min v mn, max v mx, s+v, ss+v*v, n+1)+    finish (mn,mx,s,ss,n)   = (mn, mx, av, var, stdev, n)+      where av    = s/n+            var   = (1/(n-1))*ss - (n/(n-1))*av*av+            stdev = sqrt var+++testAlgorithm :: (Num e, Ord e, Floating e)+             => String                                                  -- name of the algorithm+             -> (SparseMatrix e -> Vector e -> IO (Float, Vector e))    -- return time (ms), and result+             -> SparseMatrix e                                          -- input matrix+             -> Vector e                                                -- input vector+             -> Vector e                                                -- reference solution+             -> IO ()+testAlgorithm name f m v ref = do+    putStr name+    (t,y) <- f m v+    putStr   $ if verifyList ref y then "Ok!      " else "INVALID! "+    putStr   $ "( " ++ showFFloat (Just 2) t " ms, "+    putStrLn $ showFFloat (Just 2) (fromIntegral (2 * 1000 * sum (map length m)) / (t * 1E9)) " GFLOPS )"+++testMatrix :: String -> SparseMatrix Float -> Vector Float -> IO ()+testMatrix name sm v = do+  let w                  = length v+      (_,_,av,_,stdev,h) = stats (map (fromIntegral . length) sm)+      ref                = smvm sm v++  putStr   $ name ++ ": " ++ show w ++ "x" ++ show (round h)+  putStr   $ ", " ++ shows (round (av*h)) " non-zero elements "+  putStrLn $ "( " ++ showFFloat (Just 2) av " +/- " ++ showFFloat (Just 2) stdev " )"++  testAlgorithm "  smvm-csr:     " smvm_csr   sm v ref+  testAlgorithm "  smvm-cudpp:   " smvm_cudpp sm v ref+  putStrLn ""+++--+-- Finally, the main function+--+main :: IO ()+main = do+  dev   <- CUDA.get+  props <- CUDA.props dev+  putStrLn $ "Using device " ++ show dev ++ ": \"" ++ CUDA.deviceName props ++ "\""+  putStrLn $ "  Compute capability:  " ++ show (CUDA.computeCapability props)+  putStrLn $ "  Total global memory: " +++    showFFloat (Just 2) (fromIntegral (CUDA.totalGlobalMem props) / (1024*1024) :: Double) " GB\n"++  v1 <- randomList 512+  v2 <- randomList 2048+  m1 <- denseMat  (512,512)+  m2 <- sparseMat (20,200) (20 * 2048,2048)++  testMatrix "Dense Matrix"  m1 v1+  testMatrix "Sparse Matrix" m2 v2+
+ examples/src/smvm/smvm-csr.cu view
@@ -0,0 +1,246 @@+/*+ *  Copyright 2008-2009 NVIDIA Corporation+ *+ *  Licensed under the Apache License, Version 2.0 (the "License");+ *  you may not use this file except in compliance with the License.+ *  You may obtain a copy of the License at+ *+ *      http://www.apache.org/licenses/LICENSE-2.0+ *+ *  Unless required by applicable law or agreed to in writing, software+ *  distributed under the License is distributed on an "AS IS" BASIS,+ *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.+ *  See the License for the specific language governing permissions and+ *  limitations under the License.+ */+++#include "smvm.h"+#include "utils.h"+#include "texture.h"+++/* -----------------------------------------------------------------------------+ * Sparse-matrix dense-vector multiplication, compressed-sparse row format+ * -----------------------------------------------------------------------------+ *+ * Each row of the CSR matrix is assigned to a warp, which computes the dot+ * product of the i-th row of A with the x vector, in parallel.+ *+ *   y[i] = A[i,:] * x+ *+ * This division of work implies that the CSR index and data arrays (Aj and Ax)+ * are accessed in a contiguous manner (but generally not aligned). On the GT200+ * these accesses are coalesced, unlike kernels based on the one-row-per-thread+ * division of work. Since an entire 32-thread warp is assigned to each row,+ * many threads will remain idle when their row contains a small number of+ * elements. This code relies on implicit synchronization among threads in a+ * warp.+ *+ * Optionally, the texture cache may be used for accessing the x vector. This+ * generally shows good improvements.+ *+ * References:+ *+ *  [1] N. Bell and M. Garland. "Efficient sparse matrix-vector multiplication on+ *      CUDA." NVIDIA Technical Report NVR-2008-004, NVIDIA Corporation, Dec. 2008.+ *+ *  [2] N. Bell and M. Garland. "Implementing sparse matrix-vector+ *      multiplication on throughput-oriented processors." In Supercomputing+ *      `09: Proceedings of the 2009 Conference on High Performance Computing+ *      Networking, Storage and Analysis, pages 1-11, 2009.+ */+template <unsigned int BlockSize, typename T, bool UseCache>+__global__ static void+smvm_k+(+    T                   *d_y,+    const T             *d_x,+    const T             *d_Ax,+    const unsigned int  *d_Ap,+    const unsigned int  *d_Aj,+    const unsigned int  num_rows+)+{+    /*+     * Require at least a full warp for each row. This could be relaxed by+     * modifying the cooperative reduction step+     */+    assert(BlockSize % WARP_SIZE == 0);++    const unsigned int vectorsPerBlock = BlockSize / WARP_SIZE;+    const unsigned int num_vectors     = vectorsPerBlock * gridDim.x;+    const unsigned int thread_id       = BlockSize * blockIdx.x + threadIdx.x;+    const unsigned int vector_id       = thread_id / WARP_SIZE;+    const unsigned int thread_lane     = threadIdx.x & (WARP_SIZE-1);+    const unsigned int vector_lane     = threadIdx.x / WARP_SIZE;++    __shared__ volatile T            s_data[(vectorsPerBlock+1) * WARP_SIZE];+    __shared__ volatile unsigned int s_ptrs[vectorsPerBlock][2];++    for (unsigned int row = vector_id; row < num_rows; row += num_vectors)+    {+        /*+         * Use two threads to fetch the indices of the start and end of this+         * segment. This is a single coalesced (although unaligned) global read+         * rather than two, and hence considerably faster.+         */+        if (thread_lane < 2)+            s_ptrs[vector_lane][thread_lane] = d_Ap[row + thread_lane];++        __EMUSYNC;+        const unsigned int row_start = s_ptrs[vector_lane][0];+        const unsigned int row_end   = s_ptrs[vector_lane][1];++        /*+         * Have the threads read in all values for this row, accumulating local+         * dot-product sums. Then, reduce this cooperatively in shared memory.+         */+        T sum = 0;+        for (unsigned int j = row_start + thread_lane; j < row_end; j += WARP_SIZE)+            sum += d_Ax[j] * fetch_x<UseCache>(d_Aj[j], d_x);++        s_data[threadIdx.x] = sum;                                  __EMUSYNC;+        s_data[threadIdx.x] = sum = sum + s_data[threadIdx.x + 16]; __EMUSYNC;+        s_data[threadIdx.x] = sum = sum + s_data[threadIdx.x +  8]; __EMUSYNC;+        s_data[threadIdx.x] = sum = sum + s_data[threadIdx.x +  4]; __EMUSYNC;+        s_data[threadIdx.x] = sum = sum + s_data[threadIdx.x +  2]; __EMUSYNC;+        s_data[threadIdx.x] = sum = sum + s_data[threadIdx.x +  1]; __EMUSYNC;++#if 0+        /*+         * Alternative method (slightly slower, due to bank conflicts?)+         */+        s_data[threadIdx.x] += s_data[threadIdx.x + 16];+        s_data[threadIdx.x] += s_data[threadIdx.x +  8];+        s_data[threadIdx.x] += s_data[threadIdx.x +  4];+        s_data[threadIdx.x] += s_data[threadIdx.x +  2];+        s_data[threadIdx.x] += s_data[threadIdx.x +  1];+#endif++        /*+         * Finally, first thread writes the result for this row+         */+        if (thread_lane == 0)+            d_y[row] = s_data[threadIdx.x];+    }+}+++template <typename T, bool UseCache>+static void+smvm_dispatch+(+    T                   *d_y,+    const T             *d_x,+    const T             *d_data,+    const unsigned int  *d_ptr,+    const unsigned int  *d_indices,+    const unsigned int  num_rows,+    const unsigned int  blocks,+    const unsigned int  threads+)+{+    const unsigned int smem = 0;++    switch (threads)+    {+    case 512: smvm_k<512,T,UseCache><<<blocks,threads,smem>>>(d_y, d_x, d_data, d_ptr, d_indices, num_rows); break;+    case 256: smvm_k<256,T,UseCache><<<blocks,threads,smem>>>(d_y, d_x, d_data, d_ptr, d_indices, num_rows); break;+    case 128: smvm_k<128,T,UseCache><<<blocks,threads,smem>>>(d_y, d_x, d_data, d_ptr, d_indices, num_rows); break;+    case  64: smvm_k< 64,T,UseCache><<<blocks,threads,smem>>>(d_y, d_x, d_data, d_ptr, d_indices, num_rows); break;+    case  32: smvm_k< 32,T,UseCache><<<blocks,threads,smem>>>(d_y, d_x, d_data, d_ptr, d_indices, num_rows); break;+    default:+        assert(!"Non-exhaustive patterns in match");+    }+}+++/*+ * Select an "optimal" number of threads and blocks for the problem size. This+ * is an act of balancing resource usage: shared memory, registers, in-flight+ * threads and blocks per multiprocessor. Ultimately, this requires some+ * experimentation for every kernel, device and problem set, but we choose some+ * sensible default values.+ *+ * Additionally, each block will have at least one full warp, as required by the+ * core kernel.+ */+static void+smvm_control+(+    unsigned int        n,+    unsigned int        &blocks,+    unsigned int        &threads,+    unsigned int        maxThreads = MAX_THREADS,+    unsigned int        maxBlocks  = MAX_BLOCKS++)+{+    threads = (n < maxThreads) ? max(WARP_SIZE, ceilPow2(n)) : maxThreads;+    blocks  = (n + threads - 1) / threads;+    blocks  = min(blocks, maxBlocks);+}+++/*+ * Sparse matrix multiplication:+ *   y = A * x+ *+ * The CSR format explicitly stores column indices (indices) and non-zero values+ * (data) in row-major order, together with a third array of row pointers (ptr).+ * For an M-by-N matrix, ptr has length (M+1) and stores the offset to the start+ * of the i-th row in ptr[i]. The last entry then, corresponding to the (M+1)-st+ * row, contains the number of non-zero elements in the matrix.+ *+ * Example:+ *            | 1 7 0 0 |+ *        A = | 0 2 8 0 |+ *            | 5 0 3 9 |+ *            | 0 6 0 4 |+ *+ *      ptr = [ 0 2 4 7 9 ]+ *  indices = [ 0 1 1 2 0 2 3 1 3 ]+ *     data = [ 1 7 2 8 5 3 9 6 4 ]+ *+ * d_y          The output vector+ * d_x          The input (dense) vector to multiply against+ * d_data       The non-zero elements of the sparse, stored row-major order+ * d_ptr        Row offsets+ * d_indices    Column indices+ */+template <typename T, bool UseCache>+static void+smvm_csr+(+    T                   *d_y,+    const T             *d_x,+    const T             *d_data,+    const unsigned int  *d_ptr,+    const unsigned int  *d_indices,+    const unsigned int  num_rows+)+{+    unsigned int blocks;+    unsigned int threads;++    if (UseCache)+        bind_x(d_x);++    smvm_control(num_rows, blocks, threads);+    smvm_dispatch<T,UseCache>(d_y, d_x, d_data, d_ptr, d_indices, num_rows, blocks, threads);++    if (UseCache)+        unbind_x(d_x);+}++/* -----------------------------------------------------------------------------+ * Instances+ * ---------------------------------------------------------------------------*/++void+smvm_csr_f(float *d_y, float *d_x, float *d_data, unsigned int *d_rowPtr, unsigned int *d_colIdx, unsigned int num_rows)+{+    smvm_csr<float,true>(d_y, d_x, d_data, d_rowPtr, d_colIdx, num_rows);+}+
+ examples/src/smvm/smvm-cudpp.cu view
@@ -0,0 +1,56 @@+/* -----------------------------------------------------------------------------+ *+ * Module    : SMVM+ * Copyright : (c) 2009 Trevor L. McDonell+ * License   : BSD+ *+ * ---------------------------------------------------------------------------*/+++#include "smvm.h"+#include <cudpp.h>++template <typename T> CUDPPDatatype getType();+template <> CUDPPDatatype getType<float>() { return CUDPP_FLOAT; }+template <> CUDPPDatatype getType<unsigned int>() { return CUDPP_UINT; }+++/*+ * Sparse matrix-dense vector multiply. Hook directly into the CUDPP+ * implementation.+ */+template <typename T>+void smvm_cudpp+(+    float               *d_y,+    const float         *d_x,+    const float         *h_data,+    const unsigned int  *h_rowPtr,+    const unsigned int  *h_colIdx,+    const unsigned int  num_rows,+    const unsigned int  num_nonzeros+)+{+    CUDPPConfiguration cp;+    CUDPPHandle        sm;++    cp.datatype  = getType<T>();+    cp.options   = 0;+    cp.algorithm = CUDPP_SPMVMULT;++    cudppSparseMatrix(&sm, cp, num_nonzeros, num_rows, h_data, h_rowPtr, h_colIdx);+    cudppSparseMatrixVectorMultiply(sm, d_y, d_x);++    cudppDestroySparseMatrix(sm);+}+++// -----------------------------------------------------------------------------+// Instances+// -----------------------------------------------------------------------------++void smvm_cudpp_f(float *d_y, float *d_x, float *h_data, unsigned int *h_rowPtr, unsigned int *h_colIdx, unsigned int num_rows, unsigned int num_nonzeros)+{+    smvm_cudpp<float>(d_y, d_x, h_data, h_rowPtr, h_colIdx, num_rows, num_nonzeros);+}+
+ examples/src/smvm/smvm.h view
@@ -0,0 +1,33 @@+/* -----------------------------------------------------------------------------+ *+ * Module    : SMVM+ * Copyright : (c) 2009 Trevor L. McDonell+ * License   : BSD+ *+ * ---------------------------------------------------------------------------*/++#ifndef __SMVM_H__+#define __SMVM_H__++/*+ * Optimised for Tesla C1060 (compute 1.3)+ * Maximum performance for your card may be achieved with different values.+ *+ * http://developer.download.nvidia.com/compute/cuda/CUDA_Occupancy_calculator.xls+ */+#define MAX_THREADS             128+#define MAX_BLOCKS_PER_SM       8+#define MAX_BLOCKS              (MAX_BLOCKS_PER_SM * 30)+#define WARP_SIZE               32++#ifdef __cplusplus+extern "C" {+#endif++void smvm_csr_f(float *d_y, float *d_x, float *d_data, unsigned int *d_rowPtr, unsigned int *d_colIdx, unsigned int num_rows);+void smvm_cudpp_f(float *d_y, float *d_x, float *h_data, unsigned int *h_rowPtr, unsigned int *h_colIdx, unsigned int num_rows, unsigned int num_nonzeros);++#ifdef __cplusplus+}+#endif+#endif
+ examples/src/smvm/texture.h view
@@ -0,0 +1,94 @@+/*+ *  Copyright 2008-2009 NVIDIA Corporation+ *+ *  Licensed under the Apache License, Version 2.0 (the "License");+ *  you may not use this file except in compliance with the License.+ *  You may obtain a copy of the License at+ *+ *      http://www.apache.org/licenses/LICENSE-2.0+ *+ *  Unless required by applicable law or agreed to in writing, software+ *  distributed under the License is distributed on an "AS IS" BASIS,+ *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.+ *  See the License for the specific language governing permissions and+ *  limitations under the License.+ */+++#ifndef __TEXTURE_H__+#define __TEXTURE_H__++#include "utils.h"+#include <cuda_runtime_api.h>++/*+ * These textures are (optionally) used to cache the 'x' vector in y += A*x+ * Use int2 to pull doubles through texture cache.+ */+texture<float,1> tex_x_float;+texture<int2,1>  tex_x_double;++inline void+bind_x(const float * x)+{+    size_t offset = size_t(-1);++    CUDA_SAFE_CALL(cudaBindTexture(&offset, tex_x_float, x));+    if (offset != 0)+        assert(!"memory is not aligned, refusing to use texture cache");+}++inline void+bind_x(const double * x)+{+    size_t offset = size_t(-1);++    CUDA_SAFE_CALL(cudaBindTexture(&offset, tex_x_double, x));+    if (offset != 0)+        assert(!"memory is not aligned, refusing to use texture cache");+}++/*+ * NOTE: the parameter is unused only to distinguish the two unbind functions+ */+inline void+unbind_x(const float *)+{+    CUDA_SAFE_CALL(cudaUnbindTexture(tex_x_float));+}++inline void+unbind_x(const double *)+{+    CUDA_SAFE_CALL(cudaUnbindTexture(tex_x_double));+}++template <bool UseCache>+__inline__ __device__ float+fetch_x(const int& i, const float * x)+{+    if (UseCache) return tex1Dfetch(tex_x_float, i);+    else          return x[i];+}++#ifndef CUDA_NO_SM_13_DOUBLE_INTRINSICS+template <bool UseCache>+__inline__ __device__ double fetch_x(const int& i, const double * x)+{+#if __CUDA_ARCH__ < 130+#error "double precision require Compute Compatibility 1.3 or greater"+#endif+    if (UseCache)+    {+        int2 v = tex1Dfetch(tex_x_double, i);+        return __hiloint2double(v.y, v.x);+    }+    else+    {+        return x[i];+    }+}+#endif++#endif  // __TEXTURE_H__+
+ examples/src/sort/Makefile view
@@ -0,0 +1,19 @@+#+# Baking!+#++# ------------------------------------------------------------------------------+# Input files+# ------------------------------------------------------------------------------+EXECUTABLE      := sort++HSMAIN          := Sort.chs+CUFILES         := radix_sort.cu++USECUDPP        := 1+EXTRALIBS       := stdc++++# ------------------------------------------------------------------------------+# Haskell/CUDA build system+# ------------------------------------------------------------------------------+include ../../common/common.mk
+ examples/src/sort/Sort.chs view
@@ -0,0 +1,98 @@+{-# LANGUAGE ForeignFunctionInterface #-}+--------------------------------------------------------------------------------+--+-- Module    : Sort+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Reduce a vector of (key,value) pairs+--+--------------------------------------------------------------------------------++module Main where++#include "sort.h"++import C2HS+import RandomVector++import Data.Ord+import Data.List+import Control.Monad+import qualified Foreign.CUDA as C++++--------------------------------------------------------------------------------+-- CUDA++test_f :: (Storable a, Eq a) => [(Float,a)] -> IO Bool+test_f kv =+  let l     = length kv+      (k,v) = unzip kv+  in+  C.withListArray k $ \d_k ->+  C.withListArray v $ \d_v -> do++    sort_f d_k d_v (length kv)+    res <- liftM2 zip (C.peekListArray l d_k) (C.peekListArray l d_v)+    return (res == sortBy (comparing fst) kv)+++test_i :: (Storable a, Eq a) => [(Int,a)] -> IO Bool+test_i kv =+  let l     = length kv+      (k,v) = unzip kv+  in+  C.withListArray k $ \d_k ->+  C.withListArray v $ \d_v -> do++    sort_ui d_k d_v (length kv)+    res <- liftM2 zip (C.peekListArray l d_k) (C.peekListArray l d_v)+    return (res == sortBy (comparing fst) kv)+++{# fun unsafe sort_f+  { withDP* `C.DevicePtr Float'+  , withDP* `C.DevicePtr a'+  ,         `Int'+  } -> `()' #}+  where+    withDP p a = C.withDevicePtr p $ \p' -> a (castPtr p')++{# fun unsafe sort_ui+  { withDP* `C.DevicePtr Int'+  , withDP* `C.DevicePtr a'+  ,         `Int'+  } -> `()' #}+  where+    withDP p a = C.withDevicePtr p $ \p' -> a (castPtr p')++--+-- I don't need to learn template haskell or quick check... nah, not at all...+--+main :: IO ()+main = do+  f <- randomList  10000+  i <- randomListR 10000 (0,1000)++  putStr "Test (float,int): "+  test_f (zip f i) >>= \r -> case r of+    True -> putStrLn "Ok!"+    _    -> putStrLn "INVALID!"++  putStr "Test (float,float): "+  test_f (zip f f) >>= \r -> case r of+    True -> putStrLn "Ok!"+    _    -> putStrLn "INVALID!"++  putStr "Test (int,int): "+  test_i (zip i i) >>= \r -> case r of+    True -> putStrLn "Ok!"+    _    -> putStrLn "INVALID!"++  putStr "Test (int,float): "+  test_i (zip i f) >>= \r -> case r of+    True -> putStrLn "Ok!"+    _    -> putStrLn "INVALID!"+
+ examples/src/sort/radix_sort.cu view
@@ -0,0 +1,61 @@+/* -----------------------------------------------------------------------------+ *+ * Module    : Sort+ * Copyright : (c) 2009 Trevor L. McDonell+ * License   : BSD+ *+ *----------------------------------------------------------------------------*/++#include <cudpp.h>++#include "sort.h"+++template <typename T> CUDPPDatatype static getType();+template <> CUDPPDatatype getType<float>()        { return CUDPP_FLOAT; }+template <> CUDPPDatatype getType<unsigned int>() { return CUDPP_UINT;  }+++/*+ * In-place radix sort of values or key-value pairs. Values can be any 32-bit+ * type, as their payload is never inspected or manipulated.+ */+template <typename T>+static void+radix_sort+(+    unsigned int        length,+    T                   *d_keys,+    void                *d_vals = NULL,+    int                 bits    = 8 * sizeof(T)+)+{+    CUDPPHandle         plan;+    CUDPPConfiguration  cp;++    cp.datatype  = getType<T>();+    cp.algorithm = CUDPP_SORT_RADIX;+    cp.options   = (d_vals != NULL) ? CUDPP_OPTION_KEY_VALUE_PAIRS+                                    : CUDPP_OPTION_KEYS_ONLY;++    cudppPlan(&plan, cp, length, 1, 0);+    cudppSort(plan, d_keys, d_vals, bits, length);++    cudppDestroyPlan(plan);+}+++/* -----------------------------------------------------------------------------+ * Instances+ * ---------------------------------------------------------------------------*/++void sort_f(float *d_keys, void *d_vals, unsigned int length)+{+    radix_sort<float>(length, d_keys, d_vals);+}++void sort_ui(unsigned int *d_keys, void *d_vals, unsigned int length)+{+    radix_sort<unsigned int>(length, d_keys, d_vals);+}+
+ examples/src/sort/sort.h view
@@ -0,0 +1,23 @@+/* -----------------------------------------------------------------------------+ *+ * Module    : Sort+ * Copyright : (c) 2009 Trevor L. McDonell+ * License   : BSD+ *+ * ---------------------------------------------------------------------------*/++#ifndef __SORT_PRIV_H__+#define __SORT_PRIV_H__++#ifdef __cplusplus+extern "C" {+#endif++void sort_f(float *d_keys, void *d_vals, unsigned int length);+void sort_ui(unsigned int *d_keys, void *d_vals, unsigned int length);++#ifdef __cplusplus+}+#endif+#endif+
+ examples/src/vectorAddDrv/Makefile view
@@ -0,0 +1,18 @@+#+# Baking!+#++# ------------------------------------------------------------------------------+# Input files+# ------------------------------------------------------------------------------+EXECUTABLE	:= vectorAddDrv++HSMAIN		:= VectorAdd.hs+PTXFILES	:= vector_add.cu++USEDRVAPI	:= 1++# ------------------------------------------------------------------------------+# Haskell/CUDA build system+# ------------------------------------------------------------------------------+include ../../common/common.mk
+ examples/src/vectorAddDrv/VectorAdd.hs view
@@ -0,0 +1,121 @@+--------------------------------------------------------------------------------+--+-- Module    : VectorAdd+-- Copyright : (c) 2009 Trevor L. McDonell+-- License   : BSD+--+-- Element-wise addition of two vectors+--+--------------------------------------------------------------------------------++module Main where++-- Friends+import RandomVector++-- System+import Numeric+import Control.Monad+import Control.Exception+import Data.Array.Storable+import qualified Data.ByteString.Char8 as B+import qualified Foreign.CUDA.Driver   as CUDA+++--------------------------------------------------------------------------------+-- Reference implementation+--------------------------------------------------------------------------------++testRef :: (Num e, Storable e) => Vector e -> Vector e -> IO (Vector e)+testRef xs ys = do+  (i,j) <- getBounds xs+  res   <- newArray_ (i,j)+  forM_ [i..j] (add res)+  return res+  where+    add res idx = do+      a <- readArray xs idx+      b <- readArray ys idx+      writeArray res idx (a+b)++--------------------------------------------------------------------------------+-- CUDA+--------------------------------------------------------------------------------++--+-- Initialise the device and context. Load the PTX source code, and return a+-- reference to the kernel function.+--+initCUDA :: IO (CUDA.Context, CUDA.Fun)+initCUDA = do+  CUDA.initialise []+  dev     <- CUDA.device 0+  ctx     <- CUDA.create dev []+  ptx     <- B.readFile "data/vector_add.ptx"+  (mdl,r) <- CUDA.loadDataEx ptx [CUDA.MaxRegisters 32]+  fun     <- CUDA.getFun mdl "VecAdd"++  putStrLn $ ">> PTX JIT compilation (" ++ showFFloat (Just 2) (CUDA.jitTime r) " ms)"+  B.putStrLn (CUDA.jitInfoLog r)+  return (ctx,fun)+++--+-- Run the test+--+testCUDA :: (Num e, Storable e) => Vector e -> Vector e -> IO (Vector e)+testCUDA xs ys = do+  (m,n)   <- getBounds xs+  let len = (n-m+1)++  -- Initialise environment and copy over test data+  --+  putStrLn ">> Initialising"+  bracket initCUDA (\(ctx,_) -> CUDA.destroy ctx) $ \(_,addVec) -> do++  -- Allocate some device memory. This will be freed once the computation+  -- terminates, either normally or by exception.+  --+  putStrLn ">> Executing"+  CUDA.allocaArray len $ \dx -> do+  CUDA.allocaArray len $ \dy -> do+  CUDA.allocaArray len $ \dz -> do++  -- Copy over the data+  --+  withVector xs $ \p -> CUDA.pokeArray len p dx+  withVector ys $ \p -> CUDA.pokeArray len p dy++  -- Setup and execute the kernel (repeat test many times...)+  --+  CUDA.setParams     addVec [CUDA.VArg dx, CUDA.VArg dy, CUDA.VArg dz, CUDA.IArg len]+  CUDA.setBlockShape addVec (128,1,1)+  CUDA.launch        addVec ((len+128-1) `div` 128, 1) Nothing+  CUDA.sync++  -- Copy back result+  --+  zs <- newArray_ (m,n)+  withVector zs $ \p -> CUDA.peekArray len dz p+  return zs+++--------------------------------------------------------------------------------+-- Test & Verify+--------------------------------------------------------------------------------++main :: IO ()+main = do+  putStrLn "== Generating random vectors"+  xs  <- randomArr (1,10000) :: IO (Vector Float)+  ys  <- randomArr (1,10000) :: IO (Vector Float)++  putStrLn "== Generating reference solution"+  ref <- testRef  xs ys++  putStrLn "== Testing CUDA"+  arr <- testCUDA xs ys++  putStr   "== Validating: "+  verify ref arr >>= \rv -> putStrLn $ if rv then "Ok!" else "INVALID!"+
+ examples/src/vectorAddDrv/vector_add.cu view
@@ -0,0 +1,18 @@+/*+ * Name      : VectorAdd+ * Copyright : (c) 2009 Trevor L. McDonell+ * License   : BSD+ *+ * Element-wise addition of two (floating-point) vectors+ */+++extern "C"+__global__ void VecAdd(const float *xs, const float *ys, float *out, const unsigned int N)+{+    unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x;++    if (idx < N)+        out[idx] = xs[idx] + ys[idx];+}+