[ESIMD] Infer address space of pointer that are passed through invoke_simd to ESIMD API to generate better code on BE #14528

fineg74 · 2024-07-11T05:20:03Z

No description provided.

sarnex

Nice job optimizing this case! First pass comments below

sarnex · 2024-07-15T14:34:02Z

llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp

@@ -75,6 +76,8 @@ ModulePass *llvm::createSYCLLowerInvokeSimdPass() {
 namespace {
 // TODO support lambda and functor overloads

+ValueMap<const Value *, SmallDenseMap<uint32_t, uint32_t>> ArgAddrSpaceMap;


Maybe we could alias the map value to ArgIdxToAddrSpaceMap or something like that?

Refactored it out

sarnex · 2024-07-15T14:36:07Z

llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp

+      if (Callee) {
+        if (Callee->isDeclaration())
+          continue;
+      }


Maybe this could be

Suggested change

if (Callee) {

if (Callee->isDeclaration())

continue;

}

if (Callee && Callee->isDeclaration())

continue;

sarnex · 2024-07-15T14:39:46Z

llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp

+        ArgumentMap[i - 2] = AddressSpace;
+      }
+    }
+    if (!ArgumentMap.empty()) {


Nit: We can remove braces for one line ifs

sarnex · 2024-07-15T14:44:08Z

llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp

@@ -259,6 +262,49 @@ void markFunctionAsESIMD(Function *F) {
  }
 }

+void AdjustAddressSpace(Function *F, uint32_t ArgNo, uint32_t ArgAddrSpace) {


nit

Suggested change

void AdjustAddressSpace(Function *F, uint32_t ArgNo, uint32_t ArgAddrSpace) {

void adjustAddressSpace(Function *F, uint32_t ArgNo, uint32_t ArgAddrSpace) {

sarnex · 2024-07-15T14:49:09Z

llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp

@@ -75,6 +76,8 @@ ModulePass *llvm::createSYCLLowerInvokeSimdPass() {
 namespace {
 // TODO support lambda and functor overloads

+ValueMap<const Value *, SmallDenseMap<uint32_t, uint32_t>> ArgAddrSpaceMap;


Is there anything we could do to prevent this from being a global? Could we pass it down?

Refactored it out

sarnex · 2024-07-15T14:51:04Z

llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp

+  for (const CallInst *CI : ISCalls) {
+    SmallDenseMap<uint32_t, uint32_t> ArgumentMap;
+    for (uint32_t i = 2; i < CI->arg_size(); ++i) {
+      const Value *Arg = CI->getArgOperand(i);
+      if (Arg->getType()->isPointerTy()) {
+        uint32_t AddressSpace = Arg->getType()->getPointerAddressSpace();
+        if (AddressSpace == 4) {
+          const AddrSpaceCastInst *ASC = dyn_cast<AddrSpaceCastInst>(Arg);
+          if (!ASC)
+            continue;
+
+          AddressSpace =
+              ASC->getOperand(0)->getType()->getPointerAddressSpace();
+        }
+        ArgumentMap[i - 2] = AddressSpace;
+      }
+    }
+    if (!ArgumentMap.empty()) {
+      ArgAddrSpaceMap[CI->getArgOperand(1)] = ArgumentMap;
+    }
+  }


Can we move this to a separate function?

Refactored the logic a little bit so this part is no longer exist (its been merged to the other change to eliminate use of map

sarnex · 2024-07-15T14:53:36Z

llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp

@@ -436,6 +492,28 @@ PreservedAnalyses SYCLLowerInvokeSimdPass::run(Module &M,
      ISCalls.insert(CI);
    }
  }
+  for (const CallInst *CI : ISCalls) {
+    SmallDenseMap<uint32_t, uint32_t> ArgumentMap;
+    for (uint32_t i = 2; i < CI->arg_size(); ++i) {


Do you mind adding a comment saying arg idx 2 is the first real arg to invoke_simd calls? Thx

sarnex · 2024-07-15T14:55:47Z

llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp

+      if (Arg->getType()->isPointerTy()) {
+        uint32_t AddressSpace = Arg->getType()->getPointerAddressSpace();
+        if (AddressSpace == 4) {
+          const AddrSpaceCastInst *ASC = dyn_cast<AddrSpaceCastInst>(Arg);


Is it a common pattern that we lose the true addr space information in an addrspace cast to generic into the invoke_simd call?

For some reason we always convert pointers to address space 4 as the first step, but IGC is able to infer the correct address space if there any mention to the original address space. For invoke_simd we lose it between scalar and vector BEs so that is why vector BE has no way to infer the actual address space (all the callees of invoke_simd has parameters in address space 4 and looks like vector BE doesn't see what was used to call invoke_simd) Here I try to infer address space for pointer parameters for invoke_simd and so far I saw to patterns only : address space 3 pointers are passed directly while address space 1 pointers are casted to address space 4 and then passed to invoke_simd. I am not sure if this is the only pattern but it shouldn't be a problem to handle another patter if discovered

sarnex · 2024-07-15T15:13:17Z

llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp

+    } else {
+      for (unsigned int i = 0; i < ArgUse->getNumOperands(); ++i) {
+        if (ArgUse->getOperand(i) == Arg) {
+          const Type *Ty = ArgUse->getOperand(i)->getType();


It looks like we only use this variable to get the context, probably we could get rid of it and get the context from somewhere else

sarnex · 2024-07-15T15:14:25Z

sycl/test/invoke_simd/invoke_simd_address_space_inferral.cpp

+  sycl::free(A, q);
+
+  return 0;
+  // CHECK: addrspacecast ptr addrspace(4) %A to ptr addrspace(1)


Maybe we could add an E2E test using Jason's perf test infrastructure to lock down the lower instruction count for a case that generates the runtime checks in IGC today

Added a test although not sure of how reliable it is going to be due to large variation in instruction count between different driver versions

yeah lets try, if it fails all the time we can remove it. thanks

sycl/test/invoke_simd/invoke_simd_address_space_inferral.cpp

sarnex

lgtm, only minor comments, nice refactor!

sarnex · 2024-07-16T14:19:04Z

llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp

@@ -436,6 +495,7 @@ PreservedAnalyses SYCLLowerInvokeSimdPass::run(Module &M,
      ISCalls.insert(CI);
    }
  }
+


nit: remove whitespace

sarnex · 2024-07-16T14:20:06Z

sycl/test-e2e/ESIMD/PerformanceTests/invoke_simd_smoke.cpp

+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu


we probably want to restrict to a single GPU like we did for the rest of the perf tests

// REQUIRES: gpu-intel-dg2 && level_zero

invoke_simd works everywhere. Is it related to the fact it can produce different instruction counts for different platforms ?

yep, exactly

sarnex · 2024-07-16T14:21:14Z

sycl/test/invoke_simd/invoke_simd_address_space_inferral.cpp

+  sycl::free(A, q);
+
+  return 0;
+  // CHECK: addrspacecast ptr addrspace(4) %A to ptr addrspace(1)


yeah lets try, if it fails all the time we can remove it. thanks

…g74/llvm into InferAddrSpace

sarnex · 2024-07-17T17:11:11Z

Fail not related

sarnex · 2024-07-17T20:49:34Z

@fineg74 It looks like the test is failing in poscommit, can you please take a look?

https://github.com/intel/llvm/actions/runs/9978520401/job/27584493411

fineg74 · 2024-07-17T22:47:49Z

@fineg74 It looks like the test is failing in poscommit, can you please take a look?

https://github.com/intel/llvm/actions/runs/9978520401/job/27584493411

@sarnex , Do you know what is [Linux (Self build + shared libraries + no-assertions) / Build + LIT] ?
I am not able to reproduce it locally and it seem to pass on regular CI when I merge the latest sycl into a draft pr branch. It looks like it somehow use an old sycl-post-link or an old SYCLLowerIR library

sarnex · 2024-07-17T23:29:47Z

@fineg74 It builds the compiler using the same compiler version. Maybe build the compiler and then clone the same revision again in a different dir and use the first compiler to build it?

Also there is another postcommit fail here so I am going to revert this PR for now.

…h invoke_simd to ESIMD API to generate better code on BE" (#14607) Reverts #14528 Postcommit failures

sarnex · 2024-07-17T23:34:14Z

@fineg74 We can run postcommit in precommit by making a pull request to a new branch in intel/llvm you create that includes the sycl-devops-pr/ string so like sycl-devops-pr/testpostcommit or something, let's try that and see if we can find the issue

fineg74 · 2024-07-18T06:01:16Z

@sarnex the post commit PR with the fix is here: #14610 The only differences with this PR is in tests.

sarnex · 2024-07-18T14:11:02Z

Just replied to the PR there, thanks!

Initial implementation

76dd7ba

fineg74 temporarily deployed to WindowsCILock July 11, 2024 05:21 — with GitHub Actions Inactive

fineg74 temporarily deployed to WindowsCILock July 11, 2024 06:10 — with GitHub Actions Inactive

Improve robustnedd of thee inferral

fd28e2d

fineg74 temporarily deployed to WindowsCILock July 12, 2024 18:36 — with GitHub Actions Inactive

fineg74 had a problem deploying to WindowsCILock July 12, 2024 20:47 — with GitHub Actions Error

Add test

a246761

fineg74 marked this pull request as ready for review July 12, 2024 21:48

fineg74 requested a review from a team as a code owner July 12, 2024 21:48

fineg74 temporarily deployed to WindowsCILock July 12, 2024 21:49 — with GitHub Actions Inactive

fineg74 temporarily deployed to WindowsCILock July 12, 2024 23:12 — with GitHub Actions Inactive

Fix handling of call instruction

6dcd8c7

fineg74 had a problem deploying to WindowsCILock July 14, 2024 23:11 — with GitHub Actions Failure

fineg74 temporarily deployed to WindowsCILock July 14, 2024 23:55 — with GitHub Actions Inactive

Readd a logic that was mistakenly removed

5cf7481

fineg74 had a problem deploying to WindowsCILock July 15, 2024 01:52 — with GitHub Actions Error

Optimize the test

238250a

fineg74 temporarily deployed to WindowsCILock July 15, 2024 02:14 — with GitHub Actions Inactive

fineg74 temporarily deployed to WindowsCILock July 15, 2024 02:57 — with GitHub Actions Inactive

Fix test

6c0ab22

fineg74 temporarily deployed to WindowsCILock July 15, 2024 05:11 — with GitHub Actions Inactive

fineg74 temporarily deployed to WindowsCILock July 15, 2024 05:55 — with GitHub Actions Inactive

sarnex reviewed Jul 15, 2024

View reviewed changes

v-klochkov reviewed Jul 15, 2024

View reviewed changes

sycl/test/invoke_simd/invoke_simd_address_space_inferral.cpp Outdated Show resolved Hide resolved

Address PR comments

552b25d

fineg74 temporarily deployed to WindowsCILock July 16, 2024 00:18 — with GitHub Actions Inactive

fineg74 had a problem deploying to WindowsCILock July 16, 2024 01:01 — with GitHub Actions Error

Adjust instruction count

8088c3e

fineg74 temporarily deployed to WindowsCILock July 16, 2024 01:35 — with GitHub Actions Inactive

fineg74 temporarily deployed to WindowsCILock July 16, 2024 02:18 — with GitHub Actions Inactive

sarnex reviewed Jul 16, 2024

View reviewed changes

Address PR comments

e70b66f

fineg74 temporarily deployed to WindowsCILock July 16, 2024 17:20 — with GitHub Actions Inactive

sarnex approved these changes Jul 16, 2024

View reviewed changes

fineg74 temporarily deployed to WindowsCILock July 16, 2024 19:22 — with GitHub Actions Inactive

Fix test failure

e25108f

fineg74 temporarily deployed to WindowsCILock July 16, 2024 23:06 — with GitHub Actions Inactive

fineg74 temporarily deployed to WindowsCILock July 17, 2024 00:16 — with GitHub Actions Inactive

fineg74 added 2 commits July 16, 2024 23:38

Merge branches 'sycl' and 'InferAddrSpace' of https://github.com/fine…

8027d84

…g74/llvm into InferAddrSpace

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

90b8bee

fineg74 temporarily deployed to WindowsCILock July 17, 2024 06:39 — with GitHub Actions Inactive

fineg74 temporarily deployed to WindowsCILock July 17, 2024 07:30 — with GitHub Actions Inactive

sarnex merged commit 16e39df into intel:sycl Jul 17, 2024
13 of 14 checks passed

fineg74 deleted the InferAddrSpace branch July 17, 2024 17:13

sarnex mentioned this pull request Jul 17, 2024

Revert "[ESIMD] Infer address space of pointer that are passed through invoke_simd to ESIMD API to generate better code on BE" #14607

Merged

sarnex added a commit that referenced this pull request Jul 17, 2024

Revert "[ESIMD] Infer address space of pointer that are passed throug…

accd0b5

…h invoke_simd to ESIMD API to generate better code on BE" (#14607) Reverts #14528 Postcommit failures

fineg74 restored the InferAddrSpace branch July 18, 2024 00:59

fineg74 deleted the InferAddrSpace branch July 18, 2024 02:17

fineg74 mentioned this pull request Jul 18, 2024

[ESIMD] Infer address space of pointer that are passed through invoke_simd to ESIMD API to generate better code on BE #14610

Closed

fineg74 restored the InferAddrSpace branch July 18, 2024 14:41

	void AdjustAddressSpace(Function *F, uint32_t ArgNo, uint32_t ArgAddrSpace) {
	void adjustAddressSpace(Function *F, uint32_t ArgNo, uint32_t ArgAddrSpace) {

[ESIMD] Infer address space of pointer that are passed through invoke_simd to ESIMD API to generate better code on BE #14528

[ESIMD] Infer address space of pointer that are passed through invoke_simd to ESIMD API to generate better code on BE #14528

Conversation

fineg74 commented Jul 11, 2024 • edited Loading

sarnex 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

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

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

sarnex 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

sarnex commented Jul 17, 2024

sarnex commented Jul 17, 2024

fineg74 commented Jul 17, 2024

sarnex commented Jul 17, 2024 • edited Loading

sarnex commented Jul 17, 2024 • edited Loading

fineg74 commented Jul 18, 2024

sarnex commented Jul 18, 2024

fineg74 commented Jul 11, 2024 •

edited

Loading

sarnex commented Jul 17, 2024 •

edited

Loading

sarnex commented Jul 17, 2024 •

edited

Loading