KhronosGroup / KhronosGroup/SYCL-Docs

Spec should disallow named kernel object member variable types that cannot be kernel arguments

Open
#524 10 comments 0 reactions 0 assignees View on GitHub
clarification
Dominant language
JavaScript
Stars
158
Forks
80
Avg merge
7d 6h
Merged PRs (30d)
5

Description

We think this statement in [section 4.12.4](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:kernel.parameter.passing) "Rules for parameter passing to kernels" is not quite right:

> If the kernel is a named function object, the `operator()` member function (or other member functions that it calls) may reference member variables inside the same named function object. Any such member variables become parameters to the kernel

This wording implies that it would be OK to define a member variable that is not referenced from `operator()` (or by any of the member functions that it calls), even if that member variable is not a legal kernel argument. This is the case, for example, in this sample program:

```
#include
#include
#include

class MyKernel {
public:
MyKernel(sycl::queue q) {
prefix = "The answer: ";
answer = sycl::malloc_shared(1, q);
}

void operator()() const {
*answer = 42;
};

std::string get_answer() {
std::stringstream ss;
ss << prefix << *answer << "\n";
return ss.str();
}

private:
int *answer;
std::string prefix; // This member is not referenced from device code
};

int main() {
sycl::queue q;

MyKernel k{q};
q.single_task(k).wait();
std::cout << k.get_answer();

return 0;
}
```

The question is whether code like this is legal, even though the member variable `prefix` has a type that is not a legal kernel argument.
We think our original intent was that the entire class of the named kernel object (`MyKernel` in the example above) must be a legal kernel argument, and thus the code snippet above is not spec conformant.

In fact, DPC++ does diagnose an error for the code snippet above. My understanding is that AdaptiveCpp passes the entire object as a kernel parameter in cases like this, so I presume the code above would also be illegal in AdaptiveCpp, though I have not checked.

If we agree that this is the intent, the spec statement I quote above should be clarified.

Contributor guide

No contributing guide indexed for this repository

Research direction

Start with section 4.12.4, “Rules for parameter passing to kernels,” and compare its wording with the MyKernel example containing the unreferenced std::string member. Review the reported DPC++ diagnosis and AdaptiveCpp behavior, then clarify whether the complete named kernel object must be a legal kernel argument and update the specification wording accordingly.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
documentation
Issue type
Documentation
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
38/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.