Support for SYCL's basic and nd-range kernels

examples/concepts/cpp/Namespaces.cpp

@@ -6,7 +6,6 @@
 int x;
 
 namespace spaceA {
-	int x;
 
 	//@ context Perm(x, write);
 	//@ ensures x == 90;
@@ -28,21 +27,15 @@ namespace spaceA {
 	}
 }
 
-//@ context Perm(spaceA::x, write);
 //@ context Perm(x, write);
 int main() {
 	x = 99;
-	spaceA::x = 5;
-	//@ assert spaceA::x == 5;
   int varA = spaceA::incr();
-  //@ assert varA == 6;
-  //@ assert spaceA::x == 90;
+  //@ assert varA == 100;
+  //@ assert x == 90;
   int varB = spaceA::spaceB::incr();
   //@ assert varB == 92;
   spaceA::spaceB::doNothing();
-  int varX = spaceA::x;
-  //@ assert varX == 90;
-
-  //@ assert x == 99;
+  //@ assert x == 90;
   return 0;
 }

examples/concepts/cpp/NamespacesDoNotFindWithoutNamespacePrefix.cpp

@@ -6,10 +6,10 @@
 namespace spaceA {
 
 	namespace spaceB {
-	  int b = 10;
+	  int b();
 	}
 
 	int getB() {
-	   return b;
+	   return b();
 	}
 }

examples/concepts/sycl/kernels/BasicKernel.cpp

@@ -0,0 +1,26 @@
+#include <sycl/sycl.hpp>
+
+void main() {
+	sycl::queue myQueue;
+
+	sycl::event myEvent = myQueue.submit(
+		[&](sycl::handler& cgh) {
+			cgh.parallel_for(sycl::range<3>(6,4,2),
+				/*@
+					requires it.get_id(0) > -1;
+					ensures it.get_linear_id() > -1;
+				*/
+				[=] (sycl::item<3> it) {
+					int a = it.get_id(1);
+					//@ assert a < 6 && a >= 0;
+					int b = it.get_linear_id();
+					//@ assert b < 48 && b >= 0;
+					int c = it.get_range(0);
+					//@ assert c == 6;
+				}
+			);
+		}
+	);
+	int a = 5;
+	myEvent.wait();
+}

examples/concepts/sycl/kernels/MultipleKernelsInCommandGroup.cpp

@@ -0,0 +1,12 @@
+#include <sycl/sycl.hpp>
+
+void main() {
+	sycl::queue myQueue;
+
+	sycl::event myEvent = myQueue.submit(
+	[&](sycl::handler& cgh) {
+		// Only one kernel is allowed in a command group
+		cgh.parallel_for(sycl::range<3>(6,4,2), [=] (sycl::item<3> it) {});
+		cgh.parallel_for(sycl::range<1>(6), [=] (sycl::item<1> it) {});
+	});
+}

examples/concepts/sycl/kernels/NDRangeKernel.cpp

@@ -0,0 +1,28 @@
+#include <sycl/sycl.hpp>
+
+void main() {
+	sycl::queue myQueue;
+
+	sycl::event myEvent = myQueue.submit(
+  	[&](sycl::handler& cgh) {
+  		cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(6,4,4), sycl::range<3>(3,2,1)),
+				[=] (sycl::nd_item<3> it) {
+					int a = it.get_global_id(0);
+					int b = it.get_global_linear_id();
+					int c = it.get_global_range(0);
+
+					int d = it.get_local_id(0);
+					int e = it.get_local_linear_id();
+					int f = it.get_local_range(0);
+
+					int g = it.get_group_id(0);
+					int h = it.get_group_linear_id();
+					int i = it.get_group_range(0);
+				}
+  		);
+  	}
+  );
+
+	int a = 5;
+	myEvent.wait();
+}

examples/concepts/sycl/kernels/NegativeRange.cpp

@@ -0,0 +1,10 @@
+#include <sycl/sycl.hpp>
+
+void main() {
+	sycl::queue myQueue;
+
+	sycl::event myEvent = myQueue.submit(
+	[&](sycl::handler& cgh) {
+		cgh.parallel_for(sycl::range<2>(6, -2), [=] (sycl::item<2> it) {});
+	});
+}

examples/concepts/sycl/kernels/NoKernelInCommandGroup.cpp

@@ -0,0 +1,9 @@
+#include <sycl/sycl.hpp>
+
+void main() {
+	sycl::queue myQueue;
+
+	sycl::event myEvent = myQueue.submit(
+	// There must be a kernel inside a command group
+	[&](sycl::handler& cgh) {});
+}

examples/concepts/sycl/kernels/NonDivisibleNDRange.cpp

@@ -0,0 +1,11 @@
+#include <sycl/sycl.hpp>
+
+void main() {
+	sycl::queue myQueue;
+
+	sycl::event myEvent = myQueue.submit(
+	[&](sycl::handler& cgh) {
+		// 6 in global range is not divisible by 4 in local range
+		cgh.parallel_for(sycl::nd_range<2>(sycl::range<2>(6,4), sycl::range<2>(4,2)), [=] (sycl::nd_item<2> it) {});
+	});
+}

examples/concepts/sycl/kernels/NonMatchingRangeItem.cpp

@@ -0,0 +1,11 @@
+#include <sycl/sycl.hpp>
+
+void main() {
+	sycl::queue myQueue;
+
+	sycl::event myEvent = myQueue.submit(
+	[&](sycl::handler& cgh) {
+		// For range we need an item not nd_item
+		cgh.parallel_for(sycl::range<3>(6,4,2), [=] (sycl::nd_item<3> it) {});
+	});
+}

examples/concepts/sycl/kernels/NonMatchingRangeItemDimensions.cpp

@@ -0,0 +1,11 @@
+#include <sycl/sycl.hpp>
+
+void main() {
+	sycl::queue myQueue;
+
+	sycl::event myEvent = myQueue.submit(
+	[&](sycl::handler& cgh) {
+		// Dimensions of range and item should match
+		cgh.parallel_for(sycl::range<3>(6,4,2), [=] (sycl::item<2> it) {});
+	});
+}

examples/concepts/sycl/kernels/NonSYCLCodeInCommandGroup.cpp

@@ -0,0 +1,11 @@
+#include <sycl/sycl.hpp>
+
+void main() {
+	sycl::queue myQueue;
+
+	sycl::event myEvent = myQueue.submit(
+	[&](sycl::handler& cgh) {
+		int a = 5;
+		cgh.parallel_for(sycl::range<2>(6,2), [=] (sycl::item<2> it) {});
+	});
+}

examples/concepts/sycl/kernels/TooHighKernelDimension.cpp

@@ -0,0 +1,12 @@
+#include <sycl/sycl.hpp>
+
+void main() {
+	sycl::queue myQueue;
+
+	sycl::event myEvent = myQueue.submit(
+	[&](sycl::handler& cgh) {
+		cgh.parallel_for(sycl::range<3>(6,4,2), [=] (sycl::item<3> it) {
+			int a = it.get_id(3);
+		});
+	});
+}

examples/concepts/sycl/kernels/ZeroNDRange.cpp

@@ -0,0 +1,11 @@
+#include <sycl/sycl.hpp>
+
+void main() {
+	sycl::queue myQueue;
+
+	sycl::event myEvent = myQueue.submit(
+	[&](sycl::handler& cgh) {
+		// 6 in global range is not divisible by 0 in local range as you cannot divide by zero
+		cgh.parallel_for(sycl::nd_range<2>(sycl::range<2>(6,4), sycl::range<2>(0,2)), [=] (sycl::nd_item<2> it) {});
+	});
+}