diff --git a/CMakeLists.txt b/CMakeLists.txt index 9f621622..4c0db5ea 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -602,6 +602,14 @@ if (ENABLE_SPHINCS) list(APPEND tests sphincs) endif() +if (ENABLE_SM4_CL OR ENABLE_LMS_CL OR ENABLE_XMSS_CL) + list(APPEND src src/cl.c) +endif() + +if (ENABLE_LMS_CL OR ENABLE_XMSS_CL) + list(APPEND src src/sm3_cl.c) +endif() + if (ENABLE_KYBER) message(STATUS "ENABLE_KYBER is ON") @@ -998,7 +1006,7 @@ endif() # set(CPACK_PACKAGE_NAME "GmSSL") set(CPACK_PACKAGE_VENDOR "GmSSL develop team") -set(CPACK_PACKAGE_VERSION "3.3.0-dev.1166") +set(CPACK_PACKAGE_VERSION "3.3.0-dev.1167") set(CPACK_PACKAGE_DESCRIPTION_FILE ${PROJECT_SOURCE_DIR}/README.md) set(CPACK_NSIS_MODIFY_PATH ON) include(CPack) diff --git a/include/gmssl/version.h b/include/gmssl/version.h index 2cd3c8e0..ce0eb4dc 100644 --- a/include/gmssl/version.h +++ b/include/gmssl/version.h @@ -18,7 +18,7 @@ extern "C" { #define GMSSL_VERSION_NUM 30300 -#define GMSSL_VERSION_STR "GmSSL 3.3.0-dev.1166" +#define GMSSL_VERSION_STR "GmSSL 3.3.0-dev.1167" int gmssl_version_num(void); const char *gmssl_version_str(void); diff --git a/src/cl.c b/src/cl.c new file mode 100644 index 00000000..40f48070 --- /dev/null +++ b/src/cl.c @@ -0,0 +1,118 @@ +/* + * Copyright 2014-2026 The GmSSL Project. All Rights Reserved. + * + * Licensed under the Apache License, Version 2.0 (the License); you may + * not use this file except in compliance with the License. + * + * http://www.apache.org/licenses/LICENSE-2.0 + */ + +#include "cl.h" +#include + + +const char *gmssl_cl_error_string(cl_int err) +{ + switch (err) { + case CL_SUCCESS: return "CL_SUCCESS"; + case CL_DEVICE_NOT_FOUND: return "CL_DEVICE_NOT_FOUND"; + case CL_DEVICE_NOT_AVAILABLE: return "CL_DEVICE_NOT_AVAILABLE"; + case CL_COMPILER_NOT_AVAILABLE: return "CL_COMPILER_NOT_AVAILABLE"; + case CL_MEM_OBJECT_ALLOCATION_FAILURE: return "CL_MEM_OBJECT_ALLOCATION_FAILURE"; + case CL_OUT_OF_RESOURCES: return "CL_OUT_OF_RESOURCES"; + case CL_OUT_OF_HOST_MEMORY: return "CL_OUT_OF_HOST_MEMORY"; + case CL_PROFILING_INFO_NOT_AVAILABLE: return "CL_PROFILING_INFO_NOT_AVAILABLE"; + case CL_MEM_COPY_OVERLAP: return "CL_MEM_COPY_OVERLAP"; + case CL_IMAGE_FORMAT_MISMATCH: return "CL_IMAGE_FORMAT_MISMATCH"; + case CL_IMAGE_FORMAT_NOT_SUPPORTED: return "CL_IMAGE_FORMAT_NOT_SUPPORTED"; + case CL_BUILD_PROGRAM_FAILURE: return "CL_BUILD_PROGRAM_FAILURE"; + case CL_MAP_FAILURE: return "CL_MAP_FAILURE"; + case CL_INVALID_VALUE: return "CL_INVALID_VALUE"; + case CL_INVALID_DEVICE_TYPE: return "CL_INVALID_DEVICE_TYPE"; + case CL_INVALID_PLATFORM: return "CL_INVALID_PLATFORM"; + case CL_INVALID_DEVICE: return "CL_INVALID_DEVICE"; + case CL_INVALID_CONTEXT: return "CL_INVALID_CONTEXT"; + case CL_INVALID_QUEUE_PROPERTIES: return "CL_INVALID_QUEUE_PROPERTIES"; + case CL_INVALID_COMMAND_QUEUE: return "CL_INVALID_COMMAND_QUEUE"; + case CL_INVALID_HOST_PTR: return "CL_INVALID_HOST_PTR"; + case CL_INVALID_MEM_OBJECT: return "CL_INVALID_MEM_OBJECT"; + case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"; + case CL_INVALID_IMAGE_SIZE: return "CL_INVALID_IMAGE_SIZE"; + case CL_INVALID_SAMPLER: return "CL_INVALID_SAMPLER"; + case CL_INVALID_BINARY: return "CL_INVALID_BINARY"; + case CL_INVALID_BUILD_OPTIONS: return "CL_INVALID_BUILD_OPTIONS"; + case CL_INVALID_PROGRAM: return "CL_INVALID_PROGRAM"; + case CL_INVALID_PROGRAM_EXECUTABLE: return "CL_INVALID_PROGRAM_EXECUTABLE"; + case CL_INVALID_KERNEL_NAME: return "CL_INVALID_KERNEL_NAME"; + case CL_INVALID_KERNEL_DEFINITION: return "CL_INVALID_KERNEL_DEFINITION"; + case CL_INVALID_KERNEL: return "CL_INVALID_KERNEL"; + case CL_INVALID_ARG_INDEX: return "CL_INVALID_ARG_INDEX"; + case CL_INVALID_ARG_VALUE: return "CL_INVALID_ARG_VALUE"; + case CL_INVALID_ARG_SIZE: return "CL_INVALID_ARG_SIZE"; + case CL_INVALID_KERNEL_ARGS: return "CL_INVALID_KERNEL_ARGS"; + case CL_INVALID_WORK_DIMENSION: return "CL_INVALID_WORK_DIMENSION"; + case CL_INVALID_WORK_GROUP_SIZE: return "CL_INVALID_WORK_GROUP_SIZE"; + case CL_INVALID_WORK_ITEM_SIZE: return "CL_INVALID_WORK_ITEM_SIZE"; + case CL_INVALID_GLOBAL_OFFSET: return "CL_INVALID_GLOBAL_OFFSET"; + case CL_INVALID_EVENT_WAIT_LIST: return "CL_INVALID_EVENT_WAIT_LIST"; + case CL_INVALID_EVENT: return "CL_INVALID_EVENT"; + case CL_INVALID_OPERATION: return "CL_INVALID_OPERATION"; + case CL_INVALID_GL_OBJECT: return "CL_INVALID_GL_OBJECT"; + case CL_INVALID_BUFFER_SIZE: return "CL_INVALID_BUFFER_SIZE"; + case CL_INVALID_MIP_LEVEL: return "CL_INVALID_MIP_LEVEL"; + } + return "UNKNOWN_OPENCL_ERROR"; +} + +void gmssl_cl_error_print(const char *file, int line, cl_int err) +{ + fprintf(stderr, "%s: %d: %s\n", file, line, gmssl_cl_error_string(err)); +} + +size_t gmssl_cl_round_up(size_t a, size_t b) +{ + return (a + b - 1)/b*b; +} + +int gmssl_cl_get_gpu_device(cl_platform_id *platform, cl_device_id *device) +{ + cl_uint num_platforms; + cl_platform_id *platforms = NULL; + cl_int err; + cl_uint i; + int ret = -1; + + if (!platform || !device) { + error_print(); + return -1; + } + if ((err = clGetPlatformIDs(0, NULL, &num_platforms)) != CL_SUCCESS) { + cl_error_print(err); + return -1; + } + if (!num_platforms) { + error_print(); + return -1; + } + if (!(platforms = (cl_platform_id *)malloc(sizeof(cl_platform_id) * num_platforms))) { + error_print(); + return -1; + } + if ((err = clGetPlatformIDs(num_platforms, platforms, NULL)) != CL_SUCCESS) { + cl_error_print(err); + goto end; + } + + for (i = 0; i < num_platforms; i++) { + if (clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_GPU, 1, device, NULL) == CL_SUCCESS) { + *platform = platforms[i]; + ret = 1; + goto end; + } + } + error_print(); + +end: + free(platforms); + return ret; +} diff --git a/src/cl.h b/src/cl.h new file mode 100644 index 00000000..8560830f --- /dev/null +++ b/src/cl.h @@ -0,0 +1,35 @@ +/* + * Copyright 2014-2026 The GmSSL Project. All Rights Reserved. + * + * Licensed under the Apache License, Version 2.0 (the License); you may + * not use this file except in compliance with the License. + * + * http://www.apache.org/licenses/LICENSE-2.0 + */ + +#ifndef GMSSL_INTERNAL_CL_H +#define GMSSL_INTERNAL_CL_H + +#include +#include +#include +#ifdef MACOS +#include +#else +#include +#endif + + +const char *gmssl_cl_error_string(cl_int err); +void gmssl_cl_error_print(const char *file, int line, cl_int err); +size_t gmssl_cl_round_up(size_t a, size_t b); +int gmssl_cl_get_gpu_device(cl_platform_id *platform, cl_device_id *device); + +const char *sm3_cl_source(void); + +#define cl_error_print(e) \ + do { gmssl_cl_error_print(__FILE__, __LINE__, (e)); } while (0) + +#define KERNEL(...) #__VA_ARGS__ + +#endif diff --git a/src/lms_cl.c b/src/lms_cl.c index df325111..76d7e8f9 100644 --- a/src/lms_cl.c +++ b/src/lms_cl.c @@ -16,6 +16,7 @@ #include #include #include +#include "cl.h" static const char *lms_cl_src; @@ -24,106 +25,6 @@ static const char *lms_cl_src; #define LMS_CL_DEFAULT_MAX_LEAF_BATCH (size_t)(32768) -static char *clErrorString(cl_int err) -{ - switch (err) { - case CL_SUCCESS: return "CL_SUCCESS"; - case CL_DEVICE_NOT_FOUND: return "CL_DEVICE_NOT_FOUND"; - case CL_DEVICE_NOT_AVAILABLE: return "CL_DEVICE_NOT_AVAILABLE"; - case CL_COMPILER_NOT_AVAILABLE: return "CL_COMPILER_NOT_AVAILABLE"; - case CL_MEM_OBJECT_ALLOCATION_FAILURE: return "CL_MEM_OBJECT_ALLOCATION_FAILURE"; - case CL_OUT_OF_RESOURCES: return "CL_OUT_OF_RESOURCES"; - case CL_OUT_OF_HOST_MEMORY: return "CL_OUT_OF_HOST_MEMORY"; - case CL_PROFILING_INFO_NOT_AVAILABLE: return "CL_PROFILING_INFO_NOT_AVAILABLE"; - case CL_MEM_COPY_OVERLAP: return "CL_MEM_COPY_OVERLAP"; - case CL_IMAGE_FORMAT_MISMATCH: return "CL_IMAGE_FORMAT_MISMATCH"; - case CL_IMAGE_FORMAT_NOT_SUPPORTED: return "CL_IMAGE_FORMAT_NOT_SUPPORTED"; - case CL_BUILD_PROGRAM_FAILURE: return "CL_BUILD_PROGRAM_FAILURE"; - case CL_MAP_FAILURE: return "CL_MAP_FAILURE"; - case CL_INVALID_VALUE: return "CL_INVALID_VALUE"; - case CL_INVALID_DEVICE_TYPE: return "CL_INVALID_DEVICE_TYPE"; - case CL_INVALID_PLATFORM: return "CL_INVALID_PLATFORM"; - case CL_INVALID_DEVICE: return "CL_INVALID_DEVICE"; - case CL_INVALID_CONTEXT: return "CL_INVALID_CONTEXT"; - case CL_INVALID_QUEUE_PROPERTIES: return "CL_INVALID_QUEUE_PROPERTIES"; - case CL_INVALID_COMMAND_QUEUE: return "CL_INVALID_COMMAND_QUEUE"; - case CL_INVALID_HOST_PTR: return "CL_INVALID_HOST_PTR"; - case CL_INVALID_MEM_OBJECT: return "CL_INVALID_MEM_OBJECT"; - case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"; - case CL_INVALID_IMAGE_SIZE: return "CL_INVALID_IMAGE_SIZE"; - case CL_INVALID_SAMPLER: return "CL_INVALID_SAMPLER"; - case CL_INVALID_BINARY: return "CL_INVALID_BINARY"; - case CL_INVALID_BUILD_OPTIONS: return "CL_INVALID_BUILD_OPTIONS"; - case CL_INVALID_PROGRAM: return "CL_INVALID_PROGRAM"; - case CL_INVALID_PROGRAM_EXECUTABLE: return "CL_INVALID_PROGRAM_EXECUTABLE"; - case CL_INVALID_KERNEL_NAME: return "CL_INVALID_KERNEL_NAME"; - case CL_INVALID_KERNEL_DEFINITION: return "CL_INVALID_KERNEL_DEFINITION"; - case CL_INVALID_KERNEL: return "CL_INVALID_KERNEL"; - case CL_INVALID_ARG_INDEX: return "CL_INVALID_ARG_INDEX"; - case CL_INVALID_ARG_VALUE: return "CL_INVALID_ARG_VALUE"; - case CL_INVALID_ARG_SIZE: return "CL_INVALID_ARG_SIZE"; - case CL_INVALID_KERNEL_ARGS: return "CL_INVALID_KERNEL_ARGS"; - case CL_INVALID_WORK_DIMENSION: return "CL_INVALID_WORK_DIMENSION"; - case CL_INVALID_WORK_GROUP_SIZE: return "CL_INVALID_WORK_GROUP_SIZE"; - case CL_INVALID_WORK_ITEM_SIZE: return "CL_INVALID_WORK_ITEM_SIZE"; - case CL_INVALID_GLOBAL_OFFSET: return "CL_INVALID_GLOBAL_OFFSET"; - case CL_INVALID_EVENT_WAIT_LIST: return "CL_INVALID_EVENT_WAIT_LIST"; - case CL_INVALID_EVENT: return "CL_INVALID_EVENT"; - case CL_INVALID_OPERATION: return "CL_INVALID_OPERATION"; - case CL_INVALID_GL_OBJECT: return "CL_INVALID_GL_OBJECT"; - case CL_INVALID_BUFFER_SIZE: return "CL_INVALID_BUFFER_SIZE"; - case CL_INVALID_MIP_LEVEL: return "CL_INVALID_MIP_LEVEL"; - } - return "UNKNOWN_OPENCL_ERROR"; -} - -#define cl_error_print(e) \ - do { fprintf(stderr, "%s: %d: %s\n", __FILE__, __LINE__, clErrorString(e)); } while (0) - -static size_t lms_cl_round_up(size_t a, size_t b) -{ - return (a + b - 1)/b*b; -} - -static int lms_cl_get_device(cl_platform_id *platform, cl_device_id *device) -{ - cl_uint num_platforms; - cl_platform_id *platforms = NULL; - cl_int err; - cl_uint i; - int ret = -1; - - if ((err = clGetPlatformIDs(0, NULL, &num_platforms)) != CL_SUCCESS) { - cl_error_print(err); - return -1; - } - if (!num_platforms) { - error_print(); - return -1; - } - if (!(platforms = (cl_platform_id *)malloc(sizeof(cl_platform_id) * num_platforms))) { - error_print(); - return -1; - } - if ((err = clGetPlatformIDs(num_platforms, platforms, NULL)) != CL_SUCCESS) { - cl_error_print(err); - goto end; - } - - for (i = 0; i < num_platforms; i++) { - if (clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_GPU, 1, device, NULL) == CL_SUCCESS) { - *platform = platforms[i]; - ret = 1; - goto end; - } - } - error_print(); - -end: - free(platforms); - return ret; -} - void lms_cl_cleanup(LMS_CL_CTX *ctx) { if (ctx) { @@ -144,6 +45,7 @@ int lms_cl_init(LMS_CL_CTX *ctx) cl_int err; size_t max_work_group_size; const char *build_opts = NULL; + const char *sources[2]; if (!ctx) { error_print(); @@ -151,7 +53,7 @@ int lms_cl_init(LMS_CL_CTX *ctx) } memset(ctx, 0, sizeof(*ctx)); - if (lms_cl_get_device(&platform, &device) != 1) { + if (gmssl_cl_get_gpu_device(&platform, &device) != 1) { error_print(); return -1; } @@ -163,7 +65,9 @@ int lms_cl_init(LMS_CL_CTX *ctx) cl_error_print(err); goto end; } - if (!(ctx->program = clCreateProgramWithSource(ctx->context, 1, &lms_cl_src, NULL, &err))) { + sources[0] = sm3_cl_source(); + sources[1] = lms_cl_src; + if (!(ctx->program = clCreateProgramWithSource(ctx->context, 2, sources, NULL, &err))) { cl_error_print(err); goto end; } @@ -228,7 +132,7 @@ static int lms_cl_enqueue_leafs_tree(LMS_CL_CTX *ctx, cl_mem mem_seed, cl_mem me } local_work_size = ctx->local_work_size; - global_work_size = lms_cl_round_up(leaf_count, local_work_size); + global_work_size = gmssl_cl_round_up(leaf_count, local_work_size); if ((err = clEnqueueNDRangeKernel(ctx->queue, ctx->leafs_tree_kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL)) != CL_SUCCESS) { cl_error_print(err); @@ -256,7 +160,7 @@ static int lms_cl_enqueue_leafs_compact(LMS_CL_CTX *ctx, cl_mem mem_seed, cl_mem } local_work_size = ctx->local_work_size; - global_work_size = lms_cl_round_up(leaf_count, local_work_size); + global_work_size = gmssl_cl_round_up(leaf_count, local_work_size); if ((err = clEnqueueNDRangeKernel(ctx->queue, ctx->leafs_compact_kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL)) != CL_SUCCESS) { cl_error_print(err); @@ -318,7 +222,7 @@ int lms_cl_derive_merkle_tree(LMS_CL_CTX *ctx, cl_error_print(err); goto end; } - global_work_size = lms_cl_round_up(level_nodes, local_work_size); + global_work_size = gmssl_cl_round_up(level_nodes, local_work_size); if ((err = clEnqueueNDRangeKernel(ctx->queue, ctx->internal_nodes_kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL)) != CL_SUCCESS) { cl_error_print(err); @@ -809,126 +713,8 @@ int hss_cl_sign_init(LMS_CL_CTX *ctx, HSS_SIGN_CTX *sign_ctx, HSS_KEY *key) return 1; } -#define KERNEL(...) #__VA_ARGS__ static const char *lms_cl_src = KERNEL( -__constant uint K[64] = { - 0x79cc4519U, 0xf3988a32U, 0xe7311465U, 0xce6228cbU, - 0x9cc45197U, 0x3988a32fU, 0x7311465eU, 0xe6228cbcU, - 0xcc451979U, 0x988a32f3U, 0x311465e7U, 0x6228cbceU, - 0xc451979cU, 0x88a32f39U, 0x11465e73U, 0x228cbce6U, - 0x9d8a7a87U, 0x3b14f50fU, 0x7629ea1eU, 0xec53d43cU, - 0xd8a7a879U, 0xb14f50f3U, 0x629ea1e7U, 0xc53d43ceU, - 0x8a7a879dU, 0x14f50f3bU, 0x29ea1e76U, 0x53d43cecU, - 0xa7a879d8U, 0x4f50f3b1U, 0x9ea1e762U, 0x3d43cec5U, - 0x7a879d8aU, 0xf50f3b14U, 0xea1e7629U, 0xd43cec53U, - 0xa879d8a7U, 0x50f3b14fU, 0xa1e7629eU, 0x43cec53dU, - 0x879d8a7aU, 0x0f3b14f5U, 0x1e7629eaU, 0x3cec53d4U, - 0x79d8a7a8U, 0xf3b14f50U, 0xe7629ea1U, 0xcec53d43U, - 0x9d8a7a87U, 0x3b14f50fU, 0x7629ea1eU, 0xec53d43cU, - 0xd8a7a879U, 0xb14f50f3U, 0x629ea1e7U, 0xc53d43ceU, - 0x8a7a879dU, 0x14f50f3bU, 0x29ea1e76U, 0x53d43cecU, - 0xa7a879d8U, 0x4f50f3b1U, 0x9ea1e762U, 0x3d43cec5U, -}; - -uint rotl32(uint x, uint n) -{ - return (x << n) | (x >> (32 - n)); -} - -uint P0(uint x) -{ - return x ^ rotl32(x, 9) ^ rotl32(x, 17); -} - -uint P1(uint x) -{ - return x ^ rotl32(x, 15) ^ rotl32(x, 23); -} - -uint FF(uint x, uint y, uint z, uint j) -{ - return j < 16 ? (x ^ y ^ z) : ((x & y) | (x & z) | (y & z)); -} - -uint GG(uint x, uint y, uint z, uint j) -{ - return j < 16 ? (x ^ y ^ z) : (((y ^ z) & x) ^ z); -} - -uint load_be32(__global const uchar *p) -{ - return ((uint)p[0] << 24) | ((uint)p[1] << 16) | ((uint)p[2] << 8) | (uint)p[3]; -} - -uint load_be32_private(__private const uchar *p) -{ - return ((uint)p[0] << 24) | ((uint)p[1] << 16) | ((uint)p[2] << 8) | (uint)p[3]; -} - -void store_be32(__private uchar *p, uint x) -{ - p[0] = (uchar)(x >> 24); - p[1] = (uchar)(x >> 16); - p[2] = (uchar)(x >> 8); - p[3] = (uchar)x; -} - -void sm3_compress_words(__private uint dgst[8], __private uint W[68]) -{ - uint A = dgst[0]; - uint B = dgst[1]; - uint C = dgst[2]; - uint D = dgst[3]; - uint E = dgst[4]; - uint F = dgst[5]; - uint G = dgst[6]; - uint H = dgst[7]; - uint SS1, SS2, TT1, TT2; - uint j; - - for (j = 16; j < 68; j++) { - W[j] = P1(W[j - 16] ^ W[j - 9] ^ rotl32(W[j - 3], 15)) - ^ rotl32(W[j - 13], 7) ^ W[j - 6]; - } - - for (j = 0; j < 64; j++) { - SS1 = rotl32(rotl32(A, 12) + E + K[j], 7); - SS2 = SS1 ^ rotl32(A, 12); - TT1 = FF(A, B, C, j) + D + SS2 + (W[j] ^ W[j + 4]); - TT2 = GG(E, F, G, j) + H + SS1 + W[j]; - D = C; - C = rotl32(B, 9); - B = A; - A = TT1; - H = G; - G = rotl32(F, 19); - F = E; - E = P0(TT2); - } - - dgst[0] ^= A; - dgst[1] ^= B; - dgst[2] ^= C; - dgst[3] ^= D; - dgst[4] ^= E; - dgst[5] ^= F; - dgst[6] ^= G; - dgst[7] ^= H; -} - -void sm3_init_state(__private uint dgst[8]) -{ - dgst[0] = 0x7380166fU; - dgst[1] = 0x4914b2b9U; - dgst[2] = 0x172442d7U; - dgst[3] = 0xda8a0600U; - dgst[4] = 0xa96f30bcU; - dgst[5] = 0x163138aaU; - dgst[6] = 0xe38dee4dU; - dgst[7] = 0xb0fb0e4eU; -} - void sm3_hash_lmots_step(__global const uchar *I, uint q, ushort i, uchar j, __private const uchar in[32], __private uchar out[32]) { @@ -955,56 +741,6 @@ void sm3_hash_lmots_step(__global const uchar *I, uint q, ushort i, uchar j, } } -void sm3_compress_block_bytes(__private uint dgst[8], __private uchar block[64]) -{ - uint W[68]; - uint i; - for (i = 0; i < 16; i++) { - W[i] = load_be32_private(block + 4*i); - } - sm3_compress_words(dgst, W); -} - -void sm3_update_byte(__private uint dgst[8], __private uchar block[64], - __private uint *num, __private ulong *nblocks, uchar b) -{ - block[*num] = b; - *num += 1; - if (*num == 64) { - sm3_compress_block_bytes(dgst, block); - *nblocks += 1; - *num = 0; - } -} - -void sm3_finish_ctx(__private uint dgst[8], __private uchar block[64], - uint num, ulong nblocks, __private uchar out[32]) -{ - ulong bits; - ulong len = nblocks * 64 + num; - uint i; - - block[num++] = 0x80; - if (num > 56) { - while (num < 64) { - block[num++] = 0; - } - sm3_compress_block_bytes(dgst, block); - num = 0; - } - while (num < 56) { - block[num++] = 0; - } - bits = len * 8; - for (i = 0; i < 8; i++) { - block[56 + i] = (uchar)(bits >> (56 - 8*i)); - } - sm3_compress_block_bytes(dgst, block); - for (i = 0; i < 8; i++) { - store_be32(out + 4*i, dgst[i]); - } -} - void sm3_public_hash_update(__private uint dgst[8], __private uchar block[64], __private uint *num, __private ulong *nblocks, __private const uchar z[32]) { diff --git a/src/sm3_cl.c b/src/sm3_cl.c new file mode 100644 index 00000000..22cbf491 --- /dev/null +++ b/src/sm3_cl.c @@ -0,0 +1,211 @@ +/* + * Copyright 2014-2026 The GmSSL Project. All Rights Reserved. + * + * Licensed under the Apache License, Version 2.0 (the License); you may + * not use this file except in compliance with the License. + * + * http://www.apache.org/licenses/LICENSE-2.0 + */ + +#include "cl.h" + + +static const char *sm3_cl_src = KERNEL( + +__constant uint K[64] = { + 0x79cc4519U, 0xf3988a32U, 0xe7311465U, 0xce6228cbU, + 0x9cc45197U, 0x3988a32fU, 0x7311465eU, 0xe6228cbcU, + 0xcc451979U, 0x988a32f3U, 0x311465e7U, 0x6228cbceU, + 0xc451979cU, 0x88a32f39U, 0x11465e73U, 0x228cbce6U, + 0x9d8a7a87U, 0x3b14f50fU, 0x7629ea1eU, 0xec53d43cU, + 0xd8a7a879U, 0xb14f50f3U, 0x629ea1e7U, 0xc53d43ceU, + 0x8a7a879dU, 0x14f50f3bU, 0x29ea1e76U, 0x53d43cecU, + 0xa7a879d8U, 0x4f50f3b1U, 0x9ea1e762U, 0x3d43cec5U, + 0x7a879d8aU, 0xf50f3b14U, 0xea1e7629U, 0xd43cec53U, + 0xa879d8a7U, 0x50f3b14fU, 0xa1e7629eU, 0x43cec53dU, + 0x879d8a7aU, 0x0f3b14f5U, 0x1e7629eaU, 0x3cec53d4U, + 0x79d8a7a8U, 0xf3b14f50U, 0xe7629ea1U, 0xcec53d43U, + 0x9d8a7a87U, 0x3b14f50fU, 0x7629ea1eU, 0xec53d43cU, + 0xd8a7a879U, 0xb14f50f3U, 0x629ea1e7U, 0xc53d43ceU, + 0x8a7a879dU, 0x14f50f3bU, 0x29ea1e76U, 0x53d43cecU, + 0xa7a879d8U, 0x4f50f3b1U, 0x9ea1e762U, 0x3d43cec5U, +}; + +uint rotl32(uint x, uint n) +{ + return (x << n) | (x >> (32 - n)); +} + +uint P0(uint x) +{ + return x ^ rotl32(x, 9) ^ rotl32(x, 17); +} + +uint P1(uint x) +{ + return x ^ rotl32(x, 15) ^ rotl32(x, 23); +} + +uint FF(uint x, uint y, uint z, uint j) +{ + return j < 16 ? (x ^ y ^ z) : ((x & y) | (x & z) | (y & z)); +} + +uint GG(uint x, uint y, uint z, uint j) +{ + return j < 16 ? (x ^ y ^ z) : (((y ^ z) & x) ^ z); +} + +uint load_be32(__global const uchar *p) +{ + return ((uint)p[0] << 24) | ((uint)p[1] << 16) | ((uint)p[2] << 8) | (uint)p[3]; +} + +uint load_be32_private(__private const uchar *p) +{ + return ((uint)p[0] << 24) | ((uint)p[1] << 16) | ((uint)p[2] << 8) | (uint)p[3]; +} + +void store_be32(__private uchar *p, uint x) +{ + p[0] = (uchar)(x >> 24); + p[1] = (uchar)(x >> 16); + p[2] = (uchar)(x >> 8); + p[3] = (uchar)x; +} + +void set_be32(__private uchar *p, uint x) +{ + p[0] = (uchar)(x >> 24); + p[1] = (uchar)(x >> 16); + p[2] = (uchar)(x >> 8); + p[3] = (uchar)x; +} + +void copy32_from_global(__private uchar dst[32], __global const uchar *src) +{ + uint i; + for (i = 0; i < 32; i++) { + dst[i] = src[i]; + } +} + +void copy32(__private uchar dst[32], __private const uchar src[32]) +{ + uint i; + for (i = 0; i < 32; i++) { + dst[i] = src[i]; + } +} + +void sm3_compress_words(__private uint dgst[8], __private uint W[68]) +{ + uint A = dgst[0]; + uint B = dgst[1]; + uint C = dgst[2]; + uint D = dgst[3]; + uint E = dgst[4]; + uint F = dgst[5]; + uint G = dgst[6]; + uint H = dgst[7]; + uint SS1, SS2, TT1, TT2; + uint j; + + for (j = 16; j < 68; j++) { + W[j] = P1(W[j - 16] ^ W[j - 9] ^ rotl32(W[j - 3], 15)) + ^ rotl32(W[j - 13], 7) ^ W[j - 6]; + } + + for (j = 0; j < 64; j++) { + SS1 = rotl32(rotl32(A, 12) + E + K[j], 7); + SS2 = SS1 ^ rotl32(A, 12); + TT1 = FF(A, B, C, j) + D + SS2 + (W[j] ^ W[j + 4]); + TT2 = GG(E, F, G, j) + H + SS1 + W[j]; + D = C; + C = rotl32(B, 9); + B = A; + A = TT1; + H = G; + G = rotl32(F, 19); + F = E; + E = P0(TT2); + } + + dgst[0] ^= A; + dgst[1] ^= B; + dgst[2] ^= C; + dgst[3] ^= D; + dgst[4] ^= E; + dgst[5] ^= F; + dgst[6] ^= G; + dgst[7] ^= H; +} + +void sm3_init_state(__private uint dgst[8]) +{ + dgst[0] = 0x7380166fU; + dgst[1] = 0x4914b2b9U; + dgst[2] = 0x172442d7U; + dgst[3] = 0xda8a0600U; + dgst[4] = 0xa96f30bcU; + dgst[5] = 0x163138aaU; + dgst[6] = 0xe38dee4dU; + dgst[7] = 0xb0fb0e4eU; +} + +void sm3_compress_blocks(__private uint dgst[8], __private uchar block[64]) +{ + uint W[68]; + uint i; + for (i = 0; i < 16; i++) { + W[i] = load_be32_private(block + 4*i); + } + sm3_compress_words(dgst, W); +} + +void sm3_update_byte(__private uint dgst[8], __private uchar block[64], + __private uint *num, __private ulong *nblocks, uchar b) +{ + block[*num] = b; + *num += 1; + if (*num == 64) { + sm3_compress_blocks(dgst, block); + *nblocks += 1; + *num = 0; + } +} + +void sm3_finish_ctx(__private uint dgst[8], __private uchar block[64], + uint num, ulong nblocks, __private uchar out[32]) +{ + ulong bits; + ulong len = nblocks * 64 + num; + uint i; + + block[num++] = 0x80; + if (num > 56) { + while (num < 64) { + block[num++] = 0; + } + sm3_compress_blocks(dgst, block); + num = 0; + } + while (num < 56) { + block[num++] = 0; + } + bits = len * 8; + for (i = 0; i < 8; i++) { + block[56 + i] = (uchar)(bits >> (56 - 8*i)); + } + sm3_compress_blocks(dgst, block); + for (i = 0; i < 8; i++) { + store_be32(out + 4*i, dgst[i]); + } +} + +); + +const char *sm3_cl_source(void) +{ + return sm3_cl_src; +} diff --git a/src/sm4_cl.c b/src/sm4_cl.c index 300b170d..a5ffe429 100644 --- a/src/sm4_cl.c +++ b/src/sm4_cl.c @@ -16,68 +16,12 @@ #include #include #include +#include "cl.h" -static char *clErrorString(cl_uint err) -{ - switch (err) { - case CL_SUCCESS: return "CL_SUCCESS!"; - case CL_DEVICE_NOT_FOUND: return "CL_DEVICE_NOT_FOUND"; - case CL_DEVICE_NOT_AVAILABLE: return "CL_DEVICE_NOT_AVAILABLE"; - case CL_COMPILER_NOT_AVAILABLE: return "CL_COMPILER_NOT_AVAILABLE"; - case CL_MEM_OBJECT_ALLOCATION_FAILURE: return "CL_MEM_OBJECT_ALLOCATION_FAILURE"; - case CL_OUT_OF_RESOURCES: return "CL_OUT_OF_RESOURCES"; - case CL_OUT_OF_HOST_MEMORY: return "CL_OUT_OF_HOST_MEMORY"; - case CL_PROFILING_INFO_NOT_AVAILABLE: return "CL_PROFILING_INFO_NOT_AVAILABLE"; - case CL_MEM_COPY_OVERLAP: return "CL_MEM_COPY_OVERLAP"; - case CL_IMAGE_FORMAT_MISMATCH: return "CL_IMAGE_FORMAT_MISMATCH"; - case CL_IMAGE_FORMAT_NOT_SUPPORTED: return "CL_IMAGE_FORMAT_NOT_SUPPORTED"; - case CL_BUILD_PROGRAM_FAILURE: return "CL_BUILD_PROGRAM_FAILURE"; - case CL_MAP_FAILURE: return "CL_MAP_FAILURE"; - case CL_INVALID_VALUE: return "CL_INVALID_VALUE"; - case CL_INVALID_DEVICE_TYPE: return "CL_INVALID_DEVICE_TYPE"; - case CL_INVALID_PLATFORM: return "CL_INVALID_PLATFORM"; - case CL_INVALID_DEVICE: return "CL_INVALID_DEVICE"; - case CL_INVALID_CONTEXT: return "CL_INVALID_CONTEXT"; - case CL_INVALID_QUEUE_PROPERTIES: return "CL_INVALID_QUEUE_PROPERTIES"; - case CL_INVALID_COMMAND_QUEUE: return "CL_INVALID_COMMAND_QUEUE"; - case CL_INVALID_HOST_PTR: return "CL_INVALID_HOST_PTR"; - case CL_INVALID_MEM_OBJECT: return "CL_INVALID_MEM_OBJECT"; - case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"; - case CL_INVALID_IMAGE_SIZE: return "CL_INVALID_IMAGE_SIZE"; - case CL_INVALID_SAMPLER: return "CL_INVALID_SAMPLER"; - case CL_INVALID_BINARY: return "CL_INVALID_BINARY"; - case CL_INVALID_BUILD_OPTIONS: return "CL_INVALID_BUILD_OPTIONS"; - case CL_INVALID_PROGRAM: return "CL_INVALID_PROGRAM"; - case CL_INVALID_PROGRAM_EXECUTABLE: return "CL_INVALID_PROGRAM_EXECUTABLE"; - case CL_INVALID_KERNEL_NAME: return "CL_INVALID_KERNEL_NAME"; - case CL_INVALID_KERNEL_DEFINITION: return "CL_INVALID_KERNEL_DEFINITION"; - case CL_INVALID_KERNEL: return "CL_INVALID_KERNEL"; - case CL_INVALID_ARG_INDEX: return "CL_INVALID_ARG_INDEX"; - case CL_INVALID_ARG_VALUE: return "CL_INVALID_ARG_VALUE"; - case CL_INVALID_ARG_SIZE: return "CL_INVALID_ARG_SIZE"; - case CL_INVALID_KERNEL_ARGS: return "CL_INVALID_KERNEL_ARGS"; - case CL_INVALID_WORK_DIMENSION: return "CL_INVALID_WORK_DIMENSION"; - case CL_INVALID_WORK_GROUP_SIZE: return "CL_INVALID_WORK_GROUP_SIZE"; - case CL_INVALID_WORK_ITEM_SIZE: return "CL_INVALID_WORK_ITEM_SIZE"; - case CL_INVALID_GLOBAL_OFFSET: return "CL_INVALID_GLOBAL_OFFSET"; - case CL_INVALID_EVENT_WAIT_LIST: return "CL_INVALID_EVENT_WAIT_LIST"; - case CL_INVALID_EVENT: return "CL_INVALID_EVENT"; - case CL_INVALID_OPERATION: return "CL_INVALID_OPERATION"; - case CL_INVALID_GL_OBJECT: return "CL_INVALID_GL_OBJECT"; - case CL_INVALID_BUFFER_SIZE: return "CL_INVALID_BUFFER_SIZE"; - case CL_INVALID_MIP_LEVEL: return "CL_INVALID_MIP_LEVEL"; - } - return NULL; -} - static const char *sm4_cl_src; -#define cl_error_print(e) \ - do { fprintf(stderr, "%s: %d: %s()\n",__FILE__,__LINE__,clErrorString(e)); } while (0) - - void sm4_cl_cleanup(SM4_CL_CTX *ctx) { if (ctx) { @@ -151,7 +95,6 @@ static int sm4_cl_set_key(SM4_CL_CTX *ctx, const uint8_t key[16], int enc) { cl_platform_id platform; cl_device_id device; - cl_uint device_cnt; cl_int err; char sval[256]; size_t slen; @@ -160,12 +103,8 @@ static int sm4_cl_set_key(SM4_CL_CTX *ctx, const uint8_t key[16], int enc) memset(ctx, 0, sizeof(*ctx)); - if ((err = clGetPlatformIDs(1, &platform, NULL)) != CL_SUCCESS) { - cl_error_print(err); - return -1; - } - if ((err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, &device_cnt)) != CL_SUCCESS) { - cl_error_print(err); + if (gmssl_cl_get_gpu_device(&platform, &device) != 1) { + error_print(); return -1; } //clPrintDeviceInfo(device); @@ -188,19 +127,19 @@ static int sm4_cl_set_key(SM4_CL_CTX *ctx, const uint8_t key[16], int enc) cl_error_print(err); - // FIXME: sizeof(log) 错误？ - if ((err = clGetProgramBuildInfo(ctx->program, device, CL_PROGRAM_BUILD_LOG, sizeof(log), NULL, &loglen)) != CL_SUCCESS) { + if ((err = clGetProgramBuildInfo(ctx->program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &loglen)) != CL_SUCCESS) { cl_error_print(err); goto end; } - if (!(log = (char *)malloc(loglen))) { + if (!(log = (char *)malloc(loglen + 1))) { goto end; } - if ((err = clGetProgramBuildInfo(ctx->program, device, CL_PROGRAM_BUILD_LOG, sizeof(log), NULL, &loglen)) != CL_SUCCESS) { + if ((err = clGetProgramBuildInfo(ctx->program, device, CL_PROGRAM_BUILD_LOG, loglen, log, NULL)) != CL_SUCCESS) { cl_error_print(err); free(log); goto end; } + log[loglen] = 0; fprintf(stderr, "%s %d: %s\n", __FILE__, __LINE__, log); free(log); goto end; @@ -316,7 +255,6 @@ end: return ret; } -#define KERNEL(...) #__VA_ARGS__ static const char *sm4_cl_src = KERNEL( __constant unsigned char SBOX[256] = { diff --git a/src/xmss_cl.c b/src/xmss_cl.c index 9922a282..105f8541 100644 --- a/src/xmss_cl.c +++ b/src/xmss_cl.c @@ -16,6 +16,7 @@ #include #include #include +#include "cl.h" static const char *xmss_cl_src; @@ -36,68 +37,6 @@ static const uint8_t xmss_cl_sm3_digest_three[] = { 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, }; - -static char *clErrorString(cl_int err) -{ - switch (err) { - case CL_SUCCESS: return "CL_SUCCESS"; - case CL_DEVICE_NOT_FOUND: return "CL_DEVICE_NOT_FOUND"; - case CL_DEVICE_NOT_AVAILABLE: return "CL_DEVICE_NOT_AVAILABLE"; - case CL_COMPILER_NOT_AVAILABLE: return "CL_COMPILER_NOT_AVAILABLE"; - case CL_MEM_OBJECT_ALLOCATION_FAILURE: return "CL_MEM_OBJECT_ALLOCATION_FAILURE"; - case CL_OUT_OF_RESOURCES: return "CL_OUT_OF_RESOURCES"; - case CL_OUT_OF_HOST_MEMORY: return "CL_OUT_OF_HOST_MEMORY"; - case CL_PROFILING_INFO_NOT_AVAILABLE: return "CL_PROFILING_INFO_NOT_AVAILABLE"; - case CL_MEM_COPY_OVERLAP: return "CL_MEM_COPY_OVERLAP"; - case CL_IMAGE_FORMAT_MISMATCH: return "CL_IMAGE_FORMAT_MISMATCH"; - case CL_IMAGE_FORMAT_NOT_SUPPORTED: return "CL_IMAGE_FORMAT_NOT_SUPPORTED"; - case CL_BUILD_PROGRAM_FAILURE: return "CL_BUILD_PROGRAM_FAILURE"; - case CL_MAP_FAILURE: return "CL_MAP_FAILURE"; - case CL_INVALID_VALUE: return "CL_INVALID_VALUE"; - case CL_INVALID_DEVICE_TYPE: return "CL_INVALID_DEVICE_TYPE"; - case CL_INVALID_PLATFORM: return "CL_INVALID_PLATFORM"; - case CL_INVALID_DEVICE: return "CL_INVALID_DEVICE"; - case CL_INVALID_CONTEXT: return "CL_INVALID_CONTEXT"; - case CL_INVALID_QUEUE_PROPERTIES: return "CL_INVALID_QUEUE_PROPERTIES"; - case CL_INVALID_COMMAND_QUEUE: return "CL_INVALID_COMMAND_QUEUE"; - case CL_INVALID_HOST_PTR: return "CL_INVALID_HOST_PTR"; - case CL_INVALID_MEM_OBJECT: return "CL_INVALID_MEM_OBJECT"; - case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"; - case CL_INVALID_IMAGE_SIZE: return "CL_INVALID_IMAGE_SIZE"; - case CL_INVALID_SAMPLER: return "CL_INVALID_SAMPLER"; - case CL_INVALID_BINARY: return "CL_INVALID_BINARY"; - case CL_INVALID_BUILD_OPTIONS: return "CL_INVALID_BUILD_OPTIONS"; - case CL_INVALID_PROGRAM: return "CL_INVALID_PROGRAM"; - case CL_INVALID_PROGRAM_EXECUTABLE: return "CL_INVALID_PROGRAM_EXECUTABLE"; - case CL_INVALID_KERNEL_NAME: return "CL_INVALID_KERNEL_NAME"; - case CL_INVALID_KERNEL_DEFINITION: return "CL_INVALID_KERNEL_DEFINITION"; - case CL_INVALID_KERNEL: return "CL_INVALID_KERNEL"; - case CL_INVALID_ARG_INDEX: return "CL_INVALID_ARG_INDEX"; - case CL_INVALID_ARG_VALUE: return "CL_INVALID_ARG_VALUE"; - case CL_INVALID_ARG_SIZE: return "CL_INVALID_ARG_SIZE"; - case CL_INVALID_KERNEL_ARGS: return "CL_INVALID_KERNEL_ARGS"; - case CL_INVALID_WORK_DIMENSION: return "CL_INVALID_WORK_DIMENSION"; - case CL_INVALID_WORK_GROUP_SIZE: return "CL_INVALID_WORK_GROUP_SIZE"; - case CL_INVALID_WORK_ITEM_SIZE: return "CL_INVALID_WORK_ITEM_SIZE"; - case CL_INVALID_GLOBAL_OFFSET: return "CL_INVALID_GLOBAL_OFFSET"; - case CL_INVALID_EVENT_WAIT_LIST: return "CL_INVALID_EVENT_WAIT_LIST"; - case CL_INVALID_EVENT: return "CL_INVALID_EVENT"; - case CL_INVALID_OPERATION: return "CL_INVALID_OPERATION"; - case CL_INVALID_GL_OBJECT: return "CL_INVALID_GL_OBJECT"; - case CL_INVALID_BUFFER_SIZE: return "CL_INVALID_BUFFER_SIZE"; - case CL_INVALID_MIP_LEVEL: return "CL_INVALID_MIP_LEVEL"; - } - return "UNKNOWN_OPENCL_ERROR"; -} - -#define cl_error_print(e) \ - do { fprintf(stderr, "%s: %d: %s\n", __FILE__, __LINE__, clErrorString(e)); } while (0) - -static size_t xmss_cl_round_up(size_t a, size_t b) -{ - return (a + b - 1)/b*b; -} - static size_t xmss_cl_tree_root_offset(size_t height) { return ((size_t)1 << (height + 1)) - 2; @@ -113,44 +52,6 @@ static uint64_t xmssmt_cl_tree_index(uint64_t index, size_t height, size_t layer return (index >> ((height/layers) * layer)) % ((uint64_t)1 << (height/layers)); } -static int xmss_cl_get_device(cl_platform_id *platform, cl_device_id *device) -{ - cl_uint num_platforms; - cl_platform_id *platforms = NULL; - cl_int err; - cl_uint i; - int ret = -1; - - if ((err = clGetPlatformIDs(0, NULL, &num_platforms)) != CL_SUCCESS) { - cl_error_print(err); - return -1; - } - if (!num_platforms) { - error_print(); - return -1; - } - if (!(platforms = (cl_platform_id *)malloc(sizeof(cl_platform_id) * num_platforms))) { - error_print(); - return -1; - } - if ((err = clGetPlatformIDs(num_platforms, platforms, NULL)) != CL_SUCCESS) { - cl_error_print(err); - goto end; - } - - for (i = 0; i < num_platforms; i++) { - if (clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_GPU, 1, device, NULL) == CL_SUCCESS) { - *platform = platforms[i]; - ret = 1; - goto end; - } - } - error_print(); - -end: - free(platforms); - return ret; -} void xmss_cl_cleanup(XMSS_CL_CTX *ctx) { @@ -170,6 +71,7 @@ int xmss_cl_init(XMSS_CL_CTX *ctx) cl_device_id device; cl_int err; size_t max_work_group_size; + const char *sources[2]; if (!ctx) { error_print(); @@ -177,7 +79,7 @@ int xmss_cl_init(XMSS_CL_CTX *ctx) } memset(ctx, 0, sizeof(*ctx)); - if (xmss_cl_get_device(&platform, &device) != 1) { + if (gmssl_cl_get_gpu_device(&platform, &device) != 1) { error_print(); return -1; } @@ -189,7 +91,9 @@ int xmss_cl_init(XMSS_CL_CTX *ctx) cl_error_print(err); goto end; } - if (!(ctx->program = clCreateProgramWithSource(ctx->context, 1, &xmss_cl_src, NULL, &err))) { + sources[0] = sm3_cl_source(); + sources[1] = xmss_cl_src; + if (!(ctx->program = clCreateProgramWithSource(ctx->context, 2, sources, NULL, &err))) { cl_error_print(err); goto end; } @@ -278,7 +182,7 @@ int xmss_cl_build_tree(XMSS_CL_CTX *ctx, u32_height = (uint32_t)height; local_work_size = ctx->local_work_size; - global_work_size = xmss_cl_round_up(n, local_work_size); + global_work_size = gmssl_cl_round_up(n, local_work_size); { cl_uint arg = 0; uint32_t leaf_count = (uint32_t)n; @@ -319,7 +223,7 @@ int xmss_cl_build_tree(XMSS_CL_CTX *ctx, cl_error_print(err); goto end; } - global_work_size = xmss_cl_round_up(level_nodes, local_work_size); + global_work_size = gmssl_cl_round_up(level_nodes, local_work_size); if ((err = clEnqueueNDRangeKernel(ctx->queue, ctx->internal_nodes_kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL)) != CL_SUCCESS) { cl_error_print(err); @@ -702,166 +606,8 @@ int xmssmt_cl_sign_init(XMSS_CL_CTX *cl_ctx, XMSSMT_SIGN_CTX *sign_ctx, XMSSMT_K return 1; } -#define KERNEL(...) #__VA_ARGS__ static const char *xmss_cl_src = KERNEL( -__constant uint K[64] = { - 0x79cc4519U, 0xf3988a32U, 0xe7311465U, 0xce6228cbU, - 0x9cc45197U, 0x3988a32fU, 0x7311465eU, 0xe6228cbcU, - 0xcc451979U, 0x988a32f3U, 0x311465e7U, 0x6228cbceU, - 0xc451979cU, 0x88a32f39U, 0x11465e73U, 0x228cbce6U, - 0x9d8a7a87U, 0x3b14f50fU, 0x7629ea1eU, 0xec53d43cU, - 0xd8a7a879U, 0xb14f50f3U, 0x629ea1e7U, 0xc53d43ceU, - 0x8a7a879dU, 0x14f50f3bU, 0x29ea1e76U, 0x53d43cecU, - 0xa7a879d8U, 0x4f50f3b1U, 0x9ea1e762U, 0x3d43cec5U, - 0x7a879d8aU, 0xf50f3b14U, 0xea1e7629U, 0xd43cec53U, - 0xa879d8a7U, 0x50f3b14fU, 0xa1e7629eU, 0x43cec53dU, - 0x879d8a7aU, 0x0f3b14f5U, 0x1e7629eaU, 0x3cec53d4U, - 0x79d8a7a8U, 0xf3b14f50U, 0xe7629ea1U, 0xcec53d43U, - 0x9d8a7a87U, 0x3b14f50fU, 0x7629ea1eU, 0xec53d43cU, - 0xd8a7a879U, 0xb14f50f3U, 0x629ea1e7U, 0xc53d43ceU, - 0x8a7a879dU, 0x14f50f3bU, 0x29ea1e76U, 0x53d43cecU, - 0xa7a879d8U, 0x4f50f3b1U, 0x9ea1e762U, 0x3d43cec5U, -}; - -uint rotl32(uint x, uint n) { return (x << n) | (x >> (32 - n)); } -uint P0(uint x) { return x ^ rotl32(x, 9) ^ rotl32(x, 17); } -uint P1(uint x) { return x ^ rotl32(x, 15) ^ rotl32(x, 23); } -uint FF(uint x, uint y, uint z, uint j) { return j < 16 ? (x ^ y ^ z) : ((x & y) | (x & z) | (y & z)); } -uint GG(uint x, uint y, uint z, uint j) { return j < 16 ? (x ^ y ^ z) : (((y ^ z) & x) ^ z); } - -uint load_be32_global(__global const uchar *p) -{ - return ((uint)p[0] << 24) | ((uint)p[1] << 16) | ((uint)p[2] << 8) | (uint)p[3]; -} - -uint load_be32_private(__private const uchar *p) -{ - return ((uint)p[0] << 24) | ((uint)p[1] << 16) | ((uint)p[2] << 8) | (uint)p[3]; -} - -void store_be32(__private uchar *p, uint x) -{ - p[0] = (uchar)(x >> 24); - p[1] = (uchar)(x >> 16); - p[2] = (uchar)(x >> 8); - p[3] = (uchar)x; -} - -void set_be32(__private uchar *p, uint x) -{ - p[0] = (uchar)(x >> 24); - p[1] = (uchar)(x >> 16); - p[2] = (uchar)(x >> 8); - p[3] = (uchar)x; -} - -void copy32_from_global(__private uchar dst[32], __global const uchar *src) -{ - uint i; - for (i = 0; i < 32; i++) dst[i] = src[i]; -} - -void copy32(__private uchar dst[32], __private const uchar src[32]) -{ - uint i; - for (i = 0; i < 32; i++) dst[i] = src[i]; -} - -void sm3_compress_words(__private uint dgst[8], __private uint W[68]) -{ - uint A = dgst[0]; - uint B = dgst[1]; - uint C = dgst[2]; - uint D = dgst[3]; - uint E = dgst[4]; - uint F = dgst[5]; - uint G = dgst[6]; - uint H = dgst[7]; - uint SS1, SS2, TT1, TT2; - uint j; - - for (j = 16; j < 68; j++) { - W[j] = P1(W[j - 16] ^ W[j - 9] ^ rotl32(W[j - 3], 15)) - ^ rotl32(W[j - 13], 7) ^ W[j - 6]; - } - for (j = 0; j < 64; j++) { - SS1 = rotl32(rotl32(A, 12) + E + K[j], 7); - SS2 = SS1 ^ rotl32(A, 12); - TT1 = FF(A, B, C, j) + D + SS2 + (W[j] ^ W[j + 4]); - TT2 = GG(E, F, G, j) + H + SS1 + W[j]; - D = C; - C = rotl32(B, 9); - B = A; - A = TT1; - H = G; - G = rotl32(F, 19); - F = E; - E = P0(TT2); - } - - dgst[0] ^= A; - dgst[1] ^= B; - dgst[2] ^= C; - dgst[3] ^= D; - dgst[4] ^= E; - dgst[5] ^= F; - dgst[6] ^= G; - dgst[7] ^= H; -} - -void sm3_init_state(__private uint dgst[8]) -{ - dgst[0] = 0x7380166fU; - dgst[1] = 0x4914b2b9U; - dgst[2] = 0x172442d7U; - dgst[3] = 0xda8a0600U; - dgst[4] = 0xa96f30bcU; - dgst[5] = 0x163138aaU; - dgst[6] = 0xe38dee4dU; - dgst[7] = 0xb0fb0e4eU; -} - -void sm3_compress_block_bytes(__private uint dgst[8], __private uchar block[64]) -{ - uint W[68]; - uint i; - for (i = 0; i < 16; i++) W[i] = load_be32_private(block + 4*i); - sm3_compress_words(dgst, W); -} - -void sm3_update_byte(__private uint dgst[8], __private uchar block[64], - __private uint *num, __private ulong *nblocks, uchar b) -{ - block[*num] = b; - *num += 1; - if (*num == 64) { - sm3_compress_block_bytes(dgst, block); - *nblocks += 1; - *num = 0; - } -} - -void sm3_finish_ctx(__private uint dgst[8], __private uchar block[64], - uint num, ulong nblocks, __private uchar out[32]) -{ - ulong bits; - ulong len = nblocks * 64 + num; - uint i; - - block[num++] = 0x80; - if (num > 56) { - while (num < 64) block[num++] = 0; - sm3_compress_block_bytes(dgst, block); - num = 0; - } - while (num < 56) block[num++] = 0; - bits = len * 8; - for (i = 0; i < 8; i++) block[56 + i] = (uchar)(bits >> (56 - 8*i)); - sm3_compress_block_bytes(dgst, block); - for (i = 0; i < 8; i++) store_be32(out + 4*i, dgst[i]); -} - void sm3_hash_parts_3(__private const uchar a[32], __private const uchar b[32], __private const uchar c[32], __private uchar out[32]) {