abc/GmSSL
1
0
Fork 0
mirror of https://github.com/guanzhi/GmSSL.git synced 2026-08-08 23:13:38 +08:00
Code Issues Packages Projects Releases Wiki Activity

Update sm4_cl.c

Browse Source
This commit is contained in:
Zhi Guan
2024-04-24 17:38:19 +08:00
parent 2e6cef2999
commit 2567d86fdb
1 changed files with 9 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
src/sm4_cl.c
View File

@@ -332,13 +332,15 @@ __constant unsigned char SBOX[256] = {
__kernel void sm4_ctr32_encrypt(__global const unsigned int *rkey, __global const unsigned int *ctr, __global unsigned char *data) __kernel void sm4_ctr32_encrypt(__global const unsigned int *rkey, __global const unsigned int *ctr, __global unsigned char *data)
{ {
__local unsigned char S[256];
__local unsigned int rk[32];
unsigned int x0, x1, x2, x3, x4, i, t; unsigned int x0, x1, x2, x3, x4, i, t;
uint global_id = get_global_id(0); uint global_id = get_global_id(0);
__global unsigned int *out = (__global unsigned int *)(data + 16 * global_id); __global unsigned int *out = (__global unsigned int *)(data + 16 * global_id);
#if 0
// use local mem is little slower on Apple M2
__local unsigned char S[256];
__local unsigned int rk[32];
if (get_local_id(0) == 0) { if (get_local_id(0) == 0) {
for (i = 0; i < 256; i++) { for (i = 0; i < 256; i++) {
S[i] = SBOX[i]; S[i] = SBOX[i];
@@ -347,6 +349,10 @@ __kernel void sm4_ctr32_encrypt(__global const unsigned int *rkey, __global cons
rk[i] = rkey[i]; rk[i] = rkey[i];
} }
} }
#else
__constant unsigned char *S = SBOX;
__global const unsigned int *rk = rkey;
#endif
x0 = ctr[0]; x0 = ctr[0];
x1 = ctr[1]; x1 = ctr[1];