Skip to content

Commit 96e9248

Browse files
committed
RLE gpu decompress kernel
Signed-off-by: Robert Kruszewski <[email protected]>
1 parent e71bcbc commit 96e9248

File tree

7 files changed

+431
-50
lines changed

7 files changed

+431
-50
lines changed

java/testfiles/Cargo.lock

Lines changed: 63 additions & 26 deletions
Some generated files are not rendered by default. Learn more about customizing how changed files appear on GitHub.

vortex-gpu/kernels/dict_take.cu

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -12,16 +12,16 @@ __device__ void dict_take(
1212
ValueT *__restrict values_out
1313
) {
1414
auto i = threadIdx.x;
15-
auto block_offset = (blockIdx.x * 1024);
15+
auto block_offset = blockIdx.x * 1024;
1616

1717
auto codes = codes_array + block_offset;
1818
auto out = values_out + block_offset;
1919

2020
const int thread_ops = 32;
2121

2222
for (auto j = 0; j < thread_ops; j++) {
23-
auto idx = i * thread_ops + j;
24-
out[idx] = values[codes[idx]];
23+
auto idx = i * thread_ops + j;
24+
out[idx] = values[codes[idx]];
2525
}
2626
}
2727

@@ -33,8 +33,8 @@ __device__ void dict_take_masked(
3333
ValueT *__restrict values_out
3434
) {
3535
auto i = threadIdx.x;
36-
auto block_offset = (blockIdx.x * 1024);
37-
auto mask_block_offset = (blockIdx.x * (1024 / 32));
36+
auto block_offset = blockIdx.x * 1024;
37+
auto mask_block_offset = blockIdx.x * (1024 / 32);
3838

3939
auto codes = codes_array + block_offset;
4040
auto mask = mask_array + mask_block_offset;

0 commit comments

Comments
 (0)