[SYCL] Report Compiling error for recursion in SYCL kernel #3390

jinge90 · 2021-03-22T08:02:07Z

Signed-off-by: gejin ge.jin@intel.com
Latest compiler doesn't report error for usage of recursive code in SYCL kernel. If some recursive code is used in SYCL kernel, latest compiler can compile well but runtime error may happen, following code can compile well but crashes using "-O0" on GPU:
recursion.cpp

#include <CL/sycl.hpp>
#include <iostream>
constexpr sycl::access::mode sycl_read = sycl::access::mode::read;
constexpr sycl::access::mode sycl_write = sycl::access::mode::write;
float foo(float x, float y) {
  if (x < y)
    return x;
  return foo(x, y + 1);
}
float run_foo(float x, float y) {
  if ((x + y) > 100) return x - y;
  return foo(x, y);
}
class KernelTest;
void kernel_test() {
  sycl::queue deviceQueue;
  sycl::range<1> numOfItems{1};
  sycl::cl_float result = 0;
  {
    sycl::buffer<sycl::cl_float, 1> buffer1(&result, numOfItems);
    deviceQueue.submit([&](sycl::handler &cgh) {
      auto res_access = buffer1.get_access<sycl_write>(cgh);
      cgh.single_task<class KernelTest>(
          [=]() { res_access[0] = run_foo(5.f, -10.5f); });
    });
  }
  std::cout << "Stored in device: " << result << std::endl;
}
int main() {
  kernel_test();
  return 0;
}

According to SYCL 2020 spec(https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:language.restrictions.kernels), recursion is not allowed in SYCL kernel, we had better report compiling error for this illegal behavior. Moving runtime crash to compiling error is more friendly to users too.

Signed-off-by: gejin <ge.jin@intel.com>

jinge90 · 2021-03-22T08:05:09Z

/summary:run

…ecursive

bader · 2021-03-23T16:27:45Z

+@AaronBallman, who reviewed somewhat related patch to clang-tidy tool.
https://reviews.llvm.org/D72362#1817600

AaronBallman

As far as a narrow fix goes, this seems reasonable. However, there is some confusing code in VisitCallExpr() that should likely be cleaned up (like the dyn_cast<FunctionDecl>(Callee) check when Callee is already defined as a FunctionDecl, or an uncommented exception about pseduo destructor calls). Also, this code seems very fragile (if you don't know to call CollectKernelSet before visiting, then the visitor won't catch any recursive calls at all) and I'm not 100% convinced it isn't overly aggressive at finding recursion. However, none of this is the result of the changes in this PR, so LGTM!

bader · 2021-03-25T18:15:09Z

@jinge90, IIRC, the patch breaks some ESIMD tests, but it looks like it's stuck after recent update.
@tfzhu, could you take a look, please?

keryell · 2021-03-26T02:21:21Z

clang/test/SemaSYCL/restrict-recursion4.cpp

@@ -24,7 +24,7 @@ void *operator new(size_t);
 void usage2(myFuncDef functionPtr) {


Could it be possible to add a test like

void kernel3(void) { if constexpr (false) kernel3(); if (false) kernel3(); }

and verify it does not fire?

Hi, @keryell
I am afraid current implementation can't handle this as we are analyzing the call graph to recognize all recursions and don't know whether the recursive call is in a "dead" branch.
Thanks very much.

I think it's reasonable to not diagnose recursion in dead branches, but we should still have test coverage (with comments) that show we've explicitly considered this case.

This comment is still unaddressed, btw.

@AaronBallman :Note that my patch adds a test for this. See: #3475

jinge90 · 2021-03-26T06:46:00Z

@jinge90, IIRC, the patch breaks some ESIMD tests, but it looks like it's stuck after recent update.
@tfzhu, could you take a look, please?

Hi, @bader
Just checked the failure, it seems that compiler has decided to support "recursive call" for ESIMD kernel, if so, this PR need to support the scenario.
Thanks very much.

bader · 2021-03-26T06:49:01Z

According to my understanding some SYCL targets are able to support recursive calls, so it's quite possible that we would want to enable it for regular SYCL mode as well.
NOTE: @DenisBakhvalov is working on unifying SYCL and ESIMD modes, so if recursion support is required for ESIMD, it might be required for generic mode as well.

Signed-off-by: gejin <ge.jin@intel.com>

…ecursive

jinge90 · 2021-04-01T06:56:57Z

Previously, for all SYCL kernel(or SYCL_EXTERNAL) functions and functions called by them directly or indirectly, compiler will go through them one by one for sema check(recursion check is done in this period). The rules used in the check are same for all those functions. This mechanism seems can't meet our requirements now since different rules may be required for different types of SYCL kernels and all its callees(directly or indirectly). For example, recursion is not allowed for regular SYCL kernel according to spec but recursive function with "NoInline" attr is permitted for SYCL_ESIMD_KERNEL(https://github.com/intel/llvm-test-suite/blob/intel/SYCL/ESIMD/noinline_call_recursive.cpp#L13). So, we need to follow correct rules when checking different type of SYCL kernel and its callees. In order to achieve this, we traverse over the call graph to construct a map whose key is each SYCL kernel or SYCL_EXTERNAL function and value is the list of its direct or indirect callees.
Then, for each SYCL kernel or SYCL_EXTERNAL function, we do the check for it and all its direct and indirect callees. Before the check, we will check the SYCL kernel type to know which rules should follow.
Hi, @AaronBallman @elizabethandrews @premanandrao , could you help to review the patch?
Thanks very much.

jinge90 · 2021-04-01T07:32:54Z

/summary:run

AaronBallman

I think it should be possible to use a const FunctionDecl * just about everywhere in this patch so that we get const-correctness. Can you investigate that?

AaronBallman · 2021-04-01T11:38:32Z

clang/lib/Sema/SemaSYCL.cpp

-            CollectKernelSet(Callee, FD, VisitedSet);
-            VisitedSet.erase(Callee);
+    for (const CallGraphNode *CGN : *SYCLFuncCGN) {
+      if (FunctionDecl *Callee = dyn_cast<FunctionDecl>(CGN->getDecl())) {


Suggested change

if (FunctionDecl *Callee = dyn_cast<FunctionDecl>(CGN->getDecl())) {

if (auto *Callee = dyn_cast<FunctionDecl>(CGN->getDecl())) {

AaronBallman · 2021-04-01T11:38:54Z

clang/lib/Sema/SemaSYCL.cpp

+            if (IsCyclicSYCLFunction(Callee)) {
+              RecursiveSet.insert(Callee);
+            }


Suggested change

if (IsCyclicSYCLFunction(Callee)) {

RecursiveSet.insert(Callee);

}

if (IsCyclicSYCLFunction(Callee))

RecursiveSet.insert(Callee);

AaronBallman · 2021-04-01T11:39:44Z

clang/lib/Sema/SemaSYCL.cpp

+      return false;
+    llvm::SmallPtrSet<FunctionDecl *, 10> VisitedSet;
+    for (const CallGraphNode *CGN : *SYCLFuncCGN) {
+      if (FunctionDecl *Callee = dyn_cast<FunctionDecl>(CGN->getDecl())) {


Suggested change

if (FunctionDecl *Callee = dyn_cast<FunctionDecl>(CGN->getDecl())) {

if (auto *Callee = dyn_cast<FunctionDecl>(CGN->getDecl())) {

AaronBallman · 2021-04-01T11:40:20Z

clang/lib/Sema/SemaSYCL.cpp

+    if (!SYCLFuncCGN)
+      return false;
+    for (const CallGraphNode *CGN : *SYCLFuncCGN) {
+      if (FunctionDecl *Callee = dyn_cast<FunctionDecl>(CGN->getDecl())) {


Suggested change

if (FunctionDecl *Callee = dyn_cast<FunctionDecl>(CGN->getDecl())) {

if (auto *Callee = dyn_cast<FunctionDecl>(CGN->getDecl())) {

AaronBallman · 2021-04-01T11:40:46Z

clang/lib/Sema/SemaSYCL.cpp

+  llvm::SmallPtrSet<FunctionDecl *, 10> NonESIMDModeCheckSet;
+  llvm::SmallPtrSet<FunctionDecl *, 10> ESIMDModeCheckSet;
+  for (Decl *D : syclDeviceDecls()) {
+    if (auto SYCLKernelDec = dyn_cast<FunctionDecl>(D)) {


Suggested change

if (auto SYCLKernelDec = dyn_cast<FunctionDecl>(D)) {

if (auto *SYCLKernelDec = dyn_cast<FunctionDecl>(D)) {

AaronBallman · 2021-04-01T11:42:53Z

clang/lib/Sema/SemaSYCL.cpp

+        Marker.SetTraverseEsimdKernel(SYCLKernelDef->hasAttr<SYCLSimdAttr>());
+        auto &SYCLFunctionCheckSet =
+            (SYCLKernelDef->hasAttr<SYCLSimdAttr>() ? ESIMDModeCheckSet
+                                                    : NonESIMDModeCheckSet);


Suggested change

Marker.SetTraverseEsimdKernel(SYCLKernelDef->hasAttr<SYCLSimdAttr>());

auto &SYCLFunctionCheckSet =

(SYCLKernelDef->hasAttr<SYCLSimdAttr>() ? ESIMDModeCheckSet

: NonESIMDModeCheckSet);

bool IsSIMD = SYCLKernelDef->hasAttr<SYCLSimdAttr>();

Marker.SetTraverseEsimdKernel(IsSIMD);

auto &SYCLFunctionCheckSet = IsSIMD ? ESIMDModeCheckSet : NonESIMDModeCheckSet;

AaronBallman · 2021-04-01T11:46:17Z

clang/test/SemaSYCL/restrict-recursion4.cpp

@@ -24,7 +24,7 @@ void *operator new(size_t);
 void usage2(myFuncDef functionPtr) {


I think it's reasonable to not diagnose recursion in dead branches, but we should still have test coverage (with comments) that show we've explicitly considered this case.

AaronBallman · 2021-04-01T11:50:29Z

Adding @erichkeane as a reviewer for visibility -- he was looking at this same chunk of code recently and may have opinions on whether these changes should be generalized for other uses or not.

erichkeane

While some of this is an improvement, it seems like a lot of it is breaking the cohesion of the marker type, it seems to be gaining dependence on even more external state. I don't believe this to be the right solution to this problem.

erichkeane · 2021-04-01T13:03:25Z

clang/lib/Sema/SemaSYCL.cpp

-  // kernel set, this is used for error diagnostics.
+  // Record the mapping between each SYCL kernel or SYCL_EXTERNAL function and
+  // functions called by it.
+  llvm::DenseMap<FunctionDecl *, llvm::SmallPtrSet<FunctionDecl *, 10>>


This seems to be re-implementing a CallGraphNode. Is there a reason we need to re-store these?

Hi, @erichkeane
I didn't quite understand the question here, did you ask why the patch dropped the "KernelSet" and introduced the "SYCLKernelInvokeMap" here?
Thanks very much.

I'm basically asking 'why' on a lot of things throughout this patch. I'm unconvinced that KernelSet was necessary, and I'm even further unconvinced that SYCLKernelInvokeMap is required. I think this patch is well-intentioned but is the straw that breaks the camels back here.

The complexity of this type and section of code is too much after this patch. We've hit 'unmaintainable' unfortunately. So if you can do as @AaronBallman suggested and put together a bunch of lit tests that validate the edge-cases that you found, he and I will work on refactoring this code ensuring it fixes the issues you saw.

erichkeane · 2021-04-01T13:09:13Z

clang/lib/Sema/SemaSYCL.cpp

+      llvm::SmallPtrSet<FunctionDecl *, 10> SYCLKernelInvokeSet;
+      Marker.WalkSYCLFunctionCG(SYCLKernel, SYCLKernelInvokeSet,
+                                SYCLFunctionCyclicCheckSet);
+      Marker.SYCLKernelInvokeMap.insert(


It seems to me the ownership model is all wrong here. The Marker type I think should just own these operations and collections. We shouldn't have to initialize 4 collections for it.

erichkeane · 2021-04-01T13:10:12Z

clang/lib/Sema/SemaSYCL.cpp

+            (SYCLKernelDef->hasAttr<SYCLSimdAttr>() ? ESIMDModeCheckSet
+                                                    : NonESIMDModeCheckSet);
+        Marker.TraverseStmt(SYCLKernelDef->getBody());
+        for (FunctionDecl *SYCLFunctionDec :


This bit doesn't seem right either. From this function's perspective, all of this should be managed by the marker it self.

erichkeane

Also, this patch needs lit tests... you claim to be fixing a reproducer, but I don't see an added test, nor anything for the esimd versions you're attempting to allow.

AaronBallman · 2021-04-01T13:54:13Z

I think we should probably consider this from a pile of test cases that test the expected-good, expected-bad, and edge case functionality. @keryell brought up the dead code test case, but I suspect an uninstantiated template will run into similar problems, and I'm starting to worry about interplay with macros as well. We may also need to get the SYCL spec clarified as https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:language.restrictions.kernels just says Recursion is not allowed in a device function. which appears to obligate us to diagnose these situations even though I'm not convinced that's actually possible. For example, I don't know how we'd pick up on recursion that spans across TU boundaries.

Can you put together a bunch of really good lit tests that @erichkeane and I can look over? We can work backwards from the tests to see whether we need more refactoring to be done here or not.

jinge90 · 2021-04-01T14:12:10Z

I think we should probably consider this from a pile of test cases that test the expected-good, expected-bad, and edge case functionality. @keryell brought up the dead code test case, but I suspect an uninstantiated template will run into similar problems, and I'm starting to worry about interplay with macros as well. We may also need to get the SYCL spec clarified as https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:language.restrictions.kernels just says Recursion is not allowed in a device function. which appears to obligate us to diagnose these situations even though I'm not convinced that's actually possible. For example, I don't know how we'd pick up on recursion that spans across TU boundaries.

Can you put together a bunch of really good lit tests that @erichkeane and I can look over? We can work backwards from the tests to see whether we need more refactoring to be done here or not.

Sure, I will add some lit tests.
Thanks very much.

DenisBakhvalov · 2021-04-01T16:17:14Z

According to SYCL 2020 spec, recursion is not allowed in SYCL kernels

recursion is not allowed for normal SYCL kernel but recursive function with "noinline"
attribute is permitted in SYCL ESIMD kernel.

I'm not aware of such an ESIMD feature. @kbobrovs, are you?

kbobrovs · 2021-04-01T16:44:16Z

According to SYCL 2020 spec, recursion is not allowed in SYCL kernels

recursion is not allowed for normal SYCL kernel but recursive function with "noinline"
attribute is permitted in SYCL ESIMD kernel.

I'm not aware of such an ESIMD feature. @kbobrovs, are you?

No, this is not supported in ESIMD neither.

erichkeane · 2021-04-01T16:45:45Z

According to SYCL 2020 spec, recursion is not allowed in SYCL kernels

recursion is not allowed for normal SYCL kernel but recursive function with "noinline"
attribute is permitted in SYCL ESIMD kernel.

I'm not aware of such an ESIMD feature. @kbobrovs, are you?

No, this is not supported in ESIMD neither.

Thanks! I won't consider it then when going through the re-implementation/refactor here.

DenisBakhvalov · 2021-04-01T16:49:48Z

According to SYCL 2020 spec, recursion is not allowed in SYCL kernels

recursion is not allowed for normal SYCL kernel but recursive function with "noinline"
attribute is permitted in SYCL ESIMD kernel.

I'm not aware of such an ESIMD feature. @kbobrovs, are you?

No, this is not supported in ESIMD neither.

According to SYCL 2020 spec, recursion is not allowed in SYCL kernels

recursion is not allowed for normal SYCL kernel but recursive function with "noinline"
attribute is permitted in SYCL ESIMD kernel.

I'm not aware of such an ESIMD feature. @kbobrovs, are you?

No, this is not supported in ESIMD neither.

Thanks! I won't consider it then when going through the re-implementation/refactor here.

Yes, so I believe ESIMD should obey the same rules that exist for regular SYCL kernels regarding recursive functions, right?

@jinge90, if your earlier patch triggered errors in ESIMD tests, then I think those tests should be fixed. Do you have a list of those tests maybe?

kbobrovs · 2021-04-01T16:50:56Z

@fveselov, your PR intel/llvm-test-suite#143 adds recursive ESIMD test. This feature is not supported by ESIMD, so can you please remove it. @vladimirlaz, please add me as a code owner for ESIMD llvm-test-suite tests.

vladimirlaz · 2021-04-01T16:54:53Z

@fveselov, your PR intel/llvm-test-suite#143 adds recursive ESIMD test. This feature is not supported by ESIMD, so can you please remove it. @vladimirlaz, please add me as a code owner for ESIMD llvm-test-suite tests.

@kbobrovs I added code owners for intel/llvm-test-suite few weeks ago. Now you are code owner for ESIMD tests: https://github.com/intel/llvm-test-suite/blob/intel/.github/CODEOWNERS#L30

kbobrovs · 2021-04-01T19:43:38Z

@fveselov, your PR intel/llvm-test-suite#143 adds recursive ESIMD test. This feature is not supported by ESIMD, so can you please remove it. @vladimirlaz, please add me as a code owner for ESIMD llvm-test-suite tests.

@kbobrovs I added code owners for intel/llvm-test-suite few weeks ago. Now you are code owner for ESIMD tests: https://github.com/intel/llvm-test-suite/blob/intel/.github/CODEOWNERS#L30

Thanks!

jinge90 · 2021-04-02T00:45:27Z

According to SYCL 2020 spec, recursion is not allowed in SYCL kernels

recursion is not allowed for normal SYCL kernel but recursive function with "noinline"
attribute is permitted in SYCL ESIMD kernel.

I'm not aware of such an ESIMD feature. @kbobrovs, are you?

No, this is not supported in ESIMD neither.

According to SYCL 2020 spec, recursion is not allowed in SYCL kernels

recursion is not allowed for normal SYCL kernel but recursive function with "noinline"
attribute is permitted in SYCL ESIMD kernel.

I'm not aware of such an ESIMD feature. @kbobrovs, are you?

No, this is not supported in ESIMD neither.

Thanks! I won't consider it then when going through the re-implementation/refactor here.

Yes, so I believe ESIMD should obey the same rules that exist for regular SYCL kernels regarding recursive functions, right?

@jinge90, if your earlier patch triggered errors in ESIMD tests, then I think those tests should be fixed. Do you have a list of those tests maybe?

Hi, @DenisBakhvalov
The earlier patch triggers failure in https://github.com/jinge90/llvm-test-suite/blob/intel/SYCL/ESIMD/noinline_call_recursive.cpp#L13 and according to the comments in the test, ESIMD kernel claims to support "noinline" recursive function. If this case should be fixed, I will revert my latest patch to earlier one.
Thanks very much.

erichkeane · 2021-04-02T01:02:20Z

According to SYCL 2020 spec, recursion is not allowed in SYCL kernels

recursion is not allowed for normal SYCL kernel but recursive function with "noinline"
attribute is permitted in SYCL ESIMD kernel.

I'm not aware of such an ESIMD feature. @kbobrovs, are you?

No, this is not supported in ESIMD neither.

According to SYCL 2020 spec, recursion is not allowed in SYCL kernels

recursion is not allowed for normal SYCL kernel but recursive function with "noinline"
attribute is permitted in SYCL ESIMD kernel.

I'm not aware of such an ESIMD feature. @kbobrovs, are you?

No, this is not supported in ESIMD neither.

Thanks! I won't consider it then when going through the re-implementation/refactor here.

Yes, so I believe ESIMD should obey the same rules that exist for regular SYCL kernels regarding recursive functions, right?
@jinge90, if your earlier patch triggered errors in ESIMD tests, then I think those tests should be fixed. Do you have a list of those tests maybe?

Hi, @DenisBakhvalov
The earlier patch triggers failure in https://github.com/jinge90/llvm-test-suite/blob/intel/SYCL/ESIMD/noinline_call_recursive.cpp#L13 and according to the comments in the test, ESIMD kernel claims to support "noinline" recursive function. If this case should be fixed, I will revert my latest patch to earlier one.
Thanks very much.

FYI: Don't bother updating this patch other than adding some lit tests. We are not going to accept the patch, instead I am going to include the tests/test situation into a rewrite of most of the code here. It is horribly disorganized and intermingled with other behaviors, so we're going to clean it up in the next few days.

…ecursive

Signed-off-by: gejin <ge.jin@intel.com>

erichkeane

So at the base of it, it seems that all this does now is switch a warning to an error. I don't really have a problem with that anymore, though I'm a little odded-out by the fact that the error message already existed.

This also only steps on my re-write in 1 place (the SemaSYCL change), but that should be easy enough to merge.

jinge90 · 2021-04-02T01:29:21Z

Hi, @AaronBallman and @erichkeane
The patch is reverted to earlier one since we will remove the failed ESIMD test now. The corresponding lit tests are also updated to track whether compiling error is triggered when recursion is detected in SYCL kernel code.
As we have discussed, I also felt that the code for SYCL kernel sema check is hard to maintain now, for example:

Almost all data member of Marker are public and we can see those internal data members are directly set or initialized in may places.
Many initialization work for Marker are done by Sema::MarkDevice(), in fact, all the initialization of Marker can be done in its constructor.
Agree that we may need a factor to solve those problems.
Thanks very much.

erichkeane · 2021-04-02T01:34:39Z

Hi, @AaronBallman and @erichkeane
The patch is reverted to earlier one since we will remove the failed ESIMD test now. The corresponding lit tests are also updated to track whether compiling error is triggered when recursion is detected in SYCL kernel code.
As we have discussed, I also felt that the code for SYCL kernel sema check is hard to maintain now, for example:
1. Almost all data member of Marker are public and we can see those internal data members are directly set or initialized in may places.

2. Many initialization work for Marker are done by Sema::MarkDevice(), in fact, all the initialization of Marker can be done in its constructor.
   Agree that we may need a factor to solve those problems.
   Thanks very much.

Agreed. Seeing your change here, I don't have any problems with it, it is pretty innocuous in that it only reverts a warning to an error in a case that the device can't handle anyway. I'm OK with that.

I've also captured in my change a couple of other things, including properly not-checking if-constexpr dead branches as @keryell requested. That part likely requires a little work in the CallGraph object to make sure we don't evaluate the dead side of it. Perhaps that is another follow-up to my refactor.

Another thing I have on my list is to fix how we identify 'kernel body wrappers' that get added in the library. THAT is likely going to be another difficult, perhaps separate task.

AaronBallman

I think the changes in this patch are reasonable as far as they go.

AaronBallman · 2021-04-02T12:50:02Z

clang/test/SemaSYCL/restrict-recursion3.cpp

@@ -34,7 +34,7 @@ template <typename name, typename Func>
 __attribute__((sycl_kernel)) void kernel_single_task2(const Func &kernelFunc) {
  // expected-note@+1 {{called by 'kernel_single_task2}}
  kernelFunc();
-  // expected-warning@+1 2{{SYCL kernel cannot call a recursive function}}
+  // expected-error@+1 2{{SYCL kernel cannot call a recursive function}}


Why is this diagnostic issued twice?

I'm still curious about this.

AaronBallman · 2021-04-02T12:52:33Z

@erichkeane I was thinking about the rewrite of this logic rather than sleeping last night, and one thing we should discuss is where this kind of check lives. This seems more like an intraprocedural analysis and that suggests to me that it should either live in Analysis or potentially be considered for a clang static analyzer check rather than a frontend check. In fact, by using a CSA check, we could potentially diagnose far more cases because there's some basic cross-TU support already there, I believe.

erichkeane · 2021-04-02T13:02:38Z

So there are

@erichkeane I was thinking about the rewrite of this logic rather than sleeping last night, and one thing we should discuss is where this kind of check lives. This seems more like an intraprocedural analysis and that suggests to me that it should either live in Analysis or potentially be considered for a clang static analyzer check rather than a frontend check. In fact, by using a CSA check, we could potentially diagnose far more cases because there's some basic cross-TU support already there, I believe.

While I think putting this in the static-analyzer is a good place for it, I don't think it can live in Analysis without us doing more processing than I'd like.

I think Analysis would end up having to run on every function to determine whether it is recursive, then add it to the deferred diagnostics. I think that would be more work, right? At least here we can limit it to only things we 'know' are kernels.

AaronBallman · 2021-04-02T13:24:07Z

So there are

@erichkeane I was thinking about the rewrite of this logic rather than sleeping last night, and one thing we should discuss is where this kind of check lives. This seems more like an intraprocedural analysis and that suggests to me that it should either live in Analysis or potentially be considered for a clang static analyzer check rather than a frontend check. In fact, by using a CSA check, we could potentially diagnose far more cases because there's some basic cross-TU support already there, I believe.

While I think putting this in the static-analyzer is a good place for it, I don't think it can live in Analysis without us doing more processing than I'd like.

I think Analysis would end up having to run on every function to determine whether it is recursive, then add it to the deferred diagnostics. I think that would be more work, right? At least here we can limit it to only things we 'know' are kernels.

Analysis would run on a per-function basis, which is why I think CSA may actually be the better option.

One concern I do have with this patch is that we're strengthening a warning that's an error by default into always being an error, but the analysis we're performing isn't sound. Typically, errors are typically reported only for things where we know we are always diagnosing the problematic construct (because warnings can be disabled in the cases where the analysis is incorrect).

Pushing the check into CSA helps make it clear that this is a best effort attempt to diagnose problematic code rather than going to catch all the programmer's mistakes.

fveselov · 2021-04-07T07:46:59Z

FYI: PR#219 disabling ESIMD recursion tests is merged.

AaronBallman

Changes LGTM!

bader · 2021-04-08T12:20:17Z

/summary:run

zjin-lcf · 2022-03-08T15:31:35Z

Can you please explain that SYCL kernel cannot call a recursive function ?
I think CUDA and HIP support recursion.
Thanks

keryell · 2022-03-08T22:57:58Z

Can you please explain that SYCL kernel cannot call a recursive function ? I think CUDA and HIP support recursion. Thanks

@zjin-lcf It is forbidden in the SYCL standard to align to the minimal feature set of the accelerators envisioned by the SYCL standard.
That said, this does not prevent an implementation to support it if the target supports it. It sounds also like a good optional feature for SYCL Next. :-)
I am curious to know how do GPU handle non uniform stack allocation if for example your recursion call different functions depending on the work-item id.

zjin-lcf · 2022-03-09T00:39:44Z

I assume that recursion is not useful for most application developers porting a CPU program; otherwise, the feature would be included in SYCL Before. I will try to convert recursion to iteration.

Report Compiling error for recursion in SYCL kernel

f3e005d

Signed-off-by: gejin <ge.jin@intel.com>

jinge90 requested review from elizabethandrews and premanandrao as code owners March 22, 2021 08:02

jinge90 requested a review from bader March 22, 2021 08:02

Merge remote-tracking branch 'origin/sycl' into compiling_error_for_r…

4968d31

…ecursive

bader requested a review from AaronBallman March 23, 2021 16:26

AaronBallman previously approved these changes Mar 23, 2021

View reviewed changes

elizabethandrews previously approved these changes Mar 23, 2021

View reviewed changes

keryell reviewed Mar 26, 2021

View reviewed changes

jinge90 added 3 commits March 31, 2021 15:27

Support Recursion in ESIMD kernel

1d53420

Signed-off-by: gejin <ge.jin@intel.com>

clang format

68c5111

Signed-off-by: gejin <ge.jin@intel.com>

Merge remote-tracking branch 'origin/sycl' into compiling_error_for_r…

971f2e2

…ecursive

jinge90 dismissed stale reviews from elizabethandrews and AaronBallman via 971f2e2 April 1, 2021 06:32

AaronBallman reviewed Apr 1, 2021

View reviewed changes

AaronBallman requested a review from erichkeane April 1, 2021 11:49

erichkeane reviewed Apr 1, 2021

View reviewed changes

Merge remote-tracking branch 'origin/sycl' into compiling_error_for_r…

9b17f65

…ecursive

jinge90 requested review from AaronBallman, erichkeane and elizabethandrews April 2, 2021 01:12

revert esimd kernel recursion support

4a3eb37

Signed-off-by: gejin <ge.jin@intel.com>

erichkeane approved these changes Apr 2, 2021

View reviewed changes

AaronBallman reviewed Apr 2, 2021

View reviewed changes

bader requested a review from AaronBallman April 7, 2021 10:43

AaronBallman approved these changes Apr 7, 2021

View reviewed changes

bader merged commit d822eda into intel:sycl Apr 8, 2021

		@@ -24,7 +24,7 @@ void *operator new(size_t);
		void usage2(myFuncDef functionPtr) {

	if (FunctionDecl *Callee = dyn_cast<FunctionDecl>(CGN->getDecl())) {
	if (auto *Callee = dyn_cast<FunctionDecl>(CGN->getDecl())) {

	if (auto SYCLKernelDec = dyn_cast<FunctionDecl>(D)) {
	if (auto *SYCLKernelDec = dyn_cast<FunctionDecl>(D)) {

[SYCL] Report Compiling error for recursion in SYCL kernel #3390

[SYCL] Report Compiling error for recursion in SYCL kernel #3390

Conversation

jinge90 commented Mar 22, 2021 • edited by bader Loading

jinge90 commented Mar 22, 2021

bader commented Mar 23, 2021

AaronBallman left a comment

Choose a reason for hiding this comment

bader commented Mar 25, 2021

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

jinge90 commented Mar 26, 2021

bader commented Mar 26, 2021

jinge90 commented Apr 1, 2021

jinge90 commented Apr 1, 2021

AaronBallman left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

AaronBallman commented Apr 1, 2021

erichkeane left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

erichkeane left a comment

Choose a reason for hiding this comment

AaronBallman commented Apr 1, 2021

jinge90 commented Apr 1, 2021

DenisBakhvalov commented Apr 1, 2021

kbobrovs commented Apr 1, 2021

erichkeane commented Apr 1, 2021

DenisBakhvalov commented Apr 1, 2021

kbobrovs commented Apr 1, 2021

vladimirlaz commented Apr 1, 2021

kbobrovs commented Apr 1, 2021

jinge90 commented Apr 2, 2021

erichkeane commented Apr 2, 2021

erichkeane left a comment

Choose a reason for hiding this comment

jinge90 commented Apr 2, 2021

erichkeane commented Apr 2, 2021

AaronBallman left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

AaronBallman commented Apr 2, 2021

erichkeane commented Apr 2, 2021

AaronBallman commented Apr 2, 2021

fveselov commented Apr 7, 2021

AaronBallman left a comment

Choose a reason for hiding this comment

bader commented Apr 8, 2021

zjin-lcf commented Mar 8, 2022 • edited Loading

keryell commented Mar 8, 2022

zjin-lcf commented Mar 9, 2022

jinge90 commented Mar 22, 2021 •

edited by bader

Loading

zjin-lcf commented Mar 8, 2022 •

edited

Loading