Skip to content

Commit

Permalink
Fix inconsistencies between spec and implementation
Browse files Browse the repository at this point in the history
  • Loading branch information
Bensuo committed Mar 6, 2024
1 parent 149611d commit 5b89527
Show file tree
Hide file tree
Showing 12 changed files with 25 additions and 20 deletions.
17 changes: 11 additions & 6 deletions sycl/include/sycl/ext/oneapi/experimental/graph.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -452,13 +452,15 @@ class dynamic_parameter : public detail::dynamic_parameter_base {
: sycl::detail::kernel_param_kind_t::kind_std_layout;

public:
dynamic_parameter(experimental::command_graph<graph_state::modifiable> Graph)
/// Constructs a new dynamic parameter.
/// @param Graph The graph associated with this parameter.
/// @param Param A reference value for this parameter used for CTAD.
dynamic_parameter(experimental::command_graph<graph_state::modifiable> Graph,
const ValueT &Param)
: detail::dynamic_parameter_base(Graph), MValue() {}

dynamic_parameter(ValueT InitialValue,
experimental::command_graph<graph_state::modifiable> Graph)
: detail::dynamic_parameter_base(Graph), MValue(InitialValue) {}

/// Updates this dynamic parameter and all registered nodes with a new value.
/// @param NewValue The new value for the parameter.
void update(const ValueT &NewValue) {
MValue = NewValue;
if constexpr (IsAccessor) {
Expand All @@ -472,7 +474,10 @@ class dynamic_parameter : public detail::dynamic_parameter_base {
ValueT MValue;
};

/// Additional CTAD deduction guide.
/// Additional CTAD deduction guides.
template <typename ValueT>
dynamic_parameter(experimental::command_graph<graph_state::modifiable> Graph,
const ValueT &Param) -> dynamic_parameter<ValueT>;
template <graph_state State = graph_state::modifiable>
command_graph(const context &SyclContext, const device &SyclDevice,
const property_list &PropList) -> command_graph<State>;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Graph/Explicit/update_before_finalize.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -31,7 +31,7 @@ int main() {
Queue.memset(PtrA, 0, N * sizeof(int)).wait();
Queue.memset(PtrB, 0, N * sizeof(int)).wait();

exp_ext::dynamic_parameter<int *> InputParam(Graph);
exp_ext::dynamic_parameter InputParam(Graph, PtrA);

auto KernelNode = Graph.add([&](handler &cgh) {
// Register the input pointer, we should be using set_arg but can't
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -32,7 +32,7 @@ int main() {
BufA.set_write_back(false);
BufB.set_write_back(false);

exp_ext::dynamic_parameter InputParam(BufA.get_access(), Graph);
exp_ext::dynamic_parameter InputParam(Graph, BufA.get_access());

auto KernelNode = Graph.add([&](handler &cgh) {
// Register the input pointer, we should be using set_arg but can't
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -32,7 +32,7 @@ int main() {
Queue.memset(PtrA, 0, N * sizeof(int)).wait();
Queue.memset(PtrB, 0, N * sizeof(int)).wait();

exp_ext::dynamic_parameter<int *> InputParam(Graph);
exp_ext::dynamic_parameter InputParam(Graph, PtrA);

auto KernelNode = Graph.add([&](handler &cgh) {
// Register the input pointer, we should be using set_arg but can't
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,7 @@ int main() {
Queue.memset(PtrA, 0, N * sizeof(int)).wait();
Queue.memset(PtrB, 0, N * sizeof(int)).wait();

exp_ext::dynamic_parameter<int *> InputParam(Graph);
exp_ext::dynamic_parameter InputParam(Graph, PtrA);

auto KernelNode = Graph.add([&](handler &cgh) {
// Register the input pointer, we should be using set_arg but can't
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Graph/Explicit/update_with_indices_ptr.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -31,7 +31,7 @@ int main() {
Queue.memset(PtrA, 0, N * sizeof(int)).wait();
Queue.memset(PtrB, 0, N * sizeof(int)).wait();

exp_ext::dynamic_parameter<int *> InputParam(Graph);
exp_ext::dynamic_parameter InputParam(Graph, PtrA);

auto KernelNode = Graph.add([&](handler &cgh) {
// Register the input pointer, we should be using set_arg but can't
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -34,7 +34,7 @@ int main() {
Queue.memset(PtrB, 0, N * sizeof(int)).wait();
Queue.memset(PtrUnused, 0, N * sizeof(int)).wait();

exp_ext::dynamic_parameter<int *> InputParam(Graph);
exp_ext::dynamic_parameter InputParam(Graph, PtrA);

auto KernelNode = Graph.add([&](handler &cgh) {
// Register the input pointer, we should be using set_arg but can't
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -32,7 +32,7 @@ int main() {
Queue.memset(PtrA, 0, N * sizeof(int)).wait();
Queue.memset(PtrB, 0, N * sizeof(int)).wait();

exp_ext::dynamic_parameter<int *> InputParam(Graph);
exp_ext::dynamic_parameter InputParam(Graph, PtrA);

auto KernelNodeA = Graph.add([&](handler &cgh) {
// Register the input pointer, we should be using set_arg but can't
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -39,9 +39,9 @@ int main() {
Queue.memcpy(PtrB, HostDataB.data(), N * sizeof(int)).wait();
Queue.memset(PtrC, 0, N * sizeof(int)).wait();

exp_ext::dynamic_parameter ParamA(PtrA, Graph);
exp_ext::dynamic_parameter ParamB(PtrB, Graph);
exp_ext::dynamic_parameter ParamOut(PtrC, Graph);
exp_ext::dynamic_parameter ParamA(Graph, PtrA);
exp_ext::dynamic_parameter ParamB(Graph, PtrB);
exp_ext::dynamic_parameter ParamOut(Graph, PtrC);

auto KernelNode = Graph.add([&](handler &cgh) {
// Register the input pointers, we should be using set_arg but can't
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -34,7 +34,7 @@ int main() {
Queue.memset(PtrA, 0, N * sizeof(int)).wait();
Queue.memset(PtrB, 0, N * sizeof(int)).wait();

exp_ext::dynamic_parameter<int *> InputParam(SubGraph);
exp_ext::dynamic_parameter InputParam(SubGraph, PtrA);

auto SubKernelNode = SubGraph.add([&](handler &cgh) {
// Register the input pointer, we should be using set_arg but can't
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,7 @@ int main() {

Queue.memset(DeviceData, 0, N * sizeof(int)).wait();

exp_ext::dynamic_parameter InputParam(ScalarValue, Graph);
exp_ext::dynamic_parameter InputParam(Graph, ScalarValue);

auto KernelNode = Graph.add([&](handler &cgh) {
// Register the input scalar, we should be using set_arg but can't
Expand Down
4 changes: 2 additions & 2 deletions sycl/unittests/Extensions/CommandGraph/Update.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,7 @@ TEST_F(CommandGraphTest, UpdatableException) {
TEST_F(CommandGraphTest, DynamicParamRegister) {
// Check that registering a dynamic param with a node from a graph that was
// not passed to its constructor throws.
experimental::dynamic_parameter<int> DynamicParam(Graph);
experimental::dynamic_parameter DynamicParam(Graph, int{});

auto OtherGraph =
experimental::command_graph(Queue.get_context(), Queue.get_device());
Expand Down Expand Up @@ -72,7 +72,7 @@ TEST_F(CommandGraphTest, UpdateNodeTypeExceptions) {
int *PtrA = malloc_device<int>(16, Queue);
int *PtrB = malloc_device<int>(16, Queue);

experimental::dynamic_parameter<int> DynamicParam{Graph};
experimental::dynamic_parameter DynamicParam{Graph, int{}};

ASSERT_NO_THROW(auto NodeKernel = Graph.add([&](sycl::handler &cgh) {
DynamicParam.register_with_node(cgh, 0);
Expand Down

0 comments on commit 5b89527

Please sign in to comment.