Skip to content

Commit

Permalink
Introduce task_geometry to cut down on task constructor parameters
Browse files Browse the repository at this point in the history
  • Loading branch information
@fknorr
fknorr committed Feb 17, 2022
1 parent f106d5f commit 1cf8d68
Show file tree
Hide file tree
Showing 2 changed files with 44 additions and 43 deletions.
22 changes: 10 additions & 12 deletions include/handler.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -297,28 +297,26 @@ 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);
if(global_range.size() == 0) {
if(geometry.global_size.size() == 0) {
// TODO this can be easily supported by not creating a task in case the execution range is empty
throw std::runtime_error{"The execution range of distributed host tasks must have at least one item"};
}
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(global_range.size() == 0) {
if(geometry.global_size.size() == 0) {
// TODO unless reductions are involved, this can be easily supported by not creating a task in case the execution range is empty.
// Edge case: If the task includes reductions that specify property::reduction::initialize_to_identity, we need to create a task that sets
// the buffer state to an empty pending_reduction_state in the graph_generator. This will cause a trivial reduction_command to be generated on
// each node that reads from the reduction output buffer, initializing it to the identity value locally.
throw std::runtime_error{"The execution range of device tasks must have at least one item"};
}
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) {
Expand Down Expand Up @@ -602,8 +600,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);
Expand Down Expand Up @@ -649,8 +647,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);
}
Expand Down
65 changes: 34 additions & 31 deletions include/task.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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; }
Expand All @@ -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; }

Expand All @@ -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());
}
};
Expand Down

0 comments on commit 1cf8d68

Please sign in to comment.