Skip to content

Commit

Permalink
formatting
Browse files Browse the repository at this point in the history
  • Loading branch information
ipdemes committed Nov 14, 2023
1 parent d9b044b commit 579cfe1
Showing 1 changed file with 5 additions and 2 deletions.
7 changes: 5 additions & 2 deletions src/cunumeric/index/select.cu
Original file line number Diff line number Diff line change
Expand Up @@ -86,24 +86,27 @@ struct SelectImplBody<VariantKind::GPU, CODE, DIM> {
const size_t blocks = (out_size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;

auto stream = get_cached_stream();

if (dense && (DIM <= 1 || rect.volume() == 0)) {
auto cond_arr = create_buffer<const bool*>(condlist.size(), legate::Memory::Kind::Z_COPY_MEM);
for (uint32_t idx = 0; idx < condlist.size(); ++idx) cond_arr[idx] = condlist[idx].ptr(rect);
auto choice_arr =
create_buffer<const VAL*>(choicelist.size(), legate::Memory::Kind::Z_COPY_MEM);
for (uint32_t idx = 0; idx < choicelist.size(); ++idx)
choice_arr[idx] = choicelist[idx].ptr(rect);

VAL* outptr = out.ptr(rect);
select_kernel_dense<VAL><<<blocks, THREADS_PER_BLOCK, 0, stream>>>(
outptr, narrays, cond_arr, choice_arr, default_val, out_size);
} else {

} else { // not dense
auto cond_arr =
create_buffer<AccessorRO<bool, DIM>>(condlist.size(), legate::Memory::Kind::Z_COPY_MEM);
for (uint32_t idx = 0; idx < condlist.size(); ++idx) cond_arr[idx] = condlist[idx];

auto choice_arr =
create_buffer<AccessorRO<VAL, DIM>>(choicelist.size(), legate::Memory::Kind::Z_COPY_MEM);
for (uint32_t idx = 0; idx < choicelist.size(); ++idx) choice_arr[idx] = choicelist[idx];

if (out_size == 0) return;
select_kernel<VAL, DIM><<<blocks, THREADS_PER_BLOCK, 0, stream>>>(
out, narrays, cond_arr, choice_arr, default_val, rect, pitches, out_size, rect.volume());
Expand Down

0 comments on commit 579cfe1

Please sign in to comment.