Skip to content

GitLab

Commit 20eda944 authored by limm's avatar limm
Browse files

fixed a bug regarding PTX commands

parent e70ac2fd
Hide whitespace changes
Inline Side-by-side
Showing with 2 additions and 1 deletion
apex/contrib/csrc/group_norm/group_norm_nhwc_bwd_one_pass_kernel.cuh
View file @ 20eda944
...@@ -328,7 +328,8 @@ __global__ __launch_bounds__(THREADS_PER_BLOCK_) ...@@ -328,7 +328,8 @@ __global__ __launch_bounds__(THREADS_PER_BLOCK_)
// Mark the completion of the threadblock. // Mark the completion of the threadblock.
if( threadIdx.x == 0 ) { if( threadIdx.x == 0 ) {
asm volatile("red.release.gpu.global.add.s32 [%0], 1;" :: "l"(barrier)); int one = 1;
asm volatile("red.release.gpu.global.add.s32 [%0], 1;" :: "l"(barrier), "r"(one));
} }
// Exit if that's not the last thread block. // Exit if that's not the last thread block.
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment