module Foreign.OpenCL.Bindings.Kernel (
createKernel,
kernelContext, kernelFunctionName, kernelNumArgs,
kernelWorkGroupSize, kernelCompileWorkGroupSize, kernelLocalMemSize,
kernelPreferredWorkGroupSizeMultiple, kernelPrivateMemSize,
enqueueNDRangeKernel, enqueueTask,
KernelArg(..), setKernelArg, setKernelArgs
) where
import Control.Monad
import Foreign
import Foreign.C.String
import Foreign.C.Types
import Foreign.OpenCL.Bindings.Internal.Types
import Foreign.OpenCL.Bindings.Internal.Finalizers
import Foreign.OpenCL.Bindings.Internal.Error
import Foreign.OpenCL.Bindings.Internal.Util
import Foreign.OpenCL.Bindings.Internal.Logging as Log
createKernel :: Program
-> String
-> IO Kernel
createKernel prog name =
withForeignPtr prog $ \prog_ptr ->
withCString name $ \cstr ->
alloca $ \ep -> do
Log.debug "Invoking clCreateKernel"
kernel <- clCreateKernel prog_ptr cstr ep
checkClError_ "clCreateKernel" =<< peek ep
attachFinalizer kernel
enqueueNDRangeKernel :: CommandQueue
-> Kernel
-> [ClSize]
-> [ClSize]
-> [ClSize]
-> [Event]
-> IO Event
enqueueNDRangeKernel cq k globalWorkOffsets globalWorkSizes localWorkSizes waitEvs =
withForeignPtr cq $ \queue ->
withForeignPtr k $ \kernel ->
withArrayNull globalWorkOffsets $ \globalWorkOffsetPtr ->
withArrayNull globalWorkSizes $ \globalWorkSizePtr ->
withArrayNull localWorkSizes $ \localWorkSizePtr ->
withForeignPtrs waitEvs $ \event_ptrs ->
withArrayNullLen event_ptrs $ \n event_array ->
alloca $ \eventPtr ->
do Log.debug "Invoking clEnqueueNDRangeKernel"
checkClError_ "clEnqueueNDRangeKernel" =<<
clEnqueueNDRangeKernel
queue kernel workDim
globalWorkOffsetPtr
globalWorkSizePtr
localWorkSizePtr
(fromIntegral n)
event_array
eventPtr
attachFinalizer =<< peek eventPtr
where workDim = fromIntegral . maximum $ map length [globalWorkOffsets, globalWorkSizes, localWorkSizes]
data KernelArg where
MObjArg :: MemObject a -> KernelArg
LocalArrayArg :: Storable a => a -> Int -> KernelArg
VArg :: Storable a => a -> KernelArg
StructArg :: Storable a => [a] -> KernelArg
setKernelArg :: Kernel -> Int -> KernelArg -> IO ()
setKernelArg kernel n param =
withForeignPtr kernel $ \k ->
withPtr param $ \param_ptr -> do
Log.debug "Invoking clSetKernelArgs"
err <- clSetKernelArg k (fromIntegral n) (size param) param_ptr
case toEnum $ fromIntegral err of
InvalidArgSize -> error $ "ClInvalidArgSize occurred in call to: clSetKernelArg. Argument #"
++ show n ++ " was set to size " ++ show (size param)
InvalidArgIndex -> error $ "ClInvalidArgIndex occurred in call to: clSetKernelArg, when setting argument #"
++ show n
_ -> checkClError_ "clSetKernelArg" err
where size :: KernelArg -> ClSize
size (MObjArg mobj) = fromIntegral $ sizeOf (memobjPtr mobj)
size (VArg v) = fromIntegral $ sizeOf v
size (StructArg xs) = fromIntegral . sum $ map sizeOf xs
size (LocalArrayArg x m) = fromIntegral $ m * sizeOf x
withPtr :: KernelArg -> (Ptr () -> IO c) -> IO c
withPtr (MObjArg mobj) f = with (memobjPtr mobj) $ f . castPtr
withPtr (VArg v) f = with v $ f . castPtr
withPtr (LocalArrayArg _ _) f = f nullPtr
withPtr a@(StructArg xs) f = do
allocaBytes (fromIntegral $ size a) $ \ptr -> do
pokeElems ptr xs
f (castPtr ptr)
pokeElems :: Storable a => Ptr a -> [a] -> IO ()
pokeElems ptr (x:xs) = poke ptr x >> pokeElems (plusPtr ptr (sizeOf x)) xs
pokeElems _ [] = return ()
setKernelArgs :: Kernel -> [KernelArg] -> IO ()
setKernelArgs kernel args = zipWithM_ (setKernelArg kernel) [0..] args
enqueueTask :: CommandQueue -> Kernel -> [Event] -> IO Event
enqueueTask cq k waitEvs =
withForeignPtr cq $ \queue ->
withForeignPtr k $ \kernel ->
withForeignPtrs waitEvs $ \event_ptrs ->
withArrayNullLen event_ptrs $ \n event_array ->
alloca $ \eventPtr ->
do checkClError_ "clEnqueueTask" =<<
clEnqueueTask
queue kernel
(fromIntegral n)
event_array
eventPtr
attachFinalizer =<< peek eventPtr
kernelContext :: Kernel -> IO Context
kernelContext kernel =
getKernelInfo kernel KernelContext >>= attachRetainFinalizer
kernelFunctionName :: Kernel -> IO String
kernelFunctionName kernel = getKernelInfo kernel KernelFunctionName
kernelNumArgs :: Kernel -> IO Int
kernelNumArgs kernel = fromIntegral `fmap` (getKernelInfo kernel KernelNumArgs :: IO ClUInt)
kernelWorkGroupSize :: Kernel -> DeviceID -> IO CSize
kernelWorkGroupSize kernel device =
getKernelWorkGroupInfo kernel device KernelWorkGroupSize
kernelCompileWorkGroupSize :: Kernel -> DeviceID -> IO CSize
kernelCompileWorkGroupSize kernel device =
getKernelWorkGroupInfo kernel device KernelCompileWorkGroupSize
kernelLocalMemSize :: Kernel -> DeviceID -> IO Word64
kernelLocalMemSize kernel device =
getKernelWorkGroupInfo kernel device KernelLocalMemSize
kernelPreferredWorkGroupSizeMultiple :: Kernel -> DeviceID -> IO CSize
kernelPreferredWorkGroupSizeMultiple kernel device =
getKernelWorkGroupInfo kernel device KernelPreferredWorkGroupSizeMultiple
kernelPrivateMemSize :: Kernel -> DeviceID -> IO Word64
kernelPrivateMemSize kernel device =
getKernelWorkGroupInfo kernel device KernelPrivateMemSize
getKernelInfo kernel info =
withForeignPtr kernel $ \kernel_ptr ->
getInfo (clGetKernelInfo_ kernel_ptr) info
where
clGetKernelInfo_ =
checkClError5 "clGetKernelInfo"
clGetKernelInfo
getKernelWorkGroupInfo kernel device info =
withForeignPtr kernel $ \kernel_ptr ->
getInfo (clGetKernelWorkGroupInfo_ kernel_ptr device) info
where
clGetKernelWorkGroupInfo_ =
checkClError6 "clGetKernelWorkGroupInfo"
clGetKernelWorkGroupInfo
foreign import ccall unsafe "Foreign/OpenCL/Bindings/Kernel.chs.h clCreateKernel"
clCreateKernel :: ((Ptr (CProgram)) -> ((Ptr CChar) -> ((Ptr CInt) -> (IO (Ptr (CKernel))))))
foreign import ccall unsafe "Foreign/OpenCL/Bindings/Kernel.chs.h clEnqueueNDRangeKernel"
clEnqueueNDRangeKernel :: ((Ptr (CCommandQueue)) -> ((Ptr (CKernel)) -> (CUInt -> ((Ptr CULong) -> ((Ptr CULong) -> ((Ptr CULong) -> (CUInt -> ((Ptr (Ptr (CEvent))) -> ((Ptr (Ptr (CEvent))) -> (IO CInt))))))))))
foreign import ccall unsafe "Foreign/OpenCL/Bindings/Kernel.chs.h clSetKernelArg"
clSetKernelArg :: ((Ptr (CKernel)) -> (CUInt -> (CULong -> ((Ptr ()) -> (IO CInt)))))
foreign import ccall unsafe "Foreign/OpenCL/Bindings/Kernel.chs.h clEnqueueTask"
clEnqueueTask :: ((Ptr (CCommandQueue)) -> ((Ptr (CKernel)) -> (CUInt -> ((Ptr (Ptr (CEvent))) -> ((Ptr (Ptr (CEvent))) -> (IO CInt))))))
foreign import ccall unsafe "Foreign/OpenCL/Bindings/Kernel.chs.h clGetKernelInfo"
clGetKernelInfo :: ((Ptr (CKernel)) -> (CUInt -> (CULong -> ((Ptr ()) -> ((Ptr CULong) -> (IO CInt))))))
foreign import ccall unsafe "Foreign/OpenCL/Bindings/Kernel.chs.h clGetKernelWorkGroupInfo"
clGetKernelWorkGroupInfo :: ((Ptr (CKernel)) -> ((DeviceID) -> (CUInt -> (CULong -> ((Ptr ()) -> ((Ptr CULong) -> (IO CInt)))))))