Commits

ppavel committed 896bcac

my patches

  • Participants

Comments (0)

Files changed (3)

+# HG changeset patch
+# Parent da363388b4fef80d3feda475025ea2bf40326ae1
+
+diff -r da363388b4fe Walkthrough.hs
+--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
++++ b/Walkthrough.hs	Mon Jan 03 14:57:36 2011 +0300
+@@ -0,0 +1,66 @@
++{-# OPTIONS_GHC -framework OpenCL #-}
++import Foreign.Ptr
++import Foreign.Storable
++import System.OpenCL.Raw.V10
++import Control.Monad.Trans
++import Control.Monad
++import Foreign.Marshal.Alloc
++
++
++newtype EitherT m l r = EitherT {runEitherT :: m (Either l r)}
++instance Monad m => Monad (EitherT m l) where
++    return = EitherT . return . Right
++    (EitherT me) >>= f = EitherT $ me >>=  \e -> 
++        case e of
++            Right a  -> runEitherT $ f a
++            Left  a  -> return $ Left a
++
++instance MonadIO m => MonadIO (EitherT m l) where
++    liftIO = EitherT . liftM Right . liftIO
++
++mbToEither Nothing = Right ()
++mbToEither (Just err) = Left err
++
++defaultPlatform = nullPtr
++main = do
++    res <- runEitherT $ do
++        devices <- EitherT $ clGetDeviceIDs nullPtr clDeviceTypeAll 100 
++        let callback _ _ _ _ = return ()
++        ctx <- EitherT $ clCreateContext [] devices callback nullPtr
++        let device = head devices
++        queue <- EitherT $ clCreateCommandQueue ctx device (CommandQueueProperties 0)
++        prog <- EitherT $ clCreateProgramWithSource ctx programSrc 
++        let progCB _ _ = return ()
++        EitherT $ liftM mbToEither $ clBuildProgram prog [device]  "" progCB nullPtr 
++        kern <- EitherT $ clCreateKernel prog "square"
++        inp   <- liftIO malloc
++        outp  <- liftIO malloc
++        inpB  <- EitherT $ clCreateBuffer ctx (MemFlags $ 8 + 4)  (fromIntegral $ sizeOf inp) (castPtr (inp :: Ptr Double))
++        outpB <- EitherT $ clCreateBuffer ctx (MemFlags $ 8 + 2) (fromIntegral $ sizeOf outp) (castPtr (outp :: Ptr Double))
++        inpBStorage <- liftIO malloc
++        outBStorage <- liftIO malloc
++        liftIO $ poke inpBStorage inpB >> poke outBStorage outpB
++        EitherT $ liftM mbToEither $             poke inp 2.3 >> clSetKernelArg kern 0 (fromIntegral $ sizeOf (undefined::Ptr ())) (castPtr inpBStorage)
++        EitherT $ liftM mbToEither $                             clSetKernelArg kern 1 (fromIntegral $ sizeOf (undefined::Ptr ())) (castPtr outBStorage)
++        EitherT $ liftM mbToEither $ alloca $ \p  -> poke p 1 >> clSetKernelArg kern 2 (fromIntegral $ sizeOf (undefined::CLint)) (castPtr (p :: Ptr CLint))
++        event <- EitherT $ clEnqueueTask queue kern []
++        EitherT $ liftM mbToEither $ clWaitForEvents [event]
++        liftIO $ peek outp
++
++
++
++    print res
++    
++
++programSrc :: String
++programSrc = "\
++\__kernel square(                                                        \n \
++\   __global double* input,                                              \n \
++\   __global double* output,                                             \n \
++\   const unsigned int count)                                           \n \
++\{                                                                      \n \
++\   int i = get_global_id(0);                                           \n \
++\   if(i < count)                                                       \n \
++\       output[i] = input[i] * input[i]    ;                                \n \
++\}                                                                      \n \
++\"
+# HG changeset patch
+# Parent 15813df45d4bfbfda27e676fe410239f3bda4b37
+
+diff -r 15813df45d4b System/OpenCL/Raw/V10/Context.hs
+--- a/System/OpenCL/Raw/V10/Context.hs	Sun Jan 02 23:46:31 2011 +0300
++++ b/System/OpenCL/Raw/V10/Context.hs	Mon Jan 03 14:56:35 2011 +0300
+@@ -32,7 +32,8 @@
+         pokeArray0 nullPtr propertiesP properties
+         pokeArray devicesP devices
+         fptr <- wrapCreateContextCallback pfn_notify
+-        wrapErrorEither $ raw_clCreateContext propertiesP (fromIntegral devicesN) devicesP fptr user_dat             
++        -- FIX
++        wrapErrorEither $ raw_clCreateContext nullPtr {- propertiesP -} (fromIntegral devicesN) devicesP fptr user_dat             
+     where propertiesN = length properties
+           devicesN = length devices
+           
+diff -r 15813df45d4b System/OpenCL/Raw/V10/Kernel.hs
+--- a/System/OpenCL/Raw/V10/Kernel.hs	Sun Jan 02 23:46:31 2011 +0300
++++ b/System/OpenCL/Raw/V10/Kernel.hs	Mon Jan 03 14:56:35 2011 +0300
+@@ -3,6 +3,7 @@
+ module System.OpenCL.Raw.V10.Kernel 
+     (clCreateKernel
+     ,clCreateKernelsInProgram
++    ,clSetKernelArg
+     ,clRetainKernel
+     ,clReleaseKernel
+     ,clGetKernelInfo
+@@ -22,7 +23,8 @@
+ 
+ 
+ foreign import ccall "clCreateKernel" raw_clCreateKernel :: Program -> CString -> Ptr CLint -> IO Kernel 
+-clCreateKernel program kernel_name = wrapErrorEither $ raw_clCreateKernel program kernel_name 
++clCreateKernel :: Program -> String -> IO (Either ErrorCode Kernel)
++clCreateKernel program kernel_name = withCString kernel_name $ \kname -> wrapErrorEither $ raw_clCreateKernel program kname 
+ 
+ foreign import ccall "clCreateKernelsInProgram" raw_clCreateKernelsInProgram :: Program -> CLuint -> Ptr Kernel -> Ptr CLuint -> IO CLint 
+ clCreateKernelsInProgram :: Program -> CLuint -> IO (Either ErrorCode [Kernel])
+@@ -76,14 +78,19 @@
+ foreign import ccall "clEnqueueTask" raw_clEnqueueTask :: CommandQueue -> Kernel -> CLuint -> Ptr Event -> Ptr Event -> IO CLint
+ clEnqueueTask :: CommandQueue -> Kernel -> [Event] -> IO (Either ErrorCode Event)
+ clEnqueueTask queue kernel event_wait_listL = 
+-    allocaArray num_events_in_wait_list $ \event_wait_list ->
+-    alloca $ \event -> do
+-        pokeArray event_wait_list event_wait_listL
+-        err <- wrapError $ raw_clEnqueueTask queue kernel (fromIntegral num_events_in_wait_list) event_wait_list event 
+-        if err == Nothing
+-            then Right <$> peek event
+-            else return $ Left . fromJust $ err
+-    where num_events_in_wait_list = length event_wait_listL
++    if num_events_in_wait_list /= 0 then
++        allocaArray num_events_in_wait_list $ \p -> pokeArray p event_wait_listL >> f p
++      else
++        f nullPtr
++ where
++    f = \event_wait_list ->
++        alloca $ \event -> do
++            pokeArray event_wait_list event_wait_listL
++            err <- wrapError $ raw_clEnqueueTask queue kernel (fromIntegral num_events_in_wait_list) event_wait_list event 
++            if err == Nothing
++                then Right <$> peek event
++                else return $ Left . fromJust $ err
++    num_events_in_wait_list = length event_wait_listL
+ 
+ type NKCallbackFunction = Ptr () -> IO ()
+ foreign import ccall "wrapper" wrapNativeKernelCallback :: NKCallbackFunction -> IO (FunPtr NKCallbackFunction)
+Walkthrough.hs
+fixes