diff --git a/include/handler.h b/include/handler.h
index dcaf3a424..bbdda1932 100644
--- a/include/handler.h
+++ b/include/handler.h
@@ -297,17 +297,15 @@ namespace detail {
 			reductions.push_back(rid);
 		}
 
-		void create_host_compute_task(int dimensions, range<3> global_range, id<3> global_offset, range<3> granularity) {
+		void create_host_compute_task(task_geometry geometry) {
 			assert(task == nullptr);
-			task = detail::task::make_host_compute(tid, dimensions, global_range, global_offset, granularity, std::move(cgf), std::move(access_map),
-			    std::move(side_effect_map), std::move(reductions));
+			task = detail::task::make_host_compute(tid, geometry, std::move(cgf), std::move(access_map), std::move(side_effect_map), std::move(reductions));
 		}
 
-		void create_device_compute_task(int dimensions, range<3> global_range, id<3> global_offset, range<3> granularity, std::string debug_name) {
+		void create_device_compute_task(task_geometry geometry, std::string debug_name) {
 			assert(task == nullptr);
 			if(!side_effect_map.empty()) { throw std::runtime_error{"Side effects cannot be used in device kernels"}; }
-			task = detail::task::make_device_compute(
-			    tid, dimensions, global_range, global_offset, granularity, std::move(cgf), std::move(access_map), std::move(reductions), std::move(debug_name));
+			task = detail::task::make_device_compute(tid, geometry, std::move(cgf), std::move(access_map), std::move(reductions), std::move(debug_name));
 		}
 
 		void create_collective_task(collective_group_id cgid) {
@@ -591,8 +589,8 @@ void handler::parallel_for_kernel_and_reductions(range<Dims> global_range, id<Di
 				granularity[d] = local_range[d];
 			}
 		}
-		return dynamic_cast<detail::prepass_handler&>(*this).create_device_compute_task(
-		    Dims, detail::range_cast<3>(global_range), detail::id_cast<3>(global_offset), granularity, detail::kernel_debug_name<KernelName>());
+		const detail::task_geometry geometry{Dims, detail::range_cast<3>(global_range), detail::id_cast<3>(global_offset), granularity};
+		return dynamic_cast<detail::prepass_handler&>(*this).create_device_compute_task(geometry, detail::kernel_debug_name<KernelName>());
 	}
 
 	auto& device_handler = dynamic_cast<detail::live_pass_device_handler&>(*this);
@@ -638,8 +636,8 @@ void handler::host_task(experimental::collective_tag tag, Functor kernel) {
 template <int Dims, typename Functor>
 void handler::host_task(range<Dims> global_range, id<Dims> global_offset, Functor kernel) {
 	if(is_prepass()) {
-		dynamic_cast<detail::prepass_handler&>(*this).create_host_compute_task(
-		    Dims, detail::range_cast<3>(global_range), detail::id_cast<3>(global_offset), {1, 1, 1});
+		const detail::task_geometry geometry{Dims, detail::range_cast<3>(global_range), detail::id_cast<3>(global_offset), {1, 1, 1}};
+		dynamic_cast<detail::prepass_handler&>(*this).create_host_compute_task(geometry);
 	} else {
 		dynamic_cast<detail::live_pass_host_handler&>(*this).schedule<Dims>(kernel);
 	}
diff --git a/include/task.h b/include/task.h
index 9df8c6c6d..29d191a89 100644
--- a/include/task.h
+++ b/include/task.h
@@ -88,6 +88,13 @@ namespace detail {
 		void add_side_effect(host_object_id hoid, experimental::side_effect_order order);
 	};
 
+	struct task_geometry {
+		int dimensions = 0;
+		cl::sycl::range<3> global_size{0, 0, 0};
+		cl::sycl::id<3> global_offset{};
+		cl::sycl::range<3> granularity{1, 1, 1};
+	};
+
 	class task : public intrusive_graph_node<task> {
 	  public:
 		task_type get_type() const { return type; }
@@ -102,13 +109,15 @@ namespace detail {
 
 		const command_group_storage_base& get_command_group() const { return *cgf; }
 
-		int get_dimensions() const { return dimensions; }
+		const task_geometry& get_geometry() const { return geometry; }
+
+		int get_dimensions() const { return geometry.dimensions; }
 
-		cl::sycl::range<3> get_global_size() const { return global_size; }
+		cl::sycl::range<3> get_global_size() const { return geometry.global_size; }
 
-		cl::sycl::id<3> get_global_offset() const { return global_offset; }
+		cl::sycl::id<3> get_global_offset() const { return geometry.global_offset; }
 
-		cl::sycl::range<3> get_granularity() const { return granularity; }
+		cl::sycl::range<3> get_granularity() const { return geometry.granularity; }
 
 		const std::string& get_debug_name() const { return debug_name; }
 
@@ -129,60 +138,54 @@ namespace detail {
 		const std::vector<reduction_id>& get_reductions() const { return reductions; }
 
 		static std::unique_ptr<task> make_nop(task_id tid) {
-			return std::unique_ptr<task>(new task(tid, task_type::NOP, {}, 0, {0, 0, 0}, {}, {1, 1, 1}, nullptr, {}, {}, {}, {}));
+			return std::unique_ptr<task>(new task(tid, task_type::NOP, collective_group_id{}, task_geometry{}, nullptr, {}, {}, {}, {}));
 		}
 
-		static std::unique_ptr<task> make_host_compute(task_id tid, int dimensions, cl::sycl::range<3> global_size, cl::sycl::id<3> global_offset,
-		    cl::sycl::range<3> granularity, std::unique_ptr<command_group_storage_base> cgf, buffer_access_map access_map, side_effect_map side_effect_map,
-		    std::vector<reduction_id> reductions) {
-			return std::unique_ptr<task>(new task(tid, task_type::HOST_COMPUTE, {}, dimensions, global_size, global_offset, granularity, std::move(cgf),
-			    std::move(access_map), std::move(side_effect_map), std::move(reductions), {}));
+		static std::unique_ptr<task> make_host_compute(task_id tid, task_geometry geometry, std::unique_ptr<command_group_storage_base> cgf,
+		    buffer_access_map access_map, side_effect_map side_effect_map, std::vector<reduction_id> reductions) {
+			return std::unique_ptr<task>(new task(tid, task_type::HOST_COMPUTE, collective_group_id{}, geometry, std::move(cgf), std::move(access_map),
+			    std::move(side_effect_map), std::move(reductions), {}));
 		}
 
-		static std::unique_ptr<task> make_device_compute(task_id tid, int dimensions, cl::sycl::range<3> global_size, cl::sycl::id<3> global_offset,
-		    cl::sycl::range<3> granularity, std::unique_ptr<command_group_storage_base> cgf, buffer_access_map access_map, std::vector<reduction_id> reductions,
-		    std::string debug_name) {
-			return std::unique_ptr<task>(new task(tid, task_type::DEVICE_COMPUTE, {}, dimensions, global_size, global_offset, granularity, std::move(cgf),
-			    std::move(access_map), {}, std::move(reductions), std::move(debug_name)));
+		static std::unique_ptr<task> make_device_compute(task_id tid, task_geometry geometry, std::unique_ptr<command_group_storage_base> cgf,
+		    buffer_access_map access_map, std::vector<reduction_id> reductions, std::string debug_name) {
+			return std::unique_ptr<task>(new task(tid, task_type::DEVICE_COMPUTE, collective_group_id{}, geometry, std::move(cgf), std::move(access_map), {},
+			    std::move(reductions), std::move(debug_name)));
 		}
 
 		static std::unique_ptr<task> make_collective(task_id tid, collective_group_id cgid, size_t num_collective_nodes,
 		    std::unique_ptr<command_group_storage_base> cgf, buffer_access_map access_map, side_effect_map side_effect_map) {
-			return std::unique_ptr<task>(new task(tid, task_type::COLLECTIVE, cgid, 1, detail::range_cast<3>(cl::sycl::range<1>{num_collective_nodes}), {},
-			    {1, 1, 1}, std::move(cgf), std::move(access_map), std::move(side_effect_map), {}, {}));
+			const task_geometry geometry{1, detail::range_cast<3>(cl::sycl::range<1>{num_collective_nodes}), {}, {1, 1, 1}};
+			return std::unique_ptr<task>(
+			    new task(tid, task_type::COLLECTIVE, cgid, geometry, std::move(cgf), std::move(access_map), std::move(side_effect_map), {}, {}));
 		}
 
 		static std::unique_ptr<task> make_master_node(
 		    task_id tid, std::unique_ptr<command_group_storage_base> cgf, buffer_access_map access_map, side_effect_map side_effect_map) {
-			return std::unique_ptr<task>(new task(
-			    tid, task_type::MASTER_NODE, {}, 0, {0, 0, 0}, {}, {1, 1, 1}, std::move(cgf), std::move(access_map), std::move(side_effect_map), {}, {}));
+			return std::unique_ptr<task>(new task(tid, task_type::MASTER_NODE, collective_group_id{}, task_geometry{}, std::move(cgf), std::move(access_map),
+			    std::move(side_effect_map), {}, {}));
 		}
 
 		static std::unique_ptr<task> make_horizon_task(task_id tid) {
-			return std::unique_ptr<task>(new task(tid, task_type::HORIZON, {}, 0, {0, 0, 0}, {}, {1, 1, 1}, nullptr, {}, {}, {}, {}));
+			return std::unique_ptr<task>(new task(tid, task_type::HORIZON, collective_group_id{}, task_geometry{}, nullptr, {}, {}, {}, {}));
 		}
 
 	  private:
 		task_id tid;
 		task_type type;
 		collective_group_id cgid;
-		int dimensions;
-		cl::sycl::range<3> global_size;
-		cl::sycl::id<3> global_offset;
-		cl::sycl::range<3> granularity;
+		task_geometry geometry;
 		std::unique_ptr<command_group_storage_base> cgf;
 		buffer_access_map access_map;
 		detail::side_effect_map side_effect_map;
 		std::vector<reduction_id> reductions;
 		std::string debug_name;
 
-		task(task_id tid, task_type type, collective_group_id cgid, int dimensions, cl::sycl::range<3> global_size, cl::sycl::id<3> global_offset,
-		    cl::sycl::range<3> granularity, std::unique_ptr<command_group_storage_base> cgf, buffer_access_map access_map,
-		    detail::side_effect_map side_effect_map, std::vector<reduction_id> reductions, std::string debug_name)
-		    : tid(tid), type(type), cgid(cgid), dimensions(dimensions), global_size(global_size), global_offset(global_offset), granularity(granularity),
-		      cgf(std::move(cgf)), access_map(std::move(access_map)), side_effect_map(std::move(side_effect_map)), reductions(std::move(reductions)),
-		      debug_name(std::move(debug_name)) {
-			assert(type == task_type::HOST_COMPUTE || type == task_type::DEVICE_COMPUTE || granularity.size() == 1);
+		task(task_id tid, task_type type, collective_group_id cgid, task_geometry geometry, std::unique_ptr<command_group_storage_base> cgf,
+		    buffer_access_map access_map, detail::side_effect_map side_effect_map, std::vector<reduction_id> reductions, std::string debug_name)
+		    : tid(tid), type(type), cgid(cgid), geometry(geometry), cgf(std::move(cgf)), access_map(std::move(access_map)),
+		      side_effect_map(std::move(side_effect_map)), reductions(std::move(reductions)), debug_name(std::move(debug_name)) {
+			assert(type == task_type::HOST_COMPUTE || type == task_type::DEVICE_COMPUTE || get_granularity().size() == 1);
 			assert((type != task_type::HOST_COMPUTE && type != task_type::COLLECTIVE && type != task_type::MASTER_NODE) || side_effect_map.empty());
 		}
 	};