diff options
author | psychocrypt <psychocrypt@users.noreply.github.com> | 2017-10-20 20:57:52 +0200 |
---|---|---|
committer | psychocrypt <psychocrypt@users.noreply.github.com> | 2017-10-21 20:44:23 +0200 |
commit | 29363d5a1aede05995790d0d648f091699e9ba25 (patch) | |
tree | c271a25bc7f9217e22505c5ed3f5b9fe61d360ea /xmrstak | |
parent | 6701b0c2025987674708c9984c1d76461428c09e (diff) | |
download | xmr-stak-29363d5a1aede05995790d0d648f091699e9ba25.zip xmr-stak-29363d5a1aede05995790d0d648f091699e9ba25.tar.gz |
fix illegal memory access
remove restricted pointer
Diffstat (limited to 'xmrstak')
-rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cuda_blake.hpp | 17 |
1 files changed, 9 insertions, 8 deletions
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_blake.hpp b/xmrstak/backend/nvidia/nvcc_code/cuda_blake.hpp index 07ae169..340174c 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_blake.hpp +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_blake.hpp @@ -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; |