From 2567d86fdb9249b08e3f32cb91538336856f05db Mon Sep 17 00:00:00 2001
From: Zhi Guan <guanzhi1980@gmail.com>
Date: Wed, 24 Apr 2024 17:38:19 +0800
Subject: [PATCH] Update sm4_cl.c

---
 src/sm4_cl.c | 12 +++++++++---
 1 file changed, 9 insertions(+), 3 deletions(-)

diff --git a/src/sm4_cl.c b/src/sm4_cl.c
index 64dbfdcf..52ad3b52 100644
--- a/src/sm4_cl.c
+++ b/src/sm4_cl.c
@@ -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)
 {
-	__local unsigned char S[256];
-	__local unsigned int rk[32];
-
 	unsigned int x0, x1, x2, x3, x4, i, t;
 	uint global_id = get_global_id(0);
 	__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) {
 		for (i = 0; i < 256; 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];
 		}
 	}
+#else
+	__constant unsigned char *S = SBOX;
+	__global const unsigned int *rk = rkey;
+#endif
 
 	x0 = ctr[0];
 	x1 = ctr[1];