content: {-# LANGUAGE GADTs               #-}
  {-# LANGUAGE ImpredicativeTypes  #-}
  {-# LANGUAGE PatternGuards       #-}
  {-# LANGUAGE QuasiQuotes         #-}
  {-# LANGUAGE ScopedTypeVariables #-}
  {-# LANGUAGE TemplateHaskell     #-}
  
  -- |
  -- Module      : Data.Array.Accelerate.CUDA.CodeGen.Streaming
  --
  -- Maintainer  : Frederik Meisner Madsen <fmma@diku.dk>
  -- Stability   : experimental
  -- Portability : non-portable (GHC extensions)
  --
  
  module Data.Array.Accelerate.CUDA.CodeGen.Streaming (
  
    mkToSeq
  
  ) where
  
  import Language.C.Quote.CUDA
  import Foreign.CUDA.Analysis.Device
  
  import Data.Array.Accelerate.Error                      ( internalError )
  import Data.Array.Accelerate.Array.Representation       ( SliceIndex(..) )
  import Data.Array.Accelerate.Array.Sugar                ( Array, Shape, Elt, EltRepr )
  import Data.Array.Accelerate.CUDA.AST
  import Data.Array.Accelerate.CUDA.CodeGen.Base
  
  mkToSeq
      :: forall slix aenv sh co sl a. (Shape sl, Shape sh, Elt a)
      => SliceIndex slix
                    (EltRepr sl)
                    co
                    (EltRepr sh)
      -> DeviceProperties
      -> Gamma aenv
      -> CUDelayedAcc aenv sh a
      -> CUTranslSkel aenv (Array sl a)
  mkToSeq slix dev aenv arr
    | CUDelayed (CUExp shIn) (CUFun1 _ get) _ <- arr
    = CUTranslSkel "tostream" [cunit|
  
      $esc:("#include <accelerate_cuda.h>")
      $edecls:texIn
  
      extern "C" __global__ void
      tostream
      (
          $params:argIn,
          $params:slIn,
          $params:argOut
      )
      {
          $items:(sh .=. shIn)
  
          const int shapeSize     = $exp:(csize shOut);
          const int gridSize      = $exp:(gridSize dev);
                int ix;
  
          for ( ix =  $exp:(threadIdx dev)
              ; ix <  shapeSize
              ; ix += gridSize )
          {
              $items:(src  .=. cfromIndex shOut "ix" "tmp")
              $items:(setOut "ix" .=. get fullsrc)
          }
      }
    |]
      where
        (slIn, _, cosrc)            = cslice slix "cosrc"
        (sh, _, _)                  = locals "shIn" (undefined :: sh)
        (src, _, _)                 = locals "src" (undefined :: sl)
        fullsrc                     = reverse (combine slix (reverse src) (reverse cosrc))
        (texIn, argIn)              = environment dev aenv
        (argOut, shOut, setOut)     = writeArray "Out" (undefined :: Array sl a)
  
  combine :: SliceIndex slix sl co dim -> [a] -> [a] -> [a]
  combine SliceNil [] [] = []
  combine (SliceAll   sl) (x:xs) ys = x:(combine sl xs ys)
  combine (SliceFixed sl) xs (y:ys) = y:(combine sl xs ys)
  combine _ _ _ = $internalError "mkToSeq" "Something went wrong with the slice index."
  
  