Implemented VM mode for OpenCL RandomX.

src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.cpp

@@ -27,6 +27,15 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
+void xmrig::Blake2bHashRegistersKernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t gthreads        = threads;
+    static const size_t lthreads = 64;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
 // __kernel void blake2b_hash_registers_32(__global void *out, __global const void* in, uint inStrideBytes)
 // __kernel void blake2b_hash_registers_64(__global void *out, __global const void* in, uint inStrideBytes)
 void xmrig::Blake2bHashRegistersKernel::setArgs(cl_mem out, cl_mem in, uint32_t inStrideBytes)

src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h

@@ -37,6 +37,7 @@ class Blake2bHashRegistersKernel : public OclKernel
 public:
     inline Blake2bHashRegistersKernel(cl_program program, const char *name) : OclKernel(program, name) {}
 
+    void enqueue(cl_command_queue queue, size_t threads);
     void setArgs(cl_mem out, cl_mem in, uint32_t inStrideBytes);
 };
 

src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp

@@ -27,9 +27,32 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
+void xmrig::Blake2bInitialHashKernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t gthreads        = threads;
+    static const size_t lthreads = 64;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
 // __kernel void blake2b_initial_hash(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce)
 void xmrig::Blake2bInitialHashKernel::setArgs(cl_mem out, cl_mem blockTemplate)
 {
     setArg(0, sizeof(cl_mem), &out);
     setArg(1, sizeof(cl_mem), &blockTemplate);
 }
+
+
+void xmrig::Blake2bInitialHashKernel::setBlobSize(size_t size)
+{
+    const uint32_t s = size;
+
+    setArg(2, sizeof(uint32_t), &s);
+}
+
+
+void xmrig::Blake2bInitialHashKernel::setNonce(uint32_t nonce)
+{
+    setArg(3, sizeof(uint32_t), &nonce);
+}

src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h

@@ -37,7 +37,10 @@ class Blake2bInitialHashKernel : public OclKernel
 public:
     inline Blake2bInitialHashKernel(cl_program program) : OclKernel(program, "blake2b_initial_hash") {}
 
+    void enqueue(cl_command_queue queue, size_t threads);
     void setArgs(cl_mem out, cl_mem blockTemplate);
+    void setBlobSize(size_t size);
+    void setNonce(uint32_t nonce);
 };
 
 

src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp

@@ -27,6 +27,15 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
+void xmrig::ExecuteVmKernel::enqueue(cl_command_queue queue, size_t threads, size_t worksize)
+{
+    const size_t gthreads = (worksize == 16) ? (threads * 16) : (threads * 8);
+    const size_t lthreads = (worksize == 16) ? 32 : 16;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
 // __kernel void execute_vm(__global void* vm_states, __global void* rounding, __global void* scratchpads, __global const void* dataset_ptr, uint32_t batch_size, uint32_t num_iterations, uint32_t first, uint32_t last)
 void xmrig::ExecuteVmKernel::setArgs(cl_mem vm_states, cl_mem rounding, cl_mem scratchpads, cl_mem dataset_ptr, uint32_t batch_size)
 {
@@ -36,3 +45,21 @@ void xmrig::ExecuteVmKernel::setArgs(cl_mem vm_states, cl_mem rounding, cl_mem s
     setArg(3, sizeof(cl_mem), &dataset_ptr);
     setArg(4, sizeof(uint32_t), &batch_size);
 }
+
+
+void xmrig::ExecuteVmKernel::setFirst(uint32_t first)
+{
+    setArg(6, sizeof(uint32_t), &first);
+}
+
+
+void xmrig::ExecuteVmKernel::setIterations(uint32_t num_iterations)
+{
+    setArg(5, sizeof(uint32_t), &num_iterations);
+}
+
+
+void xmrig::ExecuteVmKernel::setLast(uint32_t last)
+{
+    setArg(7, sizeof(uint32_t), &last);
+}

src/backend/opencl/kernels/rx/ExecuteVmKernel.h

@@ -37,7 +37,11 @@ class ExecuteVmKernel : public OclKernel
 public:
     inline ExecuteVmKernel(cl_program program) : OclKernel(program, "execute_vm") {}
 
+    void enqueue(cl_command_queue queue, size_t threads, size_t worksize);
     void setArgs(cl_mem vm_states, cl_mem rounding, cl_mem scratchpads, cl_mem dataset_ptr, uint32_t batch_size);
+    void setFirst(uint32_t first);
+    void setIterations(uint32_t num_iterations);
+    void setLast(uint32_t last);
 };
 
 

src/backend/opencl/kernels/rx/FillAesKernel.cpp

@@ -27,6 +27,15 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
+void xmrig::FillAesKernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t gthreads        = threads * 4;
+    static const size_t lthreads = 64;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
 // __kernel void fillAes1Rx4_scratchpad(__global void* state, __global void* out, uint batch_size, uint rx_version)
 // __kernel void fillAes4Rx4_entropy(__global void* state, __global void* out, uint batch_size, uint rx_version)
 void xmrig::FillAesKernel::setArgs(cl_mem state, cl_mem out, uint32_t batch_size, uint32_t rx_version)

src/backend/opencl/kernels/rx/FillAesKernel.h

@@ -37,6 +37,7 @@ class FillAesKernel : public OclKernel
 public:
     inline FillAesKernel(cl_program program, const char *name) : OclKernel(program, name) {}
 
+    void enqueue(cl_command_queue queue, size_t threads);
     void setArgs(cl_mem state, cl_mem out, uint32_t batch_size, uint32_t rx_version);
 };
 

src/backend/opencl/kernels/rx/FindSharesKernel.cpp

@@ -27,9 +27,30 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
+void xmrig::FindSharesKernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t gthreads        = threads;
+    static const size_t lthreads = 64;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
 // __kernel void find_shares(__global const uint64_t* hashes, uint64_t target, uint32_t start_nonce, __global uint32_t* shares)
 void xmrig::FindSharesKernel::setArgs(cl_mem hashes, cl_mem shares)
 {
     setArg(0, sizeof(cl_mem), &hashes);
     setArg(3, sizeof(cl_mem), &shares);
 }
+
+
+void xmrig::FindSharesKernel::setTarget(uint64_t target)
+{
+    setArg(1, sizeof(uint64_t), &target);
+}
+
+
+void xmrig::FindSharesKernel::setNonce(uint32_t nonce)
+{
+    setArg(2, sizeof(uint32_t), &nonce);
+}

src/backend/opencl/kernels/rx/FindSharesKernel.h

@@ -37,7 +37,10 @@ class FindSharesKernel : public OclKernel
 public:
     inline FindSharesKernel(cl_program program) : OclKernel(program, "find_shares") {}
 
+    void enqueue(cl_command_queue queue, size_t threads);
     void setArgs(cl_mem hashes, cl_mem shares);
+    void setTarget(uint64_t target);
+    void setNonce(uint32_t nonce);
 };
 
 

src/backend/opencl/kernels/rx/HashAesKernel.cpp

@@ -27,6 +27,15 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
+void xmrig::HashAesKernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t gthreads        = threads * 4;
+    static const size_t lthreads = 64;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
 // __kernel void hashAes1Rx4(__global const void* input, __global void* hash, uint hashOffsetBytes, uint hashStrideBytes, uint batch_size)
 void xmrig::HashAesKernel::setArgs(cl_mem input, cl_mem hash, uint32_t hashStrideBytes, uint32_t batch_size)
 {

src/backend/opencl/kernels/rx/HashAesKernel.h

@@ -37,6 +37,7 @@ class HashAesKernel : public OclKernel
 public:
     inline HashAesKernel(cl_program program) : OclKernel(program, "hashAes1Rx4") {}
 
+    void enqueue(cl_command_queue queue, size_t threads);
     void setArgs(cl_mem input, cl_mem hash, uint32_t hashStrideBytes, uint32_t batch_size);
 };
 

src/backend/opencl/kernels/rx/InitVmKernel.cpp

@@ -26,6 +26,24 @@
 #include "backend/opencl/kernels/rx/InitVmKernel.h"
 #include "backend/opencl/wrappers/OclLib.h"
 
+#include "base/io/log/Log.h"
+#include <thread>
+
+
+void xmrig::InitVmKernel::enqueue(cl_command_queue queue, size_t threads, uint32_t iteration)
+{
+    setArg(3, sizeof(uint32_t), &iteration);
+
+    const size_t gthreads        = threads * 8;
+    static const size_t lthreads = 32;
+
+//    LOG_WARN("%zu %zu %u", gthreads, lthreads, iteration);
+
+//    std::this_thread::sleep_for(std::chrono::milliseconds(500));
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
 
 // __kernel void init_vm(__global const void* entropy_data, __global void* vm_states, __global uint32_t* rounding, uint32_t iteration)
 void xmrig::InitVmKernel::setArgs(cl_mem entropy_data, cl_mem vm_states, cl_mem rounding)

src/backend/opencl/kernels/rx/InitVmKernel.h

@@ -37,6 +37,7 @@ class InitVmKernel : public OclKernel
 public:
     inline InitVmKernel(cl_program program) : OclKernel(program, "init_vm") {}
 
+    void enqueue(cl_command_queue queue, size_t threads, uint32_t iteration);
     void setArgs(cl_mem entropy_data, cl_mem vm_states, cl_mem rounding);
 };