diff --git a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
index 16a10b2304312..e268f20982c98 100644
--- a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
+++ b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
@@ -29,6 +29,14 @@ namespace ext::oneapi {
 #define BITS_TYPE uint32_t
 #endif
 
+// defining `group_ballot` here to make predicate default `true`
+// need to forward declare sub_group_mask first
+struct sub_group_mask;
+template <typename Group>
+detail::enable_if_t<std::is_same<std::decay_t<Group>, sub_group>::value,
+                    sub_group_mask>
+group_ballot(Group g, bool predicate = true);
+
 struct sub_group_mask {
   friend class detail::Builder;
   using BitsType = BITS_TYPE;
diff --git a/sycl/test/check_device_code/group_ballot.cpp b/sycl/test/check_device_code/group_ballot.cpp
new file mode 100644
index 0000000000000..58f700ea15d77
--- /dev/null
+++ b/sycl/test/check_device_code/group_ballot.cpp
@@ -0,0 +1,11 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
+// expected-no-diagnostics
+
+#include <sycl/sycl.hpp>
+
+int main() {
+  sycl::queue Q;
+  Q.parallel_for(sycl::nd_range<1>{32, 32}, [=](sycl::nd_item<1> item) {
+    auto Mask = sycl::ext::oneapi::group_ballot(item.get_sub_group());
+  });
+}