diff --git a/README.md b/README.md
index 42f00b1..065cd3e 100644
--- a/README.md
+++ b/README.md
@@ -1,2 +1,24 @@
-# aes-xts-pur64
+# aes-xts-pur64
 OpenCL code for aes-xts256-plain64 encryption compatible with LUKS
+
+
+This program is free software: you can redistribute it and/or modify
+it under the terms of the GNU General Public License as published by
+the Free Software Foundation, either version 3 of the License, or
+(at your option) any later version.
+
+This code is based on work of Hashcat project and its contributors:
+https://github.com/hashcat/hashcat
+
+It modified to be more simple and OpenCLish.
+It must be compatible with most OpenCL2.0+ implementations.
+
+Code contains methods for encryption and decryption with aes-xts256-plain64 algorithm.
+And also small python test suite.
+Code tested to be compatible with LUKS encryption.
+
+It supposed to be executed inside individual work items.
+See test_aes_xts256_plain.cl for example usage.
+
+
+Copyright (C) 2025  Kirill Shakirov
diff --git a/src/aes256_xts_plain.cl b/src/aes256_xts_plain.cl
new file mode 100644
index 0000000..f2dc7aa
--- /dev/null
+++ b/src/aes256_xts_plain.cl
@@ -0,0 +1,1087 @@
+// aes-xts-pur64 is OpenCL code for aes-xts256-plain64 encryption compatible with LUKS
+//
+//     Copyright (C) 2025  Kirill Shakirov
+//
+//     This program is free software: you can redistribute it and/or modify
+//     it under the terms of the GNU General Public License as published by
+//     the Free Software Foundation, either version 3 of the License, or
+//     (at your option) any later version.
+//
+//     This program is distributed in the hope that it will be useful,
+//     but WITHOUT ANY WARRANTY; without even the implied warranty of
+//     MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+//     GNU General Public License for more details.
+//
+//     You should have received a copy of the GNU General Public License
+//     along with this program.  If not, see <https://www.gnu.org/licenses/>.
+
+
+#include "aes256_xts_plain.h"
+
+__constant const uint te0[256] =
+{
+  0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
+  0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
+  0x60303050, 0x02010103, 0xce6767a9, 0x562b2b7d,
+  0xe7fefe19, 0xb5d7d762, 0x4dababe6, 0xec76769a,
+  0x8fcaca45, 0x1f82829d, 0x89c9c940, 0xfa7d7d87,
+  0xeffafa15, 0xb25959eb, 0x8e4747c9, 0xfbf0f00b,
+  0x41adadec, 0xb3d4d467, 0x5fa2a2fd, 0x45afafea,
+  0x239c9cbf, 0x53a4a4f7, 0xe4727296, 0x9bc0c05b,
+  0x75b7b7c2, 0xe1fdfd1c, 0x3d9393ae, 0x4c26266a,
+  0x6c36365a, 0x7e3f3f41, 0xf5f7f702, 0x83cccc4f,
+  0x6834345c, 0x51a5a5f4, 0xd1e5e534, 0xf9f1f108,
+  0xe2717193, 0xabd8d873, 0x62313153, 0x2a15153f,
+  0x0804040c, 0x95c7c752, 0x46232365, 0x9dc3c35e,
+  0x30181828, 0x379696a1, 0x0a05050f, 0x2f9a9ab5,
+  0x0e070709, 0x24121236, 0x1b80809b, 0xdfe2e23d,
+  0xcdebeb26, 0x4e272769, 0x7fb2b2cd, 0xea75759f,
+  0x1209091b, 0x1d83839e, 0x582c2c74, 0x341a1a2e,
+  0x361b1b2d, 0xdc6e6eb2, 0xb45a5aee, 0x5ba0a0fb,
+  0xa45252f6, 0x763b3b4d, 0xb7d6d661, 0x7db3b3ce,
+  0x5229297b, 0xdde3e33e, 0x5e2f2f71, 0x13848497,
+  0xa65353f5, 0xb9d1d168, 0x00000000, 0xc1eded2c,
+  0x40202060, 0xe3fcfc1f, 0x79b1b1c8, 0xb65b5bed,
+  0xd46a6abe, 0x8dcbcb46, 0x67bebed9, 0x7239394b,
+  0x944a4ade, 0x984c4cd4, 0xb05858e8, 0x85cfcf4a,
+  0xbbd0d06b, 0xc5efef2a, 0x4faaaae5, 0xedfbfb16,
+  0x864343c5, 0x9a4d4dd7, 0x66333355, 0x11858594,
+  0x8a4545cf, 0xe9f9f910, 0x04020206, 0xfe7f7f81,
+  0xa05050f0, 0x783c3c44, 0x259f9fba, 0x4ba8a8e3,
+  0xa25151f3, 0x5da3a3fe, 0x804040c0, 0x058f8f8a,
+  0x3f9292ad, 0x219d9dbc, 0x70383848, 0xf1f5f504,
+  0x63bcbcdf, 0x77b6b6c1, 0xafdada75, 0x42212163,
+  0x20101030, 0xe5ffff1a, 0xfdf3f30e, 0xbfd2d26d,
+  0x81cdcd4c, 0x180c0c14, 0x26131335, 0xc3ecec2f,
+  0xbe5f5fe1, 0x359797a2, 0x884444cc, 0x2e171739,
+  0x93c4c457, 0x55a7a7f2, 0xfc7e7e82, 0x7a3d3d47,
+  0xc86464ac, 0xba5d5de7, 0x3219192b, 0xe6737395,
+  0xc06060a0, 0x19818198, 0x9e4f4fd1, 0xa3dcdc7f,
+  0x44222266, 0x542a2a7e, 0x3b9090ab, 0x0b888883,
+  0x8c4646ca, 0xc7eeee29, 0x6bb8b8d3, 0x2814143c,
+  0xa7dede79, 0xbc5e5ee2, 0x160b0b1d, 0xaddbdb76,
+  0xdbe0e03b, 0x64323256, 0x743a3a4e, 0x140a0a1e,
+  0x924949db, 0x0c06060a, 0x4824246c, 0xb85c5ce4,
+  0x9fc2c25d, 0xbdd3d36e, 0x43acacef, 0xc46262a6,
+  0x399191a8, 0x319595a4, 0xd3e4e437, 0xf279798b,
+  0xd5e7e732, 0x8bc8c843, 0x6e373759, 0xda6d6db7,
+  0x018d8d8c, 0xb1d5d564, 0x9c4e4ed2, 0x49a9a9e0,
+  0xd86c6cb4, 0xac5656fa, 0xf3f4f407, 0xcfeaea25,
+  0xca6565af, 0xf47a7a8e, 0x47aeaee9, 0x10080818,
+  0x6fbabad5, 0xf0787888, 0x4a25256f, 0x5c2e2e72,
+  0x381c1c24, 0x57a6a6f1, 0x73b4b4c7, 0x97c6c651,
+  0xcbe8e823, 0xa1dddd7c, 0xe874749c, 0x3e1f1f21,
+  0x964b4bdd, 0x61bdbddc, 0x0d8b8b86, 0x0f8a8a85,
+  0xe0707090, 0x7c3e3e42, 0x71b5b5c4, 0xcc6666aa,
+  0x904848d8, 0x06030305, 0xf7f6f601, 0x1c0e0e12,
+  0xc26161a3, 0x6a35355f, 0xae5757f9, 0x69b9b9d0,
+  0x17868691, 0x99c1c158, 0x3a1d1d27, 0x279e9eb9,
+  0xd9e1e138, 0xebf8f813, 0x2b9898b3, 0x22111133,
+  0xd26969bb, 0xa9d9d970, 0x078e8e89, 0x339494a7,
+  0x2d9b9bb6, 0x3c1e1e22, 0x15878792, 0xc9e9e920,
+  0x87cece49, 0xaa5555ff, 0x50282878, 0xa5dfdf7a,
+  0x038c8c8f, 0x59a1a1f8, 0x09898980, 0x1a0d0d17,
+  0x65bfbfda, 0xd7e6e631, 0x844242c6, 0xd06868b8,
+  0x824141c3, 0x299999b0, 0x5a2d2d77, 0x1e0f0f11,
+  0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
+};
+
+__constant const uint te1[256] =
+{
+  0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
+  0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
+  0x50603030, 0x03020101, 0xa9ce6767, 0x7d562b2b,
+  0x19e7fefe, 0x62b5d7d7, 0xe64dabab, 0x9aec7676,
+  0x458fcaca, 0x9d1f8282, 0x4089c9c9, 0x87fa7d7d,
+  0x15effafa, 0xebb25959, 0xc98e4747, 0x0bfbf0f0,
+  0xec41adad, 0x67b3d4d4, 0xfd5fa2a2, 0xea45afaf,
+  0xbf239c9c, 0xf753a4a4, 0x96e47272, 0x5b9bc0c0,
+  0xc275b7b7, 0x1ce1fdfd, 0xae3d9393, 0x6a4c2626,
+  0x5a6c3636, 0x417e3f3f, 0x02f5f7f7, 0x4f83cccc,
+  0x5c683434, 0xf451a5a5, 0x34d1e5e5, 0x08f9f1f1,
+  0x93e27171, 0x73abd8d8, 0x53623131, 0x3f2a1515,
+  0x0c080404, 0x5295c7c7, 0x65462323, 0x5e9dc3c3,
+  0x28301818, 0xa1379696, 0x0f0a0505, 0xb52f9a9a,
+  0x090e0707, 0x36241212, 0x9b1b8080, 0x3ddfe2e2,
+  0x26cdebeb, 0x694e2727, 0xcd7fb2b2, 0x9fea7575,
+  0x1b120909, 0x9e1d8383, 0x74582c2c, 0x2e341a1a,
+  0x2d361b1b, 0xb2dc6e6e, 0xeeb45a5a, 0xfb5ba0a0,
+  0xf6a45252, 0x4d763b3b, 0x61b7d6d6, 0xce7db3b3,
+  0x7b522929, 0x3edde3e3, 0x715e2f2f, 0x97138484,
+  0xf5a65353, 0x68b9d1d1, 0x00000000, 0x2cc1eded,
+  0x60402020, 0x1fe3fcfc, 0xc879b1b1, 0xedb65b5b,
+  0xbed46a6a, 0x468dcbcb, 0xd967bebe, 0x4b723939,
+  0xde944a4a, 0xd4984c4c, 0xe8b05858, 0x4a85cfcf,
+  0x6bbbd0d0, 0x2ac5efef, 0xe54faaaa, 0x16edfbfb,
+  0xc5864343, 0xd79a4d4d, 0x55663333, 0x94118585,
+  0xcf8a4545, 0x10e9f9f9, 0x06040202, 0x81fe7f7f,
+  0xf0a05050, 0x44783c3c, 0xba259f9f, 0xe34ba8a8,
+  0xf3a25151, 0xfe5da3a3, 0xc0804040, 0x8a058f8f,
+  0xad3f9292, 0xbc219d9d, 0x48703838, 0x04f1f5f5,
+  0xdf63bcbc, 0xc177b6b6, 0x75afdada, 0x63422121,
+  0x30201010, 0x1ae5ffff, 0x0efdf3f3, 0x6dbfd2d2,
+  0x4c81cdcd, 0x14180c0c, 0x35261313, 0x2fc3ecec,
+  0xe1be5f5f, 0xa2359797, 0xcc884444, 0x392e1717,
+  0x5793c4c4, 0xf255a7a7, 0x82fc7e7e, 0x477a3d3d,
+  0xacc86464, 0xe7ba5d5d, 0x2b321919, 0x95e67373,
+  0xa0c06060, 0x98198181, 0xd19e4f4f, 0x7fa3dcdc,
+  0x66442222, 0x7e542a2a, 0xab3b9090, 0x830b8888,
+  0xca8c4646, 0x29c7eeee, 0xd36bb8b8, 0x3c281414,
+  0x79a7dede, 0xe2bc5e5e, 0x1d160b0b, 0x76addbdb,
+  0x3bdbe0e0, 0x56643232, 0x4e743a3a, 0x1e140a0a,
+  0xdb924949, 0x0a0c0606, 0x6c482424, 0xe4b85c5c,
+  0x5d9fc2c2, 0x6ebdd3d3, 0xef43acac, 0xa6c46262,
+  0xa8399191, 0xa4319595, 0x37d3e4e4, 0x8bf27979,
+  0x32d5e7e7, 0x438bc8c8, 0x596e3737, 0xb7da6d6d,
+  0x8c018d8d, 0x64b1d5d5, 0xd29c4e4e, 0xe049a9a9,
+  0xb4d86c6c, 0xfaac5656, 0x07f3f4f4, 0x25cfeaea,
+  0xafca6565, 0x8ef47a7a, 0xe947aeae, 0x18100808,
+  0xd56fbaba, 0x88f07878, 0x6f4a2525, 0x725c2e2e,
+  0x24381c1c, 0xf157a6a6, 0xc773b4b4, 0x5197c6c6,
+  0x23cbe8e8, 0x7ca1dddd, 0x9ce87474, 0x213e1f1f,
+  0xdd964b4b, 0xdc61bdbd, 0x860d8b8b, 0x850f8a8a,
+  0x90e07070, 0x427c3e3e, 0xc471b5b5, 0xaacc6666,
+  0xd8904848, 0x05060303, 0x01f7f6f6, 0x121c0e0e,
+  0xa3c26161, 0x5f6a3535, 0xf9ae5757, 0xd069b9b9,
+  0x91178686, 0x5899c1c1, 0x273a1d1d, 0xb9279e9e,
+  0x38d9e1e1, 0x13ebf8f8, 0xb32b9898, 0x33221111,
+  0xbbd26969, 0x70a9d9d9, 0x89078e8e, 0xa7339494,
+  0xb62d9b9b, 0x223c1e1e, 0x92158787, 0x20c9e9e9,
+  0x4987cece, 0xffaa5555, 0x78502828, 0x7aa5dfdf,
+  0x8f038c8c, 0xf859a1a1, 0x80098989, 0x171a0d0d,
+  0xda65bfbf, 0x31d7e6e6, 0xc6844242, 0xb8d06868,
+  0xc3824141, 0xb0299999, 0x775a2d2d, 0x111e0f0f,
+  0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
+};
+
+__constant const uint te2[256] =
+{
+  0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
+  0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
+  0x30506030, 0x01030201, 0x67a9ce67, 0x2b7d562b,
+  0xfe19e7fe, 0xd762b5d7, 0xabe64dab, 0x769aec76,
+  0xca458fca, 0x829d1f82, 0xc94089c9, 0x7d87fa7d,
+  0xfa15effa, 0x59ebb259, 0x47c98e47, 0xf00bfbf0,
+  0xadec41ad, 0xd467b3d4, 0xa2fd5fa2, 0xafea45af,
+  0x9cbf239c, 0xa4f753a4, 0x7296e472, 0xc05b9bc0,
+  0xb7c275b7, 0xfd1ce1fd, 0x93ae3d93, 0x266a4c26,
+  0x365a6c36, 0x3f417e3f, 0xf702f5f7, 0xcc4f83cc,
+  0x345c6834, 0xa5f451a5, 0xe534d1e5, 0xf108f9f1,
+  0x7193e271, 0xd873abd8, 0x31536231, 0x153f2a15,
+  0x040c0804, 0xc75295c7, 0x23654623, 0xc35e9dc3,
+  0x18283018, 0x96a13796, 0x050f0a05, 0x9ab52f9a,
+  0x07090e07, 0x12362412, 0x809b1b80, 0xe23ddfe2,
+  0xeb26cdeb, 0x27694e27, 0xb2cd7fb2, 0x759fea75,
+  0x091b1209, 0x839e1d83, 0x2c74582c, 0x1a2e341a,
+  0x1b2d361b, 0x6eb2dc6e, 0x5aeeb45a, 0xa0fb5ba0,
+  0x52f6a452, 0x3b4d763b, 0xd661b7d6, 0xb3ce7db3,
+  0x297b5229, 0xe33edde3, 0x2f715e2f, 0x84971384,
+  0x53f5a653, 0xd168b9d1, 0x00000000, 0xed2cc1ed,
+  0x20604020, 0xfc1fe3fc, 0xb1c879b1, 0x5bedb65b,
+  0x6abed46a, 0xcb468dcb, 0xbed967be, 0x394b7239,
+  0x4ade944a, 0x4cd4984c, 0x58e8b058, 0xcf4a85cf,
+  0xd06bbbd0, 0xef2ac5ef, 0xaae54faa, 0xfb16edfb,
+  0x43c58643, 0x4dd79a4d, 0x33556633, 0x85941185,
+  0x45cf8a45, 0xf910e9f9, 0x02060402, 0x7f81fe7f,
+  0x50f0a050, 0x3c44783c, 0x9fba259f, 0xa8e34ba8,
+  0x51f3a251, 0xa3fe5da3, 0x40c08040, 0x8f8a058f,
+  0x92ad3f92, 0x9dbc219d, 0x38487038, 0xf504f1f5,
+  0xbcdf63bc, 0xb6c177b6, 0xda75afda, 0x21634221,
+  0x10302010, 0xff1ae5ff, 0xf30efdf3, 0xd26dbfd2,
+  0xcd4c81cd, 0x0c14180c, 0x13352613, 0xec2fc3ec,
+  0x5fe1be5f, 0x97a23597, 0x44cc8844, 0x17392e17,
+  0xc45793c4, 0xa7f255a7, 0x7e82fc7e, 0x3d477a3d,
+  0x64acc864, 0x5de7ba5d, 0x192b3219, 0x7395e673,
+  0x60a0c060, 0x81981981, 0x4fd19e4f, 0xdc7fa3dc,
+  0x22664422, 0x2a7e542a, 0x90ab3b90, 0x88830b88,
+  0x46ca8c46, 0xee29c7ee, 0xb8d36bb8, 0x143c2814,
+  0xde79a7de, 0x5ee2bc5e, 0x0b1d160b, 0xdb76addb,
+  0xe03bdbe0, 0x32566432, 0x3a4e743a, 0x0a1e140a,
+  0x49db9249, 0x060a0c06, 0x246c4824, 0x5ce4b85c,
+  0xc25d9fc2, 0xd36ebdd3, 0xacef43ac, 0x62a6c462,
+  0x91a83991, 0x95a43195, 0xe437d3e4, 0x798bf279,
+  0xe732d5e7, 0xc8438bc8, 0x37596e37, 0x6db7da6d,
+  0x8d8c018d, 0xd564b1d5, 0x4ed29c4e, 0xa9e049a9,
+  0x6cb4d86c, 0x56faac56, 0xf407f3f4, 0xea25cfea,
+  0x65afca65, 0x7a8ef47a, 0xaee947ae, 0x08181008,
+  0xbad56fba, 0x7888f078, 0x256f4a25, 0x2e725c2e,
+  0x1c24381c, 0xa6f157a6, 0xb4c773b4, 0xc65197c6,
+  0xe823cbe8, 0xdd7ca1dd, 0x749ce874, 0x1f213e1f,
+  0x4bdd964b, 0xbddc61bd, 0x8b860d8b, 0x8a850f8a,
+  0x7090e070, 0x3e427c3e, 0xb5c471b5, 0x66aacc66,
+  0x48d89048, 0x03050603, 0xf601f7f6, 0x0e121c0e,
+  0x61a3c261, 0x355f6a35, 0x57f9ae57, 0xb9d069b9,
+  0x86911786, 0xc15899c1, 0x1d273a1d, 0x9eb9279e,
+  0xe138d9e1, 0xf813ebf8, 0x98b32b98, 0x11332211,
+  0x69bbd269, 0xd970a9d9, 0x8e89078e, 0x94a73394,
+  0x9bb62d9b, 0x1e223c1e, 0x87921587, 0xe920c9e9,
+  0xce4987ce, 0x55ffaa55, 0x28785028, 0xdf7aa5df,
+  0x8c8f038c, 0xa1f859a1, 0x89800989, 0x0d171a0d,
+  0xbfda65bf, 0xe631d7e6, 0x42c68442, 0x68b8d068,
+  0x41c38241, 0x99b02999, 0x2d775a2d, 0x0f111e0f,
+  0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
+};
+
+__constant const uint te3[256] =
+{
+  0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
+  0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
+  0x30305060, 0x01010302, 0x6767a9ce, 0x2b2b7d56,
+  0xfefe19e7, 0xd7d762b5, 0xababe64d, 0x76769aec,
+  0xcaca458f, 0x82829d1f, 0xc9c94089, 0x7d7d87fa,
+  0xfafa15ef, 0x5959ebb2, 0x4747c98e, 0xf0f00bfb,
+  0xadadec41, 0xd4d467b3, 0xa2a2fd5f, 0xafafea45,
+  0x9c9cbf23, 0xa4a4f753, 0x727296e4, 0xc0c05b9b,
+  0xb7b7c275, 0xfdfd1ce1, 0x9393ae3d, 0x26266a4c,
+  0x36365a6c, 0x3f3f417e, 0xf7f702f5, 0xcccc4f83,
+  0x34345c68, 0xa5a5f451, 0xe5e534d1, 0xf1f108f9,
+  0x717193e2, 0xd8d873ab, 0x31315362, 0x15153f2a,
+  0x04040c08, 0xc7c75295, 0x23236546, 0xc3c35e9d,
+  0x18182830, 0x9696a137, 0x05050f0a, 0x9a9ab52f,
+  0x0707090e, 0x12123624, 0x80809b1b, 0xe2e23ddf,
+  0xebeb26cd, 0x2727694e, 0xb2b2cd7f, 0x75759fea,
+  0x09091b12, 0x83839e1d, 0x2c2c7458, 0x1a1a2e34,
+  0x1b1b2d36, 0x6e6eb2dc, 0x5a5aeeb4, 0xa0a0fb5b,
+  0x5252f6a4, 0x3b3b4d76, 0xd6d661b7, 0xb3b3ce7d,
+  0x29297b52, 0xe3e33edd, 0x2f2f715e, 0x84849713,
+  0x5353f5a6, 0xd1d168b9, 0x00000000, 0xeded2cc1,
+  0x20206040, 0xfcfc1fe3, 0xb1b1c879, 0x5b5bedb6,
+  0x6a6abed4, 0xcbcb468d, 0xbebed967, 0x39394b72,
+  0x4a4ade94, 0x4c4cd498, 0x5858e8b0, 0xcfcf4a85,
+  0xd0d06bbb, 0xefef2ac5, 0xaaaae54f, 0xfbfb16ed,
+  0x4343c586, 0x4d4dd79a, 0x33335566, 0x85859411,
+  0x4545cf8a, 0xf9f910e9, 0x02020604, 0x7f7f81fe,
+  0x5050f0a0, 0x3c3c4478, 0x9f9fba25, 0xa8a8e34b,
+  0x5151f3a2, 0xa3a3fe5d, 0x4040c080, 0x8f8f8a05,
+  0x9292ad3f, 0x9d9dbc21, 0x38384870, 0xf5f504f1,
+  0xbcbcdf63, 0xb6b6c177, 0xdada75af, 0x21216342,
+  0x10103020, 0xffff1ae5, 0xf3f30efd, 0xd2d26dbf,
+  0xcdcd4c81, 0x0c0c1418, 0x13133526, 0xecec2fc3,
+  0x5f5fe1be, 0x9797a235, 0x4444cc88, 0x1717392e,
+  0xc4c45793, 0xa7a7f255, 0x7e7e82fc, 0x3d3d477a,
+  0x6464acc8, 0x5d5de7ba, 0x19192b32, 0x737395e6,
+  0x6060a0c0, 0x81819819, 0x4f4fd19e, 0xdcdc7fa3,
+  0x22226644, 0x2a2a7e54, 0x9090ab3b, 0x8888830b,
+  0x4646ca8c, 0xeeee29c7, 0xb8b8d36b, 0x14143c28,
+  0xdede79a7, 0x5e5ee2bc, 0x0b0b1d16, 0xdbdb76ad,
+  0xe0e03bdb, 0x32325664, 0x3a3a4e74, 0x0a0a1e14,
+  0x4949db92, 0x06060a0c, 0x24246c48, 0x5c5ce4b8,
+  0xc2c25d9f, 0xd3d36ebd, 0xacacef43, 0x6262a6c4,
+  0x9191a839, 0x9595a431, 0xe4e437d3, 0x79798bf2,
+  0xe7e732d5, 0xc8c8438b, 0x3737596e, 0x6d6db7da,
+  0x8d8d8c01, 0xd5d564b1, 0x4e4ed29c, 0xa9a9e049,
+  0x6c6cb4d8, 0x5656faac, 0xf4f407f3, 0xeaea25cf,
+  0x6565afca, 0x7a7a8ef4, 0xaeaee947, 0x08081810,
+  0xbabad56f, 0x787888f0, 0x25256f4a, 0x2e2e725c,
+  0x1c1c2438, 0xa6a6f157, 0xb4b4c773, 0xc6c65197,
+  0xe8e823cb, 0xdddd7ca1, 0x74749ce8, 0x1f1f213e,
+  0x4b4bdd96, 0xbdbddc61, 0x8b8b860d, 0x8a8a850f,
+  0x707090e0, 0x3e3e427c, 0xb5b5c471, 0x6666aacc,
+  0x4848d890, 0x03030506, 0xf6f601f7, 0x0e0e121c,
+  0x6161a3c2, 0x35355f6a, 0x5757f9ae, 0xb9b9d069,
+  0x86869117, 0xc1c15899, 0x1d1d273a, 0x9e9eb927,
+  0xe1e138d9, 0xf8f813eb, 0x9898b32b, 0x11113322,
+  0x6969bbd2, 0xd9d970a9, 0x8e8e8907, 0x9494a733,
+  0x9b9bb62d, 0x1e1e223c, 0x87879215, 0xe9e920c9,
+  0xcece4987, 0x5555ffaa, 0x28287850, 0xdfdf7aa5,
+  0x8c8c8f03, 0xa1a1f859, 0x89898009, 0x0d0d171a,
+  0xbfbfda65, 0xe6e631d7, 0x4242c684, 0x6868b8d0,
+  0x4141c382, 0x9999b029, 0x2d2d775a, 0x0f0f111e,
+  0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
+};
+
+__constant const uint te4[256] =
+{
+  0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
+  0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
+  0x30303030, 0x01010101, 0x67676767, 0x2b2b2b2b,
+  0xfefefefe, 0xd7d7d7d7, 0xabababab, 0x76767676,
+  0xcacacaca, 0x82828282, 0xc9c9c9c9, 0x7d7d7d7d,
+  0xfafafafa, 0x59595959, 0x47474747, 0xf0f0f0f0,
+  0xadadadad, 0xd4d4d4d4, 0xa2a2a2a2, 0xafafafaf,
+  0x9c9c9c9c, 0xa4a4a4a4, 0x72727272, 0xc0c0c0c0,
+  0xb7b7b7b7, 0xfdfdfdfd, 0x93939393, 0x26262626,
+  0x36363636, 0x3f3f3f3f, 0xf7f7f7f7, 0xcccccccc,
+  0x34343434, 0xa5a5a5a5, 0xe5e5e5e5, 0xf1f1f1f1,
+  0x71717171, 0xd8d8d8d8, 0x31313131, 0x15151515,
+  0x04040404, 0xc7c7c7c7, 0x23232323, 0xc3c3c3c3,
+  0x18181818, 0x96969696, 0x05050505, 0x9a9a9a9a,
+  0x07070707, 0x12121212, 0x80808080, 0xe2e2e2e2,
+  0xebebebeb, 0x27272727, 0xb2b2b2b2, 0x75757575,
+  0x09090909, 0x83838383, 0x2c2c2c2c, 0x1a1a1a1a,
+  0x1b1b1b1b, 0x6e6e6e6e, 0x5a5a5a5a, 0xa0a0a0a0,
+  0x52525252, 0x3b3b3b3b, 0xd6d6d6d6, 0xb3b3b3b3,
+  0x29292929, 0xe3e3e3e3, 0x2f2f2f2f, 0x84848484,
+  0x53535353, 0xd1d1d1d1, 0x00000000, 0xedededed,
+  0x20202020, 0xfcfcfcfc, 0xb1b1b1b1, 0x5b5b5b5b,
+  0x6a6a6a6a, 0xcbcbcbcb, 0xbebebebe, 0x39393939,
+  0x4a4a4a4a, 0x4c4c4c4c, 0x58585858, 0xcfcfcfcf,
+  0xd0d0d0d0, 0xefefefef, 0xaaaaaaaa, 0xfbfbfbfb,
+  0x43434343, 0x4d4d4d4d, 0x33333333, 0x85858585,
+  0x45454545, 0xf9f9f9f9, 0x02020202, 0x7f7f7f7f,
+  0x50505050, 0x3c3c3c3c, 0x9f9f9f9f, 0xa8a8a8a8,
+  0x51515151, 0xa3a3a3a3, 0x40404040, 0x8f8f8f8f,
+  0x92929292, 0x9d9d9d9d, 0x38383838, 0xf5f5f5f5,
+  0xbcbcbcbc, 0xb6b6b6b6, 0xdadadada, 0x21212121,
+  0x10101010, 0xffffffff, 0xf3f3f3f3, 0xd2d2d2d2,
+  0xcdcdcdcd, 0x0c0c0c0c, 0x13131313, 0xecececec,
+  0x5f5f5f5f, 0x97979797, 0x44444444, 0x17171717,
+  0xc4c4c4c4, 0xa7a7a7a7, 0x7e7e7e7e, 0x3d3d3d3d,
+  0x64646464, 0x5d5d5d5d, 0x19191919, 0x73737373,
+  0x60606060, 0x81818181, 0x4f4f4f4f, 0xdcdcdcdc,
+  0x22222222, 0x2a2a2a2a, 0x90909090, 0x88888888,
+  0x46464646, 0xeeeeeeee, 0xb8b8b8b8, 0x14141414,
+  0xdededede, 0x5e5e5e5e, 0x0b0b0b0b, 0xdbdbdbdb,
+  0xe0e0e0e0, 0x32323232, 0x3a3a3a3a, 0x0a0a0a0a,
+  0x49494949, 0x06060606, 0x24242424, 0x5c5c5c5c,
+  0xc2c2c2c2, 0xd3d3d3d3, 0xacacacac, 0x62626262,
+  0x91919191, 0x95959595, 0xe4e4e4e4, 0x79797979,
+  0xe7e7e7e7, 0xc8c8c8c8, 0x37373737, 0x6d6d6d6d,
+  0x8d8d8d8d, 0xd5d5d5d5, 0x4e4e4e4e, 0xa9a9a9a9,
+  0x6c6c6c6c, 0x56565656, 0xf4f4f4f4, 0xeaeaeaea,
+  0x65656565, 0x7a7a7a7a, 0xaeaeaeae, 0x08080808,
+  0xbabababa, 0x78787878, 0x25252525, 0x2e2e2e2e,
+  0x1c1c1c1c, 0xa6a6a6a6, 0xb4b4b4b4, 0xc6c6c6c6,
+  0xe8e8e8e8, 0xdddddddd, 0x74747474, 0x1f1f1f1f,
+  0x4b4b4b4b, 0xbdbdbdbd, 0x8b8b8b8b, 0x8a8a8a8a,
+  0x70707070, 0x3e3e3e3e, 0xb5b5b5b5, 0x66666666,
+  0x48484848, 0x03030303, 0xf6f6f6f6, 0x0e0e0e0e,
+  0x61616161, 0x35353535, 0x57575757, 0xb9b9b9b9,
+  0x86868686, 0xc1c1c1c1, 0x1d1d1d1d, 0x9e9e9e9e,
+  0xe1e1e1e1, 0xf8f8f8f8, 0x98989898, 0x11111111,
+  0x69696969, 0xd9d9d9d9, 0x8e8e8e8e, 0x94949494,
+  0x9b9b9b9b, 0x1e1e1e1e, 0x87878787, 0xe9e9e9e9,
+  0xcececece, 0x55555555, 0x28282828, 0xdfdfdfdf,
+  0x8c8c8c8c, 0xa1a1a1a1, 0x89898989, 0x0d0d0d0d,
+  0xbfbfbfbf, 0xe6e6e6e6, 0x42424242, 0x68686868,
+  0x41414141, 0x99999999, 0x2d2d2d2d, 0x0f0f0f0f,
+  0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
+};
+
+__constant const uint td0[256] =
+{
+  0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
+  0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
+  0x2030fa55, 0xad766df6, 0x88cc7691, 0xf5024c25,
+  0x4fe5d7fc, 0xc52acbd7, 0x26354480, 0xb562a38f,
+  0xdeb15a49, 0x25ba1b67, 0x45ea0e98, 0x5dfec0e1,
+  0xc32f7502, 0x814cf012, 0x8d4697a3, 0x6bd3f9c6,
+  0x038f5fe7, 0x15929c95, 0xbf6d7aeb, 0x955259da,
+  0xd4be832d, 0x587421d3, 0x49e06929, 0x8ec9c844,
+  0x75c2896a, 0xf48e7978, 0x99583e6b, 0x27b971dd,
+  0xbee14fb6, 0xf088ad17, 0xc920ac66, 0x7dce3ab4,
+  0x63df4a18, 0xe51a3182, 0x97513360, 0x62537f45,
+  0xb16477e0, 0xbb6bae84, 0xfe81a01c, 0xf9082b94,
+  0x70486858, 0x8f45fd19, 0x94de6c87, 0x527bf8b7,
+  0xab73d323, 0x724b02e2, 0xe31f8f57, 0x6655ab2a,
+  0xb2eb2807, 0x2fb5c203, 0x86c57b9a, 0xd33708a5,
+  0x302887f2, 0x23bfa5b2, 0x02036aba, 0xed16825c,
+  0x8acf1c2b, 0xa779b492, 0xf307f2f0, 0x4e69e2a1,
+  0x65daf4cd, 0x0605bed5, 0xd134621f, 0xc4a6fe8a,
+  0x342e539d, 0xa2f355a0, 0x058ae132, 0xa4f6eb75,
+  0x0b83ec39, 0x4060efaa, 0x5e719f06, 0xbd6e1051,
+  0x3e218af9, 0x96dd063d, 0xdd3e05ae, 0x4de6bd46,
+  0x91548db5, 0x71c45d05, 0x0406d46f, 0x605015ff,
+  0x1998fb24, 0xd6bde997, 0x894043cc, 0x67d99e77,
+  0xb0e842bd, 0x07898b88, 0xe7195b38, 0x79c8eedb,
+  0xa17c0a47, 0x7c420fe9, 0xf8841ec9, 0x00000000,
+  0x09808683, 0x322bed48, 0x1e1170ac, 0x6c5a724e,
+  0xfd0efffb, 0x0f853856, 0x3daed51e, 0x362d3927,
+  0x0a0fd964, 0x685ca621, 0x9b5b54d1, 0x24362e3a,
+  0x0c0a67b1, 0x9357e70f, 0xb4ee96d2, 0x1b9b919e,
+  0x80c0c54f, 0x61dc20a2, 0x5a774b69, 0x1c121a16,
+  0xe293ba0a, 0xc0a02ae5, 0x3c22e043, 0x121b171d,
+  0x0e090d0b, 0xf28bc7ad, 0x2db6a8b9, 0x141ea9c8,
+  0x57f11985, 0xaf75074c, 0xee99ddbb, 0xa37f60fd,
+  0xf701269f, 0x5c72f5bc, 0x44663bc5, 0x5bfb7e34,
+  0x8b432976, 0xcb23c6dc, 0xb6edfc68, 0xb8e4f163,
+  0xd731dcca, 0x42638510, 0x13972240, 0x84c61120,
+  0x854a247d, 0xd2bb3df8, 0xaef93211, 0xc729a16d,
+  0x1d9e2f4b, 0xdcb230f3, 0x0d8652ec, 0x77c1e3d0,
+  0x2bb3166c, 0xa970b999, 0x119448fa, 0x47e96422,
+  0xa8fc8cc4, 0xa0f03f1a, 0x567d2cd8, 0x223390ef,
+  0x87494ec7, 0xd938d1c1, 0x8ccaa2fe, 0x98d40b36,
+  0xa6f581cf, 0xa57ade28, 0xdab78e26, 0x3fadbfa4,
+  0x2c3a9de4, 0x5078920d, 0x6a5fcc9b, 0x547e4662,
+  0xf68d13c2, 0x90d8b8e8, 0x2e39f75e, 0x82c3aff5,
+  0x9f5d80be, 0x69d0937c, 0x6fd52da9, 0xcf2512b3,
+  0xc8ac993b, 0x10187da7, 0xe89c636e, 0xdb3bbb7b,
+  0xcd267809, 0x6e5918f4, 0xec9ab701, 0x834f9aa8,
+  0xe6956e65, 0xaaffe67e, 0x21bccf08, 0xef15e8e6,
+  0xbae79bd9, 0x4a6f36ce, 0xea9f09d4, 0x29b07cd6,
+  0x31a4b2af, 0x2a3f2331, 0xc6a59430, 0x35a266c0,
+  0x744ebc37, 0xfc82caa6, 0xe090d0b0, 0x33a7d815,
+  0xf104984a, 0x41ecdaf7, 0x7fcd500e, 0x1791f62f,
+  0x764dd68d, 0x43efb04d, 0xccaa4d54, 0xe49604df,
+  0x9ed1b5e3, 0x4c6a881b, 0xc12c1fb8, 0x4665517f,
+  0x9d5eea04, 0x018c355d, 0xfa877473, 0xfb0b412e,
+  0xb3671d5a, 0x92dbd252, 0xe9105633, 0x6dd64713,
+  0x9ad7618c, 0x37a10c7a, 0x59f8148e, 0xeb133c89,
+  0xcea927ee, 0xb761c935, 0xe11ce5ed, 0x7a47b13c,
+  0x9cd2df59, 0x55f2733f, 0x1814ce79, 0x73c737bf,
+  0x53f7cdea, 0x5ffdaa5b, 0xdf3d6f14, 0x7844db86,
+  0xcaaff381, 0xb968c43e, 0x3824342c, 0xc2a3405f,
+  0x161dc372, 0xbce2250c, 0x283c498b, 0xff0d9541,
+  0x39a80171, 0x080cb3de, 0xd8b4e49c, 0x6456c190,
+  0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
+};
+
+__constant const uint td1[256] =
+{
+  0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
+  0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
+  0x552030fa, 0xf6ad766d, 0x9188cc76, 0x25f5024c,
+  0xfc4fe5d7, 0xd7c52acb, 0x80263544, 0x8fb562a3,
+  0x49deb15a, 0x6725ba1b, 0x9845ea0e, 0xe15dfec0,
+  0x02c32f75, 0x12814cf0, 0xa38d4697, 0xc66bd3f9,
+  0xe7038f5f, 0x9515929c, 0xebbf6d7a, 0xda955259,
+  0x2dd4be83, 0xd3587421, 0x2949e069, 0x448ec9c8,
+  0x6a75c289, 0x78f48e79, 0x6b99583e, 0xdd27b971,
+  0xb6bee14f, 0x17f088ad, 0x66c920ac, 0xb47dce3a,
+  0x1863df4a, 0x82e51a31, 0x60975133, 0x4562537f,
+  0xe0b16477, 0x84bb6bae, 0x1cfe81a0, 0x94f9082b,
+  0x58704868, 0x198f45fd, 0x8794de6c, 0xb7527bf8,
+  0x23ab73d3, 0xe2724b02, 0x57e31f8f, 0x2a6655ab,
+  0x07b2eb28, 0x032fb5c2, 0x9a86c57b, 0xa5d33708,
+  0xf2302887, 0xb223bfa5, 0xba02036a, 0x5ced1682,
+  0x2b8acf1c, 0x92a779b4, 0xf0f307f2, 0xa14e69e2,
+  0xcd65daf4, 0xd50605be, 0x1fd13462, 0x8ac4a6fe,
+  0x9d342e53, 0xa0a2f355, 0x32058ae1, 0x75a4f6eb,
+  0x390b83ec, 0xaa4060ef, 0x065e719f, 0x51bd6e10,
+  0xf93e218a, 0x3d96dd06, 0xaedd3e05, 0x464de6bd,
+  0xb591548d, 0x0571c45d, 0x6f0406d4, 0xff605015,
+  0x241998fb, 0x97d6bde9, 0xcc894043, 0x7767d99e,
+  0xbdb0e842, 0x8807898b, 0x38e7195b, 0xdb79c8ee,
+  0x47a17c0a, 0xe97c420f, 0xc9f8841e, 0x00000000,
+  0x83098086, 0x48322bed, 0xac1e1170, 0x4e6c5a72,
+  0xfbfd0eff, 0x560f8538, 0x1e3daed5, 0x27362d39,
+  0x640a0fd9, 0x21685ca6, 0xd19b5b54, 0x3a24362e,
+  0xb10c0a67, 0x0f9357e7, 0xd2b4ee96, 0x9e1b9b91,
+  0x4f80c0c5, 0xa261dc20, 0x695a774b, 0x161c121a,
+  0x0ae293ba, 0xe5c0a02a, 0x433c22e0, 0x1d121b17,
+  0x0b0e090d, 0xadf28bc7, 0xb92db6a8, 0xc8141ea9,
+  0x8557f119, 0x4caf7507, 0xbbee99dd, 0xfda37f60,
+  0x9ff70126, 0xbc5c72f5, 0xc544663b, 0x345bfb7e,
+  0x768b4329, 0xdccb23c6, 0x68b6edfc, 0x63b8e4f1,
+  0xcad731dc, 0x10426385, 0x40139722, 0x2084c611,
+  0x7d854a24, 0xf8d2bb3d, 0x11aef932, 0x6dc729a1,
+  0x4b1d9e2f, 0xf3dcb230, 0xec0d8652, 0xd077c1e3,
+  0x6c2bb316, 0x99a970b9, 0xfa119448, 0x2247e964,
+  0xc4a8fc8c, 0x1aa0f03f, 0xd8567d2c, 0xef223390,
+  0xc787494e, 0xc1d938d1, 0xfe8ccaa2, 0x3698d40b,
+  0xcfa6f581, 0x28a57ade, 0x26dab78e, 0xa43fadbf,
+  0xe42c3a9d, 0x0d507892, 0x9b6a5fcc, 0x62547e46,
+  0xc2f68d13, 0xe890d8b8, 0x5e2e39f7, 0xf582c3af,
+  0xbe9f5d80, 0x7c69d093, 0xa96fd52d, 0xb3cf2512,
+  0x3bc8ac99, 0xa710187d, 0x6ee89c63, 0x7bdb3bbb,
+  0x09cd2678, 0xf46e5918, 0x01ec9ab7, 0xa8834f9a,
+  0x65e6956e, 0x7eaaffe6, 0x0821bccf, 0xe6ef15e8,
+  0xd9bae79b, 0xce4a6f36, 0xd4ea9f09, 0xd629b07c,
+  0xaf31a4b2, 0x312a3f23, 0x30c6a594, 0xc035a266,
+  0x37744ebc, 0xa6fc82ca, 0xb0e090d0, 0x1533a7d8,
+  0x4af10498, 0xf741ecda, 0x0e7fcd50, 0x2f1791f6,
+  0x8d764dd6, 0x4d43efb0, 0x54ccaa4d, 0xdfe49604,
+  0xe39ed1b5, 0x1b4c6a88, 0xb8c12c1f, 0x7f466551,
+  0x049d5eea, 0x5d018c35, 0x73fa8774, 0x2efb0b41,
+  0x5ab3671d, 0x5292dbd2, 0x33e91056, 0x136dd647,
+  0x8c9ad761, 0x7a37a10c, 0x8e59f814, 0x89eb133c,
+  0xeecea927, 0x35b761c9, 0xede11ce5, 0x3c7a47b1,
+  0x599cd2df, 0x3f55f273, 0x791814ce, 0xbf73c737,
+  0xea53f7cd, 0x5b5ffdaa, 0x14df3d6f, 0x867844db,
+  0x81caaff3, 0x3eb968c4, 0x2c382434, 0x5fc2a340,
+  0x72161dc3, 0x0cbce225, 0x8b283c49, 0x41ff0d95,
+  0x7139a801, 0xde080cb3, 0x9cd8b4e4, 0x906456c1,
+  0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
+};
+
+__constant const uint td2[256] =
+{
+  0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
+  0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
+  0xfa552030, 0x6df6ad76, 0x769188cc, 0x4c25f502,
+  0xd7fc4fe5, 0xcbd7c52a, 0x44802635, 0xa38fb562,
+  0x5a49deb1, 0x1b6725ba, 0x0e9845ea, 0xc0e15dfe,
+  0x7502c32f, 0xf012814c, 0x97a38d46, 0xf9c66bd3,
+  0x5fe7038f, 0x9c951592, 0x7aebbf6d, 0x59da9552,
+  0x832dd4be, 0x21d35874, 0x692949e0, 0xc8448ec9,
+  0x896a75c2, 0x7978f48e, 0x3e6b9958, 0x71dd27b9,
+  0x4fb6bee1, 0xad17f088, 0xac66c920, 0x3ab47dce,
+  0x4a1863df, 0x3182e51a, 0x33609751, 0x7f456253,
+  0x77e0b164, 0xae84bb6b, 0xa01cfe81, 0x2b94f908,
+  0x68587048, 0xfd198f45, 0x6c8794de, 0xf8b7527b,
+  0xd323ab73, 0x02e2724b, 0x8f57e31f, 0xab2a6655,
+  0x2807b2eb, 0xc2032fb5, 0x7b9a86c5, 0x08a5d337,
+  0x87f23028, 0xa5b223bf, 0x6aba0203, 0x825ced16,
+  0x1c2b8acf, 0xb492a779, 0xf2f0f307, 0xe2a14e69,
+  0xf4cd65da, 0xbed50605, 0x621fd134, 0xfe8ac4a6,
+  0x539d342e, 0x55a0a2f3, 0xe132058a, 0xeb75a4f6,
+  0xec390b83, 0xefaa4060, 0x9f065e71, 0x1051bd6e,
+  0x8af93e21, 0x063d96dd, 0x05aedd3e, 0xbd464de6,
+  0x8db59154, 0x5d0571c4, 0xd46f0406, 0x15ff6050,
+  0xfb241998, 0xe997d6bd, 0x43cc8940, 0x9e7767d9,
+  0x42bdb0e8, 0x8b880789, 0x5b38e719, 0xeedb79c8,
+  0x0a47a17c, 0x0fe97c42, 0x1ec9f884, 0x00000000,
+  0x86830980, 0xed48322b, 0x70ac1e11, 0x724e6c5a,
+  0xfffbfd0e, 0x38560f85, 0xd51e3dae, 0x3927362d,
+  0xd9640a0f, 0xa621685c, 0x54d19b5b, 0x2e3a2436,
+  0x67b10c0a, 0xe70f9357, 0x96d2b4ee, 0x919e1b9b,
+  0xc54f80c0, 0x20a261dc, 0x4b695a77, 0x1a161c12,
+  0xba0ae293, 0x2ae5c0a0, 0xe0433c22, 0x171d121b,
+  0x0d0b0e09, 0xc7adf28b, 0xa8b92db6, 0xa9c8141e,
+  0x198557f1, 0x074caf75, 0xddbbee99, 0x60fda37f,
+  0x269ff701, 0xf5bc5c72, 0x3bc54466, 0x7e345bfb,
+  0x29768b43, 0xc6dccb23, 0xfc68b6ed, 0xf163b8e4,
+  0xdccad731, 0x85104263, 0x22401397, 0x112084c6,
+  0x247d854a, 0x3df8d2bb, 0x3211aef9, 0xa16dc729,
+  0x2f4b1d9e, 0x30f3dcb2, 0x52ec0d86, 0xe3d077c1,
+  0x166c2bb3, 0xb999a970, 0x48fa1194, 0x642247e9,
+  0x8cc4a8fc, 0x3f1aa0f0, 0x2cd8567d, 0x90ef2233,
+  0x4ec78749, 0xd1c1d938, 0xa2fe8cca, 0x0b3698d4,
+  0x81cfa6f5, 0xde28a57a, 0x8e26dab7, 0xbfa43fad,
+  0x9de42c3a, 0x920d5078, 0xcc9b6a5f, 0x4662547e,
+  0x13c2f68d, 0xb8e890d8, 0xf75e2e39, 0xaff582c3,
+  0x80be9f5d, 0x937c69d0, 0x2da96fd5, 0x12b3cf25,
+  0x993bc8ac, 0x7da71018, 0x636ee89c, 0xbb7bdb3b,
+  0x7809cd26, 0x18f46e59, 0xb701ec9a, 0x9aa8834f,
+  0x6e65e695, 0xe67eaaff, 0xcf0821bc, 0xe8e6ef15,
+  0x9bd9bae7, 0x36ce4a6f, 0x09d4ea9f, 0x7cd629b0,
+  0xb2af31a4, 0x23312a3f, 0x9430c6a5, 0x66c035a2,
+  0xbc37744e, 0xcaa6fc82, 0xd0b0e090, 0xd81533a7,
+  0x984af104, 0xdaf741ec, 0x500e7fcd, 0xf62f1791,
+  0xd68d764d, 0xb04d43ef, 0x4d54ccaa, 0x04dfe496,
+  0xb5e39ed1, 0x881b4c6a, 0x1fb8c12c, 0x517f4665,
+  0xea049d5e, 0x355d018c, 0x7473fa87, 0x412efb0b,
+  0x1d5ab367, 0xd25292db, 0x5633e910, 0x47136dd6,
+  0x618c9ad7, 0x0c7a37a1, 0x148e59f8, 0x3c89eb13,
+  0x27eecea9, 0xc935b761, 0xe5ede11c, 0xb13c7a47,
+  0xdf599cd2, 0x733f55f2, 0xce791814, 0x37bf73c7,
+  0xcdea53f7, 0xaa5b5ffd, 0x6f14df3d, 0xdb867844,
+  0xf381caaf, 0xc43eb968, 0x342c3824, 0x405fc2a3,
+  0xc372161d, 0x250cbce2, 0x498b283c, 0x9541ff0d,
+  0x017139a8, 0xb3de080c, 0xe49cd8b4, 0xc1906456,
+  0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
+};
+
+__constant const uint td3[256] =
+{
+  0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
+  0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
+  0x30fa5520, 0x766df6ad, 0xcc769188, 0x024c25f5,
+  0xe5d7fc4f, 0x2acbd7c5, 0x35448026, 0x62a38fb5,
+  0xb15a49de, 0xba1b6725, 0xea0e9845, 0xfec0e15d,
+  0x2f7502c3, 0x4cf01281, 0x4697a38d, 0xd3f9c66b,
+  0x8f5fe703, 0x929c9515, 0x6d7aebbf, 0x5259da95,
+  0xbe832dd4, 0x7421d358, 0xe0692949, 0xc9c8448e,
+  0xc2896a75, 0x8e7978f4, 0x583e6b99, 0xb971dd27,
+  0xe14fb6be, 0x88ad17f0, 0x20ac66c9, 0xce3ab47d,
+  0xdf4a1863, 0x1a3182e5, 0x51336097, 0x537f4562,
+  0x6477e0b1, 0x6bae84bb, 0x81a01cfe, 0x082b94f9,
+  0x48685870, 0x45fd198f, 0xde6c8794, 0x7bf8b752,
+  0x73d323ab, 0x4b02e272, 0x1f8f57e3, 0x55ab2a66,
+  0xeb2807b2, 0xb5c2032f, 0xc57b9a86, 0x3708a5d3,
+  0x2887f230, 0xbfa5b223, 0x036aba02, 0x16825ced,
+  0xcf1c2b8a, 0x79b492a7, 0x07f2f0f3, 0x69e2a14e,
+  0xdaf4cd65, 0x05bed506, 0x34621fd1, 0xa6fe8ac4,
+  0x2e539d34, 0xf355a0a2, 0x8ae13205, 0xf6eb75a4,
+  0x83ec390b, 0x60efaa40, 0x719f065e, 0x6e1051bd,
+  0x218af93e, 0xdd063d96, 0x3e05aedd, 0xe6bd464d,
+  0x548db591, 0xc45d0571, 0x06d46f04, 0x5015ff60,
+  0x98fb2419, 0xbde997d6, 0x4043cc89, 0xd99e7767,
+  0xe842bdb0, 0x898b8807, 0x195b38e7, 0xc8eedb79,
+  0x7c0a47a1, 0x420fe97c, 0x841ec9f8, 0x00000000,
+  0x80868309, 0x2bed4832, 0x1170ac1e, 0x5a724e6c,
+  0x0efffbfd, 0x8538560f, 0xaed51e3d, 0x2d392736,
+  0x0fd9640a, 0x5ca62168, 0x5b54d19b, 0x362e3a24,
+  0x0a67b10c, 0x57e70f93, 0xee96d2b4, 0x9b919e1b,
+  0xc0c54f80, 0xdc20a261, 0x774b695a, 0x121a161c,
+  0x93ba0ae2, 0xa02ae5c0, 0x22e0433c, 0x1b171d12,
+  0x090d0b0e, 0x8bc7adf2, 0xb6a8b92d, 0x1ea9c814,
+  0xf1198557, 0x75074caf, 0x99ddbbee, 0x7f60fda3,
+  0x01269ff7, 0x72f5bc5c, 0x663bc544, 0xfb7e345b,
+  0x4329768b, 0x23c6dccb, 0xedfc68b6, 0xe4f163b8,
+  0x31dccad7, 0x63851042, 0x97224013, 0xc6112084,
+  0x4a247d85, 0xbb3df8d2, 0xf93211ae, 0x29a16dc7,
+  0x9e2f4b1d, 0xb230f3dc, 0x8652ec0d, 0xc1e3d077,
+  0xb3166c2b, 0x70b999a9, 0x9448fa11, 0xe9642247,
+  0xfc8cc4a8, 0xf03f1aa0, 0x7d2cd856, 0x3390ef22,
+  0x494ec787, 0x38d1c1d9, 0xcaa2fe8c, 0xd40b3698,
+  0xf581cfa6, 0x7ade28a5, 0xb78e26da, 0xadbfa43f,
+  0x3a9de42c, 0x78920d50, 0x5fcc9b6a, 0x7e466254,
+  0x8d13c2f6, 0xd8b8e890, 0x39f75e2e, 0xc3aff582,
+  0x5d80be9f, 0xd0937c69, 0xd52da96f, 0x2512b3cf,
+  0xac993bc8, 0x187da710, 0x9c636ee8, 0x3bbb7bdb,
+  0x267809cd, 0x5918f46e, 0x9ab701ec, 0x4f9aa883,
+  0x956e65e6, 0xffe67eaa, 0xbccf0821, 0x15e8e6ef,
+  0xe79bd9ba, 0x6f36ce4a, 0x9f09d4ea, 0xb07cd629,
+  0xa4b2af31, 0x3f23312a, 0xa59430c6, 0xa266c035,
+  0x4ebc3774, 0x82caa6fc, 0x90d0b0e0, 0xa7d81533,
+  0x04984af1, 0xecdaf741, 0xcd500e7f, 0x91f62f17,
+  0x4dd68d76, 0xefb04d43, 0xaa4d54cc, 0x9604dfe4,
+  0xd1b5e39e, 0x6a881b4c, 0x2c1fb8c1, 0x65517f46,
+  0x5eea049d, 0x8c355d01, 0x877473fa, 0x0b412efb,
+  0x671d5ab3, 0xdbd25292, 0x105633e9, 0xd647136d,
+  0xd7618c9a, 0xa10c7a37, 0xf8148e59, 0x133c89eb,
+  0xa927eece, 0x61c935b7, 0x1ce5ede1, 0x47b13c7a,
+  0xd2df599c, 0xf2733f55, 0x14ce7918, 0xc737bf73,
+  0xf7cdea53, 0xfdaa5b5f, 0x3d6f14df, 0x44db8678,
+  0xaff381ca, 0x68c43eb9, 0x24342c38, 0xa3405fc2,
+  0x1dc37216, 0xe2250cbc, 0x3c498b28, 0x0d9541ff,
+  0xa8017139, 0x0cb3de08, 0xb4e49cd8, 0x56c19064,
+  0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
+};
+
+__constant const uint td4[256] =
+{
+  0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
+  0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
+  0xbfbfbfbf, 0x40404040, 0xa3a3a3a3, 0x9e9e9e9e,
+  0x81818181, 0xf3f3f3f3, 0xd7d7d7d7, 0xfbfbfbfb,
+  0x7c7c7c7c, 0xe3e3e3e3, 0x39393939, 0x82828282,
+  0x9b9b9b9b, 0x2f2f2f2f, 0xffffffff, 0x87878787,
+  0x34343434, 0x8e8e8e8e, 0x43434343, 0x44444444,
+  0xc4c4c4c4, 0xdededede, 0xe9e9e9e9, 0xcbcbcbcb,
+  0x54545454, 0x7b7b7b7b, 0x94949494, 0x32323232,
+  0xa6a6a6a6, 0xc2c2c2c2, 0x23232323, 0x3d3d3d3d,
+  0xeeeeeeee, 0x4c4c4c4c, 0x95959595, 0x0b0b0b0b,
+  0x42424242, 0xfafafafa, 0xc3c3c3c3, 0x4e4e4e4e,
+  0x08080808, 0x2e2e2e2e, 0xa1a1a1a1, 0x66666666,
+  0x28282828, 0xd9d9d9d9, 0x24242424, 0xb2b2b2b2,
+  0x76767676, 0x5b5b5b5b, 0xa2a2a2a2, 0x49494949,
+  0x6d6d6d6d, 0x8b8b8b8b, 0xd1d1d1d1, 0x25252525,
+  0x72727272, 0xf8f8f8f8, 0xf6f6f6f6, 0x64646464,
+  0x86868686, 0x68686868, 0x98989898, 0x16161616,
+  0xd4d4d4d4, 0xa4a4a4a4, 0x5c5c5c5c, 0xcccccccc,
+  0x5d5d5d5d, 0x65656565, 0xb6b6b6b6, 0x92929292,
+  0x6c6c6c6c, 0x70707070, 0x48484848, 0x50505050,
+  0xfdfdfdfd, 0xedededed, 0xb9b9b9b9, 0xdadadada,
+  0x5e5e5e5e, 0x15151515, 0x46464646, 0x57575757,
+  0xa7a7a7a7, 0x8d8d8d8d, 0x9d9d9d9d, 0x84848484,
+  0x90909090, 0xd8d8d8d8, 0xabababab, 0x00000000,
+  0x8c8c8c8c, 0xbcbcbcbc, 0xd3d3d3d3, 0x0a0a0a0a,
+  0xf7f7f7f7, 0xe4e4e4e4, 0x58585858, 0x05050505,
+  0xb8b8b8b8, 0xb3b3b3b3, 0x45454545, 0x06060606,
+  0xd0d0d0d0, 0x2c2c2c2c, 0x1e1e1e1e, 0x8f8f8f8f,
+  0xcacacaca, 0x3f3f3f3f, 0x0f0f0f0f, 0x02020202,
+  0xc1c1c1c1, 0xafafafaf, 0xbdbdbdbd, 0x03030303,
+  0x01010101, 0x13131313, 0x8a8a8a8a, 0x6b6b6b6b,
+  0x3a3a3a3a, 0x91919191, 0x11111111, 0x41414141,
+  0x4f4f4f4f, 0x67676767, 0xdcdcdcdc, 0xeaeaeaea,
+  0x97979797, 0xf2f2f2f2, 0xcfcfcfcf, 0xcececece,
+  0xf0f0f0f0, 0xb4b4b4b4, 0xe6e6e6e6, 0x73737373,
+  0x96969696, 0xacacacac, 0x74747474, 0x22222222,
+  0xe7e7e7e7, 0xadadadad, 0x35353535, 0x85858585,
+  0xe2e2e2e2, 0xf9f9f9f9, 0x37373737, 0xe8e8e8e8,
+  0x1c1c1c1c, 0x75757575, 0xdfdfdfdf, 0x6e6e6e6e,
+  0x47474747, 0xf1f1f1f1, 0x1a1a1a1a, 0x71717171,
+  0x1d1d1d1d, 0x29292929, 0xc5c5c5c5, 0x89898989,
+  0x6f6f6f6f, 0xb7b7b7b7, 0x62626262, 0x0e0e0e0e,
+  0xaaaaaaaa, 0x18181818, 0xbebebebe, 0x1b1b1b1b,
+  0xfcfcfcfc, 0x56565656, 0x3e3e3e3e, 0x4b4b4b4b,
+  0xc6c6c6c6, 0xd2d2d2d2, 0x79797979, 0x20202020,
+  0x9a9a9a9a, 0xdbdbdbdb, 0xc0c0c0c0, 0xfefefefe,
+  0x78787878, 0xcdcdcdcd, 0x5a5a5a5a, 0xf4f4f4f4,
+  0x1f1f1f1f, 0xdddddddd, 0xa8a8a8a8, 0x33333333,
+  0x88888888, 0x07070707, 0xc7c7c7c7, 0x31313131,
+  0xb1b1b1b1, 0x12121212, 0x10101010, 0x59595959,
+  0x27272727, 0x80808080, 0xecececec, 0x5f5f5f5f,
+  0x60606060, 0x51515151, 0x7f7f7f7f, 0xa9a9a9a9,
+  0x19191919, 0xb5b5b5b5, 0x4a4a4a4a, 0x0d0d0d0d,
+  0x2d2d2d2d, 0xe5e5e5e5, 0x7a7a7a7a, 0x9f9f9f9f,
+  0x93939393, 0xc9c9c9c9, 0x9c9c9c9c, 0xefefefef,
+  0xa0a0a0a0, 0xe0e0e0e0, 0x3b3b3b3b, 0x4d4d4d4d,
+  0xaeaeaeae, 0x2a2a2a2a, 0xf5f5f5f5, 0xb0b0b0b0,
+  0xc8c8c8c8, 0xebebebeb, 0xbbbbbbbb, 0x3c3c3c3c,
+  0x83838383, 0x53535353, 0x99999999, 0x61616161,
+  0x17171717, 0x2b2b2b2b, 0x04040404, 0x7e7e7e7e,
+  0xbabababa, 0x77777777, 0xd6d6d6d6, 0x26262626,
+  0xe1e1e1e1, 0x69696969, 0x14141414, 0x63636363,
+  0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d,
+};
+
+
+typedef union {
+    uint i;
+    uchar4 char_v;
+} __attribute__((aligned(32))) ui_uch4_union;
+
+inline uint hc_swap32_S (const uint v)
+{
+  ui_uch4_union r;
+  r.char_v = shuffle(((ui_uch4_union)v).char_v, (uchar4)(3, 2, 1, 0));
+  // r = ((v & 0xff000000) >> 24)
+  //   | ((v & 0x00ff0000) >>  8)
+  //   | ((v & 0x0000ff00) <<  8)
+  //   | ((v & 0x000000ff) << 24);
+
+  return r.i;
+}
+
+// 128 bit key
+inline void aes128_InvertKey (uint *ks)
+{
+  uint temp;
+
+  temp = ks[ 0]; ks[ 0] = ks[40]; ks[40] = temp;
+  temp = ks[ 1]; ks[ 1] = ks[41]; ks[41] = temp;
+  temp = ks[ 2]; ks[ 2] = ks[42]; ks[42] = temp;
+  temp = ks[ 3]; ks[ 3] = ks[43]; ks[43] = temp;
+  temp = ks[ 4]; ks[ 4] = ks[36]; ks[36] = temp;
+  temp = ks[ 5]; ks[ 5] = ks[37]; ks[37] = temp;
+  temp = ks[ 6]; ks[ 6] = ks[38]; ks[38] = temp;
+  temp = ks[ 7]; ks[ 7] = ks[39]; ks[39] = temp;
+  temp = ks[ 8]; ks[ 8] = ks[32]; ks[32] = temp;
+  temp = ks[ 9]; ks[ 9] = ks[33]; ks[33] = temp;
+  temp = ks[10]; ks[10] = ks[34]; ks[34] = temp;
+  temp = ks[11]; ks[11] = ks[35]; ks[35] = temp;
+  temp = ks[12]; ks[12] = ks[28]; ks[28] = temp;
+  temp = ks[13]; ks[13] = ks[29]; ks[29] = temp;
+  temp = ks[14]; ks[14] = ks[30]; ks[30] = temp;
+  temp = ks[15]; ks[15] = ks[31]; ks[31] = temp;
+  temp = ks[16]; ks[16] = ks[24]; ks[24] = temp;
+  temp = ks[17]; ks[17] = ks[25]; ks[25] = temp;
+  temp = ks[18]; ks[18] = ks[26]; ks[26] = temp;
+  temp = ks[19]; ks[19] = ks[27]; ks[27] = temp;
+
+  ks[ 4] = td0[te1[(ks[ 4] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[ 4] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[ 4] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[ 4] >>  0) & 0xff] & 0xff];
+  ks[ 5] = td0[te1[(ks[ 5] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[ 5] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[ 5] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[ 5] >>  0) & 0xff] & 0xff];
+  ks[ 6] = td0[te1[(ks[ 6] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[ 6] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[ 6] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[ 6] >>  0) & 0xff] & 0xff];
+  ks[ 7] = td0[te1[(ks[ 7] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[ 7] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[ 7] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[ 7] >>  0) & 0xff] & 0xff];
+  ks[ 8] = td0[te1[(ks[ 8] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[ 8] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[ 8] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[ 8] >>  0) & 0xff] & 0xff];
+  ks[ 9] = td0[te1[(ks[ 9] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[ 9] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[ 9] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[ 9] >>  0) & 0xff] & 0xff];
+  ks[10] = td0[te1[(ks[10] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[10] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[10] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[10] >>  0) & 0xff] & 0xff];
+  ks[11] = td0[te1[(ks[11] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[11] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[11] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[11] >>  0) & 0xff] & 0xff];
+  ks[12] = td0[te1[(ks[12] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[12] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[12] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[12] >>  0) & 0xff] & 0xff];
+  ks[13] = td0[te1[(ks[13] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[13] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[13] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[13] >>  0) & 0xff] & 0xff];
+  ks[14] = td0[te1[(ks[14] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[14] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[14] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[14] >>  0) & 0xff] & 0xff];
+  ks[15] = td0[te1[(ks[15] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[15] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[15] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[15] >>  0) & 0xff] & 0xff];
+  ks[16] = td0[te1[(ks[16] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[16] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[16] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[16] >>  0) & 0xff] & 0xff];
+  ks[17] = td0[te1[(ks[17] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[17] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[17] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[17] >>  0) & 0xff] & 0xff];
+  ks[18] = td0[te1[(ks[18] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[18] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[18] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[18] >>  0) & 0xff] & 0xff];
+  ks[19] = td0[te1[(ks[19] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[19] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[19] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[19] >>  0) & 0xff] & 0xff];
+  ks[20] = td0[te1[(ks[20] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[20] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[20] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[20] >>  0) & 0xff] & 0xff];
+  ks[21] = td0[te1[(ks[21] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[21] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[21] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[21] >>  0) & 0xff] & 0xff];
+  ks[22] = td0[te1[(ks[22] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[22] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[22] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[22] >>  0) & 0xff] & 0xff];
+  ks[23] = td0[te1[(ks[23] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[23] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[23] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[23] >>  0) & 0xff] & 0xff];
+  ks[24] = td0[te1[(ks[24] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[24] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[24] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[24] >>  0) & 0xff] & 0xff];
+  ks[25] = td0[te1[(ks[25] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[25] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[25] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[25] >>  0) & 0xff] & 0xff];
+  ks[26] = td0[te1[(ks[26] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[26] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[26] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[26] >>  0) & 0xff] & 0xff];
+  ks[27] = td0[te1[(ks[27] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[27] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[27] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[27] >>  0) & 0xff] & 0xff];
+  ks[28] = td0[te1[(ks[28] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[28] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[28] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[28] >>  0) & 0xff] & 0xff];
+  ks[29] = td0[te1[(ks[29] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[29] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[29] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[29] >>  0) & 0xff] & 0xff];
+  ks[30] = td0[te1[(ks[30] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[30] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[30] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[30] >>  0) & 0xff] & 0xff];
+  ks[31] = td0[te1[(ks[31] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[31] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[31] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[31] >>  0) & 0xff] & 0xff];
+  ks[32] = td0[te1[(ks[32] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[32] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[32] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[32] >>  0) & 0xff] & 0xff];
+  ks[33] = td0[te1[(ks[33] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[33] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[33] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[33] >>  0) & 0xff] & 0xff];
+  ks[34] = td0[te1[(ks[34] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[34] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[34] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[34] >>  0) & 0xff] & 0xff];
+  ks[35] = td0[te1[(ks[35] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[35] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[35] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[35] >>  0) & 0xff] & 0xff];
+  ks[36] = td0[te1[(ks[36] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[36] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[36] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[36] >>  0) & 0xff] & 0xff];
+  ks[37] = td0[te1[(ks[37] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[37] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[37] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[37] >>  0) & 0xff] & 0xff];
+  ks[38] = td0[te1[(ks[38] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[38] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[38] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[38] >>  0) & 0xff] & 0xff];
+  ks[39] = td0[te1[(ks[39] >> 24) & 0xff] & 0xff] ^ td1[te1[(ks[39] >> 16) & 0xff] & 0xff] ^ td2[te1[(ks[39] >>  8) & 0xff] & 0xff] ^ td3[te1[(ks[39] >>  0) & 0xff] & 0xff];
+}
+
+inline void aes128_ExpandKey (uint *ks, const uint *ukey)
+{
+  ks[ 0] = ukey[0];
+  ks[ 1] = ukey[1];
+  ks[ 2] = ukey[2];
+  ks[ 3] = ukey[3];
+  ks[ 4] = ks[ 0] ^ 0x01000000
+                  ^ (te2[(ks[ 3] >> 16) & 0xff] & 0xff000000)
+                  ^ (te3[(ks[ 3] >>  8) & 0xff] & 0x00ff0000)
+                  ^ (te0[(ks[ 3] >>  0) & 0xff] & 0x0000ff00)
+                  ^ (te1[(ks[ 3] >> 24) & 0xff] & 0x000000ff);
+  ks[ 5] = ks[ 1] ^ ks[ 4];
+  ks[ 6] = ks[ 2] ^ ks[ 5];
+  ks[ 7] = ks[ 3] ^ ks[ 6];
+  ks[ 8] = ks[ 4] ^ 0x02000000
+                  ^ (te2[(ks[ 7] >> 16) & 0xff] & 0xff000000)
+                  ^ (te3[(ks[ 7] >>  8) & 0xff] & 0x00ff0000)
+                  ^ (te0[(ks[ 7] >>  0) & 0xff] & 0x0000ff00)
+                  ^ (te1[(ks[ 7] >> 24) & 0xff] & 0x000000ff);
+  ks[ 9] = ks[ 5] ^ ks[ 8];
+  ks[10] = ks[ 6] ^ ks[ 9];
+  ks[11] = ks[ 7] ^ ks[10];
+  ks[12] = ks[ 8] ^ 0x04000000
+                  ^ (te2[(ks[11] >> 16) & 0xff] & 0xff000000)
+                  ^ (te3[(ks[11] >>  8) & 0xff] & 0x00ff0000)
+                  ^ (te0[(ks[11] >>  0) & 0xff] & 0x0000ff00)
+                  ^ (te1[(ks[11] >> 24) & 0xff] & 0x000000ff);
+  ks[13] = ks[ 9] ^ ks[12];
+  ks[14] = ks[10] ^ ks[13];
+  ks[15] = ks[11] ^ ks[14];
+  ks[16] = ks[12] ^ 0x08000000
+                  ^ (te2[(ks[15] >> 16) & 0xff] & 0xff000000)
+                  ^ (te3[(ks[15] >>  8) & 0xff] & 0x00ff0000)
+                  ^ (te0[(ks[15] >>  0) & 0xff] & 0x0000ff00)
+                  ^ (te1[(ks[15] >> 24) & 0xff] & 0x000000ff);
+  ks[17] = ks[13] ^ ks[16];
+  ks[18] = ks[14] ^ ks[17];
+  ks[19] = ks[15] ^ ks[18];
+  ks[20] = ks[16] ^ 0x10000000
+                  ^ (te2[(ks[19] >> 16) & 0xff] & 0xff000000)
+                  ^ (te3[(ks[19] >>  8) & 0xff] & 0x00ff0000)
+                  ^ (te0[(ks[19] >>  0) & 0xff] & 0x0000ff00)
+                  ^ (te1[(ks[19] >> 24) & 0xff] & 0x000000ff);
+  ks[21] = ks[17] ^ ks[20];
+  ks[22] = ks[18] ^ ks[21];
+  ks[23] = ks[19] ^ ks[22];
+  ks[24] = ks[20] ^ 0x20000000
+                  ^ (te2[(ks[23] >> 16) & 0xff] & 0xff000000)
+                  ^ (te3[(ks[23] >>  8) & 0xff] & 0x00ff0000)
+                  ^ (te0[(ks[23] >>  0) & 0xff] & 0x0000ff00)
+                  ^ (te1[(ks[23] >> 24) & 0xff] & 0x000000ff);
+  ks[25] = ks[21] ^ ks[24];
+  ks[26] = ks[22] ^ ks[25];
+  ks[27] = ks[23] ^ ks[26];
+  ks[28] = ks[24] ^ 0x40000000
+                  ^ (te2[(ks[27] >> 16) & 0xff] & 0xff000000)
+                  ^ (te3[(ks[27] >>  8) & 0xff] & 0x00ff0000)
+                  ^ (te0[(ks[27] >>  0) & 0xff] & 0x0000ff00)
+                  ^ (te1[(ks[27] >> 24) & 0xff] & 0x000000ff);
+  ks[29] = ks[25] ^ ks[28];
+  ks[30] = ks[26] ^ ks[29];
+  ks[31] = ks[27] ^ ks[30];
+  ks[32] = ks[28] ^ 0x80000000
+                  ^ (te2[(ks[31] >> 16) & 0xff] & 0xff000000)
+                  ^ (te3[(ks[31] >>  8) & 0xff] & 0x00ff0000)
+                  ^ (te0[(ks[31] >>  0) & 0xff] & 0x0000ff00)
+                  ^ (te1[(ks[31] >> 24) & 0xff] & 0x000000ff);
+  ks[33] = ks[29] ^ ks[32];
+  ks[34] = ks[30] ^ ks[33];
+  ks[35] = ks[31] ^ ks[34];
+  ks[36] = ks[32] ^ 0x1b000000
+                  ^ (te2[(ks[35] >> 16) & 0xff] & 0xff000000)
+                  ^ (te3[(ks[35] >>  8) & 0xff] & 0x00ff0000)
+                  ^ (te0[(ks[35] >>  0) & 0xff] & 0x0000ff00)
+                  ^ (te1[(ks[35] >> 24) & 0xff] & 0x000000ff);
+  ks[37] = ks[33] ^ ks[36];
+  ks[38] = ks[34] ^ ks[37];
+  ks[39] = ks[35] ^ ks[38];
+  ks[40] = ks[36] ^ 0x36000000
+                  ^ (te2[(ks[39] >> 16) & 0xff] & 0xff000000)
+                  ^ (te3[(ks[39] >>  8) & 0xff] & 0x00ff0000)
+                  ^ (te0[(ks[39] >>  0) & 0xff] & 0x0000ff00)
+                  ^ (te1[(ks[39] >> 24) & 0xff] & 0x000000ff);
+  ks[41] = ks[37] ^ ks[40];
+  ks[42] = ks[38] ^ ks[41];
+  ks[43] = ks[39] ^ ks[42];
+}
+
+
+inline void aes128_set_encrypt_key (uint *ks, const uint *ukey)
+{
+  uint ukey_s[4];
+
+  ukey_s[0] = hc_swap32_S (ukey[0]);
+  ukey_s[1] = hc_swap32_S (ukey[1]);
+  ukey_s[2] = hc_swap32_S (ukey[2]);
+  ukey_s[3] = hc_swap32_S (ukey[3]);
+
+  aes128_ExpandKey (ks, ukey_s);
+}
+
+inline void aes128_set_decrypt_key (uint *ks, const uint *ukey)
+{
+  uint ukey_s[4];
+
+  ukey_s[0] = hc_swap32_S (ukey[0]);
+  ukey_s[1] = hc_swap32_S (ukey[1]);
+  ukey_s[2] = hc_swap32_S (ukey[2]);
+  ukey_s[3] = hc_swap32_S (ukey[3]);
+
+  aes128_ExpandKey (ks, ukey_s);
+
+  aes128_InvertKey (ks);
+}
+
+inline void aes128_encrypt (const uint *ks, const uint *in, uint *out)
+{
+  uint4 s = (uint4)(hc_swap32_S(in[0]), hc_swap32_S(in[1]), hc_swap32_S(in[2]), hc_swap32_S(in[3]));
+  s ^= vload4(0, ks);
+  
+  uint4 t;
+
+  t.s0 = te0[s.s0 >> 24] ^ te1[(s.s1 >> 16) & 0xff] ^ te2[(s.s2 >>  8) & 0xff] ^ te3[s.s3 & 0xff] ^ ks[ 4];
+  t.s1 = te0[s.s1 >> 24] ^ te1[(s.s2 >> 16) & 0xff] ^ te2[(s.s3 >>  8) & 0xff] ^ te3[s.s0 & 0xff] ^ ks[ 5];
+  t.s2 = te0[s.s2 >> 24] ^ te1[(s.s3 >> 16) & 0xff] ^ te2[(s.s0 >>  8) & 0xff] ^ te3[s.s1 & 0xff] ^ ks[ 6];
+  t.s3 = te0[s.s3 >> 24] ^ te1[(s.s0 >> 16) & 0xff] ^ te2[(s.s1 >>  8) & 0xff] ^ te3[s.s2 & 0xff] ^ ks[ 7];
+  s.s0 = te0[t.s0 >> 24] ^ te1[(t.s1 >> 16) & 0xff] ^ te2[(t.s2 >>  8) & 0xff] ^ te3[t.s3 & 0xff] ^ ks[ 8];
+  s.s1 = te0[t.s1 >> 24] ^ te1[(t.s2 >> 16) & 0xff] ^ te2[(t.s3 >>  8) & 0xff] ^ te3[t.s0 & 0xff] ^ ks[ 9];
+  s.s2 = te0[t.s2 >> 24] ^ te1[(t.s3 >> 16) & 0xff] ^ te2[(t.s0 >>  8) & 0xff] ^ te3[t.s1 & 0xff] ^ ks[10];
+  s.s3 = te0[t.s3 >> 24] ^ te1[(t.s0 >> 16) & 0xff] ^ te2[(t.s1 >>  8) & 0xff] ^ te3[t.s2 & 0xff] ^ ks[11];
+  t.s0 = te0[s.s0 >> 24] ^ te1[(s.s1 >> 16) & 0xff] ^ te2[(s.s2 >>  8) & 0xff] ^ te3[s.s3 & 0xff] ^ ks[12];
+  t.s1 = te0[s.s1 >> 24] ^ te1[(s.s2 >> 16) & 0xff] ^ te2[(s.s3 >>  8) & 0xff] ^ te3[s.s0 & 0xff] ^ ks[13];
+  t.s2 = te0[s.s2 >> 24] ^ te1[(s.s3 >> 16) & 0xff] ^ te2[(s.s0 >>  8) & 0xff] ^ te3[s.s1 & 0xff] ^ ks[14];
+  t.s3 = te0[s.s3 >> 24] ^ te1[(s.s0 >> 16) & 0xff] ^ te2[(s.s1 >>  8) & 0xff] ^ te3[s.s2 & 0xff] ^ ks[15];
+  s.s0 = te0[t.s0 >> 24] ^ te1[(t.s1 >> 16) & 0xff] ^ te2[(t.s2 >>  8) & 0xff] ^ te3[t.s3 & 0xff] ^ ks[16];
+  s.s1 = te0[t.s1 >> 24] ^ te1[(t.s2 >> 16) & 0xff] ^ te2[(t.s3 >>  8) & 0xff] ^ te3[t.s0 & 0xff] ^ ks[17];
+  s.s2 = te0[t.s2 >> 24] ^ te1[(t.s3 >> 16) & 0xff] ^ te2[(t.s0 >>  8) & 0xff] ^ te3[t.s1 & 0xff] ^ ks[18];
+  s.s3 = te0[t.s3 >> 24] ^ te1[(t.s0 >> 16) & 0xff] ^ te2[(t.s1 >>  8) & 0xff] ^ te3[t.s2 & 0xff] ^ ks[19];
+  t.s0 = te0[s.s0 >> 24] ^ te1[(s.s1 >> 16) & 0xff] ^ te2[(s.s2 >>  8) & 0xff] ^ te3[s.s3 & 0xff] ^ ks[20];
+  t.s1 = te0[s.s1 >> 24] ^ te1[(s.s2 >> 16) & 0xff] ^ te2[(s.s3 >>  8) & 0xff] ^ te3[s.s0 & 0xff] ^ ks[21];
+  t.s2 = te0[s.s2 >> 24] ^ te1[(s.s3 >> 16) & 0xff] ^ te2[(s.s0 >>  8) & 0xff] ^ te3[s.s1 & 0xff] ^ ks[22];
+  t.s3 = te0[s.s3 >> 24] ^ te1[(s.s0 >> 16) & 0xff] ^ te2[(s.s1 >>  8) & 0xff] ^ te3[s.s2 & 0xff] ^ ks[23];
+  s.s0 = te0[t.s0 >> 24] ^ te1[(t.s1 >> 16) & 0xff] ^ te2[(t.s2 >>  8) & 0xff] ^ te3[t.s3 & 0xff] ^ ks[24];
+  s.s1 = te0[t.s1 >> 24] ^ te1[(t.s2 >> 16) & 0xff] ^ te2[(t.s3 >>  8) & 0xff] ^ te3[t.s0 & 0xff] ^ ks[25];
+  s.s2 = te0[t.s2 >> 24] ^ te1[(t.s3 >> 16) & 0xff] ^ te2[(t.s0 >>  8) & 0xff] ^ te3[t.s1 & 0xff] ^ ks[26];
+  s.s3 = te0[t.s3 >> 24] ^ te1[(t.s0 >> 16) & 0xff] ^ te2[(t.s1 >>  8) & 0xff] ^ te3[t.s2 & 0xff] ^ ks[27];
+  t.s0 = te0[s.s0 >> 24] ^ te1[(s.s1 >> 16) & 0xff] ^ te2[(s.s2 >>  8) & 0xff] ^ te3[s.s3 & 0xff] ^ ks[28];
+  t.s1 = te0[s.s1 >> 24] ^ te1[(s.s2 >> 16) & 0xff] ^ te2[(s.s3 >>  8) & 0xff] ^ te3[s.s0 & 0xff] ^ ks[29];
+  t.s2 = te0[s.s2 >> 24] ^ te1[(s.s3 >> 16) & 0xff] ^ te2[(s.s0 >>  8) & 0xff] ^ te3[s.s1 & 0xff] ^ ks[30];
+  t.s3 = te0[s.s3 >> 24] ^ te1[(s.s0 >> 16) & 0xff] ^ te2[(s.s1 >>  8) & 0xff] ^ te3[s.s2 & 0xff] ^ ks[31];
+  s.s0 = te0[t.s0 >> 24] ^ te1[(t.s1 >> 16) & 0xff] ^ te2[(t.s2 >>  8) & 0xff] ^ te3[t.s3 & 0xff] ^ ks[32];
+  s.s1 = te0[t.s1 >> 24] ^ te1[(t.s2 >> 16) & 0xff] ^ te2[(t.s3 >>  8) & 0xff] ^ te3[t.s0 & 0xff] ^ ks[33];
+  s.s2 = te0[t.s2 >> 24] ^ te1[(t.s3 >> 16) & 0xff] ^ te2[(t.s0 >>  8) & 0xff] ^ te3[t.s1 & 0xff] ^ ks[34];
+  s.s3 = te0[t.s3 >> 24] ^ te1[(t.s0 >> 16) & 0xff] ^ te2[(t.s1 >>  8) & 0xff] ^ te3[t.s2 & 0xff] ^ ks[35];
+  t.s0 = te0[s.s0 >> 24] ^ te1[(s.s1 >> 16) & 0xff] ^ te2[(s.s2 >>  8) & 0xff] ^ te3[s.s3 & 0xff] ^ ks[36];
+  t.s1 = te0[s.s1 >> 24] ^ te1[(s.s2 >> 16) & 0xff] ^ te2[(s.s3 >>  8) & 0xff] ^ te3[s.s0 & 0xff] ^ ks[37];
+  t.s2 = te0[s.s2 >> 24] ^ te1[(s.s3 >> 16) & 0xff] ^ te2[(s.s0 >>  8) & 0xff] ^ te3[s.s1 & 0xff] ^ ks[38];
+  t.s3 = te0[s.s3 >> 24] ^ te1[(s.s0 >> 16) & 0xff] ^ te2[(s.s1 >>  8) & 0xff] ^ te3[s.s2 & 0xff] ^ ks[39];
+
+  out[0] = (te4[(t.s0 >> 24) & 0xff] & 0xff000000)
+         ^ (te4[(t.s1 >> 16) & 0xff] & 0x00ff0000)
+         ^ (te4[(t.s2 >>  8) & 0xff] & 0x0000ff00)
+         ^ (te4[(t.s3 >>  0) & 0xff] & 0x000000ff)
+         ^ ks[40];
+
+  out[1] = (te4[(t.s1 >> 24) & 0xff] & 0xff000000)
+         ^ (te4[(t.s2 >> 16) & 0xff] & 0x00ff0000)
+         ^ (te4[(t.s3 >>  8) & 0xff] & 0x0000ff00)
+         ^ (te4[(t.s0 >>  0) & 0xff] & 0x000000ff)
+         ^ ks[41];
+
+  out[2] = (te4[(t.s2 >> 24) & 0xff] & 0xff000000)
+         ^ (te4[(t.s3 >> 16) & 0xff] & 0x00ff0000)
+         ^ (te4[(t.s0 >>  8) & 0xff] & 0x0000ff00)
+         ^ (te4[(t.s1 >>  0) & 0xff] & 0x000000ff)
+         ^ ks[42];
+
+  out[3] = (te4[(t.s3 >> 24) & 0xff] & 0xff000000)
+         ^ (te4[(t.s0 >> 16) & 0xff] & 0x00ff0000)
+         ^ (te4[(t.s1 >>  8) & 0xff] & 0x0000ff00)
+         ^ (te4[(t.s2 >>  0) & 0xff] & 0x000000ff)
+         ^ ks[43];
+
+  out[0] = hc_swap32_S (out[0]);
+  out[1] = hc_swap32_S (out[1]);
+  out[2] = hc_swap32_S (out[2]);
+  out[3] = hc_swap32_S (out[3]);
+}
+
+inline void aes128_decrypt (const uint *ks, const uint *in, uint *out)
+{
+
+  uint4 s = (uint4)(hc_swap32_S(in[0]), hc_swap32_S(in[1]), hc_swap32_S(in[2]), hc_swap32_S(in[3]));
+  s ^= vload4(0, ks);
+
+  uint4 t;
+
+  t.s0 = td0[s.s0 >> 24] ^ td1[(s.s3 >> 16) & 0xff] ^ td2[(s.s2 >>  8) & 0xff] ^ td3[s.s1 & 0xff] ^ ks[ 4];
+  t.s1 = td0[s.s1 >> 24] ^ td1[(s.s0 >> 16) & 0xff] ^ td2[(s.s3 >>  8) & 0xff] ^ td3[s.s2 & 0xff] ^ ks[ 5];
+  t.s2 = td0[s.s2 >> 24] ^ td1[(s.s1 >> 16) & 0xff] ^ td2[(s.s0 >>  8) & 0xff] ^ td3[s.s3 & 0xff] ^ ks[ 6];
+  t.s3 = td0[s.s3 >> 24] ^ td1[(s.s2 >> 16) & 0xff] ^ td2[(s.s1 >>  8) & 0xff] ^ td3[s.s0& 0xff] ^ ks[ 7];
+  s.s0= td0[t.s0 >> 24] ^ td1[(t.s3 >> 16) & 0xff] ^ td2[(t.s2 >>  8) & 0xff] ^ td3[t.s1 & 0xff] ^ ks[ 8];
+  s.s1 = td0[t.s1 >> 24] ^ td1[(t.s0 >> 16) & 0xff] ^ td2[(t.s3 >>  8) & 0xff] ^ td3[t.s2 & 0xff] ^ ks[ 9];
+  s.s2 = td0[t.s2 >> 24] ^ td1[(t.s1 >> 16) & 0xff] ^ td2[(t.s0 >>  8) & 0xff] ^ td3[t.s3 & 0xff] ^ ks[10];
+  s.s3 = td0[t.s3 >> 24] ^ td1[(t.s2 >> 16) & 0xff] ^ td2[(t.s1 >>  8) & 0xff] ^ td3[t.s0 & 0xff] ^ ks[11];
+  t.s0 = td0[s.s0 >> 24] ^ td1[(s.s3 >> 16) & 0xff] ^ td2[(s.s2 >>  8) & 0xff] ^ td3[s.s1 & 0xff] ^ ks[12];
+  t.s1 = td0[s.s1 >> 24] ^ td1[(s.s0 >> 16) & 0xff] ^ td2[(s.s3 >>  8) & 0xff] ^ td3[s.s2 & 0xff] ^ ks[13];
+  t.s2 = td0[s.s2 >> 24] ^ td1[(s.s1 >> 16) & 0xff] ^ td2[(s.s0 >>  8) & 0xff] ^ td3[s.s3 & 0xff] ^ ks[14];
+  t.s3 = td0[s.s3 >> 24] ^ td1[(s.s2 >> 16) & 0xff] ^ td2[(s.s1 >>  8) & 0xff] ^ td3[s.s0& 0xff] ^ ks[15];
+  s.s0= td0[t.s0 >> 24] ^ td1[(t.s3 >> 16) & 0xff] ^ td2[(t.s2 >>  8) & 0xff] ^ td3[t.s1 & 0xff] ^ ks[16];
+  s.s1 = td0[t.s1 >> 24] ^ td1[(t.s0 >> 16) & 0xff] ^ td2[(t.s3 >>  8) & 0xff] ^ td3[t.s2 & 0xff] ^ ks[17];
+  s.s2 = td0[t.s2 >> 24] ^ td1[(t.s1 >> 16) & 0xff] ^ td2[(t.s0 >>  8) & 0xff] ^ td3[t.s3 & 0xff] ^ ks[18];
+  s.s3 = td0[t.s3 >> 24] ^ td1[(t.s2 >> 16) & 0xff] ^ td2[(t.s1 >>  8) & 0xff] ^ td3[t.s0 & 0xff] ^ ks[19];
+  t.s0 = td0[s.s0 >> 24] ^ td1[(s.s3 >> 16) & 0xff] ^ td2[(s.s2 >>  8) & 0xff] ^ td3[s.s1 & 0xff] ^ ks[20];
+  t.s1 = td0[s.s1 >> 24] ^ td1[(s.s0 >> 16) & 0xff] ^ td2[(s.s3 >>  8) & 0xff] ^ td3[s.s2 & 0xff] ^ ks[21];
+  t.s2 = td0[s.s2 >> 24] ^ td1[(s.s1 >> 16) & 0xff] ^ td2[(s.s0 >>  8) & 0xff] ^ td3[s.s3 & 0xff] ^ ks[22];
+  t.s3 = td0[s.s3 >> 24] ^ td1[(s.s2 >> 16) & 0xff] ^ td2[(s.s1 >>  8) & 0xff] ^ td3[s.s0& 0xff] ^ ks[23];
+  s.s0= td0[t.s0 >> 24] ^ td1[(t.s3 >> 16) & 0xff] ^ td2[(t.s2 >>  8) & 0xff] ^ td3[t.s1 & 0xff] ^ ks[24];
+  s.s1 = td0[t.s1 >> 24] ^ td1[(t.s0 >> 16) & 0xff] ^ td2[(t.s3 >>  8) & 0xff] ^ td3[t.s2 & 0xff] ^ ks[25];
+  s.s2 = td0[t.s2 >> 24] ^ td1[(t.s1 >> 16) & 0xff] ^ td2[(t.s0 >>  8) & 0xff] ^ td3[t.s3 & 0xff] ^ ks[26];
+  s.s3 = td0[t.s3 >> 24] ^ td1[(t.s2 >> 16) & 0xff] ^ td2[(t.s1 >>  8) & 0xff] ^ td3[t.s0 & 0xff] ^ ks[27];
+  t.s0 = td0[s.s0 >> 24] ^ td1[(s.s3 >> 16) & 0xff] ^ td2[(s.s2 >>  8) & 0xff] ^ td3[s.s1 & 0xff] ^ ks[28];
+  t.s1 = td0[s.s1 >> 24] ^ td1[(s.s0 >> 16) & 0xff] ^ td2[(s.s3 >>  8) & 0xff] ^ td3[s.s2 & 0xff] ^ ks[29];
+  t.s2 = td0[s.s2 >> 24] ^ td1[(s.s1 >> 16) & 0xff] ^ td2[(s.s0 >>  8) & 0xff] ^ td3[s.s3 & 0xff] ^ ks[30];
+  t.s3 = td0[s.s3 >> 24] ^ td1[(s.s2 >> 16) & 0xff] ^ td2[(s.s1 >>  8) & 0xff] ^ td3[s.s0& 0xff] ^ ks[31];
+  s.s0= td0[t.s0 >> 24] ^ td1[(t.s3 >> 16) & 0xff] ^ td2[(t.s2 >>  8) & 0xff] ^ td3[t.s1 & 0xff] ^ ks[32];
+  s.s1 = td0[t.s1 >> 24] ^ td1[(t.s0 >> 16) & 0xff] ^ td2[(t.s3 >>  8) & 0xff] ^ td3[t.s2 & 0xff] ^ ks[33];
+  s.s2 = td0[t.s2 >> 24] ^ td1[(t.s1 >> 16) & 0xff] ^ td2[(t.s0 >>  8) & 0xff] ^ td3[t.s3 & 0xff] ^ ks[34];
+  s.s3 = td0[t.s3 >> 24] ^ td1[(t.s2 >> 16) & 0xff] ^ td2[(t.s1 >>  8) & 0xff] ^ td3[t.s0 & 0xff] ^ ks[35];
+  t.s0 = td0[s.s0 >> 24] ^ td1[(s.s3 >> 16) & 0xff] ^ td2[(s.s2 >>  8) & 0xff] ^ td3[s.s1 & 0xff] ^ ks[36];
+  t.s1 = td0[s.s1 >> 24] ^ td1[(s.s0 >> 16) & 0xff] ^ td2[(s.s3 >>  8) & 0xff] ^ td3[s.s2 & 0xff] ^ ks[37];
+  t.s2 = td0[s.s2 >> 24] ^ td1[(s.s1 >> 16) & 0xff] ^ td2[(s.s0 >>  8) & 0xff] ^ td3[s.s3 & 0xff] ^ ks[38];
+  t.s3 = td0[s.s3 >> 24] ^ td1[(s.s2 >> 16) & 0xff] ^ td2[(s.s1 >>  8) & 0xff] ^ td3[s.s0& 0xff] ^ ks[39];
+
+  out[0] = (td4[(t.s0 >> 24) & 0xff] & 0xff000000)
+         ^ (td4[(t.s3 >> 16) & 0xff] & 0x00ff0000)
+         ^ (td4[(t.s2 >>  8) & 0xff] & 0x0000ff00)
+         ^ (td4[(t.s1 >>  0) & 0xff] & 0x000000ff)
+         ^ ks[40];
+
+  out[1] = (td4[(t.s1 >> 24) & 0xff] & 0xff000000)
+         ^ (td4[(t.s0 >> 16) & 0xff] & 0x00ff0000)
+         ^ (td4[(t.s3 >>  8) & 0xff] & 0x0000ff00)
+         ^ (td4[(t.s2 >>  0) & 0xff] & 0x000000ff)
+         ^ ks[41];
+
+  out[2] = (td4[(t.s2 >> 24) & 0xff] & 0xff000000)
+         ^ (td4[(t.s1 >> 16) & 0xff] & 0x00ff0000)
+         ^ (td4[(t.s0 >>  8) & 0xff] & 0x0000ff00)
+         ^ (td4[(t.s3 >>  0) & 0xff] & 0x000000ff)
+         ^ ks[42];
+
+  out[3] = (td4[(t.s3 >> 24) & 0xff] & 0xff000000)
+         ^ (td4[(t.s2 >> 16) & 0xff] & 0x00ff0000)
+         ^ (td4[(t.s1 >>  8) & 0xff] & 0x0000ff00)
+         ^ (td4[(t.s0 >>  0) & 0xff] & 0x000000ff)
+         ^ ks[43];
+
+  out[0] = hc_swap32_S (out[0]);
+  out[1] = hc_swap32_S (out[1]);
+  out[2] = hc_swap32_S (out[2]);
+  out[3] = hc_swap32_S (out[3]);
+}
+
+// Galois field multiplication
+inline void xts_mul2 (uint *in, uint *out)
+{
+  const uint c = in[3] >> 31;
+
+  out[3] = (in[3] << 1) | (in[2] >> 31);
+  out[2] = (in[2] << 1) | (in[1] >> 31);
+  out[1] = (in[1] << 1) | (in[0] >> 31);
+  out[0] = (in[0] << 1);
+
+  out[0] ^= c * 0x87;
+}
+
+
+// ks - expanded key
+// sec_n - sector number (uint[4])
+// block_n - encryption block number
+
+inline void aes_xts256_gen_tweak (const uint *ks, const uint *sec_n, const uint block_n, uint *out)
+{
+  aes128_encrypt (ks, sec_n, out);
+  //Double the tweak value block_n-1 times (since first block is 0)
+  for (uint i = 0; i<block_n; i++) xts_mul2(out, out);
+}
+
+// ks - expanded key (uint[44])
+// T - tweak (uint[4])
+inline void aes_xts256_enc_block (const uint *ks, const uint *T, const uint *in, uint *out)
+{
+  vstore4(*(uint4*)in^*(uint4*)T, 0, out);
+  aes128_encrypt (ks, out, out);
+  vstore4(*(uint4*)out^*(uint4*)T, 0, out);
+}
+
+inline void aes_xts256_dec_block (const uint *ks, const uint *T, const uint *in, uint *out)
+{
+  vstore4(*(uint4*)in^*(uint4*)T, 0, out);
+  aes128_decrypt (ks, out, out);
+  vstore4(*(uint4*)out^*(uint4*)T, 0, out);
+}
+
diff --git a/src/aes256_xts_plain.h b/src/aes256_xts_plain.h
new file mode 100644
index 0000000..302ab1a
--- /dev/null
+++ b/src/aes256_xts_plain.h
@@ -0,0 +1,29 @@
+// aes-xts-pur64 is OpenCL code for aes-xts256-plain64 encryption compatible with LUKS
+//
+//     Copyright (C) 2025  Kirill Shakirov
+//
+//     This program is free software: you can redistribute it and/or modify
+//     it under the terms of the GNU General Public License as published by
+//     the Free Software Foundation, either version 3 of the License, or
+//     (at your option) any later version.
+//
+//     This program is distributed in the hope that it will be useful,
+//     but WITHOUT ANY WARRANTY; without even the implied warranty of
+//     MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+//     GNU General Public License for more details.
+//
+//     You should have received a copy of the GNU General Public License
+//     along with this program.  If not, see <https://www.gnu.org/licenses/>.
+
+
+inline void aes128_InvertKey (uint *ks);
+inline void aes128_ExpandKey (uint *ks, const uint *ukey);
+inline void aes128_set_encrypt_key (uint *ks, const uint *ukey);
+inline void aes128_set_decrypt_key (uint *ks, const uint *ukey);
+inline void aes128_encrypt (const uint *ks, const uint *in, uint *out);
+inline void aes128_decrypt (const uint *ks, const uint *in, uint *out);
+
+inline void xts_mul2 (uint *in, uint *out);
+inline void aes_xts256_gen_tweak (const uint *ks, const uint *sec_n, const uint block_n, uint *out);
+inline void aes_xts256_enc_block (const uint *ks, const uint *T, const uint *in, uint *out);
+inline void aes_xts256_dec_block (const uint *ks, const uint *T, const uint *in, uint *out);
diff --git a/src/test_aes_xts256_plain.cl b/src/test_aes_xts256_plain.cl
new file mode 100644
index 0000000..9e95001
--- /dev/null
+++ b/src/test_aes_xts256_plain.cl
@@ -0,0 +1,98 @@
+// aes-xts-pur64 is OpenCL code for aes-xts256-plain64 encryption compatible with LUKS
+//
+//     Copyright (C) 2025  Kirill Shakirov
+//
+//     This program is free software: you can redistribute it and/or modify
+//     it under the terms of the GNU General Public License as published by
+//     the Free Software Foundation, either version 3 of the License, or
+//     (at your option) any later version.
+//
+//     This program is distributed in the hope that it will be useful,
+//     but WITHOUT ANY WARRANTY; without even the implied warranty of
+//     MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+//     GNU General Public License for more details.
+//
+//     You should have received a copy of the GNU General Public License
+//     along with this program.  If not, see <https://www.gnu.org/licenses/>.
+
+
+#include "aes256_xts_plain.cl"
+
+__kernel void encrypt_data(__global const ulong* g_Ti, __global const uint* g_Tj, 
+                           __global const uint8* g_key,
+                           __global const uint4* g_u_data,
+                           __global uint* g_enc_data)
+{
+  const size_t g_id = get_global_id(0);
+  uint d_ks[44];
+  uint t_ks[44];
+  uint tweak[4];
+  uint enc_key[8];
+  uint u_data[4];
+  uint enc_data[4] = { 0 };
+
+  uint sec_n[4];
+  ulong Ti = g_Ti[g_id];
+  sec_n[0] = ((uint*)&Ti)[0];
+  sec_n[1] = ((uint*)&Ti)[1];
+  sec_n[2] = 0;
+  sec_n[3] = 0;
+
+  uint Tj = g_Tj[g_id];
+
+  vstore8(*g_key, 0, enc_key);
+  vstore4(g_u_data[g_id], 0, u_data);
+
+  // printf("Ti: %lu\\n", Ti);
+  // printf("Tj: %u\\n", Tj);
+  // printf("enc_key: %v8u\\n", *(uint8*)enc_key);
+  // printf("uenc_data: %v4u\\n", *(uint4*)uenc_data);
+
+  //calculate tweak value
+  aes128_set_encrypt_key (t_ks, &enc_key[4]);
+  aes_xts256_gen_tweak (t_ks, sec_n, Tj, tweak);
+
+  // encrypt data
+  aes128_set_encrypt_key (d_ks, enc_key);
+  aes_xts256_enc_block (d_ks, tweak, u_data, enc_data);
+  // printf("enc_data: %v4u\\n", *(uint4*)enc_data);
+  vstore4(*(uint4*)enc_data, g_id, g_enc_data);
+}
+
+
+__kernel void decrypt_data(__global const ulong* g_Ti, __global const uint* g_Tj, 
+                           __global const uint8* g_key,
+                           __global const uint4* g_enc_data,
+                           __global uint* g_u_data)
+{
+  const size_t g_id = get_global_id(0);
+  uint d_ks[44];
+  uint t_ks[44];
+  uint tweak[4];
+  uint enc_key[8];
+  uint enc_data[4];
+  uint u_data[4] = { 0 };
+
+  uint sec_n[4];
+  ulong Ti = g_Ti[g_id];
+  sec_n[0] = ((uint*)&Ti)[0];
+  sec_n[1] = ((uint*)&Ti)[1];
+  sec_n[2] = 0;
+  sec_n[3] = 0;
+
+  uint Tj = g_Tj[g_id];
+
+  vstore8(*g_key, 0, enc_key);
+  vstore4(g_enc_data[g_id], 0, enc_data);
+
+
+  //calculate tweak value
+  aes128_set_encrypt_key (t_ks, &enc_key[4]);
+  aes_xts256_gen_tweak (t_ks, sec_n, Tj, tweak);
+
+  // decrypt data
+  aes128_set_decrypt_key (d_ks, enc_key);
+  aes_xts256_dec_block (d_ks, tweak, enc_data, u_data);
+  // printf("enc_data: %v4u\\n", *(uint4*)enc_data);
+  vstore4(*(uint4*)u_data, g_id, g_u_data);
+}
diff --git a/tests/data/test_vectors.json b/tests/data/test_vectors.json
new file mode 100644
index 0000000..1859bd7
--- /dev/null
+++ b/tests/data/test_vectors.json
@@ -0,0 +1,1541 @@
+{
+  "encryption_key": "8adb7b7e8a722df091ecea988a4ad223 4836636a102ceb688b3985f89bf40002",
+  "vectors": [
+    {
+      "sector_number": 0,
+      "block_number": 8,
+      "unencrypted_data": "d3ea0be50eca14b514b12ab7d472edf6",
+      "encrypted_data": "95836ed018322d24ba10fd477676ab23"
+    },
+    {
+      "sector_number": 1,
+      "block_number": 29,
+      "unencrypted_data": "014b6c2c3eddb765a0ced561591e4470",
+      "encrypted_data": "d24e9a021a310f3c783c72ea103e2c44"
+    },
+    {
+      "sector_number": 2,
+      "block_number": 13,
+      "unencrypted_data": "4fbaf9d9a68900ab6c61c33b1c20f8a3",
+      "encrypted_data": "4b593062aaff46d64f9e0c94cb9dd087"
+    },
+    {
+      "sector_number": 3,
+      "block_number": 6,
+      "unencrypted_data": "7a56a4c24df9139a14a1155343142682",
+      "encrypted_data": "0030b6826bebaa26c6f64daf23a9decb"
+    },
+    {
+      "sector_number": 4,
+      "block_number": 7,
+      "unencrypted_data": "4e726ffc7974e1a70fac70e41dd4a233",
+      "encrypted_data": "b0b70a3fcbc6cba961c30f6877c2b21d"
+    },
+    {
+      "sector_number": 5,
+      "block_number": 9,
+      "unencrypted_data": "a861716b3a356e2af7b24c203c6207e8",
+      "encrypted_data": "e7d69284eebad1efcbd691e66da78c80"
+    },
+    {
+      "sector_number": 6,
+      "block_number": 19,
+      "unencrypted_data": "05ec11b2131aa32dc6a0a8328d3b6fb5",
+      "encrypted_data": "015b67994c5477b05a9744e306fa79e2"
+    },
+    {
+      "sector_number": 7,
+      "block_number": 4,
+      "unencrypted_data": "55afd34c70a65b25051c0060e33bad80",
+      "encrypted_data": "3674508687ce2c21e0314b1406faa338"
+    },
+    {
+      "sector_number": 8,
+      "block_number": 26,
+      "unencrypted_data": "087a551a8bc6317210c074c6ec2d0854",
+      "encrypted_data": "013dd186def79b4eba1a1baa40f18a62"
+    },
+    {
+      "sector_number": 9,
+      "block_number": 15,
+      "unencrypted_data": "73846193ce3e94fb29ffcd90f3a52768",
+      "encrypted_data": "4380feef9431ee21a97eebed2da3eba6"
+    },
+    {
+      "sector_number": 10,
+      "block_number": 31,
+      "unencrypted_data": "fea11ac940dc334f76fecdc5b836fa61",
+      "encrypted_data": "bdbdc35ec31ded539398f7373b27de51"
+    },
+    {
+      "sector_number": 11,
+      "block_number": 3,
+      "unencrypted_data": "58c183144092e08e1dcc4a306333976b",
+      "encrypted_data": "857e264c91279c19db480322bc92dd58"
+    },
+    {
+      "sector_number": 12,
+      "block_number": 21,
+      "unencrypted_data": "426acffc40bc5a8a3a5e6d8d28f23617",
+      "encrypted_data": "e6ed4f70f5302cb5a064d3cd3e6b18db"
+    },
+    {
+      "sector_number": 13,
+      "block_number": 25,
+      "unencrypted_data": "84e2f6c532756a424acd203d3a2f9296",
+      "encrypted_data": "6b0eb7e0ae220b4a8143831fa633a8c7"
+    },
+    {
+      "sector_number": 14,
+      "block_number": 28,
+      "unencrypted_data": "fde028b1671a5e305aa61351e8470a27",
+      "encrypted_data": "364be9c8e726a4e0798c5375648890f8"
+    },
+    {
+      "sector_number": 15,
+      "block_number": 17,
+      "unencrypted_data": "9174a6822fb2e99d5b4e9ed0baa15fa3",
+      "encrypted_data": "761fda433102677be5590c410b9fbe04"
+    },
+    {
+      "sector_number": 16,
+      "block_number": 29,
+      "unencrypted_data": "fd8f08fdfc1e24fd4a58218c60858951",
+      "encrypted_data": "28ae7835e1b5bf4204b2b1ae4c48111d"
+    },
+    {
+      "sector_number": 17,
+      "block_number": 4,
+      "unencrypted_data": "357d51963b96ff473045b2245071cd34",
+      "encrypted_data": "5b84ac6e59179599b4ffae8dec65aee0"
+    },
+    {
+      "sector_number": 18,
+      "block_number": 20,
+      "unencrypted_data": "44d62ddafccee00f254ab95eeeb40841",
+      "encrypted_data": "c49755aa441df623ca59fab46dd159e5"
+    },
+    {
+      "sector_number": 19,
+      "block_number": 6,
+      "unencrypted_data": "243890d823906bd0d49da90900ee5a36",
+      "encrypted_data": "3255a45a66466e9616d4c493b94518be"
+    },
+    {
+      "sector_number": 20,
+      "block_number": 10,
+      "unencrypted_data": "a82f0547065c0889ba3ca4488f58e8c5",
+      "encrypted_data": "c95fb0f9f8ec920290e0ffe02fa84408"
+    },
+    {
+      "sector_number": 21,
+      "block_number": 16,
+      "unencrypted_data": "8e04f39c627acb95ff35d248ecda38bf",
+      "encrypted_data": "7592ea222a4a23a01c6dbb04e5906ea3"
+    },
+    {
+      "sector_number": 22,
+      "block_number": 27,
+      "unencrypted_data": "5c07c5d9ef211dd32ed801d0c9b8945f",
+      "encrypted_data": "96009d57320bd709d299bf85c93e53b7"
+    },
+    {
+      "sector_number": 23,
+      "block_number": 31,
+      "unencrypted_data": "e58573bf531cde0f2dc25a2826578c36",
+      "encrypted_data": "0de156c807a87b0477ca7ebff0a4a28c"
+    },
+    {
+      "sector_number": 24,
+      "block_number": 29,
+      "unencrypted_data": "bdae5bb8d7a982c29bb31773ad352d31",
+      "encrypted_data": "e293db066238b93e03972be0f373c1d8"
+    },
+    {
+      "sector_number": 25,
+      "block_number": 26,
+      "unencrypted_data": "354dedd18c9ac9f90456674561a8f67f",
+      "encrypted_data": "88a64c5939609e79c09a338bf0f761f1"
+    },
+    {
+      "sector_number": 26,
+      "block_number": 2,
+      "unencrypted_data": "d58b400d447b5ef8294f709239c8c84d",
+      "encrypted_data": "f2a04ee47c44c369f979d79d0ebc992e"
+    },
+    {
+      "sector_number": 27,
+      "block_number": 19,
+      "unencrypted_data": "392a4baf4f5eb07617c8c187e6329c38",
+      "encrypted_data": "e5771eaaa0e872f40260bb7e0385c91a"
+    },
+    {
+      "sector_number": 28,
+      "block_number": 7,
+      "unencrypted_data": "25c6a688b83f95e81d01f78c39470198",
+      "encrypted_data": "41a15b68e5e14121f28d9e93e4a32ef7"
+    },
+    {
+      "sector_number": 29,
+      "block_number": 1,
+      "unencrypted_data": "4e51f3047601e0528e7de2b60f21c997",
+      "encrypted_data": "7a85e8033d4bf53947544edebbfec0f9"
+    },
+    {
+      "sector_number": 30,
+      "block_number": 18,
+      "unencrypted_data": "f76cfa104ed473f9275405865a6f3664",
+      "encrypted_data": "edd3aea7aeed3bed3833376d34731474"
+    },
+    {
+      "sector_number": 31,
+      "block_number": 6,
+      "unencrypted_data": "41185e90f84dfd9cc29f9d5b08c7979e",
+      "encrypted_data": "158e76cd9c149bcc6ddd3c7a4f4c4288"
+    },
+    {
+      "sector_number": 32,
+      "block_number": 24,
+      "unencrypted_data": "da9660145c6039ea2442dd153e279e49",
+      "encrypted_data": "54145e84bcb80d8bc842a692adaac808"
+    },
+    {
+      "sector_number": 33,
+      "block_number": 8,
+      "unencrypted_data": "93d1c1ced542e321a7da32b97e2c0e00",
+      "encrypted_data": "36df2e63321248cdd67cbb4bcd0dc2f4"
+    },
+    {
+      "sector_number": 34,
+      "block_number": 24,
+      "unencrypted_data": "27c165b852c7c25368a87512c7005b4e",
+      "encrypted_data": "e1c9cf4b2bffc69bf52ebd57d66dc582"
+    },
+    {
+      "sector_number": 35,
+      "block_number": 24,
+      "unencrypted_data": "5109d6983d4d1cac217d093faa978db9",
+      "encrypted_data": "f6c3b9597714a17165b9844f0c2a3d47"
+    },
+    {
+      "sector_number": 36,
+      "block_number": 30,
+      "unencrypted_data": "94500966482e4afbe0e7a8d9469438a3",
+      "encrypted_data": "6ec009328795390f0aedcff8cfc3e9fa"
+    },
+    {
+      "sector_number": 37,
+      "block_number": 30,
+      "unencrypted_data": "b90006f89d78ef4b96daf92750b014f9",
+      "encrypted_data": "444419b43d4e666143399a95c7608f17"
+    },
+    {
+      "sector_number": 38,
+      "block_number": 10,
+      "unencrypted_data": "a941547d3ea0c76c4964d3f0ddd5c9bd",
+      "encrypted_data": "d846ac36f25f59ff13b54297535d6d85"
+    },
+    {
+      "sector_number": 39,
+      "block_number": 20,
+      "unencrypted_data": "6ea4ef458b9725b2d7d104555e54d48a",
+      "encrypted_data": "eac2c7c27f6afcdf4afd463240670d8b"
+    },
+    {
+      "sector_number": 40,
+      "block_number": 22,
+      "unencrypted_data": "ee326c147f75a237b923dee239c63453",
+      "encrypted_data": "cf412f898710dd01880b956f2c753755"
+    },
+    {
+      "sector_number": 41,
+      "block_number": 22,
+      "unencrypted_data": "fed892111cc39ddec868af67869c87fe",
+      "encrypted_data": "29ce4b70a372ed9c62b85b6c21cd4324"
+    },
+    {
+      "sector_number": 42,
+      "block_number": 25,
+      "unencrypted_data": "f4237b350633f69d59376ed271b7acd8",
+      "encrypted_data": "4770cd915cfef0756d8de58a83bea2f7"
+    },
+    {
+      "sector_number": 43,
+      "block_number": 24,
+      "unencrypted_data": "2120e1fe0f479575ec6c6a29f558d9c2",
+      "encrypted_data": "dd580ce74310e94274329f1cd5bafd48"
+    },
+    {
+      "sector_number": 44,
+      "block_number": 26,
+      "unencrypted_data": "332eaf008878d51f2c38c7458cb41f1a",
+      "encrypted_data": "527878193fbf34dbd3b8c639401619d3"
+    },
+    {
+      "sector_number": 45,
+      "block_number": 26,
+      "unencrypted_data": "a1123dfe6c9b2076ddb73e83f6155874",
+      "encrypted_data": "bd1beabccacc4e132c2b0592d5035cd9"
+    },
+    {
+      "sector_number": 46,
+      "block_number": 2,
+      "unencrypted_data": "073a0b5ef21700fe03a3cf33dfedf4ad",
+      "encrypted_data": "d3937742e28fcfc96edc78608acb0cb2"
+    },
+    {
+      "sector_number": 47,
+      "block_number": 18,
+      "unencrypted_data": "2b9655647befe8c6c799439a08b56d78",
+      "encrypted_data": "8023ebcf4565eeeb637ca2a3dd953d2b"
+    },
+    {
+      "sector_number": 48,
+      "block_number": 26,
+      "unencrypted_data": "f510db005dc34c01b0cba128385a05e8",
+      "encrypted_data": "127faff97818d652ab8d378ca3a84b9c"
+    },
+    {
+      "sector_number": 49,
+      "block_number": 25,
+      "unencrypted_data": "6f90fcc6a307ce545ed7983d9cfee06b",
+      "encrypted_data": "90b03589059d8c094ea879c4a24ba702"
+    },
+    {
+      "sector_number": 50,
+      "block_number": 26,
+      "unencrypted_data": "e52dcb678bf55d037a971fd420d82ae2",
+      "encrypted_data": "d924baa7f2cb1490d2c3c84419ac345a"
+    },
+    {
+      "sector_number": 51,
+      "block_number": 20,
+      "unencrypted_data": "3cc01043d95202475b6911f7fbce02b9",
+      "encrypted_data": "d25a4dffbc3f9916c48033f7a1951729"
+    },
+    {
+      "sector_number": 52,
+      "block_number": 11,
+      "unencrypted_data": "cb070f3d37efbfdb8c8fdd30f38b0924",
+      "encrypted_data": "b71a7c1461a2484bee148f124f01e314"
+    },
+    {
+      "sector_number": 53,
+      "block_number": 7,
+      "unencrypted_data": "8468085691a55a1441470d10f38dded3",
+      "encrypted_data": "b85103d61460ea70c6bf39502324d210"
+    },
+    {
+      "sector_number": 54,
+      "block_number": 26,
+      "unencrypted_data": "33af89a86e6a8669304f98a47438006a",
+      "encrypted_data": "1e221625f3682a41b3524196fb1a7941"
+    },
+    {
+      "sector_number": 55,
+      "block_number": 14,
+      "unencrypted_data": "69d26b473026de8b1f6b28d28744a6fc",
+      "encrypted_data": "fa5167e40b2dd380da6410ade9fc57b5"
+    },
+    {
+      "sector_number": 56,
+      "block_number": 11,
+      "unencrypted_data": "f2c7f75c77f0630108e9d0f4ed16dc60",
+      "encrypted_data": "ab80f6a51154f19b55bb4f001af3df8a"
+    },
+    {
+      "sector_number": 57,
+      "block_number": 17,
+      "unencrypted_data": "997000d6c53b07bfac551a4529bbd752",
+      "encrypted_data": "adf1123b14aca6ea04b596175bf1cb4e"
+    },
+    {
+      "sector_number": 58,
+      "block_number": 20,
+      "unencrypted_data": "971855efc177b9461a845f372866fcfe",
+      "encrypted_data": "934993ffca16c7337a5ef67331770d0e"
+    },
+    {
+      "sector_number": 59,
+      "block_number": 12,
+      "unencrypted_data": "4a7c325294ae6f6f7741cdd36783650a",
+      "encrypted_data": "42e9f1339612b1a766efe89520cd1256"
+    },
+    {
+      "sector_number": 60,
+      "block_number": 15,
+      "unencrypted_data": "3df635c49ec8169c75b61b83db35062f",
+      "encrypted_data": "aac03737598a1ee1d015aa6f413b3125"
+    },
+    {
+      "sector_number": 61,
+      "block_number": 7,
+      "unencrypted_data": "edf01b6380c145d2b3b121e64cc42b91",
+      "encrypted_data": "7e36b10bd5b5e2c77695a11b5d6c8024"
+    },
+    {
+      "sector_number": 62,
+      "block_number": 4,
+      "unencrypted_data": "173640469d98c2f6681a83342cd729a7",
+      "encrypted_data": "e8921deb23e44b3a93a520a18c58e358"
+    },
+    {
+      "sector_number": 63,
+      "block_number": 11,
+      "unencrypted_data": "a4d13cf230550c72837ae0cf2db41098",
+      "encrypted_data": "312b77dd3e8743129620eecffb55db56"
+    },
+    {
+      "sector_number": 64,
+      "block_number": 23,
+      "unencrypted_data": "b52d5ff4f6bed2b7d9b1e534984a62c6",
+      "encrypted_data": "53e69c93d11f649fb39cb997e0d45107"
+    },
+    {
+      "sector_number": 65,
+      "block_number": 9,
+      "unencrypted_data": "847ea45f2acf3700da90f6226c12535a",
+      "encrypted_data": "ffb611f0d990f4619a23358fe6de4ab4"
+    },
+    {
+      "sector_number": 66,
+      "block_number": 27,
+      "unencrypted_data": "e223e028b0f8c536d26c5bde1528da53",
+      "encrypted_data": "ed7aa1337a16f0a8f484f83800677dd1"
+    },
+    {
+      "sector_number": 67,
+      "block_number": 10,
+      "unencrypted_data": "af54fce3ed030850df177916fe75f243",
+      "encrypted_data": "1ff1ee1cc230932d7a4a8b0d4b5d4734"
+    },
+    {
+      "sector_number": 68,
+      "block_number": 31,
+      "unencrypted_data": "0bd5291626a63c58a62585ceef33c192",
+      "encrypted_data": "0f66c173e8680bce2453a8d83e070cfa"
+    },
+    {
+      "sector_number": 69,
+      "block_number": 2,
+      "unencrypted_data": "88bf199c9bcd5971d69ad17700997540",
+      "encrypted_data": "998b31606a4d2d7ea6a1e59cd383f6c3"
+    },
+    {
+      "sector_number": 70,
+      "block_number": 8,
+      "unencrypted_data": "7fe3182d5310638bd5b9bbb85eaf5e3b",
+      "encrypted_data": "6ba55063e1abdb61e511e997474a3760"
+    },
+    {
+      "sector_number": 71,
+      "block_number": 21,
+      "unencrypted_data": "4d7683bf7258c36ae2df3abdc8aecb20",
+      "encrypted_data": "651d6ba2a7940d9f0be0053a41cdd9ac"
+    },
+    {
+      "sector_number": 72,
+      "block_number": 28,
+      "unencrypted_data": "551c8e4e06d02540e2b105ec6ce2063f",
+      "encrypted_data": "d28339146f2afef14f31445375c46f92"
+    },
+    {
+      "sector_number": 73,
+      "block_number": 10,
+      "unencrypted_data": "a7e2b62634e4184083b0b8253db5c404",
+      "encrypted_data": "8f5e090ccd002bdbbc8e8522e433a6d1"
+    },
+    {
+      "sector_number": 74,
+      "block_number": 19,
+      "unencrypted_data": "6a8410487ef17aa1edaf4987041e5680",
+      "encrypted_data": "3d395355cf859523e1c94fb79d060fb7"
+    },
+    {
+      "sector_number": 75,
+      "block_number": 24,
+      "unencrypted_data": "1604142d63da2eb54f6d94107cfe1e53",
+      "encrypted_data": "9f490cb4516af73b6be94a481040631e"
+    },
+    {
+      "sector_number": 76,
+      "block_number": 26,
+      "unencrypted_data": "645fad5d0c27f633a4023b88a22fc8ae",
+      "encrypted_data": "191aedd3f2e8c6a76229d3b1ffb50efe"
+    },
+    {
+      "sector_number": 77,
+      "block_number": 28,
+      "unencrypted_data": "844ff4ab226bda69b7abfd5144008491",
+      "encrypted_data": "1529fa6db99f61e780b83915fec68ba1"
+    },
+    {
+      "sector_number": 78,
+      "block_number": 1,
+      "unencrypted_data": "f5c60f95268b5ec1b50cd306eaa34eb9",
+      "encrypted_data": "20fba6b8a2c8906f49f5f98e19890341"
+    },
+    {
+      "sector_number": 79,
+      "block_number": 14,
+      "unencrypted_data": "62ea3cebc36246debbd3fb2de5266c0e",
+      "encrypted_data": "8a9d38e30ed97e62b02a8617739eeeb4"
+    },
+    {
+      "sector_number": 80,
+      "block_number": 26,
+      "unencrypted_data": "c6be1978b2d24e870c213107d831631f",
+      "encrypted_data": "43b437e3ba9cf97d90baf3856a298d62"
+    },
+    {
+      "sector_number": 81,
+      "block_number": 2,
+      "unencrypted_data": "66ca89b5e879e9e273311300c7b39470",
+      "encrypted_data": "1a992059339a87ddeecb358b7733a73f"
+    },
+    {
+      "sector_number": 82,
+      "block_number": 26,
+      "unencrypted_data": "3fd863e9691f0c1ed749d45d513d5d12",
+      "encrypted_data": "0e36d5d6f6651ce5f214837165192e79"
+    },
+    {
+      "sector_number": 83,
+      "block_number": 17,
+      "unencrypted_data": "ea0597cf05074ccab6664bb5339bcd46",
+      "encrypted_data": "e4593c68da998e9ffeeb5256fb16b768"
+    },
+    {
+      "sector_number": 84,
+      "block_number": 29,
+      "unencrypted_data": "064fbc04a416c1a5148efcdeebd461a3",
+      "encrypted_data": "c306437a3d93789944a7171829c3eeee"
+    },
+    {
+      "sector_number": 85,
+      "block_number": 24,
+      "unencrypted_data": "e399ecfc71e8474fd0c52aeac96145e2",
+      "encrypted_data": "7c33715fe8d5dcaa09534beddbaa8a37"
+    },
+    {
+      "sector_number": 86,
+      "block_number": 9,
+      "unencrypted_data": "8ff0896d4cb63c07eab44b325e894ca5",
+      "encrypted_data": "17f395646fefa4f6e3202498b0005bb6"
+    },
+    {
+      "sector_number": 87,
+      "block_number": 22,
+      "unencrypted_data": "e2c5a53f38a8d63f19ded21bede92c60",
+      "encrypted_data": "9062d3be5ac83c05c415f2ab06d1ea41"
+    },
+    {
+      "sector_number": 88,
+      "block_number": 31,
+      "unencrypted_data": "c58ba253ffc435332aa692f4b434f87a",
+      "encrypted_data": "1358894f85f845ed649766c8a36e7b8b"
+    },
+    {
+      "sector_number": 89,
+      "block_number": 15,
+      "unencrypted_data": "4e17d1706449521a8288f06fb15309fe",
+      "encrypted_data": "e174d4a9de8ab40e3129dec3744050a4"
+    },
+    {
+      "sector_number": 90,
+      "block_number": 16,
+      "unencrypted_data": "4e9d8a08e76aeb3b95ca668aa4f56913",
+      "encrypted_data": "dec7d51f72b27ed28becf31357a0a9b2"
+    },
+    {
+      "sector_number": 91,
+      "block_number": 28,
+      "unencrypted_data": "961b344b5b80eeb28c645ff9cdfa2c21",
+      "encrypted_data": "7fbb1e82e16ee8f124a5a17a247d9937"
+    },
+    {
+      "sector_number": 92,
+      "block_number": 28,
+      "unencrypted_data": "2039ed59972eec524b60c3ee44882bea",
+      "encrypted_data": "66ea9600b1d86feaa67d91627c8b462a"
+    },
+    {
+      "sector_number": 93,
+      "block_number": 18,
+      "unencrypted_data": "7eae86fdd9cc986d73d880e81676d655",
+      "encrypted_data": "65550fe08fa93bfce363e280a4964103"
+    },
+    {
+      "sector_number": 94,
+      "block_number": 22,
+      "unencrypted_data": "549963e64e383190b7adf9e35ebbc625",
+      "encrypted_data": "1d50e2b90828b23a5d8db4cab415cb3e"
+    },
+    {
+      "sector_number": 95,
+      "block_number": 17,
+      "unencrypted_data": "986cd60628a4353586d57c39e5949ae8",
+      "encrypted_data": "da0f3c88b3593f6f2b2ec4a1afe99793"
+    },
+    {
+      "sector_number": 96,
+      "block_number": 24,
+      "unencrypted_data": "ea8945819548d67c4f0d4c0aa9e3b702",
+      "encrypted_data": "754c0627656faefc2932bac92b47297e"
+    },
+    {
+      "sector_number": 97,
+      "block_number": 18,
+      "unencrypted_data": "b35bfa78af377fe2af86944025bc932a",
+      "encrypted_data": "81d894a1932cb7a0823944c5bac130a6"
+    },
+    {
+      "sector_number": 98,
+      "block_number": 2,
+      "unencrypted_data": "14a23eaaf2fd57dd5424b83298a562c1",
+      "encrypted_data": "17f5bf28940175d3cda3834e962af2ec"
+    },
+    {
+      "sector_number": 99,
+      "block_number": 28,
+      "unencrypted_data": "f05082b23351d871451fd104c022369b",
+      "encrypted_data": "1fadf6c9146f425e1a3aee61fb79a5e8"
+    },
+    {
+      "sector_number": 100,
+      "block_number": 12,
+      "unencrypted_data": "bdfe09c954fe07c741ace7034ffe4766",
+      "encrypted_data": "18fe61fa18511e23d0033876384e611b"
+    },
+    {
+      "sector_number": 101,
+      "block_number": 16,
+      "unencrypted_data": "4d6f19925608d7e091e6f8f22985f501",
+      "encrypted_data": "d8e3a29a679043569c3a4be0173ccc1f"
+    },
+    {
+      "sector_number": 102,
+      "block_number": 10,
+      "unencrypted_data": "b3c0069d0c6c3b4f3c4343cd22434f08",
+      "encrypted_data": "7234d956d6b22b755dab2746fa89b347"
+    },
+    {
+      "sector_number": 103,
+      "block_number": 18,
+      "unencrypted_data": "4a59964139721ae4adda5e076d31f4fe",
+      "encrypted_data": "209cdcce82edfc8579363c36d6f1d1b9"
+    },
+    {
+      "sector_number": 104,
+      "block_number": 24,
+      "unencrypted_data": "5af071775e1eb1f10e799ac767fab1c8",
+      "encrypted_data": "e9e7738a4f5c968ba61160e1e46603ad"
+    },
+    {
+      "sector_number": 105,
+      "block_number": 22,
+      "unencrypted_data": "984832fed82366062358be3c0e5ad207",
+      "encrypted_data": "e80cfa162ba8d7d5f4976aa94ab82bf2"
+    },
+    {
+      "sector_number": 106,
+      "block_number": 20,
+      "unencrypted_data": "0ddf4c03973147edccd4c4e67aab0f78",
+      "encrypted_data": "03cd627ff305b0733d724603e114a7ff"
+    },
+    {
+      "sector_number": 107,
+      "block_number": 3,
+      "unencrypted_data": "7838a06d19b3e834ba00f45483cc7a99",
+      "encrypted_data": "718c69154a8af11f03bfe3650695b045"
+    },
+    {
+      "sector_number": 108,
+      "block_number": 18,
+      "unencrypted_data": "f363b75dacdafeb50b626de94f7abe6c",
+      "encrypted_data": "d32716be5611c01a44b6b98d41417a21"
+    },
+    {
+      "sector_number": 109,
+      "block_number": 26,
+      "unencrypted_data": "60e0d19fbcb3dd506d3e5d02d5494ce1",
+      "encrypted_data": "b784d15a4e2a7c1469d88d554abfa9bb"
+    },
+    {
+      "sector_number": 110,
+      "block_number": 21,
+      "unencrypted_data": "49a985a6ab3675692ed41441c35978a3",
+      "encrypted_data": "8280aa45e4cf3031ec58b033184a2994"
+    },
+    {
+      "sector_number": 111,
+      "block_number": 22,
+      "unencrypted_data": "7a7f39cf974c4139626c50448edd919e",
+      "encrypted_data": "61aba957d0c1fc4ede161e1daebf11f9"
+    },
+    {
+      "sector_number": 112,
+      "block_number": 11,
+      "unencrypted_data": "21fb4c698906f3b07959c58d0556fe74",
+      "encrypted_data": "63e1296ef96e4856a4998269ed487cb3"
+    },
+    {
+      "sector_number": 113,
+      "block_number": 26,
+      "unencrypted_data": "f2b5efc39427af3d150fccdc9fd12cb8",
+      "encrypted_data": "2cfe5533f8eaf2869ec92cf90f289ab5"
+    },
+    {
+      "sector_number": 114,
+      "block_number": 10,
+      "unencrypted_data": "2ff4e4905ef5555bf3515f308119a31e",
+      "encrypted_data": "90039493f7fd11f2e947daf713ed0d66"
+    },
+    {
+      "sector_number": 115,
+      "block_number": 22,
+      "unencrypted_data": "9e839d866d6ddc990b03fb74091a1284",
+      "encrypted_data": "858c0c71c226ab754c0b65c43ddb29bc"
+    },
+    {
+      "sector_number": 116,
+      "block_number": 16,
+      "unencrypted_data": "ce2ab50d63d4d84ebb48c0104e3a5993",
+      "encrypted_data": "ac74081f0495e3d1f9ba7daeb6f34fab"
+    },
+    {
+      "sector_number": 117,
+      "block_number": 23,
+      "unencrypted_data": "7659113834255a8893bf8b1c48ff54dc",
+      "encrypted_data": "d4b1b2898842854a863a43cfcbefb0eb"
+    },
+    {
+      "sector_number": 118,
+      "block_number": 10,
+      "unencrypted_data": "fe166f004b89cb507035969bfe5ec897",
+      "encrypted_data": "a89a5412d590839e8d85ece4d6e24875"
+    },
+    {
+      "sector_number": 119,
+      "block_number": 4,
+      "unencrypted_data": "05da0144f3f10ac9cb5e35bbd9b92931",
+      "encrypted_data": "1c226913b46840ebdc5c05566f274c87"
+    },
+    {
+      "sector_number": 120,
+      "block_number": 0,
+      "unencrypted_data": "4646d94304f00def64620ca929f3d092",
+      "encrypted_data": "d291aff19e777ac9109561d19e455384"
+    },
+    {
+      "sector_number": 121,
+      "block_number": 16,
+      "unencrypted_data": "588d35895d1ea9596514cd5c0ca86c56",
+      "encrypted_data": "1b9ce6cc7c2ef0ca233fe451d88b361f"
+    },
+    {
+      "sector_number": 122,
+      "block_number": 28,
+      "unencrypted_data": "4c0112e0073a513d09b3b8f383c5167c",
+      "encrypted_data": "4a1104539c7f035284de7acba3ca05c7"
+    },
+    {
+      "sector_number": 123,
+      "block_number": 0,
+      "unencrypted_data": "661be9fd512704515b494c8e71af1dcf",
+      "encrypted_data": "a74a0bfba9719de56348aa0fe45ee93b"
+    },
+    {
+      "sector_number": 124,
+      "block_number": 7,
+      "unencrypted_data": "ea09e391141af1a7a755320c0588bc4f",
+      "encrypted_data": "7aaf9e8137442a7ac19ca499f369beb0"
+    },
+    {
+      "sector_number": 125,
+      "block_number": 6,
+      "unencrypted_data": "7235fbd387c48a9cdd5407fe034668c9",
+      "encrypted_data": "95c5533209ca975bc888c2eaba15d70d"
+    },
+    {
+      "sector_number": 126,
+      "block_number": 7,
+      "unencrypted_data": "eef9fdd5efcdf3344dd273b42702ce47",
+      "encrypted_data": "599a1d741f8a215fa4c21a86c2a6e730"
+    },
+    {
+      "sector_number": 127,
+      "block_number": 26,
+      "unencrypted_data": "7d1db873d56081dc6cbfae8ef56109cf",
+      "encrypted_data": "e2aa85c2a17706b75e59f721d080e361"
+    },
+    {
+      "sector_number": 128,
+      "block_number": 17,
+      "unencrypted_data": "bcb8cf3610e868ebee3ecd2ee87ab204",
+      "encrypted_data": "5f4c97170c11aaf28dd87fc730ff8943"
+    },
+    {
+      "sector_number": 129,
+      "block_number": 8,
+      "unencrypted_data": "6eccdadb16488cbf3b6c69b95e072ee0",
+      "encrypted_data": "d2e457322629acd467bdfdcb76c17468"
+    },
+    {
+      "sector_number": 130,
+      "block_number": 16,
+      "unencrypted_data": "9468381475a6dc11ad5cb3c9c1310211",
+      "encrypted_data": "bd016ace54488736346461052f1fa54d"
+    },
+    {
+      "sector_number": 131,
+      "block_number": 24,
+      "unencrypted_data": "ee3a27feb2a565662506f55f7755cdcf",
+      "encrypted_data": "1b0e9e16c4373a6d91baa7119c659f5b"
+    },
+    {
+      "sector_number": 132,
+      "block_number": 8,
+      "unencrypted_data": "4da9ea76a139631d70605579ca03b76c",
+      "encrypted_data": "63fcab8c1320c967a0665d8a1b2f39d4"
+    },
+    {
+      "sector_number": 133,
+      "block_number": 7,
+      "unencrypted_data": "7d6dac02f16d2d9b2ebcd90721639c4d",
+      "encrypted_data": "75c38dfe05321b06134b54f25ebee108"
+    },
+    {
+      "sector_number": 134,
+      "block_number": 27,
+      "unencrypted_data": "9ef2d86e40ecbeb4d8cebbf41f445978",
+      "encrypted_data": "66271312a6fe37f9e03e8bb8a36308ac"
+    },
+    {
+      "sector_number": 135,
+      "block_number": 12,
+      "unencrypted_data": "00a1ae724716c7815513a96883c31c52",
+      "encrypted_data": "ba42c3ed54b32f2a70b8b7d5ca230396"
+    },
+    {
+      "sector_number": 136,
+      "block_number": 14,
+      "unencrypted_data": "d95c83fc6c140259d16a03e94c784577",
+      "encrypted_data": "1b16d7fb609aff302609e75447b0a823"
+    },
+    {
+      "sector_number": 137,
+      "block_number": 7,
+      "unencrypted_data": "ac224bb661a7bb6940de9abe75fd2a15",
+      "encrypted_data": "8a8a9c932e58affb776e8033e47f93ac"
+    },
+    {
+      "sector_number": 138,
+      "block_number": 11,
+      "unencrypted_data": "eb515f7cb0b5b068cdfb2dc1d430b63b",
+      "encrypted_data": "b5a42fd886c8a040f66c1cec2b25480c"
+    },
+    {
+      "sector_number": 139,
+      "block_number": 28,
+      "unencrypted_data": "2c6e09e3112b4a1fbd14b3f8d9aea507",
+      "encrypted_data": "e343a4094bfc5105e50ab2845fd4a96d"
+    },
+    {
+      "sector_number": 140,
+      "block_number": 3,
+      "unencrypted_data": "7eabb3de2c136f33a03eb6ba6708d45f",
+      "encrypted_data": "4ae6e233a9352e1c681138dfbb62b4ea"
+    },
+    {
+      "sector_number": 141,
+      "block_number": 30,
+      "unencrypted_data": "f3d81c83f6d1239e1fcde08e12dc9f57",
+      "encrypted_data": "cf27b8517f3c0ac217f57b2d0a650bc6"
+    },
+    {
+      "sector_number": 142,
+      "block_number": 28,
+      "unencrypted_data": "ece3708c2f58ec46f719680bca1e3f8f",
+      "encrypted_data": "263df753aa4f3a1b1c4f69cda22dc6ad"
+    },
+    {
+      "sector_number": 143,
+      "block_number": 9,
+      "unencrypted_data": "49d83924fa6aed2ed4c25df8706f3b52",
+      "encrypted_data": "ec28aae9b3a6430c3bab20001957fa03"
+    },
+    {
+      "sector_number": 144,
+      "block_number": 22,
+      "unencrypted_data": "7c08a111f676e2ebf1f6b6215816368a",
+      "encrypted_data": "f1dd67882561a92df554ff0f5f5cf027"
+    },
+    {
+      "sector_number": 145,
+      "block_number": 13,
+      "unencrypted_data": "129f212746e7ec826538ff71960b3d34",
+      "encrypted_data": "04fe5f43d9ba95caa035da86ddd71c1a"
+    },
+    {
+      "sector_number": 146,
+      "block_number": 1,
+      "unencrypted_data": "fd1299a15f1bde1127c1dda56a32694a",
+      "encrypted_data": "6822c9f60824424e35d7a3cabc3d2d95"
+    },
+    {
+      "sector_number": 147,
+      "block_number": 6,
+      "unencrypted_data": "3dd28261115604b39c0cdb0aaed3743c",
+      "encrypted_data": "23db7642d202004f545dc768f276f7c9"
+    },
+    {
+      "sector_number": 148,
+      "block_number": 3,
+      "unencrypted_data": "377a1e05f8783c6e13a08b2e1fa5d406",
+      "encrypted_data": "f7eb2ede33ff909d8e78aee16e924286"
+    },
+    {
+      "sector_number": 149,
+      "block_number": 1,
+      "unencrypted_data": "ea19d0ada7e6ca6e1ce4b21a5aac9bd8",
+      "encrypted_data": "392f766e7a2217737252a8243fe2474a"
+    },
+    {
+      "sector_number": 150,
+      "block_number": 30,
+      "unencrypted_data": "a4b97e6c810da0deaca230e66fe2c548",
+      "encrypted_data": "0330bceedd1da0aed03305eeb66c262f"
+    },
+    {
+      "sector_number": 151,
+      "block_number": 15,
+      "unencrypted_data": "6af354c1d9329c7f11c24424ac4d7f27",
+      "encrypted_data": "ae2238557ca04f2e6bd52fad7bd67646"
+    },
+    {
+      "sector_number": 152,
+      "block_number": 5,
+      "unencrypted_data": "76fcae6f10e89b0639d5f11f4a21355e",
+      "encrypted_data": "3edcb1450c71b7c671489685beb04010"
+    },
+    {
+      "sector_number": 153,
+      "block_number": 15,
+      "unencrypted_data": "66bb7b364c4fd0ea2d557dfec6cb0094",
+      "encrypted_data": "51149891df05bc28a290fc1e74e33259"
+    },
+    {
+      "sector_number": 154,
+      "block_number": 19,
+      "unencrypted_data": "ab8beeaa8fb5d0853d6ea99027986126",
+      "encrypted_data": "61da8241106577241cf29584dfbb030a"
+    },
+    {
+      "sector_number": 155,
+      "block_number": 28,
+      "unencrypted_data": "b98c5db43242a054d45714b913d097ee",
+      "encrypted_data": "b522dfe07487fdfda80fd3bc20419072"
+    },
+    {
+      "sector_number": 156,
+      "block_number": 26,
+      "unencrypted_data": "b84583cb40b2441347ce78182112dc6b",
+      "encrypted_data": "778afc3f771004dbda82c0a629cf5b8f"
+    },
+    {
+      "sector_number": 157,
+      "block_number": 4,
+      "unencrypted_data": "e032664d4b02b8d5ed01e90b51de7150",
+      "encrypted_data": "39a12fa568b9569f217ce22980af91e0"
+    },
+    {
+      "sector_number": 158,
+      "block_number": 24,
+      "unencrypted_data": "f4ef29722c30a830f368d62145cfa02a",
+      "encrypted_data": "dfa2d25b9293459c2651e46c8ecf9e60"
+    },
+    {
+      "sector_number": 159,
+      "block_number": 9,
+      "unencrypted_data": "4d8d1f711fd16a8eb50246ea0d2de22e",
+      "encrypted_data": "ca1390e1facd8f948f48272873ad40a2"
+    },
+    {
+      "sector_number": 160,
+      "block_number": 23,
+      "unencrypted_data": "382b0df720f987f8fcc732f13da28125",
+      "encrypted_data": "3ee1c3ac57e5c114612eee47affcfae4"
+    },
+    {
+      "sector_number": 161,
+      "block_number": 2,
+      "unencrypted_data": "825bd3d144129d42398e0ee3da227772",
+      "encrypted_data": "cfb82be77bf6a6bf5d4da470f6f3c884"
+    },
+    {
+      "sector_number": 162,
+      "block_number": 17,
+      "unencrypted_data": "39d955dc59f3c31f8e2bd1821551d624",
+      "encrypted_data": "e02cc86ef34c90335eb1678416a114c8"
+    },
+    {
+      "sector_number": 163,
+      "block_number": 23,
+      "unencrypted_data": "2ab2b78ddd256193b9f20a3924b5ac68",
+      "encrypted_data": "87f3fd5babe7c6b4fdd32d7d96c60c14"
+    },
+    {
+      "sector_number": 164,
+      "block_number": 10,
+      "unencrypted_data": "f6f57be1117b9601c2f7db0637d03e12",
+      "encrypted_data": "fdaa1686ab4b76e539344d565a766b8b"
+    },
+    {
+      "sector_number": 165,
+      "block_number": 4,
+      "unencrypted_data": "9c1be66f3a3fc43e879482b7a84cfad9",
+      "encrypted_data": "a59544d0efc54fc8d8d3db7081f4f224"
+    },
+    {
+      "sector_number": 166,
+      "block_number": 1,
+      "unencrypted_data": "93fd172d55a5350c3e15611d8a21baee",
+      "encrypted_data": "84ab8e1ff75cccc752eae4dd623c2053"
+    },
+    {
+      "sector_number": 167,
+      "block_number": 30,
+      "unencrypted_data": "99b52424fd4ee8e1eeac6e78fdaf9a06",
+      "encrypted_data": "6fe09e3d9e81dc0cac800b9d3928caa5"
+    },
+    {
+      "sector_number": 168,
+      "block_number": 24,
+      "unencrypted_data": "6ab578a0bb51cd9c72cda4b586b88f88",
+      "encrypted_data": "15708024187b7f825beed751dd18a408"
+    },
+    {
+      "sector_number": 169,
+      "block_number": 2,
+      "unencrypted_data": "765c9373c4537ca96c45de7998f7304b",
+      "encrypted_data": "7e6927b4429c0ccc071a8757bc7cf1d4"
+    },
+    {
+      "sector_number": 170,
+      "block_number": 3,
+      "unencrypted_data": "071f9a88738ed5d89c6023de973ec53e",
+      "encrypted_data": "ac9afa95c9c7744f19c8b005b7cb3b7d"
+    },
+    {
+      "sector_number": 171,
+      "block_number": 16,
+      "unencrypted_data": "8fa81f891d7a6d2e517da28688029f7d",
+      "encrypted_data": "5898a670fe0e9752805576a747ad0acd"
+    },
+    {
+      "sector_number": 172,
+      "block_number": 22,
+      "unencrypted_data": "7d0f6be04b8cdb232b05534ac5383ffa",
+      "encrypted_data": "5a9c1c0f2b64b3c1654026f39133996f"
+    },
+    {
+      "sector_number": 173,
+      "block_number": 12,
+      "unencrypted_data": "6769882db22d714e2b2212143b8d3450",
+      "encrypted_data": "c3db3b599a8c9dff5f3f12627dc67229"
+    },
+    {
+      "sector_number": 174,
+      "block_number": 21,
+      "unencrypted_data": "f9960336593bf4618d6b72d2b00ebf7e",
+      "encrypted_data": "5e1bb6a6ab6d0c6d12ac962db7a72fb7"
+    },
+    {
+      "sector_number": 175,
+      "block_number": 11,
+      "unencrypted_data": "019d14158852e8e2b90bee6d8c471b0f",
+      "encrypted_data": "6d5bf35726487dfa1d996ac24f1e32f8"
+    },
+    {
+      "sector_number": 176,
+      "block_number": 8,
+      "unencrypted_data": "2b6e6b7770ef26026463e27afa5e3913",
+      "encrypted_data": "ab7aa0e2067db3449615d29fb7f8d26d"
+    },
+    {
+      "sector_number": 177,
+      "block_number": 17,
+      "unencrypted_data": "be89b535fd6b0cafc2d9c507db0694c3",
+      "encrypted_data": "6a81b7af63be2719b23b6301d16f169d"
+    },
+    {
+      "sector_number": 178,
+      "block_number": 23,
+      "unencrypted_data": "ae27b96f40e6b9b31f98330999bf179b",
+      "encrypted_data": "d5450c1c0c1a7377f13d54566116be24"
+    },
+    {
+      "sector_number": 179,
+      "block_number": 9,
+      "unencrypted_data": "3d876b17534754795538c1afbe182c72",
+      "encrypted_data": "4ae6d433483b03694dd6b84fd7f21864"
+    },
+    {
+      "sector_number": 180,
+      "block_number": 19,
+      "unencrypted_data": "4f1ca8062dc1649a94973a5e0099d609",
+      "encrypted_data": "9a21e11410982667248383e59711b870"
+    },
+    {
+      "sector_number": 181,
+      "block_number": 23,
+      "unencrypted_data": "030c355eff56a1ee7b056fbdfa30dbb6",
+      "encrypted_data": "9d1583f84d41e36ed61161e7714aa3cc"
+    },
+    {
+      "sector_number": 182,
+      "block_number": 10,
+      "unencrypted_data": "8bac3dd30b0fbd69aca12e06f34b9de8",
+      "encrypted_data": "ff7552649998d28945516c5b2f1423b8"
+    },
+    {
+      "sector_number": 183,
+      "block_number": 8,
+      "unencrypted_data": "508b2559cb51d75cdb4aaafafa220d26",
+      "encrypted_data": "e188d290263d97a2906c8f55f1aa3e8c"
+    },
+    {
+      "sector_number": 184,
+      "block_number": 31,
+      "unencrypted_data": "d9e6129bbad1908f56cae40efb7fba29",
+      "encrypted_data": "51dd9479c7a43e4eb8c343350c5e5118"
+    },
+    {
+      "sector_number": 185,
+      "block_number": 31,
+      "unencrypted_data": "43a81b3d7d2511cbbdbfeafb98d7b045",
+      "encrypted_data": "4f8b84742d34a4be2730de11e1f8c040"
+    },
+    {
+      "sector_number": 186,
+      "block_number": 10,
+      "unencrypted_data": "36b5f14987481f078b31e2ec95fa5d29",
+      "encrypted_data": "7d6ca4c7c142cd8b363a4a7661212ce9"
+    },
+    {
+      "sector_number": 187,
+      "block_number": 20,
+      "unencrypted_data": "ac491020b9b4e88e396041af77d3f0f1",
+      "encrypted_data": "dd443b3a70831bc6601204c8c31119a3"
+    },
+    {
+      "sector_number": 188,
+      "block_number": 24,
+      "unencrypted_data": "529babbfa3e97d76333e975370810517",
+      "encrypted_data": "12c66d2f66a7ddf69df79edfd4766a38"
+    },
+    {
+      "sector_number": 189,
+      "block_number": 24,
+      "unencrypted_data": "aa4dd50f276e3d7eb5cf213e9b5b626f",
+      "encrypted_data": "4f4adbbe29b3c592df98e77dd86be66b"
+    },
+    {
+      "sector_number": 190,
+      "block_number": 28,
+      "unencrypted_data": "ff0d826bf0dcea746564c0117e1a3a15",
+      "encrypted_data": "22db8385a9e9d64ab46986adf659329b"
+    },
+    {
+      "sector_number": 191,
+      "block_number": 19,
+      "unencrypted_data": "9fe7f10c5fb28dfeed87c7d6ee654c13",
+      "encrypted_data": "01dd96aeba1bd365bebbb4b10bb59899"
+    },
+    {
+      "sector_number": 192,
+      "block_number": 26,
+      "unencrypted_data": "2dbb16ad32fcac48d7597b1a84edcd84",
+      "encrypted_data": "7334dbc2268991a0c57f8e911ae35380"
+    },
+    {
+      "sector_number": 193,
+      "block_number": 1,
+      "unencrypted_data": "1b1eb2da4ba4896ecda4f9ea53b0b64b",
+      "encrypted_data": "4d5545a54ffccd5d0a389cf9a5336118"
+    },
+    {
+      "sector_number": 194,
+      "block_number": 29,
+      "unencrypted_data": "b734a42b99315fbd7db0b2e36722b30b",
+      "encrypted_data": "c3e112a30f997ec8f05862504e9ded0e"
+    },
+    {
+      "sector_number": 195,
+      "block_number": 1,
+      "unencrypted_data": "46b6fe6ea1aa3b5f25b9afebb8f8d916",
+      "encrypted_data": "fcfdb0db3c05a5faff37b344938e348e"
+    },
+    {
+      "sector_number": 196,
+      "block_number": 23,
+      "unencrypted_data": "54615518ea268e35a3ed8cbc9adb8ad6",
+      "encrypted_data": "5b34c516b0dd49bf2a7dd59a7aec9c8a"
+    },
+    {
+      "sector_number": 197,
+      "block_number": 18,
+      "unencrypted_data": "8d3fc2ca3ce97cada68d18eb07d89a9a",
+      "encrypted_data": "edcb5bc5b4a3c03edb496c4e65562ad3"
+    },
+    {
+      "sector_number": 198,
+      "block_number": 3,
+      "unencrypted_data": "26a02332e3bdb1d8dcf88c1d4ddc173c",
+      "encrypted_data": "e11457ec6ac59effcb576c4957ad70e0"
+    },
+    {
+      "sector_number": 199,
+      "block_number": 28,
+      "unencrypted_data": "105b7dc14544e0e8a496842c67433290",
+      "encrypted_data": "43932a27af3044397b74553a71fa5f10"
+    },
+    {
+      "sector_number": 200,
+      "block_number": 17,
+      "unencrypted_data": "5b341e4ee068e603254d99cf3828dfda",
+      "encrypted_data": "0ce70815c5c2487ab50b005b7927609f"
+    },
+    {
+      "sector_number": 201,
+      "block_number": 13,
+      "unencrypted_data": "bf3d2bdaabc4d588caf34b5f99fd26c6",
+      "encrypted_data": "f600e0c48a594144f8fc70c32bc50f97"
+    },
+    {
+      "sector_number": 202,
+      "block_number": 10,
+      "unencrypted_data": "4dd305bb9720eb694a3c533387b58bea",
+      "encrypted_data": "1e89f029a48da19be10d1e18e7614f85"
+    },
+    {
+      "sector_number": 203,
+      "block_number": 14,
+      "unencrypted_data": "3b4b0d7518f2a6800f65ba50d38d02b5",
+      "encrypted_data": "d36437b8c8e5a80b6e671dc0341790b2"
+    },
+    {
+      "sector_number": 204,
+      "block_number": 20,
+      "unencrypted_data": "eea4f8ce5cb60fbc066020186b0d904d",
+      "encrypted_data": "60a38241113b97f23bd494f2a368efc0"
+    },
+    {
+      "sector_number": 205,
+      "block_number": 24,
+      "unencrypted_data": "bc9ed07b6df60e9015615194dff587b8",
+      "encrypted_data": "5d63303f18e05ce72f97b49dfcb8e81e"
+    },
+    {
+      "sector_number": 206,
+      "block_number": 16,
+      "unencrypted_data": "45fb5a1f5cb3788b013973f60431900a",
+      "encrypted_data": "56308497eb2900450d824e0734b0bf4d"
+    },
+    {
+      "sector_number": 207,
+      "block_number": 27,
+      "unencrypted_data": "8f29090f618e2321426eed3b502d1111",
+      "encrypted_data": "d78dc2325813ad5ecb8f91865d4f59c5"
+    },
+    {
+      "sector_number": 208,
+      "block_number": 0,
+      "unencrypted_data": "a277bae998bc7d0ea307a16ef92bd72e",
+      "encrypted_data": "171eef747e511146715f288ecf741108"
+    },
+    {
+      "sector_number": 209,
+      "block_number": 5,
+      "unencrypted_data": "3c578aed5d2ca92b8af4342cc8a8ec71",
+      "encrypted_data": "9bc24f24705ed474891321b8c89f017b"
+    },
+    {
+      "sector_number": 210,
+      "block_number": 23,
+      "unencrypted_data": "c3c008e028468a6e4447a58cf83667a0",
+      "encrypted_data": "f695e5c42e68618fc3ad9eb7c2112495"
+    },
+    {
+      "sector_number": 211,
+      "block_number": 28,
+      "unencrypted_data": "a9fd41d4f8a9016ba0bc46b7e72ceda8",
+      "encrypted_data": "042f227d54550ef1cd7bdfbad20da9a1"
+    },
+    {
+      "sector_number": 212,
+      "block_number": 25,
+      "unencrypted_data": "1d237378c41f99900d48adcabb5c2d2b",
+      "encrypted_data": "e0450ed82f77c698445e3097bc397b8e"
+    },
+    {
+      "sector_number": 213,
+      "block_number": 10,
+      "unencrypted_data": "85d6d1ee12f8710f24c4088161e79e34",
+      "encrypted_data": "dffe7abd6d2bc3ec9ad7790f02082151"
+    },
+    {
+      "sector_number": 214,
+      "block_number": 13,
+      "unencrypted_data": "3e888e3bc0eceef69aef91e54df6cb7b",
+      "encrypted_data": "b146ef17fba8cdb0e3bc76db1d4fd283"
+    },
+    {
+      "sector_number": 215,
+      "block_number": 1,
+      "unencrypted_data": "a207f43447877ea81035765e36cbf050",
+      "encrypted_data": "efa05cc30e1d165b71f4ee280dd42af5"
+    },
+    {
+      "sector_number": 216,
+      "block_number": 8,
+      "unencrypted_data": "138f273409181de72fed7fd490877cf7",
+      "encrypted_data": "dfba3bd4528e94cee1f9ba9fda930389"
+    },
+    {
+      "sector_number": 217,
+      "block_number": 23,
+      "unencrypted_data": "d56ca8ace3d650d041c92850bf987376",
+      "encrypted_data": "717117a32194c71d697ce6c382c16321"
+    },
+    {
+      "sector_number": 218,
+      "block_number": 20,
+      "unencrypted_data": "89037976cfa4342149eb69ef87ef4b41",
+      "encrypted_data": "ffabb47acab3c003db03f2a3fbf3bce6"
+    },
+    {
+      "sector_number": 219,
+      "block_number": 11,
+      "unencrypted_data": "8155c102be703f4d1867cd32ff6a1a24",
+      "encrypted_data": "3cf9fd855def01c0d86c199c5792873a"
+    },
+    {
+      "sector_number": 220,
+      "block_number": 0,
+      "unencrypted_data": "fe0634d85a93da562980301da740967f",
+      "encrypted_data": "ef89096647a9ce9f1d077b6a27193419"
+    },
+    {
+      "sector_number": 221,
+      "block_number": 8,
+      "unencrypted_data": "4685030f5a3f2059efd3d4fd082577e6",
+      "encrypted_data": "faa83d0ca92efae4ebdcfa4252171609"
+    },
+    {
+      "sector_number": 222,
+      "block_number": 17,
+      "unencrypted_data": "a16e1abbaec4354c62b8d7cccb2fb0e8",
+      "encrypted_data": "9801f1e9f10395852ea2112d77be0234"
+    },
+    {
+      "sector_number": 223,
+      "block_number": 10,
+      "unencrypted_data": "19141cec5d755ac4dba5a4418fb09329",
+      "encrypted_data": "79b478141883a100d567306c7cec6b1f"
+    },
+    {
+      "sector_number": 224,
+      "block_number": 12,
+      "unencrypted_data": "37f10b71fec4bc57673604ecfec8b614",
+      "encrypted_data": "83d2718a723ec2da0aff01756142e558"
+    },
+    {
+      "sector_number": 225,
+      "block_number": 17,
+      "unencrypted_data": "e361987df14f3a975d33c571a079e168",
+      "encrypted_data": "3b652328da1288f9b5b17154bf7e8fc9"
+    },
+    {
+      "sector_number": 226,
+      "block_number": 9,
+      "unencrypted_data": "b24b44626528b842c8635af4022a6dcd",
+      "encrypted_data": "dd0ef71c595bab25c686ccc42f87d0ee"
+    },
+    {
+      "sector_number": 227,
+      "block_number": 7,
+      "unencrypted_data": "543d00bd4321f225ab4ec2b92dd67a00",
+      "encrypted_data": "a23b6be06ebd7c068f3cbe0855b01ae5"
+    },
+    {
+      "sector_number": 228,
+      "block_number": 22,
+      "unencrypted_data": "dcb72e04daf29dc6b89077f89859b3ce",
+      "encrypted_data": "f5114ef6437a8b5243efadf1c2a7d9ca"
+    },
+    {
+      "sector_number": 229,
+      "block_number": 30,
+      "unencrypted_data": "5dee4e46acd6ef65418eb5406ca44c27",
+      "encrypted_data": "1f351dad2d6e4653f9db9a1166d817a2"
+    },
+    {
+      "sector_number": 230,
+      "block_number": 6,
+      "unencrypted_data": "8404034776c1ed07bb7df4487d702a30",
+      "encrypted_data": "7034fd7c5792102798eafafc0b282414"
+    },
+    {
+      "sector_number": 231,
+      "block_number": 25,
+      "unencrypted_data": "b16e831ac17c3c9a765e8f9a57dfb580",
+      "encrypted_data": "e674790855105777e3f56847450d3563"
+    },
+    {
+      "sector_number": 232,
+      "block_number": 6,
+      "unencrypted_data": "8ef2576cc965f99713ff503f65ab1dfa",
+      "encrypted_data": "f6dc48f9777516fb540d9fcd642628b5"
+    },
+    {
+      "sector_number": 233,
+      "block_number": 6,
+      "unencrypted_data": "d655ea9e31affe15c5256bb5aec94e67",
+      "encrypted_data": "04fe1e95ddf8f90f00658b3f887c4ca7"
+    },
+    {
+      "sector_number": 234,
+      "block_number": 20,
+      "unencrypted_data": "4a78c8399cbc8a27ebd59f6023ff3e15",
+      "encrypted_data": "29fff7403faf172c691c5c9f31973164"
+    },
+    {
+      "sector_number": 235,
+      "block_number": 24,
+      "unencrypted_data": "584238ec4b1a27fd81dd45f0a1fdc17e",
+      "encrypted_data": "a46325f703dadc4c9b6f55e6da096348"
+    },
+    {
+      "sector_number": 236,
+      "block_number": 8,
+      "unencrypted_data": "6b1f97948c76accd08377df3a58ce623",
+      "encrypted_data": "070b1dd240c31ac1a3223c9872210a96"
+    },
+    {
+      "sector_number": 237,
+      "block_number": 19,
+      "unencrypted_data": "17837e559c80341a6390c35faf87c6e4",
+      "encrypted_data": "42d161c25052625028339100d53e2b36"
+    },
+    {
+      "sector_number": 238,
+      "block_number": 3,
+      "unencrypted_data": "070a99f8ee57581d5e7c5e99f053c639",
+      "encrypted_data": "b82afcb69871f70b30726776aa72c3ec"
+    },
+    {
+      "sector_number": 239,
+      "block_number": 0,
+      "unencrypted_data": "49e01486c6675db7ab5d5167c253e6e1",
+      "encrypted_data": "d034983492015083d733f433432db23c"
+    },
+    {
+      "sector_number": 240,
+      "block_number": 17,
+      "unencrypted_data": "ac8faa3047bb88e1294d3b3e141047a9",
+      "encrypted_data": "d2e52bd1b06e007cac174432042b6a3a"
+    },
+    {
+      "sector_number": 241,
+      "block_number": 16,
+      "unencrypted_data": "94760d56e1ab5e8733236bec3797e365",
+      "encrypted_data": "fbc14ea36cd525682c44205ca40a6e6a"
+    },
+    {
+      "sector_number": 242,
+      "block_number": 12,
+      "unencrypted_data": "51df5fc508a9ee92d30e12fae7eef6ab",
+      "encrypted_data": "3b496d306cf8858f5df006f1860bab35"
+    },
+    {
+      "sector_number": 243,
+      "block_number": 29,
+      "unencrypted_data": "ccb203858d338d19223f61757b160885",
+      "encrypted_data": "dc1e42385fbc5fbd57817ba0fde7eaf0"
+    },
+    {
+      "sector_number": 244,
+      "block_number": 26,
+      "unencrypted_data": "f4a7470e0b2a886c8187f3743f177ec8",
+      "encrypted_data": "23cbdbc33aca3213aa4a77d6aa12d1c1"
+    },
+    {
+      "sector_number": 245,
+      "block_number": 18,
+      "unencrypted_data": "6b6f3216fef7bd5b3336cdc7fc0d9d7f",
+      "encrypted_data": "3d90ce3113728312165d3316b8f313ff"
+    },
+    {
+      "sector_number": 246,
+      "block_number": 18,
+      "unencrypted_data": "f6f36d2564ab3d070b9388b6df4ec4c1",
+      "encrypted_data": "8e5b0a4041db81f888dccff8450a288b"
+    },
+    {
+      "sector_number": 247,
+      "block_number": 9,
+      "unencrypted_data": "467322e6f8691fd68127ca177f3850ab",
+      "encrypted_data": "cebdb413c6ea6e9fea900baeba83a1e8"
+    },
+    {
+      "sector_number": 248,
+      "block_number": 12,
+      "unencrypted_data": "abf57a9be88422944c9067b3ad6af552",
+      "encrypted_data": "f722574fee66c0490e181892ab6e3f35"
+    },
+    {
+      "sector_number": 249,
+      "block_number": 7,
+      "unencrypted_data": "ddd2d8bf6200497d4defb95809762e69",
+      "encrypted_data": "3c36243b0dbf24f4c60080a4b8b62ab8"
+    },
+    {
+      "sector_number": 250,
+      "block_number": 21,
+      "unencrypted_data": "0d2d59e7a86f99b6999c83cecf66b9be",
+      "encrypted_data": "cfbf275a520e90bc528c00f8e6efb8c4"
+    },
+    {
+      "sector_number": 251,
+      "block_number": 24,
+      "unencrypted_data": "e7af41394d0a9caaaeeb7f7c04e4f86a",
+      "encrypted_data": "f89cf04b66ebaed7f0494b55b4e9ec46"
+    },
+    {
+      "sector_number": 252,
+      "block_number": 7,
+      "unencrypted_data": "2854cad3eb45defdcdeeb30bb7268a1e",
+      "encrypted_data": "e779300c026bb7184c6e2d09db986558"
+    },
+    {
+      "sector_number": 253,
+      "block_number": 11,
+      "unencrypted_data": "e69a86b5f8dd5d5b99ca8287504e0bf6",
+      "encrypted_data": "93007b5d6602cfdf9eac5d21ea9dd8e9"
+    },
+    {
+      "sector_number": 254,
+      "block_number": 17,
+      "unencrypted_data": "3fd6f1383fb6d2625982b618389d677d",
+      "encrypted_data": "bc9280046e2f2aa971b6c54fa0a16e30"
+    },
+    {
+      "sector_number": 255,
+      "block_number": 14,
+      "unencrypted_data": "83fd26fc8a443b33eb61c630c4bc587b",
+      "encrypted_data": "f9a54311016875127f011b119a638c7c"
+    }
+  ]
+}
\ No newline at end of file
diff --git a/tests/gen_test_vectors.py b/tests/gen_test_vectors.py
new file mode 100644
index 0000000..c81b44a
--- /dev/null
+++ b/tests/gen_test_vectors.py
@@ -0,0 +1,81 @@
+# aes-xts-pur64 is OpenCL code for aes-xts256-plain64 encryption compatible with LUKS
+#
+#     Copyright (C) 2025  Kirill Shakirov
+#
+#     This program is free software: you can redistribute it and/or modify
+#     it under the terms of the GNU General Public License as published by
+#     the Free Software Foundation, either version 3 of the License, or
+#     (at your option) any later version.
+#
+#     This program is distributed in the hope that it will be useful,
+#     but WITHOUT ANY WARRANTY; without even the implied warranty of
+#     MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+#     GNU General Public License for more details.
+#
+#     You should have received a copy of the GNU General Public License
+#     along with this program.  If not, see <https://www.gnu.org/licenses/>.
+
+import json
+import subprocess
+import nyanger.simple.static as nya_stat
+import random
+
+# Init logger
+log = nya_stat.get_logger("nyan")
+
+LUKS_VOL_FILE_NAME = "./data/test_vol.img"
+UNENCRYPT_DATA_FILE_NAME= "./data/unencrypt.img"
+OUT_FILE_NAME = "./data/test_vectors.json"
+KEY_FILE_NAME = "./data/master.key"
+VECTORS_NUM=256
+
+def read_metadata(file_name: str) -> dict:
+    #cryptsetup luksDump --dump-json-metadata /dev/loop0
+    luks_cmd: list[str] = ["cryptsetup", "luksDump", "--dump-json-metadata", file_name]
+
+    result = subprocess.run(luks_cmd, capture_output=True, encoding="UTF-8")
+    if result.returncode == 0 and result.stdout is not None:
+        metadata = json.loads(result.stdout)
+        return metadata
+    else:
+        raise Exception(f"Error executing 'cryptsetup' binary! {result.stderr}")
+
+
+def main():
+    log.info("Start :3")
+    metad = read_metadata(LUKS_VOL_FILE_NAME)
+    log.info(f"metadata:\n{metad}")
+
+    segments_offset_bytes = int(metad["segments"]["0"]["offset"])
+
+    with open(LUKS_VOL_FILE_NAME, 'rb') as luks_file:
+        luks_file.seek(segments_offset_bytes)
+        enc_data = luks_file.read(512*VECTORS_NUM)
+
+    with open(UNENCRYPT_DATA_FILE_NAME, 'rb') as data_file:
+        uenc_data = data_file.read(512*VECTORS_NUM)
+
+    with open(KEY_FILE_NAME, "rb") as key_file:
+        key_data = key_file.read(32)
+
+    data_list = []
+    for sec_n in range(VECTORS_NUM):
+        block_num = random.randint(0, 31)
+        _offset = ((sec_n*512)+(block_num*16))
+        entry = {
+            "sector_number": sec_n,
+            "block_number": block_num,
+            "unencrypted_data": uenc_data[_offset:_offset+16].hex(),
+            "encrypted_data": enc_data[_offset:_offset+16].hex()
+        }
+        data_list.append(entry)
+
+    with open(OUT_FILE_NAME, 'wt') as out_file:
+        json.dump({"encryption_key": key_data.hex(sep=' ', bytes_per_sep=16), "vectors": data_list},
+                  out_file, indent="  ")
+
+    log.info(" --- Finish --- ")
+
+
+if __name__ == '__main__':
+    main()
\ No newline at end of file
diff --git a/tests/pyproject.toml b/tests/pyproject.toml
new file mode 100644
index 0000000..6e8ffe2
--- /dev/null
+++ b/tests/pyproject.toml
@@ -0,0 +1,9 @@
+[project]
+name = "aes-xts-pur64"
+version = "0.1.0"
+description = "Tests fro aes-xts-pur64"
+requires-python = ">=3.13"
+dependencies = [
+    "nyanger==0.9.2",
+    "pyopencl==2025.2.7"
+]
diff --git a/tests/test_vectors.py b/tests/test_vectors.py
new file mode 100644
index 0000000..df7efe2
--- /dev/null
+++ b/tests/test_vectors.py
@@ -0,0 +1,227 @@
+# aes-xts-pur64 is OpenCL code for aes-xts256-plain64 encryption compatible with LUKS
+#
+#     Copyright (C) 2025  Kirill Shakirov
+#
+#     This program is free software: you can redistribute it and/or modify
+#     it under the terms of the GNU General Public License as published by
+#     the Free Software Foundation, either version 3 of the License, or
+#     (at your option) any later version.
+#
+#     This program is distributed in the hope that it will be useful,
+#     but WITHOUT ANY WARRANTY; without even the implied warranty of
+#     MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+#     GNU General Public License for more details.
+#
+#     You should have received a copy of the GNU General Public License
+#     along with this program.  If not, see <https://www.gnu.org/licenses/>.
+
+
+import json
+import time
+import argparse
+from nyanger.simple.log_writers.file_writer import FileWriter
+from nyanger.simple.log_writers.console_writer import ConsoleWriter
+import nyanger.simple.static as nya_stat
+import pyopencl as cl
+
+# Init logger
+log: nya_stat.Nyanger
+
+def load_program(include_dirs: list[str], file_name: str, ctx: cl.Context) -> cl.Program:
+    with open(file_name, 'rt') as ft:
+        prog_src = ft.read()
+        opt_list = [f"-I {d}" for d in include_dirs]
+        prg = cl.Program(ctx, prog_src).build(options=opt_list)
+        return prg
+
+
+def test(cl_src_file: str, cl_incl_dir: str, vectors_file: str):
+    log.info("Starting tests...")
+    total_errors = 0
+    start_time = time.time()
+
+    log.info(f"Loading test vectors from {vectors_file}")
+
+    with open(vectors_file, "rt") as vec_file:
+        test_json = json.load(vec_file)
+
+    b_key = bytes.fromhex(test_json["encryption_key"])
+    test_vectors = test_json["vectors"]
+    test_vectors_num = len(test_vectors)
+
+    log.info(f"Data key: {b_key[:16].hex()}")
+    log.info(f"Tweak key: {b_key[16:].hex()}")
+
+    b_u_data = bytearray(test_vectors_num * 16)
+    b_enc_data = bytearray(test_vectors_num * 16)
+    b_sec_nums = bytearray(test_vectors_num * 8)
+    b_blk_nums = bytearray(test_vectors_num * 4)
+
+    for i, _vect in enumerate(test_vectors):
+        of16 = i * 16
+        of8 = i * 8
+        of4 = i * 4
+        b_u_data[of16:of16 + 16] = bytes.fromhex(_vect["unencrypted_data"])
+        b_enc_data[of16:of16 + 16] = bytes.fromhex(_vect["encrypted_data"])
+        b_sec_nums[of8:of8 + 8] = _vect["sector_number"].to_bytes(length=8, byteorder="little")
+        b_blk_nums[of4:of4 + 4] = _vect["block_number"].to_bytes(length=4, byteorder="little")
+
+    log.info(f"{test_vectors_num} test vectors loaded.")
+
+    log.info("Getting OCL devices...")
+    # get OCL devices
+    cl_platforms = cl.get_platforms()
+    log.info(f"{len(cl_platforms)} platforms found.")
+    cl_devices = []
+    for platform in cl_platforms:
+        _dev = platform.get_devices(cl.device_type.ALL)
+        log.info(f"Platform \"{platform.name}\" have {len(_dev)} devices: {[_d.name for _d in _dev]}")
+        cl_devices.extend(_dev)
+
+    print()
+
+    for cl_device in cl_devices:
+        log.info(f"Testing on {cl.device_type.to_string(cl_device.type)} device: {cl_device.name}")
+        cl_ctx = cl.Context(devices=[cl_device])
+        cl_queue = cl.CommandQueue(cl_ctx)
+        cl_prg = load_program([cl_incl_dir], cl_src_file, cl_ctx)
+
+        encrypt_data_kernel = cl_prg.encrypt_data
+        decrypt_data_kernel = cl_prg.decrypt_data
+
+        mf = cl.mem_flags
+
+        clb_key = cl.Buffer(cl_ctx, mf.READ_ONLY, size=32)
+        clb_sec_nums = cl.Buffer(cl_ctx, mf.READ_ONLY, size=8 * test_vectors_num)
+        clb_blk_nums = cl.Buffer(cl_ctx, mf.READ_ONLY, size=4 * test_vectors_num)
+        clb_u_data = cl.Buffer(cl_ctx, mf.READ_ONLY, size=16 * test_vectors_num)
+        clb_enc_data = cl.Buffer(cl_ctx, mf.WRITE_ONLY, size=16 * test_vectors_num)
+
+        log.info("Executing encrypt kernel...")
+
+        cl.enqueue_copy(cl_queue, clb_sec_nums, b_sec_nums)
+        cl.enqueue_copy(cl_queue, clb_blk_nums, b_blk_nums)
+        cl.enqueue_copy(cl_queue, clb_key, b_key)
+        cl.enqueue_copy(cl_queue, clb_u_data, b_u_data)
+
+        knl_e = encrypt_data_kernel(cl_queue, (test_vectors_num,), None,
+                                    clb_sec_nums, clb_blk_nums, clb_key, clb_u_data, clb_enc_data)
+
+        b_test_enc_data = bytearray(test_vectors_num * 16)
+        cl.enqueue_copy(cl_queue, b_test_enc_data, clb_enc_data, wait_for=[knl_e])
+
+        log.info("Executing decrypt kernel...")
+
+        cl.enqueue_copy(cl_queue, clb_sec_nums, b_sec_nums)
+        cl.enqueue_copy(cl_queue, clb_blk_nums, b_blk_nums)
+        cl.enqueue_copy(cl_queue, clb_key, b_key)
+        cl.enqueue_copy(cl_queue, clb_u_data, b_enc_data)
+
+        knl_e = decrypt_data_kernel(cl_queue, (test_vectors_num,), None,
+                                    clb_sec_nums, clb_blk_nums, clb_key, clb_u_data, clb_enc_data)
+
+        b_test_u_data = bytearray(test_vectors_num * 16)
+        cl.enqueue_copy(cl_queue, b_test_u_data, clb_enc_data, wait_for=[knl_e])
+
+
+        cl_queue.finish()
+
+        log.info("Comparing encryption results...")
+        right = 0
+        wrong = 0
+        for i in range(test_vectors_num):
+            of16 = i * 16
+            if b_enc_data[of16:of16 + 16] == b_test_enc_data[of16:of16 + 16]:
+                right += 1
+            else:
+                wrong += 1
+                log.error(f"Missmatch of test vector {i} !")
+                log.error(f"  expected:{b_enc_data[of16:of16 + 16].hex()}")
+                log.error(f"    actual:{b_test_enc_data[of16:of16 + 16].hex()}")
+
+        if wrong > 0:
+            log.error(f"Test failed, total of {wrong} mismatch results!")
+        else:
+            log.info("All match! Success!")
+        total_errors += wrong
+
+
+        log.info("Comparing decryption results...")
+        right = 0
+        wrong = 0
+        for i in range(test_vectors_num):
+            of16 = i * 16
+            if b_u_data[of16:of16 + 16] == b_test_u_data[of16:of16 + 16]:
+                right += 1
+            else:
+                wrong += 1
+                log.error(f"Missmatch of test vector {i} !")
+                log.error(f"  expected:{b_u_data[of16:of16 + 16].hex()}")
+                log.error(f"    actual:{b_test_u_data[of16:of16 + 16].hex()}")
+
+        if wrong > 0:
+            log.error(f"Test failed, total of {wrong} mismatch results!\n")
+        else:
+            log.info("All match! Success!\n")
+        total_errors += wrong
+
+    run_time = time.time() - start_time
+    if total_errors > 0:
+        log.error(f"--- Finish in {run_time} seconds with {total_errors} failed tests ---\n\n")
+    else:
+        log.info(f"--- Finish in {run_time} seconds without errors ---\n\n")
+
+
+
+def cmd_parse():
+    # Initialize arguments parser
+    parser = argparse.ArgumentParser(
+        prog="test_vectors.py",
+        description="This test suite for OpenCL AES-XTS implementation",
+        epilog="Have a nice day!")
+
+    parser.add_argument("-s", "--cl-src-file",
+                        action="store",
+                        default="../src/test_aes_xts256_plain.cl",
+                        help="OpenCL src file with defined test kernels.",
+                        required=False)
+    parser.add_argument("-d", "--cl-incl-dir",
+                        action="store",
+                        default="../src",
+                        help="OpenCL include dir with aes256_xts_plain.cl src file.",
+                        required=False)
+    parser.add_argument("-v", "--test-vectors",
+                        action="store",
+                        default="./data/test_vectors.json",
+                        help="Path to json file with test vectors.",
+                        required=False)
+
+    parser.add_argument("-l", "--log-file",
+                        action="store",
+                        default=None,
+                        help="Path to optional log file.",
+                        required=False)
+
+    return parser.parse_args()
+
+
+def main():
+    argumets = cmd_parse()
+    # Init logger
+    global log
+    if argumets.log_file is not None:
+        log = nya_stat.get_logger("nyan", log_writers=[ConsoleWriter(), FileWriter(argumets.log_file)])
+    else:
+        log = nya_stat.get_logger("nyan")
+
+    log.start()
+
+    test(cl_src_file=argumets.cl_src_file,
+         cl_incl_dir=argumets.cl_incl_dir,
+         vectors_file=argumets.test_vectors)
+
+    log.stop()
+
+
+if __name__ == '__main__':
+    main()