Enable end to end test with Linalg input and Vulkan backend #415

silee2 · 2022-10-26T01:00:56Z

This PR enables the first end to end test targeting GPU.
It has multiple related sub parts.

Extend insert-gpu-allocs, set-spirv-capabilities and set-spirv-abi-attrs to support both OpenCL and Vulkan SPIRV
Change how insert-gpu-allocs, set-spirv-capabilities and set-spirv-abi-attrs passes gets registered.
Add test support on Windows
Extend imex-runner.py - Support comments in pass pipeline file, add options to print IR before and after passes, add support for mlir-vulkan-runner in addition to mlir-cpu-runner
Other compile issues uncovered by building on Windows.
Add end to end test case.
Change some defaults for gpu to spirv conversion pass. Now spirv memory space mapping is enabled and targets OpenCL.
OpenCL pipeline is not completely end to end but includes all passes implemented so far. Upcoming passes are also added to pass pipeline but commented out for now.

Please review these guidelines to help with the review process:

Have you provided a meaningful PR description?
Have you added a test, a reproducer, or a reference to an issue with a reproducer?
Have you tested your changes locally for CPU and GPU devices?
Have you made sure that new changes do not introduce compiler warnings?
If this PR is a work in progress, are you filing the PR as a draft?
Have you organized your commits logically and ensured each can be built by itself?

silee2 · 2022-10-26T21:33:21Z

I'll be taking PTO on Thur 10/17, so response would be delayed.

Now memref's integer space id is converted to spirv opencl space by default. Fix bug in imex-runner.py. MLIR's textual pass pipeline is sensitive to space for mulit-option passes. space is used as a separator. Add gpu to spirv pass and gpu to gpux pass to OpenCL pass pipeline.

…pu-to-spirv. Extend OpenCL pipeline to spirv serialization.

silee2 · 2022-10-26T23:52:58Z

CI seems to have some LLVM caching issue after LLVM SHA update. CI throws error that some files are not found.

Add remaining part of pipeline as a comment (disabled for now).

Works for both separate llvm install and llvm external projects build.

fschlimb

Nice!
A few more comments in the code would be helpful.

fschlimb · 2022-10-27T09:28:38Z

@diptorupd some changes to the cmake machinery, maybe this also touches #289

fschlimb · 2022-10-27T09:30:16Z

6. Add end to end test case.

end-to-end would start with PTensor :)

drprajap · 2022-10-28T19:47:19Z

lib/Transforms/SetSPIRVAbiAttribute.cpp

-      if (!mlir::gpu::GPUDialect::isKernel(gpuFunc) ||
-          gpuFunc->getAttr(attrName))
-        continue;
+    if (m_clientAPI == "opencl") {


it seems only getEntryPointABIAttr argument differs between opencl and vulkan, can this code be restructured to avoid code duplication?

Dewei-Wang-sh

lgtm overall

imex-runner.py.in is not a valid python file.

available.

Hardcode84 · 2022-10-31T22:23:21Z

lib/Conversion/PTensorToLinalg/PTensorToLinalg.cpp

@@ -270,18 +270,17 @@ static BodyType buildTrivial(::mlir::Type typ) {
              builder.create<IOP>(loc, lhs, rhs).getResult());
        return;
      } else
-        assert("Found integer type but binary op not defined for integers" ==
-               nullptr);
+        assert(0 &&


you can also do assert(!"something") or just llvm_unreachable("something")

Hardcode84 · 2022-10-31T23:01:06Z

Did it picked up cache from main? You could just change LLVM_CACHE_NUMBER to force cache rebuild

fschlimb · 2022-11-02T09:16:09Z

Did it picked up cache from main? You could just change LLVM_CACHE_NUMBER to force cache rebuild

Every branch has its own cache. So the CI will not reuse the cache of a given branch for the merge commit.
Only the cache from main can be used on other branches. That's why it is desirable to have main use the same LLVM SHA as refactor and its forks.
Changing the LLVM_CACHE_NUMBER will not help with this issue. It is only useful if you need to force a rebuild of the cache (for example because you think it's corrupted).

silee2 added 25 commits October 7, 2022 09:10

Temp save.

af9ab57

Temp save.

421ee75

Run pre-commit

a191334

Temp save.

56c2da8

Temp save.

0aa2826

Temp save.

2cdc105

Merge remote-tracking branch 'upstream/refactor' into end2end

7bb6afe

Hack to enable end2end from Linalg to Vulkan with linalg_addf.mlir

6c86be8

Make imex-runner.py more flexible with runners.

ad05757

Add options to InsertGPUAllocs

9988f7d

Fix test issues on Windows.

4e7fd75

Temp save.

796a55a

Re-enable insert gpu allocs tests.

97812b3

Do some cleanup

ca12c98

Add opencl pipeline.

f453874

Separate Vulkan and OpenCL test cases.

cfc6e92

Merge remote-tracking branch 'upstream/refactor' into share_pipeline

f802bec

Enable Vulkan end to end test only if mlir vulkan runner is enabled.

316fb9d

Remove unused files and run pre-commit.

f721e09

Fix compile issues on Windows.

655645b

Set path for tools used in imex-runner.py

236dbf9

Add missing dependency

88be36c

Add checks for client-api=vulkan

bb22b54

Run pre-commit.

7d08857

Fix boolean issue.

d529465

silee2 requested a review from nbpatel October 26, 2022 01:45

silee2 marked this pull request as ready for review October 26, 2022 21:27

silee2 requested review from fschlimb, Dewei-Wang-sh and drprajap October 26, 2022 21:28

silee2 added 3 commits October 26, 2022 16:13

Update LLVM SHA to match main branch.

dd0de7e

Change misleading pass name convert-gpux-to-spirv into imex-convert-g…

01e7a7d

…pu-to-spirv. Extend OpenCL pipeline to spirv serialization.

silee2 added 2 commits October 26, 2022 16:53

Merge remote-tracking branch 'upstream/refactor' into end2end_pr

8b7894e

Extend OpenCL pipeline to convert-func-to-llvm pass.

9a3d519

Add remaining part of pipeline as a comment (disabled for now).

silee2 mentioned this pull request Oct 27, 2022

Add dummy file. #417

Closed

6 tasks

silee2 added 3 commits October 26, 2022 20:10

Implement correct way to check if mlir-vulkan-runner is enabled.

a796e5c

Works for both separate llvm install and llvm external projects build.

Add small comment.

e8dac3e

Update comment.

8807b0e

fschlimb approved these changes Oct 27, 2022

View reviewed changes

drprajap reviewed Oct 28, 2022

View reviewed changes

Dewei-Wang-sh approved these changes Oct 31, 2022

View reviewed changes

silee2 added 2 commits October 31, 2022 10:40

Rename file to avoid confusion.

ea9d8a2

imex-runner.py.in is not a valid python file.

Offer mlir-vulkan-runner in imex-runner.py only if mlir-vulkan-runner is

46709fa

available.

silee2 mentioned this pull request Oct 31, 2022

Fix GPUXToSpirv Pass to iterate for all gpuModules #426

Merged

6 tasks

silee2 added 4 commits October 31, 2022 14:46

Merge remote-tracking branch 'upstream/refactor' into end2end_pr

485c44e

Update LLVM SHA

8c50f76

Reduce LLVM distribution size. Ported intel#423 from main.

89a107e

Add mlir-cpu-runner to LLVM distribution list.

cdbc469

Hardcode84 reviewed Oct 31, 2022

View reviewed changes

Update LLVM SHA

64fcbc4

silee2 merged commit a1c8eab into intel:refactor Nov 1, 2022

mshahneo pushed a commit to mshahneo/mlir-extensions that referenced this pull request May 1, 2024

Adding pvc target to CI (intel#415)

3d9f15c

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Enable end to end test with Linalg input and Vulkan backend #415

Enable end to end test with Linalg input and Vulkan backend #415

silee2 commented Oct 26, 2022 •

edited

Loading

silee2 commented Oct 26, 2022

silee2 commented Oct 26, 2022

fschlimb left a comment

fschlimb commented Oct 27, 2022

fschlimb commented Oct 27, 2022

drprajap Oct 28, 2022

Dewei-Wang-sh left a comment

Hardcode84 Oct 31, 2022

Hardcode84 commented Oct 31, 2022

fschlimb commented Nov 2, 2022 •

edited

Loading

Enable end to end test with Linalg input and Vulkan backend #415

Enable end to end test with Linalg input and Vulkan backend #415

Conversation

silee2 commented Oct 26, 2022 • edited Loading

silee2 commented Oct 26, 2022

silee2 commented Oct 26, 2022

fschlimb left a comment

Choose a reason for hiding this comment

fschlimb commented Oct 27, 2022

fschlimb commented Oct 27, 2022

drprajap Oct 28, 2022

Choose a reason for hiding this comment

Dewei-Wang-sh left a comment

Choose a reason for hiding this comment

Hardcode84 Oct 31, 2022

Choose a reason for hiding this comment

Hardcode84 commented Oct 31, 2022

fschlimb commented Nov 2, 2022 • edited Loading

silee2 commented Oct 26, 2022 •

edited

Loading

fschlimb commented Nov 2, 2022 •

edited

Loading