testing-tensor (empty) → 0.1.0
raw patch · 14 files changed
+2185/−0 lines, 14 filesdep +QuickCheckdep +arraydep +base
Dependencies added: QuickCheck, array, base, carray, fft, fin, random, tasty, tasty-hunit, tasty-quickcheck, testing-tensor, transformers, vec, vector
Files
- CHANGELOG.md +5/−0
- LICENSE +29/−0
- src/Test/Tensor.hs +617/−0
- src/Test/Tensor/TestValue.hs +119/−0
- test-cbits/test-cudnn.c +230/−0
- test/Main.hs +32/−0
- test/TestSuite/Test/Convolution.hs +241/−0
- test/TestSuite/Test/Convolution/CUDNN.hs +354/−0
- test/TestSuite/Test/Convolution/Examples3B1B.hs +99/−0
- test/TestSuite/Test/Convolution/FFT.hs +136/−0
- test/TestSuite/Test/QuickCheck.hs +108/−0
- test/TestSuite/Test/StdOps.hs +45/−0
- test/TestSuite/Util/TestKernel.hs +62/−0
- testing-tensor.cabal +108/−0
+ CHANGELOG.md view
@@ -0,0 +1,5 @@+# Revision history for tmp++## 0.1.0.0 -- YYYY-mm-dd++* First version. Released on an unsuspecting world.
+ LICENSE view
@@ -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.
+ src/Test/Tensor.hs view
@@ -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
+ src/Test/Tensor/TestValue.hs view
@@ -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)
+ test-cbits/test-cudnn.c view
@@ -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;+}
+ test/Main.hs view
@@ -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+ ]+ ]
+ test/TestSuite/Test/Convolution.hs view
@@ -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
+ test/TestSuite/Test/Convolution/CUDNN.hs view
@@ -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)
+ test/TestSuite/Test/Convolution/Examples3B1B.hs view
@@ -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+ ]
+ test/TestSuite/Test/Convolution/FFT.hs view
@@ -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+ ]+++
+ test/TestSuite/Test/QuickCheck.hs view
@@ -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
+ test/TestSuite/Test/StdOps.hs view
@@ -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
+ test/TestSuite/Util/TestKernel.hs view
@@ -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+
+ testing-tensor.cabal view
@@ -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