Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Compiler crashes when using the same name for two different arguments in two different scopes #27

Open
ghost opened this issue Apr 4, 2016 · 0 comments

Comments

@ghost
Copy link

ghost commented Apr 4, 2016

I'm on the latest HLC compiler from the branch hsail-stable-3.7. I compiled with -O2.

The following kernel fails to compile.

__attribute__((noinline)) static uint Op(uint in)
{
    return 2*in;
}

__kernel void Bug1(uint in, __global uint *out)
{
    *out = Op(in);
}

The HLC compiler crashes with the following error message:

ERROR:  The following command failed with return code 134.
        ./llc -O2 -march=hsail64 -filetype=asm -o /tmp/cloc13948/temp.hsail /tmp/cloc13948/temp.opt.bc

llc: /home/dgeier/Documents/swarm64/code/hsa/HLC-HSAIL-Development-LLVM/lib/CodeGen/SelectionDAG/SelectionDAG.cpp:3346: llvm::SDValue llvm::SelectionDAG::getNode(unsigned int, llvm::SDLoc, llvm::EVT, llvm::SDValue, llvm::SDValue, const llvm::SDNodeFlags*): Assertion `N1.getValueType() == N2.getValueType() && N1.getValueType() == VT && "Binary operator types must match!"' failed.
0  llc             0x00000000022348d2 llvm::sys::PrintStackTrace(llvm::raw_ostream&) + 59
1  llc             0x0000000002234c26
2  llc             0x00000000022336b1
3  libpthread.so.0 0x00007fe6272b1d10
4  libc.so.6       0x00007fe626451267 gsignal + 55
5  libc.so.6       0x00007fe626452eca abort + 362
6  libc.so.6       0x00007fe62644a03d
7  libc.so.6       0x00007fe62644a0f2
8  llc             0x0000000002091357 llvm::SelectionDAG::getNode(unsigned int, llvm::SDLoc, llvm::EVT, llvm::SDValue, llvm::SDValue, llvm::SDNodeFlags const*) + 2285
9  llc             0x00000000017e2ba1 llvm::HSAILTargetLowering::getArgStore(llvm::SelectionDAG&, llvm::SDLoc, llvm::EVT, llvm::Type*, unsigned int, llvm::SDValue, llvm::SDValue, llvm::SDValue, unsigned int, llvm::SDValue, llvm::AAMDNodes const&, unsigned long) const + 741
10 llc             0x00000000017e360d llvm::HSAILTargetLowering::LowerArgument(llvm::SDValue, llvm::SDValue, bool, llvm::SmallVectorImpl<llvm::ISD::InputArg> const*, llvm::SmallVectorImpl<llvm::ISD::OutputArg> const*, llvm::SDLoc, llvm::SelectionDAG&, llvm::SmallVectorImpl<llvm::SDValue>*, unsigned int&, llvm::Type*, unsigned int, char const*, llvm::SDValue, llvm::SmallVectorImpl<llvm::SDValue> const*, bool, llvm::AAMDNodes const&, unsigned long) const + 2131
11 llc             0x00000000017e4ca5 llvm::HSAILTargetLowering::LowerCall(llvm::TargetLowering::CallLoweringInfo&, llvm::SmallVectorImpl<llvm::SDValue>&) const + 4339
12 llc             0x00000000020e5da8 llvm::TargetLowering::LowerCallTo(llvm::TargetLowering::CallLoweringInfo&) const + 4450
13 llc             0x00000000020dc4e2 llvm::SelectionDAGBuilder::lowerInvokable(llvm::TargetLowering::CallLoweringInfo&, llvm::MachineBasicBlock*) + 444
14 llc             0x00000000020dca61 llvm::SelectionDAGBuilder::LowerCallTo(llvm::ImmutableCallSite, llvm::SDValue, bool, llvm::MachineBasicBlock*) + 951
15 llc             0x00000000020df148 llvm::SelectionDAGBuilder::visitCall(llvm::CallInst const&) + 2294
16 llc             0x00000000020be4d6 llvm::SelectionDAGBuilder::visit(unsigned int, llvm::User const&) + 1256
17 llc             0x00000000020bdf49 llvm::SelectionDAGBuilder::visit(llvm::Instruction const&) + 155
18 llc             0x000000000210c037 llvm::SelectionDAGISel::SelectBasicBlock(llvm::ilist_iterator<llvm::Instruction const>, llvm::ilist_iterator<llvm::Instruction const>, bool&) + 83
19 llc             0x000000000210f3d0 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) + 2866
20 llc             0x000000000210b261 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) + 1049
21 llc             0x00000000017d2fd9
22 llc             0x0000000001a8061d llvm::MachineFunctionPass::runOnFunction(llvm::Function&) + 95
23 llc             0x0000000001e0cc9d llvm::FPPassManager::runOnFunction(llvm::Function&) + 305
24 llc             0x0000000001e0ce3c llvm::FPPassManager::runOnModule(llvm::Module&) + 112
25 llc             0x0000000001e0d1b4
26 llc             0x0000000001e0d894 llvm::legacy::PassManagerImpl::run(llvm::Module&) + 252
27 llc             0x0000000001e0dacd llvm::legacy::PassManager::run(llvm::Module&) + 39
28 llc             0x0000000000cb0376
29 llc             0x0000000000caf30a main + 257
30 libc.so.6       0x00007fe62643ca40 __libc_start_main + 240
31 llc             0x0000000000cae049 _start + 41
Stack dump:
0.  Program arguments: ./llc -O2 -march=hsail64 -filetype=asm -o /tmp/cloc13948/temp.hsail /tmp/cloc13948/temp.opt.bc 
1.  Running pass 'Function Pass Manager' on module '/tmp/cloc13948/temp.opt.bc'.
2.  Running pass 'HSAIL DAG->DAG Instruction Selection' on function '@__OpenCL_Bug1_kernel'
/opt/amd/bin/cloc.sh: line 333: 13966 Aborted                 (core dumped) $HSA_LLVM_PATH/$CMD_LLC -o $TMPDIR/$FNAME.hsail $TMPDIR/$FNAME.opt.bc

Weirdly, after changing the name of the in argument of Op() to in2, the error disappears. The following kernel compilers successfully.

__attribute__((noinline)) static uint Op(uint in2)
{
    return 2*in2;
}

__kernel void Bug1(uint in, __global uint *out)
{
    *out = Op(in);
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

0 participants