[SYCL][ESIMD]Limit propagation of ESIMD attributes #7191

fineg74 · 2022-10-26T18:45:43Z

No description provided.

fineg74 · 2022-10-26T18:47:05Z

Complementary test PR: intel/llvm-test-suite#1352

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

kbobrovs · 2022-11-01T17:11:33Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

 #define DEBUG_TYPE "LowerESIMDKernelAttrs"

 using namespace llvm;

+namespace {


This allocator is used in other places - please factor this out to a separate source to avoid code duplication.

kbobrovs · 2022-11-01T17:20:22Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

@@ -30,7 +63,28 @@ SYCLFixupESIMDKernelWrapperMDPass::run(Module &M, ModuleAnalysisManager &MAM) {
      sycl::utils::traverseCallgraphUp(
          &F,
          [&](Function *GraphNode) {
-            if (!llvm::esimd::isESIMD(*GraphNode)) {
+            if (!llvm::esimd::isESIMD(*GraphNode) &&
+                llvm::esimd::isKernel(*GraphNode)) {


This filtering leaves out SYCL_EXTERNAL functions, which does not seem correct.

kbobrovs · 2022-11-01T17:25:05Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

+              if (!Name.contains("sycl::_V1::detail::RoundedRangeKernel<")) {
+                return;
+              }
+


looks like the GraphNode->setMetadata... call below is reached only when GraphNode is a RoundedRangeKernel, which means only RoundedRangeKernel gets the marker, all others are left out. Previous logic was that all functions were marked. This needs additional investigation. I'll also give this a thought.

Originally this pass was introduced to resolve a problem of mixing ESIMD and NonESIMD kernels and back then the issue was determined as
The immediate cause of the issue is that ESIMD kernel is wrapped with RoundedRangeKernel and sycl_explicit_esimd attribute is not propagated to the wrapper class. As a result ModuleSplitter in sycl-post-link creates a spurious split for the wrapper and later creates a spurious non-esimd module and pulls called functions into it. As a result non-esimd module would contain esimd functions and later fails to link. It looks like that for named ESIMD kernels the attributes are properly propagated and therefore the issue doesn't occur. (#6557)
This bug basically says that marking all functions that call ESIMD functions is also incorrect as it results in a crash in some circumstances. Therefore the fix here is to limit marking to RoundedRangeKernel which was the original source of the problem. So in this sense the fix does what it is expected to do.

What this patch does is it excludes all the functions in the caller chain from given ESIMD function F up to RoundedRangeKernel from marking them as ESIMD, instead, it only marks the kernel. Previously all the function in the call chain were marked.
Now, this patch is supposed to fix internal LLVM error on some function met in the IR:
Assertion `isa(Val) && "cast() argument of incompatible type!"' failed.

So

All the functions in a call chain ending at ESIMD function must be ESIMD functions, so not marking them is incorrect.

I think the patch just masks the problem by skipping marking the problematic function (which in turn can spawn other problems).

I think we need to come up with a patch that fixes the root cause of the above assert.

It needs more review/investigation.

kbobrovs · 2022-11-03T04:35:38Z

llvm/include/llvm/SYCLLowerIR/ESIMD/Allocator.h

@@ -0,0 +1,53 @@
+//===--------- Allocator.h - Allocator for demangler


I should've expressed my proposal clearer, sorry. This is just implementation detail of demangling code, not a general purpose service, so I would very much prefer this to reside in ESIMDUtils.cpp/ESIMDUtils.hpp

kbobrovs · 2022-11-10T19:24:56Z

llvm/lib/SYCLLowerIR/SYCLUtils.cpp

+
+using namespace llvm::esimd;
+bool isAddressArgumentInvokeSIMD(const CallInst *CI) {
+  constexpr char INVOKE_SIMD_PREF[] = "_Z33__regcall3____builtin_invoke_simd";


please factor out. There is another user:
llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp:constexpr char INVOKE_SIMD_PREF[] = "_Z33__regcall3____builtin_invoke_simd";

kbobrovs · 2022-11-10T19:26:17Z

llvm/lib/SYCLLowerIR/SYCLUtils.cpp


 namespace llvm {
 namespace sycl {
 namespace utils {
+
+using namespace llvm::esimd;
+bool isAddressArgumentInvokeSIMD(const CallInst *CI) {


Suggested change

bool isAddressArgumentInvokeSIMD(const CallInst *CI) {

bool isInvokeSimdBuiltinCall(const CallInst *CI) {

kbobrovs · 2022-11-10T19:26:41Z

llvm/lib/SYCLLowerIR/SYCLUtils.cpp

+  return false;
+}
+
+bool filterFunctionPointer(Value *address) {


please add comment what this function does, what it returns.
I also thought the deduceFunction from LowerInvokeSimd.cpp could be used directly (factored out first, of course).

My understanding is that deduceFunction is doing the opposite: i.e. given an argument of invoke_simd, it tries to find the actual function that invoke_simd tries to call.

kbobrovs · 2022-11-10T19:37:18Z

llvm/lib/SYCLLowerIR/SYCLUtils.cpp

@@ -43,6 +101,13 @@ void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction ActionF,
        } else {
          // ... non-call is OK - add using function to the worklist
          if (auto *I = dyn_cast<Instruction>(FCall)) {
+            if (auto *SI = dyn_cast<StoreInst>(I)) {


as we talked, such filtering within traverseCallgraphUp should be opaque, and filtering function should be passed as parameter rather than hardcoded, as there are multiple users of traverseCallgraphUp, but we need this invoke_simd behavior only for attribute markup. Somthing like

void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction ActionF, SmallPtrSetImpl<Function *> &FunctionsVisited, bool ErrorOnNonCallUse, std::function<bool(const Instruction *I, const Function *F)> && FuncUseFilter);

and the LowerESIMDKernelAttrs pass would invoke the traverseCallgraphUp with proper FuncUseFilter argument.

Done. However, I believe we are going to encounter similar problems in other cases where traverseCallgraphUp is used with invoke_simd and therefore embedding the filtering functionality to traverseCallgraphUp would make more sense

kbobrovs · 2022-11-22T03:46:23Z

llvm/include/llvm/SYCLLowerIR/SYCLUtils.h

-void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction NodeF,
-                         SmallPtrSetImpl<Function *> &Visited,
-                         bool ErrorOnNonCallUse);
+void traverseCallgraphUp(


Please add a comment how functionFilter parameter is used.

kbobrovs · 2022-11-22T03:48:34Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

+#include "llvm/Demangle/Demangle.h"
+#include "llvm/Demangle/ItaniumDemangle.h"
+#include "llvm/IR/Module.h"
+#include "llvm/Pass.h"


are these still needed?

kbobrovs · 2022-11-22T04:02:14Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

 #define DEBUG_TYPE "LowerESIMDKernelAttrs"

 using namespace llvm;
+using namespace llvm::esimd;


Suggested change

using namespace llvm::esimd;

kbobrovs · 2022-11-22T04:02:24Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

@@ -37,7 +39,7 @@ SYCLFixupESIMDKernelWrapperMDPass::run(Module &M, ModuleAnalysisManager &MAM) {
              Modified = true;
            }
          },
-          false);
+          false, filterFunctionPointer);


Suggested change

false, filterFunctionPointer);

false, llvm::esimd::filterFunctionPointer);

kbobrovs · 2022-11-22T04:04:49Z

llvm/lib/SYCLLowerIR/SYCLUtils.cpp

@@ -43,6 +47,13 @@ void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction ActionF,
        } else {
          // ... non-call is OK - add using function to the worklist
          if (auto *I = dyn_cast<Instruction>(FCall)) {
+            if (auto *SI = dyn_cast<StoreInst>(I)) {


please add comment why this special branch is needed if (auto *SI = dyn_cast<StoreInst>(I))

kbobrovs · 2022-11-22T04:13:09Z

llvm/lib/SYCLLowerIR/SYCLUtils.cpp

@@ -43,6 +47,13 @@ void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction ActionF,
        } else {
          // ... non-call is OK - add using function to the worklist
          if (auto *I = dyn_cast<Instruction>(FCall)) {
+            if (auto *SI = dyn_cast<StoreInst>(I)) {


The StoreInst check should be moved into the filter as it has nothing to do with this generic traversal functionality. That's why I suggested std::function<bool(const Instruction *I, const Function *F)>, not std::function<bool(const Value *)>

Suggested change

if (auto *SI = dyn_cast<StoreInst>(I)) {

if (!functionFilter(I, CurF)) {

continue;

}

kbobrovs · 2022-11-22T04:14:40Z

llvm/lib/SYCLLowerIR/SYCLUtils.cpp

+              Value *addr = SI->getPointerOperand();
+              if (!functionFilter(addr)) {
+                continue;
+              }
+            }
+


Suggested change

Value *addr = SI->getPointerOperand();

if (!functionFilter(addr)) {

continue;

}

}

kbobrovs · 2022-11-22T04:18:53Z

llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp

+// Tracks the use of a function pointer being stored in a memory.
+// Returns false if the function pointer is used as an argument for invoke_simd
+// function call, true otherwise.
+bool filterFunctionPointer(Value *address) {


Since this function remains quite specific and is not reused anywhere, it should be moved into the place of use - LowerESIMDKernelAttrs.cpp.

Please improve comment. It is not clear what is memory, what is function pointer, what is the incoming argument.

I added a note below that this should have different signature.

AlexeySachkov · 2022-11-23T09:03:53Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

+  if (address == nullptr) {
+    return true;
+  }


Note: LLVM Coding Standard says don’t use braces on simple single-statement bodies of if/else/loop statements

I always thought this is not strict - we prefer to omit the braces to avoid unnecessary line noise.

AlexeySachkov · 2022-11-23T09:04:22Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

+  if (F && F->getName().startswith(esimd::INVOKE_SIMD_PREF)) {
+    return true;
+  }
+  return false;


Suggested change

if (F && F->getName().startswith(esimd::INVOKE_SIMD_PREF)) {

return true;

}

return false;

return F && F->getName().startswith(esimd::INVOKE_SIMD_PREF);

AlexeySachkov · 2022-11-23T09:06:05Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

+  // if the instruction is to store address of a function, check if it is later
+  // used by InvokeSimd.
+  if (auto *SI = dyn_cast<StoreInst>(I)) {
+    const Value *addr = SI->getPointerOperand();


Suggested change

const Value *addr = SI->getPointerOperand();

const Value *Addr = SI->getPointerOperand();

AlexeySachkov · 2022-11-23T09:06:48Z

llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp

@@ -73,10 +72,6 @@ ModulePass *llvm::createSYCLLowerInvokeSimdPass() {

 namespace {
 // TODO support lambda and functor overloads


Do we need to move this TODO as well or is it unrelated to the removed code below?

TODO is unrelated to the removed code

kbobrovs · 2022-11-28T22:32:36Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

+// Checks the use of a function address being stored in a memory.
+// Returns false if the function address is used as an argument for
+// invoke_simd function call, true otherwise.
+bool checkFunctionAddressUse(const Value *address) {


comment still not updated

BTW, why not just inline it into filterInvokeSimdUse, since it is not general purpose function anyway?

It uses recursion to follow the use of a function. filterInvokeSimdUse was introduced to take StoreInst check out of generic traverseCallgraphUp and inlining will complicate the recursion.

kbobrovs · 2022-11-28T22:35:36Z

llvm/lib/SYCLLowerIR/SYCLUtils.cpp

@@ -43,6 +47,10 @@ void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction ActionF,
        } else {
          // ... non-call is OK - add using function to the worklist
          if (auto *I = dyn_cast<Instruction>(FCall)) {
+            if (!functionFilter(I)) {


Function* should be part of the interface, as we are deciding whether a use of particular function should be filtered or not. What if there are multiple function pointers used by I?

kbobrovs · 2022-11-29T19:45:52Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

@@ -25,8 +25,8 @@ bool isInvokeSimdBuiltinCall(const CallInst *CI) {
  return F && F->getName().startswith(esimd::INVOKE_SIMD_PREF);
 }

-// Checks the use of a function address being stored in a memory.
-// Returns false if the function address is used as an argument for
+// Checks the use of a function.


There is no "function" per se in any of the parameters. So the comment is not clear.

llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h

Co-authored-by: kbobrovs <[email protected]>

llvm/include/llvm/SYCLLowerIR/SYCLUtils.h

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

kbobrovs · 2022-11-30T02:29:56Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

+  // used by InvokeSimd.
+  if (auto *SI = dyn_cast<StoreInst>(I)) {
+    const Value *Addr = SI->getPointerOperand();
+    return checkFunctionAddressUse(Addr);


Suggested change

return checkFunctionAddressUse(Addr);

return !isInvokeSimdArgumentSource(Addr, F);

kbobrovs · 2022-11-30T02:47:01Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

+// location is used diectly or indirectly as an argument for invoke_simd.
+// Returns false if the function is used as an argument for
+// invoke_simd function call, true otherwise.
+bool checkFunctionAddressUse(const Value *address) {


Suggested change

bool checkFunctionAddressUse(const Value *address) {

bool isInvokeSimdArgumentSource(const Value *address) {

kbobrovs · 2022-11-30T02:47:16Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

+// invoke_simd function call, true otherwise.
+bool checkFunctionAddressUse(const Value *address) {
+  if (address == nullptr)
+    return true;


Suggested change

return true;

return false;

kbobrovs · 2022-11-30T02:48:43Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

+    } else if (const auto *CI = dyn_cast<CallInst>(V)) {
+      // if __builtin_invoke_simd uses the pointer, do not traverse the function
+      if (isInvokeSimdBuiltinCall(CI))
+        return false;


Suggested change

return false;

return true;

kbobrovs · 2022-11-30T03:25:52Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

+// location is used diectly or indirectly as an argument for invoke_simd.
+// Returns false if the function is used as an argument for
+// invoke_simd function call, true otherwise.
+bool checkFunctionAddressUse(const Value *address) {


Looking into the code deeper, I think this part should be reworked. The overall approach can be:

Loads = find all loads from `address` for (Load : Loads) { Uses = collectUsesSkipThroughCasts(Load); for (Use : Uses) { if (Use is parameter to invoke_simd) { return true; } } return false; }

Current logic is very different. We also don't need recursion, so this can be inlined intp the caller.

I believe the suggested rework is a simplification of the current approach and we are going to miss even more cases that we may currently miss.
For example what if we store the value to a different address, we are not going to track this new address and are going to miss the valid code flow that may lead to invoke_simd. The purpose of recursion was to track such chain of load or load/store manipulations especially for -O0 when no optimizations that eliminate redundant loads/stores are applied.
For example following code will fail:

FuncType SIMD_CALLEE_PTR = SIMD_CALLEE; .... FuncType SIMD_CALLEE_PTR1 = SIMD_CALLEE_PTR; .... invoke_simd(sg, SIMD_CALLEE_PTR1 , ...);

I think original checkFunctionAddressUse is borrowed from https://github.com/intel/llvm/blob/sycl/llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp#L102, whose purpose is quite different from what we need here. collectPossibleStoredVals is used to conservatively detect if optimization is possible, checkFunctionAddressUse must tell with 100% reliability that given address:

either contains only given function address or cannot contain it at all, anything in between should be an assert (as we lose reliability in this case)

if it contains only given function address, then the value stored (the function pointer) either can't appear as an argument of __builtin_invoke_simd.
In general case this is not possible, e.g. when user program uses some complex data flow on the function pointer before calling invoke_simd.

The more I think of this, the more I'm getting inclined to stop adding a call graph edge on any use of a function other than call.
In our case, if a user function can call ESIMD function via a pointer, we should require that this function is marked with SYCL_ESIMD_FUNCTION explicitly.

Not sure I follow the conclusion. The problem that this fix is trying to fix is to mark all callers of an ESIMD function up the chain with ESIMD until it reaches either kernel or invoke_simd. So it is invoked only for functions explicitly marked as SYCL_ESIMD_FUNCTION. If we allow function pointers as arguments for invoke_simd we would need some data flow analysis, similar to what checkFunctionAddressUse does. Yes it has its limitations and doesn't cover all the cases.

As we talked, the idea is:

When propagating sycl_explicit_simd attribute we never go up through non-call use of function pointers, which automatically covers invoke_simd

Users must explicitly add sycl_explicit_simd attribute if function calls a sycl_explicit_simd function through pointer.

Co-authored-by: kbobrovs <[email protected]>

llvm/include/llvm/SYCLLowerIR/SYCLUtils.h

Co-authored-by: kbobrovs <[email protected]>

…esimdSplitFix

pvchupin · 2022-12-01T02:56:32Z

hip fails are due to driver fail.

[16528323.699738] amdgpu: Runlist is getting oversubscribed. Expect reduced ROCm performance.
[16528323.892286] amdgpu: Runlist is getting oversubscribed. Expect reduced ROCm performance.
[16528325.741668] amdgpu: Runlist is getting oversubscribed. Expect reduced ROCm performance.
[16528327.333944] amdgpu: Runlist is getting oversubscribed. Expect reduced ROCm performance.
[16528720.285095] static-buffer-d[1433031]: segfault at 39 ip 00007fe85cea9871 sp 00007fffd57bed80 error 4 in libamdhip64.so.5.4.50400[7fe85cdf0000+37e000]

I've rebooted runner.

kbobrovs · 2022-12-01T22:49:59Z

llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h

@@ -68,5 +74,35 @@ void collectUsesLookThroughCastsAndZeroGEPs(const Value *V,
 /// in it. Returns nullptr if failed to do so.
 Type *getVectorTyOrNull(StructType *STy);

+// Simplest possible implementation of an allocator for the Itanium demangler


I think this change (and other ones related to demangling) is no longer needed

Upon second thought, it is good to have the allocator defined in shared header, as others may need to use it too.

kbobrovs · 2022-12-05T08:39:55Z

@AlexeySachkov, do you approve?

AlexeySachkov · 2022-12-05T08:46:25Z

@AlexeySachkov, do you approve?

My comments were mostly minor about coding style, I don't have objections against merging this PR, feel free to go ahead

AlexeySachkov

Formal approval (I forgot that there are two files owned by dpcpp-tools-reviewers)

Limit propagation of ESIMD attributes to RoundedRangeKernel

5767c50

fineg74 requested a review from a team as a code owner October 26, 2022 18:45

fineg74 mentioned this pull request Oct 26, 2022

[SYCL][ESIMD]Add test to validate correct split intel/llvm-test-suite#1352

Merged

v-klochkov reviewed Oct 27, 2022

View reviewed changes

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp Outdated Show resolved Hide resolved

fineg74 changed the title ~~Limit propagation of ESIMD attributes to RoundedRangeKernel~~ [SYCL][ESIMD]Limit propagation of ESIMD attributes to RoundedRangeKernel Oct 27, 2022

v-klochkov previously approved these changes Oct 31, 2022

View reviewed changes

kbobrovs reviewed Nov 1, 2022

View reviewed changes

fineg74 added 3 commits November 1, 2022 14:23

Merge remote-tracking branch 'origin/sycl' into esimdSplitFix

56854ed

Refactorcode to reduce code duplication

38c1541

Fix compilation issue on RHEL

a0a9c53

kbobrovs reviewed Nov 3, 2022

View reviewed changes

fineg74 and others added 3 commits November 2, 2022 23:29

Move allocator to a utils file

dedf2ed

Merge remote-tracking branch 'origin/sycl' into esimdSplitFix

497aa64

Prevent traverseCallgraphUp from traversing beyond invoke_simd call

a1aadf5

fineg74 requested a review from a team as a code owner November 10, 2022 18:35

kbobrovs reviewed Nov 10, 2022

View reviewed changes

fineg74 changed the title ~~[SYCL][ESIMD]Limit propagation of ESIMD attributes to RoundedRangeKernel~~ [SYCL][ESIMD]Limit propagation of ESIMD attributes Nov 10, 2022

Address PR comments

8f089a6

fineg74 requested a review from rolandschulz as a code owner November 10, 2022 20:18

Address PR comments

9d33f38

kbobrovs reviewed Nov 22, 2022

View reviewed changes

Address PR comments

209532f

AlexeySachkov reviewed Nov 23, 2022

View reviewed changes

Address PR comments

992fcda

kbobrovs reviewed Nov 28, 2022

View reviewed changes

fineg74 added 2 commits November 29, 2022 09:38

Address PR comments

04737a2

Address PR comments

cf0b167

kbobrovs reviewed Nov 29, 2022

View reviewed changes

Address PR comments

9b59931

kbobrovs reviewed Nov 29, 2022

View reviewed changes

llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h Outdated Show resolved Hide resolved

fineg74 and others added 2 commits November 29, 2022 15:56

Update llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h

ac0422b

Co-authored-by: kbobrovs <[email protected]>

Fix clang-format issue

0ae282d

kbobrovs reviewed Nov 30, 2022

View reviewed changes

fineg74 and others added 4 commits November 29, 2022 20:09

Update llvm/include/llvm/SYCLLowerIR/SYCLUtils.h

87d0f46

Co-authored-by: kbobrovs <[email protected]>

Update llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

a073628

Co-authored-by: kbobrovs <[email protected]>

Update llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

f1cad2f

Co-authored-by: kbobrovs <[email protected]>

Fix clang-format issue.

53fd818

kbobrovs reviewed Nov 30, 2022

View reviewed changes

llvm/include/llvm/SYCLLowerIR/SYCLUtils.h Outdated Show resolved Hide resolved

fineg74 and others added 4 commits November 30, 2022 11:36

Update llvm/include/llvm/SYCLLowerIR/SYCLUtils.h

5a4eae3

Co-authored-by: kbobrovs <[email protected]>

Fix handling of StoreInst

70ae7a0

Merge branch 'esimdSplitFix' of https://github.com/fineg74/llvm into …

2618034

…esimdSplitFix

Address PR comments

a7c301d

kbobrovs reviewed Dec 1, 2022

View reviewed changes

kbobrovs approved these changes Dec 1, 2022

View reviewed changes

v-klochkov approved these changes Dec 1, 2022

View reviewed changes

rolandschulz approved these changes Dec 2, 2022

View reviewed changes

AlexeySachkov approved these changes Dec 5, 2022

View reviewed changes

kbobrovs merged commit 4d3c150 into intel:sycl Dec 5, 2022

fineg74 deleted the esimdSplitFix branch December 27, 2022 01:46

		@@ -0,0 +1,53 @@
		//===--------- Allocator.h - Allocator for demangler

	bool isAddressArgumentInvokeSIMD(const CallInst *CI) {
	bool isInvokeSimdBuiltinCall(const CallInst *CI) {

	false, filterFunctionPointer);
	false, llvm::esimd::filterFunctionPointer);

	const Value *addr = SI->getPointerOperand();
	const Value *Addr = SI->getPointerOperand();

		@@ -73,10 +72,6 @@ ModulePass *llvm::createSYCLLowerInvokeSimdPass() {

		namespace {
		// TODO support lambda and functor overloads

	return checkFunctionAddressUse(Addr);
	return !isInvokeSimdArgumentSource(Addr, F);

	bool checkFunctionAddressUse(const Value *address) {
	bool isInvokeSimdArgumentSource(const Value *address) {

[SYCL][ESIMD]Limit propagation of ESIMD attributes #7191

[SYCL][ESIMD]Limit propagation of ESIMD attributes #7191

Uh oh!

Conversation

fineg74 commented Oct 26, 2022

Uh oh!

fineg74 commented Oct 26, 2022

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kbobrovs Nov 10, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kbobrovs Nov 28, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

kbobrovs Nov 10, 2022 •

edited

Loading

kbobrovs Nov 28, 2022 •

edited

Loading

kbobrovs Nov 30, 2022 •

edited

Loading

fineg74 Nov 30, 2022 •

edited

Loading

kbobrovs Dec 1, 2022 •

edited

Loading