API tools faq
paste
image28

ethash.cl nvidia LHR possible unlock for card with 2:1 TFlop 16bit performance over 32bit

image28
Feb 28th, 2022 (edited)
977
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
C 17.58 KB | None | 0 0
raw download clone embed print report
  1. /*
  2. Description:
  3. The unlock works by removing unions, which use the largest variable type in the union which was 64bits.
  4. 64bits had the speed clamped on LHR gpu's.
  5. Some other LHR gpu's also clamp 16bit performance (this wont work on them).
  6. For those gpu's see my other file on pastebin.
  7.  
  8. Still room for more improvements.
  9.  
  10. Cheers,
  11. imagekiwi/image28
  12. */
  13.  
  14. #define OPENCL_PLATFORM_UNKNOWN 0
  15. #define OPENCL_PLATFORM_AMD 1
  16. #define OPENCL_PLATFORM_CLOVER 2
  17. #define OPENCL_PLATFORM_NVIDIA 3
  18. #define OPENCL_PLATFORM_INTEL 4
  19. /*
  20. #define WORKSIZE 256
  21. #define MAX_OUTPUTS 4
  22. #define ACCESSES 64
  23. #define PLATFORM 1
  24. #define COMPUTE 1
  25. */
  26.  
  27. #define CHUNKSIZE 64
  28.  
  29. #ifdef cl_clang_storage_class_specifiers
  30. #pragma OPENCL EXTENSION cl_clang_storage_class_specifiers : enable
  31. #endif
  32.  
  33. #if defined(cl_amd_media_ops)
  34.     #if PLATFORM == OPENCL_PLATFORM_CLOVER
  35.         uint2 amd_bitalign(uint2 src0, uint2 src1, uint2 src2)
  36.         {
  37.             uint2 dst;
  38.             __asm(
  39.                 "v_alignbit_b32 %0, %2, %3, %4\n"
  40.                 "v_alignbit_b32 %1, %5, %6, %7"
  41.                 : "=v"(dst.x), "=v"(dst.y)
  42.                 : "v"(src0.x), "v"(src1.x), "v"(src2.x), "v"(src0.y), "v"(src1.y), "v"(src2.y));
  43.             return dst;
  44.         }
  45.     #endif
  46.  
  47.     #pragma OPENCL EXTENSION cl_amd_media_ops : enable
  48. #elif defined(cl_nv_pragma_unroll)
  49.     uint amd_bitalign(uint src0, uint src1, uint src2)
  50.     {
  51.         uint dest;
  52.         asm("shf.r.wrap.b32 %0, %2, %1, %3;" : "=r"(dest) : "r"(src0), "r"(src1), "r"(src2));
  53.         return dest;
  54.     }
  55. #else
  56.     #define amd_bitalign(src0, src1, src2) \
  57.         ((uint)(((((ulong)(src0)) << 32) | (ulong)(src1)) >> ((src2)&31)))
  58. #endif
  59.  
  60. #define EndianSwap(n) (rotate(n & 0x00FF00FF, 24U)|(rotate(n, 8U) & 0x00FF00FF)
  61.  
  62. #if WORKSIZE % 4 != 0
  63.     #error "WORKSIZE has to be a multiple of 4"
  64. #endif
  65.  
  66. #define FNV_PRIME 0x01000193U // 2^24+403
  67.  
  68. static __constant uint2 const Keccak_f1600_RC[24] = {
  69.     (uint2)(0x00000001, 0x00000000),
  70.     (uint2)(0x00008082, 0x00000000),
  71.     (uint2)(0x0000808a, 0x80000000),
  72.     (uint2)(0x80008000, 0x80000000),
  73.     (uint2)(0x0000808b, 0x00000000),
  74.     (uint2)(0x80000001, 0x00000000),
  75.     (uint2)(0x80008081, 0x80000000),
  76.     (uint2)(0x00008009, 0x80000000),
  77.     (uint2)(0x0000008a, 0x00000000),
  78.     (uint2)(0x00000088, 0x00000000),
  79.     (uint2)(0x80008009, 0x00000000),
  80.     (uint2)(0x8000000a, 0x00000000),
  81.     (uint2)(0x8000808b, 0x00000000),
  82.     (uint2)(0x0000008b, 0x80000000),
  83.     (uint2)(0x00008089, 0x80000000),
  84.     (uint2)(0x00008003, 0x80000000),
  85.     (uint2)(0x00008002, 0x80000000),
  86.     (uint2)(0x00000080, 0x80000000),
  87.     (uint2)(0x0000800a, 0x00000000),
  88.     (uint2)(0x8000000a, 0x80000000),
  89.     (uint2)(0x80008081, 0x80000000),
  90.     (uint2)(0x00008080, 0x80000000),
  91.     (uint2)(0x80000001, 0x00000000),
  92.     (uint2)(0x80008008, 0x80000000),
  93. };
  94.  
  95. #ifdef cl_amd_media_ops
  96.     #define ROTL32_1(x, y) amd_bitalign((x), (x).s10, 32 - (y))
  97.     #define ROTL32_2(x, y) amd_bitalign((x).s10, (x), 32 - (y))
  98. #else
  99.     #define ROTL32_1(x, y) rotate(x, y)
  100.     #define ROTL32_2(x, y) ROTL32_1(x, (y) + 32)
  101. #endif
  102.  
  103. #define KECCAKF_1600_RND(a, i, outsz)                                      \
  104.     do                                                                     \
  105.     {                                                                      \
  106.         const uint2 m0 = a[0] ^ a[5] ^ a[10] ^ a[15] ^ a[20] ^             \
  107.                          ROTL32_1(a[2] ^ a[7] ^ a[12] ^ a[17] ^ a[22], 1); \
  108.         const uint2 m1 = a[1] ^ a[6] ^ a[11] ^ a[16] ^ a[21] ^             \
  109.                          ROTL32_1(a[3] ^ a[8] ^ a[13] ^ a[18] ^ a[23], 1); \
  110.         const uint2 m2 = a[2] ^ a[7] ^ a[12] ^ a[17] ^ a[22] ^             \
  111.                          ROTL32_1(a[4] ^ a[9] ^ a[14] ^ a[19] ^ a[24], 1); \
  112.         const uint2 m3 = a[3] ^ a[8] ^ a[13] ^ a[18] ^ a[23] ^             \
  113.                          ROTL32_1(a[0] ^ a[5] ^ a[10] ^ a[15] ^ a[20], 1); \
  114.         const uint2 m4 = a[4] ^ a[9] ^ a[14] ^ a[19] ^ a[24] ^             \
  115.                          ROTL32_1(a[1] ^ a[6] ^ a[11] ^ a[16] ^ a[21], 1); \
  116.                                                                            \
  117.         const uint2 tmp = a[1] ^ m0;                                       \
  118.                                                                            \
  119.         a[0] ^= m4;                                                        \
  120.         a[5] ^= m4;                                                        \
  121.         a[10] ^= m4;                                                       \
  122.         a[15] ^= m4;                                                       \
  123.         a[20] ^= m4;                                                       \
  124.                                                                            \
  125.         a[6] ^= m0;                                                        \
  126.         a[11] ^= m0;                                                       \
  127.         a[16] ^= m0;                                                       \
  128.         a[21] ^= m0;                                                       \
  129.                                                                            \
  130.         a[2] ^= m1;                                                        \
  131.         a[7] ^= m1;                                                        \
  132.         a[12] ^= m1;                                                       \
  133.         a[17] ^= m1;                                                       \
  134.         a[22] ^= m1;                                                       \
  135.                                                                            \
  136.         a[3] ^= m2;                                                        \
  137.         a[8] ^= m2;                                                        \
  138.         a[13] ^= m2;                                                       \
  139.         a[18] ^= m2;                                                       \
  140.         a[23] ^= m2;                                                       \
  141.                                                                            \
  142.         a[4] ^= m3;                                                        \
  143.         a[9] ^= m3;                                                        \
  144.         a[14] ^= m3;                                                       \
  145.         a[19] ^= m3;                                                       \
  146.         a[24] ^= m3;                                                       \
  147.                                                                            \
  148.         a[1] = ROTL32_2(a[6], 12);                                         \
  149.         a[6] = ROTL32_1(a[9], 20);                                         \
  150.         a[9] = ROTL32_2(a[22], 29);                                        \
  151.         a[22] = ROTL32_2(a[14], 7);                                        \
  152.         a[14] = ROTL32_1(a[20], 18);                                       \
  153.         a[20] = ROTL32_2(a[2], 30);                                        \
  154.         a[2] = ROTL32_2(a[12], 11);                                        \
  155.         a[12] = ROTL32_1(a[13], 25);                                       \
  156.         a[13] = ROTL32_1(a[19], 8);                                        \
  157.         a[19] = ROTL32_2(a[23], 24);                                       \
  158.         a[23] = ROTL32_2(a[15], 9);                                        \
  159.         a[15] = ROTL32_1(a[4], 27);                                        \
  160.         a[4] = ROTL32_1(a[24], 14);                                        \
  161.         a[24] = ROTL32_1(a[21], 2);                                        \
  162.         a[21] = ROTL32_2(a[8], 23);                                        \
  163.         a[8] = ROTL32_2(a[16], 13);                                        \
  164.         a[16] = ROTL32_2(a[5], 4);                                         \
  165.         a[5] = ROTL32_1(a[3], 28);                                         \
  166.         a[3] = ROTL32_1(a[18], 21);                                        \
  167.         a[18] = ROTL32_1(a[17], 15);                                       \
  168.         a[17] = ROTL32_1(a[11], 10);                                       \
  169.         a[11] = ROTL32_1(a[7], 6);                                         \
  170.         a[7] = ROTL32_1(a[10], 3);                                         \
  171.         a[10] = ROTL32_1(tmp, 1);                                          \
  172.                                                                            \
  173.         uint2 m5 = a[0];                                                   \
  174.         uint2 m6 = a[1];                                                   \
  175.         a[0] = bitselect(a[0] ^ a[2], a[0], a[1]);                         \
  176.         a[0] ^= as_uint2(Keccak_f1600_RC[i]);                              \
  177.         if (outsz > 1)                                                     \
  178.         {                                                                  \
  179.             a[1] = bitselect(a[1] ^ a[3], a[1], a[2]);                     \
  180.             a[2] = bitselect(a[2] ^ a[4], a[2], a[3]);                     \
  181.             a[3] = bitselect(a[3] ^ m5, a[3], a[4]);                       \
  182.             a[4] = bitselect(a[4] ^ m6, a[4], m5);                         \
  183.             if (outsz > 4)                                                 \
  184.             {                                                              \
  185.                 m5 = a[5];                                                 \
  186.                 m6 = a[6];                                                 \
  187.                 a[5] = bitselect(a[5] ^ a[7], a[5], a[6]);                 \
  188.                 a[6] = bitselect(a[6] ^ a[8], a[6], a[7]);                 \
  189.                 a[7] = bitselect(a[7] ^ a[9], a[7], a[8]);                 \
  190.                 a[8] = bitselect(a[8] ^ m5, a[8], a[9]);                   \
  191.                 a[9] = bitselect(a[9] ^ m6, a[9], m5);                     \
  192.                 if (outsz > 8)                                             \
  193.                 {                                                          \
  194.                     m5 = a[10];                                            \
  195.                     m6 = a[11];                                            \
  196.                     a[10] = bitselect(a[10] ^ a[12], a[10], a[11]);        \
  197.                     a[11] = bitselect(a[11] ^ a[13], a[11], a[12]);        \
  198.                     a[12] = bitselect(a[12] ^ a[14], a[12], a[13]);        \
  199.                     a[13] = bitselect(a[13] ^ m5, a[13], a[14]);           \
  200.                     a[14] = bitselect(a[14] ^ m6, a[14], m5);              \
  201.                     m5 = a[15];                                            \
  202.                     m6 = a[16];                                            \
  203.                     a[15] = bitselect(a[15] ^ a[17], a[15], a[16]);        \
  204.                     a[16] = bitselect(a[16] ^ a[18], a[16], a[17]);        \
  205.                     a[17] = bitselect(a[17] ^ a[19], a[17], a[18]);        \
  206.                     a[18] = bitselect(a[18] ^ m5, a[18], a[19]);           \
  207.                     a[19] = bitselect(a[19] ^ m6, a[19], m5);              \
  208.                     m5 = a[20];                                            \
  209.                     m6 = a[21];                                            \
  210.                     a[20] = bitselect(a[20] ^ a[22], a[20], a[21]);        \
  211.                     a[21] = bitselect(a[21] ^ a[23], a[21], a[22]);        \
  212.                     a[22] = bitselect(a[22] ^ a[24], a[22], a[23]);        \
  213.                     a[23] = bitselect(a[23] ^ m5, a[23], a[24]);           \
  214.                     a[24] = bitselect(a[24] ^ m6, a[24], m5);              \
  215.                 }                                                          \
  216.             }                                                              \
  217.         }                                                                  \
  218.     } while (0)
  219.  
  220.  
  221. #define KECCAK_PROCESS(st, in_size, out_size)    \
  222.     do                                           \
  223.     {                                            \
  224.         uchar r=0;                               \
  225.         uchar os=0;                                      \
  226.         do                                       \
  227.         {                                        \
  228.             os=(r < 23 ? 25 : (out_size));  \
  229.             KECCAKF_1600_RND(st, r, os);         \
  230.             r++;                                 \
  231.         }while(r<24);                            \
  232.     } while (0)
  233.  
  234. #define fnv(x, y) ((x)*FNV_PRIME ^ (y))
  235. #define fnv_reduce(v) fnv(fnv(fnv(v.x, v.y), v.z), v.w)
  236.  
  237. #define MIX(x)                                                          \
  238.     do                                                                  \
  239.     {                                                                   \
  240.         *(local_buffer) = fnv(init0 ^ (a + x), *(imix+x)) % dag_size;   \
  241.         mix = fnv(mix, g_dag_uint[(buffer[lane_idx]*4)+ids[1]]);        \
  242.         mem_fence(CLK_LOCAL_MEM_FENCE);                                 \
  243.     } while (0)
  244.  
  245. // NOTE: This struct must match the one defined in CLMiner.cpp
  246. struct __attribute__((packed)) __attribute__((aligned(128))) SearchResults
  247. {
  248.     uint count;
  249.     uint hashCount;
  250.     volatile uint abort;
  251.     uint gid[MAX_OUTPUTS];
  252. };
  253.  
  254. __attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) __kernel void search(
  255.     __global struct SearchResults* g_output, __constant uint2 const* g_header,
  256.     __global ulong8 const* _g_dag0, __global ulong8 const* _g_dag1, uint dag_size,
  257.     ulong start_nonce, ulong target)
  258. {
  259.     if (g_output->abort)
  260.         return;
  261.  
  262.     const uchar ids[] = {           \
  263.     (uchar) get_local_id(0),        \
  264.     (uchar)(get_local_id(0) & 3),       \
  265.     (uchar)(get_local_id(0) >> 2),      \
  266.     (uchar)(get_local_id(0) >> 2 << 2)};
  267.     const uint gid = get_global_id(0);
  268.     __global uint8 const* g_dag_uint = (__global uint8 const*)_g_dag0;
  269.  
  270.     __local uint sharebuf[1024];
  271.     __local uint buffer[256];
  272.     __local uint *local_buffer=&buffer[ids[0]];
  273.     __local ulong8 *ulong8_buffer=&sharebuf[ids[2]*16]; // (write buffer 64 bytes )
  274.     __local ulong4 *ulong4_buffer=&sharebuf[ids[2]*16]; // (read buffer 32 bytes )
  275.     __local uint8 *uint8_buffer=(uint)&sharebuf[ids[2]*16]; // (read buffer 32 bytes )
  276.     __local uint2 *uint2_buffer=(uint)&sharebuf[ids[2]*16]; // (write buffer 16 bytes )
  277.     __local ushort16 *ushort16_buffer=&sharebuffer[ids[2]*16];
  278.     __local uint *uint_buffer=(uint)&sharebuf[ids[2]*16]; // ( read buffer )
  279.  
  280.     uint2 state[25]; // would this be faster on my gpu as a ushort4
  281.     ulong8 *convert=&state;
  282.     ulong4 *convert2=&state;
  283.     uchar a,x,lane;
  284.     char tid=0;
  285.     ushort init0;
  286.     ushort16 mix;
  287.     ushort *imix=&mix;
  288.  
  289.     *(convert)=(ulong8)(0);
  290.     *(convert+1)=*(convert);
  291.     *(convert+2)=*(convert);
  292.     state[0] = g_header[0];
  293.     state[1] = g_header[1];
  294.     state[2] = g_header[2];
  295.     state[3] = g_header[3];
  296.     state[4] = as_uint2(start_nonce + gid);
  297.     state[5] = as_uint2(0x0000000000000001UL);
  298.     state[8] = as_uint2(0x8000000000000000UL);
  299.     state[24] = state[23];
  300.  
  301.     KECCAK_PROCESS(state,5, 8);
  302.  
  303.     for(tid=0; tid < 4; tid++)
  304.     {
  305.         barrier(CLK_LOCAL_MEM_FENCE);
  306.         if ( ids[1] == tid-1 )
  307.             *(convert2+2) = *(ulong4_buffer);
  308.  
  309.         if ( tid == ids[1])
  310.             *(ulong8_buffer)=*(convert);
  311.         barrier(CLK_LOCAL_MEM_FENCE);
  312.  
  313.         mix = *(ushort16_buffer+(ids[1]&1));
  314.         init0 = *(uint_buffer);
  315.  
  316.         a=0;lane=0;
  317.         barrier(CLK_LOCAL_MEM_FENCE);
  318.         #pragma unroll 1
  319.         do
  320.         {
  321.             const uchar lane_idx = ids[3] + lane;
  322.             #pragma unroll 8
  323.             for (x = 0; x < 8; ++x)
  324.                 MIX(x);
  325.  
  326.             lane=(lane+1)&3;
  327.             a += 8;
  328.         }while(a < ACCESSES);
  329.  
  330.         barrier(CLK_LOCAL_MEM_FENCE);
  331.         *(ushort16_buffer+ids[1]) = (ushort)(fnv_reduce(mix.lo), fnv_reduce(mix.hi));
  332.     }
  333.  
  334.     if ( ids[1] == 3 )
  335.         *(convert2+2) = *(ulong4_buffer);
  336.  
  337.     *(convert2+3)=(ulong4)(0);
  338.     *(convert2+4)=*(convert2+3);
  339.     *(convert2+5)=*(convert2+3);
  340.     state[12] = as_uint2(0x0000000000000001UL);
  341.     state[16] = as_uint2(0x8000000000000000UL);
  342.     state[24] = state[23];
  343.  
  344.     KECCAK_PROCESS(state,12, 1);
  345.  
  346.     if (get_local_id(0) == 0)
  347.     {
  348.         atomic_inc(&g_output->hashCount);
  349.     }
  350.  
  351.     // weird
  352.     if (as_ulong(as_uchar8(state[0]).s76543210) <= target)
  353.     {
  354.         atomic_inc(&g_output->abort);
  355.         uint slot = min(MAX_OUTPUTS - 1u, atomic_inc(&g_output->count));
  356.         g_output->gid[slot] = gid;
  357.     }
  358. }
  359.  
  360. typedef union _Node
  361. {
  362.     uint dwords[16];
  363.     uint2 qwords[8];
  364.     uint4 dqwords[4];
  365. } Node;
  366.  
  367. static void SHA3_512(uint2* s)
  368. {
  369.     uint2 state[25];
  370.     ulong8 *convert=&state;
  371.     ulong8 *result=s;
  372.  
  373.     *(convert)=*(result);
  374.     *(convert+1)=(ulong8)(0);
  375.     *(convert+2)=(ulong8)(0);
  376.     state[8] = (uint2)(0x00000001, 0x80000000);
  377.     state[24] = (uint2)(0);
  378.  
  379.     KECCAK_PROCESS(state,8, 8);
  380.  
  381.     *(result)=*(convert);
  382. }
  383.  
  384. __kernel void GenerateDAG(uint start, __global const uint16* _Cache, __global uint16* _DAG0,
  385.     __global uint16* _DAG1, uint light_size)
  386. {
  387.     __global const Node* Cache = (__global const Node*)_Cache;
  388.     const uint gid = get_global_id(0);
  389.     uint NodeIdx = start + gid;
  390.     const uint thread_id = gid & 3;
  391.  
  392.     __local Node sharebuf[WORKSIZE];
  393.     __local uint indexbuf[WORKSIZE];
  394.     __local Node* dagNode = sharebuf + (get_local_id(0) / 4) * 4;
  395.     __local uint* indexes = indexbuf + (get_local_id(0) / 4) * 4;
  396.     __global const Node* parentNode;
  397.  
  398.     Node DAGNode = Cache[NodeIdx % light_size];
  399.  
  400.     DAGNode.dwords[0] ^= NodeIdx;
  401.     SHA3_512(DAGNode.qwords);
  402.  
  403.     dagNode[thread_id] = DAGNode;
  404.     barrier(CLK_LOCAL_MEM_FENCE);
  405.  
  406.     for (uint i = 0; i < 256; ++i)
  407.     {
  408.         uint ParentIdx = fnv(NodeIdx ^ i, dagNode[thread_id].dwords[i & 15]) % light_size;
  409.         indexes[thread_id] = ParentIdx;
  410.         barrier(CLK_LOCAL_MEM_FENCE);
  411.  
  412.         for (uint t = 0; t < 4; ++t)
  413.         {
  414.             uint parentIndex = indexes[t];
  415.             parentNode = Cache + parentIndex;
  416.  
  417.             dagNode[t].dqwords[thread_id] = fnv(dagNode[t].dqwords[thread_id], parentNode->dqwords[thread_id]);
  418.             barrier(CLK_LOCAL_MEM_FENCE);
  419.         }
  420.     }
  421.     DAGNode = dagNode[thread_id];
  422.  
  423.     SHA3_512(DAGNode.qwords);
  424.  
  425.     __global Node* DAG;
  426.     DAG = (__global Node *) _DAG0;
  427.     DAG[NodeIdx] = DAGNode;
  428. }
Add Comment
Please, Sign In to add comment
Public Pastes
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy.  OK, I Understand
Not a member of Pastebin yet?
Sign Up, it unlocks many cool features!