Skip to content
This repository has been archived by the owner on Mar 21, 2024. It is now read-only.

Commit

Permalink
Fix invalid PTX issue
Browse files Browse the repository at this point in the history
  • Loading branch information
miscco committed Apr 13, 2023
1 parent c31d239 commit 895ffbe
Showing 1 changed file with 4 additions and 6 deletions.
10 changes: 4 additions & 6 deletions .upstream-tests/test/cuda/barrier_native_handle.pass.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -20,19 +20,17 @@
__device__
void test()
{
shared_memory_selector<cuda::barrier<cuda::thread_scope_block>, constructor_initializer> sel;
SHARED cuda::barrier<cuda::thread_scope_block>* b;
b = sel.construct();
init(b, 2);
__shared__ cuda::barrier<cuda::thread_scope_block> b;
init(&b, 2);

uint64_t token;
asm volatile ("mbarrier.arrive.b64 %0, [%1];"
: "=l"(token)
: "l"(cuda::device::barrier_native_handle(*b))
: "l"(cuda::device::barrier_native_handle(b))
: "memory");
(void)token;

b->arrive_and_wait();
b.arrive_and_wait();
}

int main(int argc, char ** argv)
Expand Down

0 comments on commit 895ffbe

Please sign in to comment.