OpGroupWaitEvents doesn't require the Group capability

Fixes #595

Author: karolherbst

llvm-spirv/lib/SPIRV/libSPIRV/SPIRVInstruction.h

@@ -2479,8 +2479,11 @@ class SPIRVGroupInstBase : public SPIRVInstTemplateBase {
 
 #define _SPIRV_OP(x, ...)                                                      \
   typedef SPIRVInstTemplate<SPIRVGroupInstBase, Op##x, __VA_ARGS__> SPIRV##x;
-// Group instructions
-_SPIRV_OP(GroupWaitEvents, false, 4)
+// Group instructions.
+// Even though GroupWaitEvents has Group in its name, it doesn't require the
+// Group capability
+typedef SPIRVInstTemplate<SPIRVInstTemplateBase, OpGroupWaitEvents, false, 4>
+    SPIRVGroupWaitEvents;
 _SPIRV_OP(GroupAll, true, 5)
 _SPIRV_OP(GroupAny, true, 5)
 _SPIRV_OP(GroupBroadcast, true, 6)

llvm-spirv/test/event_no_group_cap.cl

@@ -0,0 +1,12 @@
+__kernel void test_fn( const __global char *src)
+{
+	wait_group_events(0, NULL);
+}
+// RUN: %clang_cc1 -triple spir64 -x cl -cl-std=CL2.0 -finclude-default-header -O0 -emit-llvm-bc %s -o %t.bc
+// RUN: llvm-spirv %t.bc -spirv-text -o %t.spt
+// RUN: FileCheck < %t.spt %s
+// RUN: llvm-spirv %t.bc -o %t.spv
+// RUN: spirv-val %t.spv
+
+// CHECK-NOT:Capability Groups
+// CHECK:GroupWaitEvents