Skip to content

Commit

Permalink
Reintroduce unittest for range update errors
Browse files Browse the repository at this point in the history
  • Loading branch information
EwanC committed Oct 31, 2024
1 parent c02b5c0 commit 9bd9cee
Show file tree
Hide file tree
Showing 23 changed files with 274 additions and 239 deletions.
15 changes: 7 additions & 8 deletions sycl/doc/design/CommandGraph.md
Original file line number Diff line number Diff line change
Expand Up @@ -291,18 +291,17 @@ from a `std::unique_ptr` to a `std::shared_ptr` so that multiple nodes and the
the overhead of having to allocate and free copies of the CG when a new active
CG is selected.

The `dynamic_command_group_impl` class contains weak pointers to the nodes which
have been created with it, so that when a new active CG is selected it can
propagate the change to those nodes. The `node_impl` class also contains a
reference to the dynamic command-group that created it, so that when the graph
is finalized each node can use the list of kernels in its dynamic command-group
as part of the `urCommandBufferAppendKernelLaunchExp` call to pass the possible
alternative kernels.
The `dynamic_command_group_impl` class contains a list of weak pointers to the
nodes which have been created with it, so that when a new active CG is selected
it can propagate the change to those nodes. The `dynamic_parameter_impl` class
also contains a list of weak pointers, but to the `dynamic_command_group_impl`
instances of any dynamic command-groups where they are used. This allows
updating the dynamic parameter to propagate to dynamic command-group nodes.

The `sycl::detail::CGExecKernel` class has been added to, so that if the
object was created from an element in the dynamic command-group list, the class
stores a vector of weak pointers to the other alternative command-groups created
from the same dynamic command-group object. This allows the DPC++ scheduler to
from the same dynamic command-group object. This allows the SYCL runtime to
access the list of alternative kernels when calling the UR API to append a
kernel command to a command-buffer.

Expand Down
1 change: 0 additions & 1 deletion sycl/source/detail/graph_impl.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -562,7 +562,6 @@ graph_impl::add(std::shared_ptr<dynamic_command_group_impl> &DynCGImpl,

// Track the dynamic command-group used inside the node object
DynCGImpl->MNodes.push_back(NodeImpl);
NodeImpl->MDynCG = DynCGImpl;

return NodeImpl;
}
Expand Down
9 changes: 3 additions & 6 deletions sycl/source/detail/graph_impl.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -99,8 +99,6 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
/// Stores the executable graph impl associated with this node if it is a
/// subgraph node.
std::shared_ptr<exec_graph_impl> MSubGraphImpl;
/// Dynamic command-group object used in node, if any.
std::shared_ptr<dynamic_command_group_impl> MDynCG;

/// Used for tracking visited status during cycle checks.
bool MVisited = false;
Expand Down Expand Up @@ -158,7 +156,7 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
: enable_shared_from_this(Other), MSuccessors(Other.MSuccessors),
MPredecessors(Other.MPredecessors), MCGType(Other.MCGType),
MNodeType(Other.MNodeType), MCommandGroup(Other.getCGCopy()),
MSubGraphImpl(Other.MSubGraphImpl), MDynCG(Other.MDynCG) {}
MSubGraphImpl(Other.MSubGraphImpl) {}

/// Copy-assignment operator. This will perform a deep-copy of the
/// command group object associated with this node.
Expand All @@ -170,7 +168,6 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
MNodeType = Other.MNodeType;
MCommandGroup = Other.getCGCopy();
MSubGraphImpl = Other.MSubGraphImpl;
MDynCG = Other.MDynCG;
}
return *this;
}
Expand Down Expand Up @@ -420,7 +417,7 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
throw sycl::exception(sycl::errc::invalid,
"Cannot update execution range of a node with an "
"execution range of different dimensions than what "
"the node was originall created with.");
"the node was original created with.");
}

NDRDesc = sycl::detail::NDRDescT{ExecutionRange};
Expand All @@ -441,7 +438,7 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
throw sycl::exception(sycl::errc::invalid,
"Cannot update execution range of a node with an "
"execution range of different dimensions than what "
"the node was originall created with.");
"the node was original created with.");
}

NDRDesc = sycl::detail::NDRDescT{ExecutionRange};
Expand Down
15 changes: 8 additions & 7 deletions sycl/test-e2e/Graph/Update/dyn_cgf_accessor.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -14,9 +14,7 @@

int main() {
queue Queue{};
const size_t N = 1024;
std::vector<int> HostData(N, 0);
buffer Buf{HostData};
buffer<int> Buf{sycl::range<1>(Size)};
Buf.set_write_back(false);
auto Acc = Buf.get_access();

Expand All @@ -28,22 +26,25 @@ int main() {
int PatternA = 42;
auto CGFA = [&](handler &CGH) {
CGH.require(Acc);
CGH.parallel_for(N, [=](item<1> Item) { Acc[Item.get_id()] = PatternA; });
CGH.parallel_for(Size,
[=](item<1> Item) { Acc[Item.get_id()] = PatternA; });
};

int PatternB = 0xA;
auto CGFB = [&](handler &CGH) {
CGH.require(Acc);
CGH.parallel_for(N, [=](item<1> Item) { Acc[Item.get_id()] = PatternB; });
CGH.parallel_for(Size,
[=](item<1> Item) { Acc[Item.get_id()] = PatternB; });
};

auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CGFA, CGFB});
auto DynamicCGNode = Graph.add(DynamicCG);
auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});

Queue.ext_oneapi_graph(ExecGraph).wait();
std::vector<int> HostData(Size, 0);
Queue.copy(Acc, HostData.data()).wait();
for (size_t i = 0; i < N; i++) {
for (size_t i = 0; i < Size; i++) {
assert(HostData[i] == PatternA);
}

Expand All @@ -52,7 +53,7 @@ int main() {
Queue.ext_oneapi_graph(ExecGraph).wait();

Queue.copy(Acc, HostData.data()).wait();
for (size_t i = 0; i < N; i++) {
for (size_t i = 0; i < Size; i++) {
assert(HostData[i] == PatternB);
}

Expand Down
33 changes: 19 additions & 14 deletions sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -15,10 +15,8 @@

int main() {
queue Queue{};
const size_t N = 1024;
int *Ptr = (int *)sycl::malloc_device<int>(N, Queue);
std::vector<int> HostData(N, 0);
buffer Buf{HostData};
int *Ptr = (int *)sycl::malloc_device<int>(Size, Queue);
buffer<int> Buf{sycl::range<1>(Size)};
Buf.set_write_back(false);

exp_ext::command_graph Graph{
Expand All @@ -28,19 +26,21 @@ int main() {

auto RootNode = Graph.add([&](handler &CGH) {
auto Acc = Buf.get_access<access::mode::write>(CGH);
CGH.parallel_for(N, [=](item<1> Item) { Acc[Item.get_id()] = 1; });
CGH.parallel_for(Size, [=](item<1> Item) { Acc[Item.get_id()] = 1; });
});

int PatternA = 42;
auto CGFA = [&](handler &CGH) {
auto Acc = Buf.get_access<access::mode::read_write>(CGH);
CGH.parallel_for(N, [=](item<1> Item) { Acc[Item.get_id()] += PatternA; });
CGH.parallel_for(Size,
[=](item<1> Item) { Acc[Item.get_id()] += PatternA; });
};

int PatternB = 0xA;
auto CGFB = [&](handler &CGH) {
auto Acc = Buf.get_access<access::mode::read_write>(CGH);
CGH.parallel_for(N, [=](item<1> Item) { Acc[Item.get_id()] += PatternB; });
CGH.parallel_for(Size,
[=](item<1> Item) { Acc[Item.get_id()] += PatternB; });
};

auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CGFA, CGFB});
Expand All @@ -49,23 +49,28 @@ int main() {
auto LeafNode = Graph.add([&](handler &CGH) {
auto Acc = Buf.get_access<access::mode::read>(CGH);
CGH.parallel_for(
N, [=](item<1> Item) { Ptr[Item.get_id()] = Acc[Item.get_id()]; });
Size, [=](item<1> Item) { Ptr[Item.get_id()] = Acc[Item.get_id()]; });
});
auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});

Queue.ext_oneapi_graph(ExecGraph).wait();
Queue.copy(Ptr, HostData.data(), N).wait();
for (size_t i = 0; i < N; i++) {
assert(HostData[i] == (PatternA + 1));

std::vector<int> HostData(Size, 0);
Queue.copy(Ptr, HostData.data(), Size).wait();

int Ref = PatternA + 1;
for (size_t i = 0; i < Size; i++) {
assert(HostData[i] == Ref);
}

DynamicCG.set_active_cgf(1);
ExecGraph.update(DynamicCGNode);

Queue.ext_oneapi_graph(ExecGraph).wait();
Queue.copy(Ptr, HostData.data(), N).wait();
for (size_t i = 0; i < N; i++) {
assert(HostData[i] == (PatternB + 1));
Queue.copy(Ptr, HostData.data(), Size).wait();
Ref = PatternB + 1;
for (size_t i = 0; i < Size; i++) {
assert(HostData[i] == Ref);
}

sycl::free(Ptr, Queue);
Expand Down
31 changes: 17 additions & 14 deletions sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps2.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -16,11 +16,9 @@

int main() {
queue Queue{};
const size_t N = 1024;
int *Ptr = (int *)sycl::malloc_device<int>(N, Queue);
std::vector<int> HostData(N, 0);
buffer<int> BufA{sycl::range<1>(N)};
buffer<int> BufB{sycl::range<1>(N)};
int *Ptr = (int *)sycl::malloc_device<int>(Size, Queue);
buffer<int> BufA{sycl::range<1>(Size)};
buffer<int> BufB{sycl::range<1>(Size)};
BufA.set_write_back(false);
BufB.set_write_back(false);

Expand All @@ -34,7 +32,7 @@ int main() {
auto RootNode = Graph.add([&](handler &CGH) {
auto AccA = BufA.get_access<access::mode::write>(CGH);
auto AccB = BufB.get_access<access::mode::write>(CGH);
CGH.parallel_for(N, [=](item<1> Item) {
CGH.parallel_for(Size, [=](item<1> Item) {
AccA[Item.get_id()] = InitA;
AccB[Item.get_id()] = InitB;
});
Expand All @@ -43,13 +41,15 @@ int main() {
int PatternA = 42;
auto CGFA = [&](handler &CGH) {
auto AccA = BufA.get_access<access::mode::read_write>(CGH);
CGH.parallel_for(N, [=](item<1> Item) { AccA[Item.get_id()] += PatternA; });
CGH.parallel_for(Size,
[=](item<1> Item) { AccA[Item.get_id()] += PatternA; });
};

int PatternB = 0xA;
auto CGFB = [&](handler &CGH) {
auto AccB = BufB.get_access<access::mode::read_write>(CGH);
CGH.parallel_for(N, [=](item<1> Item) { AccB[Item.get_id()] += PatternB; });
CGH.parallel_for(Size,
[=](item<1> Item) { AccB[Item.get_id()] += PatternB; });
};

auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CGFA, CGFB});
Expand All @@ -58,25 +58,28 @@ int main() {
auto LeafNode = Graph.add([&](handler &CGH) {
auto AccA = BufA.get_access<access::mode::read>(CGH);
auto AccB = BufB.get_access<access::mode::read>(CGH);
CGH.parallel_for(N, [=](item<1> Item) {
CGH.parallel_for(Size, [=](item<1> Item) {
Ptr[Item.get_id()] = AccA[Item.get_id()] + AccB[Item.get_id()];
});
});
auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});

Queue.ext_oneapi_graph(ExecGraph).wait();
Queue.copy(Ptr, HostData.data(), N).wait();
for (size_t i = 0; i < N; i++) {

std::vector<int> HostData(Size, 0);
Queue.copy(Ptr, HostData.data(), Size).wait();
for (size_t i = 0; i < Size; i++) {
assert(HostData[i] == (InitA + InitB + PatternA));
}

DynamicCG.set_active_cgf(1);
ExecGraph.update(DynamicCGNode);

Queue.ext_oneapi_graph(ExecGraph).wait();
Queue.copy(Ptr, HostData.data(), N).wait();
for (size_t i = 0; i < N; i++) {
assert(HostData[i] == (InitA + InitB + PatternB));
Queue.copy(Ptr, HostData.data(), Size).wait();
int Ref = InitA + InitB + PatternB;
for (size_t i = 0; i < Size; i++) {
assert(HostData[i] == Ref);
}

sycl::free(Ptr, Queue);
Expand Down
16 changes: 8 additions & 8 deletions sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -26,17 +26,13 @@ int main(int, char **argv) {
kernel kernel = getKernel(
KernelBundle, "_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_");

const size_t N = 1024;

exp_ext::command_graph Graph{
Queue.get_context(),
Queue.get_device(),
{exp_ext::property::graph::assume_buffer_outlives_graph{}}};
std::vector<int> HostDataA(N, 0);
std::vector<int> HostDataB(N, 0);

buffer BufA{HostDataA};
buffer BufB{HostDataB};
buffer<int> BufA{sycl::range<1>(Size)};
buffer<int> BufB{sycl::range<1>(Size)};
BufA.set_write_back(false);
BufB.set_write_back(false);

Expand All @@ -60,9 +56,12 @@ int main(int, char **argv) {
auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});

Queue.ext_oneapi_graph(ExecGraph).wait();

std::vector<int> HostDataA(Size, 0);
std::vector<int> HostDataB(Size, 0);
Queue.copy(BufA.get_access(), HostDataA.data()).wait();
Queue.copy(BufB.get_access(), HostDataB.data()).wait();
for (size_t i = 0; i < N; i++) {
for (size_t i = 0; i < Size; i++) {
assert(HostDataA[i] == i);
assert(HostDataB[i] == 0);
}
Expand All @@ -71,9 +70,10 @@ int main(int, char **argv) {
ExecGraph.update(DynamicCGNode);

Queue.ext_oneapi_graph(ExecGraph).wait();

Queue.copy(BufA.get_access(), HostDataA.data()).wait();
Queue.copy(BufB.get_access(), HostDataB.data()).wait();
for (size_t i = 0; i < N; i++) {
for (size_t i = 0; i < Size; i++) {
assert(HostDataA[i] == i);
assert(HostDataB[i] == i);
}
Expand Down
Loading

0 comments on commit 9bd9cee

Please sign in to comment.