Merge pull request #2147 from ekuznetsov139/master

jean-m-cyr · web-flow · commit ce52c74021b6 · 2021-02-09T19:42:43.000-05:00
Avoiding buffer overflow in SHA3
diff --git a/libethash-cuda/ethash_cuda_miner_kernel.cu b/libethash-cuda/ethash_cuda_miner_kernel.cu
@@ -52,11 +52,13 @@ __global__ void ethash_calculate_dag_item(uint32_t start)
     uint32_t const node_index = start + blockIdx.x * blockDim.x + threadIdx.x;
     if (((node_index >> 1) & (~1)) >= d_dag_size)
         return;
-
-    hash128_t dag_node;
+    union {
+       hash128_t dag_node;
+       uint2 dag_node_mem[25];
+    };
     copy(dag_node.uint4s, d_light[node_index % d_light_size].uint4s, 4);
     dag_node.words[0] ^= node_index;
-    SHA3_512(dag_node.uint2s);
+    SHA3_512(dag_node_mem);
 
     const int thread_id = threadIdx.x & 3;
 
@@ -78,7 +80,7 @@ __global__ void ethash_calculate_dag_item(uint32_t start)
             }
         }
     }
-    SHA3_512(dag_node.uint2s);
+    SHA3_512(dag_node_mem);
     hash64_t* dag_nodes = (hash64_t*)d_dag;
     copy(dag_nodes[node_index].uint4s, dag_node.uint4s, 4);
 }

Original file line number	Diff line number	Diff line change
`@@ -52,11 +52,13 @@ __global__ void ethash_calculate_dag_item(uint32_t start)`
`52`	`52`	`uint32_t const node_index = start + blockIdx.x * blockDim.x + threadIdx.x;`
`53`	`53`	`if (((node_index >> 1) & (~1)) >= d_dag_size)`
`54`	`54`	`return;`
`55`		`-`
`56`		`- hash128_t dag_node;`
	`55`	`+ union {`
	`56`	`+ hash128_t dag_node;`
	`57`	`+ uint2 dag_node_mem[25];`
	`58`	`+ };`
`57`	`59`	`copy(dag_node.uint4s, d_light[node_index % d_light_size].uint4s, 4);`
`58`	`60`	`dag_node.words[0] ^= node_index;`
`59`		`- SHA3_512(dag_node.uint2s);`
	`61`	`+ SHA3_512(dag_node_mem);`
`60`	`62`
`61`	`63`	`const int thread_id = threadIdx.x & 3;`
`62`	`64`
`@@ -78,7 +80,7 @@ __global__ void ethash_calculate_dag_item(uint32_t start)`
`78`	`80`	`}`
`79`	`81`	`}`
`80`	`82`	`}`
`81`		`- SHA3_512(dag_node.uint2s);`
	`83`	`+ SHA3_512(dag_node_mem);`
`82`	`84`	`hash64_t* dag_nodes = (hash64_t*)d_dag;`
`83`	`85`	`copy(dag_nodes[node_index].uint4s, dag_node.uint4s, 4);`
`84`	`86`	`}`