diff --git a/CMakeLists.txt b/CMakeLists.txt
index acb62108..60651d62 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -8,6 +8,8 @@ option(WITH_TLS             "Enable OpenSSL support" ON)
 option(WITH_EMBEDDED_CONFIG "Enable internal embedded JSON config" OFF)
 option(WITH_CUDA            "Enable CUDA support" ON)
 option(WITH_OPENCL          "Enable OpenCL support" ON)
+option(WITH_OPENCL_DOUBLE_THREADS  "Enable dual threads for OpenCL jobs" ON)
+option(WITH_CUDA_DOUBLE_THREADS  "Enable dual threads for CUDA jobs" ON)
 
 include (CheckIncludeFile)
 include (cmake/cpu.cmake)
@@ -373,6 +375,9 @@ endif(ARCH STREQUAL "arm" OR ARCH STREQUAL "aarch64")
 if(WITH_OPENCL)
     add_definitions(-DWITH_OPENCL)
     find_package(OpenCL REQUIRED)
+    if(WITH_OPENCL_DOUBLE_THREADS)
+        add_definitions(-DPARALLEL_OPENCL)
+    endif()
     include_directories(${OpenCL_INCLUDE_DIR})
     add_library(opencl_hasher MODULE ${SOURCE_OPENCL_HASHER})
     set_target_properties(opencl_hasher
@@ -388,7 +393,7 @@ endif()
 if(WITH_CUDA)
     add_definitions(-DWITH_CUDA)
     find_package(CUDA REQUIRED)
-    if(NOT WIN32)
+    if(WITH_CUDA_DOUBLE_THREADS)
         add_definitions(-DPARALLEL_CUDA)
     endif()
     set(
diff --git a/src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLHasher.cpp b/src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLHasher.cpp
index d31bf69b..41d46c04 100755
--- a/src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLHasher.cpp
+++ b/src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLHasher.cpp
@@ -363,7 +363,14 @@ bool OpenCLHasher::setupDeviceInfo(OpenCLDeviceInfo *device, double intensity) {
         return false;
     }
 
-    device->queue = clCreateCommandQueue(device->context, device->device, CL_QUEUE_PROFILING_ENABLE, &error);
+    device->queue[0] = clCreateCommandQueue(device->context, device->device, NULL, &error);
+    if (error != CL_SUCCESS) {
+        device->error = error;
+        device->errorMessage = "Error getting device command queue.";
+        return false;
+    }
+
+    device->queue[1] = clCreateCommandQueue(device->context, device->device, NULL, &error);
     if (error != CL_SUCCESS) {
         device->error = error;
         device->errorMessage = "Error getting device command queue.";
@@ -398,26 +405,44 @@ bool OpenCLHasher::setupDeviceInfo(OpenCLDeviceInfo *device, double intensity) {
         return false;
     }
 
-    device->kernelPrehash = clCreateKernel(device->program, "prehash", &error);
+    device->kernelPrehash[0] = clCreateKernel(device->program, "prehash", &error);
     if (error != CL_SUCCESS) {
         device->error = error;
         device->errorMessage = "Error creating opencl prehash kernel for device.";
         return false;
     }
-    device->kernelFillBlocks = clCreateKernel(device->program, "fill_blocks", &error);
+    device->kernelPrehash[1] = clCreateKernel(device->program, "prehash", &error);
+    if (error != CL_SUCCESS) {
+        device->error = error;
+        device->errorMessage = "Error creating opencl prehash kernel for device.";
+        return false;
+    }
+    device->kernelFillBlocks[0] = clCreateKernel(device->program, "fill_blocks", &error);
     if (error != CL_SUCCESS) {
         device->error = error;
         device->errorMessage = "Error creating opencl main kernel for device.";
         return false;
     }
-    device->kernelPosthash = clCreateKernel(device->program, "posthash", &error);
+    device->kernelFillBlocks[1] = clCreateKernel(device->program, "fill_blocks", &error);
+    if (error != CL_SUCCESS) {
+        device->error = error;
+        device->errorMessage = "Error creating opencl main kernel for device.";
+        return false;
+    }
+    device->kernelPosthash[0] = clCreateKernel(device->program, "posthash", &error);
+    if (error != CL_SUCCESS) {
+        device->error = error;
+        device->errorMessage = "Error creating opencl posthash kernel for device.";
+        return false;
+    }
+    device->kernelPosthash[1] = clCreateKernel(device->program, "posthash", &error);
     if (error != CL_SUCCESS) {
         device->error = error;
         device->errorMessage = "Error creating opencl posthash kernel for device.";
         return false;
     }
 
-    device->profileInfo.threads_per_chunk = (uint32_t) (device->maxAllocableMemSize / device->profileInfo.profile->memSize);
+    device->profileInfo.threads_per_chunk = ((uint32_t) (device->maxAllocableMemSize / device->profileInfo.profile->memSize) / 2) * 2; // make it divisible by 2 to allow 2 hashes per wavefront
     size_t chunk_size = device->profileInfo.threads_per_chunk * device->profileInfo.profile->memSize;
 
     if (chunk_size == 0) {
@@ -650,7 +675,7 @@ bool OpenCLHasher::setupDeviceInfo(OpenCLDeviceInfo *device, double intensity) {
         refs[i] = device->profileInfo.profile->blockRefs[i * 3 + 1];
     }
 
-    error=clEnqueueWriteBuffer(device->queue, device->arguments.refs, CL_TRUE, 0, device->profileInfo.profile->blockRefsSize * sizeof(uint32_t), refs, 0, NULL, NULL);
+    error=clEnqueueWriteBuffer(device->queue[0], device->arguments.refs, CL_TRUE, 0, device->profileInfo.profile->blockRefsSize * sizeof(uint32_t), refs, 0, NULL, NULL);
     if(error != CL_SUCCESS) {
         device->error = error;
         device->errorMessage = "Error writing to gpu memory.";
@@ -668,7 +693,7 @@ bool OpenCLHasher::setupDeviceInfo(OpenCLDeviceInfo *device, double intensity) {
             }
         }
 
-        error=clEnqueueWriteBuffer(device->queue, device->arguments.idxs, CL_TRUE, 0, device->profileInfo.profile->blockRefsSize * sizeof(uint32_t), idxs, 0, NULL, NULL);
+        error=clEnqueueWriteBuffer(device->queue[0], device->arguments.idxs, CL_TRUE, 0, device->profileInfo.profile->blockRefsSize * sizeof(uint32_t), idxs, 0, NULL, NULL);
         if(error != CL_SUCCESS) {
             device->error = error;
             device->errorMessage = "Error writing to gpu memory.";
@@ -678,37 +703,60 @@ bool OpenCLHasher::setupDeviceInfo(OpenCLDeviceInfo *device, double intensity) {
         free(idxs);
     }
 
-    error=clEnqueueWriteBuffer(device->queue, device->arguments.segments, CL_TRUE, 0, device->profileInfo.profile->segCount * 3 * sizeof(uint32_t), device->profileInfo.profile->segments, 0, NULL, NULL);
+    error=clEnqueueWriteBuffer(device->queue[0], device->arguments.segments, CL_TRUE, 0, device->profileInfo.profile->segCount * 3 * sizeof(uint32_t), device->profileInfo.profile->segments, 0, NULL, NULL);
     if(error != CL_SUCCESS) {
         device->error = error;
         device->errorMessage = "Error writing to gpu memory.";
         return false;
     }
 
-	clSetKernelArg(device->kernelFillBlocks, 0, sizeof(device->arguments.memoryChunk_0), &device->arguments.memoryChunk_0);
-	clSetKernelArg(device->kernelFillBlocks, 1, sizeof(device->arguments.memoryChunk_1), &device->arguments.memoryChunk_1);
-	clSetKernelArg(device->kernelFillBlocks, 2, sizeof(device->arguments.memoryChunk_2), &device->arguments.memoryChunk_2);
-	clSetKernelArg(device->kernelFillBlocks, 3, sizeof(device->arguments.memoryChunk_3), &device->arguments.memoryChunk_3);
-	clSetKernelArg(device->kernelFillBlocks, 4, sizeof(device->arguments.memoryChunk_4), &device->arguments.memoryChunk_4);
-	clSetKernelArg(device->kernelFillBlocks, 5, sizeof(device->arguments.memoryChunk_5), &device->arguments.memoryChunk_5);
-    clSetKernelArg(device->kernelFillBlocks, 8, sizeof(device->arguments.refs), &device->arguments.refs);
-    if(device->profileInfo.profile->succesiveIdxs == 0)
-        clSetKernelArg(device->kernelFillBlocks, 9, sizeof(device->arguments.idxs), &device->arguments.idxs);
-    else
-        clSetKernelArg(device->kernelFillBlocks, 9, sizeof(cl_mem), NULL);
-	clSetKernelArg(device->kernelFillBlocks, 10, sizeof(device->arguments.segments), &device->arguments.segments);
-    clSetKernelArg(device->kernelFillBlocks, 11, sizeof(int32_t), &device->profileInfo.profile->memSize);
-    clSetKernelArg(device->kernelFillBlocks, 12, sizeof(int32_t), &device->profileInfo.profile->thrCost);
-    clSetKernelArg(device->kernelFillBlocks, 13, sizeof(int32_t), &device->profileInfo.profile->segSize);
-    clSetKernelArg(device->kernelFillBlocks, 14, sizeof(int32_t), &device->profileInfo.profile->segCount);
-    clSetKernelArg(device->kernelFillBlocks, 15, sizeof(int32_t), &device->profileInfo.threads_per_chunk);
-
-    clSetKernelArg(device->kernelPrehash, 2, sizeof(int32_t), &device->profileInfo.profile->memCost);
-    clSetKernelArg(device->kernelPrehash, 3, sizeof(int32_t), &device->profileInfo.profile->thrCost);
     int passes = device->profileInfo.profile->segCount / (4 * device->profileInfo.profile->thrCost);
-    clSetKernelArg(device->kernelPrehash, 4, sizeof(int32_t), &passes);
-    clSetKernelArg(device->kernelPrehash, 6, sizeof(int32_t), &device->profileInfo.profile->saltLen);
 
+    clSetKernelArg(device->kernelFillBlocks[0], 0, sizeof(device->arguments.memoryChunk_0), &device->arguments.memoryChunk_0);
+	clSetKernelArg(device->kernelFillBlocks[0], 1, sizeof(device->arguments.memoryChunk_1), &device->arguments.memoryChunk_1);
+	clSetKernelArg(device->kernelFillBlocks[0], 2, sizeof(device->arguments.memoryChunk_2), &device->arguments.memoryChunk_2);
+	clSetKernelArg(device->kernelFillBlocks[0], 3, sizeof(device->arguments.memoryChunk_3), &device->arguments.memoryChunk_3);
+	clSetKernelArg(device->kernelFillBlocks[0], 4, sizeof(device->arguments.memoryChunk_4), &device->arguments.memoryChunk_4);
+	clSetKernelArg(device->kernelFillBlocks[0], 5, sizeof(device->arguments.memoryChunk_5), &device->arguments.memoryChunk_5);
+    clSetKernelArg(device->kernelFillBlocks[0], 8, sizeof(device->arguments.refs), &device->arguments.refs);
+    if(device->profileInfo.profile->succesiveIdxs == 0)
+        clSetKernelArg(device->kernelFillBlocks[0], 9, sizeof(device->arguments.idxs), &device->arguments.idxs);
+    else
+        clSetKernelArg(device->kernelFillBlocks[0], 9, sizeof(cl_mem), NULL);
+	clSetKernelArg(device->kernelFillBlocks[0], 10, sizeof(device->arguments.segments), &device->arguments.segments);
+    clSetKernelArg(device->kernelFillBlocks[0], 11, sizeof(int32_t), &device->profileInfo.profile->memSize);
+    clSetKernelArg(device->kernelFillBlocks[0], 12, sizeof(int32_t), &device->profileInfo.profile->thrCost);
+    clSetKernelArg(device->kernelFillBlocks[0], 13, sizeof(int32_t), &device->profileInfo.profile->segSize);
+    clSetKernelArg(device->kernelFillBlocks[0], 14, sizeof(int32_t), &device->profileInfo.profile->segCount);
+    clSetKernelArg(device->kernelFillBlocks[0], 15, sizeof(int32_t), &device->profileInfo.threads_per_chunk);
+
+    clSetKernelArg(device->kernelPrehash[0], 2, sizeof(int32_t), &device->profileInfo.profile->memCost);
+    clSetKernelArg(device->kernelPrehash[0], 3, sizeof(int32_t), &device->profileInfo.profile->thrCost);
+    clSetKernelArg(device->kernelPrehash[0], 4, sizeof(int32_t), &passes);
+    clSetKernelArg(device->kernelPrehash[0], 6, sizeof(int32_t), &device->profileInfo.profile->saltLen);
+
+    clSetKernelArg(device->kernelFillBlocks[1], 0, sizeof(device->arguments.memoryChunk_0), &device->arguments.memoryChunk_0);
+    clSetKernelArg(device->kernelFillBlocks[1], 1, sizeof(device->arguments.memoryChunk_1), &device->arguments.memoryChunk_1);
+    clSetKernelArg(device->kernelFillBlocks[1], 2, sizeof(device->arguments.memoryChunk_2), &device->arguments.memoryChunk_2);
+    clSetKernelArg(device->kernelFillBlocks[1], 3, sizeof(device->arguments.memoryChunk_3), &device->arguments.memoryChunk_3);
+    clSetKernelArg(device->kernelFillBlocks[1], 4, sizeof(device->arguments.memoryChunk_4), &device->arguments.memoryChunk_4);
+    clSetKernelArg(device->kernelFillBlocks[1], 5, sizeof(device->arguments.memoryChunk_5), &device->arguments.memoryChunk_5);
+    clSetKernelArg(device->kernelFillBlocks[1], 8, sizeof(device->arguments.refs), &device->arguments.refs);
+    if(device->profileInfo.profile->succesiveIdxs == 0)
+        clSetKernelArg(device->kernelFillBlocks[1], 9, sizeof(device->arguments.idxs), &device->arguments.idxs);
+    else
+        clSetKernelArg(device->kernelFillBlocks[1], 9, sizeof(cl_mem), NULL);
+    clSetKernelArg(device->kernelFillBlocks[1], 10, sizeof(device->arguments.segments), &device->arguments.segments);
+    clSetKernelArg(device->kernelFillBlocks[1], 11, sizeof(int32_t), &device->profileInfo.profile->memSize);
+    clSetKernelArg(device->kernelFillBlocks[1], 12, sizeof(int32_t), &device->profileInfo.profile->thrCost);
+    clSetKernelArg(device->kernelFillBlocks[1], 13, sizeof(int32_t), &device->profileInfo.profile->segSize);
+    clSetKernelArg(device->kernelFillBlocks[1], 14, sizeof(int32_t), &device->profileInfo.profile->segCount);
+    clSetKernelArg(device->kernelFillBlocks[1], 15, sizeof(int32_t), &device->profileInfo.threads_per_chunk);
+
+    clSetKernelArg(device->kernelPrehash[1], 2, sizeof(int32_t), &device->profileInfo.profile->memCost);
+    clSetKernelArg(device->kernelPrehash[1], 3, sizeof(int32_t), &device->profileInfo.profile->thrCost);
+    clSetKernelArg(device->kernelPrehash[1], 4, sizeof(int32_t), &passes);
+    clSetKernelArg(device->kernelPrehash[1], 6, sizeof(int32_t), &device->profileInfo.profile->saltLen);
     return true;
 }
 
@@ -724,29 +772,29 @@ bool opencl_kernel_prehasher(void *memory, int threads, Argon2Profile *profile,
     size_t total_work_items = sessions * 4 * ceil(threads / hashes_per_block);
     size_t local_work_items = sessions * 4;
 
-    device->deviceLock.lock();
+    gpumgmt_thread->lock();
 
-    error = clEnqueueWriteBuffer(device->queue, device->arguments.preseedMemory[gpumgmt_thread->threadId],
+    error = clEnqueueWriteBuffer(device->queue[gpumgmt_thread->threadId], device->arguments.preseedMemory[gpumgmt_thread->threadId],
                                  CL_FALSE, 0, gpumgmt_thread->hashData.inSize, memory, 0, NULL, NULL);
     if (error != CL_SUCCESS) {
         device->error = error;
         device->errorMessage = "Error writing to gpu memory.";
-        device->deviceLock.unlock();
+        gpumgmt_thread->unlock();
         return false;
     }
 
     int inSizeInInt = gpumgmt_thread->hashData.inSize / 4;
-    clSetKernelArg(device->kernelPrehash, 0, sizeof(device->arguments.preseedMemory[gpumgmt_thread->threadId]), &device->arguments.preseedMemory[gpumgmt_thread->threadId]);
-    clSetKernelArg(device->kernelPrehash, 1, sizeof(device->arguments.seedMemory[gpumgmt_thread->threadId]), &device->arguments.seedMemory[gpumgmt_thread->threadId]);
-    clSetKernelArg(device->kernelPrehash, 5, sizeof(int), &inSizeInInt);
-    clSetKernelArg(device->kernelPrehash, 7, sizeof(int), &threads);
-    clSetKernelArg(device->kernelPrehash, 8, sessions * sizeof(cl_ulong) * 76, NULL); // (preseed size is 16 ulongs = 128 bytes)
+    clSetKernelArg(device->kernelPrehash[gpumgmt_thread->threadId], 0, sizeof(device->arguments.preseedMemory[gpumgmt_thread->threadId]), &device->arguments.preseedMemory[gpumgmt_thread->threadId]);
+    clSetKernelArg(device->kernelPrehash[gpumgmt_thread->threadId], 1, sizeof(device->arguments.seedMemory[gpumgmt_thread->threadId]), &device->arguments.seedMemory[gpumgmt_thread->threadId]);
+    clSetKernelArg(device->kernelPrehash[gpumgmt_thread->threadId], 5, sizeof(int), &inSizeInInt);
+    clSetKernelArg(device->kernelPrehash[gpumgmt_thread->threadId], 7, sizeof(int), &threads);
+    clSetKernelArg(device->kernelPrehash[gpumgmt_thread->threadId], 8, sessions * sizeof(cl_ulong) * 76, NULL); // (preseed size is 16 ulongs = 128 bytes)
 
-    error=clEnqueueNDRangeKernel(device->queue, device->kernelPrehash, 1, NULL, &total_work_items, &local_work_items, 0, NULL, NULL);
+    error=clEnqueueNDRangeKernel(device->queue[gpumgmt_thread->threadId], device->kernelPrehash[gpumgmt_thread->threadId], 1, NULL, &total_work_items, &local_work_items, 0, NULL, NULL);
     if(error != CL_SUCCESS) {
         device->error = error;
         device->errorMessage = "Error running the kernel.";
-        device->deviceLock.unlock();
+        gpumgmt_thread->unlock();
         return false;
     }
 
@@ -764,15 +812,16 @@ void *opencl_kernel_filler(int threads, Argon2Profile *profile, void *user_data)
 
     size_t shared_mem = 2 * profile->thrCost * ARGON2_QWORDS_IN_BLOCK;
 
-    clSetKernelArg(device->kernelFillBlocks, 6, sizeof(device->arguments.seedMemory[gpumgmt_thread->threadId]), &device->arguments.seedMemory[gpumgmt_thread->threadId]);
-    clSetKernelArg(device->kernelFillBlocks, 7, sizeof(device->arguments.outMemory[gpumgmt_thread->threadId]), &device->arguments.outMemory[gpumgmt_thread->threadId]);
-    clSetKernelArg(device->kernelFillBlocks, 16, sizeof(cl_ulong) * shared_mem, NULL);
+    clSetKernelArg(device->kernelFillBlocks[gpumgmt_thread->threadId], 6, sizeof(device->arguments.seedMemory[gpumgmt_thread->threadId]), &device->arguments.seedMemory[gpumgmt_thread->threadId]);
+    clSetKernelArg(device->kernelFillBlocks[gpumgmt_thread->threadId], 7, sizeof(device->arguments.outMemory[gpumgmt_thread->threadId]), &device->arguments.outMemory[gpumgmt_thread->threadId]);
+    clSetKernelArg(device->kernelFillBlocks[gpumgmt_thread->threadId], 16, sizeof(int), &(gpumgmt_thread->threadsIdx));
+    clSetKernelArg(device->kernelFillBlocks[gpumgmt_thread->threadId], 17, sizeof(cl_ulong) * shared_mem, NULL);
 
-    error=clEnqueueNDRangeKernel(device->queue, device->kernelFillBlocks, 1, NULL, &total_work_items, &local_work_items, 0, NULL, NULL);
+    error=clEnqueueNDRangeKernel(device->queue[gpumgmt_thread->threadId], device->kernelFillBlocks[gpumgmt_thread->threadId], 1, NULL, &total_work_items, &local_work_items, 0, NULL, NULL);
     if(error != CL_SUCCESS) {
         device->error = error;
         device->errorMessage = "Error running the kernel.";
-        device->deviceLock.unlock();
+        gpumgmt_thread->unlock();
         return NULL;
     }
 
@@ -788,36 +837,36 @@ bool opencl_kernel_posthasher(void *memory, int threads, Argon2Profile *profile,
     size_t total_work_items = threads * 4;
     size_t local_work_items = 4;
 
-    clSetKernelArg(device->kernelPosthash, 0, sizeof(device->arguments.hashMemory[gpumgmt_thread->threadId]), &device->arguments.hashMemory[gpumgmt_thread->threadId]);
-    clSetKernelArg(device->kernelPosthash, 1, sizeof(device->arguments.outMemory[gpumgmt_thread->threadId]), &device->arguments.outMemory[gpumgmt_thread->threadId]);
-    clSetKernelArg(device->kernelPosthash, 2, sizeof(device->arguments.preseedMemory[gpumgmt_thread->threadId]), &device->arguments.preseedMemory[gpumgmt_thread->threadId]);
-    clSetKernelArg(device->kernelPosthash, 3, sizeof(cl_ulong) * 60, NULL);
+    clSetKernelArg(device->kernelPosthash[gpumgmt_thread->threadId], 0, sizeof(device->arguments.hashMemory[gpumgmt_thread->threadId]), &device->arguments.hashMemory[gpumgmt_thread->threadId]);
+    clSetKernelArg(device->kernelPosthash[gpumgmt_thread->threadId], 1, sizeof(device->arguments.outMemory[gpumgmt_thread->threadId]), &device->arguments.outMemory[gpumgmt_thread->threadId]);
+    clSetKernelArg(device->kernelPosthash[gpumgmt_thread->threadId], 2, sizeof(device->arguments.preseedMemory[gpumgmt_thread->threadId]), &device->arguments.preseedMemory[gpumgmt_thread->threadId]);
+    clSetKernelArg(device->kernelPosthash[gpumgmt_thread->threadId], 3, sizeof(cl_ulong) * 60, NULL);
 
-    error=clEnqueueNDRangeKernel(device->queue, device->kernelPosthash, 1, NULL, &total_work_items, &local_work_items, 0, NULL, NULL);
+    error=clEnqueueNDRangeKernel(device->queue[gpumgmt_thread->threadId], device->kernelPosthash[gpumgmt_thread->threadId], 1, NULL, &total_work_items, &local_work_items, 0, NULL, NULL);
     if(error != CL_SUCCESS) {
         device->error = error;
         device->errorMessage = "Error running the kernel.";
-        device->deviceLock.unlock();
+        gpumgmt_thread->unlock();
         return false;
     }
 
-    error = clEnqueueReadBuffer(device->queue, device->arguments.hashMemory[gpumgmt_thread->threadId], CL_FALSE, 0, threads * (xmrig::ARGON2_HASHLEN + 4), memory, 0, NULL, NULL);
+    error = clEnqueueReadBuffer(device->queue[gpumgmt_thread->threadId], device->arguments.hashMemory[gpumgmt_thread->threadId], CL_FALSE, 0, threads * (xmrig::ARGON2_HASHLEN + 4), memory, 0, NULL, NULL);
     if (error != CL_SUCCESS) {
         device->error = error;
         device->errorMessage = "Error reading gpu memory.";
-        device->deviceLock.unlock();
+        gpumgmt_thread->unlock();
         return false;
     }
 
-    error=clFinish(device->queue);
+    error=clFinish(device->queue[gpumgmt_thread->threadId]);
     if(error != CL_SUCCESS) {
         device->error = error;
         device->errorMessage = "Error flushing GPU queue.";
-        device->deviceLock.unlock();
+        gpumgmt_thread->unlock();
         return false;
     }
 
-    device->deviceLock.unlock();
+    gpumgmt_thread->unlock();
 
     return true;
 }
@@ -831,9 +880,24 @@ void OpenCLHasher::buildThreadData() {
             OpenCLGpuMgmtThreadData &thread_data = m_threadData[i * 2 + threadId];
             thread_data.device = device;
             thread_data.threadId = threadId;
+
+#ifdef PARALLEL_OPENCL
+            if(threadId == 0) {
+                thread_data.threadsIdx = 0;
+                thread_data.threads = device->profileInfo.threads / 2;
+            }
+            else {
+                thread_data.threadsIdx = device->profileInfo.threads / 2;
+                thread_data.threads = device->profileInfo.threads - thread_data.threadsIdx;
+            }
+#else
+            thread_data.threadsIdx = 0;
+            thread_data.threads = device->profileInfo.threads;
+#endif
+
             thread_data.argon2 = new Argon2(opencl_kernel_prehasher, opencl_kernel_filler, opencl_kernel_posthasher,
                                             nullptr, &thread_data);
-            thread_data.argon2->setThreads(device->profileInfo.threads);
+            thread_data.argon2->setThreads(thread_data.threads);
             thread_data.hashData.outSize = xmrig::ARGON2_HASHLEN + 4;
         }
     }
@@ -851,7 +915,7 @@ int OpenCLHasher::compute(int threadIdx, uint8_t *input, size_t size, uint8_t *o
     }
 
     uint32_t *nonce = ((uint32_t *)(((uint8_t*)threadData.hashData.input) + 39));
-    (*nonce) += threadData.device->profileInfo.threads;
+    (*nonce) += threadData.threads;
 
     return hashCount;
 }
@@ -878,11 +942,15 @@ void OpenCLHasher::cleanup() {
             clReleaseMemObject((*it)->arguments.hashMemory[0]);
             clReleaseMemObject((*it)->arguments.hashMemory[1]);
 
-            clReleaseKernel((*it)->kernelPrehash);
-            clReleaseKernel((*it)->kernelFillBlocks);
-            clReleaseKernel((*it)->kernelPosthash);
+            clReleaseKernel((*it)->kernelPrehash[0]);
+            clReleaseKernel((*it)->kernelPrehash[1]);
+            clReleaseKernel((*it)->kernelFillBlocks[0]);
+            clReleaseKernel((*it)->kernelFillBlocks[1]);
+            clReleaseKernel((*it)->kernelPosthash[0]);
+            clReleaseKernel((*it)->kernelPosthash[1]);
 			clReleaseProgram((*it)->program);
-			clReleaseCommandQueue((*it)->queue);
+            clReleaseCommandQueue((*it)->queue[0]);
+            clReleaseCommandQueue((*it)->queue[1]);
 			clReleaseContext((*it)->context);
 		}
         clReleaseDevice((*it)->device);
@@ -892,13 +960,8 @@ void OpenCLHasher::cleanup() {
 }
 
 size_t OpenCLHasher::parallelism(int workerIdx) {
-    // there are 2 computing threads per device, so divide by 2 to get device index
-    workerIdx /= 2;
-
-    if(workerIdx < 0 || workerIdx > m_enabledDevices.size())
-        return 0;
-
-    return m_enabledDevices[workerIdx]->profileInfo.threads;
+    OpenCLGpuMgmtThreadData &threadData = m_threadData[workerIdx];
+    return threadData.threads;
 }
 
 size_t OpenCLHasher::deviceCount() {
diff --git a/src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLHasher.h b/src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLHasher.h
index 3b688272..b6929047 100755
--- a/src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLHasher.h
+++ b/src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLHasher.h
@@ -51,12 +51,12 @@ struct OpenCLDeviceInfo {
     cl_platform_id platform;
     cl_device_id device;
     cl_context context;
-    cl_command_queue queue;
+    cl_command_queue queue[2];
 
     cl_program program;
-    cl_kernel kernelPrehash;
-    cl_kernel kernelFillBlocks;
-    cl_kernel kernelPosthash;
+    cl_kernel kernelPrehash[2];
+    cl_kernel kernelFillBlocks[2];
+    cl_kernel kernelPosthash[2];
 
     int deviceIndex;
 
@@ -75,10 +75,23 @@ struct OpenCLDeviceInfo {
 };
 
 struct OpenCLGpuMgmtThreadData {
+    void lock() {
+#ifndef PARALLEL_OPENCL
+        device->deviceLock.lock();
+#endif
+    }
+
+    void unlock() {
+#ifndef PARALLEL_OPENCL
+        device->deviceLock.unlock();
+#endif
+    }
     int threadId;
     OpenCLDeviceInfo *device;
     Argon2 *argon2;
     HashData hashData;
+    int threads;
+    int threadsIdx;
 };
 
 class OpenCLHasher : public Hasher {
diff --git a/src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLKernel.cpp b/src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLKernel.cpp
index fad3b886..0f14e1cf 100644
--- a/src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLKernel.cpp
+++ b/src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLKernel.cpp
@@ -767,11 +767,13 @@ __kernel void fill_blocks(__global ulong *chunk_0,
                         int seg_length,
                         int seg_count,
 						int threads_per_chunk,
+                        int thread_idx,
                         __local ulong *scratchpad) { // lanes * BLOCK_SIZE_ULONG
     ulong4 tmp;
 	ulong a, b, c, d;
 
 	int hash_base = get_group_id(0) * 2;
+	int mem_hash = hash_base + thread_idx;
 	int local_id = get_local_id(0);
 
     int hash_idx = (local_id / THREADS_PER_LANE) % 2;
@@ -789,8 +791,8 @@ __kernel void fill_blocks(__global ulong *chunk_0,
 	chunks[3] = (ulong)chunk_3;
 	chunks[4] = (ulong)chunk_4;
 	chunks[5] = (ulong)chunk_5;
-	int chunk_index = hash_base / threads_per_chunk;
-	int chunk_offset = hash_base - chunk_index * threads_per_chunk;
+	int chunk_index = mem_hash / threads_per_chunk;
+	int chunk_offset = mem_hash - chunk_index * threads_per_chunk;
 	__global ulong *memory = (__global ulong *)chunks[chunk_index] + chunk_offset * (memsize / 8);
 
 	int i1_0 = offsets_round_1[id][0];