[SYCL] Explicit fill broken when using nested template types

[j-stephan]

Example code:

#include <CL/sycl.hpp>

template <typename S>
struct parent
{
    template <typename T>
    struct foo
    {
        T val;
    };

    using Mask = foo<bool>;
    int some_val;
};


auto main() -> int
{
    using type = typename parent<int>::Mask;
    auto queue = cl::sycl::queue{cl::sycl::default_selector{}};
    auto buf = cl::sycl::buffer<type>{cl::sycl::range<1>{1024}};

    queue.submit([&](cl::sycl::handler& cgh)
    {
        auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
        cgh.fill(acc, type{true});
    });
    queue.wait();

    return 0;
}

Command line:

clang++ -std=c++2a -fsycl bool.cpp -ftemplate-backtrace-limit=0 -o bool -lOpenCL

Error message:

In file included from <built-in>:1:
/tmp/bool-29f3c3.h:44:43: error: no member named 'parent' in the global namespace
template <> struct KernelInfo<::__fill< ::parent<int>::foo<_Bool>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>> {
                                        ~~^
/tmp/bool-29f3c3.h:44:53: error: expected '(' for function-style cast or type construction
template <> struct KernelInfo<::__fill< ::parent<int>::foo<_Bool>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>> {
                                                 ~~~^
/tmp/bool-29f3c3.h:44:60: error: use of undeclared identifier '_Bool'
template <> struct KernelInfo<::__fill< ::parent<int>::foo<_Bool>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>> {

[j-stephan]

I tried to work around this by implementing a simple loop in a single_task which produced the same error. So this seems to be tied to the nested templates and not fill.

[alexbatashev]

@j-stephan please, try to eliminate bool from your example. If it doesn't resolve your issue, check if PR #464 helps you.

[j-stephan]

Replacing bool with int produces the following error message:

In file included from <built-in>:1:
/tmp/bool-664cf7.h:39:43: error: no member named 'parent' in the global namespace
template <> struct KernelInfo<::__fill< ::parent<int>::foo<int>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>> {
                                        ~~^
/tmp/bool-664cf7.h:39:53: error: expected '(' for function-style cast or type construction
template <> struct KernelInfo<::__fill< ::parent<int>::foo<int>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>> {
                                                 ~~~^
/tmp/bool-664cf7.h:39:20: error: too many template arguments for class template 'KernelInfo'
template <> struct KernelInfo<::__fill< ::parent<int>::foo<int>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>> {
                   ^                                             ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/jan/software/sycl/intel/lib/clang/9.0.0/include/CL/sycl/detail/kernel_desc.hpp:55:40: note: template is declared here
template <class KernelNameType> struct KernelInfo {
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~        ^
In file included from <built-in>:1:
/tmp/bool-664cf7.h:39:185: error: expected unqualified-id
template <> struct KernelInfo<::__fill< ::parent<int>::foo<int>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>> {
                                                                                                                                                                                        ^
4 errors generated.

I will try out the pull request and report back.

[j-stephan]

Applying the PR didn't change anything, the error message is identical to the one in my previous comment.

[idubinov]

As far as I understand, the problem is in the integration header generation. The function SYCLIntegrationHeader::emitForwardClassDecls in SemaSYCL.cpp does not add enough information to integration header.
This problem refers to the types (TemplateArgument::ArgKind::Type), not template template parameters (TemplateArgument::ArgKind::Template) and PR #464 won't solve the problem.
In the integration header there is forward declaration about struct foo (template struct foo) in global space. There is not any information about parent structure in the integration header.

forward declaration block:

// Forward declarations of templated kernel function types:
template <typename T> struct foo;
template <typename DataT, int Dimensions, cl::sycl::access::mode AccessMode, cl::sycl::access::target AccessTarget, cl::sycl::access::placeholder IsPlaceholder> class __fill;

Kernel specialization:

// Specializations of KernelInfo for kernel function types:
template <> struct KernelInfo<::__fill< ::parent<int>::foo<_Bool>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>> {
<...>
}

there is not any info about parent structure

[j-stephan]

Is there an ETA for a fix? Unfortunately I can't easily work around this issue in the original codebase without investing a few weeks.

[ilyastepykin]

We have the following issue here: fill uses the type of its parameter as part of the kernel name. Since this type is a nested class, the compiler doesn't know how to handle it, so it doesn't emit declaration for the outer class into the integration header. This is why these errors appear.
But we can't emit a proper forward declaration for a nested class as it's not allowed by the language. There seems to be a compiler diagnostic for such cases, but it failed to catch this particular one.

@j-stephan There should be a simple workaround for this issue: as long as the nested class doesn't appear in kernel name, you won't get the errors. So you can re-write fill with parallel_for using something else to name the kernel. You can also use -fsycl-unnamed-lambda compiler option. It's an extension which allows to use kernel lambdas without naming them. In this case you can still use fill with nested types as it doesn't require forward declarations.

BTW, according to the SYCL specification, fill's type T must be an integral scalar value or a SYCL vector type. But in the example, the type is neither one of them. So it make sense to use parallel_for instead.

[j-stephan]

BTW, according to the SYCL specification, fill's type T must be an integral scalar value or a SYCL vector type. But in the example, the type is neither one of them. So it make sense to use parallel_for instead.

Thanks, I didn't notice I wasn't conforming to the specification here. I will change that ASAP.

There should be a simple workaround for this issue:

Thanks for the workaround options, I'll keep them in mind. The compiler flag looks pretty useful.

For the original codebase I could fortunately apply a very simple "fix" by reinterpreting the type as uint8_t and then doing a memset with fill.

[keryell]

I have the feeling that the SYCL specification should allow any trivially copyable data type.

[ilyastepykin]

Looks like the restriction to use only scalar/vector types comes from OpenCL's clEnqueueFillBuffer function. For other trivially copyable types fill can be done via a fallback to regular parallel_for. Right now the same thing is done for 2D and 3D buffers as there is no corresponding function in OpenCL.
I guess it make sense to allow fill with any trivially copyable types.

[TApplencourt]

Sorry for the necromancy, but this seems to have be fixed. At least your example compile in run in my version of the compiler

[bader]

@TApplencourt, thanks for confirming. I'm closing this issue.
@j-stephan, please, let us know if you still have issues with building your example.