Skip to content
Snippets Groups Projects
Commit 336f2325 authored by fireice-uk's avatar fireice-uk Committed by GitHub
Browse files

Merge pull request #57 from psychocrypt/fix-nvidiaBackendCrash

fix illegal memory access
parents 610f4f0f 29363d5a
No related branches found
No related tags found
No related merge requests found
......@@ -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;
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment