// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s

// Checks the execution scope, memory scope and memory semantics operands that
// group_barrier passes to __spirv_ControlBarrier.

#include <sycl/sycl.hpp>

void test(sycl::queue Q) {
  Q.parallel_for<class barrier_kernel>(
      sycl::nd_range<1>{8, 4}, [=](sycl::nd_item<1> It) {
        // Execution scope is Workgroup (2), memory scope defaults to the
        // group's fence_scope, i.e. Workgroup (2).
        // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
        // 2, i32 noundef 2, i32 noundef 400)
        sycl::group_barrier(It.get_group());
        // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
        // 2, i32 noundef 4, i32 noundef 400)
        sycl::group_barrier(It.get_group(), sycl::memory_scope::work_item);
        // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
        // 2, i32 noundef 3, i32 noundef 400)
        sycl::group_barrier(It.get_group(), sycl::memory_scope::sub_group);
        // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
        // 2, i32 noundef 2, i32 noundef 400)
        sycl::group_barrier(It.get_group(), sycl::memory_scope::work_group);
        // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
        // 2, i32 noundef 1, i32 noundef 400)
        sycl::group_barrier(It.get_group(), sycl::memory_scope::device);
        // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
        // 2, i32 noundef 0, i32 noundef 400)
        sycl::group_barrier(It.get_group(), sycl::memory_scope::system);

        // Execution scope is Subgroup (3), memory scope defaults to the
        // sub-group's fence_scope, i.e. Subgroup (3).
        // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
        // 3, i32 noundef 3, i32 noundef 400)
        sycl::group_barrier(It.get_sub_group());
      });
}