[ROCm] Adding ROCDL Dialect. #146

deven-amd · 2019-09-26T00:42:34Z

This PR introduces the ROCDL Dialect (i.e. the ROCDL ops + the code to lower those ROCDL ops to LLVM intrinsics/functions). Think of ROCDL Dialect as analogous to the NVVM Dialect, but for AMD GPUs. This PR contains just the essentials needed to get a simple example up and running. We expect to make further additions to the ROCDL Dialect.

This is the first of 3 PRs we will be filing. Once this PR is merged it will be followed by

a PR to add a pass that lowers GPU Dialect to ROCDL Dialect
a PR to add a "mlir-rocm-runner" utility

Note: This PR is a follow-up on PR #63. It builds on the work done by @whchung in that PR, and addresses most of the review comments in that PR.

@whchung @joker-eph

joker-eph

Looks great! Thanks

include/mlir/Dialect/LLVMIR/ROCDLDialect.h

lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp

lib/Target/LLVMIR/ConvertToROCDLIR.cpp

joker-eph · 2019-09-26T01:25:29Z

The test failed on MacOS:

FAIL: MLIR :: Dialect/LLVMIR/rocdl.mlir (130 of 293)
******************** TEST 'MLIR :: Dialect/LLVMIR/rocdl.mlir' FAILED ********************
Script:
--
: 'RUN: at line 1';   /tmpfs/src/build/bin/mlir-translate -mlir-to-rocdlir /Volumes/BuildData/tmpfs/src/github/mlir/test/Dialect/LLVMIR/rocdl.mlir | /tmpfs/src/build/bin/FileCheck /Volumes/BuildData/tmpfs/src/github/mlir/test/Dialect/LLVMIR/rocdl.mlir
--
Exit Code: 2

Command Output (stderr):
--
Stack dump:
0.	Program arguments: /tmpfs/src/build/bin/mlir-translate -mlir-to-rocdlir /Volumes/BuildData/tmpfs/src/github/mlir/test/Dialect/LLVMIR/rocdl.mlir 
0  mlir-translate           0x00000001073501a8 llvm::sys::PrintStackTrace(llvm::raw_ostream&) + 40
1  mlir-translate           0x000000010734f208 llvm::sys::RunSignalHandlers() + 248
2  mlir-translate           0x00000001073507c8 SignalHandler(int) + 264
3  libsystem_platform.dylib 0x00007fff5570df5a _sigtramp + 26
4  libsystem_platform.dylib 0x0000000000001218 _sigtramp + 18446603339082707672
5  mlir-translate           0x00000001071ff2ed llvm::CallInst::Create(llvm::FunctionType*, llvm::Value*, llvm::ArrayRef<llvm::Value*>, llvm::ArrayRef<llvm::OperandBundleDefT<llvm::Value*> >, llvm::Twine const&, llvm::Instruction*) + 333
6  mlir-translate           0x00000001071ff042 llvm::IRBuilder<llvm::ConstantFolder, llvm::IRBuilderDefaultInserter>::CreateCall(llvm::FunctionType*, llvm::Value*, llvm::ArrayRef<llvm::Value*>, llvm::Twine const&, llvm::MDNode*) + 66
7  mlir-translate           0x0000000107200aed (anonymous namespace)::createDeviceFunctionCall(llvm::IRBuilder<llvm::ConstantFolder, llvm::IRBuilderDefaultInserter>&, llvm::StringRef, int) + 189
8  mlir-translate           0x0000000107200105 (anonymous namespace)::ModuleTranslation::convertOperation(mlir::Operation&, llvm::IRBuilder<llvm::ConstantFolder, llvm::IRBuilderDefaultInserter>&) + 853
9  mlir-translate           0x0000000107367ae7 mlir::LLVM::ModuleTranslation::convertBlock(mlir::Block&, bool) + 695
10 mlir-translate           0x0000000107368dc5 mlir::LLVM::ModuleTranslation::convertOneFunction(mlir::FuncOp) + 2005
11 mlir-translate           0x0000000107369470 mlir::LLVM::ModuleTranslation::convertFunctions() + 1536
12 mlir-translate           0x00000001071ffba4 mlir::translateModuleToROCDLIR(mlir::ModuleOp) + 372
13 mlir-translate           0x0000000107200b89 std::__1::__function::__func<$_0, std::__1::allocator<$_0>, mlir::LogicalResult (mlir::ModuleOp, llvm::raw_ostream&)>::operator()(mlir::ModuleOp&&, llvm::raw_ostream&) + 25
14 mlir-translate           0x00000001072a618d std::__1::__function::__func<mlir::TranslationParser::TranslationParser(llvm::cl::Option&)::$_1, std::__1::allocator<mlir::TranslationParser::TranslationParser(llvm::cl::Option&)::$_1>, mlir::LogicalResult (std::__1::unique_ptr<llvm::MemoryBuffer, std::__1::default_delete<llvm::MemoryBuffer> >, llvm::raw_ostream&, mlir::MLIRContext*)>::operator()(std::__1::unique_ptr<llvm::MemoryBuffer, std::__1::default_delete<llvm::MemoryBuffer> >&&, llvm::raw_ostream&, mlir::MLIRContext*&&) + 189
15 mlir-translate           0x0000000107250663 main + 1155
16 libdyld.dylib            0x00007fff553ff015 start + 1
FileCheck error: '-' is empty.
FileCheck command line:  /tmpfs/src/build/bin/FileCheck /Volumes/BuildData/tmpfs/src/github/mlir/test/Dialect/LLVMIR/rocdl.mlir

deven-amd · 2019-09-26T01:36:41Z

hmm...same failure on Ubuntu as well, but test is passing for me locally

[1/1] Running lit suite /home/rocm-user/llvm-project/llvm/projects/mlir/test/Dialect/LLVMIR
Testing Time: 0.22s
  Expected Passes    : 7

looking into it further

deven-amd · 2019-09-26T03:10:47Z

Still do not know the cause of the CI failures. My local runs pass, even the one I tried on a machine with no ROCm libraries installed!

Hoping the CI failure was an anomaly and the re-run for the new commit passes.

deven-amd · 2019-09-26T03:15:45Z

links to the CI failure logs from the first run
(just in case we need to refer to them again)

Mac : https://source.cloud.google.com/results/invocations/f5be02e6-c7da-4c2e-b72a-0dc7c946b568/log

Ubuntu : https://source.cloud.google.com/results/invocations/0ceb0d9a-d6c2-455a-b2ec-f55181683180/log

joker-eph · 2019-09-26T05:35:47Z

Hoping the CI failure was an anomaly

I'm not a big believer in the cosmic ray effect here ;)

Did you try the same configuration as the CI? (Release mode?)

deven-amd · 2019-09-26T10:56:46Z

Default is Release mode..right?

Here are the commands I am using to do the build and run

cmake -G Ninja ../llvm -DLLVM_BUILD_EXAMPLES=ON -DLLVM_ENABLE_CXX1Y=Y -DLLVM_TARGETS_TO_BUILD="host"
cmake --build . --target check-mlir

I see that the -DLLVM_ENABLE_CXX1Y=Y is no longer present in the instructions provided on README.md...wonder if that is the difference. Let me try that out.

deven-amd · 2019-09-26T11:31:48Z

that is not it either...

[1584/1584] Running the MLIR regression tests
Testing Time: 4.08s
  Expected Passes    : 289
  Unsupported Tests  : 4

deven-amd · 2019-09-26T11:32:37Z

How do I go about replicating exactly what the Ubuntu CI job is running?

deven-amd · 2019-09-26T12:27:49Z

able to reproduce the error locally now.

It is a Debug vs Release build issue (default build is Debug ... )

looking into it

deven-amd · 2019-09-26T15:11:19Z

found my bug, and fixed it.

@joker-eph , I have squashed all my changes into one commit, and pushed it out. Please re-approve.

thanks

deven

deven-amd · 2019-09-26T18:56:35Z

This PR is missing a test file

test/Target/rocdl.mlir is missing
- also the contents of test/Dialect/LLVMIR/rocdl.mlir contents should be moved here
test/Dialect/LLVMIR/rocdl.mlir need updating (use mlir-opt instead of mlir-transform)

will push out these changes as a separate commit in the next PR (Lowering GPU Dialect to ROCDL Dialect), which will be filed once this one is merged.

@joker-eph , let me know if you would rather have that change in this PR

joker-eph · 2019-09-27T00:02:31Z

I rather have these in the current PR if you can push?

test/Dialect/LLVMIR/rocdl.mlir

This PR introduces the ROCDL Dialect (i.e. the ROCDL ops + the code to lower those ROCDL ops to LLWM intrinsics/functions). Think of ROCDL Dialect as analogous to the NVVM Dialect, but for AMD GPUs. This PR contains just the essentials needed to get a simple example up and running. We expect to make further additions to the ROCDL Dialect. This is the first of 3 PRs we will be filing. Once this PR is merged it will be followed by * a PR to add a pass that lowers GPU Dialect to ROCDL Dialect * a PR to add a "mlir-rocm-runner" utility

deven-amd · 2019-09-27T00:41:18Z

I rather have these in the current PR if you can push?

done. Please re-approve. thanks.

This commit introduces the ROCDL Dialect (i.e. the ROCDL ops + the code to lower those ROCDL ops to LLWM intrinsics/functions). Think of ROCDL Dialect as analogous to the NVVM Dialect, but for AMD GPUs. This patch contains just the essentials needed to get a simple example up and running. We expect to make further additions to the ROCDL Dialect. This is the first of 3 commits, the follow-up will be: * add a pass that lowers GPU Dialect to ROCDL Dialect * add a "mlir-rocm-runner" utility Closes #146 COPYBARA_INTEGRATE_REVIEW=tensorflow/mlir#146 from deven-amd:deven-rocdl-dialect e78e8005c75a78912631116c78dc844fcc4b0de9 PiperOrigin-RevId: 271511259

This commit introduces the ROCDL Dialect (i.e. the ROCDL ops + the code to lower those ROCDL ops to LLWM intrinsics/functions). Think of ROCDL Dialect as analogous to the NVVM Dialect, but for AMD GPUs. This patch contains just the essentials needed to get a simple example up and running. We expect to make further additions to the ROCDL Dialect. This is the first of 3 commits, the follow-up will be: * add a pass that lowers GPU Dialect to ROCDL Dialect * add a "mlir-rocm-runner" utility Closes tensorflow/mlir#146 COPYBARA_INTEGRATE_REVIEW=tensorflow/mlir#146 from deven-amd:deven-rocdl-dialect e78e8005c75a78912631116c78dc844fcc4b0de9 PiperOrigin-RevId: 271511259

googlebot added the cla: yes label Sep 26, 2019

joker-eph added the kokoro:run label Sep 26, 2019

kokoro-team removed the kokoro:run label Sep 26, 2019

joker-eph reviewed Sep 26, 2019

View reviewed changes

include/mlir/Dialect/LLVMIR/ROCDLDialect.h Show resolved Hide resolved

lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp Show resolved Hide resolved

lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp Outdated Show resolved Hide resolved

lib/Target/LLVMIR/ConvertToROCDLIR.cpp Show resolved Hide resolved

whchung mentioned this pull request Sep 26, 2019

Optimise the gridDim.n * blockDim.m idiom ROCm/hip#1468

Merged

joker-eph added kokoro:run ready to pull labels Sep 26, 2019

kokoro-team removed the kokoro:run label Sep 26, 2019

joker-eph reviewed Sep 27, 2019

View reviewed changes

test/Dialect/LLVMIR/rocdl.mlir Outdated Show resolved Hide resolved

mlir-copybara-bot closed this in 182ba06 Sep 27, 2019

deven-amd mentioned this pull request Sep 27, 2019

[ROCm] Adding pass to lower GPU Dialect to ROCDL Dialect. #154

Closed

Uh oh!

[ROCm] Adding ROCDL Dialect. #146

[ROCm] Adding ROCDL Dialect. #146

Uh oh!

Conversation

deven-amd commented Sep 26, 2019 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

joker-eph left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

joker-eph commented Sep 26, 2019

Uh oh!

deven-amd commented Sep 26, 2019

Uh oh!

deven-amd commented Sep 26, 2019

Uh oh!

deven-amd commented Sep 26, 2019

Uh oh!

joker-eph commented Sep 26, 2019

Uh oh!

deven-amd commented Sep 26, 2019

Uh oh!

deven-amd commented Sep 26, 2019

Uh oh!

deven-amd commented Sep 26, 2019

Uh oh!

deven-amd commented Sep 26, 2019

Uh oh!

deven-amd commented Sep 26, 2019

Uh oh!

deven-amd commented Sep 26, 2019

Uh oh!

joker-eph commented Sep 27, 2019

Uh oh!

Uh oh!

deven-amd commented Sep 27, 2019

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

deven-amd commented Sep 26, 2019 •

edited

Loading