Skip to content

xor ID into key_root instead of dereferencing #29

@webmaster128

Description

@webmaster128

In a fork of this project (other coin), the following change made the difference between working an non-working (no GPU results after thousands of %) on a NVIDIA Corporation GeForce GTX 1080, OpenCL 1.2.

diff --git a/src/opencl/entry.cl b/src/opencl/entry.cl
index fbd7a4f..47fb738 100644
--- a/src/opencl/entry.cl
+++ b/src/opencl/entry.cl
@@ -23,12 +23,20 @@ __kernel void generate_pubkey(
        ulong max_address_value,
        uchar generate_key_type
 ) {
-       int const thread = get_global_id (0);
        uchar key_material[32];
        for (size_t i = 0; i < 32; i++) {
                key_material[i] = key_root[i];
        }
-       *((size_t *) key_material) += thread;
+
+       const uint64_t thread_id = get_global_id(0);
+       key_material[24] ^= (thread_id >> (7*8)) & 0xFF;
+       key_material[25] ^= (thread_id >> (6*8)) & 0xFF;
+       key_material[26] ^= (thread_id >> (5*8)) & 0xFF;
+       key_material[27] ^= (thread_id >> (4*8)) & 0xFF;
+       key_material[28] ^= (thread_id >> (3*8)) & 0xFF;
+       key_material[29] ^= (thread_id >> (2*8)) & 0xFF;
+       key_material[30] ^= (thread_id >> (1*8)) & 0xFF;
+       key_material[31] ^= (thread_id >> (0*8)) & 0xFF;

My suggestion is that when dereferencing and assigning to key_material, the compiler does not understand that key_material depends on get_global_id.

Do you want to have this patch on master to avoid future surprises? Unfortunatly I cannot reproduce master by the patch because of #28.

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions