Constexpr recursive function not working anymore

[Luigi-Crisci]

Describe the bug When I try to compile a kernel with a recursive constexpr function for the CUDA backend I get

error: SYCL kernel cannot call a recursive function

It works well with previous release 2021-12 (https://github.com/intel/llvm/tree/2021-12).

To Reproduce Simple example code:

constexpr int fact(int n){
        return n > 1 ? 1 : fact(n-1);
}

int main(){

        sycl::queue Q{sycl::gpu_selector()};

        Q.submit([&](sycl::handler& cgh) {
                auto global_range = sycl::range<3>(1, 4, 64);
                auto local_range = sycl::range<3>(1, 4, 32);
                cgh.parallel_for(sycl::nd_range<3>(global_range, local_range), [=](sycl::nd_item<3> it) {
                                constexpr int tmp = fact(10);

                });
        }).wait();

}

Environment:

• OS: Ubuntu 20.04 (Docker on Ubuntu 20.04)
• Target device and vendor: Nvidia Tesla V100
• DPC++ version: Release 2022-06 (4043dda356af59d3d88607037955a0728dc0f466)

[rolandschulz]

I don't think this is a bug. Recursive function calls aren't allowed in device code. And there is no exception for constexpr. I think it is a reasonable request, that constexpr should be an exception form that rule. But that makes it a feature request.

[Luigi-Crisci]

Yeah I know recursive functions are now allowed in device code, but it is strange that in a previous release (https://github.com/intel/llvm/tree/2021-12) they work without any issue.

[keryell]

It looks like an oversight of the SYCL specification not being precise enough on the SYCL spirit. While in C++ constexpr does not force a compile-time execution, at least a use-case similar here to a constinit or consteval should be allowed, broadening the range of accepted C++ code in kernel. Typically this kind of code is used in libraries using metaprogramming and should be allowed. I have opened https://github.com/KhronosGroup/SYCL-Docs/issues/267 accordingly.