Conceal (CCX) support

src/backend/opencl/cl/cn/algorithm.cl

@@ -16,16 +16,17 @@
 #define ALGO_CN_HEAVY_XHV   15
 #define ALGO_CN_PICO_0      16
 #define ALGO_CN_PICO_TLO    17
-#define ALGO_RX_0           18
-#define ALGO_RX_WOW         19
-#define ALGO_RX_LOKI        20
-#define ALGO_RX_ARQMA       21
-#define ALGO_RX_SFX         22
-#define ALGO_RX_KEVA        23
-#define ALGO_AR2_CHUKWA     24
-#define ALGO_AR2_WRKZ       25
-#define ALGO_ASTROBWT_DERO  26
-#define ALGO_KAWPOW_RVN     27
+#define ALGO_CN_CCX         18
+#define ALGO_RX_0           19
+#define ALGO_RX_WOW         20
+#define ALGO_RX_LOKI        21
+#define ALGO_RX_ARQMA       22
+#define ALGO_RX_SFX         23
+#define ALGO_RX_KEVA        24
+#define ALGO_AR2_CHUKWA     25
+#define ALGO_AR2_WRKZ       26
+#define ALGO_ASTROBWT_DERO  27
+#define ALGO_KAWPOW_RVN     28
 
 #define FAMILY_UNKNOWN      0
 #define FAMILY_CN           1

src/backend/opencl/cl/cn/cryptonight.cl

@@ -253,11 +253,34 @@ __kernel void cn1(__global ulong *input, __global uint4 *Scratchpad, __global ul
     {
         uint idx0 = a[0];
 
+#       if (ALGO == ALGO_CN_CCX)
+        float4 conc_var = (float4)(0.0f);
+        const uint4 conc_t = (uint4)(0x807FFFFFU);
+        const uint4 conc_u = (uint4)(0x40000000U);
+        const uint4 conc_v = (uint4)(0x4DFFFFFFU);
+#       endif
+
         #pragma unroll CN_UNROLL
         for (int i = 0; i < ITERATIONS; ++i) {
             ulong c[2];
 
             ((uint4 *)c)[0] = Scratchpad[IDX((idx0 & MASK) >> 4)];
+
+#           if (ALGO == ALGO_CN_CCX)
+            {
+                float4 r = convert_float4_rte(((int4 *)c)[0]) + conc_var;
+                r = r * r * r;
+                r = as_float4((as_uint4(r) & conc_t) | conc_u);
+
+                float4 c_old = conc_var;
+                conc_var += r;
+
+                c_old = as_float4((as_uint4(c_old) & conc_t) | conc_u);
+
+                ((int4 *)c)[0] ^= convert_int4_rtz(c_old * as_float4(conc_v));
+            }
+#           endif
+
             ((uint4 *)c)[0] = AES_Round_Two_Tables(AES0, AES1, ((uint4 *)c)[0], ((uint4 *)a)[0]);
 
             Scratchpad[IDX((idx0 & MASK) >> 4)] = b_x ^ ((uint4 *)c)[0];