content: {-# LANGUAGE TypeFamilies #-}
  
  -- | A representation with flat parallelism via GPU-oriented kernels.
  module Futhark.IR.GPU
    ( GPU,
  
      -- * Module re-exports
      module Futhark.IR.Prop,
      module Futhark.IR.Traversals,
      module Futhark.IR.Pretty,
      module Futhark.IR.Syntax,
      module Futhark.IR.GPU.Op,
      module Futhark.IR.GPU.Sizes,
      module Futhark.IR.SOACS.SOAC,
    )
  where
  
  import Futhark.Builder
  import Futhark.Construct
  import Futhark.IR.Aliases (Aliases)
  import Futhark.IR.GPU.Op
  import Futhark.IR.GPU.Sizes
  import Futhark.IR.Pretty
  import Futhark.IR.Prop
  import Futhark.IR.SOACS.SOAC hiding (HistOp (..))
  import Futhark.IR.Syntax
  import Futhark.IR.Traversals
  import Futhark.IR.TypeCheck qualified as TC
  
  -- | The phantom data type for the kernels representation.
  data GPU
  
  instance RepTypes GPU where
    type OpC GPU = HostOp SOAC
  
  instance ASTRep GPU where
    expTypesFromPat = pure . expExtTypesFromPat
  
  instance TC.Checkable GPU where
    checkOp = typeCheckGPUOp Nothing
      where
        -- GHC 9.2 goes into an infinite loop without the type annotation.
        typeCheckGPUOp ::
          Maybe SegLevel ->
          HostOp SOAC (Aliases GPU) ->
          TC.TypeM GPU ()
        typeCheckGPUOp lvl =
          typeCheckHostOp (typeCheckGPUOp . Just) lvl typeCheckSOAC
  
  instance Buildable GPU where
    mkBody = Body ()
    mkExpPat idents _ = basicPat idents
    mkExpDec _ _ = ()
    mkLetNames = simpleMkLetNames
  
  instance BuilderOps GPU
  
  instance PrettyRep GPU
  
  instance HasSegOp GPU where
    type SegOpLevel GPU = SegLevel
    asSegOp (SegOp op) = Just op
    asSegOp _ = Nothing
    segOp = SegOp
  
  -- Note [GPU Terminology]
  --
  -- For lack of a better spot to put it, this Note summarises the
  -- terminology used for GPU concepts in the Futhark compiler. The
  -- terminology is based on CUDA terminology, and tries to match it as
  -- closely as possible. However, this was not always the case (issue
  -- #2062), so you may find some code that uses e.g. OpenCL
  -- terminology. In most cases there is no ambiguity, but there are a
  -- few instances where the same term is used for different things.
  -- Please fix any instances you find.
  --
  -- The terminology is as follows:
  --
  -- Host: Essentially the CPU; whatever is controlling the GPU.
  --
  -- Kernel: A GPU program that can be launched from the host.
  --
  -- Grid: The geometry of the thread blocks launched for a kernel. The
  -- size of a grid is always in terms of the number of thread blocks
  -- ("grid size"). A grid can have up to 3 dimensions, although we do
  -- not make much use of it - and not at all prior to code generation.
  --
  -- Thread block: Just as in CUDA. "Workgroup" in OpenCL. Abbretiation:
  -- tblock. Never just call this "block"; there are too many things
  -- called "block". Must match the dimensionality of the grid.
  --
  -- Thread: Just as in CUDA.  "Workitem" in OpenCL.
  --
  -- Global thread identifier: A globally unique number for a thread
  -- along one dimension. Abbreviation: gtid. We also use this term for
  -- the identifiers bound by SegOps. In OpenCL, corresponds to
  -- get_global_id(). (Except when we virtualise the thread space.)
  --
  -- Local thread identifier: A locally unique number (within the thread
  -- block) for each thread. Abbreviation: ltid. In OpenCL, corresponds
  -- to get_local_id().  In CUDA, corresponds to threadIdx.
  --
  -- Thread block identifier: A number unique to each thread block in a
  -- single dimension.  In CUDA, corresponds to blockIdx.
  --
  -- Local memory: Thread-local private memory. In CUDA, this is
  -- sometimes put in registers (if you are very careful in how you use
  -- it). In OpenCL, this is called "private memory", and "local memory"
  -- is something else entirely.
  --
  -- Shared memory: Just as in CUDA. Fast scratchpad memory accessible
  -- to all threads within the same thread block. In OpenCL, this is
  -- "local memory".
  --
  -- Device memory: Sometimes also called "global memory"; this is the
  -- big-but-slow memory on the GPU.
  