diff --git a/CHANGELOG.md b/CHANGELOG.md
new file mode 100644
--- /dev/null
+++ b/CHANGELOG.md
@@ -0,0 +1,5 @@
+# Revision history for tmp
+
+## 0.1.0.0 -- YYYY-mm-dd
+
+* First version. Released on an unsuspecting world.
diff --git a/LICENSE b/LICENSE
new file mode 100644
--- /dev/null
+++ b/LICENSE
@@ -0,0 +1,29 @@
+Copyright (c) 2025, Well-Typed LLP
+
+
+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 copyright holder 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.
diff --git a/src/Test/Tensor.hs b/src/Test/Tensor.hs
new file mode 100644
--- /dev/null
+++ b/src/Test/Tensor.hs
@@ -0,0 +1,617 @@
+-- | Tensors (n-dimensional arrays)
+--
+-- This is an implementation of tensors that emphasizes simplicify above all; it
+-- is meant for use in QuickCheck tests.
+--
+-- Intended for qualified import.
+--
+-- > import Test.Tensor (Tensor)
+-- > import Test.Tensor qualified as Tensor
+module Test.Tensor (
+    -- * Definition
+    Tensor(..)
+  , getScalar
+  , getTensor
+    -- ** Convenience constructors
+  , scalar
+  , dim1
+  , dim2
+  , dim3
+  , dim4
+  , dim5
+  , dim6
+  , dim7
+  , dim8
+  , dim9
+    -- * Size
+  , Size
+  , size
+  , sizeAtLeast
+    -- * Standard operations
+  , zipWith
+  , replicate
+  , rotate
+  , distrib
+  , transpose
+  , foreach
+  , foreachWith
+    -- * Subtensors
+  , subs
+  , subsWithStride
+  , convolve
+  , convolveWithStride
+  , padWith
+  , padWith'
+    -- * Conversions
+  , Lists
+  , toLists
+  , fromLists
+  , fromList
+    -- * QuickCheck support
+    -- ** Generation
+  , arbitraryOfSize
+    -- ** Shrinking
+  , shrinkWith
+  , shrinkWith'
+  , shrinkElem
+    -- *** Axes
+  , Axe(..)
+  , allAxes
+  , axeWith
+  , axeSize
+    -- *** Zeroing
+  , Zero(..)
+  , zero
+  , zeroWith
+    -- * FFI
+  , toStorable
+  , fromStorable
+  , unsafeWithCArray
+  , unsafeFromCArray
+  , unsafeFromPrealloc
+  , unsafeFromPrealloc_
+  ) where
+
+import Prelude hiding (zipWith, replicate)
+
+import Control.Monad.Trans.State (StateT(..), evalStateT)
+import Data.Bifunctor
+import Data.Foldable (foldl')
+import Data.Foldable qualified as Foldable
+import Data.List qualified as L
+import Data.Maybe (catMaybes)
+import Data.Ord
+import Data.Proxy
+import Data.Type.Nat
+import Data.Vec.Lazy (Vec(..))
+import Data.Vec.Lazy qualified as Vec
+import Data.Vector.Storable qualified as Storable (Vector)
+import Data.Vector.Storable qualified as Vector
+import Foreign hiding (rotate)
+import GHC.Show (appPrec1, showSpace)
+import GHC.Stack
+import Numeric.Natural
+import Test.QuickCheck (Arbitrary(..), Arbitrary1(..), Gen)
+import Test.QuickCheck qualified as QC
+
+{-------------------------------------------------------------------------------
+  Definition
+-------------------------------------------------------------------------------}
+
+data Tensor n a where
+  Scalar :: a -> Tensor Z a
+  Tensor :: [Tensor n a] -> Tensor (S n) a
+
+deriving stock instance Eq a => Eq (Tensor n a)
+
+deriving stock instance Functor     (Tensor n)
+deriving stock instance Traversable (Tensor n)
+deriving stock instance Foldable    (Tensor n)
+
+getScalar :: Tensor Z a -> a
+getScalar (Scalar x) = x
+
+getTensor :: Tensor (S n) a -> [Tensor n a]
+getTensor (Tensor xs) = xs
+
+{-------------------------------------------------------------------------------
+  Size
+-------------------------------------------------------------------------------}
+
+type Size n = Vec n Int
+
+-- | Analogue of 'List.length'
+size :: Tensor n a -> Size n
+size (Scalar _)  = VNil
+size (Tensor xs) = L.length xs ::: size (L.head xs)
+
+-- | Check that each dimension has at least the specified size
+sizeAtLeast :: Size n -> Tensor n a -> Bool
+sizeAtLeast sz = and . Foldable.toList . Vec.zipWith (<=) sz . size
+
+{-------------------------------------------------------------------------------
+  Standard operations
+-------------------------------------------------------------------------------}
+
+-- | Analogue of 'List.zipWith'
+zipWith :: (a -> b -> c) -> Tensor n a -> Tensor n b -> Tensor n c
+zipWith f (Scalar a)  (Scalar b)  = Scalar (f a b)
+zipWith f (Tensor as) (Tensor bs) = Tensor $ L.zipWith (zipWith f) as bs
+
+-- | Analogue of 'List.replicate'
+replicate :: Size n -> a -> Tensor n a
+replicate VNil       x = Scalar x
+replicate (n ::: ns) x = Tensor $ L.replicate n (replicate ns x)
+
+-- | Analogue of 'List.reverse'
+--
+-- This amounts to a 180 degrees rotation of the tensor.
+rotate :: Tensor n a -> Tensor n a
+rotate (Scalar x)  = Scalar x
+rotate (Tensor xs) = Tensor $ map rotate (L.reverse xs)
+
+-- | Distribute '[]' over 'Tensor'
+--
+-- Collects values in corresponding in all tensors.
+distrib :: [Tensor n a] -> Tensor n [a]
+distrib = \case
+    []   -> error "distrib: empty list"
+    t:ts -> go ((:[]) <$> t) ts
+  where
+    go :: Tensor n [a] -> [Tensor n a] -> Tensor n [a]
+    go acc []     = reverse <$> acc
+    go acc (t:ts) = go (zipWith (:) t acc) ts
+
+-- | Transpose
+--
+-- This is essentially a special case of 'distrib'.
+transpose :: Tensor Nat2 a -> Tensor Nat2 a
+transpose = fromLists . L.transpose . toLists
+
+-- | Map element over the first dimension of the tensor
+foreach :: Tensor (S n) a -> (Tensor n a -> Tensor m b) -> Tensor (S m) b
+foreach (Tensor as) f = Tensor (Prelude.map f as)
+
+-- | Variation of 'foreach' with an auxiliary list
+foreachWith ::
+    Tensor (S n) a
+ -> [x]
+ -> (Tensor n a -> x -> Tensor m b)
+ -> Tensor (S m) b
+foreachWith (Tensor as) xs f = Tensor (L.zipWith f as xs)
+
+{-------------------------------------------------------------------------------
+  Subtensors
+-------------------------------------------------------------------------------}
+
+-- | Subtensors of the specified size
+subs :: SNatI n => Size n -> Tensor n a -> Tensor n (Tensor n a)
+subs = subsWithStride (pure 1)
+
+-- | Generalization of 'subs' with non-default stride
+subsWithStride :: Vec n Int -> Size n -> Tensor n a -> Tensor n (Tensor n a)
+subsWithStride VNil       VNil       (Scalar x)  = Scalar (Scalar x)
+subsWithStride (s ::: ss) (n ::: ns) (Tensor xs) = Tensor [
+      Tensor <$> distrib selected
+    | selected <- everyNth s $ consecutive n (map (subsWithStride ss ns) xs)
+    ]
+
+-- | Convolution
+--
+-- See 'padWith' for adjusting boundary conditions.
+convolve ::
+     (SNatI n, Num a)
+  => Tensor n a  -- ^ Kernel
+  -> Tensor n a  -- ^ Input
+  -> Tensor n a
+convolve = convolveWithStride (pure 1)
+
+-- | Generalization of 'convolve' when using a non-default stride
+convolveWithStride :: forall n a.
+     Num a
+  => Vec n Int   -- ^ Stride
+  -> Tensor n a  -- ^ Kernel
+  -> Tensor n a  -- ^ Input
+  -> Tensor n a
+convolveWithStride stride kernel input =
+    aux <$> subsWithStride stride (size kernel) input
+  where
+    aux :: Tensor n a -> a
+    aux = foldl' (+) 0 . zipWith (*) kernel
+
+{-------------------------------------------------------------------------------
+  Padding
+-------------------------------------------------------------------------------}
+
+-- | Add uniform padding
+padWith :: SNatI n => a -> Int -> Tensor n a -> Tensor n a
+padWith padding n = padWith' padding (pure (n, n))
+
+-- | Generalization of 'padWith' with different padding per dimension
+padWith' :: forall n a. a -> Vec n (Int, Int) -> Tensor n a -> Tensor n a
+padWith' padding paddingSize tensor =
+    go paddingSize newSize tensor
+  where
+    newSize :: Size n
+    newSize = Vec.zipWith (\(b, a) n -> n + b + a) paddingSize (size tensor)
+
+    go :: forall m. Vec m (Int, Int) -> Size m -> Tensor m a -> Tensor m a
+    go VNil                     VNil       (Scalar x)  = Scalar x
+    go ((before, after) ::: ps) (_ ::: ns) (Tensor xs) = Tensor $ concat [
+          L.replicate before $ replicate ns padding
+        , map (go ps ns) xs
+        , L.replicate after $ replicate ns padding
+        ]
+
+{-------------------------------------------------------------------------------
+  QuickCheck support
+-------------------------------------------------------------------------------}
+
+arbitraryOfSize :: Size n -> Gen a -> Gen (Tensor n a)
+arbitraryOfSize sz = sequence . replicate sz
+
+data Axe (n :: Nat) where
+  -- | Axe some elements from the current dimension
+  --
+  -- We record which elements to drop as an @(offset, length)@ pair.
+  AxeHere :: (Int, Int) -> Axe (S n)
+
+  -- | Axe some elements from a nested dimension
+  --
+  -- In order to keep the tensor square, we must apply the same axe for every
+  -- element of the /current/ dimension
+  AxeNested :: Axe n -> Axe (S n)
+
+deriving instance Show (Axe n)
+
+-- | How many elements are removed by this axe?
+--
+-- Examples:
+--
+-- > axeSize (2 ::: 100 ::: VNil) (AxeHere (0, 1))               == 100
+-- > axeSize (2 ::: 100 ::: VNil) (AxeNested (AxeHere (0, 99)))  == 198
+axeSize :: Size n -> Axe n -> Int
+axeSize = flip go
+  where
+    go ::  Axe n -> Size n -> Int
+    go (AxeHere (_, len)) (_ ::: ns) = len * L.foldl' (*) 1 ns
+    go (AxeNested axe)    (n ::: ns) = n * go axe ns
+
+-- | All possible ways to axe some elements
+--
+-- This is adopted from the implementation of 'shrinkList' (in a way, an 'Axe'
+-- is an explanation of the decisions made by 'shrinkList', generalized to
+-- multiple dimensions).
+--
+-- Axes are sorted to remove as many elements as early as possible.
+allAxes :: Size n -> [Axe n]
+allAxes = \sz ->
+    L.sortBy (flip $ comparing (axeSize sz)) $ go sz
+  where
+    go :: Size n -> [Axe n]
+    go VNil       = []
+    go (n ::: ns) = concat [
+          concat [
+              L.map AxeHere (removes 0 k n)
+            | k <- takeWhile (> 0) (iterate (`div` 2) (n `div` 2))
+            ]
+        , L.map AxeNested (go ns)
+        ]
+
+    removes :: Int -> Int -> Int -> [(Int, Int)]
+    removes offset k n
+      | k > n     = []
+      | otherwise = (offset, k) : removes (offset + k) k (n - k)
+
+-- | Remove elements from the tensor (shrink dimensions)
+axeWith :: Axe n -> Tensor n a -> Tensor n a
+axeWith (AxeHere (offset, len)) (Tensor xss) = Tensor $
+    before <> after
+  where
+    (before, dropFrom) = L.splitAt offset xss
+    (_dropped, after)  = L.splitAt len dropFrom
+axeWith (AxeNested axe) (Tensor xss) = Tensor $
+    L.map (axeWith axe) xss
+
+-- | Zero element
+data Zero a where
+  Zero :: Eq a => a -> Zero a
+
+-- | Default 'Zero'
+zero :: (Num a, Eq a) => Zero a
+zero = Zero 0
+
+-- | Zero elements in the tensor (leaving dimensions the same)
+--
+-- Returns 'Nothing' if the specified region was already zero everywhere.
+zeroWith :: forall n a. Zero a -> Axe n -> Tensor n a -> Maybe (Tensor n a)
+zeroWith (Zero z) = \axe tensor ->
+    case go axe (size tensor) tensor of
+      (_, False)      -> Nothing
+      (tensor', True) -> Just tensor'
+  where
+    -- Additionally returns if anything changed
+    go :: forall n'. Axe n' -> Size n' -> Tensor n' a -> (Tensor n' a, Bool)
+    go (AxeHere (offset, len)) (_ ::: ns) (Tensor xss) = (
+          Tensor $ before <> L.replicate len (replicate ns z) <> after
+        , any (/= z) (Tensor dropped)
+        )
+      where
+         (before, dropFrom) = L.splitAt offset xss
+         (dropped, after)   = L.splitAt len dropFrom
+    go (AxeNested axe) (_ ::: ns) (Tensor xss) =
+        bimap Tensor or $ L.unzip $ L.map (go axe ns) xss
+
+-- | Shrink tensor
+shrinkWith ::
+     Maybe (Zero a)  -- ^ Optional zero element (see 'shrinkElem')
+  -> (a -> [a])      -- ^ Shrink individual elements
+  -> Tensor n a -> [Tensor n a]
+shrinkWith mZero f xs = shrinkWith' (allAxes (size xs)) mZero f xs
+
+-- | Generalization of 'shrinkWith'
+shrinkWith' :: forall n a.
+     [Axe n]         -- ^ Shrink the size of the tensor (see 'allAxes')
+  -> Maybe (Zero a)  -- ^ Optional zero element (see 'shrinkElem')
+  -> (a -> [a])      -- ^ Shrink elements of the tensor
+  -> Tensor n a -> [Tensor n a]
+shrinkWith' axes mZero f xss = concat [
+      [axeWith axe xss | axe <- axes]
+    , shrinkElem mZero f xss
+    ]
+
+-- | Shrink an element of the tensor, leaving the size of the tensor unchanged
+--
+-- If a zero element is specified, we will first try to replace entire regions
+-- of the tensor by zeroes; this can dramatically speed up shrinking.
+shrinkElem :: forall n a.
+     Maybe (Zero a)  -- ^ Optional zero element
+  -> (a -> [a])      -- ^ Shrink individual elements
+  -> Tensor n a -> [Tensor n a]
+shrinkElem mZero f tensor = concat [
+      case mZero of
+        Nothing -> []
+        Just z  -> catMaybes [
+            zeroWith z axe tensor
+          | axe <- allAxes overallSize
+          , axeSize overallSize axe > 1
+          ]
+    , shrinkOne tensor
+    ]
+  where
+    overallSize :: Size n
+    overallSize = size tensor
+
+    shrinkOne :: forall n'. Tensor n' a -> [Tensor n' a]
+    shrinkOne (Scalar x)   = Scalar <$> f x
+    shrinkOne (Tensor xss) = [
+          Tensor $ before ++ [xs'] ++ after
+        | (before, xs, after) <- pickOne xss
+        , xs' <- shrinkOne xs
+        ]
+
+instance (SNatI n, Arbitrary a, Num a, Eq a) => Arbitrary (Tensor n a) where
+  arbitrary = liftArbitrary arbitrary
+  shrink    = shrinkWith (Just (Zero 0)) shrink
+
+-- | Lift generators and shrinkers
+--
+-- NOTE: Since we cannot put any constraints on the type of the elements here,
+-- we cannot use any zero elements. Using 'shrink' (or 'shrinkWith' directly)
+-- might result in faster shrinking.
+instance SNatI n => Arbitrary1 (Tensor n) where
+  liftArbitrary g = QC.sized $ \n -> do
+      sz :: Size n <- liftArbitrary $ QC.choose (1, 1 + n)
+      arbitraryOfSize sz g
+
+  liftShrink f = shrinkWith Nothing f
+
+{-------------------------------------------------------------------------------
+  FFI
+-------------------------------------------------------------------------------}
+
+-- | Translate to storable vector
+--
+-- The tensor is laid out in order specified (outer dimensions before inner).
+toStorable :: Storable a => Tensor n a -> Storable.Vector a
+toStorable = Vector.fromList . Foldable.toList
+
+-- | Translate from storable vector
+--
+-- Throws an exception if the vector does not contain enough elements.
+fromStorable ::
+     (HasCallStack, Storable a)
+  => Size n -> Storable.Vector a -> Tensor n a
+fromStorable sz = fromList sz . Vector.toList
+
+-- | Get pointer to elements of the tensor
+--
+-- See 'toStorable' for discussion of the layout.
+--
+-- The data should not be modified through the pointer, and the pointer should
+-- not be used outside its scope.
+unsafeWithCArray :: Storable a => Tensor n a -> (Ptr a -> IO r) -> IO r
+unsafeWithCArray tensor = Vector.unsafeWith (toStorable tensor)
+
+-- | Construct tensor from C array
+--
+-- The data should not be modified through the pointer after the tensor has
+-- been constructed.
+unsafeFromCArray :: Storable a => Size n -> ForeignPtr a -> Tensor n a
+unsafeFromCArray sz fptr =
+    fromStorable sz $ Vector.unsafeFromForeignPtr0 fptr n
+  where
+    n :: Int
+    n = L.foldl' (*) 1 sz
+
+-- | Construct tensor from preallocated C array
+--
+-- Allocates sufficient memory to hold the elements of the tensor; writing more
+-- data will result in invalid memory access. The pointer should not be used
+-- outside its scope.
+unsafeFromPrealloc ::
+     Storable a
+  => Size n -> (Ptr a -> IO r) -> IO (Tensor n a, r)
+unsafeFromPrealloc sz k = do
+    fptr <- mallocForeignPtrArray n
+    res  <- withForeignPtr fptr k
+    return (unsafeFromCArray sz fptr, res)
+  where
+    n :: Int
+    n = L.foldl' (*) 1 sz
+
+-- | Like 'unsafeFromPrealloc' but without an additional return value
+unsafeFromPrealloc_ ::
+     Storable a
+  => Size n -> (Ptr a -> IO ()) -> IO (Tensor n a)
+unsafeFromPrealloc_ sz = fmap fst . unsafeFromPrealloc sz
+
+{-------------------------------------------------------------------------------
+  Convenience constructors
+-------------------------------------------------------------------------------}
+
+scalar :: a -> Tensor Nat0 a
+scalar = fromLists
+
+dim1 :: [a] -> Tensor Nat1 a
+dim1 = fromLists
+
+dim2 :: [[a]] -> Tensor Nat2 a
+dim2 = fromLists
+
+dim3 :: [[[a]]] -> Tensor Nat3 a
+dim3 = fromLists
+
+dim4 :: [[[[a]]]] -> Tensor Nat4 a
+dim4 = fromLists
+
+dim5 :: [[[[[a]]]]] -> Tensor Nat5 a
+dim5 = fromLists
+
+dim6 :: [[[[[[a]]]]]] -> Tensor Nat6 a
+dim6 = fromLists
+
+dim7 :: [[[[[[[a]]]]]]] -> Tensor Nat7 a
+dim7 = fromLists
+
+dim8 :: [[[[[[[[a]]]]]]]] -> Tensor Nat8 a
+dim8 = fromLists
+
+dim9 :: [[[[[[[[[a]]]]]]]]] -> Tensor Nat9 a
+dim9 = fromLists
+
+{-------------------------------------------------------------------------------
+  Conversions
+
+  This is primarily useful for specify tensor constants.
+-------------------------------------------------------------------------------}
+
+type family Lists n a where
+  Lists Z     a = a
+  Lists (S n) a = [Lists n a]
+
+toLists :: Tensor n a -> Lists n a
+toLists (Scalar x)  = x
+toLists (Tensor xs) = map toLists xs
+
+fromLists :: SNatI n => Lists n a -> Tensor n a
+fromLists = go snat
+  where
+    go :: SNat n -> Lists n a -> Tensor n a
+    go SZ = Scalar
+    go SS = Tensor . map (go snat)
+
+-- | Inverse to 'Foldable.toList'
+--
+-- Throws a pure exception if the list does not contain enough elements.
+fromList :: forall n a. Size n -> [a] -> Tensor n a
+fromList sz xs =
+    checkEnoughElems . flip evalStateT xs $ sequenceA (replicate sz genElem)
+  where
+    genElem :: StateT [a] Maybe a
+    genElem = StateT L.uncons
+
+    checkEnoughElems :: Maybe (Tensor n a) -> Tensor n a
+    checkEnoughElems Nothing  = error "fromList: insufficient elements"
+    checkEnoughElems (Just t) = t
+
+{-------------------------------------------------------------------------------
+  Show instance
+-------------------------------------------------------------------------------}
+
+showLists :: Show a => Proxy a -> SNat n -> (Show (Lists n a) => r) -> r
+showLists _ SZ      k = k
+showLists p (SS' n) k = showLists p n k
+
+showConstructor :: Int -> SNat n -> ShowS
+showConstructor p sn
+  | n' == 0            = showString "scalar"
+  | 1 <= n' && n' <= 9 = showString "dim" . shows n'
+  | otherwise          = showString "fromLists @"
+                       . explicitShowsPrec p (snatToNat sn)
+  where
+    n' :: Natural
+    n' = snatToNatural sn
+
+instance Show a => Show (Tensor n a) where
+  showsPrec p tensor = showLists (Proxy @a) (tensorSNat tensor) $
+      showParen (p >= appPrec1) $
+          showConstructor appPrec1 (tensorSNat tensor)
+        . showSpace
+        . showsPrec appPrec1 (toLists tensor)
+
+{-------------------------------------------------------------------------------
+  Internal auxiliary: SNat
+-------------------------------------------------------------------------------}
+
+tensorSNatI :: Tensor n a -> (SNatI n => r) -> r
+tensorSNatI (Scalar _)  k = k
+tensorSNatI (Tensor xs) k = tensorSNatI (L.head xs) k
+
+tensorSNat :: Tensor n a -> SNat n
+tensorSNat tensor = tensorSNatI tensor snat
+
+{-------------------------------------------------------------------------------
+  Internal auxiliary: lists
+-------------------------------------------------------------------------------}
+
+-- | Consecutive elements
+--
+-- >    consecutive 3 [1..5]
+-- > == [[1,2,3],[2,3,4],[3,4,5]]
+consecutive :: Int -> [a] -> [[a]]
+consecutive n = L.takeWhile ((== n) . length) . fmap (L.take n) . L.tails
+
+-- | Every nth element of the list
+--
+-- Examples
+--
+-- > everyNth 1 [0..9] == [0,2,3,4,5,6,7,8,9]
+-- > everyNth 2 [0..9] == [0,2,4,6,8]
+-- > everyNth 3 [0..9] == [0,3,6,9]
+everyNth :: forall a. Int -> [a] -> [a]
+everyNth n = \xs ->
+    if n > 0
+      then go xs
+      else error "everyNth: n should be strictly positive"
+  where
+    go :: [a] -> [a]
+    go []     = []
+    go (x:xs) = x : go (drop (n - 1) xs)
+
+-- | Single out an element from the list
+--
+-- >    pickOne [1..4]
+-- > == [ ( []      , 1 , [2,3,4] )
+-- >    , ( [1]     , 2 , [3,4]   )
+-- >    , ( [1,2]   , 3 , [4]     )
+-- >    , ( [1,2,3] , 4 , []      )
+-- >    ]
+pickOne :: forall a. [a] -> [([a], a, [a])]
+pickOne = \case
+    []   -> error "pickOne: empty list"
+    x:xs -> go [] x xs
+  where
+    go :: [a] -> a -> [a] -> [([a], a, [a])]
+    go acc x []     = [(reverse acc, x, [])]
+    go acc x (y:ys) = (reverse acc, x, (y:ys)) : go (x:acc) y ys
diff --git a/src/Test/Tensor/TestValue.hs b/src/Test/Tensor/TestValue.hs
new file mode 100644
--- /dev/null
+++ b/src/Test/Tensor/TestValue.hs
@@ -0,0 +1,119 @@
+-- | Test values
+--
+-- Intended for unqualified import.
+module Test.Tensor.TestValue (
+    TestValue -- opaque
+  ) where
+
+import Data.List (sort)
+import System.Random (Random)
+import Test.QuickCheck
+import Text.Printf (printf)
+
+{-------------------------------------------------------------------------------
+  Definition
+-------------------------------------------------------------------------------}
+
+-- | Test values
+--
+-- Test values are suitable for use in QuickCheck tests involving floating
+-- point numbers, if you want to ignore rounding errors.
+newtype TestValue = TestValue Float
+  deriving newtype (Num, Fractional, Real, Random)
+
+-- | Test values are equipped with a crude equality
+--
+-- >               (==)
+-- > --------------------
+-- > 1.0    1.1    False
+-- > 1.00   1.01   True
+-- > 10     11     False
+-- > 10.0   10.1   True
+-- > 100    110    False
+-- > 100    101    True
+instance Eq TestValue where
+  TestValue x == TestValue y = nearlyEqual x y
+
+-- | Show instance
+--
+-- We have more precision available for smaller values, so we show more
+-- decimals. However, larger values the show instance does not reflect the
+-- precision: @1000@ and @1001@ are shown as @1000@ and @1001@, even though
+-- they are considered to be equal.
+--
+-- > show @TestValue 0     == "0"     -- True zero
+-- > show @TestValue 1     == "1"     -- True one
+-- > show @TestValue 0.001 == "0.00"
+-- > show @TestValue 0.009 == "0.01"
+-- > show @TestValue 1.001 == "1.0"
+-- > show @TestValue 11    == "11"
+instance Show TestValue where
+  show (TestValue x)
+    | x == 0    = "0"
+    | x == 1    = "1"
+    | x <  1    = printf "%0.2f" x
+    | x <  10   = printf "%0.1f" x
+    | otherwise = printf "%0.0f" x
+
+-- | Arbitrary instance
+--
+-- The definition of 'arbitrary' simply piggy-backs on the definition for
+-- 'Float', but in shrinking we avoid generating nearly equal values, and prefer
+-- values closer to integral values. Compare:
+--
+-- >    shrink @TestValue 100.1
+-- > == [0,50,75,88,94,97]
+--
+-- versus
+--
+-- >    shrink @Float 100.1
+-- > == [100.0,0.0,50.0,75.0,88.0,94.0,97.0,99.0,0.0,50.1,75.1,87.6,93.9,97.0,98.6,99.4,99.8,100.0]
+instance Arbitrary TestValue where
+  arbitrary = TestValue <$> arbitrary
+
+  shrink (TestValue x)
+    | x == 0          = []
+    | nearlyEqual x 0 = [0]
+    | otherwise       = case sort (shrink x) of
+                          []   -> []
+                          y:ys -> aux y ys
+    where
+      aux :: Float -> [Float] -> [TestValue]
+      aux y []
+        | nearlyEqual y x = []
+        | otherwise       = [TestValue y]
+      aux y (z:zs)
+        | nearlyEqual y z = if decimalPart y < decimalPart z
+                              then aux y zs
+                              else aux z zs
+        | otherwise       = TestValue y : aux z zs
+
+instance Ord TestValue where
+  compare (TestValue x) (TestValue y)
+    | nearlyEqual x y = EQ
+    | x < y           = LT
+    | otherwise       = GT
+
+{-------------------------------------------------------------------------------
+  Internal auxiliary
+-------------------------------------------------------------------------------}
+
+-- | Compare for near equality
+--
+-- Adapted from <https://stackoverflow.com/a/32334103/742991>
+nearlyEqual :: Float -> Float -> Bool
+nearlyEqual a b
+  | a == b    = True
+  | otherwise = diff < max abs_th (epsilon * norm)
+  where
+    diff, norm :: Float
+    diff = abs (a - b)
+    norm = abs a + abs b
+
+    -- Define precision
+    abs_th, epsilon :: Float
+    epsilon = 0.01
+    abs_th  = 0.01
+
+decimalPart :: Float -> Float
+decimalPart x = x - fromIntegral (floor x :: Int)
diff --git a/test-cbits/test-cudnn.c b/test-cbits/test-cudnn.c
new file mode 100644
--- /dev/null
+++ b/test-cbits/test-cudnn.c
@@ -0,0 +1,230 @@
+#include "test-cudnn.h"
+
+#include <assert.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#include <cudnn.h>
+
+int test_cudnn_binding_version(void) {
+    return 1;
+}
+
+int test_cudnn_library_version(void) {
+    return CUDNN_VERSION;
+}
+
+// #define DEBUG 1
+
+/**
+ * Relevant references:
+ *
+ * - https://www.goldsborough.me/cuda/ml/cudnn/c++/2017/10/01/14-37-23-convolutions_with_cudnn/
+ * - https://docs.nvidia.com/deeplearning/cudnn/backend/latest/index.html
+ *   (in particular the `ops` and `cnn` libraries)
+ * - https://docs.nvidia.com/deeplearning/performance/dl-performance-convolutional/index.html
+ * - https://cs231n.github.io/convolutional-networks/
+ *
+ * Note on kernels: the kernel should have as many channels as the input
+ * ("input features"), but we can have multiple kernels ("output features").
+ * The result will be of size batch size * output features * height * width.
+ *
+ * There is a helpful diagram at
+ * https://docs.nvidia.com/deeplearning/performance/dl-performance-convolutional/index.html#convo-intro
+ * that illustrates this.
+ */
+
+#define checkCUDNN(expression)                                                 \
+  {                                                                            \
+    cudnnStatus_t status = (expression);                                       \
+    if (status != CUDNN_STATUS_SUCCESS) {                                      \
+      printf("Error on line %d: %s\n", __LINE__, cudnnGetErrorString(status)); \
+      exit(EXIT_FAILURE);                                                      \
+    }                                                                          \
+  }
+
+float* test_cudnn_convolve(
+  cudnnConvolutionMode_t mode,
+  int vertical_stride, int horizontal_stride,
+  int num_kernels, int kernel_height, int kernel_width,
+  float* kernel,
+  int num_images, int input_channels, int input_height, int input_width,
+  float* input,
+  int* output_height, int* output_width
+) {
+    cudnnHandle_t cudnn;
+    checkCUDNN(cudnnCreate(&cudnn));
+
+#ifdef DEBUG
+    printf("mode = %d, vertical_stride = %d, horizontal_stride = %d\n", mode, vertical_stride, horizontal_stride);
+    printf("num_kernels = %d, kernel_height = %d, kernel_width = %d\n", num_kernels, kernel_height, kernel_width);
+    printf("num_images = %d, input_channels = %d, input_height = %d, input_width = %d\n", num_images, input_channels, input_height, input_width);
+#endif
+
+    /**
+     * Configure convolution
+     */
+
+    cudnnConvolutionDescriptor_t convolution_descriptor;
+    checkCUDNN(cudnnCreateConvolutionDescriptor(&convolution_descriptor));
+    checkCUDNN(cudnnSetConvolution2dDescriptor(convolution_descriptor,
+      /* pad_h             */ 0,
+      /* pad_w             */ 0,
+      /* vertical stride   */ vertical_stride,
+      /* horizontal stride */ horizontal_stride,
+      /* dilation_h        */ 1, // No dilation
+      /* dilation_w        */ 1,
+      /* mode              */ mode,
+      /* computeType       */ CUDNN_DATA_FLOAT));
+
+    /**
+     * Setup input
+     */
+
+    cudnnTensorDescriptor_t input_descriptor;
+    checkCUDNN(cudnnCreateTensorDescriptor(&input_descriptor));
+    checkCUDNN(cudnnSetTensor4dDescriptor(input_descriptor,
+      /* format   */ CUDNN_TENSOR_NCHW,
+      /* dataType */ CUDNN_DATA_FLOAT,
+      /* n        */ num_images,
+      /* c        */ input_channels,
+      /* h        */ input_height,
+      /* w        */ input_width));
+
+    cudnnFilterDescriptor_t kernel_descriptor;
+    checkCUDNN(cudnnCreateFilterDescriptor(&kernel_descriptor));
+    checkCUDNN(cudnnSetFilter4dDescriptor(kernel_descriptor,
+      /* dataType        */ CUDNN_DATA_FLOAT,
+      /* format          */ CUDNN_TENSOR_NCHW,
+      /* output channels */ num_kernels,
+      /* input channels  */ input_channels,
+      /* h               */ kernel_height,
+      /* w               */ kernel_width));
+
+    /**
+     * Setup output
+     */
+
+    int num_output_images = -1;
+    int output_channels   = -1;
+
+    checkCUDNN(cudnnGetConvolution2dForwardOutputDim(
+      /* convDesc        */ convolution_descriptor,
+      /* inputTensorDesc */ input_descriptor,
+      /* filterDesc      */ kernel_descriptor,
+      /* n               */ &num_output_images,
+      /* c               */ &output_channels,
+      /* h               */ output_height,
+      /* w               */ output_width));
+
+#ifdef DEBUG
+    printf("num_output_images = %d, output_channels = %d, output_height = %d, output_width = %d\n", num_output_images, output_channels, *output_height, *output_width);
+#endif
+
+    assert(num_output_images == num_images);
+    assert(output_channels   == num_kernels);
+
+    cudnnTensorDescriptor_t output_descriptor;
+    checkCUDNN(cudnnCreateTensorDescriptor(&output_descriptor));
+    checkCUDNN(cudnnSetTensor4dDescriptor(output_descriptor,
+      /* format   */ CUDNN_TENSOR_NCHW,
+      /* dataType */ CUDNN_DATA_FLOAT,
+      /* n        */ num_output_images,
+      /* c        */ output_channels,
+      /* h        */ *output_height,
+      /* w        */ *output_width));
+
+    /**
+     * Prepare convolution
+     */
+
+    cudnnConvolutionFwdAlgoPerf_t convolution_algorithm_perf;
+    int returned_algo_count;
+    checkCUDNN(cudnnFindConvolutionForwardAlgorithm(cudnn,
+      /* xDesc              */ input_descriptor,
+      /* wDesc              */ kernel_descriptor,
+      /* convDesc           */ convolution_descriptor,
+      /* yDesc              */ output_descriptor,
+      /* requestedAlgoCount */ 1,
+      /* returnedAlgoCount  */ &returned_algo_count,
+      /* perfResults        */ &convolution_algorithm_perf));
+    cudnnConvolutionFwdAlgo_t convolution_algorithm = convolution_algorithm_perf.algo;
+
+    size_t workspace_bytes = 0;
+    checkCUDNN(cudnnGetConvolutionForwardWorkspaceSize(cudnn,
+      /* xDesc       */ input_descriptor,
+      /* wDesc       */ kernel_descriptor,
+      /* convDesc    */ convolution_descriptor,
+      /* yDesc       */ output_descriptor,
+      /* algo        */ convolution_algorithm,
+      /* sizeInBytes */ &workspace_bytes));
+
+    /**
+     * Allocate device memory
+     */
+
+    int input_bytes  = num_images  * input_channels  * input_height  * input_width  * sizeof(float);
+    int output_bytes = num_output_images * output_channels * (*output_height) * (*output_width) * sizeof(float);
+    int kernel_bytes = num_kernels * input_channels * kernel_height * kernel_width * sizeof(float);
+
+    void*  d_workspace = NULL;
+    float* d_input     = NULL;
+    float* d_output    = NULL;
+    float* d_kernel    = NULL;
+
+    cudaMalloc((void**) &d_workspace, workspace_bytes);
+    cudaMalloc((void**) &d_input, input_bytes);
+    cudaMalloc((void**) &d_output, output_bytes);
+    cudaMalloc((void**) &d_kernel, kernel_bytes);
+
+    /**
+     * Initialize memory
+     *
+     * Everything up to this point has been completely independent from the
+     * specific choice of input and kernel (apart from their size).
+     */
+
+    cudaMemcpy(d_input, input, input_bytes, cudaMemcpyHostToDevice);
+    cudaMemcpy(d_kernel, kernel, kernel_bytes, cudaMemcpyHostToDevice);
+    cudaMemset(d_output, 0, output_bytes);
+
+    /**
+     * Execute the convolution
+     */
+
+    float alpha = 1, beta = 0; // no blending
+    checkCUDNN(cudnnConvolutionForward(cudnn,
+      /* alpha                */ &alpha,
+      /* xDesc                */ input_descriptor,
+      /* x                    */ d_input,
+      /* wDesc                */ kernel_descriptor,
+      /* w                    */ d_kernel,
+      /* convDesc             */ convolution_descriptor,
+      /* algo                 */ convolution_algorithm,
+      /* workSpace            */ d_workspace,
+      /* workSpaceSizeInBytes */ workspace_bytes,
+      /* beta                 */ &beta,
+      /* yDesc                */ output_descriptor,
+      /* y                    */ d_output));
+
+    /**
+     * Copy results back to host and deallocate resources
+     */
+
+    float* output = (float*) malloc(output_bytes);
+    cudaMemcpy(output, d_output, output_bytes, cudaMemcpyDeviceToHost);
+
+    cudaFree(d_workspace);
+    cudaFree(d_input);
+    cudaFree(d_output);
+    cudaFree(d_kernel);
+
+    cudnnDestroyConvolutionDescriptor(convolution_descriptor);
+    cudnnDestroyTensorDescriptor(input_descriptor);
+    cudnnDestroyFilterDescriptor(kernel_descriptor);
+    cudnnDestroyTensorDescriptor(output_descriptor);
+
+    checkCUDNN(cudnnDestroy(cudnn));
+
+    return output;
+}
diff --git a/test/Main.hs b/test/Main.hs
new file mode 100644
--- /dev/null
+++ b/test/Main.hs
@@ -0,0 +1,32 @@
+{-# LANGUAGE CPP #-}
+
+module Main (main) where
+
+import Test.Tasty
+
+import TestSuite.Test.Convolution qualified as Convolution
+import TestSuite.Test.QuickCheck  qualified as QuickCheck
+import TestSuite.Test.StdOps      qualified as StdOps
+
+#ifdef TEST_FFT
+import TestSuite.Test.Convolution.FFT qualified as Convolution.FFT
+#endif
+
+#ifdef TEST_CUDNN
+import TestSuite.Test.Convolution.CUDNN qualified as Convolution.CUDNN
+#endif
+
+main :: IO ()
+main = defaultMain $ testGroup "testing-tensor" [
+      testGroup "Convolutions" [
+          QuickCheck.tests
+        , StdOps.tests
+        , Convolution.tests
+#ifdef TEST_FFT
+        , Convolution.FFT.tests
+#endif
+#ifdef TEST_CUDNN
+        , Convolution.CUDNN.tests
+#endif
+      ]
+    ]
diff --git a/test/TestSuite/Test/Convolution.hs b/test/TestSuite/Test/Convolution.hs
new file mode 100644
--- /dev/null
+++ b/test/TestSuite/Test/Convolution.hs
@@ -0,0 +1,241 @@
+module TestSuite.Test.Convolution (tests) where
+
+import Data.List qualified as L
+import Data.Type.Nat
+import Data.Vec.Lazy (Vec(..))
+import Test.Tasty
+import Test.Tasty.HUnit
+import Test.Tasty.QuickCheck
+
+import Test.Tensor (Tensor)
+import Test.Tensor qualified as Tensor
+import Test.Tensor.TestValue
+
+import TestSuite.Test.Convolution.Examples3B1B
+
+{-------------------------------------------------------------------------------
+  List of tests
+-------------------------------------------------------------------------------}
+
+tests :: TestTree
+tests = testGroup "TestSuite.Test.Convolution.Prop" [
+      testGroup "Examples" [
+          testCase "rotate"       example_rotate
+        , testCase "distrib_dim2" example_distrib_dim2
+        , testCase "subs_dim1"    example_subs_dim1
+        , testCase "subs_dim2"    example_subs_dim2
+        , testCase "subs_dim3"    example_subs_dim3
+        , testCase "padWith"      example_padWith
+        , testCase "padWith'"     example_padWith'
+        ]
+    , testGroup "3B1B" [
+          testCase "simple"                example_3b1b_simple
+        , testCase "movingAverage"         example_3b1b_movingAverage
+        , testCase "movingWeightedAverage" example_3b1b_movingWeightedAverage
+        , testCase "weightedDice"          example_3b1b_weightedDice
+        ]
+    , testGroup "Properties" [
+          testProperty "distrib_dim0"            prop_distrib_dim0
+        , testProperty "distrib_dim1"            prop_distrib_dim1
+        , testProperty "distrib_dim1_nonUniform" prop_distrib_dim1_nonUniform
+        ]
+    ]
+
+{-------------------------------------------------------------------------------
+  Examples
+-------------------------------------------------------------------------------}
+
+example_rotate :: Assertion
+example_rotate =
+    assertEqual "" expected $
+      Tensor.rotate (Tensor.dim2 [ [1,2,3], [4,5,6] ])
+  where
+    expected :: Tensor Nat2 Integer
+    expected = Tensor.dim2 [ [6,5,4], [3,2,1] ]
+
+example_distrib_dim2 :: Assertion
+example_distrib_dim2 =
+    assertEqual "" expected $
+      Tensor.distrib input
+  where
+    input :: [Tensor Nat2 Int]
+    input = [
+          Tensor.dim2 [[111, 112, 113, 114], [121, 122, 123, 124], [131, 132, 133, 134]]
+        , Tensor.dim2 [[211, 212, 213, 214], [221, 222, 223, 224], [231, 232, 233, 234]]
+        , Tensor.dim2 [[311, 312, 313, 314], [321, 322, 323, 324], [331, 332, 333, 334]]
+        , Tensor.dim2 [[411, 412, 413, 414], [421, 422, 423, 424], [431, 432, 433, 434]]
+        , Tensor.dim2 [[511, 512, 513, 514], [521, 522, 523, 524], [531, 532, 533, 534]]
+        ]
+
+    expected :: Tensor Nat2 [Int]
+    expected = Tensor.dim2 [
+          [ [111,211,311,411,511]
+          , [112,212,312,412,512]
+          , [113,213,313,413,513]
+          , [114,214,314,414,514]
+          ]
+        , [ [121,221,321,421,521]
+          , [122,222,322,422,522]
+          , [123,223,323,423,523]
+          , [124,224,324,424,524]
+          ]
+        , [ [131,231,331,431,531]
+          , [132,232,332,432,532]
+          , [133,233,333,433,533]
+          , [134,234,334,434,534]
+          ]
+        ]
+
+example_subs_dim1 :: Assertion
+example_subs_dim1 =
+    assertEqual "" expected $
+      Tensor.subs (2 ::: VNil) $
+        Tensor.dim1 [1,2,3]
+  where
+    expected :: Tensor Nat1 (Tensor Nat1 Int)
+    expected = Tensor.dim1 [ Tensor.dim1 [1,2], Tensor.dim1 [2,3] ]
+
+example_subs_dim2 :: Assertion
+example_subs_dim2 =
+    assertEqual "" expected $
+      Tensor.subs (2 ::: 2 ::: VNil) $
+        Tensor.dim2 [[11,12,13],[21,22,23],[31,32,33]]
+  where
+    expected :: Tensor Nat2 (Tensor Nat2 Int)
+    expected = Tensor.dim2 [
+          [ Tensor.dim2 [[11,12],[21,22]], Tensor.dim2 [[12,13],[22,23]] ]
+        , [ Tensor.dim2 [[21,22],[31,32]], Tensor.dim2 [[22,23],[32,33]] ]
+        ]
+
+example_subs_dim3 :: Assertion
+example_subs_dim3 =
+    assertEqual "" expected $
+      Tensor.subs (2 ::: 2 ::: 2 ::: VNil) $
+        Tensor.dim3 [
+            [[111,112,113],[121,122,123],[131,132,133]]
+          , [[211,212,213],[221,222,223],[231,232,233]]
+          , [[311,312,313],[321,322,323],[331,332,333]]
+          ]
+  where
+    expected :: Tensor Nat3 (Tensor Nat3 Int)
+    expected = Tensor.dim3 [
+            [ [ Tensor.dim3 [[[111,112],[121,122]],[[211,212],[221,222]]]
+              , Tensor.dim3 [[[112,113],[122,123]],[[212,213],[222,223]]]
+              ]
+            , [ Tensor.dim3 [[[121,122],[131,132]],[[221,222],[231,232]]]
+              , Tensor.dim3 [[[122,123],[132,133]],[[222,223],[232,233]]]
+              ]
+            ]
+          , [ [ Tensor.dim3 [[[211,212],[221,222]],[[311,312],[321,322]]]
+              , Tensor.dim3 [[[212,213],[222,223]],[[312,313],[322,323]]]
+              ]
+            , [ Tensor.dim3 [[[221,222],[231,232]],[[321,322],[331,332]]]
+              , Tensor.dim3 [[[222,223],[232,233]],[[322,323],[332,333]]]
+              ]
+            ]
+        ]
+
+example_padWith :: Assertion
+example_padWith =
+    assertEqual "" expected $
+      Tensor.padWith 0 2 $ Tensor.dim2 [ [1, 2, 3], [4, 5, 6] ]
+  where
+    expected :: Tensor Nat2 Int
+    expected = Tensor.dim2 [
+          [ 0, 0, 0, 0, 0, 0, 0 ]
+        , [ 0, 0, 0, 0, 0, 0, 0 ]
+        , [ 0, 0, 1, 2, 3, 0, 0 ]
+        , [ 0, 0, 4, 5, 6, 0, 0 ]
+        , [ 0, 0, 0, 0, 0, 0, 0 ]
+        , [ 0, 0, 0, 0, 0, 0, 0 ]
+        ]
+
+example_padWith' :: Assertion
+example_padWith' =
+    assertEqual "" expected $
+      Tensor.padWith' 0 ((1, 1) ::: (2, 3) ::: VNil) (Tensor.dim2 [[1]])
+  where
+    expected :: Tensor Nat2 Int
+    expected = Tensor.dim2 [
+          [0,0,0,0,0,0]
+        , [0,0,1,0,0,0]
+        , [0,0,0,0,0,0]
+        ]
+
+{-------------------------------------------------------------------------------
+  Examples from the 3B1B video
+-------------------------------------------------------------------------------}
+
+example_3b1b ::
+     Tensor Nat1 TestValue -- ^ Input (padded)
+  -> Tensor Nat1 TestValue -- ^ Kernel
+  -> Tensor Nat1 TestValue -- ^ Expected result
+  -> Assertion
+example_3b1b input kernel result =
+    assertEqual "" result $
+      Tensor.convolve kernel input
+
+example_3b1b_simple :: Assertion
+example_3b1b_simple =
+    example_3b1b
+      (Tensor.padWith 0 2 $ Tensor.dim1 simpleInput)
+      (Tensor.dim1 simpleKernel)
+      (Tensor.dim1 simpleResult)
+
+example_3b1b_weightedDice :: Assertion
+example_3b1b_weightedDice =
+    example_3b1b
+      (Tensor.padWith 0 5 $ Tensor.dim1 weightedDiceInput)
+      (Tensor.dim1 weightedDiceKernel)
+      (Tensor.dim1 weightedDiceResult)
+
+example_3b1b_movingAverage :: Assertion
+example_3b1b_movingAverage =
+    example_3b1b
+      (Tensor.padWith 0 2 $ Tensor.dim1 movingAverageInput)
+      (Tensor.dim1 movingAverageKernel)
+      (Tensor.dim1 movingAverageResult)
+
+example_3b1b_movingWeightedAverage :: Assertion
+example_3b1b_movingWeightedAverage =
+    example_3b1b
+      (Tensor.padWith 0 2 $ Tensor.dim1 movingAverageInput)
+      (Tensor.dim1 movingWeightedAverageKernel)
+      (Tensor.dim1 movingWeightedAverageResult)
+
+{-------------------------------------------------------------------------------
+  Properties
+-------------------------------------------------------------------------------}
+
+-- | Distribute over a list of 0-D tensor is the identity
+prop_distrib_dim0 :: NonEmptyList Int -> Property
+prop_distrib_dim0 (getNonEmpty -> xs) =
+        Tensor.toLists (Tensor.distrib (map Tensor.scalar xs))
+    === xs
+
+-- | Distribute over a list of 1-D tensor is 'transpose'
+prop_distrib_dim1 :: NonEmptyList (NonEmptyList Int) -> Property
+prop_distrib_dim1 (getSameLength -> xss) =
+    counterexample ("input: " ++ show xss) $
+          Tensor.toLists (Tensor.distrib (map Tensor.dim1 xss))
+      === L.transpose xss
+
+-- | Counterpart to 'prop_distrib_dim1': this is only true for same-size lists
+prop_distrib_dim1_nonUniform :: NonEmptyList (NonEmptyList Int) -> Property
+prop_distrib_dim1_nonUniform (getNonEmpty2 -> xss) =
+    expectFailure $
+          Tensor.toLists (Tensor.distrib (map Tensor.dim1 xss))
+      === L.transpose xss
+
+{-------------------------------------------------------------------------------
+  Auxiliary
+-------------------------------------------------------------------------------}
+
+getNonEmpty2 :: NonEmptyList (NonEmptyList a) -> [[a]]
+getNonEmpty2 = map getNonEmpty . getNonEmpty
+
+getSameLength :: NonEmptyList (NonEmptyList a) -> [[a]]
+getSameLength = aux . getNonEmpty2
+  where
+    aux :: [[a]] -> [[a]]
+    aux xss = map (take (minimum $ map length xss)) xss
diff --git a/test/TestSuite/Test/Convolution/CUDNN.hs b/test/TestSuite/Test/Convolution/CUDNN.hs
new file mode 100644
--- /dev/null
+++ b/test/TestSuite/Test/Convolution/CUDNN.hs
@@ -0,0 +1,354 @@
+module TestSuite.Test.Convolution.CUDNN (tests) where
+
+import Data.List qualified as L
+import Data.Type.Nat
+import Data.Vec.Lazy (Vec(..))
+import Foreign
+import Foreign.C
+import System.IO.Unsafe (unsafePerformIO)
+import Test.Tasty
+import Test.Tasty.HUnit
+import Test.Tasty.QuickCheck
+
+import Test.Tensor (Tensor(..))
+import Test.Tensor qualified as Tensor
+import Test.Tensor.TestValue
+
+import TestSuite.Test.Convolution.Examples3B1B
+import TestSuite.Util.TestKernel
+
+{-------------------------------------------------------------------------------
+  Lists of tests
+-------------------------------------------------------------------------------}
+
+tests :: TestTree
+tests = testGroup "TestSuite.Test.Convolution.CUDNN" [
+      testGroup "Sanity" [
+            testCase "bindingVersion" test_bindingVersion
+          , testCase "libraryVersion" test_libraryVersion
+        ]
+    , testGroup "Examples" [
+          testCase "weightedMovingAverage" example_weightedMovingAverage
+        ]
+    , testGroup "Properties" [
+          testGroup "matchesModel" [
+              testGroup "1d" [
+                  testProperty "kernelSize2" $ prop_matchesModel_1d @Nat2
+                , testProperty "kernelSize3" $ prop_matchesModel_1d @Nat3
+                , testProperty "kernelSize4" $ prop_matchesModel_1d @Nat4
+                ]
+            , testProperty "4d" prop_matchesModel
+            ]
+        ]
+        , testProperty "mode" prop_mode
+    ]
+
+{-------------------------------------------------------------------------------
+  Sanity checks
+-------------------------------------------------------------------------------}
+
+-- | Confirm that basic FFI interaction works as expected
+test_bindingVersion :: Assertion
+test_bindingVersion =
+    assertEqual "" 1 $
+      c_test_cudnn_binding_version
+
+-- | Confirm cuDNN version (we expect at least version 9.0)
+test_libraryVersion :: Assertion
+test_libraryVersion =
+    if c_test_cudnn_library_version >= 90000
+      then return ()
+      else assertFailure "Expect cuDNN version 9.0 or higher"
+
+{-------------------------------------------------------------------------------
+  Examples
+-------------------------------------------------------------------------------}
+
+example_weightedMovingAverage :: Assertion
+example_weightedMovingAverage =
+    assertEqual "" (Tensor.dim1 $ movingWeightedAverageResult @TestValue) $
+      convolveCUDNN_1d
+        (Tensor.dim1 movingWeightedAverageKernel)
+        (Tensor.padWith 0 2 $ Tensor.dim1 $ movingAverageInput @TestValue)
+
+{-------------------------------------------------------------------------------
+  Properties
+
+  NOTE: cuDNN does not like it when the size of the image is smaller than
+  the size of the kernel.
+-------------------------------------------------------------------------------}
+
+-- | Compare our implementation against cuDNN, 1D case
+prop_matchesModel_1d :: forall w.
+     SNatI w
+  => TestKernel '[w] TestValue  -- ^ Kernel
+  -> Tensor Nat1 TestValue      -- ^ Input
+  -> Property
+prop_matchesModel_1d (testKernel -> kernel) input =
+    Tensor.sizeAtLeast (minWidth ::: VNil) input ==>
+          convolveCUDNN_1d kernel input
+      === Tensor.convolve kernel input
+  where
+    minWidth :: Int
+    minWidth = fromIntegral $ snatToNatural (snat @w)
+
+-- | Compare our implementation against cuDNN, general case
+prop_matchesModel :: ConvolutionParams TestValue -> Property
+prop_matchesModel params =
+        convolveCUDNN c_mode_cross_correlation stride kernels input
+    === convolve_cuDNN_style params
+  where
+    ConvolutionParams{stride, input, kernels} = params
+
+prop_mode :: ConvolutionParams TestValue -> Property
+prop_mode params =
+        convolveCUDNN
+          c_mode_cross_correlation
+          stride
+          kernels
+          input
+    === convolveCUDNN
+          c_mode_convolution
+          stride
+          ( Tensor.foreach kernels $ \outputFeature ->
+              Tensor.foreach outputFeature $ \inputFeature ->
+                Tensor.rotate inputFeature
+          )
+          input
+  where
+    ConvolutionParams{stride, input, kernels} = params
+
+{-------------------------------------------------------------------------------
+  Model
+-------------------------------------------------------------------------------}
+
+-- | cuDNN-style convolutions, but using our implementation
+convolve_cuDNN_style :: forall a.
+     (Fractional a, Real a)
+  => ConvolutionParams a -> Tensor Nat4 a
+convolve_cuDNN_style params =
+    Tensor.foreach input $ \(Tensor channels) -> Tensor [
+        fmap (L.foldl' (+) 0) . Tensor.distrib $
+          zipWith (Tensor.convolveWithStride stride') inputFeatures channels
+      | Tensor inputFeatures <- Tensor.getTensor kernels
+      ]
+  where
+    ConvolutionParams{stride = (sv, sh), input, kernels} = params
+
+    stride' :: Vec Nat2 Int
+    stride' = sv ::: sh ::: VNil
+
+-- | Convolution parameters
+--
+-- Although both the input and the output are 4D tensors, their structure is
+-- different:
+--
+-- * The input is NCHW:
+--   - N images
+--   - each image has C channels
+--   - height H and width W
+--
+-- * The output is KCRS:
+--   - K "output features"
+--   - C "input features"
+--   - height R and width S
+--
+-- For every input image we compute K output images ("channels"); each output
+-- image results from applying C 2D kernels to each channel, adding up the
+-- results. The result is an N*K*H*W tensor.
+data ConvolutionParams a = ConvolutionParams {
+      stride  :: (Int, Int)
+    , input   :: Tensor Nat4 a
+    , kernels :: Tensor Nat4 a
+    }
+  deriving stock (Show)
+
+instance (Arbitrary a, Num a, Eq a) => Arbitrary (ConvolutionParams a) where
+  arbitrary = sized $ \n -> do
+      numImages      <- choose (1, max 1 n)
+      inputFeatures  <- choose (1, 3)
+      outputFeatures <- choose (1, 3)
+      kernelHeight   <- choose (1, 5)
+      kernelWidth    <- choose (1, 5)
+      inputHeight    <- choose (kernelHeight, max kernelHeight n)
+      inputWidth     <- choose (kernelWidth,  max kernelWidth  n)
+
+      let inputSize :: Tensor.Size Nat4
+          inputSize = numImages
+                  ::: inputFeatures
+                  ::: inputHeight
+                  ::: inputWidth
+                  ::: VNil
+
+      let kernelSize :: Tensor.Size Nat4
+          kernelSize = outputFeatures
+                   ::: inputFeatures
+                   ::: kernelHeight
+                   ::: kernelWidth
+                   ::: VNil
+
+      stride  <- (,) <$> choose (1, 5) <*> choose (1, 5)
+      input   <- Tensor.arbitraryOfSize inputSize arbitrary
+      kernels <- Tensor.arbitraryOfSize kernelSize arbitrary
+
+      return ConvolutionParams {stride, input, kernels}
+
+  -- Shrinking is a bit complicated, because we need to maintain consistency
+  -- between the kernels and the input
+  shrink params = concat [
+        -- Shrink stride
+        [ params{stride = (sv', sh')}
+        | (sv', sh') <- shrink stride
+        , sv' > 0
+        , sh' > 0
+        ]
+
+        -- Shrink input size
+      , [ params{input = input', kernels = kernels'}
+        | axe <- Tensor.allAxes (Tensor.size input)
+        , let input'   = Tensor.axeWith axe input
+        , let kernels' = adjustKernels axe kernels
+
+          -- Image should not be smaller than the kernel
+        , let (_ ::: _ ::: ih ::: iw ::: VNil) = Tensor.size input'
+        , let (_ ::: _ ::: kh ::: kw ::: VNil) = Tensor.size kernels'
+        , ih >= kh
+        , iw >= kw
+        ]
+
+        -- Shrink the kernel
+      , [ params{input = input', kernels = kernels'}
+        | axe <- Tensor.allAxes (Tensor.size kernels)
+        , let kernels' = Tensor.axeWith axe kernels
+        , let input'   = adjustInput axe input
+        ]
+
+        -- Shrink input elements
+      , [ params{input = images'}
+        | images' <- Tensor.shrinkElem (Just Tensor.zero) shrink input
+        ]
+
+        -- Shrink kernel element
+      , [ params{kernels = outputFeatures'}
+        | outputFeatures' <- Tensor.shrinkElem Nothing shrink kernels
+        ]
+      ]
+    where
+      ConvolutionParams{stride, input, kernels} = params
+
+      -- Adjust each kernel after we axe some part of the input
+      adjustKernels :: Tensor.Axe Nat4 -> Tensor Nat4 a -> Tensor Nat4 a
+      adjustKernels (Tensor.AxeHere _) =
+          -- We dropped some images; kernel is unaffected
+          id
+      adjustKernels axe@(Tensor.AxeNested (Tensor.AxeHere _)) =
+          -- We dropped some input channels; also drop the corresponding
+          -- input features from the kernel
+          Tensor.axeWith axe
+      adjustKernels _otherwise =
+          -- We reduced image height or width; kernel is unaffected
+          -- (though we must check that the image is large enough now)
+          id
+
+      -- Adjust the input after we axe some of the kernels
+      adjustInput :: Tensor.Axe Nat4 -> Tensor Nat4 a -> Tensor Nat4 a
+      adjustInput (Tensor.AxeHere _) =
+          -- We dropped some output features; input is unaffected
+          id
+      adjustInput axe@(Tensor.AxeNested (Tensor.AxeHere _)) =
+          -- We dropped some input features; drop the corresponding channels
+          Tensor.axeWith axe
+      adjustInput _otherwise =
+          -- We shrunk the kernel size (height or width), input is unaffected
+          id
+
+{-------------------------------------------------------------------------------
+  Compute convolution using cuDNN
+-------------------------------------------------------------------------------}
+
+convolveCUDNN_1d :: forall a.
+     (Fractional a, Real a)
+  => Tensor Nat1 a -> Tensor Nat1 a -> Tensor Nat1 a
+convolveCUDNN_1d kernel input = extract1d $
+    convolveCUDNN
+      c_mode_cross_correlation
+      (1, 1)
+      (Tensor [Tensor [Tensor [kernel]]])
+      (Tensor [Tensor [Tensor [input]]])
+  where
+    extract1d :: Tensor Nat4 a -> Tensor Nat1 a
+    extract1d (Tensor [Tensor [Tensor [output]]]) = output
+    extract1d _ = error "convolveCUDNN_1d: unexpected output"
+
+convolveCUDNN ::
+     (Fractional a, Real a)
+  => CudnnConvolutionMode
+  -> (Int, Int)    -- ^ vertical and horizontal stride
+  -> Tensor Nat4 a -- ^ kernel
+  -> Tensor Nat4 a -- ^ input
+  -> Tensor Nat4 a
+convolveCUDNN mode (sv, sh) kernels input = unsafePerformIO $
+    Tensor.unsafeWithCArray (realToFrac <$> kernels) $ \kernelsPtr ->
+    Tensor.unsafeWithCArray (realToFrac <$> input)   $ \inputPtr   ->
+    alloca $ \outputHeightPtr ->
+    alloca $ \outputWidthPtr  -> do
+      outputPtr <-
+        c_test_cudnn_convolve
+          mode
+          (fromIntegral sv)
+          (fromIntegral sh)
+          (fromIntegral k)
+          (fromIntegral kh)
+          (fromIntegral kw)
+          kernelsPtr
+          (fromIntegral n)
+          (fromIntegral c)
+          (fromIntegral ih)
+          (fromIntegral iw)
+          inputPtr
+          outputHeightPtr
+          outputWidthPtr
+      oh <- fromIntegral <$> peek outputHeightPtr
+      ow <- fromIntegral <$> peek outputWidthPtr
+      outputFPtr <- newForeignPtr finalizerFree outputPtr
+      let outputSize = n ::: k ::: oh ::: ow ::: VNil
+      return $ realToFrac <$> Tensor.unsafeFromCArray outputSize outputFPtr
+  where
+    n ::: c ::: ih ::: iw ::: VNil = Tensor.size input
+    k ::: _ ::: kh ::: kw ::: VNil = Tensor.size kernels
+
+{-------------------------------------------------------------------------------
+  FFI imports
+-------------------------------------------------------------------------------}
+
+type CudnnConvolutionMode = CInt
+
+foreign import capi unsafe "test-cudnn.h test_cudnn_binding_version"
+  c_test_cudnn_binding_version :: Int
+
+foreign import capi unsafe "test-cudnn.h test_cudnn_library_version"
+  c_test_cudnn_library_version :: Int
+
+foreign import capi unsafe "cudnn.h value CUDNN_CONVOLUTION"
+  c_mode_convolution :: CudnnConvolutionMode
+
+foreign import capi unsafe "cudnn.h value CUDNN_CROSS_CORRELATION"
+  c_mode_cross_correlation :: CudnnConvolutionMode
+
+foreign import capi unsafe "test-cudnn.h test_cudnn_convolve"
+  c_test_cudnn_convolve ::
+       CudnnConvolutionMode
+    -> CInt       -- ^ vertical_stride
+    -> CInt       -- ^ horizontal_stride
+    -> CInt       -- ^ num_kernels
+    -> CInt       -- ^ kernel_height
+    -> CInt       -- ^ kernel_width
+    -> Ptr Float  -- ^ kernel
+    -> CInt       -- ^ num_images
+    -> CInt       -- ^ input_channels
+    -> CInt       -- ^ input_height
+    -> CInt       -- ^ input_width
+    -> Ptr Float  -- ^ input
+    -> Ptr CInt   -- ^ output_height
+    -> Ptr CInt   -- ^ output_width
+    -> IO (Ptr Float)
diff --git a/test/TestSuite/Test/Convolution/Examples3B1B.hs b/test/TestSuite/Test/Convolution/Examples3B1B.hs
new file mode 100644
--- /dev/null
+++ b/test/TestSuite/Test/Convolution/Examples3B1B.hs
@@ -0,0 +1,99 @@
+-- | Examples from the 3Blue1Brown video on convolutions
+--
+-- See "But what is a convolution?", <https://www.youtube.com/watch?v=KuXjwB4LzSA>
+module TestSuite.Test.Convolution.Examples3B1B (
+    -- * Simple example
+    simpleInput
+  , simpleKernel
+  , simpleResult
+    -- * Weighted dice
+  , weightedDiceInput
+  , weightedDiceKernel
+  , weightedDiceResult
+    -- * Moving average
+  , movingAverageInput
+  , movingAverageKernel
+  , movingWeightedAverageKernel
+  , movingAverageResult
+  , movingWeightedAverageResult
+  ) where
+
+{-------------------------------------------------------------------------------
+  Simple example
+
+  In this example the input/kernel distinction is somewhat artificial.
+  We rotate the kernel.
+-------------------------------------------------------------------------------}
+
+simpleInput :: Num a => [a]
+simpleInput = [1, 2, 3]
+
+simpleKernel :: Num a => [a]
+simpleKernel = reverse [4, 5, 6]
+
+simpleResult :: Num a => [a]
+simpleResult = [4, 13, 28, 27, 18]
+
+{-------------------------------------------------------------------------------
+  Weighted dice
+
+  Same comments as for the simple example apply.
+-------------------------------------------------------------------------------}
+
+weightedDiceInput :: Fractional a => [a]
+weightedDiceInput = [0.03, 0.11, 0.23, 0.29, 0.23, 0.11]
+
+weightedDiceKernel :: Fractional a => [a]
+weightedDiceKernel = reverse [0.46, 0.20, 0.12, 0.09, 0.07, 0.05]
+
+weightedDiceResult :: Fractional a => [a]
+weightedDiceResult = [
+      0.01 -- 2
+    , 0.06 -- 3
+    , 0.13 -- 4
+    , 0.20 -- 5
+    , 0.21 -- 6
+    , 0.16 -- 7
+    , 0.10 -- 8
+    , 0.07 -- 9
+    , 0.04 -- 10
+    , 0.02 -- 11
+    , 0.01 -- 12
+    ]
+
+{-------------------------------------------------------------------------------
+  Moving average
+-------------------------------------------------------------------------------}
+
+movingAverageInput :: Fractional a => [a]
+movingAverageInput = concat [
+      replicate 5 0.1
+    , replicate 5 1.0
+    , replicate 5 0.1
+    , replicate 5 1.0
+    , replicate 5 0.1
+    ]
+
+movingAverageKernel :: Fractional a => [a]
+movingAverageKernel = [0.2, 0.2, 0.2, 0.2, 0.2]
+
+movingWeightedAverageKernel :: Fractional a => [a]
+movingWeightedAverageKernel = [0.1, 0.2, 0.4, 0.2, 0.1]
+
+movingAverageResult :: Fractional a => [a]
+movingAverageResult = [
+      0.06, 0.08, 0.10, 0.28, 0.46
+    , 0.64, 0.82, 1.00, 0.82, 0.64
+    , 0.46, 0.28, 0.10, 0.28, 0.46
+    , 0.64, 0.82, 1.00, 0.82, 0.64
+    , 0.46, 0.28, 0.10, 0.08, 0.06
+    ]
+
+movingWeightedAverageResult :: Fractional a => [a]
+movingWeightedAverageResult = [
+      0.07, 0.09, 0.10, 0.19, 0.37
+    , 0.73, 0.91, 1.00, 0.91, 0.73
+    , 0.37, 0.19, 0.10, 0.19, 0.37
+    , 0.73, 0.91, 1.00, 0.91, 0.73
+    , 0.37, 0.19, 0.10, 0.09, 0.07
+    ]
diff --git a/test/TestSuite/Test/Convolution/FFT.hs b/test/TestSuite/Test/Convolution/FFT.hs
new file mode 100644
--- /dev/null
+++ b/test/TestSuite/Test/Convolution/FFT.hs
@@ -0,0 +1,136 @@
+-- | Test against a reference implementation using fast fourier transforms
+--
+-- We do this only for 1D tensors.
+module TestSuite.Test.Convolution.FFT (tests) where
+
+import Data.Array.CArray (CArray)
+import Data.Array.IArray (IArray)
+import Data.Array.IArray qualified as IA
+import Data.Complex (Complex)
+import Data.Ix (Ix)
+import Data.Type.Nat
+import Math.FFT qualified as FFT
+import Test.Tasty
+import Test.Tasty.HUnit
+import Test.Tasty.QuickCheck
+
+import Test.Tensor qualified as Tensor
+import Test.Tensor.TestValue
+
+import TestSuite.Test.Convolution.Examples3B1B
+import TestSuite.Util.TestKernel
+
+{-------------------------------------------------------------------------------
+  List of testse
+-------------------------------------------------------------------------------}
+
+tests :: TestTree
+tests = testGroup "Test.Convolution.FFT" [
+      testGroup "Examples" [
+          testCase "weightedMovingAverage" example_weightedMovingAverage
+        ]
+    , testGroup "Properties" [
+          testGroup "matchesModel" [
+              testProperty "kernelSize3" $ prop_matchesModel @Nat3
+            , testProperty "kernelSize4" $ prop_matchesModel @Nat4
+            , testProperty "kernelSize5" $ prop_matchesModel @Nat5
+            ]
+        ]
+    ]
+
+{-------------------------------------------------------------------------------
+  Examples
+-------------------------------------------------------------------------------}
+
+example_weightedMovingAverage :: Assertion
+example_weightedMovingAverage =
+    assertEqual "" (movingWeightedAverageResult @TestValue) $
+      removePadding 2 $
+        convolveFFT
+          movingWeightedAverageKernel
+          (movingAverageInput @TestValue)
+
+{-------------------------------------------------------------------------------
+  Properties
+-------------------------------------------------------------------------------}
+
+-- | Compare our implementation against FFT implementation
+prop_matchesModel :: forall n.
+     TestKernel '[n] TestValue  -- ^ Kernel
+  -> NonEmptyList TestValue     -- ^ Input
+  -> Property
+prop_matchesModel (testKernel -> kernel) (getNonEmpty -> input) =
+        convolveFFT (reverse $ Tensor.toLists kernel) input
+    === ( Tensor.toLists $
+            Tensor.convolve
+              kernel
+              (Tensor.padWith 0 (length kernel - 1) $ Tensor.dim1 input)
+        )
+
+{-------------------------------------------------------------------------------
+  Convolution implementation using FFT
+
+  FFT requires an input of even length, so if the input has odd length, we add
+  an additional zero padding byte.
+-------------------------------------------------------------------------------}
+
+-- | Compute convolution using FFT
+convolveFFT :: forall a. (Fractional a, Real a) => [a] -> [a] -> [a]
+convolveFFT kernel input_ =
+    adjustOutput needOddAdjustment $ map realToFrac $ IA.elems inv
+  where
+    needOddAdjustment :: Bool
+    needOddAdjustment = odd (length input_ + length kernel - 1)
+
+    input :: [a]
+    input = adjustInput needOddAdjustment input_
+
+    n, m :: Int
+    n = length input
+    m = length kernel
+
+    arrInput, arrKernel :: CArray Int Double
+    arrInput  = paddedArrayFromList (m + n - 1) (map realToFrac input)
+    arrKernel = paddedArrayFromList (m + n - 1) (map realToFrac kernel)
+
+    dftInput, dftKernel, dftMult :: CArray Int (Complex Double)
+    dftInput  = FFT.dftRC arrInput
+    dftKernel = FFT.dftRC arrKernel
+    dftMult   = zipArraySameBounds (*) dftInput dftKernel
+
+    inv :: CArray Int Double
+    inv = FFT.dftCR dftMult
+
+adjustInput :: Num a => Bool -> [a] -> [a]
+adjustInput True  = (:) 0
+adjustInput False = id
+
+adjustOutput :: Bool -> [a] -> [a]
+adjustOutput True  = drop 1
+adjustOutput False = id
+
+removePadding :: Int -> [a] -> [a]
+removePadding n xs = take (length xs - 2 * n) (drop n xs)
+
+{-------------------------------------------------------------------------------
+  Internal auxiliary: arrays
+-------------------------------------------------------------------------------}
+
+paddedArrayFromList :: forall a e. (IArray a e, Num e)
+  => Int  -- ^ Decided length of the array
+  -> [e]  -- ^ List to initialize the array from
+  -> a Int e
+paddedArrayFromList len xs = IA.listArray (0, len - 1) (xs ++ repeat 0)
+
+zipArraySameBounds ::
+     (IArray a x, IArray a y, IArray a z, Ix i)
+  => (x -> y -> z)
+  -> a i x -> a i y -> a i z
+zipArraySameBounds f xs ys =
+    IA.listArray (IA.bounds xs) [
+        f (xs IA.! i) (ys IA.! i)
+      | i <- IA.indices xs
+      ]
+
+
+
diff --git a/test/TestSuite/Test/QuickCheck.hs b/test/TestSuite/Test/QuickCheck.hs
new file mode 100644
--- /dev/null
+++ b/test/TestSuite/Test/QuickCheck.hs
@@ -0,0 +1,108 @@
+-- | Meta-tests: test the Tensor QuickCheck infrastructure
+module TestSuite.Test.QuickCheck (tests) where
+
+import Data.Foldable qualified as Foldable
+import Data.Type.Nat
+import Data.Vec.Lazy (Vec(..))
+import Test.Tasty
+import Test.Tasty.HUnit
+import Test.Tasty.QuickCheck
+
+import Test.Tensor (Tensor)
+import Test.Tensor qualified as Tensor
+
+{-------------------------------------------------------------------------------
+  List of tests
+-------------------------------------------------------------------------------}
+
+tests :: TestTree
+tests = testGroup "TestSuite.Test.QuickCheck" [
+      testGroup "Examples" [
+          testCase "shrinkWith" example_shrinkWith
+        ]
+    , testGroup "Properties" [
+          testProperty "allAxes_shrinkList" prop_allAxes_shrinkList
+        , testProperty "axeSize" prop_axeSize
+        , testProperty "length_zeroWith" prop_length_zeroWith
+        ]
+    ]
+
+{-------------------------------------------------------------------------------
+  Examples
+-------------------------------------------------------------------------------}
+
+example_shrinkWith :: Assertion
+example_shrinkWith =
+    assertEqual "" expected $
+      Tensor.shrinkWith
+        (Just $ Tensor.Zero (-1))
+        (const [0])
+        (Tensor.dim2 [[1,2,3], [4,5,6]])
+  where
+    expected :: [Tensor.Tensor Nat2 Int]
+    expected = [
+          -- Shrink outer dimension
+          Tensor.dim2 [[4,5,6]]
+        , Tensor.dim2 [[1,2,3]]
+          -- Shrink inner dimension
+        , Tensor.dim2 [[2,3],[5,6]]
+        , Tensor.dim2 [[1,3],[4,6]]
+        , Tensor.dim2 [[1,2],[4,5]]
+          -- Zero outer dimension
+        , Tensor.dim2 [[-1,-1,-1],[4,5,6]]
+        , Tensor.dim2 [[1,2,3],[-1,-1,-1]]
+          -- Zero inner dimension
+        , Tensor.dim2 [[-1,2,3],[-1,5,6]]
+        , Tensor.dim2 [[1,-1,3],[4,-1,6]]
+        , Tensor.dim2 [[1,2,-1],[4,5,-1]]
+          -- Shrink one of the elements
+        , Tensor.dim2 [[0,2,3],[4,5,6]]
+        , Tensor.dim2 [[1,0,3],[4,5,6]]
+        , Tensor.dim2 [[1,2,0],[4,5,6]]
+        , Tensor.dim2 [[1,2,3],[0,5,6]]
+        , Tensor.dim2 [[1,2,3],[4,0,6]]
+        , Tensor.dim2 [[1,2,3],[4,5,0]]
+        ]
+
+{-------------------------------------------------------------------------------
+  Properties
+-------------------------------------------------------------------------------}
+
+-- | 'allAxes' essentially reifies the decisions made by 'shrinkList'
+prop_allAxes_shrinkList :: NonEmptyList Int -> Property
+prop_allAxes_shrinkList (getNonEmpty -> xs) =
+    counterexample ("tensor: " ++ show tensor) $
+    counterexample ("size: " ++ show size) $
+          filter (not . null) (shrinkList (const []) xs)
+      === [ Foldable.toList $ Tensor.axeWith axe tensor
+          | axe <- Tensor.allAxes size
+          ]
+  where
+    tensor :: Tensor Nat1 Int
+    tensor = Tensor.fromList (length xs ::: VNil) xs
+
+    size :: Tensor.Size Nat1
+    size = Tensor.size tensor
+
+prop_axeSize :: Tensor Nat2 Int -> Property
+prop_axeSize tensor = conjoin [
+      counterexample ("axe: " ++ show axe) $
+            length (Tensor.axeWith axe tensor)
+        === length tensor - Tensor.axeSize size axe
+    | axe <- Tensor.allAxes size
+    ]
+  where
+    size :: Tensor.Size Nat2
+    size = Tensor.size tensor
+
+prop_length_zeroWith :: Tensor Nat2 Int -> Property
+prop_length_zeroWith tensor = conjoin [
+      counterexample ("axe: " ++ show axe) $
+        case Tensor.zeroWith Tensor.zero axe tensor of
+          Nothing      -> property True
+          Just tensor' -> length tensor' === length tensor
+    | axe <- Tensor.allAxes size
+    ]
+  where
+    size :: Tensor.Size Nat2
+    size = Tensor.size tensor
diff --git a/test/TestSuite/Test/StdOps.hs b/test/TestSuite/Test/StdOps.hs
new file mode 100644
--- /dev/null
+++ b/test/TestSuite/Test/StdOps.hs
@@ -0,0 +1,45 @@
+module TestSuite.Test.StdOps (tests) where
+
+import Data.Foldable qualified as Foldable
+import Data.Type.Nat
+import Test.Tasty
+import Test.Tasty.QuickCheck
+
+import Test.Tensor (Tensor)
+import Test.Tensor qualified as Tensor
+
+{-------------------------------------------------------------------------------
+  List of tests
+-------------------------------------------------------------------------------}
+
+tests :: TestTree
+tests = testGroup "TestSuite.Test.StdOps" [
+      testGroup "properties" [
+            testGroup "fromList_toList" [
+                testProperty "dim0" $ prop_fromList_toList @Nat0
+              , testProperty "dim1" $ prop_fromList_toList @Nat1
+              , testProperty "dim2" $ prop_fromList_toList @Nat2
+              , testProperty "dim3" $
+                  withMaxSuccess 100 $ -- random 3D tensors get large quick
+                    prop_fromList_toList @Nat3
+              ]
+          , testProperty "distrib_transpose" $ prop_distrib_transpose
+        ]
+    ]
+
+{-------------------------------------------------------------------------------
+  Properties
+-------------------------------------------------------------------------------}
+
+prop_fromList_toList :: SNatI n => Tensor n Int -> Property
+prop_fromList_toList tensor =
+        Tensor.fromList (Tensor.size tensor) (Foldable.toList tensor)
+    === tensor
+
+prop_distrib_transpose :: Tensor Nat2 Int -> Property
+prop_distrib_transpose tensor =
+        (restructure . Tensor.distrib . Tensor.getTensor $ tensor)
+    === (Tensor.transpose $ tensor)
+  where
+    restructure :: Tensor Nat1 [Int] -> Tensor Nat2 Int
+    restructure = Tensor.fromLists . Tensor.toLists
diff --git a/test/TestSuite/Util/TestKernel.hs b/test/TestSuite/Util/TestKernel.hs
new file mode 100644
--- /dev/null
+++ b/test/TestSuite/Util/TestKernel.hs
@@ -0,0 +1,62 @@
+-- | Test kernels
+--
+-- For testing purposes it's very useful to be able to specify at the type level
+-- the exact size of kernel we want (not just its dimension).
+--
+-- Notes:
+--
+-- * Use sites will always pick a specific size, so we're not worried here about
+--   stuck type families etc.
+-- * Size we specify the size of the kernel at the type level, the size of the
+--   kernel does not shrink (in the 'Arbitrary' instance).
+--
+-- Intended for unqualified import.
+module TestSuite.Util.TestKernel (
+    TestKernel -- opaque
+  , testKernel
+  ) where
+
+import Data.Kind
+import Data.Type.Nat
+import Data.Vec.Lazy (Vec(..))
+import Data.Vec.Lazy qualified as Vec
+import Test.QuickCheck
+
+import Test.Tensor (Tensor(..))
+
+{-------------------------------------------------------------------------------
+  Definition
+-------------------------------------------------------------------------------}
+
+data TestKernel :: [Nat] -> Type -> Type where
+  TKZ :: a -> TestKernel '[] a
+  TKS :: Vec n (TestKernel ns a) -> TestKernel (n : ns) a
+
+instance Show a => Show (TestKernel ns a) where
+  show = show . testKernel
+
+{-------------------------------------------------------------------------------
+  Conversion
+-------------------------------------------------------------------------------}
+
+type family Length (as :: [k]) where
+  Length '[]    = Z
+  Length (x:xs) = S (Length xs)
+
+testKernel :: TestKernel ns a -> Tensor (Length ns) a
+testKernel (TKZ x)  = Scalar x
+testKernel (TKS xs) = Tensor $ map testKernel (Vec.toList xs)
+
+{-------------------------------------------------------------------------------
+  Arbitrary instance
+-------------------------------------------------------------------------------}
+
+instance Arbitrary a => Arbitrary (TestKernel '[] a) where
+  arbitrary      = TKZ <$> arbitrary
+  shrink (TKZ x) = TKZ <$> shrink x
+
+instance (SNatI n, Arbitrary (TestKernel ns a))
+      => Arbitrary (TestKernel (n : ns) a) where
+  arbitrary       = TKS <$> liftArbitrary arbitrary
+  shrink (TKS xs) = TKS <$> shrink xs
+
diff --git a/testing-tensor.cabal b/testing-tensor.cabal
new file mode 100644
--- /dev/null
+++ b/testing-tensor.cabal
@@ -0,0 +1,108 @@
+cabal-version:   3.0
+name:            testing-tensor
+version:         0.1.0
+license:         BSD-3-Clause
+license-file:    LICENSE
+author:          Edsko de Vries
+maintainer:      edsko@well-typed.com
+category:        Testing
+build-type:      Simple
+synopsis:        Pure implementation of tensors, for use in tests.
+description:     This is a pure Haskell implementation of tensors, emphasizing
+                 simplify over all else. It is intended to be used as a model
+                 in tests.
+extra-doc-files: CHANGELOG.md
+tested-with:     GHC ==9.2.8
+                 GHC ==9.4.8
+                 GHC ==9.6.6
+                 GHC ==9.8.4
+                 GHC ==9.10.1
+
+source-repository head
+  type:     git
+  location: https://github.com/well-typed/testing-tensor
+
+common lang
+  build-depends:    base >= 4.16 && < 5
+  default-language: GHC2021
+
+  ghc-options:
+      -Wall
+      -Wprepositive-qualified-module
+      -Wunused-packages
+      -Widentities
+      -Wno-unticked-promoted-constructors
+
+  default-extensions:
+      CApiFFI
+      DataKinds
+      DerivingStrategies
+      LambdaCase
+      TypeFamilies
+      ViewPatterns
+
+library
+  import:          lang
+  hs-source-dirs:  src
+
+  exposed-modules:
+      Test.Tensor
+      Test.Tensor.TestValue
+
+  build-depends:
+    , fin          >= 0.3  && < 0.4
+    , QuickCheck   >= 2.15 && < 2.16
+    , random       >= 1.2  && < 1.4
+    , transformers >= 0.5  && < 0.7
+    , vec          >= 0.5  && < 0.6
+    , vector       >= 0.13 && < 0.14
+
+test-suite testing-tensor-test
+  import:         lang
+  type:           exitcode-stdio-1.0
+  hs-source-dirs: test
+  main-is:        Main.hs
+  build-depends:  testing-tensor
+
+  build-depends:
+    , tasty            >= 1.5  && < 1.6
+    , tasty-hunit      >= 0.10 && < 0.11
+    , tasty-quickcheck >= 0.11 && < 0.12
+
+  -- inherited dependencies
+  build-depends:
+    , fin
+    , QuickCheck
+    , vec
+
+  other-modules:
+      TestSuite.Test.Convolution
+      TestSuite.Test.Convolution.Examples3B1B
+      TestSuite.Test.QuickCheck
+      TestSuite.Test.StdOps
+      TestSuite.Util.TestKernel
+
+  if flag(test-fft)
+    cpp-options:   -DTEST_FFT
+    other-modules: TestSuite.Test.Convolution.FFT
+    build-depends:
+      , array  >= 0.5 && < 0.6
+      , carray >= 0.1 && < 0.2
+      , fft    >= 0.1 && < 0.2
+
+  if flag(test-cudnn)
+    cpp-options:     -DTEST_CUDNN
+    other-modules:   TestSuite.Test.Convolution.CUDNN
+    include-dirs:    test-cbits
+    c-sources:       test-cbits/test-cudnn.c
+    extra-libraries: cudart cudnn
+
+Flag test-fft
+  description: Test against an FFT implementation
+  default: False
+  manual: True
+
+Flag test-cudnn
+  description: Test against cuDNN
+  default: False
+  manual: True
