luna - wolfssl/wolfcrypt/src/port/cuda/aes-cuda.cu

   1/* aes.cu
   2 *
   3 * Copyright (C) 2006-2026 wolfSSL Inc.
   4 *
   5 * This file is part of wolfSSL.
   6 *
   7 * wolfSSL is free software; you can redistribute it and/or modify
   8 * it under the terms of the GNU General Public License as published by
   9 * the Free Software Foundation; either version 3 of the License, or
  10 * (at your option) any later version.
  11 *
  12 * wolfSSL is distributed in the hope that it will be useful,
  13 * but WITHOUT ANY WARRANTY; without even the implied warranty of
  14 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
  15 * GNU General Public License for more details.
  16 *
  17 * You should have received a copy of the GNU General Public License
  18 * along with this program; if not, write to the Free Software
  19 * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1335, USA
  20 */
  21
  22/*
  23
  24DESCRIPTION
  25This library provides the interfaces to the Advanced Encryption Standard (AES)
  26for encrypting and decrypting data. AES is the standard known for a symmetric
  27block cipher mechanism that uses n-bit binary string parameter key with 128-bits,
  28192-bits, and 256-bits of key sizes.
  29
  30*/
  31#ifdef HAVE_CONFIG_H
  32    #include <config.h>
  33#endif
  34
  35#include <wolfssl/wolfcrypt/settings.h>
  36#include <wolfssl/wolfcrypt/error-crypt.h>
  37
  38#include <wolfssl/wolfcrypt/aes.h>
  39
  40#ifdef WOLFSSL_AESNI
  41#include <wmmintrin.h>
  42#include <emmintrin.h>
  43#include <smmintrin.h>
  44#endif /* WOLFSSL_AESNI */
  45
  46#include <wolfssl/wolfcrypt/cpuid.h>
  47
  48#ifdef WOLF_CRYPTO_CB
  49    #include <wolfssl/wolfcrypt/cryptocb.h>
  50#endif
  51
  52#ifdef WOLFSSL_SECO_CAAM
  53#include <wolfssl/wolfcrypt/port/caam/wolfcaam.h>
  54#endif
  55
  56#ifdef WOLFSSL_IMXRT_DCP
  57    #include <wolfssl/wolfcrypt/port/nxp/dcp_port.h>
  58#endif
  59#if defined(WOLFSSL_SE050) && defined(WOLFSSL_SE050_CRYPT)
  60    #include <wolfssl/wolfcrypt/port/nxp/se050_port.h>
  61#endif
  62
  63#if defined(WOLFSSL_AES_SIV)
  64    #include <wolfssl/wolfcrypt/cmac.h>
  65#endif /* WOLFSSL_AES_SIV */
  66
  67#if defined(WOLFSSL_HAVE_PSA) && !defined(WOLFSSL_PSA_NO_AES)
  68    #include <wolfssl/wolfcrypt/port/psa/psa.h>
  69#endif
  70
  71
  72#include <wolfssl/wolfcrypt/logging.h>
  73
  74#ifdef NO_INLINE
  75    #include <wolfssl/wolfcrypt/misc.h>
  76#else
  77    #define WOLFSSL_MISC_INCLUDED
  78    #define WOLFSSL_HAVE_MIN
  79    #define WOLFSSL_HAVE_MAX
  80/*    #include <wolfcrypt/src/misc.c> */
  81#endif
  82/* This routine performs a left circular arithmetic shift of <x> by <y> value. */
  83
  84extern "C" {
  85
  86#if !defined(WOLFSSL_TI_CRYPT)
  87
  88#define rotlFixed(x, y) ( (x << y) | (x >> (sizeof(x) * 8 - y)) )
  89
  90/* This routine performs a right circular arithmetic shift of <x> by <y> value. */
  91#define rotrFixed(x, y) ( (x >> y) | (x << (sizeof(x) * 8 - y)) )
  92
  93#ifdef WC_RC2
  94
  95/* This routine performs a left circular arithmetic shift of <x> by <y> value */
  96static WC_INLINE word16 rotlFixed16(word16 x, word16 y)
  97{
  98    return (x << y) | (x >> (sizeof(x) * 8 - y));
  99}
 100
 101
 102/* This routine performs a right circular arithmetic shift of <x> by <y> value */
 103static WC_INLINE word16 rotrFixed16(word16 x, word16 y)
 104{
 105    return (x >> y) | (x << (sizeof(x) * 8 - y));
 106}
 107
 108#endif /* WC_RC2 */
 109
 110/* This routine performs a byte swap of 32-bit word value. */
 111#if defined(__CCRX__) && !defined(NO_INLINE) /* shortest version for CC-RX */
 112    #define ByteReverseWord32(value, outRef) ( *outRef = _builtin_revl(value) )
 113#else
 114__device__
 115static WC_INLINE word32 ByteReverseWord32(word32 value)
 116{
 117#ifdef PPC_INTRINSICS
 118    /* PPC: load reverse indexed instruction */
 119    return (word32)__lwbrx(&value,0);
 120#elif defined(__ICCARM__)
 121    return (word32)__REV(value);
 122#elif defined(KEIL_INTRINSICS)
 123    return (word32)__rev(value);
 124#elif defined(__CCRX__)
 125    return (word32)_builtin_revl(value);
 126#elif defined(WOLF_ALLOW_BUILTIN) && \
 127        defined(__GNUC_PREREQ) && __GNUC_PREREQ(4, 3)
 128    return (word32)__builtin_bswap32(value);
 129#elif defined(WOLFSSL_BYTESWAP32_ASM) && defined(__GNUC__) && \
 130      defined(__aarch64__)
 131    __asm__ volatile (
 132        "REV32 %0, %0  \n"
 133        : "+r" (value)
 134        :
 135    );
 136    return value;
 137#elif defined(WOLFSSL_BYTESWAP32_ASM) && defined(__GNUC__) && \
 138      (defined(__thumb__) || defined(__arm__))
 139    __asm__ volatile (
 140        "REV %0, %0  \n"
 141        : "+r" (value)
 142        :
 143    );
 144    return value;
 145#elif defined(FAST_ROTATE)
 146    /* 5 instructions with rotate instruction, 9 without */
 147    return (rotrFixed(value, 8U) & 0xff00ff00) |
 148           (rotlFixed(value, 8U) & 0x00ff00ff);
 149#else
 150    /* 6 instructions with rotate instruction, 8 without */
 151    value = ((value & 0xFF00FF00) >> 8) | ((value & 0x00FF00FF) << 8);
 152    return rotlFixed(value, 16U);
 153#endif
 154}
 155#endif /* ! (__CCRX__ && !NO_INLINE) */
 156
 157#if defined(STM32_CRYPTO)
 158#elif defined(HAVE_COLDFIRE_SEC)
 159#elif defined(FREESCALE_LTC)
 160#elif defined(FREESCALE_MMCAU)
 161#elif defined(WOLFSSL_PIC32MZ_CRYPT)
 162#elif defined(WOLFSSL_NRF51_AES)
 163#elif defined(WOLFSSL_ESP32_CRYPT) && \
 164     !defined(NO_WOLFSSL_ESP32_CRYPT_AES)
 165#elif defined(WOLFSSL_AESNI)
 166#elif (defined(WOLFSSL_IMX6_CAAM) && !defined(NO_IMX6_CAAM_AES) \
 167        && !defined(WOLFSSL_QNX_CAAM)) || \
 168      ((defined(WOLFSSL_AFALG) || defined(WOLFSSL_DEVCRYPTO_AES)) && \
 169        defined(HAVE_AESCCM))
 170#elif defined(WOLFSSL_AFALG)
 171    /* implemented in wolfcrypt/src/port/af_alg/afalg_aes.c */
 172
 173#elif defined(WOLFSSL_DEVCRYPTO_AES)
 174    /* implemented in wolfcrypt/src/port/devcrypto/devcrypto_aes.c */
 175
 176#elif defined(WOLFSSL_SCE) && !defined(WOLFSSL_SCE_NO_AES)
 177#elif defined(WOLFSSL_KCAPI_AES)
 178#elif defined(WOLFSSL_HAVE_PSA) && !defined(WOLFSSL_PSA_NO_AES)
 179/* implemented in wolfcrypt/src/port/psa/psa_aes.c */
 180
 181#else
 182
 183    /* using wolfCrypt software implementation */
 184    #define NEED_AES_TABLES
 185#endif
 186
 187#if !defined(NO_AES) && !defined(WOLFSSL_TI_CRYPT) && !defined(WOLFSSL_ARMASM) && \
 188    defined(NEED_AES_TABLES) && (defined(HAVE_AES_CBC) || defined(WOLFSSL_AES_DIRECT) || defined(HAVE_AESCCM) || defined(HAVE_AESGCM)) && \
 189    defined(HAVE_CUDA)
 190
 191#define GETBYTE(x, y) (word32)((byte)((x) >> (8 * (y))))
 192
 193#ifndef WC_CACHE_LINE_SZ
 194    #if defined(__x86_64__) || defined(_M_X64) || \
 195       (defined(__ILP32__) && (__ILP32__ >= 1))
 196        #define WC_CACHE_LINE_SZ 64
 197    #else
 198        /* default cache line size */
 199        #define WC_CACHE_LINE_SZ 32
 200    #endif
 201#endif
 202
 203#if WC_CACHE_LINE_SZ == 128
 204    #define WC_CACHE_LINE_BITS      5
 205    #define WC_CACHE_LINE_MASK_HI   0xe0
 206    #define WC_CACHE_LINE_MASK_LO   0x1f
 207    #define WC_CACHE_LINE_ADD       0x20
 208#elif WC_CACHE_LINE_SZ == 64
 209    #define WC_CACHE_LINE_BITS      4
 210    #define WC_CACHE_LINE_MASK_HI   0xf0
 211    #define WC_CACHE_LINE_MASK_LO   0x0f
 212    #define WC_CACHE_LINE_ADD       0x10
 213#elif WC_CACHE_LINE_SZ == 32
 214    #define WC_CACHE_LINE_BITS      3
 215    #define WC_CACHE_LINE_MASK_HI   0xf8
 216    #define WC_CACHE_LINE_MASK_LO   0x07
 217    #define WC_CACHE_LINE_ADD       0x08
 218#elif WC_CACHE_LINE_SZ == 16
 219    #define WC_CACHE_LINE_BITS      2
 220    #define WC_CACHE_LINE_MASK_HI   0xfc
 221    #define WC_CACHE_LINE_MASK_LO   0x03
 222    #define WC_CACHE_LINE_ADD       0x04
 223#else
 224    #error Cache line size not supported
 225#endif
 226
 227#ifndef WOLFSSL_AES_SMALL_TABLES
 228__device__
 229const FLASH_QUALIFIER word32 Te_CUDA[4][256] = {
 230{
 231    0xc66363a5U, 0xf87c7c84U, 0xee777799U, 0xf67b7b8dU,
 232    0xfff2f20dU, 0xd66b6bbdU, 0xde6f6fb1U, 0x91c5c554U,
 233    0x60303050U, 0x02010103U, 0xce6767a9U, 0x562b2b7dU,
 234    0xe7fefe19U, 0xb5d7d762U, 0x4dababe6U, 0xec76769aU,
 235    0x8fcaca45U, 0x1f82829dU, 0x89c9c940U, 0xfa7d7d87U,
 236    0xeffafa15U, 0xb25959ebU, 0x8e4747c9U, 0xfbf0f00bU,
 237    0x41adadecU, 0xb3d4d467U, 0x5fa2a2fdU, 0x45afafeaU,
 238    0x239c9cbfU, 0x53a4a4f7U, 0xe4727296U, 0x9bc0c05bU,
 239    0x75b7b7c2U, 0xe1fdfd1cU, 0x3d9393aeU, 0x4c26266aU,
 240    0x6c36365aU, 0x7e3f3f41U, 0xf5f7f702U, 0x83cccc4fU,
 241    0x6834345cU, 0x51a5a5f4U, 0xd1e5e534U, 0xf9f1f108U,
 242    0xe2717193U, 0xabd8d873U, 0x62313153U, 0x2a15153fU,
 243    0x0804040cU, 0x95c7c752U, 0x46232365U, 0x9dc3c35eU,
 244    0x30181828U, 0x379696a1U, 0x0a05050fU, 0x2f9a9ab5U,
 245    0x0e070709U, 0x24121236U, 0x1b80809bU, 0xdfe2e23dU,
 246    0xcdebeb26U, 0x4e272769U, 0x7fb2b2cdU, 0xea75759fU,
 247    0x1209091bU, 0x1d83839eU, 0x582c2c74U, 0x341a1a2eU,
 248    0x361b1b2dU, 0xdc6e6eb2U, 0xb45a5aeeU, 0x5ba0a0fbU,
 249    0xa45252f6U, 0x763b3b4dU, 0xb7d6d661U, 0x7db3b3ceU,
 250    0x5229297bU, 0xdde3e33eU, 0x5e2f2f71U, 0x13848497U,
 251    0xa65353f5U, 0xb9d1d168U, 0x00000000U, 0xc1eded2cU,
 252    0x40202060U, 0xe3fcfc1fU, 0x79b1b1c8U, 0xb65b5bedU,
 253    0xd46a6abeU, 0x8dcbcb46U, 0x67bebed9U, 0x7239394bU,
 254    0x944a4adeU, 0x984c4cd4U, 0xb05858e8U, 0x85cfcf4aU,
 255    0xbbd0d06bU, 0xc5efef2aU, 0x4faaaae5U, 0xedfbfb16U,
 256    0x864343c5U, 0x9a4d4dd7U, 0x66333355U, 0x11858594U,
 257    0x8a4545cfU, 0xe9f9f910U, 0x04020206U, 0xfe7f7f81U,
 258    0xa05050f0U, 0x783c3c44U, 0x259f9fbaU, 0x4ba8a8e3U,
 259    0xa25151f3U, 0x5da3a3feU, 0x804040c0U, 0x058f8f8aU,
 260    0x3f9292adU, 0x219d9dbcU, 0x70383848U, 0xf1f5f504U,
 261    0x63bcbcdfU, 0x77b6b6c1U, 0xafdada75U, 0x42212163U,
 262    0x20101030U, 0xe5ffff1aU, 0xfdf3f30eU, 0xbfd2d26dU,
 263    0x81cdcd4cU, 0x180c0c14U, 0x26131335U, 0xc3ecec2fU,
 264    0xbe5f5fe1U, 0x359797a2U, 0x884444ccU, 0x2e171739U,
 265    0x93c4c457U, 0x55a7a7f2U, 0xfc7e7e82U, 0x7a3d3d47U,
 266    0xc86464acU, 0xba5d5de7U, 0x3219192bU, 0xe6737395U,
 267    0xc06060a0U, 0x19818198U, 0x9e4f4fd1U, 0xa3dcdc7fU,
 268    0x44222266U, 0x542a2a7eU, 0x3b9090abU, 0x0b888883U,
 269    0x8c4646caU, 0xc7eeee29U, 0x6bb8b8d3U, 0x2814143cU,
 270    0xa7dede79U, 0xbc5e5ee2U, 0x160b0b1dU, 0xaddbdb76U,
 271    0xdbe0e03bU, 0x64323256U, 0x743a3a4eU, 0x140a0a1eU,
 272    0x924949dbU, 0x0c06060aU, 0x4824246cU, 0xb85c5ce4U,
 273    0x9fc2c25dU, 0xbdd3d36eU, 0x43acacefU, 0xc46262a6U,
 274    0x399191a8U, 0x319595a4U, 0xd3e4e437U, 0xf279798bU,
 275    0xd5e7e732U, 0x8bc8c843U, 0x6e373759U, 0xda6d6db7U,
 276    0x018d8d8cU, 0xb1d5d564U, 0x9c4e4ed2U, 0x49a9a9e0U,
 277    0xd86c6cb4U, 0xac5656faU, 0xf3f4f407U, 0xcfeaea25U,
 278    0xca6565afU, 0xf47a7a8eU, 0x47aeaee9U, 0x10080818U,
 279    0x6fbabad5U, 0xf0787888U, 0x4a25256fU, 0x5c2e2e72U,
 280    0x381c1c24U, 0x57a6a6f1U, 0x73b4b4c7U, 0x97c6c651U,
 281    0xcbe8e823U, 0xa1dddd7cU, 0xe874749cU, 0x3e1f1f21U,
 282    0x964b4bddU, 0x61bdbddcU, 0x0d8b8b86U, 0x0f8a8a85U,
 283    0xe0707090U, 0x7c3e3e42U, 0x71b5b5c4U, 0xcc6666aaU,
 284    0x904848d8U, 0x06030305U, 0xf7f6f601U, 0x1c0e0e12U,
 285    0xc26161a3U, 0x6a35355fU, 0xae5757f9U, 0x69b9b9d0U,
 286    0x17868691U, 0x99c1c158U, 0x3a1d1d27U, 0x279e9eb9U,
 287    0xd9e1e138U, 0xebf8f813U, 0x2b9898b3U, 0x22111133U,
 288    0xd26969bbU, 0xa9d9d970U, 0x078e8e89U, 0x339494a7U,
 289    0x2d9b9bb6U, 0x3c1e1e22U, 0x15878792U, 0xc9e9e920U,
 290    0x87cece49U, 0xaa5555ffU, 0x50282878U, 0xa5dfdf7aU,
 291    0x038c8c8fU, 0x59a1a1f8U, 0x09898980U, 0x1a0d0d17U,
 292    0x65bfbfdaU, 0xd7e6e631U, 0x844242c6U, 0xd06868b8U,
 293    0x824141c3U, 0x299999b0U, 0x5a2d2d77U, 0x1e0f0f11U,
 294    0x7bb0b0cbU, 0xa85454fcU, 0x6dbbbbd6U, 0x2c16163aU,
 295},
 296{
 297    0xa5c66363U, 0x84f87c7cU, 0x99ee7777U, 0x8df67b7bU,
 298    0x0dfff2f2U, 0xbdd66b6bU, 0xb1de6f6fU, 0x5491c5c5U,
 299    0x50603030U, 0x03020101U, 0xa9ce6767U, 0x7d562b2bU,
 300    0x19e7fefeU, 0x62b5d7d7U, 0xe64dababU, 0x9aec7676U,
 301    0x458fcacaU, 0x9d1f8282U, 0x4089c9c9U, 0x87fa7d7dU,
 302    0x15effafaU, 0xebb25959U, 0xc98e4747U, 0x0bfbf0f0U,
 303    0xec41adadU, 0x67b3d4d4U, 0xfd5fa2a2U, 0xea45afafU,
 304    0xbf239c9cU, 0xf753a4a4U, 0x96e47272U, 0x5b9bc0c0U,
 305    0xc275b7b7U, 0x1ce1fdfdU, 0xae3d9393U, 0x6a4c2626U,
 306    0x5a6c3636U, 0x417e3f3fU, 0x02f5f7f7U, 0x4f83ccccU,
 307    0x5c683434U, 0xf451a5a5U, 0x34d1e5e5U, 0x08f9f1f1U,
 308    0x93e27171U, 0x73abd8d8U, 0x53623131U, 0x3f2a1515U,
 309    0x0c080404U, 0x5295c7c7U, 0x65462323U, 0x5e9dc3c3U,
 310    0x28301818U, 0xa1379696U, 0x0f0a0505U, 0xb52f9a9aU,
 311    0x090e0707U, 0x36241212U, 0x9b1b8080U, 0x3ddfe2e2U,
 312    0x26cdebebU, 0x694e2727U, 0xcd7fb2b2U, 0x9fea7575U,
 313    0x1b120909U, 0x9e1d8383U, 0x74582c2cU, 0x2e341a1aU,
 314    0x2d361b1bU, 0xb2dc6e6eU, 0xeeb45a5aU, 0xfb5ba0a0U,
 315    0xf6a45252U, 0x4d763b3bU, 0x61b7d6d6U, 0xce7db3b3U,
 316    0x7b522929U, 0x3edde3e3U, 0x715e2f2fU, 0x97138484U,
 317    0xf5a65353U, 0x68b9d1d1U, 0x00000000U, 0x2cc1ededU,
 318    0x60402020U, 0x1fe3fcfcU, 0xc879b1b1U, 0xedb65b5bU,
 319    0xbed46a6aU, 0x468dcbcbU, 0xd967bebeU, 0x4b723939U,
 320    0xde944a4aU, 0xd4984c4cU, 0xe8b05858U, 0x4a85cfcfU,
 321    0x6bbbd0d0U, 0x2ac5efefU, 0xe54faaaaU, 0x16edfbfbU,
 322    0xc5864343U, 0xd79a4d4dU, 0x55663333U, 0x94118585U,
 323    0xcf8a4545U, 0x10e9f9f9U, 0x06040202U, 0x81fe7f7fU,
 324    0xf0a05050U, 0x44783c3cU, 0xba259f9fU, 0xe34ba8a8U,
 325    0xf3a25151U, 0xfe5da3a3U, 0xc0804040U, 0x8a058f8fU,
 326    0xad3f9292U, 0xbc219d9dU, 0x48703838U, 0x04f1f5f5U,
 327    0xdf63bcbcU, 0xc177b6b6U, 0x75afdadaU, 0x63422121U,
 328    0x30201010U, 0x1ae5ffffU, 0x0efdf3f3U, 0x6dbfd2d2U,
 329    0x4c81cdcdU, 0x14180c0cU, 0x35261313U, 0x2fc3ececU,
 330    0xe1be5f5fU, 0xa2359797U, 0xcc884444U, 0x392e1717U,
 331    0x5793c4c4U, 0xf255a7a7U, 0x82fc7e7eU, 0x477a3d3dU,
 332    0xacc86464U, 0xe7ba5d5dU, 0x2b321919U, 0x95e67373U,
 333    0xa0c06060U, 0x98198181U, 0xd19e4f4fU, 0x7fa3dcdcU,
 334    0x66442222U, 0x7e542a2aU, 0xab3b9090U, 0x830b8888U,
 335    0xca8c4646U, 0x29c7eeeeU, 0xd36bb8b8U, 0x3c281414U,
 336    0x79a7dedeU, 0xe2bc5e5eU, 0x1d160b0bU, 0x76addbdbU,
 337    0x3bdbe0e0U, 0x56643232U, 0x4e743a3aU, 0x1e140a0aU,
 338    0xdb924949U, 0x0a0c0606U, 0x6c482424U, 0xe4b85c5cU,
 339    0x5d9fc2c2U, 0x6ebdd3d3U, 0xef43acacU, 0xa6c46262U,
 340    0xa8399191U, 0xa4319595U, 0x37d3e4e4U, 0x8bf27979U,
 341    0x32d5e7e7U, 0x438bc8c8U, 0x596e3737U, 0xb7da6d6dU,
 342    0x8c018d8dU, 0x64b1d5d5U, 0xd29c4e4eU, 0xe049a9a9U,
 343    0xb4d86c6cU, 0xfaac5656U, 0x07f3f4f4U, 0x25cfeaeaU,
 344    0xafca6565U, 0x8ef47a7aU, 0xe947aeaeU, 0x18100808U,
 345    0xd56fbabaU, 0x88f07878U, 0x6f4a2525U, 0x725c2e2eU,
 346    0x24381c1cU, 0xf157a6a6U, 0xc773b4b4U, 0x5197c6c6U,
 347    0x23cbe8e8U, 0x7ca1ddddU, 0x9ce87474U, 0x213e1f1fU,
 348    0xdd964b4bU, 0xdc61bdbdU, 0x860d8b8bU, 0x850f8a8aU,
 349    0x90e07070U, 0x427c3e3eU, 0xc471b5b5U, 0xaacc6666U,
 350    0xd8904848U, 0x05060303U, 0x01f7f6f6U, 0x121c0e0eU,
 351    0xa3c26161U, 0x5f6a3535U, 0xf9ae5757U, 0xd069b9b9U,
 352    0x91178686U, 0x5899c1c1U, 0x273a1d1dU, 0xb9279e9eU,
 353    0x38d9e1e1U, 0x13ebf8f8U, 0xb32b9898U, 0x33221111U,
 354    0xbbd26969U, 0x70a9d9d9U, 0x89078e8eU, 0xa7339494U,
 355    0xb62d9b9bU, 0x223c1e1eU, 0x92158787U, 0x20c9e9e9U,
 356    0x4987ceceU, 0xffaa5555U, 0x78502828U, 0x7aa5dfdfU,
 357    0x8f038c8cU, 0xf859a1a1U, 0x80098989U, 0x171a0d0dU,
 358    0xda65bfbfU, 0x31d7e6e6U, 0xc6844242U, 0xb8d06868U,
 359    0xc3824141U, 0xb0299999U, 0x775a2d2dU, 0x111e0f0fU,
 360    0xcb7bb0b0U, 0xfca85454U, 0xd66dbbbbU, 0x3a2c1616U,
 361},
 362{
 363    0x63a5c663U, 0x7c84f87cU, 0x7799ee77U, 0x7b8df67bU,
 364    0xf20dfff2U, 0x6bbdd66bU, 0x6fb1de6fU, 0xc55491c5U,
 365    0x30506030U, 0x01030201U, 0x67a9ce67U, 0x2b7d562bU,
 366    0xfe19e7feU, 0xd762b5d7U, 0xabe64dabU, 0x769aec76U,
 367    0xca458fcaU, 0x829d1f82U, 0xc94089c9U, 0x7d87fa7dU,
 368    0xfa15effaU, 0x59ebb259U, 0x47c98e47U, 0xf00bfbf0U,
 369    0xadec41adU, 0xd467b3d4U, 0xa2fd5fa2U, 0xafea45afU,
 370    0x9cbf239cU, 0xa4f753a4U, 0x7296e472U, 0xc05b9bc0U,
 371    0xb7c275b7U, 0xfd1ce1fdU, 0x93ae3d93U, 0x266a4c26U,
 372    0x365a6c36U, 0x3f417e3fU, 0xf702f5f7U, 0xcc4f83ccU,
 373    0x345c6834U, 0xa5f451a5U, 0xe534d1e5U, 0xf108f9f1U,
 374    0x7193e271U, 0xd873abd8U, 0x31536231U, 0x153f2a15U,
 375    0x040c0804U, 0xc75295c7U, 0x23654623U, 0xc35e9dc3U,
 376    0x18283018U, 0x96a13796U, 0x050f0a05U, 0x9ab52f9aU,
 377    0x07090e07U, 0x12362412U, 0x809b1b80U, 0xe23ddfe2U,
 378    0xeb26cdebU, 0x27694e27U, 0xb2cd7fb2U, 0x759fea75U,
 379    0x091b1209U, 0x839e1d83U, 0x2c74582cU, 0x1a2e341aU,
 380    0x1b2d361bU, 0x6eb2dc6eU, 0x5aeeb45aU, 0xa0fb5ba0U,
 381    0x52f6a452U, 0x3b4d763bU, 0xd661b7d6U, 0xb3ce7db3U,
 382    0x297b5229U, 0xe33edde3U, 0x2f715e2fU, 0x84971384U,
 383    0x53f5a653U, 0xd168b9d1U, 0x00000000U, 0xed2cc1edU,
 384    0x20604020U, 0xfc1fe3fcU, 0xb1c879b1U, 0x5bedb65bU,
 385    0x6abed46aU, 0xcb468dcbU, 0xbed967beU, 0x394b7239U,
 386    0x4ade944aU, 0x4cd4984cU, 0x58e8b058U, 0xcf4a85cfU,
 387    0xd06bbbd0U, 0xef2ac5efU, 0xaae54faaU, 0xfb16edfbU,
 388    0x43c58643U, 0x4dd79a4dU, 0x33556633U, 0x85941185U,
 389    0x45cf8a45U, 0xf910e9f9U, 0x02060402U, 0x7f81fe7fU,
 390    0x50f0a050U, 0x3c44783cU, 0x9fba259fU, 0xa8e34ba8U,
 391    0x51f3a251U, 0xa3fe5da3U, 0x40c08040U, 0x8f8a058fU,
 392    0x92ad3f92U, 0x9dbc219dU, 0x38487038U, 0xf504f1f5U,
 393    0xbcdf63bcU, 0xb6c177b6U, 0xda75afdaU, 0x21634221U,
 394    0x10302010U, 0xff1ae5ffU, 0xf30efdf3U, 0xd26dbfd2U,
 395    0xcd4c81cdU, 0x0c14180cU, 0x13352613U, 0xec2fc3ecU,
 396    0x5fe1be5fU, 0x97a23597U, 0x44cc8844U, 0x17392e17U,
 397    0xc45793c4U, 0xa7f255a7U, 0x7e82fc7eU, 0x3d477a3dU,
 398    0x64acc864U, 0x5de7ba5dU, 0x192b3219U, 0x7395e673U,
 399    0x60a0c060U, 0x81981981U, 0x4fd19e4fU, 0xdc7fa3dcU,
 400    0x22664422U, 0x2a7e542aU, 0x90ab3b90U, 0x88830b88U,
 401    0x46ca8c46U, 0xee29c7eeU, 0xb8d36bb8U, 0x143c2814U,
 402    0xde79a7deU, 0x5ee2bc5eU, 0x0b1d160bU, 0xdb76addbU,
 403    0xe03bdbe0U, 0x32566432U, 0x3a4e743aU, 0x0a1e140aU,
 404    0x49db9249U, 0x060a0c06U, 0x246c4824U, 0x5ce4b85cU,
 405    0xc25d9fc2U, 0xd36ebdd3U, 0xacef43acU, 0x62a6c462U,
 406    0x91a83991U, 0x95a43195U, 0xe437d3e4U, 0x798bf279U,
 407    0xe732d5e7U, 0xc8438bc8U, 0x37596e37U, 0x6db7da6dU,
 408    0x8d8c018dU, 0xd564b1d5U, 0x4ed29c4eU, 0xa9e049a9U,
 409    0x6cb4d86cU, 0x56faac56U, 0xf407f3f4U, 0xea25cfeaU,
 410    0x65afca65U, 0x7a8ef47aU, 0xaee947aeU, 0x08181008U,
 411    0xbad56fbaU, 0x7888f078U, 0x256f4a25U, 0x2e725c2eU,
 412    0x1c24381cU, 0xa6f157a6U, 0xb4c773b4U, 0xc65197c6U,
 413    0xe823cbe8U, 0xdd7ca1ddU, 0x749ce874U, 0x1f213e1fU,
 414    0x4bdd964bU, 0xbddc61bdU, 0x8b860d8bU, 0x8a850f8aU,
 415    0x7090e070U, 0x3e427c3eU, 0xb5c471b5U, 0x66aacc66U,
 416    0x48d89048U, 0x03050603U, 0xf601f7f6U, 0x0e121c0eU,
 417    0x61a3c261U, 0x355f6a35U, 0x57f9ae57U, 0xb9d069b9U,
 418    0x86911786U, 0xc15899c1U, 0x1d273a1dU, 0x9eb9279eU,
 419    0xe138d9e1U, 0xf813ebf8U, 0x98b32b98U, 0x11332211U,
 420    0x69bbd269U, 0xd970a9d9U, 0x8e89078eU, 0x94a73394U,
 421    0x9bb62d9bU, 0x1e223c1eU, 0x87921587U, 0xe920c9e9U,
 422    0xce4987ceU, 0x55ffaa55U, 0x28785028U, 0xdf7aa5dfU,
 423    0x8c8f038cU, 0xa1f859a1U, 0x89800989U, 0x0d171a0dU,
 424    0xbfda65bfU, 0xe631d7e6U, 0x42c68442U, 0x68b8d068U,
 425    0x41c38241U, 0x99b02999U, 0x2d775a2dU, 0x0f111e0fU,
 426    0xb0cb7bb0U, 0x54fca854U, 0xbbd66dbbU, 0x163a2c16U,
 427},
 428{
 429    0x6363a5c6U, 0x7c7c84f8U, 0x777799eeU, 0x7b7b8df6U,
 430    0xf2f20dffU, 0x6b6bbdd6U, 0x6f6fb1deU, 0xc5c55491U,
 431    0x30305060U, 0x01010302U, 0x6767a9ceU, 0x2b2b7d56U,
 432    0xfefe19e7U, 0xd7d762b5U, 0xababe64dU, 0x76769aecU,
 433    0xcaca458fU, 0x82829d1fU, 0xc9c94089U, 0x7d7d87faU,
 434    0xfafa15efU, 0x5959ebb2U, 0x4747c98eU, 0xf0f00bfbU,
 435    0xadadec41U, 0xd4d467b3U, 0xa2a2fd5fU, 0xafafea45U,
 436    0x9c9cbf23U, 0xa4a4f753U, 0x727296e4U, 0xc0c05b9bU,
 437    0xb7b7c275U, 0xfdfd1ce1U, 0x9393ae3dU, 0x26266a4cU,
 438    0x36365a6cU, 0x3f3f417eU, 0xf7f702f5U, 0xcccc4f83U,
 439    0x34345c68U, 0xa5a5f451U, 0xe5e534d1U, 0xf1f108f9U,
 440    0x717193e2U, 0xd8d873abU, 0x31315362U, 0x15153f2aU,
 441    0x04040c08U, 0xc7c75295U, 0x23236546U, 0xc3c35e9dU,
 442    0x18182830U, 0x9696a137U, 0x05050f0aU, 0x9a9ab52fU,
 443    0x0707090eU, 0x12123624U, 0x80809b1bU, 0xe2e23ddfU,
 444    0xebeb26cdU, 0x2727694eU, 0xb2b2cd7fU, 0x75759feaU,
 445    0x09091b12U, 0x83839e1dU, 0x2c2c7458U, 0x1a1a2e34U,
 446    0x1b1b2d36U, 0x6e6eb2dcU, 0x5a5aeeb4U, 0xa0a0fb5bU,
 447    0x5252f6a4U, 0x3b3b4d76U, 0xd6d661b7U, 0xb3b3ce7dU,
 448    0x29297b52U, 0xe3e33eddU, 0x2f2f715eU, 0x84849713U,
 449    0x5353f5a6U, 0xd1d168b9U, 0x00000000U, 0xeded2cc1U,
 450    0x20206040U, 0xfcfc1fe3U, 0xb1b1c879U, 0x5b5bedb6U,
 451    0x6a6abed4U, 0xcbcb468dU, 0xbebed967U, 0x39394b72U,
 452    0x4a4ade94U, 0x4c4cd498U, 0x5858e8b0U, 0xcfcf4a85U,
 453    0xd0d06bbbU, 0xefef2ac5U, 0xaaaae54fU, 0xfbfb16edU,
 454    0x4343c586U, 0x4d4dd79aU, 0x33335566U, 0x85859411U,
 455    0x4545cf8aU, 0xf9f910e9U, 0x02020604U, 0x7f7f81feU,
 456    0x5050f0a0U, 0x3c3c4478U, 0x9f9fba25U, 0xa8a8e34bU,
 457    0x5151f3a2U, 0xa3a3fe5dU, 0x4040c080U, 0x8f8f8a05U,
 458    0x9292ad3fU, 0x9d9dbc21U, 0x38384870U, 0xf5f504f1U,
 459    0xbcbcdf63U, 0xb6b6c177U, 0xdada75afU, 0x21216342U,
 460    0x10103020U, 0xffff1ae5U, 0xf3f30efdU, 0xd2d26dbfU,
 461    0xcdcd4c81U, 0x0c0c1418U, 0x13133526U, 0xecec2fc3U,
 462    0x5f5fe1beU, 0x9797a235U, 0x4444cc88U, 0x1717392eU,
 463    0xc4c45793U, 0xa7a7f255U, 0x7e7e82fcU, 0x3d3d477aU,
 464    0x6464acc8U, 0x5d5de7baU, 0x19192b32U, 0x737395e6U,
 465    0x6060a0c0U, 0x81819819U, 0x4f4fd19eU, 0xdcdc7fa3U,
 466    0x22226644U, 0x2a2a7e54U, 0x9090ab3bU, 0x8888830bU,
 467    0x4646ca8cU, 0xeeee29c7U, 0xb8b8d36bU, 0x14143c28U,
 468    0xdede79a7U, 0x5e5ee2bcU, 0x0b0b1d16U, 0xdbdb76adU,
 469    0xe0e03bdbU, 0x32325664U, 0x3a3a4e74U, 0x0a0a1e14U,
 470    0x4949db92U, 0x06060a0cU, 0x24246c48U, 0x5c5ce4b8U,
 471    0xc2c25d9fU, 0xd3d36ebdU, 0xacacef43U, 0x6262a6c4U,
 472    0x9191a839U, 0x9595a431U, 0xe4e437d3U, 0x79798bf2U,
 473    0xe7e732d5U, 0xc8c8438bU, 0x3737596eU, 0x6d6db7daU,
 474    0x8d8d8c01U, 0xd5d564b1U, 0x4e4ed29cU, 0xa9a9e049U,
 475    0x6c6cb4d8U, 0x5656faacU, 0xf4f407f3U, 0xeaea25cfU,
 476    0x6565afcaU, 0x7a7a8ef4U, 0xaeaee947U, 0x08081810U,
 477    0xbabad56fU, 0x787888f0U, 0x25256f4aU, 0x2e2e725cU,
 478    0x1c1c2438U, 0xa6a6f157U, 0xb4b4c773U, 0xc6c65197U,
 479    0xe8e823cbU, 0xdddd7ca1U, 0x74749ce8U, 0x1f1f213eU,
 480    0x4b4bdd96U, 0xbdbddc61U, 0x8b8b860dU, 0x8a8a850fU,
 481    0x707090e0U, 0x3e3e427cU, 0xb5b5c471U, 0x6666aaccU,
 482    0x4848d890U, 0x03030506U, 0xf6f601f7U, 0x0e0e121cU,
 483    0x6161a3c2U, 0x35355f6aU, 0x5757f9aeU, 0xb9b9d069U,
 484    0x86869117U, 0xc1c15899U, 0x1d1d273aU, 0x9e9eb927U,
 485    0xe1e138d9U, 0xf8f813ebU, 0x9898b32bU, 0x11113322U,
 486    0x6969bbd2U, 0xd9d970a9U, 0x8e8e8907U, 0x9494a733U,
 487    0x9b9bb62dU, 0x1e1e223cU, 0x87879215U, 0xe9e920c9U,
 488    0xcece4987U, 0x5555ffaaU, 0x28287850U, 0xdfdf7aa5U,
 489    0x8c8c8f03U, 0xa1a1f859U, 0x89898009U, 0x0d0d171aU,
 490    0xbfbfda65U, 0xe6e631d7U, 0x4242c684U, 0x6868b8d0U,
 491    0x4141c382U, 0x9999b029U, 0x2d2d775aU, 0x0f0f111eU,
 492    0xb0b0cb7bU, 0x5454fca8U, 0xbbbbd66dU, 0x16163a2cU,
 493}
 494};
 495
 496
 497__device__
 498static word32 GetTable(const word32* t, byte o)
 499{
 500    word32 e = 0;
 501#if WC_CACHE_LINE_SZ == 64
 502  byte hi = o & 0xf0;
 503  byte lo = o & 0x0f;
 504
 505  e  = t[lo + 0x00] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 506  e |= t[lo + 0x10] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 507  e |= t[lo + 0x20] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 508  e |= t[lo + 0x30] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 509  e |= t[lo + 0x40] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 510  e |= t[lo + 0x50] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 511  e |= t[lo + 0x60] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 512  e |= t[lo + 0x70] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 513  e |= t[lo + 0x80] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 514  e |= t[lo + 0x90] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 515  e |= t[lo + 0xa0] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 516  e |= t[lo + 0xb0] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 517  e |= t[lo + 0xc0] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 518  e |= t[lo + 0xd0] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 519  e |= t[lo + 0xe0] & ((word32)0 - (((word32)hi - 0x01) >> 31)); hi -= 0x10;
 520  e |= t[lo + 0xf0] & ((word32)0 - (((word32)hi - 0x01) >> 31));
 521#else
 522  int i;
 523  byte hi = o & WC_CACHE_LINE_MASK_HI;
 524  byte lo = o & WC_CACHE_LINE_MASK_LO;
 525
 526  for (i = 0; i < 256; i += (1 << WC_CACHE_LINE_BITS)) {
 527      e |= t[lo + i] & ((word32)0 - (((word32)hi - 0x01) >> 31));
 528      hi -= WC_CACHE_LINE_ADD;
 529  }
 530#endif
 531  return e;
 532}
 533
 534__device__
 535static void GetTable_Multi(const word32* t, word32* t0, byte o0,
 536  word32* t1, byte o1, word32* t2, byte o2, word32* t3, byte o3)
 537{
 538  word32 e0 = 0;
 539  word32 e1 = 0;
 540  word32 e2 = 0;
 541  word32 e3 = 0;
 542  byte hi0 = o0 & WC_CACHE_LINE_MASK_HI;
 543  byte lo0 = o0 & WC_CACHE_LINE_MASK_LO;
 544  byte hi1 = o1 & WC_CACHE_LINE_MASK_HI;
 545  byte lo1 = o1 & WC_CACHE_LINE_MASK_LO;
 546  byte hi2 = o2 & WC_CACHE_LINE_MASK_HI;
 547  byte lo2 = o2 & WC_CACHE_LINE_MASK_LO;
 548  byte hi3 = o3 & WC_CACHE_LINE_MASK_HI;
 549  byte lo3 = o3 & WC_CACHE_LINE_MASK_LO;
 550  int i;
 551
 552  for (i = 0; i < 256; i += (1 << WC_CACHE_LINE_BITS)) {
 553      e0 |= t[lo0 + i] & ((word32)0 - (((word32)hi0 - 0x01) >> 31));
 554      hi0 -= WC_CACHE_LINE_ADD;
 555      e1 |= t[lo1 + i] & ((word32)0 - (((word32)hi1 - 0x01) >> 31));
 556      hi1 -= WC_CACHE_LINE_ADD;
 557      e2 |= t[lo2 + i] & ((word32)0 - (((word32)hi2 - 0x01) >> 31));
 558      hi2 -= WC_CACHE_LINE_ADD;
 559      e3 |= t[lo3 + i] & ((word32)0 - (((word32)hi3 - 0x01) >> 31));
 560      hi3 -= WC_CACHE_LINE_ADD;
 561  }
 562  *t0 = e0;
 563  *t1 = e1;
 564  *t2 = e2;
 565  *t3 = e3;
 566}
 567
 568/* load 4 Te Tables into cache by cache line stride */
 569__device__
 570static WARN_UNUSED_RESULT WC_INLINE word32 PreFetchTe(void)
 571{
 572#ifndef WOLFSSL_AES_TOUCH_LINES
 573    word32 x = 0;
 574    int i,j;
 575
 576    for (i = 0; i < 4; i++) {
 577        /* 256 elements, each one is 4 bytes */
 578        for (j = 0; j < 256; j += WC_CACHE_LINE_SZ/4) {
 579            x &= Te_CUDA[i][j];
 580        }
 581    }
 582    return x;
 583#else
 584    return 0;
 585#endif
 586}
 587#else
 588__device__ static const byte Tsbox[256] = {
 589    0x63U, 0x7cU, 0x77U, 0x7bU, 0xf2U, 0x6bU, 0x6fU, 0xc5U,
 590    0x30U, 0x01U, 0x67U, 0x2bU, 0xfeU, 0xd7U, 0xabU, 0x76U,
 591    0xcaU, 0x82U, 0xc9U, 0x7dU, 0xfaU, 0x59U, 0x47U, 0xf0U,
 592    0xadU, 0xd4U, 0xa2U, 0xafU, 0x9cU, 0xa4U, 0x72U, 0xc0U,
 593    0xb7U, 0xfdU, 0x93U, 0x26U, 0x36U, 0x3fU, 0xf7U, 0xccU,
 594    0x34U, 0xa5U, 0xe5U, 0xf1U, 0x71U, 0xd8U, 0x31U, 0x15U,
 595    0x04U, 0xc7U, 0x23U, 0xc3U, 0x18U, 0x96U, 0x05U, 0x9aU,
 596    0x07U, 0x12U, 0x80U, 0xe2U, 0xebU, 0x27U, 0xb2U, 0x75U,
 597    0x09U, 0x83U, 0x2cU, 0x1aU, 0x1bU, 0x6eU, 0x5aU, 0xa0U,
 598    0x52U, 0x3bU, 0xd6U, 0xb3U, 0x29U, 0xe3U, 0x2fU, 0x84U,
 599    0x53U, 0xd1U, 0x00U, 0xedU, 0x20U, 0xfcU, 0xb1U, 0x5bU,
 600    0x6aU, 0xcbU, 0xbeU, 0x39U, 0x4aU, 0x4cU, 0x58U, 0xcfU,
 601    0xd0U, 0xefU, 0xaaU, 0xfbU, 0x43U, 0x4dU, 0x33U, 0x85U,
 602    0x45U, 0xf9U, 0x02U, 0x7fU, 0x50U, 0x3cU, 0x9fU, 0xa8U,
 603    0x51U, 0xa3U, 0x40U, 0x8fU, 0x92U, 0x9dU, 0x38U, 0xf5U,
 604    0xbcU, 0xb6U, 0xdaU, 0x21U, 0x10U, 0xffU, 0xf3U, 0xd2U,
 605    0xcdU, 0x0cU, 0x13U, 0xecU, 0x5fU, 0x97U, 0x44U, 0x17U,
 606    0xc4U, 0xa7U, 0x7eU, 0x3dU, 0x64U, 0x5dU, 0x19U, 0x73U,
 607    0x60U, 0x81U, 0x4fU, 0xdcU, 0x22U, 0x2aU, 0x90U, 0x88U,
 608    0x46U, 0xeeU, 0xb8U, 0x14U, 0xdeU, 0x5eU, 0x0bU, 0xdbU,
 609    0xe0U, 0x32U, 0x3aU, 0x0aU, 0x49U, 0x06U, 0x24U, 0x5cU,
 610    0xc2U, 0xd3U, 0xacU, 0x62U, 0x91U, 0x95U, 0xe4U, 0x79U,
 611    0xe7U, 0xc8U, 0x37U, 0x6dU, 0x8dU, 0xd5U, 0x4eU, 0xa9U,
 612    0x6cU, 0x56U, 0xf4U, 0xeaU, 0x65U, 0x7aU, 0xaeU, 0x08U,
 613    0xbaU, 0x78U, 0x25U, 0x2eU, 0x1cU, 0xa6U, 0xb4U, 0xc6U,
 614    0xe8U, 0xddU, 0x74U, 0x1fU, 0x4bU, 0xbdU, 0x8bU, 0x8aU,
 615    0x70U, 0x3eU, 0xb5U, 0x66U, 0x48U, 0x03U, 0xf6U, 0x0eU,
 616    0x61U, 0x35U, 0x57U, 0xb9U, 0x86U, 0xc1U, 0x1dU, 0x9eU,
 617    0xe1U, 0xf8U, 0x98U, 0x11U, 0x69U, 0xd9U, 0x8eU, 0x94U,
 618    0x9bU, 0x1eU, 0x87U, 0xe9U, 0xceU, 0x55U, 0x28U, 0xdfU,
 619    0x8cU, 0xa1U, 0x89U, 0x0dU, 0xbfU, 0xe6U, 0x42U, 0x68U,
 620    0x41U, 0x99U, 0x2dU, 0x0fU, 0xb0U, 0x54U, 0xbbU, 0x16U
 621};
 622
 623#define AES_XTIME(x)    ((byte)((byte)((x) << 1) ^ ((0 - ((x) >> 7)) & 0x1b)))
 624
 625#define col_mul(t, i2, i3, ia, ib) \
 626  ( GETBYTE(t, ia) ^ GETBYTE(t, ib) ^ GETBYTE(t, i3) ^ AES_XTIME(GETBYTE(t, i2) ^ GETBYTE(t, i3)) )
 627
 628#define GetTable(t, o)  t[o]
 629#define GetTable8(t, o) t[o]
 630#define GetTable_Multi(t, t0, o0, t1, o1, t2, o2, t3, o3)  \
 631  *(t0) = (t)[o0]; *(t1) = (t)[o1]; *(t2) = (t)[o2]; *(t3) = (t)[o3]
 632#define XorTable_Multi(t, t0, o0, t1, o1, t2, o2, t3, o3)  \
 633  *(t0) ^= (t)[o0]; *(t1) ^= (t)[o1]; *(t2) ^= (t)[o2]; *(t3) ^= (t)[o3]
 634#define GetTable8_4(t, o0, o1, o2, o3) \
 635  (((word32)(t)[o0] << 24) | ((word32)(t)[o1] << 16) |   \
 636   ((word32)(t)[o2] <<  8) | ((word32)(t)[o3] <<  0))
 637
 638/* load sbox into cache by cache line stride */
 639#ifndef WOLFSSL_AES_TOUCH_LINES
 640    #define PreFetchSBox(x) { \
 641    x = 0; \
 642    int i; \
 643    for (i = 0; i < 256; i += WC_CACHE_LINE_SZ/4) { \
 644        x &= Tsbox[i]; \
 645    } \
 646    }
 647#else
 648    #define PreFetchSBox(x) ( x = 0 )
 649#endif
 650#endif
 651
 652#if !defined(WC_AES_BITSLICED)
 653/* Encrypt a block using AES.
 654 *
 655 * @param [in]  aes       AES object.
 656 * @param [in]  inBlock   Block to encrypt.
 657 * @param [out] outBlock  Encrypted block.
 658 * @param [in]  r         Rounds divided by 2.
 659 * @param [in]  sz        Number of blocks to encrypt
 660 */
 661__global__ void AesEncrypt_C_CUDA(word32* rkBase, const byte* inBlockBase, byte* outBlockBase,
 662        word32 r, word32 sz)
 663{
 664    word32 s0, s1, s2, s3;
 665    word32 t0, t1, t2, t3;
 666    word32 sBox;
 667    int idx = blockIdx.x * blockDim.x + threadIdx.x;
 668    int stride = blockDim.x * gridDim.x;
 669    const byte* inBlock = inBlockBase;
 670    byte* outBlock = outBlockBase;
 671    word32* rk;
 672
 673    for (int i = idx; i < sz; i += stride) {
 674        rk = rkBase;
 675        inBlock = inBlockBase + i * 4 * sizeof(s0);
 676        outBlock = outBlockBase + i * 4 * sizeof(s0);
 677
 678        /*
 679         * map byte array block to cipher state
 680         * and add initial round key:
 681         */
 682        XMEMCPY(&s0, inBlock,                  sizeof(s0));
 683        XMEMCPY(&s1, inBlock +     sizeof(s0), sizeof(s1));
 684        XMEMCPY(&s2, inBlock + 2 * sizeof(s0), sizeof(s2));
 685        XMEMCPY(&s3, inBlock + 3 * sizeof(s0), sizeof(s3));
 686
 687#ifdef LITTLE_ENDIAN_ORDER
 688        s0 = ByteReverseWord32(s0);
 689        s1 = ByteReverseWord32(s1);
 690        s2 = ByteReverseWord32(s2);
 691        s3 = ByteReverseWord32(s3);
 692#endif
 693
 694        /* AddRoundKey */
 695        s0 ^= rk[0];
 696        s1 ^= rk[1];
 697        s2 ^= rk[2];
 698        s3 ^= rk[3];
 699
 700#ifndef WOLFSSL_AES_SMALL_TABLES
 701#ifndef WC_NO_CACHE_RESISTANT
 702        s0 |= PreFetchTe();
 703#endif
 704
 705#ifndef WOLFSSL_AES_TOUCH_LINES
 706#define ENC_ROUND_T_S(o)                                                       \
 707        t0 = GetTable(Te_CUDA[0], GETBYTE(s0, 3)) ^ GetTable(Te_CUDA[1], GETBYTE(s1, 2)) ^   \
 708             GetTable(Te_CUDA[2], GETBYTE(s2, 1)) ^ GetTable(Te_CUDA[3], GETBYTE(s3, 0)) ^   \
 709             rk[(o)+4];                                                            \
 710        t1 = GetTable(Te_CUDA[0], GETBYTE(s1, 3)) ^ GetTable(Te_CUDA[1], GETBYTE(s2, 2)) ^   \
 711             GetTable(Te_CUDA[2], GETBYTE(s3, 1)) ^ GetTable(Te_CUDA[3], GETBYTE(s0, 0)) ^   \
 712             rk[(o)+5];                                                            \
 713        t2 = GetTable(Te_CUDA[0], GETBYTE(s2, 3)) ^ GetTable(Te_CUDA[1], GETBYTE(s3, 2)) ^   \
 714             GetTable(Te_CUDA[2], GETBYTE(s0, 1)) ^ GetTable(Te_CUDA[3], GETBYTE(s1, 0)) ^   \
 715             rk[(o)+6];                                                            \
 716        t3 = GetTable(Te_CUDA[0], GETBYTE(s3, 3)) ^ GetTable(Te_CUDA[1], GETBYTE(s0, 2)) ^   \
 717             GetTable(Te_CUDA[2], GETBYTE(s1, 1)) ^ GetTable(Te_CUDA[3], GETBYTE(s2, 0)) ^   \
 718             rk[(o)+7]
 719#define ENC_ROUND_S_T(o)                                                       \
 720        s0 = GetTable(Te_CUDA[0], GETBYTE(t0, 3)) ^ GetTable(Te_CUDA[1], GETBYTE(t1, 2)) ^   \
 721             GetTable(Te_CUDA[2], GETBYTE(t2, 1)) ^ GetTable(Te_CUDA[3], GETBYTE(t3, 0)) ^   \
 722             rk[(o)+0];                                                            \
 723        s1 = GetTable(Te_CUDA[0], GETBYTE(t1, 3)) ^ GetTable(Te_CUDA[1], GETBYTE(t2, 2)) ^   \
 724             GetTable(Te_CUDA[2], GETBYTE(t3, 1)) ^ GetTable(Te_CUDA[3], GETBYTE(t0, 0)) ^   \
 725             rk[(o)+1];                                                            \
 726        s2 = GetTable(Te_CUDA[0], GETBYTE(t2, 3)) ^ GetTable(Te_CUDA[1], GETBYTE(t3, 2)) ^   \
 727             GetTable(Te_CUDA[2], GETBYTE(t0, 1)) ^ GetTable(Te_CUDA[3], GETBYTE(t1, 0)) ^   \
 728             rk[(o)+2];                                                            \
 729        s3 = GetTable(Te_CUDA[0], GETBYTE(t3, 3)) ^ GetTable(Te_CUDA[1], GETBYTE(t0, 2)) ^   \
 730             GetTable(Te_CUDA[2], GETBYTE(t1, 1)) ^ GetTable(Te_CUDA[3], GETBYTE(t2, 0)) ^   \
 731             rk[(o)+3]
 732#else
 733#define ENC_ROUND_T_S(o)                                                       \
 734        GetTable_Multi(Te_CUDA[0], &t0, GETBYTE(s0, 3), &t1, GETBYTE(s1, 3),            \
 735                              &t2, GETBYTE(s2, 3), &t3, GETBYTE(s3, 3));           \
 736        XorTable_Multi(Te_CUDA[1], &t0, GETBYTE(s1, 2), &t1, GETBYTE(s2, 2),            \
 737                              &t2, GETBYTE(s3, 2), &t3, GETBYTE(s0, 2));           \
 738        XorTable_Multi(Te_CUDA[2], &t0, GETBYTE(s2, 1), &t1, GETBYTE(s3, 1),            \
 739                              &t2, GETBYTE(s0, 1), &t3, GETBYTE(s1, 1));           \
 740        XorTable_Multi(Te_CUDA[3], &t0, GETBYTE(s3, 0), &t1, GETBYTE(s0, 0),            \
 741                              &t2, GETBYTE(s1, 0), &t3, GETBYTE(s2, 0));           \
 742        t0 ^= rk[(o)+4]; t1 ^= rk[(o)+5]; t2 ^= rk[(o)+6]; t3 ^= rk[(o)+7];
 743
 744#define ENC_ROUND_S_T(o)                                                       \
 745        GetTable_Multi(Te_CUDA[0], &s0, GETBYTE(t0, 3), &s1, GETBYTE(t1, 3),            \
 746                              &s2, GETBYTE(t2, 3), &s3, GETBYTE(t3, 3));           \
 747        XorTable_Multi(Te_CUDA[1], &s0, GETBYTE(t1, 2), &s1, GETBYTE(t2, 2),            \
 748                              &s2, GETBYTE(t3, 2), &s3, GETBYTE(t0, 2));           \
 749        XorTable_Multi(Te_CUDA[2], &s0, GETBYTE(t2, 1), &s1, GETBYTE(t3, 1),            \
 750                              &s2, GETBYTE(t0, 1), &s3, GETBYTE(t1, 1));           \
 751        XorTable_Multi(Te_CUDA[3], &s0, GETBYTE(t3, 0), &s1, GETBYTE(t0, 0),            \
 752                              &s2, GETBYTE(t1, 0), &s3, GETBYTE(t2, 0));           \
 753        s0 ^= rk[(o)+0]; s1 ^= rk[(o)+1]; s2 ^= rk[(o)+2]; s3 ^= rk[(o)+3];
 754#endif
 755
 756#ifndef WOLFSSL_AES_NO_UNROLL
 757    /* Unroll the loop. */
 758                           ENC_ROUND_T_S( 0);
 759        ENC_ROUND_S_T( 8); ENC_ROUND_T_S( 8);
 760        ENC_ROUND_S_T(16); ENC_ROUND_T_S(16);
 761        ENC_ROUND_S_T(24); ENC_ROUND_T_S(24);
 762        ENC_ROUND_S_T(32); ENC_ROUND_T_S(32);
 763        if (r > 5) {
 764            ENC_ROUND_S_T(40); ENC_ROUND_T_S(40);
 765            if (r > 6) {
 766                ENC_ROUND_S_T(48); ENC_ROUND_T_S(48);
 767            }
 768        }
 769        rk += r * 8;
 770#else
 771        /*
 772         * Nr - 1 full rounds:
 773         */
 774
 775        for (;;) {
 776            ENC_ROUND_T_S(0);
 777
 778            rk += 8;
 779            if (--r == 0) {
 780                break;
 781            }
 782
 783            ENC_ROUND_S_T(0);
 784        }
 785#endif
 786
 787        /*
 788         * apply last round and
 789         * map cipher state to byte array block:
 790         */
 791
 792#ifndef WOLFSSL_AES_TOUCH_LINES
 793        s0 =
 794            (GetTable(Te_CUDA[2], GETBYTE(t0, 3)) & 0xff000000) ^
 795            (GetTable(Te_CUDA[3], GETBYTE(t1, 2)) & 0x00ff0000) ^
 796            (GetTable(Te_CUDA[0], GETBYTE(t2, 1)) & 0x0000ff00) ^
 797            (GetTable(Te_CUDA[1], GETBYTE(t3, 0)) & 0x000000ff) ^
 798            rk[0];
 799        s1 =
 800            (GetTable(Te_CUDA[2], GETBYTE(t1, 3)) & 0xff000000) ^
 801            (GetTable(Te_CUDA[3], GETBYTE(t2, 2)) & 0x00ff0000) ^
 802            (GetTable(Te_CUDA[0], GETBYTE(t3, 1)) & 0x0000ff00) ^
 803            (GetTable(Te_CUDA[1], GETBYTE(t0, 0)) & 0x000000ff) ^
 804            rk[1];
 805        s2 =
 806            (GetTable(Te_CUDA[2], GETBYTE(t2, 3)) & 0xff000000) ^
 807            (GetTable(Te_CUDA[3], GETBYTE(t3, 2)) & 0x00ff0000) ^
 808            (GetTable(Te_CUDA[0], GETBYTE(t0, 1)) & 0x0000ff00) ^
 809            (GetTable(Te_CUDA[1], GETBYTE(t1, 0)) & 0x000000ff) ^
 810            rk[2];
 811        s3 =
 812            (GetTable(Te_CUDA[2], GETBYTE(t3, 3)) & 0xff000000) ^
 813            (GetTable(Te_CUDA[3], GETBYTE(t0, 2)) & 0x00ff0000) ^
 814            (GetTable(Te_CUDA[0], GETBYTE(t1, 1)) & 0x0000ff00) ^
 815            (GetTable(Te_CUDA[1], GETBYTE(t2, 0)) & 0x000000ff) ^
 816            rk[3];
 817#else
 818    {
 819        word32 u0;
 820        word32 u1;
 821        word32 u2;
 822        word32 u3;
 823
 824        s0 = rk[0]; s1 = rk[1]; s2 = rk[2]; s3 = rk[3];
 825        GetTable_Multi(Te_CUDA[2], &u0, GETBYTE(t0, 3), &u1, GETBYTE(t1, 3),
 826                              &u2, GETBYTE(t2, 3), &u3, GETBYTE(t3, 3));
 827        s0 ^= u0 & 0xff000000; s1 ^= u1 & 0xff000000;
 828        s2 ^= u2 & 0xff000000; s3 ^= u3 & 0xff000000;
 829        GetTable_Multi(Te_CUDA[3], &u0, GETBYTE(t1, 2), &u1, GETBYTE(t2, 2),
 830                              &u2, GETBYTE(t3, 2), &u3, GETBYTE(t0, 2));
 831        s0 ^= u0 & 0x00ff0000; s1 ^= u1 & 0x00ff0000;
 832        s2 ^= u2 & 0x00ff0000; s3 ^= u3 & 0x00ff0000;
 833        GetTable_Multi(Te_CUDA[0], &u0, GETBYTE(t2, 1), &u1, GETBYTE(t3, 1),
 834                              &u2, GETBYTE(t0, 1), &u3, GETBYTE(t1, 1));
 835        s0 ^= u0 & 0x0000ff00; s1 ^= u1 & 0x0000ff00;
 836        s2 ^= u2 & 0x0000ff00; s3 ^= u3 & 0x0000ff00;
 837        GetTable_Multi(Te_CUDA[1], &u0, GETBYTE(t3, 0), &u1, GETBYTE(t0, 0),
 838                              &u2, GETBYTE(t1, 0), &u3, GETBYTE(t2, 0));
 839        s0 ^= u0 & 0x000000ff; s1 ^= u1 & 0x000000ff;
 840        s2 ^= u2 & 0x000000ff; s3 ^= u3 & 0x000000ff;
 841    }
 842#endif
 843#else
 844#ifndef WC_NO_CACHE_RESISTANT
 845        PreFetchSBox(sBox);
 846        s0 |= sBox;
 847#endif
 848
 849        r *= 2;
 850        /* Two rounds at a time */
 851        for (rk += 4; r > 1; r--, rk += 4) {
 852            t0 =
 853                ((word32)GetTable8(Tsbox, GETBYTE(s0, 3)) << 24) ^
 854                ((word32)GetTable8(Tsbox, GETBYTE(s1, 2)) << 16) ^
 855                ((word32)GetTable8(Tsbox, GETBYTE(s2, 1)) <<  8) ^
 856                ((word32)GetTable8(Tsbox, GETBYTE(s3, 0)));
 857            t1 =
 858                ((word32)GetTable8(Tsbox, GETBYTE(s1, 3)) << 24) ^
 859                ((word32)GetTable8(Tsbox, GETBYTE(s2, 2)) << 16) ^
 860                ((word32)GetTable8(Tsbox, GETBYTE(s3, 1)) <<  8) ^
 861                ((word32)GetTable8(Tsbox, GETBYTE(s0, 0)));
 862            t2 =
 863                ((word32)GetTable8(Tsbox, GETBYTE(s2, 3)) << 24) ^
 864                ((word32)GetTable8(Tsbox, GETBYTE(s3, 2)) << 16) ^
 865                ((word32)GetTable8(Tsbox, GETBYTE(s0, 1)) <<  8) ^
 866                ((word32)GetTable8(Tsbox, GETBYTE(s1, 0)));
 867            t3 =
 868                ((word32)GetTable8(Tsbox, GETBYTE(s3, 3)) << 24) ^
 869                ((word32)GetTable8(Tsbox, GETBYTE(s0, 2)) << 16) ^
 870                ((word32)GetTable8(Tsbox, GETBYTE(s1, 1)) <<  8) ^
 871                ((word32)GetTable8(Tsbox, GETBYTE(s2, 0)));
 872
 873            s0 =
 874                (col_mul(t0, 3, 2, 0, 1) << 24) ^
 875                (col_mul(t0, 2, 1, 0, 3) << 16) ^
 876                (col_mul(t0, 1, 0, 2, 3) <<  8) ^
 877                (col_mul(t0, 0, 3, 2, 1)      ) ^
 878                rk[0];
 879            s1 =
 880                (col_mul(t1, 3, 2, 0, 1) << 24) ^
 881                (col_mul(t1, 2, 1, 0, 3) << 16) ^
 882                (col_mul(t1, 1, 0, 2, 3) <<  8) ^
 883                (col_mul(t1, 0, 3, 2, 1)      ) ^
 884                rk[1];
 885            s2 =
 886                (col_mul(t2, 3, 2, 0, 1) << 24) ^
 887                (col_mul(t2, 2, 1, 0, 3) << 16) ^
 888                (col_mul(t2, 1, 0, 2, 3) <<  8) ^
 889                (col_mul(t2, 0, 3, 2, 1)      ) ^
 890                rk[2];
 891            s3 =
 892                (col_mul(t3, 3, 2, 0, 1) << 24) ^
 893                (col_mul(t3, 2, 1, 0, 3) << 16) ^
 894                (col_mul(t3, 1, 0, 2, 3) <<  8) ^
 895                (col_mul(t3, 0, 3, 2, 1)      ) ^
 896                rk[3];
 897        }
 898
 899        t0 =
 900            ((word32)GetTable8(Tsbox, GETBYTE(s0, 3)) << 24) ^
 901            ((word32)GetTable8(Tsbox, GETBYTE(s1, 2)) << 16) ^
 902            ((word32)GetTable8(Tsbox, GETBYTE(s2, 1)) <<  8) ^
 903            ((word32)GetTable8(Tsbox, GETBYTE(s3, 0)));
 904        t1 =
 905            ((word32)GetTable8(Tsbox, GETBYTE(s1, 3)) << 24) ^
 906            ((word32)GetTable8(Tsbox, GETBYTE(s2, 2)) << 16) ^
 907            ((word32)GetTable8(Tsbox, GETBYTE(s3, 1)) <<  8) ^
 908            ((word32)GetTable8(Tsbox, GETBYTE(s0, 0)));
 909        t2 =
 910            ((word32)GetTable8(Tsbox, GETBYTE(s2, 3)) << 24) ^
 911            ((word32)GetTable8(Tsbox, GETBYTE(s3, 2)) << 16) ^
 912            ((word32)GetTable8(Tsbox, GETBYTE(s0, 1)) <<  8) ^
 913            ((word32)GetTable8(Tsbox, GETBYTE(s1, 0)));
 914        t3 =
 915            ((word32)GetTable8(Tsbox, GETBYTE(s3, 3)) << 24) ^
 916            ((word32)GetTable8(Tsbox, GETBYTE(s0, 2)) << 16) ^
 917            ((word32)GetTable8(Tsbox, GETBYTE(s1, 1)) <<  8) ^
 918            ((word32)GetTable8(Tsbox, GETBYTE(s2, 0)));
 919        s0 = t0 ^ rk[0];
 920        s1 = t1 ^ rk[1];
 921        s2 = t2 ^ rk[2];
 922        s3 = t3 ^ rk[3];
 923#endif
 924
 925        /* write out */
 926#ifdef LITTLE_ENDIAN_ORDER
 927        s0 = ByteReverseWord32(s0);
 928        s1 = ByteReverseWord32(s1);
 929        s2 = ByteReverseWord32(s2);
 930        s3 = ByteReverseWord32(s3);
 931#endif
 932
 933        XMEMCPY(outBlock,                  &s0, sizeof(s0));
 934        XMEMCPY(outBlock +     sizeof(s0), &s1, sizeof(s1));
 935        XMEMCPY(outBlock + 2 * sizeof(s0), &s2, sizeof(s2));
 936        XMEMCPY(outBlock + 3 * sizeof(s0), &s3, sizeof(s3));
 937    }
 938}
 939
 940void AesEncrypt_C(Aes* aes, const byte* inBlock, byte* outBlock,
 941        word32 r)
 942{
 943    byte *inBlock_GPU = NULL;
 944    byte *outBlock_GPU = NULL;
 945    word32* rk_GPU = NULL;
 946    cudaError_t ret = cudaSuccess;
 947
 948#ifdef WC_C_DYNAMIC_FALLBACK
 949    if ( ret == cudaSuccess )
 950        ret = cudaMalloc(&rk_GPU, sizeof(aes->key_C_fallback));
 951    if ( ret == cudaSuccess )
 952        ret = cudaMemcpy(rk_GPU, aes->key_C_fallback, sizeof(aes->key_C_fallback), cudaMemcpyDefault);
 953#else
 954    if ( ret == cudaSuccess )
 955        ret = cudaMalloc(&rk_GPU, sizeof(aes->key));
 956    if ( ret == cudaSuccess )
 957        ret = cudaMemcpy(rk_GPU, aes->key, sizeof(aes->key), cudaMemcpyDefault);
 958#endif
 959
 960    if ( ret == cudaSuccess )
 961        ret = cudaMalloc(&inBlock_GPU, WC_AES_BLOCK_SIZE);
 962    if ( ret == cudaSuccess )
 963        ret = cudaMemcpy(inBlock_GPU, inBlock, WC_AES_BLOCK_SIZE, cudaMemcpyDefault);
 964
 965    if ( ret == cudaSuccess )
 966        ret = cudaMalloc(&outBlock_GPU, WC_AES_BLOCK_SIZE);
 967
 968    if ( ret == cudaSuccess )
 969        AesEncrypt_C_CUDA<<<1,1>>>(rk_GPU, inBlock_GPU, outBlock_GPU, r, 1);
 970
 971    if ( ret == cudaSuccess )
 972        ret = cudaMemcpy(outBlock, outBlock_GPU, WC_AES_BLOCK_SIZE, cudaMemcpyDefault);
 973
 974    cudaFree(inBlock_GPU);
 975    cudaFree(outBlock_GPU);
 976    cudaFree(rk_GPU);
 977}
 978
 979#if defined(HAVE_AES_ECB) && !(defined(WOLFSSL_IMX6_CAAM) && \
 980    !defined(NO_IMX6_CAAM_AES) && !defined(WOLFSSL_QNX_CAAM))
 981/* Encrypt a number of blocks using AES.
 982 *
 983 * @param [in]  aes  AES object.
 984 * @param [in]  in   Block to encrypt.
 985 * @param [out] out  Encrypted block.
 986 * @param [in]  sz   Number of blocks to encrypt.
 987 */
 988void AesEncryptBlocks_C(Aes* aes, const byte* in, byte* out, word32 sz)
 989{
 990    byte *in_GPU = NULL;
 991    byte *out_GPU = NULL;
 992    word32* rk_GPU = NULL;
 993    cudaError_t ret = cudaSuccess;
 994
 995#ifdef WC_C_DYNAMIC_FALLBACK
 996    if ( ret == cudaSuccess )
 997        ret = cudaMalloc(&rk_GPU, sizeof(aes->key_C_fallback));
 998    if ( ret == cudaSuccess )
 999        ret = cudaMemcpy(rk_GPU, aes->key_C_fallback, sizeof(aes->key_C_fallback), cudaMemcpyDefault);
1000#else
1001    if ( ret == cudaSuccess )
1002        ret = cudaMalloc(&rk_GPU, sizeof(aes->key));
1003    if ( ret == cudaSuccess )
1004        ret = cudaMemcpy(rk_GPU, aes->key, sizeof(aes->key), cudaMemcpyDefault);
1005#endif
1006
1007    if ( ret == cudaSuccess )
1008        ret = cudaMalloc(&in_GPU, sz);
1009    if ( ret == cudaSuccess )
1010        ret = cudaMemcpy(in_GPU, in, sz, cudaMemcpyDefault);
1011
1012    if ( ret == cudaSuccess )
1013        ret = cudaMalloc(&out_GPU, sz);
1014
1015    if ( ret == cudaSuccess ) {
1016        int blockSize = 256;
1017        int numBlocks = (sz / WC_AES_BLOCK_SIZE + blockSize - 1) / blockSize;
1018        AesEncrypt_C_CUDA<<<numBlocks,blockSize>>>(rk_GPU, in_GPU, out_GPU, aes->rounds >> 1, sz / WC_AES_BLOCK_SIZE);
1019    }
1020
1021    if ( ret == cudaSuccess )
1022        ret = cudaMemcpy(out, out_GPU, sz, cudaMemcpyDefault);
1023
1024    cudaFree(in_GPU);
1025    cudaFree(out_GPU);
1026    cudaFree(rk_GPU);
1027}
1028#endif
1029
1030#else
1031
1032/* Encrypt a block using AES.
1033 *
1034 * @param [in]  aes       AES object.
1035 * @param [in]  inBlock   Block to encrypt.
1036 * @param [out] outBlock  Encrypted block.
1037 * @param [in]  r         Rounds divided by 2.
1038 */
1039__global__
1040void AesEncrypt_C_CUDA(Aes* aes, const byte* inBlock, byte* outBlock,
1041        word32 r)
1042{
1043    bs_word state[AES_BLOCK_BITS];
1044
1045    (void)r;
1046
1047    XMEMCPY(state, inBlock, WC_AES_BLOCK_SIZE);
1048    XMEMSET(((byte*)state) + WC_AES_BLOCK_SIZE, 0, sizeof(state) - WC_AES_BLOCK_SIZE);
1049
1050    bs_encrypt(state, aes->bs_key, aes->rounds);
1051
1052    XMEMCPY(outBlock, state, WC_AES_BLOCK_SIZE);
1053}
1054
1055void AesEncrypt_C(Aes* aes, const byte* inBlock, byte* outBlock,
1056        word32 r)
1057{
1058    AesEncrypt_C_CUDA<<<1,1>>>(aes, inBlock, outBlock, r);
1059}
1060
1061#if defined(HAVE_AES_ECB) && !(defined(WOLFSSL_IMX6_CAAM) && \
1062    !defined(NO_IMX6_CAAM_AES) && !defined(WOLFSSL_QNX_CAAM))
1063/* Encrypt a number of blocks using AES.
1064 *
1065 * @param [in]  aes  AES object.
1066 * @param [in]  in   Block to encrypt.
1067 * @param [out] out  Encrypted block.
1068 * @param [in]  sz   Number of blocks to encrypt.
1069 */
1070void AesEncryptBlocks_C(Aes* aes, const byte* in, byte* out, word32 sz)
1071{
1072    bs_word state[AES_BLOCK_BITS];
1073
1074    while (sz >= BS_BLOCK_SIZE) {
1075        XMEMCPY(state, in, BS_BLOCK_SIZE);
1076        bs_encrypt(state, aes->bs_key, aes->rounds);
1077        XMEMCPY(out, state, BS_BLOCK_SIZE);
1078        sz  -= BS_BLOCK_SIZE;
1079        in  += BS_BLOCK_SIZE;
1080        out += BS_BLOCK_SIZE;
1081    }
1082    if (sz > 0) {
1083        XMEMCPY(state, in, sz);
1084        XMEMSET(((byte*)state) + sz, 0, sizeof(state) - sz);
1085        bs_encrypt(state, aes->bs_key, aes->rounds);
1086        XMEMCPY(out, state, sz);
1087    }
1088}
1089#endif
1090
1091#endif /* !WC_AES_BITSLICED */
1092
1093#endif /* HAVE_CUDA */
1094
1095#endif /* !WOLFSSL_TI_CRYPT */
1096
1097} /* extern "C" */