def consumer_release(self, state: PipelineState):
        """
        TMA consumer release conditionally signals the empty buffer to the producer.
        """
        # Only 1 thread per warp group signals the empty buffer.
        if_generate(
            cute.arch.thread_idx()[0] % 128 == 0,
            lambda: self.sync_object_array_empty.arrive(state.index, self.consumer_mask),
        )

        TMA consumer release conditionally signals the empty buffer to the producer.
        
flash_attn/cute/pipeline.py