From e7c0a28f9b5dcb3684b22f7753da9f0990fa46a9 Mon Sep 17 00:00:00 2001 From: Zhi Guan Date: Wed, 24 Jun 2026 10:23:18 +0800 Subject: [PATCH] Add LMS/HSS CL --- CMakeLists.txt | 30 +- include/gmssl/lms_cl.h | 63 +++ include/gmssl/version.h | 2 +- src/lms_cl.c | 1168 +++++++++++++++++++++++++++++++++++++++ src/tlcp.c | 38 +- tests/lms_cltest.c | 289 ++++++++++ tools/hsskeygen.c | 25 + tools/hsssign.c | 32 ++ tools/lmskeygen.c | 25 + tools/lmssign.c | 25 + 10 files changed, 1677 insertions(+), 20 deletions(-) create mode 100644 include/gmssl/lms_cl.h create mode 100644 src/lms_cl.c create mode 100644 tests/lms_cltest.c diff --git a/CMakeLists.txt b/CMakeLists.txt index 8a69ac7f..5e517797 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -112,6 +112,7 @@ option(ENABLE_SM3_SSE "Enable SM3 SSE assembly implementation" ${GMSSL_DEFAULT_E option(ENABLE_SM4_CTR_AESNI_AVX "Enable SM4 CTR AESNI+AVX assembly implementation" OFF) option(ENABLE_SM4_CL "Enable SM4 OpenCL" OFF) +option(ENABLE_LMS_CL "Enable LMS OpenCL" OFF) option(ENABLE_INTEL_RDRAND "Enable Intel RDRAND instructions" OFF) @@ -538,6 +539,19 @@ if (ENABLE_LMS) list(APPEND tests lms) endif() +if (ENABLE_LMS_CL) + if (NOT ENABLE_LMS) + message(FATAL_ERROR "ENABLE_LMS_CL requires ENABLE_LMS") + endif() + message(STATUS "ENABLE_LMS_CL is ON") + add_definitions(-DENABLE_LMS_CL) + if (CMAKE_SYSTEM_NAME STREQUAL "Darwin") + add_definitions(-DMACOS) # to include + endif() + list(APPEND src src/lms_cl.c) + list(APPEND tests lms_cl) +endif() + if (ENABLE_XMSS) message(STATUS "ENABLE_XMSS is ON") @@ -761,10 +775,6 @@ elseif (APPLE) target_link_libraries(gmssl dl) endif() target_link_libraries(gmssl "-framework Security") - if (ENABLE_SM4_CL) - # FIXME: different rules for cl and OpenCL framework - target_link_libraries(gmssl "-framework OpenCL") - endif() #target_link_libraries(gmssl "-framework CoreFoundation") # rand_apple.c CFRelease() elseif (MINGW) target_link_libraries(gmssl PRIVATE wsock32) @@ -774,6 +784,16 @@ else() endif() endif() +if (ENABLE_SM4_CL OR ENABLE_LMS_CL) + if (APPLE) + target_link_libraries(gmssl "-framework OpenCL") + else() + find_package(OpenCL REQUIRED) + target_include_directories(gmssl PRIVATE ${OpenCL_INCLUDE_DIRS}) + target_link_libraries(gmssl ${OpenCL_LIBRARIES}) + endif() +endif() + set(CMAKE_INSTALL_RPATH "${CMAKE_INSTALL_PREFIX}/lib") @@ -947,7 +967,7 @@ endif() # set(CPACK_PACKAGE_NAME "GmSSL") set(CPACK_PACKAGE_VENDOR "GmSSL develop team") -set(CPACK_PACKAGE_VERSION "3.3.0-dev.1163") +set(CPACK_PACKAGE_VERSION "3.3.0-dev.1164") set(CPACK_PACKAGE_DESCRIPTION_FILE ${PROJECT_SOURCE_DIR}/README.md) set(CPACK_NSIS_MODIFY_PATH ON) include(CPack) diff --git a/include/gmssl/lms_cl.h b/include/gmssl/lms_cl.h new file mode 100644 index 00000000..07fac267 --- /dev/null +++ b/include/gmssl/lms_cl.h @@ -0,0 +1,63 @@ +/* + * 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_LMS_CL_H +#define GMSSL_LMS_CL_H + +#include +#include +#include +#include +#include +#ifdef MACOS +#include +#else +#include +#endif + +#ifdef __cplusplus +extern "C" { +#endif + + +typedef struct { + cl_context context; + cl_command_queue queue; + cl_program program; + cl_kernel leafs_tree_kernel; + cl_kernel leafs_compact_kernel; + cl_kernel internal_nodes_kernel; + size_t local_work_size; + size_t max_leaf_batch; +} LMS_CL_CTX; + + +int lms_cl_init(LMS_CL_CTX *ctx); +void lms_cl_cleanup(LMS_CL_CTX *ctx); + +int lms_cl_derive_merkle_tree(LMS_CL_CTX *ctx, + const lms_sm3_digest_t seed, const uint8_t I[16], int height, lms_sm3_digest_t *tree); +int lms_cl_derive_merkle_root(LMS_CL_CTX *ctx, + const lms_sm3_digest_t seed, const uint8_t I[16], int height, lms_sm3_digest_t root); + +int lms_cl_key_generate_ex(LMS_CL_CTX *ctx, LMS_KEY *key, int lms_type, + const lms_sm3_digest_t seed, const uint8_t I[16], int cache_tree); +int lms_cl_key_generate(LMS_CL_CTX *ctx, LMS_KEY *key, int lms_type); +int lms_cl_private_key_from_bytes(LMS_CL_CTX *ctx, LMS_KEY *key, const uint8_t **in, size_t *inlen); + +int hss_cl_key_generate(LMS_CL_CTX *ctx, HSS_KEY *key, const int *lms_types, size_t levels); +int hss_cl_private_key_from_bytes(LMS_CL_CTX *ctx, HSS_KEY *key, const uint8_t **in, size_t *inlen); +int hss_cl_key_update(LMS_CL_CTX *ctx, HSS_KEY *key); +int hss_cl_sign_init(LMS_CL_CTX *ctx, HSS_SIGN_CTX *sign_ctx, HSS_KEY *key); + + +#ifdef __cplusplus +} +#endif +#endif diff --git a/include/gmssl/version.h b/include/gmssl/version.h index 42009989..de56287f 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.1163" +#define GMSSL_VERSION_STR "GmSSL 3.3.0-dev.1164" int gmssl_version_num(void); const char *gmssl_version_str(void); diff --git a/src/lms_cl.c b/src/lms_cl.c new file mode 100644 index 00000000..df325111 --- /dev/null +++ b/src/lms_cl.c @@ -0,0 +1,1168 @@ +/* + * 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 +#include +#include +#include +#include +#include +#include +#include +#include + + +static const char *lms_cl_src; + +#define LMS_CL_DEFAULT_LOCAL_WORK_SIZE 64 +#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) { + if (ctx->internal_nodes_kernel) clReleaseKernel(ctx->internal_nodes_kernel); + if (ctx->leafs_compact_kernel) clReleaseKernel(ctx->leafs_compact_kernel); + if (ctx->leafs_tree_kernel) clReleaseKernel(ctx->leafs_tree_kernel); + if (ctx->program) clReleaseProgram(ctx->program); + if (ctx->queue) clReleaseCommandQueue(ctx->queue); + if (ctx->context) clReleaseContext(ctx->context); + memset(ctx, 0, sizeof(*ctx)); + } +} + +int lms_cl_init(LMS_CL_CTX *ctx) +{ + cl_platform_id platform; + cl_device_id device; + cl_int err; + size_t max_work_group_size; + const char *build_opts = NULL; + + if (!ctx) { + error_print(); + return -1; + } + memset(ctx, 0, sizeof(*ctx)); + + if (lms_cl_get_device(&platform, &device) != 1) { + error_print(); + return -1; + } + if (!(ctx->context = clCreateContext(NULL, 1, &device, NULL, NULL, &err))) { + cl_error_print(err); + goto end; + } + if (!(ctx->queue = clCreateCommandQueue(ctx->context, device, 0, &err))) { + cl_error_print(err); + goto end; + } + if (!(ctx->program = clCreateProgramWithSource(ctx->context, 1, &lms_cl_src, NULL, &err))) { + cl_error_print(err); + goto end; + } + if ((err = clBuildProgram(ctx->program, 1, &device, build_opts, NULL, NULL)) != CL_SUCCESS) { + char *log = NULL; + size_t loglen = 0; + + cl_error_print(err); + (void)clGetProgramBuildInfo(ctx->program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &loglen); + if (loglen && (log = (char *)malloc(loglen + 1))) { + if (clGetProgramBuildInfo(ctx->program, device, CL_PROGRAM_BUILD_LOG, loglen, log, NULL) == CL_SUCCESS) { + log[loglen] = 0; + fprintf(stderr, "%s\n", log); + } + free(log); + } + goto end; + } + if (!(ctx->leafs_tree_kernel = clCreateKernel(ctx->program, "lms_leafs_tree", &err))) { + cl_error_print(err); + goto end; + } + if (!(ctx->leafs_compact_kernel = clCreateKernel(ctx->program, "lms_leafs_compact", &err))) { + cl_error_print(err); + goto end; + } + if (!(ctx->internal_nodes_kernel = clCreateKernel(ctx->program, "lms_internal_nodes", &err))) { + cl_error_print(err); + goto end; + } + + ctx->local_work_size = LMS_CL_DEFAULT_LOCAL_WORK_SIZE; + if (clGetKernelWorkGroupInfo(ctx->leafs_tree_kernel, device, CL_KERNEL_WORK_GROUP_SIZE, + sizeof(max_work_group_size), &max_work_group_size, NULL) == CL_SUCCESS + && max_work_group_size > 0 && ctx->local_work_size > max_work_group_size) { + ctx->local_work_size = max_work_group_size; + } + ctx->max_leaf_batch = LMS_CL_DEFAULT_MAX_LEAF_BATCH; + return 1; + +end: + lms_cl_cleanup(ctx); + return -1; +} + +static int lms_cl_enqueue_leafs_tree(LMS_CL_CTX *ctx, cl_mem mem_seed, cl_mem mem_I, + uint32_t leaf_offset, uint32_t leaf_count, uint32_t h, cl_mem mem_tree) +{ + cl_int err; + cl_uint arg = 0; + size_t global_work_size; + size_t local_work_size; + + if ((err = clSetKernelArg(ctx->leafs_tree_kernel, arg++, sizeof(cl_mem), &mem_seed)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->leafs_tree_kernel, arg++, sizeof(cl_mem), &mem_I)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->leafs_tree_kernel, arg++, sizeof(leaf_offset), &leaf_offset)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->leafs_tree_kernel, arg++, sizeof(leaf_count), &leaf_count)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->leafs_tree_kernel, arg++, sizeof(h), &h)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->leafs_tree_kernel, arg++, sizeof(cl_mem), &mem_tree)) != CL_SUCCESS) { + cl_error_print(err); + return -1; + } + + local_work_size = ctx->local_work_size; + global_work_size = lms_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); + return -1; + } + return 1; +} + +static int lms_cl_enqueue_leafs_compact(LMS_CL_CTX *ctx, cl_mem mem_seed, cl_mem mem_I, + uint32_t leaf_offset, uint32_t leaf_count, uint32_t h, cl_mem mem_leafs) +{ + cl_int err; + cl_uint arg = 0; + size_t global_work_size; + size_t local_work_size; + + if ((err = clSetKernelArg(ctx->leafs_compact_kernel, arg++, sizeof(cl_mem), &mem_seed)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->leafs_compact_kernel, arg++, sizeof(cl_mem), &mem_I)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->leafs_compact_kernel, arg++, sizeof(leaf_offset), &leaf_offset)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->leafs_compact_kernel, arg++, sizeof(leaf_count), &leaf_count)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->leafs_compact_kernel, arg++, sizeof(h), &h)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->leafs_compact_kernel, arg++, sizeof(cl_mem), &mem_leafs)) != CL_SUCCESS) { + cl_error_print(err); + return -1; + } + + local_work_size = ctx->local_work_size; + global_work_size = lms_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); + return -1; + } + return 1; +} + +int lms_cl_derive_merkle_tree(LMS_CL_CTX *ctx, + const lms_sm3_digest_t seed, const uint8_t I[16], int height, lms_sm3_digest_t *tree) +{ + int ret = -1; + cl_int err; + cl_mem mem_seed = NULL; + cl_mem mem_I = NULL; + cl_mem mem_tree = NULL; + size_t n; + size_t tree_nodes; + size_t tree_size; + size_t level_nodes; + size_t global_work_size; + size_t local_work_size; + uint32_t level_first; + + if (!ctx || !seed || !I || !tree || height < 0 || height > LMS_MAX_HEIGHT) { + error_print(); + return -1; + } + + n = (size_t)1 << height; + tree_nodes = 2*n - 1; + tree_size = tree_nodes * sizeof(lms_sm3_digest_t); + + if (!(mem_seed = clCreateBuffer(ctx->context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, 32, (void *)seed, &err))) { + cl_error_print(err); + goto end; + } + if (!(mem_I = clCreateBuffer(ctx->context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, 16, (void *)I, &err))) { + cl_error_print(err); + goto end; + } + if (!(mem_tree = clCreateBuffer(ctx->context, CL_MEM_READ_WRITE, tree_size, NULL, &err))) { + cl_error_print(err); + goto end; + } + + if (lms_cl_enqueue_leafs_tree(ctx, mem_seed, mem_I, 0, (uint32_t)n, (uint32_t)height, mem_tree) != 1) { + goto end; + } + + local_work_size = ctx->local_work_size; + for (level_nodes = n/2; level_nodes > 0; level_nodes >>= 1) { + cl_uint arg = 0; + level_first = (uint32_t)level_nodes; + if ((err = clSetKernelArg(ctx->internal_nodes_kernel, arg++, sizeof(cl_mem), &mem_I)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->internal_nodes_kernel, arg++, sizeof(level_first), &level_first)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->internal_nodes_kernel, arg++, sizeof(uint32_t), &level_first)) != CL_SUCCESS + || (err = clSetKernelArg(ctx->internal_nodes_kernel, arg++, sizeof(cl_mem), &mem_tree)) != CL_SUCCESS) { + cl_error_print(err); + goto end; + } + global_work_size = lms_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); + goto end; + } + } + + if ((err = clEnqueueReadBuffer(ctx->queue, mem_tree, CL_TRUE, 0, tree_size, tree, 0, NULL, NULL)) != CL_SUCCESS) { + cl_error_print(err); + goto end; + } + ret = 1; + +end: + if (mem_tree) clReleaseMemObject(mem_tree); + if (mem_I) clReleaseMemObject(mem_I); + if (mem_seed) clReleaseMemObject(mem_seed); + return ret; +} + +static void lms_cl_hash_internal(const uint8_t I[16], uint32_t r, + const lms_sm3_digest_t left, const lms_sm3_digest_t right, lms_sm3_digest_t out) +{ + SM3_CTX sm3_ctx; + uint8_t rbytes[4]; + static const uint8_t D_INTR[2] = { 0x83, 0x83 }; + + PUTU32(rbytes, r); + sm3_init(&sm3_ctx); + sm3_update(&sm3_ctx, I, 16); + sm3_update(&sm3_ctx, rbytes, 4); + sm3_update(&sm3_ctx, D_INTR, 2); + sm3_update(&sm3_ctx, left, 32); + sm3_update(&sm3_ctx, right, 32); + sm3_finish(&sm3_ctx, out); +} + +int lms_cl_derive_merkle_root(LMS_CL_CTX *ctx, + const lms_sm3_digest_t seed, const uint8_t I[16], int height, lms_sm3_digest_t root) +{ + int ret = -1; + cl_int err; + cl_mem mem_seed = NULL; + cl_mem mem_I = NULL; + cl_mem mem_leafs = NULL; + lms_sm3_digest_t *leafs = NULL; + lms_sm3_digest_t stack[LMS_MAX_HEIGHT + 1]; + size_t n; + size_t batch; + size_t q; + size_t i; + int num = 0; + + if (!ctx || !seed || !I || !root || height < 0 || height > LMS_MAX_HEIGHT) { + error_print(); + return -1; + } + + n = (size_t)1 << height; + batch = ctx->max_leaf_batch ? ctx->max_leaf_batch : LMS_CL_DEFAULT_MAX_LEAF_BATCH; + if (batch > n) { + batch = n; + } + if (!(leafs = (lms_sm3_digest_t *)malloc(sizeof(lms_sm3_digest_t) * batch))) { + error_print(); + return -1; + } + + if (!(mem_seed = clCreateBuffer(ctx->context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, 32, (void *)seed, &err))) { + cl_error_print(err); + goto end; + } + if (!(mem_I = clCreateBuffer(ctx->context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, 16, (void *)I, &err))) { + cl_error_print(err); + goto end; + } + if (!(mem_leafs = clCreateBuffer(ctx->context, CL_MEM_WRITE_ONLY, + sizeof(lms_sm3_digest_t) * batch, NULL, &err))) { + cl_error_print(err); + goto end; + } + + for (q = 0; q < n; ) { + size_t count = n - q; + if (count > batch) { + count = batch; + } + if (lms_cl_enqueue_leafs_compact(ctx, mem_seed, mem_I, + (uint32_t)q, (uint32_t)count, (uint32_t)height, mem_leafs) != 1) { + goto end; + } + if ((err = clEnqueueReadBuffer(ctx->queue, mem_leafs, CL_TRUE, 0, + sizeof(lms_sm3_digest_t) * count, leafs, 0, NULL, NULL)) != CL_SUCCESS) { + cl_error_print(err); + goto end; + } + for (i = 0; i < count; i++, q++) { + uint32_t r = (uint32_t)(n + q); + size_t qbits = q; + + memcpy(stack[num], leafs[i], 32); + num++; + while (qbits & 1) { + r >>= 1; + lms_cl_hash_internal(I, r, stack[num - 2], stack[num - 1], stack[num - 2]); + num--; + qbits >>= 1; + } + } + } + if (num != 1) { + error_print(); + goto end; + } + memcpy(root, stack[0], 32); + ret = 1; + +end: + if (mem_leafs) clReleaseMemObject(mem_leafs); + if (mem_I) clReleaseMemObject(mem_I); + if (mem_seed) clReleaseMemObject(mem_seed); + if (leafs) free(leafs); + return ret; +} + +int lms_cl_key_generate_ex(LMS_CL_CTX *ctx, LMS_KEY *key, int lms_type, + const lms_sm3_digest_t seed, const uint8_t I[16], int cache_tree) +{ + size_t h, n; + + if (!ctx || !key || !seed || !I) { + error_print(); + return -1; + } + if (lms_type_to_height(lms_type, &h) != 1) { + error_print(); + return -1; + } + n = (size_t)1 << h; + + memset(key, 0, sizeof(LMS_KEY)); + key->public_key.lms_type = lms_type; + key->public_key.lmots_type = LMOTS_SM3_N32_W8; + memcpy(key->public_key.I, I, 16); + memcpy(key->seed, seed, 32); + + if (cache_tree) { + if (!(key->tree = (lms_sm3_digest_t *)malloc(sizeof(lms_sm3_digest_t) * (2*n - 1)))) { + error_print(); + return -1; + } + if (lms_cl_derive_merkle_tree(ctx, key->seed, key->public_key.I, (int)h, key->tree) != 1) { + lms_key_cleanup(key); + error_print(); + return -1; + } + memcpy(key->public_key.root, key->tree[0], 32); + } else { + if (lms_cl_derive_merkle_root(ctx, key->seed, key->public_key.I, (int)h, key->public_key.root) != 1) { + lms_key_cleanup(key); + error_print(); + return -1; + } + } + key->q = 0; + return 1; +} + +int lms_cl_key_generate(LMS_CL_CTX *ctx, LMS_KEY *key, int lms_type) +{ + lms_sm3_digest_t seed; + uint8_t I[16]; + int cache_tree = 1; + + if (rand_bytes(seed, sizeof(seed)) != 1) { + error_print(); + return -1; + } + if (rand_bytes(I, sizeof(I)) != 1) { + error_print(); + return -1; + } + if (lms_cl_key_generate_ex(ctx, key, lms_type, seed, I, cache_tree) != 1) { + error_print(); + return -1; + } + return 1; +} + +int lms_cl_private_key_from_bytes(LMS_CL_CTX *ctx, LMS_KEY *key, const uint8_t **in, size_t *inlen) +{ + size_t height; + size_t n; + + if (!ctx || !key || !in || !(*in) || !inlen) { + error_print(); + return -1; + } + if (*inlen < LMS_PRIVATE_KEY_SIZE) { + error_print(); + return -1; + } + if (lms_public_key_from_bytes(key, in, inlen) != 1) { + error_print(); + return -1; + } + + memcpy(key->seed, *in, 32); + *in += 32; + *inlen -= 32; + + key->q = GETU32(*in); + *in += 4; + *inlen -= 4; + + if (lms_type_to_height(key->public_key.lms_type, &height) != 1) { + error_print(); + goto err; + } + if (key->q >= ((uint32_t)1 << height)) { + error_print(); + goto err; + } + + n = (size_t)1 << height; + if (!(key->tree = (lms_sm3_digest_t *)malloc(sizeof(lms_sm3_digest_t) * (2*n - 1)))) { + error_print(); + goto err; + } + if (lms_cl_derive_merkle_tree(ctx, key->seed, key->public_key.I, (int)height, key->tree) != 1) { + error_print(); + goto err; + } + memcpy(key->public_key.root, key->tree[0], 32); + return 1; + +err: + lms_key_cleanup(key); + return -1; +} + +int hss_cl_key_generate(LMS_CL_CTX *ctx, HSS_KEY *key, const int *lms_types, size_t levels) +{ + int ret = -1; + lms_sm3_digest_t seed; + uint8_t I[16]; + LMS_SIGN_CTX sign_ctx; + uint8_t buf[LMS_SIGNATURE_MAX_SIZE]; + int cache_tree = 1; + size_t i; + + if (!ctx || !key || !lms_types) { + error_print(); + return -1; + } + if (levels < 1 || levels > HSS_MAX_LEVELS) { + error_print(); + return -1; + } + for (i = 0; i < levels; i++) { + if (!lms_type_name(lms_types[i])) { + error_print(); + return -1; + } + } + + memset(key, 0, sizeof(*key)); + memset(&sign_ctx, 0, sizeof(sign_ctx)); + key->levels = (uint32_t)levels; + + if (rand_bytes(seed, sizeof(seed)) != 1) { + error_print(); + goto end; + } + if (rand_bytes(I, sizeof(I)) != 1) { + error_print(); + goto end; + } + if (lms_cl_key_generate_ex(ctx, &key->lms_key[0], lms_types[0], seed, I, cache_tree) != 1) { + error_print(); + goto end; + } + + for (i = 1; i < levels; i++) { + uint8_t *p = buf; + size_t len = 0; + + if (rand_bytes(seed, sizeof(seed)) != 1) { + error_print(); + goto end; + } + if (rand_bytes(I, sizeof(I)) != 1) { + error_print(); + goto end; + } + if (lms_cl_key_generate_ex(ctx, &key->lms_key[i], lms_types[i], seed, I, cache_tree) != 1) { + error_print(); + goto end; + } + if (lms_public_key_to_bytes(&key->lms_key[i], &p, &len) != 1) { + error_print(); + goto end; + } + if (lms_sign_init(&sign_ctx, &key->lms_key[i - 1]) != 1 + || lms_sign_update(&sign_ctx, buf, len) != 1 + || lms_sign_finish(&sign_ctx, buf, &len) != 1) { + error_print(); + goto end; + } + key->lms_sig[i - 1] = sign_ctx.lms_sig; + } + + ret = 1; +end: + gmssl_secure_clear(seed, sizeof(seed)); + gmssl_secure_clear(&sign_ctx, sizeof(sign_ctx)); + if (ret != 1) hss_key_cleanup(key); + return ret; +} + +int hss_cl_private_key_from_bytes(LMS_CL_CTX *ctx, HSS_KEY *key, const uint8_t **in, size_t *inlen) +{ + size_t i; + + if (!ctx || !key || !in || !(*in) || !inlen) { + error_print(); + return -1; + } + if (*inlen < 4) { + error_print(); + return -1; + } + + key->levels = GETU32(*in); + if (key->levels < 1 || key->levels > HSS_MAX_LEVELS) { + error_print(); + return -1; + } + *in += 4; + *inlen -= 4; + + if (lms_cl_private_key_from_bytes(ctx, &key->lms_key[0], in, inlen) != 1) { + error_print(); + goto err; + } + + for (i = 1; i < key->levels; i++) { + LMS_SIGN_CTX sign_ctx; + uint8_t buf[LMS_PUBLIC_KEY_SIZE]; + uint8_t *p = buf; + size_t len = 0; + + if (lms_cl_private_key_from_bytes(ctx, &key->lms_key[i], in, inlen) != 1) { + error_print(); + goto err; + } + if (lms_signature_from_bytes(&key->lms_sig[i - 1], in, inlen) != 1) { + error_print(); + goto err; + } + + if (lms_public_key_to_bytes(&key->lms_key[i], &p, &len) != 1) { + error_print(); + goto err; + } + if (lms_verify_init_ex(&sign_ctx, &key->lms_key[i - 1], &key->lms_sig[i - 1]) != 1 + || lms_verify_update(&sign_ctx, buf, len) != 1 + || lms_verify_finish(&sign_ctx) != 1) { + error_print(); + goto err; + } + } + return 1; + +err: + hss_key_cleanup(key); + return -1; +} + +int hss_cl_key_update(LMS_CL_CTX *ctx, HSS_KEY *key) +{ + int level; + LMS_KEY *lms_key; + size_t count; + + if (!ctx || !key) { + error_print(); + return -1; + } + + for (level = key->levels; level > 0; level--) { + lms_key = &key->lms_key[level - 1]; + if (lms_key_remaining_signs(lms_key, &count) != 1) { + error_print(); + return -1; + } + if (count > 0) { + break; + } + } + if (level >= (int)key->levels) { + error_print(); + return -1; + } + if (level == 0) { + return 0; + } + + for (; level < (int)key->levels; level++) { + int lms_type = key->lms_key[level].public_key.lms_type; + LMS_SIGN_CTX sign_ctx; + uint8_t buf[LMS_PUBLIC_KEY_SIZE]; + uint8_t *p = buf; + size_t len = 0; + + lms_key_cleanup(&key->lms_key[level]); + + if (lms_cl_key_generate(ctx, &key->lms_key[level], lms_type) != 1) { + error_print(); + return -1; + } + if (lms_public_key_to_bytes(&key->lms_key[level], &p, &len) != 1) { + error_print(); + return -1; + } + if (lms_sign_init(&sign_ctx, &key->lms_key[level - 1]) != 1) { + error_print(); + return -1; + } + if (lms_sign_update(&sign_ctx, buf, len) != 1) { + error_print(); + return -1; + } + if (lms_sign_finish_ex(&sign_ctx, &key->lms_sig[level - 1]) != 1) { + error_print(); + return -1; + } + } + + if (key->update_callback) { + if (key->update_callback(key) != 1) { + error_print(); + return -1; + } + } + return 1; +} + +int hss_cl_sign_init(LMS_CL_CTX *ctx, HSS_SIGN_CTX *sign_ctx, HSS_KEY *key) +{ + size_t count; + size_t i; + + if (!ctx || !sign_ctx || !key) { + error_print(); + return -1; + } + if (key->levels < 1 || key->levels > HSS_MAX_LEVELS) { + error_print(); + return -1; + } + + memset(sign_ctx, 0, sizeof(*sign_ctx)); + + if (lms_sign_init(&sign_ctx->lms_sign_ctx, &key->lms_key[key->levels - 1]) != 1) { + error_print(); + return -1; + } + + sign_ctx->levels = key->levels; + + for (i = 0; i < key->levels - 1; i++) { + sign_ctx->lms_public_keys[i] = key->lms_key[i + 1].public_key; + sign_ctx->lms_sigs[i] = key->lms_sig[i]; + } + + if (lms_key_remaining_signs(&key->lms_key[key->levels - 1], &count) != 1) { + error_print(); + return -1; + } + if (count == 0) { + if (hss_cl_key_update(ctx, key) != 1) { + error_print(); + return -1; + } + } + + 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]) +{ + uint dgst[8]; + uint W[68]; + uint k; + + sm3_init_state(dgst); + W[0] = load_be32(I); + W[1] = load_be32(I + 4); + W[2] = load_be32(I + 8); + W[3] = load_be32(I + 12); + W[4] = q; + W[5] = ((uint)i << 16) | ((uint)j << 8) | (uint)in[0]; + for (k = 0; k < 7; k++) { + W[6 + k] = load_be32_private(in + 1 + 4*k); + } + W[13] = ((uint)in[29] << 24) | ((uint)in[30] << 16) | ((uint)in[31] << 8) | 0x80U; + W[14] = 0; + W[15] = 440; + sm3_compress_words(dgst, W); + for (k = 0; k < 8; k++) { + store_be32(out + 4*k, dgst[k]); + } +} + +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]) +{ + uint i; + for (i = 0; i < 32; i++) { + sm3_update_byte(dgst, block, num, nblocks, z[i]); + } +} + +void lms_compute_leaf(__global const uchar *seed, __global const uchar *I, + uint q, uint h, __private uchar leaf[32]) +{ + uchar z[32]; + uchar tmp[32]; + uint dgst[8]; + uchar block[64]; + uint num = 0; + ulong nblocks = 0; + uint i, j, k; + uint r = (1U << h) + q; + + sm3_init_state(dgst); + for (i = 0; i < 16; i++) { + sm3_update_byte(dgst, block, &num, &nblocks, I[i]); + } + sm3_update_byte(dgst, block, &num, &nblocks, (uchar)(q >> 24)); + sm3_update_byte(dgst, block, &num, &nblocks, (uchar)(q >> 16)); + sm3_update_byte(dgst, block, &num, &nblocks, (uchar)(q >> 8)); + sm3_update_byte(dgst, block, &num, &nblocks, (uchar)q); + sm3_update_byte(dgst, block, &num, &nblocks, 0x80); + sm3_update_byte(dgst, block, &num, &nblocks, 0x80); + + for (i = 0; i < 34; i++) { + for (k = 0; k < 32; k++) { + tmp[k] = seed[k]; + } + sm3_hash_lmots_step(I, q, (ushort)i, 0xff, tmp, z); + for (j = 0; j < 255; j++) { + for (k = 0; k < 32; k++) { + tmp[k] = z[k]; + } + sm3_hash_lmots_step(I, q, (ushort)i, (uchar)j, tmp, z); + } + sm3_public_hash_update(dgst, block, &num, &nblocks, z); + } + sm3_finish_ctx(dgst, block, num, nblocks, tmp); + + sm3_init_state(dgst); + { + uint W[68]; + W[0] = load_be32(I); + W[1] = load_be32(I + 4); + W[2] = load_be32(I + 8); + W[3] = load_be32(I + 12); + W[4] = r; + W[5] = 0x82820000U | ((uint)tmp[0] << 8) | (uint)tmp[1]; + for (k = 0; k < 7; k++) { + W[6 + k] = load_be32_private(tmp + 2 + 4*k); + } + W[13] = ((uint)tmp[30] << 24) | ((uint)tmp[31] << 16) | 0x8000U; + W[14] = 0; + W[15] = 432; + sm3_compress_words(dgst, W); + for (k = 0; k < 8; k++) { + store_be32(leaf + 4*k, dgst[k]); + } + } +} + +void sm3_hash_internal(__global const uchar *I, uint r, + __global const uchar *left, __global const uchar *right, __global uchar *out) +{ + uint dgst[8]; + uint W[68]; + uint k; + + sm3_init_state(dgst); + W[0] = load_be32(I); + W[1] = load_be32(I + 4); + W[2] = load_be32(I + 8); + W[3] = load_be32(I + 12); + W[4] = r; + W[5] = 0x83830000U | ((uint)left[0] << 8) | (uint)left[1]; + for (k = 0; k < 7; k++) { + W[6 + k] = load_be32(left + 2 + 4*k); + } + W[13] = ((uint)left[30] << 24) | ((uint)left[31] << 16) + | ((uint)right[0] << 8) | (uint)right[1]; + W[14] = load_be32(right + 2); + W[15] = load_be32(right + 6); + sm3_compress_words(dgst, W); + + W[0] = load_be32(right + 10); + W[1] = load_be32(right + 14); + W[2] = load_be32(right + 18); + W[3] = load_be32(right + 22); + W[4] = load_be32(right + 26); + W[5] = ((uint)right[30] << 24) | ((uint)right[31] << 16) | 0x8000U; + for (k = 6; k < 15; k++) { + W[k] = 0; + } + W[15] = 688; + sm3_compress_words(dgst, W); + for (k = 0; k < 8; k++) { + out[4*k] = (uchar)(dgst[k] >> 24); + out[4*k + 1] = (uchar)(dgst[k] >> 16); + out[4*k + 2] = (uchar)(dgst[k] >> 8); + out[4*k + 3] = (uchar)dgst[k]; + } +} + +__kernel void lms_leafs_tree(__global const uchar *seed, __global const uchar *I, + uint leaf_offset, uint leaf_count, uint h, __global uchar *tree) +{ + uint id = get_global_id(0); + uint q; + uint r; + uchar leaf[32]; + uint i; + + if (id >= leaf_count) { + return; + } + q = leaf_offset + id; + r = (1U << h) + q; + lms_compute_leaf(seed, I, q, h, leaf); + for (i = 0; i < 32; i++) { + tree[(r - 1)*32 + i] = leaf[i]; + } +} + +__kernel void lms_leafs_compact(__global const uchar *seed, __global const uchar *I, + uint leaf_offset, uint leaf_count, uint h, __global uchar *leafs) +{ + uint id = get_global_id(0); + uchar leaf[32]; + uint i; + + if (id >= leaf_count) { + return; + } + lms_compute_leaf(seed, I, leaf_offset + id, h, leaf); + for (i = 0; i < 32; i++) { + leafs[id*32 + i] = leaf[i]; + } +} + +__kernel void lms_internal_nodes(__global const uchar *I, uint level_first, + uint level_count, __global uchar *tree) +{ + uint id = get_global_id(0); + uint r; + + if (id >= level_count) { + return; + } + r = level_first + id; + sm3_hash_internal(I, r, tree + (2*r - 1)*32, tree + (2*r)*32, tree + (r - 1)*32); +} + +); diff --git a/src/tlcp.c b/src/tlcp.c index 0fd1201e..f243c598 100644 --- a/src/tlcp.c +++ b/src/tlcp.c @@ -91,22 +91,32 @@ static int tlcp_cipher_suite_is_ecdhe(int cipher_suite) -ServerKeyExchange +服务器Certificate + +常规采用ECC_, ECDHE_ 套件时，服务器证书是双证书证书链 + +opaque ASN.1Cert<1..2^24-1>; + +struct { + ASN.1Cert certificate<0..2^24-1>; +} Certificate; + +但是采用IBC_ 套件时，服务器证书消息变为 + +opaque ASN1.1IBCParam<1..2^24-1>; + +struct { + opaque ibc_id<1..2^16-1>; + ASN1.1IBCParam ibc_parameter; +} Certificate; + +其中ibc_id的负载数据为服务器的域名 +ibc_parameter 是SM9_ENC_MASTER_PUBLIC_KEY 的der编码 + 即sm9_enc_master_public_key_to_der的编码 + +客户端密钥交换的方式和ECC_套件类似，用sm9_encrypt(ibc_parameter, ibc_id)加密pre_master_secret，然后把密文放到ClientKeyExchange中. -select (KeyExchangeAlgorithm) { - case ECC: - digitall-signed struct { - opaque client_random[32]; - opaque server_random[32]; - opaque ASN1.Cert<1..2^24-1>; - } signed_params; - case ECDHE: - ServerECDHEParams params; - digitally-signed struct { - opaque client_random[32]; - opaque server_random[32]; - ServerECDHEParams params; diff --git a/tests/lms_cltest.c b/tests/lms_cltest.c new file mode 100644 index 00000000..c21d8ff8 --- /dev/null +++ b/tests/lms_cltest.c @@ -0,0 +1,289 @@ +/* + * 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 +#include +#include +#include +#ifdef _WIN32 +#include +#else +#include +#endif +#include +#include +#include + + +static int test_lms_cl_derive_merkle_tree(void) +{ + LMS_CL_CTX ctx; + lms_sm3_digest_t seed = {0}; + uint8_t I[16] = {0}; + int h = 5; + size_t n = (size_t)1 << h; + lms_sm3_digest_t *tree = NULL; + lms_sm3_digest_t *cl_tree = NULL; + lms_sm3_digest_t root; + int ret = -1; + + if (lms_cl_init(&ctx) != 1) { + fprintf(stderr, "%s: OpenCL unavailable, skipped\n", __FUNCTION__); + return 1; + } + if (!(tree = (lms_sm3_digest_t *)malloc(sizeof(lms_sm3_digest_t) * (2*n - 1)))) { + error_print(); + goto end; + } + if (!(cl_tree = (lms_sm3_digest_t *)malloc(sizeof(lms_sm3_digest_t) * (2*n - 1)))) { + error_print(); + goto end; + } + + lms_derive_merkle_tree(seed, I, h, tree); + if (lms_cl_derive_merkle_tree(&ctx, seed, I, h, cl_tree) != 1) { + error_print(); + goto end; + } + if (memcmp(tree, cl_tree, sizeof(lms_sm3_digest_t) * (2*n - 1)) != 0) { + error_print(); + goto end; + } + if (lms_cl_derive_merkle_root(&ctx, seed, I, h, root) != 1) { + error_print(); + goto end; + } + if (memcmp(tree[0], root, 32) != 0) { + error_print(); + goto end; + } + + printf("%s() ok\n", __FUNCTION__); + ret = 1; +end: + if (tree) free(tree); + if (cl_tree) free(cl_tree); + lms_cl_cleanup(&ctx); + return ret; +} + +static int test_lms_cl_key_generate(void) +{ + LMS_CL_CTX ctx; + lms_sm3_digest_t seed = {1}; + uint8_t I[16] = {2}; + LMS_KEY key; + LMS_KEY cl_key; + LMS_KEY key2; + uint8_t keybuf[LMS_PRIVATE_KEY_SIZE]; + uint8_t *p = keybuf; + const uint8_t *cp = keybuf; + size_t keylen = 0; + int ret = -1; + + memset(&key, 0, sizeof(key)); + memset(&cl_key, 0, sizeof(cl_key)); + memset(&key2, 0, sizeof(key2)); + + if (lms_cl_init(&ctx) != 1) { + fprintf(stderr, "%s: OpenCL unavailable, skipped\n", __FUNCTION__); + return 1; + } + if (lms_key_generate_ex(&key, LMS_SM3_M32_H5, seed, I, 1) != 1) { + error_print(); + goto end; + } + if (lms_cl_key_generate_ex(&ctx, &cl_key, LMS_SM3_M32_H5, seed, I, 1) != 1) { + error_print(); + goto end; + } + if (memcmp(&key.public_key, &cl_key.public_key, sizeof(LMS_PUBLIC_KEY)) != 0) { + error_print(); + goto end; + } + if (lms_private_key_to_bytes(&key, &p, &keylen) != 1) { + error_print(); + goto end; + } + if (keylen != LMS_PRIVATE_KEY_SIZE) { + error_print(); + goto end; + } + if (lms_cl_private_key_from_bytes(&ctx, &key2, &cp, &keylen) != 1 || keylen != 0) { + error_print(); + goto end; + } + if (memcmp(&key.public_key, &key2.public_key, sizeof(LMS_PUBLIC_KEY)) != 0 + || memcmp(key.tree, key2.tree, sizeof(lms_sm3_digest_t) * ((1 << 5)*2 - 1)) != 0) { + error_print(); + goto end; + } + printf("%s() ok\n", __FUNCTION__); + ret = 1; +end: + lms_key_cleanup(&key); + lms_key_cleanup(&cl_key); + lms_key_cleanup(&key2); + lms_cl_cleanup(&ctx); + return ret; +} + +static int test_hss_cl_key_generate(void) +{ + LMS_CL_CTX ctx; + int lms_types[] = { + LMS_SM3_M32_H5, + LMS_SM3_M32_H5, + }; + HSS_KEY key; + HSS_KEY key2; + HSS_SIGN_CTX sign_ctx; + HSS_SIGN_CTX verify_ctx; + uint8_t keybuf[HSS_PRIVATE_KEY_MAX_SIZE]; + uint8_t sig[HSS_SIGNATURE_MAX_SIZE]; + uint8_t *p = keybuf; + const uint8_t *cp = keybuf; + uint8_t msg[] = "abc"; + size_t keylen = 0; + size_t siglen = 0; + int ret = -1; + + memset(&key, 0, sizeof(key)); + memset(&key2, 0, sizeof(key2)); + memset(&sign_ctx, 0, sizeof(sign_ctx)); + memset(&verify_ctx, 0, sizeof(verify_ctx)); + + if (lms_cl_init(&ctx) != 1) { + fprintf(stderr, "%s: OpenCL unavailable, skipped\n", __FUNCTION__); + return 1; + } + if (hss_cl_key_generate(&ctx, &key, lms_types, sizeof(lms_types)/sizeof(lms_types[0])) != 1) { + error_print(); + goto end; + } + if (hss_private_key_to_bytes(&key, &p, &keylen) != 1) { + error_print(); + goto end; + } + if (hss_cl_private_key_from_bytes(&ctx, &key2, &cp, &keylen) != 1 || keylen != 0) { + error_print(); + goto end; + } + if (hss_public_key_equ(&key, &key2) != 1) { + error_print(); + goto end; + } + key2.lms_key[1].q = 31; + if (hss_cl_sign_init(&ctx, &sign_ctx, &key2) != 1 + || hss_sign_update(&sign_ctx, msg, sizeof(msg)) != 1 + || hss_sign_finish(&sign_ctx, sig, &siglen) != 1) { + error_print(); + goto end; + } + if (key2.lms_key[0].q != 2 || key2.lms_key[1].q != 0) { + error_print(); + goto end; + } + if (hss_verify_init(&verify_ctx, &key, sig, siglen) != 1 + || hss_verify_update(&verify_ctx, msg, sizeof(msg)) != 1 + || hss_verify_finish(&verify_ctx) != 1) { + error_print(); + goto end; + } + + printf("%s() ok\n", __FUNCTION__); + ret = 1; +end: + hss_key_cleanup(&key); + hss_key_cleanup(&key2); + lms_cl_cleanup(&ctx); + return ret; +} + +#if ENABLE_TEST_SPEED +static double get_seconds(void) +{ +#ifdef _WIN32 + LARGE_INTEGER freq; + LARGE_INTEGER count; + + QueryPerformanceFrequency(&freq); + QueryPerformanceCounter(&count); + return (double)count.QuadPart/(double)freq.QuadPart; +#else + struct timeval tv; + + gettimeofday(&tv, NULL); + return (double)tv.tv_sec + (double)tv.tv_usec/1000000; +#endif +} + +static int speed_lms_cl_derive_merkle_root(void) +{ + LMS_CL_CTX ctx; + lms_sm3_digest_t seed = {0}; + uint8_t I[16] = {0}; + lms_sm3_digest_t root; + lms_sm3_digest_t cl_root; + double begin; + double seconds; + int h = 10; + int cl_h = 15; + + if (lms_cl_init(&ctx) != 1) { + fprintf(stderr, "%s: OpenCL unavailable, skipped\n", __FUNCTION__); + return 1; + } + + begin = get_seconds(); + lms_derive_merkle_root(seed, I, h, root); + seconds = get_seconds() - begin; + fprintf(stderr, "%s: CPU H%d %.3f seconds\n", __FUNCTION__, h, seconds); + + begin = get_seconds(); + if (lms_cl_derive_merkle_root(&ctx, seed, I, h, cl_root) != 1) { + error_print(); + lms_cl_cleanup(&ctx); + return -1; + } + seconds = get_seconds() - begin; + fprintf(stderr, "%s: OpenCL H%d %.3f seconds\n", __FUNCTION__, h, seconds); + if (memcmp(root, cl_root, 32) != 0) { + error_print(); + lms_cl_cleanup(&ctx); + return -1; + } + + begin = get_seconds(); + if (lms_cl_derive_merkle_root(&ctx, seed, I, cl_h, cl_root) != 1) { + error_print(); + lms_cl_cleanup(&ctx); + return -1; + } + seconds = get_seconds() - begin; + fprintf(stderr, "%s: OpenCL H%d %.3f seconds\n", __FUNCTION__, cl_h, seconds); + lms_cl_cleanup(&ctx); + return 1; +} +#endif + +int main(void) +{ + if (test_lms_cl_derive_merkle_tree() != 1) goto err; + if (test_lms_cl_key_generate() != 1) goto err; + if (test_hss_cl_key_generate() != 1) goto err; +#if ENABLE_TEST_SPEED + if (speed_lms_cl_derive_merkle_root() != 1) goto err; +#endif + printf("%s all tests passed\n", __FILE__); + return 0; +err: + error_print(); + return 1; +} diff --git a/tools/hsskeygen.c b/tools/hsskeygen.c index 62532440..4035ac2b 100644 --- a/tools/hsskeygen.c +++ b/tools/hsskeygen.c @@ -15,6 +15,9 @@ #include #include #include +#ifdef ENABLE_LMS_CL +#include +#endif static const char *usage = "-lms_types types -out file [-pubout file] [-verbose]\n"; @@ -55,8 +58,15 @@ int hsskeygen_main(int argc, char **argv) uint8_t *pout = out; uint8_t *ppubout = pubout; size_t outlen = 0, puboutlen = 0; +#ifdef ENABLE_LMS_CL + LMS_CL_CTX cl_ctx; + int cl_initialized = 0; +#endif memset(&key, 0, sizeof(key)); +#ifdef ENABLE_LMS_CL + memset(&cl_ctx, 0, sizeof(cl_ctx)); +#endif argc--; argv++; @@ -129,10 +139,22 @@ bad: goto end; } +#ifdef ENABLE_LMS_CL + if (lms_cl_init(&cl_ctx) != 1) { + error_print(); + goto end; + } + cl_initialized = 1; + if (hss_cl_key_generate(&cl_ctx, &key, lms_types_val, levels) != 1) { + error_print(); + goto end; + } +#else if (hss_key_generate(&key, lms_types_val, levels) != 1) { error_print(); return -1; } +#endif if (verbose) { hss_public_key_print(stderr, 0, 0, "hss_public_key", &key); } @@ -157,6 +179,9 @@ bad: ret = 0; end: +#ifdef ENABLE_LMS_CL + if (cl_initialized) lms_cl_cleanup(&cl_ctx); +#endif hss_key_cleanup(&key); gmssl_secure_clear(out, outlen); if (outfile && outfp) fclose(outfp); diff --git a/tools/hsssign.c b/tools/hsssign.c index 0c8b0c51..fd6ac64b 100644 --- a/tools/hsssign.c +++ b/tools/hsssign.c @@ -15,6 +15,9 @@ #include #include #include +#ifdef ENABLE_LMS_CL +#include +#endif static const char *usage = "-key file [-in file] [-out file] [-verbose]\n"; @@ -75,8 +78,15 @@ int hsssign_main(int argc, char **argv) HSS_SIGN_CTX ctx; uint8_t sig[HSS_SIGNATURE_MAX_SIZE]; size_t siglen; +#ifdef ENABLE_LMS_CL + LMS_CL_CTX cl_ctx; + int cl_initialized = 0; +#endif memset(&key, 0, sizeof(key)); +#ifdef ENABLE_LMS_CL + memset(&cl_ctx, 0, sizeof(cl_ctx)); +#endif argc--; argv++; @@ -136,10 +146,22 @@ bad: fprintf(stderr, "%s: read private key failure\n", prog); goto end; } +#ifdef ENABLE_LMS_CL + if (lms_cl_init(&cl_ctx) != 1) { + error_print(); + goto end; + } + cl_initialized = 1; + if (hss_cl_private_key_from_bytes(&cl_ctx, &key, &cp, &keylen) != 1) { + error_print(); + goto end; + } +#else if (hss_private_key_from_bytes(&key, &cp, &keylen) != 1) { error_print(); goto end; } +#endif if (keylen) { error_print(); goto end; @@ -154,10 +176,17 @@ bad: goto end; } +#ifdef ENABLE_LMS_CL + if (hss_cl_sign_init(&cl_ctx, &ctx, &key) != 1) { + error_print(); + goto end; + } +#else if (hss_sign_init(&ctx, &key) != 1) { error_print(); goto end; } +#endif while (1) { uint8_t buf[1024]; @@ -185,6 +214,9 @@ bad: ret = 0; end: +#ifdef ENABLE_LMS_CL + if (cl_initialized) lms_cl_cleanup(&cl_ctx); +#endif hss_key_cleanup(&key); gmssl_secure_clear(keybuf, sizeof(keybuf)); gmssl_secure_clear(&ctx, sizeof(ctx)); diff --git a/tools/lmskeygen.c b/tools/lmskeygen.c index c0d7232c..52461668 100644 --- a/tools/lmskeygen.c +++ b/tools/lmskeygen.c @@ -15,6 +15,9 @@ #include #include #include +#ifdef ENABLE_LMS_CL +#include +#endif static const char *usage = "-lms_type type -out file [-pubout file] [-verbose]\n"; @@ -49,8 +52,15 @@ int lmskeygen_main(int argc, char **argv) uint8_t *pout = out; uint8_t *ppubout = pubout; size_t outlen = 0, puboutlen = 0; +#ifdef ENABLE_LMS_CL + LMS_CL_CTX cl_ctx; + int cl_initialized = 0; +#endif memset(&key, 0, sizeof(key)); +#ifdef ENABLE_LMS_CL + memset(&cl_ctx, 0, sizeof(cl_ctx)); +#endif argc--; argv++; @@ -110,10 +120,22 @@ bad: goto end; } +#ifdef ENABLE_LMS_CL + if (lms_cl_init(&cl_ctx) != 1) { + error_print(); + goto end; + } + cl_initialized = 1; + if (lms_cl_key_generate(&cl_ctx, &key, lms_type_val) != 1) { + error_print(); + goto end; + } +#else if (lms_key_generate(&key, lms_type_val) != 1) { error_print(); return -1; } +#endif if (verbose) { lms_public_key_print(stderr, 0, 0, "lms_public_key", &key); } @@ -138,6 +160,9 @@ bad: ret = 0; end: +#ifdef ENABLE_LMS_CL + if (cl_initialized) lms_cl_cleanup(&cl_ctx); +#endif lms_key_cleanup(&key); gmssl_secure_clear(out, outlen); if (outfile && outfp) fclose(outfp); diff --git a/tools/lmssign.c b/tools/lmssign.c index d8eb93ea..3e889c8a 100644 --- a/tools/lmssign.c +++ b/tools/lmssign.c @@ -15,6 +15,9 @@ #include #include #include +#ifdef ENABLE_LMS_CL +#include +#endif static const char *usage = "-key file [-in file] [-out file] [-verbose]\n"; @@ -75,8 +78,15 @@ int lmssign_main(int argc, char **argv) LMS_SIGN_CTX ctx; uint8_t sig[LMS_SIGNATURE_MAX_SIZE]; size_t siglen; +#ifdef ENABLE_LMS_CL + LMS_CL_CTX cl_ctx; + int cl_initialized = 0; +#endif memset(&key, 0, sizeof(key)); +#ifdef ENABLE_LMS_CL + memset(&cl_ctx, 0, sizeof(cl_ctx)); +#endif argc--; argv++; @@ -136,10 +146,22 @@ bad: fprintf(stderr, "%s: read private key failure\n", prog); goto end; } +#ifdef ENABLE_LMS_CL + if (lms_cl_init(&cl_ctx) != 1) { + error_print(); + goto end; + } + cl_initialized = 1; + if (lms_cl_private_key_from_bytes(&cl_ctx, &key, &cp, &keylen) != 1) { + error_print(); + goto end; + } +#else if (lms_private_key_from_bytes(&key, &cp, &keylen) != 1) { error_print(); goto end; } +#endif if (keylen) { error_print(); return -1; @@ -185,6 +207,9 @@ bad: ret = 0; end: +#ifdef ENABLE_LMS_CL + if (cl_initialized) lms_cl_cleanup(&cl_ctx); +#endif lms_key_cleanup(&key); gmssl_secure_clear(keybuf, sizeof(keybuf)); gmssl_secure_clear(&ctx, sizeof(ctx));