diff --git a/CHANGELOG b/CHANGELOG --- a/CHANGELOG +++ b/CHANGELOG @@ -1,5 +1,10 @@ # Changelog +- 0.3.3 (2026-11-14) + * Adds four-block kernels producing interleaved keystream blocks, + as used by ChaCha8Rand, and roughly twice as fast as the + single-block kernels. + - 0.3.2 (2026-11-14) * Adds Crypto.Cipher.ChaCha, with ChaCha keystreams of 8, 12, or 20 rounds. diff --git a/bench/Main.hs b/bench/Main.hs --- a/bench/Main.hs +++ b/bench/Main.hs @@ -12,7 +12,10 @@ import qualified Crypto.Cipher.ChaCha20 as ChaCha20 import qualified Data.ByteString as BS import qualified Data.ByteString.Base16 as B16 +import qualified Data.ByteString.Internal as BI +import qualified Data.ByteString.Unsafe as BU import Data.Maybe (fromJust) +import Foreign.Ptr (castPtr) import GHC.Generics deriving instance Generic ChaCha20.Error @@ -23,6 +26,7 @@ main = defaultMain [ suite , keystream + , interleaved ] plain :: BS.ByteString @@ -59,3 +63,15 @@ , bench "ChaCha20" $ nf (ChaCha.keystream ChaCha.ChaCha20 key 1 non) 1024 ] + +interleaved :: Benchmark +interleaved = + bgroup "ppad-chacha (interleaved, 1024B)" [ + bench "ChaCha8" $ nfIO (inter ChaCha.ChaCha8 ChaCha.FeedForwardKey) + , bench "ChaCha20" $ nfIO (inter ChaCha.ChaCha20 ChaCha.FeedForwardAll) + ] + where + inter r ff = BI.create 1024 $ \dst -> + BU.unsafeUseAsCString key $ \kp -> + BU.unsafeUseAsCString non $ \np -> + ChaCha._interleaved_into r ff (castPtr kp) 1 (castPtr np) dst 4 diff --git a/cbits/chacha20_native.c b/cbits/chacha20_native.c --- a/cbits/chacha20_native.c +++ b/cbits/chacha20_native.c @@ -188,6 +188,68 @@ } } +/* + * ChaCha NEON kernel using inter-block parallelism: four consecutive + * blocks at once, with register x[i] holding word i of each block. A + * quarter-round is then one set of element-wise operations on four + * registers, with no lane rotation, and the four quarter-rounds of + * each half-round are independent chains. The registers, stored in + * order, give the blocks interleaved word by word. + */ +static inline void chacha20_core4v(uint32x4_t x[16], int drounds) { + for (int i = 0; i < drounds; i++) { + QUARTER(x[0], x[4], x[ 8], x[12]); + QUARTER(x[1], x[5], x[ 9], x[13]); + QUARTER(x[2], x[6], x[10], x[14]); + QUARTER(x[3], x[7], x[11], x[15]); + QUARTER(x[0], x[5], x[10], x[15]); + QUARTER(x[1], x[6], x[11], x[12]); + QUARTER(x[2], x[7], x[ 8], x[13]); + QUARTER(x[3], x[4], x[ 9], x[14]); + } +} + +/* + * Write 'ngroups' groups of four consecutive keystream blocks to + * 'out', 256 bytes per group, using 'drounds' double rounds and + * starting at block 'counter'. Within a group the blocks are + * interleaved word by word: word i of block j is at byte offset + * 16 * i + 4 * j. If 'key_ff_only' is nonzero, the final addition of + * the input state is applied to the key words only, omitting the + * public constant, counter, and nonce words, as in ChaCha8Rand. + * 'out' must not overlap 'key' or 'nonce'. + */ +void chacha20_interleaved_native(const uint8_t key[32], uint32_t counter, + const uint8_t nonce[12], uint8_t *out, + size_t ngroups, int drounds, + int key_ff_only) { + uint32_t k[8], n[3]; + memcpy(k, key, 32); + memcpy(n, nonce, 12); + + uint32x4_t s[16]; + for (int i = 0; i < 4; i++) s[i] = vdupq_n_u32(chacha_constants[i]); + for (int i = 0; i < 8; i++) s[4 + i] = vdupq_n_u32(k[i]); + for (int i = 0; i < 3; i++) s[13 + i] = vdupq_n_u32(n[i]); + static const uint32_t lanes[4] = { 0, 1, 2, 3 }; + const uint32x4_t inc = vld1q_u32(lanes); + + for (size_t g = 0; g < ngroups; g++) { + s[12] = vaddq_u32(vdupq_n_u32(counter), inc); + uint32x4_t x[16]; + for (int i = 0; i < 16; i++) x[i] = s[i]; + chacha20_core4v(x, drounds); + + uint8_t *o = out + 256 * g; + for (int i = 0; i < 16; i++) { + int add = !key_ff_only || (i >= 4 && i < 12); + uint32x4_t v = add ? vaddq_u32(x[i], s[i]) : x[i]; + vst1q_u8(o + 16 * i, vreinterpretq_u8_u32(v)); + } + counter += 4; + } +} + int chacha20_native_available(void) { return 1; } @@ -577,6 +639,135 @@ } } +/* ChaCha core over eight consecutive blocks, one state word per + * register, as in the four-block SSE2 kernel on wider registers. */ +__attribute__((target("avx2"))) +static inline void chacha20_core8v_avx2(__m256i x[16], int drounds) { + for (int i = 0; i < drounds; i++) { + QUARTER_AVX2(x[0], x[4], x[ 8], x[12]); + QUARTER_AVX2(x[1], x[5], x[ 9], x[13]); + QUARTER_AVX2(x[2], x[6], x[10], x[14]); + QUARTER_AVX2(x[3], x[7], x[11], x[15]); + QUARTER_AVX2(x[0], x[5], x[10], x[15]); + QUARTER_AVX2(x[1], x[6], x[11], x[12]); + QUARTER_AVX2(x[2], x[7], x[ 8], x[13]); + QUARTER_AVX2(x[3], x[4], x[ 9], x[14]); + } +} + +/* + * Write as many pairs of interleaved four-block groups as 'ngroups' + * allows, eight blocks per pass: the low 128-bit lane of each + * register holds the first group's word, the high lane the second's. + * Returns the number of groups written. + */ +__attribute__((target("avx2"))) +static size_t chacha20_interleaved_avx2(const uint8_t key[32], + uint32_t counter, + const uint8_t nonce[12], + uint8_t *out, size_t ngroups, + int drounds, int key_ff_only) { + uint32_t k[8], n[3]; + memcpy(k, key, 32); + memcpy(n, nonce, 12); + + __m256i s[16]; + for (int i = 0; i < 4; i++) + s[i] = _mm256_set1_epi32((int)chacha_constants[i]); + for (int i = 0; i < 8; i++) s[4 + i] = _mm256_set1_epi32((int)k[i]); + for (int i = 0; i < 3; i++) s[13 + i] = _mm256_set1_epi32((int)n[i]); + const __m256i inc = _mm256_setr_epi32(0, 1, 2, 3, 4, 5, 6, 7); + + size_t g = 0; + for (; g + 2 <= ngroups; g += 2) { + s[12] = _mm256_add_epi32(_mm256_set1_epi32((int)counter), inc); + __m256i x[16]; + for (int i = 0; i < 16; i++) x[i] = s[i]; + chacha20_core8v_avx2(x, drounds); + + uint8_t *o = out + 256 * g; + for (int i = 0; i < 16; i++) { + int add = !key_ff_only || (i >= 4 && i < 12); + __m256i v = add ? _mm256_add_epi32(x[i], s[i]) : x[i]; + _mm_storeu_si128((__m128i *)(o + 16 * i), + _mm256_castsi256_si128(v)); + _mm_storeu_si128((__m128i *)(o + 256 + 16 * i), + _mm256_extracti128_si256(v, 1)); + } + counter += 8; + } + return g; +} + +/* + * ChaCha SSE2 kernel using inter-block parallelism, as in the NEON + * kernel of the same name: register x[i] holds word i of each of + * four consecutive blocks. + */ +static inline void chacha20_core4v(__m128i x[16], int drounds) { + for (int i = 0; i < drounds; i++) { + QUARTER(x[0], x[4], x[ 8], x[12]); + QUARTER(x[1], x[5], x[ 9], x[13]); + QUARTER(x[2], x[6], x[10], x[14]); + QUARTER(x[3], x[7], x[11], x[15]); + QUARTER(x[0], x[5], x[10], x[15]); + QUARTER(x[1], x[6], x[11], x[12]); + QUARTER(x[2], x[7], x[ 8], x[13]); + QUARTER(x[3], x[4], x[ 9], x[14]); + } +} + +/* + * Write 'ngroups' groups of four consecutive keystream blocks to + * 'out', 256 bytes per group, using 'drounds' double rounds and + * starting at block 'counter'. Within a group the blocks are + * interleaved word by word: word i of block j is at byte offset + * 16 * i + 4 * j. If 'key_ff_only' is nonzero, the final addition of + * the input state is applied to the key words only, omitting the + * public constant, counter, and nonce words, as in ChaCha8Rand. + * 'out' must not overlap 'key' or 'nonce'. + */ +void chacha20_interleaved_native(const uint8_t key[32], uint32_t counter, + const uint8_t nonce[12], uint8_t *out, + size_t ngroups, int drounds, + int key_ff_only) { + uint32_t k[8], n[3]; + memcpy(k, key, 32); + memcpy(n, nonce, 12); + + /* pairs of groups via AVX2 where available, the rest via SSE2 */ + if (ngroups >= 2 && have_avx2()) { + size_t done = chacha20_interleaved_avx2(key, counter, nonce, out, + ngroups, drounds, + key_ff_only); + counter += (uint32_t)(4 * done); + out += 256 * done; + ngroups -= done; + } + + __m128i s[16]; + for (int i = 0; i < 4; i++) + s[i] = _mm_set1_epi32((int)chacha_constants[i]); + for (int i = 0; i < 8; i++) s[4 + i] = _mm_set1_epi32((int)k[i]); + for (int i = 0; i < 3; i++) s[13 + i] = _mm_set1_epi32((int)n[i]); + const __m128i inc = _mm_setr_epi32(0, 1, 2, 3); + + for (size_t g = 0; g < ngroups; g++) { + s[12] = _mm_add_epi32(_mm_set1_epi32((int)counter), inc); + __m128i x[16]; + for (int i = 0; i < 16; i++) x[i] = s[i]; + chacha20_core4v(x, drounds); + + uint8_t *o = out + 256 * g; + for (int i = 0; i < 16; i++) { + int add = !key_ff_only || (i >= 4 && i < 12); + __m128i v = add ? _mm_add_epi32(x[i], s[i]) : x[i]; + _mm_storeu_si128((__m128i *)(o + 16 * i), v); + } + counter += 4; + } +} + int chacha20_native_available(void) { return 1; } @@ -598,6 +789,14 @@ size_t inlen, int drounds) { (void)key; (void)counter; (void)nonce; (void)in; (void)out; (void)inlen; (void)drounds; +} + +void chacha20_interleaved_native(const uint8_t *key, uint32_t counter, + const uint8_t *nonce, uint8_t *out, + size_t ngroups, int drounds, + int key_ff_only) { + (void)key; (void)counter; (void)nonce; (void)out; + (void)ngroups; (void)drounds; (void)key_ff_only; } int chacha20_native_available(void) { diff --git a/lib/Crypto/Cipher/ChaCha.hs b/lib/Crypto/Cipher/ChaCha.hs --- a/lib/Crypto/Cipher/ChaCha.hs +++ b/lib/Crypto/Cipher/ChaCha.hs @@ -25,20 +25,26 @@ , Error(..) -- internal + , FeedForward(..) , _keystream_into + , _interleaved_into + , _interleaved_pure_into ) where import Control.Monad.ST (runST) +import Data.Bits ((.|.), (.<<.), (.>>.)) import Crypto.Cipher.ChaCha20 (Error(..)) import qualified Crypto.Cipher.ChaCha20 as ChaCha20 import qualified Crypto.Cipher.ChaCha20.Native as Native import qualified Data.ByteString as BS import qualified Data.ByteString.Internal as BI import qualified Data.ByteString.Unsafe as BU +import Data.Foldable (for_) import Data.Word (Word8, Word32, Word64) import Foreign.ForeignPtr (withForeignPtr) import Foreign.Marshal.Utils (copyBytes) import Foreign.Ptr (Ptr, castPtr, plusPtr) +import Foreign.Storable (pokeByteOff) -- utils ---------------------------------------------------------------------- @@ -130,3 +136,93 @@ (ChaCha20._parse_nonce non) (BS.replicate len 0) BU.unsafeUseAsCString out $ \src -> copyBytes dst (castPtr src) len + +-- interleaved keystream ------------------------------------------------------ + +-- | Which words of the input state are added to a block's output. +data FeedForward = + FeedForwardAll -- ^ all words, as standard + | FeedForwardKey -- ^ the key words only, omitting the public + -- constant, counter, and nonce words, as in + -- ChaCha8Rand + deriving (Eq, Show) + +-- | Write groups of four consecutive keystream blocks to the +-- destination pointer, interleaved word by word: within each +-- 256-byte group, word i of block j is at byte offset 16 * i + 4 * j. +-- This is the layout that four-way SIMD produces naturally. +-- +-- Unchecked: the key and nonce pointers must address 32 and 12 +-- bytes respectively, the destination must have room for 256 bytes +-- per group and must not overlap the key or nonce, and the counter +-- must leave room for four blocks per group. +_interleaved_into + :: Rounds -- ^ number of rounds + -> FeedForward -- ^ final addition of the input state + -> Ptr Word8 -- ^ 256-bit key + -> Word32 -- ^ 32-bit counter + -> Ptr Word8 -- ^ 96-bit nonce + -> Ptr Word8 -- ^ destination + -> Int -- ^ number of four-block groups + -> IO () +_interleaved_into r ff kp counter np dst groups + | Native.chacha20_native_available = + Native.interleaved_into (drounds r) (ff == FeedForwardKey) kp + counter np dst groups + | otherwise = _interleaved_pure_into r ff kp counter np dst groups + +-- | A pure Haskell '_interleaved_into', with the same preconditions. +_interleaved_pure_into + :: Rounds + -> FeedForward + -> Ptr Word8 + -> Word32 + -> Ptr Word8 + -> Ptr Word8 + -> Int + -> IO () +_interleaved_pure_into r ff kp counter np dst groups = do + key <- BS.packCStringLen (castPtr kp, 32) + non <- BS.packCStringLen (castPtr np, 12) + let !ks = runST $ + ChaCha20._cipher_rounds (drounds r) + (ChaCha20._parse_key key) counter (ChaCha20._parse_nonce non) + (BS.replicate (256 * groups) 0) + -- the input word that the final addition omits + omit i blk = case ff of + FeedForwardAll -> 0 + FeedForwardKey + | i < 4 -> constant i + | i == 12 -> counter + fi blk + | i > 12 -> le32_at non (4 * (i - 13)) + | otherwise -> 0 + for_ [0 .. groups - 1] $ \g -> + for_ [0 .. 15] $ \i -> + for_ [0 .. 3] $ \j -> do + let !blk = 4 * g + j + !v = le32_at ks (64 * blk + 4 * i) - omit i blk + poke_le32 (dst `plusPtr` (256 * g + 16 * i + 4 * j)) v + where + constant :: Int -> Word32 + constant i = case i of + 0 -> 0x61707865 + 1 -> 0x3320646e + 2 -> 0x79622d32 + _ -> 0x6b206574 + +-- little-endian word at the given (in-bounds) offset +le32_at :: BS.ByteString -> Int -> Word32 +le32_at bs o = + fi (BU.unsafeIndex bs o) + .|. (fi (BU.unsafeIndex bs (o + 1)) .<<. 8) + .|. (fi (BU.unsafeIndex bs (o + 2)) .<<. 16) + .|. (fi (BU.unsafeIndex bs (o + 3)) .<<. 24) +{-# INLINE le32_at #-} + +poke_le32 :: Ptr Word8 -> Word32 -> IO () +poke_le32 p w = do + pokeByteOff p 0 (fi w :: Word8) + pokeByteOff p 1 (fi (w .>>. 8) :: Word8) + pokeByteOff p 2 (fi (w .>>. 16) :: Word8) + pokeByteOff p 3 (fi (w .>>. 24) :: Word8) +{-# INLINE poke_le32 #-} diff --git a/lib/Crypto/Cipher/ChaCha20/Native.hs b/lib/Crypto/Cipher/ChaCha20/Native.hs --- a/lib/Crypto/Cipher/ChaCha20/Native.hs +++ b/lib/Crypto/Cipher/ChaCha20/Native.hs @@ -14,6 +14,7 @@ , block , cipher , keystream_into + , interleaved_into ) where import qualified Data.ByteString as BS @@ -36,6 +37,11 @@ :: Ptr Word8 -> Word32 -> Ptr Word8 -> Ptr Word8 -> Ptr Word8 -> CSize -> CInt -> IO () +foreign import ccall unsafe "chacha20_interleaved_native" + c_chacha20_interleaved + :: Ptr Word8 -> Word32 -> Ptr Word8 -> Ptr Word8 -> CSize -> CInt + -> CInt -> IO () + foreign import ccall unsafe "chacha20_native_available" c_chacha20_native_available :: IO CInt @@ -97,3 +103,16 @@ keystream_into drounds kp counter np dst len = do fillBytes dst 0 len c_chacha20_cipher kp counter np dst dst (fi len) (fi drounds) + +-- | Write 'groups' groups of four keystream blocks, interleaved word +-- by word, using the given number of double rounds, for the 32-byte +-- key and 12-byte nonce at the given pointers, to the destination, +-- which must not overlap them. If the flag is set, the final +-- addition of the input state covers the key words only. The +-- counter must leave room for four blocks per group. +interleaved_into + :: Int -> Bool -> Ptr Word8 -> Word32 -> Ptr Word8 -> Ptr Word8 -> Int + -> IO () +interleaved_into drounds key_only kp counter np dst groups = + c_chacha20_interleaved kp counter np dst (fi groups) (fi drounds) + (if key_only then 1 else 0) diff --git a/ppad-chacha.cabal b/ppad-chacha.cabal --- a/ppad-chacha.cabal +++ b/ppad-chacha.cabal @@ -1,6 +1,6 @@ cabal-version: 3.0 name: ppad-chacha -version: 0.3.2 +version: 0.3.3 synopsis: A fast ChaCha20 stream cipher license: MIT license-file: LICENSE diff --git a/test/Main.hs b/test/Main.hs --- a/test/Main.hs +++ b/test/Main.hs @@ -17,11 +17,11 @@ import Data.Foldable (for_) import Data.Maybe (fromJust) import qualified Data.Primitive.PrimArray as PA -import Data.Word (Word32) -import Foreign.Ptr (castPtr) +import Data.Word (Word8, Word32) +import Foreign.Ptr (Ptr, castPtr) import Test.Tasty import qualified Test.Tasty.HUnit as H -import Vectors (V(..), vectors) +import Vectors (V(..), vectors, chacha8rand_seed, chacha8rand_iteration) main :: IO () main = defaultMain $ testGroup "ppad-chacha" [ @@ -42,6 +42,11 @@ , keystream_dispatch , keystream_into ] + , testGroup "interleaved" [ + interleaved_chacha8rand + , interleaved_layout + , interleaved_dispatch + ] ] quarter :: TestTree @@ -287,3 +292,55 @@ BU.unsafeUseAsCString block_non $ \np -> Reduced._keystream_into r (castPtr kp) 7 (castPtr np) dst 1000 H.assertEqual (show r) ref out + +-- interleaved + +type Interleave = + Reduced.Rounds -> Reduced.FeedForward -> Ptr Word8 -> Word32 + -> Ptr Word8 -> Ptr Word8 -> Int -> IO () + +interleave + :: Interleave -> Reduced.Rounds -> Reduced.FeedForward -> BS.ByteString + -> Word32 -> BS.ByteString -> Int -> IO BS.ByteString +interleave f r ff key con non groups = + BI.create (256 * groups) $ \dst -> + BU.unsafeUseAsCString key $ \kp -> + BU.unsafeUseAsCString non $ \np -> + f r ff (castPtr kp) con (castPtr np) dst groups + +-- ChaCha8Rand's iteration is sixteen ChaCha8 blocks, interleaved, with +-- the key-only final addition; its first 992 bytes are output +interleaved_chacha8rand :: TestTree +interleaved_chacha8rand = H.testCase "ChaCha8Rand iteration" $ do + out <- interleave Reduced._interleaved_into Reduced.ChaCha8 + Reduced.FeedForwardKey chacha8rand_seed 0 (BS.replicate 12 0) 4 + H.assertEqual mempty (hex chacha8rand_iteration) (BS.take 992 out) + +-- with the full final addition, the output is the standard keystream +-- with each group of four blocks transposed +interleaved_layout :: TestTree +interleaved_layout = H.testCase "layout (vs keystream)" $ + for_ [Reduced.ChaCha8, Reduced.ChaCha12, Reduced.ChaCha20] $ \r -> do + let Right ks = Reduced.keystream r block_key 7 block_non (256 * 3) + word b i = BS.take 4 (BS.drop (64 * b + 4 * i) ks) + ex = BS.concat [ word (4 * g + j) i + | g <- [0 .. 2], i <- [0 .. 15], j <- [0 .. 3] ] + out <- interleave Reduced._interleaved_into r Reduced.FeedForwardAll + block_key 7 block_non 3 + H.assertEqual (show r) ex out + +-- the dispatching '_interleaved_into' agrees with the pure Haskell +-- implementation; the second counter puts the final block of six +-- groups at the maximum counter value +interleaved_dispatch :: TestTree +interleaved_dispatch = H.testCase "dispatch vs pure" $ + for_ [Reduced.ChaCha8, Reduced.ChaCha12, Reduced.ChaCha20] $ \r -> + for_ [Reduced.FeedForwardAll, Reduced.FeedForwardKey] $ \ff -> + for_ [block_key, BS.replicate 32 0xff] $ \key -> + for_ [1, maxBound - 23] $ \con -> + for_ [1 .. 6] $ \groups -> do + out <- interleave Reduced._interleaved_into r ff key con block_non + groups + ref <- interleave Reduced._interleaved_pure_into r ff key con + block_non groups + H.assertEqual (show (r, ff, con, groups)) ref out diff --git a/test/Vectors.hs b/test/Vectors.hs --- a/test/Vectors.hs +++ b/test/Vectors.hs @@ -8,6 +8,8 @@ module Vectors ( V(..) , vectors + , chacha8rand_seed + , chacha8rand_iteration ) where import Crypto.Cipher.ChaCha (Rounds(..)) @@ -239,4 +241,44 @@ , "e5fbc34e60a1d9a9db17345b0a402736853bf910b060bdf1f897b6290f01d138" , "ae2c4c90225ba9ea14d518f55929dea098ca7a6ccfe61227053c84e49a4a3332" ]) + ] + +-- the first iteration of ChaCha8Rand output (992 bytes) for this seed, +-- from the sample output in etc/chacha8rand.md +chacha8rand_seed :: BS.ByteString +chacha8rand_seed = "ABCDEFGHIJKLMNOPQRSTUVWXYZ123456" + +chacha8rand_iteration :: BS.ByteString +chacha8rand_iteration = mconcat [ + "a516463d06b673b73cbc6aa622af60117c288d41d999258cd65cdc7e037ee07e" + , "ad161c2de09eaacf79eaeb8fef0e090e3e5b8b1271d2823c4fa35212c1dd5a9c" + , "a6ee6c7d61bb79df6a739d1f5953d5361fe01e184ed1f0ee3684e50a76fc9b08" + , "68d22acc592be5d9ba8a1b4b44b42feb6166c492a6c8954f2ce6ca173232c6c3" + , "7e2e4e7f36b4eb91c69beca0c6f24c7820beea345cec345ca8da7055518f0a4f" + , "f26b3d11b4dc35fcbc5455644ca40d5be1d921dba33d966df300e55031fcaeee" + , "a55037da3d92372db0c84b6d627a0d3849eed7e3ed68afee7c713b88956635f4" + , "a495243921906a841ba630965410858e93e4ba5d5402dc18433d5aa6f09f8f0f" + , "80f00f19f765f0cc30336739aad176fd3364ba6c9332d29517bd0c07d156746c" + , "62658cffdacf2a466afc346d86abaf5b88290a03782f860c3d16c307e4839ad3" + , "bf2ef2457b2b0ac0a9b16624c6074356d472c0b024047e25fe286c49995eb56f" + , "e0d45c8fa87398ae73370dc62a3657468e3ef2cd1ef9831cc08753c19207dc6f" + , "5c2bfd0da3d2da36b7bd95052993b5a4c502cce43489e14da7e315f004d6c0cd" + , "2103d89af6dba0fb87de39d1a3bee8605687f41e854d8ad14af315227c446663" + , "ee07d0d3972e6805b24ad5c678890e4cc2ed12676a9f1ecfd3cf804c41391406" + , "ad0e5c74e2b6a8d1e810c4458d91a731ec6e2190ad61ccab1aa732202dd94040" + , "68cd40fb6ff6d23c7a859572c051d0dc7e52abd9bc5cb5eaacda1b78ce1d4718" + , "5272dc44d18cf0f7d1407f3db1e004585bd3b2e446a4b15c068a13d228a7d4e6" + , "d8da60ca403e2205686aac0632ec612db8174c87562369ab1cde7676415409c3" + , "2456223237ce1a4f388398130851a9fb112e75520f207f99fa2162e8afaa1611" + , "1935a1b25c3bce07148345bc72bc562978eb406192b788414deaadd4bf6dca56" + , "e50c344923c2e37f8a5f67379c8fc03521d55eeffbc7e11175bcc14e46c8ad98" + , "f803123dc7b263d1f3f3a243e01e768cb7d7eccc6a9db92470032f11aa313e79" + , "395128192adc878e256e09f704ae47420fe26f9222314f51e9a7d245f0b36fdc" + , "ba8ed1cdce30cb15740290f6ecfdc7cb21a08bdc96c6b53fe674d2c8174466d1" + , "787245ea45e4f70557b69d1bcabb20f999cb22dab450190c92e209aff1ba75f8" + , "38f85042b8d7d3be6041d70f08e898f103a77e9b1da5edc9f6f89b4355ef09f7" + , "fc16f1ebfe740cd246756d14eb6856308797d810ecf39a82bc1d557b69f9b815" + , "c9b8648e6c3c82fcbc403b18e8855534681358d671414b67f7e970d61cd83412" + , "195ea5d81052500edca0ecee698d25e87ef6bae852c4d4059955a41601e3dbe8" + , "006f17b1e18cf01c49cb1eb8a4d0f7cc0e432c6b13ea3f3071c8069c136c1d86" ]