Skip to content

Commit 9663929

Browse files
chore(cuda): keccakf memcpy handling of padding (#2167)
Currently the invalid memory access is never used, but it's still better to address it so `compute-analyzer` does not complain.
1 parent 426091b commit 9663929

File tree

1 file changed

+10
-3
lines changed

1 file changed

+10
-3
lines changed

extensions/keccak256/circuit/cuda/src/keccakf.cu

Lines changed: 10 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -180,8 +180,15 @@ __global__ void keccakf_kernel(
180180
uint8_t *chunk = input + blk * KECCAK_RATE_BYTES;
181181
bool is_last_block = (blk == num_blocks - 1);
182182
for (int round = 0; round < NUM_ABSORB_ROUNDS; round++) {
183-
uint8_t i_bytes[8];
184-
memcpy(&i_bytes, chunk + round * 8, sizeof(uint64_t));
183+
uint8_t i_bytes[8] = {};
184+
int round_base = round * sizeof(uint64_t);
185+
// For the last block, zero-initialize and only copy available bytes
186+
if (is_last_block && round_base < last_block_len) {
187+
size_t bytes_to_copy = min(sizeof(uint64_t), last_block_len - round_base);
188+
memcpy(i_bytes, chunk + round_base, bytes_to_copy);
189+
} else if (!is_last_block) {
190+
memcpy(i_bytes, chunk + round_base, sizeof(uint64_t));
191+
}
185192

186193
// Handle Keccak spec padding
187194
if (is_last_block) {
@@ -244,4 +251,4 @@ extern "C" int _keccakf_kernel(
244251
bitwise_num_bits
245252
);
246253
return CHECK_KERNEL();
247-
}
254+
}

0 commit comments

Comments
 (0)