github-jwinterm/xmr-stak
1
0
Fork 0
mirror of https://github.com/jwinterm/xmr-stak.git synced 2026-10-17 05:41:41 +02:00
Code Issues Packages Projects Releases Wiki Activity

fix illegal memory access

Browse Source 
remove restricted pointer
This commit is contained in:
psychocrypt
2017-10-21 20:44:23 +02:00
parent 6701b0c202
commit 29363d5a1a
1 changed files with 9 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files
xmrstak/backend/nvidia/nvcc_code/cuda_blake.hpp
+9 -8
View File
@@ -2,7 +2,8 @@
typedef struct {
uint32_t h[8], s[4], t[2];
int buflen, nullt;
uint32_t buflen;
int nullt;
uint8_t buf[64];
} blake_state;
@@ -50,7 +51,7 @@ __constant__ uint32_t d_blake_cst[16]
0xC0AC29B7, 0xC97C50DD, 0x3F84D5B5, 0xB5470917
};
__device__ void cn_blake_compress(blake_state * __restrict__ S, const uint8_t * __restrict__ block)
__device__ void cn_blake_compress(blake_state * S, const uint8_t * block)
{
uint32_t v[16], m[16], i;
@@ -89,12 +90,12 @@ __device__ void cn_blake_compress(blake_state * __restrict__ S, const uint8_t *
for (i = 0; i < 8; ++i) S->h[i] ^= S->s[i % 4];
}
__device__ void cn_blake_update(blake_state * __restrict__ S, const uint8_t * __restrict__ data, uint64_t datalen)
__device__ void cn_blake_update(blake_state * S, const uint8_t * data, uint64_t datalen)
{
int left = S->buflen >> 3;
int fill = 64 - left;
uint32_t left = S->buflen >> 3;
uint32_t fill = 64 - left;
if (left && (((datalen >> 3) & 0x3F) >= (unsigned) fill))
if (left && (((datalen >> 3) & 0x3F) >= fill))
{
memcpy((void *) (S->buf + left), (void *) data, fill);
S->t[0] += 512;
@@ -125,7 +126,7 @@ __device__ void cn_blake_update(blake_state * __restrict__ S, const uint8_t * __
}
}
__device__ void cn_blake_final(blake_state * __restrict__ S, uint8_t * __restrict__ digest)
__device__ void cn_blake_final(blake_state * S, uint8_t * digest)
{
const uint8_t padding[] =
{
@@ -177,7 +178,7 @@ __device__ void cn_blake_final(blake_state * __restrict__ S, uint8_t * __restric
U32TO8(digest + 28, S->h[7]);
}
__device__ void cn_blake(const uint8_t * __restrict__ in, uint64_t inlen, uint8_t * __restrict__ out)
__device__ void cn_blake(const uint8_t * in, uint64_t inlen, uint8_t * out)
{
blake_state bs;
blake_state *S = (blake_state *)&bs;