Remove explicit kernel names from examples

examples/convolution/convolution.cc

@@ -65,7 +65,7 @@ int main(int argc, char* argv[]) {
 		celerity::accessor gauss{gaussian_mat_buf, cgh, celerity::access::all{}, celerity::read_only};
 		celerity::accessor out{image_tmp_buf, cgh, celerity::access::one_to_one{}, celerity::write_only, celerity::no_init};
 
-		cgh.parallel_for<class gaussian_blur>(celerity::range<2>(image_height, image_width), [=, fs = filter_size](celerity::item<2> item) {
+		cgh.parallel_for(celerity::range<2>(image_height, image_width), [=, fs = filter_size](celerity::item<2> item) {
 			using sycl::float3;
 			if(is_on_boundary(celerity::range<2>(image_height, image_width), fs, item)) {
 				out[item] = float3(0.f, 0.f, 0.f);
@@ -89,7 +89,7 @@ int main(int argc, char* argv[]) {
 		celerity::accessor in{image_tmp_buf, cgh, celerity::access::neighborhood{1, 1}, celerity::read_only};
 		celerity::accessor out{image_output_buf, cgh, celerity::access::one_to_one{}, celerity::write_only, celerity::no_init};
 
-		cgh.parallel_for<class sharpen>(celerity::range<2>(image_height, image_width), [=, fs = filter_size](celerity::item<2> item) {
+		cgh.parallel_for(celerity::range<2>(image_height, image_width), [=, fs = filter_size](celerity::item<2> item) {
 			using sycl::float3;
 			if(is_on_boundary(celerity::range<2>(image_height, image_width), fs, item)) {
 				out[item] = float3(0.f, 0.f, 0.f);

examples/distr_io/distr_io.cc

@@ -119,7 +119,7 @@ int main(int argc, char* argv[]) {
 		q.submit([&](celerity::handler& cgh) {
 			celerity::accessor a{in, cgh, celerity::access::one_to_one{}, celerity::read_only};
 			celerity::accessor b{out, cgh, transposed, celerity::write_only, celerity::no_init};
-			cgh.parallel_for<class transpose>(celerity::range<2>{N, N}, [=](celerity::item<2> item) {
+			cgh.parallel_for(celerity::range<2>{N, N}, [=](celerity::item<2> item) {
 				auto id = item.get_id();
 				b[{id[1], id[0]}] = a[id];
 			});

examples/matmul/matmul.cc

@@ -13,7 +13,7 @@ void set_identity(celerity::distr_queue queue, celerity::buffer<T, 2> mat, bool
 	queue.submit([&](celerity::handler& cgh) {
 		celerity::accessor dw{mat, cgh, celerity::access::one_to_one{}, celerity::write_only, celerity::no_init};
 		const auto range = mat.get_range();
-		cgh.parallel_for<class set_identity_kernel>(range, [=](celerity::item<2> item) {
+		cgh.parallel_for(range, [=](celerity::item<2> item) {
 			if(!reverse) {
 				dw[item] = item[0] == item[1];
 			} else {
@@ -36,7 +36,7 @@ void multiply(celerity::distr_queue queue, celerity::buffer<T, 2> mat_a, celerit
 		celerity::local_accessor<T, 2> scratch_a{{group_size, group_size}, cgh};
 		celerity::local_accessor<T, 2> scratch_b{{group_size, group_size}, cgh};
 
-		cgh.parallel_for<class mat_mul>(celerity::nd_range<2>{{MAT_SIZE, MAT_SIZE}, {group_size, group_size}}, [=](celerity::nd_item<2> item) {
+		cgh.parallel_for(celerity::nd_range<2>{{MAT_SIZE, MAT_SIZE}, {group_size, group_size}}, [=](celerity::nd_item<2> item) {
 			T sum{};
 			const auto lid = item.get_local_id();
 			for(size_t j = 0; j < MAT_SIZE; j += group_size) {

examples/reduction/reduction.cc

@@ -77,7 +77,7 @@ int main(int argc, char* argv[]) {
 		celerity::accessor srgb_255_acc{srgb_255_buf, cgh, celerity::access::one_to_one{}, celerity::read_only};
 		celerity::accessor rgb_acc{lab_buf, cgh, celerity::access::one_to_one{}, celerity::write_only, celerity::no_init};
 		auto minmax_r = celerity::reduction(minmax_buf, cgh, minmax_identity, minmax, celerity::property::reduction::initialize_to_identity{});
-		cgh.parallel_for<class linearize_and_accumulate>(image_size, minmax_r, [=](celerity::item<2> item, auto& minmax) {
+		cgh.parallel_for(image_size, minmax_r, [=](celerity::item<2> item, auto& minmax) {
 			const auto rgb = srgb_to_rgb(srgb_255_acc[item].convert<float>() / 255.0f);
 			rgb_acc[item] = rgb;
 			for(int i = 0; i < 3; ++i) {
@@ -95,7 +95,7 @@ int main(int argc, char* argv[]) {
 		celerity::accessor rgb_acc{lab_buf, cgh, celerity::access::one_to_one{}, celerity::read_only};
 		celerity::accessor minmax_acc{minmax_buf, cgh, celerity::access::all{}, celerity::read_only};
 		celerity::accessor srgb_255_acc{srgb_255_buf, cgh, celerity::access::one_to_one{}, celerity::write_only, celerity::no_init};
-		cgh.parallel_for<class correct_and_compress>(image_size, [=](celerity::item<2> item) {
+		cgh.parallel_for(image_size, [=](celerity::item<2> item) {
 			const auto min = minmax_acc[0][0], max = minmax_acc[0][1];
 			auto rgb = rgb_acc[item];
 			for(int i = 0; i < 3; ++i) {

examples/syncing/syncing.cc

@@ -11,7 +11,7 @@ int main(int argc, char* argv[]) {
 	// Initialize buffer in a distributed device kernel
 	queue.submit([&](celerity::handler& cgh) {
 		celerity::accessor b{buf, cgh, celerity::access::one_to_one{}, celerity::write_only, celerity::no_init};
-		cgh.parallel_for<class write_linear_id>(buf.get_range(), [=](celerity::item<1> item) { b[item] = item.get_linear_id(); });
+		cgh.parallel_for(buf.get_range(), [=](celerity::item<1> item) { b[item] = item.get_linear_id(); });
 	});
 
 	// Process values on the host

examples/wave_sim/wave_sim.cc

@@ -8,7 +8,7 @@
 void setup_wave(celerity::distr_queue& queue, celerity::buffer<float, 2> u, sycl::float2 center, float amplitude, sycl::float2 sigma) {
 	queue.submit([&](celerity::handler& cgh) {
 		celerity::accessor dw_u{u, cgh, celerity::access::one_to_one{}, celerity::write_only, celerity::no_init};
-		cgh.parallel_for<class setup_wave>(u.get_range(), [=, c = center, a = amplitude, s = sigma](celerity::item<2> item) {
+		cgh.parallel_for(u.get_range(), [=, c = center, a = amplitude, s = sigma](celerity::item<2> item) {
 			const float dx = item[1] - c.x();
 			const float dy = item[0] - c.y();
 			dw_u[item] = a * sycl::exp(-(dx * dx / (2.f * s.x() * s.x()) + dy * dy / (2.f * s.y() * s.y())));
@@ -19,7 +19,7 @@ void setup_wave(celerity::distr_queue& queue, celerity::buffer<float, 2> u, sycl
 void zero(celerity::distr_queue& queue, celerity::buffer<float, 2> buf) {
 	queue.submit([&](celerity::handler& cgh) {
 		celerity::accessor dw_buf{buf, cgh, celerity::access::one_to_one{}, celerity::write_only, celerity::no_init};
-		cgh.parallel_for<class zero>(buf.get_range(), [=](celerity::item<2> item) { dw_buf[item] = 0.f; });
+		cgh.parallel_for(buf.get_range(), [=](celerity::item<2> item) { dw_buf[item] = 0.f; });
 	});
 }
 
@@ -35,14 +35,14 @@ struct update_config {
 	static constexpr float c = 1.f;
 };
 
-template <typename T, typename Config, typename KernelName>
+template <typename T, typename Config>
 void step(celerity::distr_queue& queue, celerity::buffer<T, 2> up, celerity::buffer<T, 2> u, float dt, sycl::float2 delta) {
 	queue.submit([&](celerity::handler& cgh) {
 		celerity::accessor rw_up{up, cgh, celerity::access::one_to_one{}, celerity::read_write};
 		celerity::accessor r_u{u, cgh, celerity::access::neighborhood{1, 1}, celerity::read_only};
 
 		const auto size = up.get_range();
-		cgh.parallel_for<KernelName>(size, [=](celerity::item<2> item) {
+		cgh.parallel_for(size, [=](celerity::item<2> item) {
 			const size_t py = item[0] < size[0] - 1 ? item[0] + 1 : item[0];
 			const size_t my = item[0] > 0 ? item[0] - 1 : item[0];
 			const size_t px = item[1] < size[1] - 1 ? item[1] + 1 : item[1];
@@ -56,11 +56,11 @@ void step(celerity::distr_queue& queue, celerity::buffer<T, 2> up, celerity::buf
 }
 
 void initialize(celerity::distr_queue& queue, celerity::buffer<float, 2> up, celerity::buffer<float, 2> u, float dt, sycl::float2 delta) {
-	step<float, init_config, class initialize>(queue, up, u, dt, delta);
+	step<float, init_config>(queue, up, u, dt, delta);
 }
 
 void update(celerity::distr_queue& queue, celerity::buffer<float, 2> up, celerity::buffer<float, 2> u, float dt, sycl::float2 delta) {
-	step<float, update_config, class update>(queue, up, u, dt, delta);
+	step<float, update_config>(queue, up, u, dt, delta);
 }
 
 void stream_open(celerity::distr_queue& queue, size_t N, size_t num_samples, celerity::experimental::host_object<std::ofstream> os) {