KhronosGroup / KhronosGroup/SYCL-Docs
Spec should disallow named kernel object member variable types that cannot be kernel arguments
- 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