Jokeren · July 30, 2025 21:57
diff --git a/README.md b/README.md
diff --git a/test.cu b/test.cu
 #include <cuda/ptx>
 #include <cuda/std/cstdint>
 #include <cstdio>

 using u64 = cuda::std::uint64_t;
 using u32 = cuda::std::uint32_t;

 __device__ void wait(u64 *barrier_generic) {
  while (!cuda::ptx::mbarrier_test_wait(barrier_generic,
                                        static_cast<u64>(0))) {
    // busy‐loop
  }
 }

 __global__ void repro() {
  extern __shared__ u64 smem[];

  if (threadIdx.x == 0) {
      u32 count = 1;
      cuda::ptx::mbarrier_init(smem, count);
  }
  __syncthreads();

  wait(smem);

  if (threadIdx.x == 0) {
    printf("Barrier released\n");
  }
 }

 int main() {
  // one block of 32 threads, plus 8 bytes of shared storage
  repro<<<1, 32, sizeof(cuda::std::uint64_t)>>>();
  cudaDeviceSynchronize();
  return 0;
 }
 ~                
diff --git a/test.py b/test.py
 from triton.experimental import gluon
 from triton.experimental.gluon import language as gl
 from triton.experimental.gluon.language.nvidia.hopper import mbarrier


 @gluon.jit
 def foo(b):
    mbarrier.wait(b, 0)


 @gluon.jit
 def bar(b):
    foo(b)


 @gluon.jit
 def repro():
    b = gl.allocate_shared_memory(gl.int64, [1], mbarrier.MBarrierLayout())
    mbarrier.init(b, count=1)
    bar(b)
    mbarrier.invalidate(b)


 repro[(1, )]()
	#include <cuda/ptx>
	#include <cuda/std/cstdint>
	#include <cstdio>

	using u64 = cuda::std::uint64_t;
	using u32 = cuda::std::uint32_t;

	__device__ void wait(u64 *barrier_generic) {
	while (!cuda::ptx::mbarrier_test_wait(barrier_generic,
	static_cast<u64>(0))) {
	// busy‐loop
	}
	}

	__global__ void repro() {
	extern __shared__ u64 smem[];

	if (threadIdx.x == 0) {
	u32 count = 1;
	cuda::ptx::mbarrier_init(smem, count);
	}
	__syncthreads();

	wait(smem);

	if (threadIdx.x == 0) {
	printf("Barrier released\n");
	}
	}

	int main() {
	// one block of 32 threads, plus 8 bytes of shared storage
	repro<<<1, 32, sizeof(cuda::std::uint64_t)>>>();
	cudaDeviceSynchronize();
	return 0;
	}
	~
	from triton.experimental import gluon
	from triton.experimental.gluon import language as gl
	from triton.experimental.gluon.language.nvidia.hopper import mbarrier


	@gluon.jit
	def foo(b):
	mbarrier.wait(b, 0)


	@gluon.jit
	def bar(b):
	foo(b)


	@gluon.jit
	def repro():
	b = gl.allocate_shared_memory(gl.int64, [1], mbarrier.MBarrierLayout())
	mbarrier.init(b, count=1)
	bar(b)
	mbarrier.invalidate(b)


	repro[(1, )]()