From 734f142b4787c6ed9c2fc0b89a1155f84f0b0da8 Mon Sep 17 00:00:00 2001 From: SChernykh Date: Thu, 28 May 2020 20:27:25 +0200 Subject: [PATCH] KawPow fix for retarted AMD OpenCL compiler --- src/backend/opencl/cl/kawpow/kawpow.cl | 4 ++-- src/backend/opencl/cl/kawpow/kawpow_cl.h | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/src/backend/opencl/cl/kawpow/kawpow.cl b/src/backend/opencl/cl/kawpow/kawpow.cl index 81ab56ed..6821fb37 100644 --- a/src/backend/opencl/cl/kawpow/kawpow.cl +++ b/src/backend/opencl/cl/kawpow/kawpow.cl @@ -201,8 +201,8 @@ __kernel void progpow_search(__global dag_t const* g_dag, __global uint* job_blo // initialize mix for all lanes fill_mix(share[group_id].uint32s, lane_id, mix); -#pragma unroll 1 - for (uint32_t loop = 0; loop < PROGPOW_CNT_DAG; ++loop) + #pragma unroll 2 + for (uint32_t loop = 0; loop < PROGPOW_CNT_DAG; ++loop) { // global load if(lane_id == (loop % PROGPOW_LANES)) diff --git a/src/backend/opencl/cl/kawpow/kawpow_cl.h b/src/backend/opencl/cl/kawpow/kawpow_cl.h index 4bb45c0a..265857a5 100644 --- a/src/backend/opencl/cl/kawpow/kawpow_cl.h +++ b/src/backend/opencl/cl/kawpow/kawpow_cl.h @@ -143,7 +143,7 @@ static char kawpow_cl[5803] = { 0x65,0x32,0x5b,0x31,0x5d,0x3b,0x0a,0x7d,0x0a,0x62,0x61,0x72,0x72,0x69,0x65,0x72,0x28,0x43,0x4c,0x4b,0x5f,0x4c,0x4f,0x43,0x41,0x4c,0x5f,0x4d,0x45,0x4d,0x5f,0x46, 0x45,0x4e,0x43,0x45,0x29,0x3b,0x0a,0x66,0x69,0x6c,0x6c,0x5f,0x6d,0x69,0x78,0x28,0x73,0x68,0x61,0x72,0x65,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x2e, 0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x2c,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x2c,0x6d,0x69,0x78,0x29,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e, - 0x72,0x6f,0x6c,0x6c,0x20,0x31,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6c,0x6f,0x6f,0x70,0x3d,0x30,0x3b,0x20,0x6c,0x6f,0x6f, + 0x72,0x6f,0x6c,0x6c,0x20,0x32,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6c,0x6f,0x6f,0x70,0x3d,0x30,0x3b,0x20,0x6c,0x6f,0x6f, 0x70,0x3c,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x43,0x4e,0x54,0x5f,0x44,0x41,0x47,0x3b,0x20,0x2b,0x2b,0x6c,0x6f,0x6f,0x70,0x29,0x0a,0x7b,0x0a,0x69,0x66,0x28, 0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x3d,0x3d,0x28,0x6c,0x6f,0x6f,0x70,0x20,0x25,0x20,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x29,0x29, 0x0a,0x73,0x68,0x61,0x72,0x65,0x5b,0x30,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x3d,0x6d,0x69,0x78,0x5b,