mirror of
https://github.com/Jimmy-Z/bfCL.git
synced 2025-06-18 19:15:52 -04:00
74 lines
1.9 KiB
Common Lisp
74 lines
1.9 KiB
Common Lisp
|
|
__kernel void sha1_16_test(
|
|
__global const uint32_t *in,
|
|
__global uint32_t *out)
|
|
{
|
|
unsigned offset = get_global_id(0) * BLOCKS_PER_ITEM * 16 / sizeof(uint32_t);
|
|
#if BLOCKS_PER_ITEM != 1
|
|
for(unsigned i = 0; i < BLOCKS_PER_ITEM; ++i){
|
|
#endif
|
|
uint32_t buf[4];
|
|
buf[0] = in[offset + 0];
|
|
buf[1] = in[offset + 1];
|
|
buf[2] = in[offset + 2];
|
|
buf[3] = in[offset + 3];
|
|
sha1_16((unsigned char*)buf);
|
|
out[offset + 0] = buf[0];
|
|
out[offset + 1] = buf[1];
|
|
out[offset + 2] = buf[2];
|
|
out[offset + 3] = buf[3];
|
|
#if BLOCKS_PER_ITEM != 1
|
|
offset += 16;
|
|
}
|
|
#endif
|
|
}
|
|
|
|
#define AES_BLOCK_SIZE 16
|
|
|
|
__kernel void aes_enc_128_test(
|
|
__global const uint32_t *key,
|
|
__global const uint32_t *in,
|
|
__global uint32_t *out)
|
|
{
|
|
uint32_t rk[RK_LEN];
|
|
rk[0] = key[0]; rk[1] = key[1]; rk[2] = key[2]; rk[3] = key[3];
|
|
aes_set_key_enc_128(rk);
|
|
unsigned offset = get_global_id(0) * BLOCKS_PER_ITEM * AES_BLOCK_SIZE / 4;
|
|
in += offset; out += offset;
|
|
#if BLOCKS_PER_ITEM != 1
|
|
for (unsigned i = 0; i < BLOCKS_PER_ITEM; ++i) {
|
|
#endif
|
|
uint32_t buf[4];
|
|
buf[0] = in[0]; buf[1] = in[1]; buf[2] = in[2]; buf[3] = in[3];
|
|
aes_encrypt_128(rk, buf);
|
|
out[0] = buf[0]; out[1] = buf[1]; out[2] = buf[2]; out[3] = buf[3];
|
|
#if BLOCKS_PER_ITEM != 1
|
|
offset += AES_BLOCK_SIZE / 4;
|
|
}
|
|
#endif
|
|
}
|
|
|
|
__kernel void aes_dec_128_test(
|
|
__global const uint32_t *key,
|
|
__global const uint32_t *in,
|
|
__global uint32_t *out)
|
|
{
|
|
uint32_t rk[RK_LEN];
|
|
rk[0] = key[0]; rk[1] = key[1]; rk[2] = key[2]; rk[3] = key[3];
|
|
aes_set_key_dec_128(rk);
|
|
unsigned offset = get_global_id(0) * BLOCKS_PER_ITEM * AES_BLOCK_SIZE / 4;
|
|
in += offset; out += offset;
|
|
#if BLOCKS_PER_ITEM != 1
|
|
for (unsigned i = 0; i < BLOCKS_PER_ITEM; ++i) {
|
|
#endif
|
|
uint32_t buf[4];
|
|
buf[0] = in[0]; buf[1] = in[1]; buf[2] = in[2]; buf[3] = in[3];
|
|
aes_decrypt_128(rk, buf);
|
|
out[0] = buf[0]; out[1] = buf[1]; out[2] = buf[2]; out[3] = buf[3];
|
|
#if BLOCKS_PER_ITEM != 1
|
|
offset += AES_BLOCK_SIZE / 4;
|
|
}
|
|
#endif
|
|
}
|
|
|