int main(int argc, char *argv[]) {
big_number *res = bn_create(1000, NULL);
int i;
for (i = 0; i < 50; i++) {
big_number *tmp1 = bn_create(500, text[i]);
big_number *tmp2 = bn_add(tmp1, res);
bn_free(tmp1);
bn_free(res);
res = tmp2;
}
bn_print(res);
return 0;
}
status_t element_printf(const char *msg, element_t e)
{
if(e->isInitialized == TRUE) {
printf("%s", msg);
if(e->type == ZR)
bn_print(e->bn);
else if(e->type == G1)
g1_print(e->g1);
else if(e->type == G2)
g2_print(e->g2);
else if(e->type == GT)
gt_print(e->gt);
return ELEMENT_OK;
}
return ELEMENT_INVALID_RESULT;
}
int ecdsa_set_curve(u32 type)
{
if (ecdsa_get_params(type, ec_p, ec_a, ec_b, ec_N, ec_G.x, ec_G.y) < 0)
return -1;
bn_print("p", ec_p, 20);
bn_print("a", ec_a, 20);
bn_print("b", ec_b, 20);
bn_print("N", ec_N, 21);
bn_print("Gx", ec_G.x, 20);
bn_print("Gy", ec_G.y, 20);
bn_to_mon(ec_a, ec_p, 20);
bn_to_mon(ec_b, ec_p, 20);
point_to_mon(&ec_G);
return 0;
}
static void tests_relic_ecdh(void)
{
/* The following is an example for doing an elliptic-curve Diffie-Hellman
key exchange.
*/
/* Select an elliptic curve configuration */
if (ec_param_set_any() == STS_OK) {
#if (TEST_RELIC_SHOW_OUTPUT == 1)
ec_param_print();
#endif
bn_t privateA;
ec_t publicA;
uint8_t sharedKeyA[MD_LEN];
bn_t privateB;
ec_t publicB;
uint8_t sharedKeyB[MD_LEN];
bn_null(privateA);
ec_null(publicA);
bn_new(privateA);
ec_new(publicA);
bn_null(privateB);
ec_null(publicB);
bn_new(privateB);
ec_new(publicB);
/* User A generates private/public key pair */
TEST_ASSERT_EQUAL_INT(STS_OK, cp_ecdh_gen(privateA, publicA));
#if (TEST_RELIC_SHOW_OUTPUT == 1)
printf("User A\n");
printf("======\n");
printf("private key: ");
bn_print(privateA);
printf("\npublic key: ");
ec_print(publicA);
printf("\n");
#endif
/* User B generates private/public key pair */
TEST_ASSERT_EQUAL_INT(STS_OK, cp_ecdh_gen(privateB, publicB));
#if (TEST_RELIC_SHOW_OUTPUT == 1)
printf("User B\n");
printf("======\n");
printf("private key: ");
bn_print(privateB);
printf("\npublic key: ");
ec_print(publicB);
printf("\n");
#endif
/* In a protocol you would exchange the public keys now */
/* User A calculates shared secret */
TEST_ASSERT_EQUAL_INT(STS_OK, cp_ecdh_key(sharedKeyA, MD_LEN, privateA, publicB));
#if (TEST_RELIC_SHOW_OUTPUT == 1)
printf("\nshared key computed by user A: ");
print_mem(sharedKeyA, MD_LEN);
#endif
/* User B calculates shared secret */
TEST_ASSERT_EQUAL_INT(STS_OK, cp_ecdh_key(sharedKeyB, MD_LEN, privateB, publicA));
#if (TEST_RELIC_SHOW_OUTPUT == 1)
printf("\nshared key computed by user B: ");
print_mem(sharedKeyB, MD_LEN);
#endif
/* The secrets should be the same now */
TEST_ASSERT_EQUAL_INT(CMP_EQ, util_cmp_const(sharedKeyA, sharedKeyB, MD_LEN));
bn_free(privateA);
ec_free(publicA);
bn_free(privateB);
ec_free(publicB);
#if (TEST_RELIC_SHOW_OUTPUT == 1)
printf("\nRELIC EC-DH test successful\n");
#endif
}
}
acpi_status
bn_add_device(
BM_HANDLE device_handle,
void **context)
{
acpi_status status = AE_OK;
BM_DEVICE *device = NULL;
BN_CONTEXT *button = NULL;
FUNCTION_TRACE("bn_add_device");
ACPI_DEBUG_PRINT ((ACPI_DB_INFO, "Adding button device [%02x].\n", device_handle));
if (!context || *context) {
ACPI_DEBUG_PRINT ((ACPI_DB_ERROR, "Invalid context.\n"));
return_ACPI_STATUS(AE_BAD_PARAMETER);
}
/*
* Get information on this device.
*/
status = bm_get_device_info( device_handle, &device );
if (ACPI_FAILURE(status)) {
return_ACPI_STATUS(status);
}
/*
* Allocate a new BN_CONTEXT structure.
*/
button = acpi_os_callocate(sizeof(BN_CONTEXT));
if (!button) {
return_ACPI_STATUS(AE_NO_MEMORY);
}
button->device_handle = device->handle;
button->acpi_handle = device->acpi_handle;
/*
* Power Button?
* -------------
* Either fixed-feature or generic (namespace) types.
*/
if (strncmp(device->id.hid, BN_HID_POWER_BUTTON,
sizeof(BM_DEVICE_HID)) == 0) {
if (device->id.type == BM_TYPE_FIXED_BUTTON) {
button->type = BN_TYPE_POWER_BUTTON_FIXED;
/* Register for fixed-feature events. */
status = acpi_install_fixed_event_handler(
ACPI_EVENT_POWER_BUTTON, bn_notify_fixed,
(void*)button);
}
else {
button->type = BN_TYPE_POWER_BUTTON;
}
}
/*
* Sleep Button?
* -------------
* Either fixed-feature or generic (namespace) types.
*/
else if (strncmp( device->id.hid, BN_HID_SLEEP_BUTTON,
sizeof(BM_DEVICE_HID)) == 0) {
if (device->id.type == BM_TYPE_FIXED_BUTTON) {
button->type = BN_TYPE_SLEEP_BUTTON_FIXED;
/* Register for fixed-feature events. */
status = acpi_install_fixed_event_handler(
ACPI_EVENT_SLEEP_BUTTON, bn_notify_fixed,
(void*)button);
}
else {
button->type = BN_TYPE_SLEEP_BUTTON;
}
}
/*
* LID Switch?
* -----------
*/
else if (strncmp( device->id.hid, BN_HID_LID_SWITCH,
sizeof(BM_DEVICE_HID)) == 0) {
button->type = BN_TYPE_LID_SWITCH;
}
status = bn_osl_add_device(button);
if (ACPI_FAILURE(status)) {
goto end;
}
*context = button;
bn_print(button);
end:
if (ACPI_FAILURE(status)) {
acpi_os_free(button);
}
return_ACPI_STATUS(status);
}
int main(int argc, char **argv){
printf("Check OpenCL environtment\n");
cl_platform_id platid;
cl_device_id devid;
cl_int res;
size_t param;
/* Query OpenCL, get some information about the returned device */
clGetPlatformIDs(1u, &platid, NULL);
clGetDeviceIDs(platid, CL_DEVICE_TYPE_ALL, 1, &devid, NULL);
cl_char vendor_name[1024] = {0};
cl_char device_name[1024] = {0};
clGetDeviceInfo(devid, CL_DEVICE_VENDOR, sizeof(vendor_name), vendor_name, NULL);
clGetDeviceInfo(devid, CL_DEVICE_NAME, sizeof(device_name), device_name, NULL);
printf("Connecting to OpenCL device:\t%s %s\n", vendor_name, device_name);
clGetDeviceInfo(devid, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), ¶m, NULL);
printf("CL_DEVICE_MAX_COMPUTE_UNITS\t%d\n", param);
clGetDeviceInfo(devid, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), ¶m, NULL);
printf("CL_DEVICE_MAX_WORK_GROUP_SIZE\t%u\n", param);
clGetDeviceInfo(devid, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), ¶m, NULL);
printf("CL_DEVICE_LOCAL_MEM_SIZE\t%ub\n", param);
/* Check if kernel source exists, we compile argv[1] passed kernel */
if(argv[1] == NULL) { printf("\nUsage: %s kernel_source.cl kernel_function\n", argv[0]); exit(1); }
char *kernel_source;
if(load_program_source(argv[1], &kernel_source)) return 1;
printf("Building from OpenCL source: \t%s\n", argv[1]);
printf("Compile/query OpenCL_program:\t%s\n", argv[2]);
/* Create context and kernel program */
cl_context context = clCreateContext(0, 1, &devid, NULL, NULL, NULL);
cl_program pro = clCreateProgramWithSource(context, 1, (const char **)&kernel_source, NULL, NULL);
res = clBuildProgram(pro, 1, &devid, "-cl-fast-relaxed-math", NULL, NULL);
if(res != CL_SUCCESS){
printf("clBuildProgram failed: %d\n", res); char buf[0x10000];
clGetProgramBuildInfo(pro, devid, CL_PROGRAM_BUILD_LOG, 0x10000, buf, NULL);
printf("\n%s\n", buf); return(-1); }
cl_kernel kernelobj = clCreateKernel(pro, argv[2], &res); check_return(res);
/* Get the maximum work-group size for executing the kernel on the device */
size_t global, local;
res = clGetKernelWorkGroupInfo(kernelobj, devid, CL_KERNEL_WORK_GROUP_SIZE, sizeof(int), &local, NULL); check_return(res);
printf("CL_KERNEL_WORK_GROUP_SIZE\t%u\n", local);
res = clGetKernelWorkGroupInfo(kernelobj, devid, CL_KERNEL_LOCAL_MEM_SIZE, sizeof(cl_ulong), ¶m, NULL); check_return(res);
printf("CL_KERNEL_LOCAL_MEM_SIZE\t%ub\n", param);
cl_command_queue cmd_queue = clCreateCommandQueue(context, devid, CL_QUEUE_PROFILING_ENABLE, NULL);
if(cmd_queue == NULL) { printf("Compute device setup failed\n"); return(-1); }
local = 4;
int n = 2 * local; //num_group * local workgroup size
global = n;
int num_groups= global / local,
allocated_local= sizeof(data) * local +
sizeof(debug) * local;
data *DP __attribute__ ((aligned(16)));
DP = calloc(n, sizeof(data) *1);
debug *dbg __attribute__ ((aligned(16)));
dbg = calloc(n, sizeof(debug));
printf("global:%d, local:%d, (should be):%d groups\n", global, local, num_groups);
printf("structs size: %db, %db, %db\n", sizeof(data), sizeof(Elliptic_Curve), sizeof(inv256));
printf("sets:%d, total of %db needed, allocated _local: %db\n", n, n * sizeof(cl_uint4) *5 *4, allocated_local);
cl_mem cl_DP, cl_EC, cl_INV, DEBUG;
cl_DP = clCreateBuffer(context, CL_MEM_ALLOC_HOST_PTR, n * sizeof(data), NULL, &res); check_return(res);
cl_EC = clCreateBuffer(context, CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY, 1 * sizeof(Elliptic_Curve), NULL, &res); check_return(res); //_constant address space
cl_INV= clCreateBuffer(context, CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY, 1 * sizeof(u8) * 0x80, NULL, &res); check_return(res);
DEBUG = clCreateBuffer(context, CL_MEM_ALLOC_HOST_PTR | CL_MEM_WRITE_ONLY, n * sizeof(debug), NULL, &res); check_return(res);
Elliptic_Curve EC;
/*
Curve domain parameters, (test vectors)
-------------------------------------------------------------------------------------
p: c1c627e1638fdc8e24299bb041e4e23af4bb5427 is prime
a: c1c627e1638fdc8e24299bb041e4e23af4bb5424 divisor g = 62980
b: 877a6d84155a1de374b72d9f9d93b36bb563b2ab divisor g = 227169643
Gx: 010aff82b3ac72569ae645af3b527be133442131 divisor g = 32209245
Gy: 46b8ec1e6d71e5ecb549614887d57a287df573cc divisor g = 972
precomputed_per_curve_constants:
U: c1c627e1638fdc8e24299bb041e4e23af4bb5425
V: 3e39d81e9c702371dbd6644fbe1b1dc50b44abd9
already prepared mod p to test:
a: 07189f858e3f723890a66ec1079388ebd2ed509c
b: 6043379beb0dade6eed1e9d6de64f4a0c50639d4
gx: 5ef84aacf4f0ea6752f572d0741f40049f354dca
gy: 418c695435af6b3d4d7cbb72967395016ef67239
resulting point:
P.x: 01718f862ebe9423bd661a65355aa1c86ba330f8 program MUST got this point !!
P.y: 557e8ed53ffbfe2c990a121967b340f62e0e4fe2
taken mod p:
P.x: 41da1a8f74ff8d3f1ce20ef3e9d8865c96014fe3
P.y: 73ca143c9badedf2d9d3c7573307115ccfe04f13
*/
u8 *t;
t = _x_to_u8_buffer("c1c627e1638fdc8e24299bb041e4e23af4bb5427"); memcpy(EC.p, t, 20);
t = _x_to_u8_buffer("07189f858e3f723890a66ec1079388ebd2ed509c"); memcpy(EC.a, t, 20);
t = _x_to_u8_buffer("6043379beb0dade6eed1e9d6de64f4a0c50639d4"); memcpy(EC.b, t, 20);
t = _x_to_u8_buffer("5ef84aacf4f0ea6752f572d0741f40049f354dca"); memcpy(EC.Gx, t, 20);
t = _x_to_u8_buffer("418c695435af6b3d4d7cbb72967395016ef67239"); memcpy(EC.Gy, t, 20);
t = _x_to_u8_buffer("c1c627e1638fdc8e24299bb041e4e23af4bb5425"); memcpy(EC.U, t, 20);
t = _x_to_u8_buffer("3e39d81e9c702371dbd6644fbe1b1dc50b44abd9"); memcpy(EC.V, t, 20);
/* we need to map buffer now to load some k into data */
DP = clEnqueueMapBuffer(cmd_queue, cl_DP, CL_TRUE, CL_MAP_WRITE, 0, n * sizeof(data), 0, NULL, NULL, &res); check_return(res);
t = _x_to_u8_buffer("00542d46e7b3daac8aeb81e533873aabd6d74bb710");
for(u8 i = 0; i < n; i++) memcpy(DP[i].k, t, 21);
free(t);
//d for(u8 i = 0; i < n; i++) bn_print("", DP[i].k, 21, 1);
/* we can alter just a byte into a chosen k to verify that we'll get a different point! */
//DP[2].k[2] = 0x09;
//no res = clEnqueueWriteBuffer(cmd_queue, cl_DP, CL_TRUE, 0, n * sizeof(data), &DP, 0, NULL, NULL); check_return(res);
res = clEnqueueWriteBuffer(cmd_queue, cl_EC, CL_TRUE, 0, 1 * sizeof(Elliptic_Curve), &EC, 0, NULL, NULL); check_return(res);
res = clEnqueueWriteBuffer(cmd_queue, cl_INV, CL_TRUE, 0, 1 * sizeof(u8) * 0x80, &inv256, 0, NULL, NULL); check_return(res);
res = clSetKernelArg(kernelobj, 0, sizeof(cl_mem), &cl_DP); /* i/o buffer */
res|= clSetKernelArg(kernelobj, 1, sizeof(data) * local *1, NULL); //allocate space for __local in kernel (just this!) one * localsize
res|= clSetKernelArg(kernelobj, 2, sizeof(cl_mem), &cl_EC);
res|= clSetKernelArg(kernelobj, 3, sizeof(cl_mem), &cl_INV);
res|= clSetKernelArg(kernelobj, 4, sizeof(debug) * local *1, NULL); //allocate space for __local in kernel (just this!) one * localsize
res|= clSetKernelArg(kernelobj, 5, sizeof(cl_mem), &DEBUG); //this used to debug kernel output
check_return(res);
// printf("n:%d, total of %db needed, allocated _local: %db\n", n, n * sizeof(debug), allocated_local);
cl_event NDRangeEvent;
cl_ulong start, end;
/* Execute NDrange */
res = clEnqueueNDRangeKernel(cmd_queue, kernelobj, 1, NULL, &global, &local, 0, NULL, &NDRangeEvent); check_return(res);
// res = clEnqueueNDRangeKernel(cmd_queue, kernelobj, 1, NULL, &global, NULL, 0, NULL, &NDRangeEvent); check_return(res);
printf("Read back, Mapping buffer:\t%db\n", n * sizeof(data));
DP = clEnqueueMapBuffer(cmd_queue, cl_DP, CL_TRUE, CL_MAP_READ, 0, n * sizeof(data), 0, NULL, NULL, &res); check_return(res);
dbg =clEnqueueMapBuffer(cmd_queue, DEBUG, CL_TRUE, CL_MAP_READ, 0, n * sizeof(debug), 0, NULL, NULL, &res); check_return(res);
/* using clEnqueueReadBuffer template */
// res = clEnqueueReadBuffer(cmd_queue, ST, CL_TRUE, 0, sets * sizeof(cl_uint8), dbg, 0, NULL, NULL); check_return(res);
clFlush(cmd_queue);
clFinish(cmd_queue);
/* get NDRange execution time with internal ocl profiler */
res = clGetEventProfilingInfo(NDRangeEvent, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
res|= clGetEventProfilingInfo(NDRangeEvent, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
check_return(res);
printf("kernel execution time:\t\t%.2f ms\n", (float) ((end - start) /1000000)); //relative to NDRange call
printf("number of computes/sec:\t%.2f\n", (float) global *1000000 /((end - start)));
printf("i,\tgid\tlid0\tlsize0\tgid0/lsz0,\tgsz0,\tn_gr0,\tlid5,\toffset\n");
for(int i = 0; i < n; i++) {
// if(i %local == 0) {
printf("%d \t", i);
//printf("%u\t%u\t%u\t%u\t| %2u, %2u, %2u, %u\n", *p, *(p +1), *(p +2), *(p +3), *(p +4), *(p +5), *(p +6), *(p +7));
/* silence this doubled debug info
printf("%u\t%u\t%u\t%u\t| %2u, %2u, %2u, %u\n",
dbg[i].data[0], dbg[i].data[1], dbg[i].data[2], dbg[i].data[3],
dbg[i].data[4], dbg[i].data[5], dbg[i].data[6], dbg[i].data[7]);
*/
//printf("%d %d\n", P[i].dig, P[i].c);
bn_print("", DP[i].k, 21, 1);
bn_print("", DP[i].rx, 20, 0); bn_print(" ", DP[i].ry, 20, 1);
printf("%u(/%u) = %u*%u(/%u) +%u, offset:%u, stride:%u\n",
DP[i].pad[0], DP[i].pad[1], DP[i].pad[2], DP[i].pad[3],
DP[i].pad[4], DP[i].pad[5], DP[i].pad[6], DP[i].pad[7]);
// }
}
/* Release OpenCL stuff, free the rest */
clReleaseMemObject(cl_DP);
clReleaseMemObject(cl_EC);
clReleaseMemObject(cl_INV);
clReleaseMemObject(DEBUG);
clReleaseKernel(kernelobj);
clReleaseProgram(pro);
clReleaseCommandQueue(cmd_queue);
clReleaseContext(context);
free(kernel_source);
puts("Done!");
return 0;
}
