Skip to content
This repository was archived by the owner on Oct 9, 2019. It is now read-only.

Adds handling for llvm.memmove-intrinsic #209

Merged

Conversation

doe300
Copy link
Contributor

@doe300 doe300 commented Apr 29, 2017

Closes #205.

Produces following SPIR-V output for the example given in #205:

119734787 65536 393230 22 0 
2 Capability Addresses 
2 Capability Kernel 
2 Capability Vector16 
2 Capability Int8 
5 ExtInstImport 1 "OpenCL.std"
3 MemoryModel 1 2 
6 EntryPoint 6 12 "test_struct"
3 Source 3 102000 
7 Name 3 "struct.SomeStruct"
3 Name 13 "in"
3 Name 14 "out"
4 Decorate 21 FuncParamAttr 5 
2 DecorationGroup 21 
4 Decorate 13 FuncParamAttr 6 
4 GroupDecorate 21 13 14 
4 TypeInt 6 32 0 
4 TypeInt 8 8 0 
4 Constant 6 7 60 
4 Constant 6 20 128 
2 TypeVoid 2 
3 TypeFloat 4 32 
4 TypeVector 5 4 16 
4 TypeArray 9 8 7 
5 TypeStruct 3 5 6 9 
4 TypePointer 10 5 3 
5 TypeFunction 11 2 10 10 
4 TypePointer 16 5 8 

5 Function 2 12 0 11 
3 FunctionParameter 10 13 
3 FunctionParameter 10 14 

2 Label 15 
4 Variable 2 19 7 
4 Bitcast 16 17 13 
4 Bitcast 16 18 14 
3 LifetimeStart 19 128 
6 CopyMemorySized 19 17 20 2 64 
6 CopyMemorySized 18 19 20 2 64 
3 LifetimeStop 19 128 
1 Return 

1 FunctionEnd 

This result is "correct" in the sense, that the code does the right thing, but:

  • The temporary Variable is of type void and has no size specified.
  • We probably should add LifetimeStart and LifeTimeStop before/after the CopyMemorySized instruction.
  • According to the SPIR-V specification, section 2.4:

All OpVariable instructions in a function must be in the first block in the function. These instructions, together with any immediately preceding OpLine instructions, must be the first instructions in that block. (Note the validation rules prevent OpPhi instructions in the first block of a function.)

@yxsamliu
Copy link
Contributor

You can add SPIRVFunction::addVariable to make sure the variable instruction is properly ordered, then modify SPIRVModuleImpl::addVariable to call it.

@doe300
Copy link
Contributor Author

doe300 commented May 13, 2017

So, this version does everything correct, as far as I can tell and is ready to be pulled.

SPIRVVariable *SPIRVFunction::addVariable(SPIRVVariable *Var)
{
//according to the SPIR-V specs, the Variable (with Storage Class Function) need to be at the beginning of the first block of the function, see section 2.4 (Logical Layout of a Module)
return BBVec.at(0)->addVariable(Var);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

needs to check if BBVec is empty. If so add a BB.

assert(Var && "Invalid variable");
Module->add(Var);
Var->setParent(this);
InstVec.insert(InstVec.begin(), Var);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This will cause auto vars in source code emitted in reverse order and break existing lit tests.

Should find the first non-Variable instruction and insert before that.

@yxsamliu
Copy link
Contributor

Please check the Travis CI build status. It should pass.

@doe300
Copy link
Contributor Author

doe300 commented May 16, 2017

Made the requested changes. Also now the tests compile.

I'm not quite sure though about adding the new BasicBlock in SPIRVFunction::addVariable (line 70). Do I have to specify an ID? If so, where do I get it from?

Copy link
Contributor

@bader bader left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It would be great if you add a regression test validating correct behavior. You can use the test case from the issue #205.
Just, a side note: backward translation will not be able to recover original llvm intrinsic.

Var->setParent(this);
auto It = InstVec.begin();
while(It != InstVec.end())
{
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Minor style comment.
Please, follow the existing style - move { to the line above.
I think the whole loop can be replaced with standard find algorithm.

@@ -67,6 +67,16 @@ SPIRVFunctionParameter::foreachAttr(
}
}

SPIRVVariable *SPIRVFunction::addVariable(SPIRVVariable *Var)
{
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please, move the brace to the line above.

{
//according to the SPIR-V specs, the Variable (with Storage Class Function) need to be at the beginning of the first block of the function, see section 2.4 (Logical Layout of a Module)
if(getNumBasicBlock() == 0)
{
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please, move the brace to the line above.

@doe300
Copy link
Contributor Author

doe300 commented May 17, 2017

Applied the style-changes and added a test case.

Is this test-case any good? If not, how should I write it?


; CHECK-NOT: llvm.memmove

; CHECK: Variable [[voidPtr:[0-9]+]] [[mem:[0-9]+]] 7
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There is no need in voidPtr variable - it's not used in the test. Suggest replacing with {{[0-9]+}}.

@doe300
Copy link
Contributor Author

doe300 commented May 17, 2017

The tests in DEBUG mode on travis CI fail with "Object type must be an integer type scalar" in SPIRVInstruction.h:1882.

I don't understand, why the pointer-type needs to be an integer. Especially, because the next line seems to accept the void* type.

@bader
Copy link
Contributor

bader commented May 18, 2017

It looks like a bug in the assertion statement.

Here is the quote from the spec:
"Pointer is a pointer to the object whose lifetime is starting. Its type must be an OpTypePointer with Storage Class Function.

Size must be 0 if Pointer is a pointer to a non-void type or the Addresses capability is not being used. If Size is non-zero, it is the number of bytes of memory whose lifetime is starting. Its type must be an integer type scalar. It is treated as unsigned; if its type has Signedness of 1, its sign bit cannot be set."

"Size" parameter type must be integer. "Pointer" can point to non-integer type values.

Could you fix the assertion statement, please?

Copy link
Contributor

@bader bader left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

One minor comment.

@@ -1877,7 +1877,6 @@ class SPIRVLifetime : public SPIRVInstruction {
assert(Obj->getStorageClass() == StorageClassFunction &&
"Invalid storage class");
assert(Obj->getType()->isTypePointer() &&
Obj->getType()->getPointerElementType()->isTypeInt() &&
"Object type must be an integer type scalar");
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please, update the message to avoid confusion.
"Objects type must be a pointer".

@@ -67,6 +67,14 @@ SPIRVFunctionParameter::foreachAttr(
}
}

SPIRVVariable *SPIRVFunction::addVariable(SPIRVVariable *Var) {
//according to the SPIR-V specs, the Variable (with Storage Class Function) need to be at the beginning of the first block of the function, see section 2.4 (Logical Layout of a Module)
if(getNumBasicBlock() == 0) {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

insert space after if

@yxsamliu
Copy link
Contributor

There is still travis-ci build failure. Looking at the result:

FAIL: LLVM :: SPIRV/llvm.memmove.ll (8805 of 12438)
******************** TEST 'LLVM :: SPIRV/llvm.memmove.ll' FAILED ********************
Script:

/home/travis/build/KhronosGroup/SPIRV-LLVM/build/./bin/llvm-as /home/travis/build/KhronosGroup/SPIRV-LLVM/test/SPIRV/llvm.memmove.ll -o /home/travis/build/KhronosGroup/SPIRV-LLVM/build/test/SPIRV/Output/llvm.memmove.ll.tmp.bc
/home/travis/build/KhronosGroup/SPIRV-LLVM/build/./bin/llvm-spirv /home/travis/build/KhronosGroup/SPIRV-LLVM/build/test/SPIRV/Output/llvm.memmove.ll.tmp.bc -spirv-text -o - | /home/travis/build/KhronosGroup/SPIRV-LLVM/build/./bin/FileCheck /home/travis/build/KhronosGroup/SPIRV-LLVM/test/SPIRV/llvm.memmove.ll

Exit Code: 2

Command Output (stderr):

llvm-spirv: /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/SPIRV/libSPIRV/SPIRVInstruction.h:1880: virtual void SPIRV::SPIRVLifetime<256>::validate() const [OC = 256]: Assertion `Obj->getType()->isTypePointer() && "Objects type must be a pointer"' failed.
Stack dump:
0. Program arguments: /home/travis/build/KhronosGroup/SPIRV-LLVM/build/./bin/llvm-spirv /home/travis/build/KhronosGroup/SPIRV-LLVM/build/test/SPIRV/Output/llvm.memmove.ll.tmp.bc -spirv-text -o -

  1. Running pass 'LLVMToSPIRV' on module '/home/travis/build/KhronosGroup/SPIRV-LLVM/build/test/SPIRV/Output/llvm.memmove.ll.tmp.bc'.
    #0 0x80b21e llvm::sys::PrintStackTrace(_IO_FILE*) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/Support/Unix/Signals.inc:422:15
    Result IDs start at 0 in SPIR-V binary produced by llvm-spirv #1 0x80bfdb PrintStackTraceSignalHandler(void*) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/Support/Unix/Signals.inc:481:1
    [SPIRV] Have SPIRV IDs start at 1 #2 0x80e704 SignalHandler(int) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/Support/Unix/Signals.inc:198:60
    Generated Linkage Attribute OpDecorate doesn't contain a literal string #3 0x2adbb96de330 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x10330)
    Add support for OpDecorate LinkageAttributes name #4 0x2adbba35bc37 gsignal (/lib/x86_64-linux-gnu/libc.so.6+0x36c37)
    Capabilities required by OpDecorate are not added to SPIR-V #5 0x2adbba35f028 abort (/lib/x86_64-linux-gnu/libc.so.6+0x3a028)
    get_local_id(X) marked as LinkageAttributes export #6 0x2adbba354bf6 (/lib/x86_64-linux-gnu/libc.so.6+0x2fbf6)
    Fix SPV_VERSION definition #7 0x2adbba354ca2 (/lib/x86_64-linux-gnu/libc.so.6+0x2fca2)
    Update spirv.h to latest available version #8 0x4a85f6 SPIRV::SPIRVLifetime<(spv::Op)256>::validate() const /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/SPIRV/libSPIRV/SPIRVInstruction.h:1881:10
    Fix instructions order when writing the binary module #9 0x498c4d SPIRV::SPIRVLifetime<(spv::Op)256>::SPIRVLifetime(unsigned int, unsigned int, SPIRV::SPIRVBasicBlock*) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/SPIRV/libSPIRV/SPIRVInstruction.h:1861:5
    OpTypeImage reports wrong number of operands #10 0x498b6d SPIRV::SPIRVModuleImpl::addLifetimeInst(spv::Op, SPIRV::SPIRVValue*, unsigned int, SPIRV::SPIRVBasicBlock*) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/SPIRV/libSPIRV/SPIRVModule.cpp:1125:12
    [SPIRV] github issue#5 Capabilities required by OpDecorate are not ad… #11 0x5919d9 SPIRV::LLVMToSPIRV::transIntrinsicInst(llvm::IntrinsicInst*, SPIRV::SPIRVBasicBlock*) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/SPIRV/SPIRVWriter.cpp:1313:5
    OpSource should allow optional string operand for source code #12 0x58f1b7 SPIRV::LLVMToSPIRV::transValueWithoutDecoration(llvm::Value*, SPIRV::SPIRVBasicBlock*, bool) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/SPIRV/SPIRVWriter.cpp:1157:22
    [Feature Request] OpDecorationGroup doesn't group more than one decoration. #13 0x58a1b3 SPIRV::LLVMToSPIRV::transValue(llvm::Value*, SPIRV::SPIRVBasicBlock*, bool) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/SPIRV/SPIRVWriter.cpp:812:13
    Replace eraseUselessFunctions with a more specific one #14 0x5938b2 SPIRV::LLVMToSPIRV::transFunction(llvm::Function*) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/SPIRV/SPIRVWriter.cpp:1486:7
    Debug information issue. #15 0x593d8b SPIRV::LLVMToSPIRV::translate() /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/SPIRV/SPIRVWriter.cpp:1530:5
    Garbage passed as an argument to llvm::Value::setName #16 0x5a047f SPIRV::LLVMToSPIRV::runOnModule(llvm::Module&) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/SPIRV/SPIRVWriter.cpp:178:5
    Khronos/spirv 3.6.1 out #17 0x73da7c (anonymous namespace)::MPPassManager::runOnModule(llvm::Module&) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/IR/LegacyPassManager.cpp:1616:23
    [SPIRV] Fix capabilities dependencies #18 0x73d65e llvm::legacy::PassManagerImpl::run(llvm::Module&) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/IR/LegacyPassManager.cpp:1723:16
    SPIR-V generator is crashed on a simple device execution kernel. #19 0x73e041 llvm::legacy::PassManager::run(llvm::Module&) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/IR/LegacyPassManager.cpp:1756:10
    SPIR-V reader translate some conversion/bitcast instructions incorrectly to OCL builtin functions #20 0x5967d1 llvm::WriteSPIRV(llvm::Module*, llvm::raw_ostream&, std::string&) /home/travis/build/KhronosGroup/SPIRV-LLVM/lib/SPIRV/SPIRVWriter.cpp:1857:3
    Fix vector convert #21 0x409828 convertLLVMToSPIRV() /home/travis/build/KhronosGroup/SPIRV-LLVM/tools/llvm-spirv/llvm-spirv.cpp:154:8
    Duplicate OpTypeImage definition #22 0x408fd8 main /home/travis/build/KhronosGroup/SPIRV-LLVM/tools/llvm-spirv/llvm-spirv.cpp:312:12
    Missing SPIRVSubTarget.h file #23 0x2adbba346f45 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x21f45)
    [SPIRV] Fix typo in include of SPIRVSubtarget.h #24 0x4088e4 _start (/home/travis/build/KhronosGroup/SPIRV-LLVM/build/bin/llvm-spirv+0x4088e4)

Can you fix that?

@doe300
Copy link
Contributor Author

doe300 commented May 21, 2017

I can't figure out, why this test fails. So, besides removing the test I don't know of any fix.

@@ -1303,6 +1303,28 @@ LLVMToSPIRV::transIntrinsicInst(IntrinsicInst *II, SPIRVBasicBlock *BB) {
transValue(II->getOperand(2), BB),
getMemoryAccess(cast<MemIntrinsic>(II)),
BB);
case Intrinsic::memmove: {
//Since memory areas for source and dest can overlap, 1. copy to temporary, 2. copy to destination
SPIRVType *Ty = transType(II->getType());
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

should be II->getOperand(0)->getType()

@@ -0,0 +1,50 @@
; RUN: llvm-as %s -o %t.bc
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

can we make it a transcodeing test to make sure the generated SPIRV can be translated back to LLVM IR?

@yxsamliu
Copy link
Contributor

I think I know why. The type passed to addVariable is wrong. I've added inline comments about that.

@yxsamliu
Copy link
Contributor

Also, can you add an assertion to SPIRVVariable::validate to assert the type is pointer type? Thanks.

@doe300
Copy link
Contributor Author

doe300 commented May 23, 2017

Applied the requested changes, the trancoding-test fails, since the method-signatures do not match for the llvm.memcpy calls:

Fails to verify module: Call parameter type does not match function signature!
  %1 = alloca i8
 i8 addrspace(1)*  call void @llvm.memcpy.p1i8.p1i8.i32(i8* %1, i8 addrspace(1)* %2, i32 128, i32 64, i1 false)
Call parameter type does not match function signature!
  %1 = alloca i8
 i8 addrspace(1)*  call void @llvm.memcpy.p1i8.p1i8.i32(i8 addrspace(1)* %3, i8* %1, i32 128, i32 64, i1 false)

With the single function declaration for llvm.memcpy:

declare void @llvm.memcpy.p1i8.p1i8.i32(i8 addrspace(1)* nocapture, i8 addrspace(1)* nocapture readonly, i32, i32, i1) #0

As far as i can tell, there should be 2 declarations, something like:

declare void @llvm.memcpy.p1i8.p1i8.i32(i8*, i8 addrspace(1)* nocapture readonly, i32, i32, i1)
declare void @llvm.memcpy.p1i8.p1i8.i32(i8 addrspace(1)* nocapture, i8* readonly, i32, i32, i1)

@yxsamliu
Copy link
Contributor

It seems to be more complicated than expected. One issue is the type of the temporary variable. It should be either a pointer to the object type to be moved, or a pointer to an array of chars which have the same size of the memmov operation.

Another issue is that adding an existing type to SPIR-V module may cause problem. The SPIR-V module does not have a fold set to guarantee uniqueness of types, so the user has to make sure no duplicate types are added to SPIR-V module.

For this type of transformation, it may be easier to write a pass like SPIRVLowerBool to transform memmov to memcpy in LLVM IR.

@AlexeySotkin
Copy link
Contributor

Using -O3 we transform two memcpy calls to single memmove call. For translation to SPIR-V purposes it looks reasonable to transform the memmove back to memcpy.

Maybe we just should not use -O3 in the first place?
Moreover I don't see any benefits of using -O3, if we end up with the same SPIR-V code as we had with -O0.

@yxsamliu
Copy link
Contributor

That can be a workaround if users are using Clang. I think it is better to be able to lower memmov.

@doe300
Copy link
Contributor Author

doe300 commented Jun 6, 2017

I need the other optimizations run with -O3, so disabling optimizations is not an option.
I rewrote the pull-request using a lowering pass as suggested by @yxsamliu .

auto Src = I.getRawSource();
auto SrcTy = Src->getType();
ConstantInt *Length = nullptr;
if(isa<ConstantInt>(I.getLength())) //the length (in bytes) could be non-constant
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

space between if and (

virtual void visitMemMoveInst(MemMoveInst &I) {
IRBuilder<> Builder(I.getParent());
Builder.SetInsertPoint(&I);
auto Dest = I.getRawDest();
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

For pointers better use auto *. Same as below.

auto SrcTy = Src->getType();
ConstantInt *Length = nullptr;
if(isa<ConstantInt>(I.getLength())) //the length (in bytes) could be non-constant
Length = ConstantInt::get(Type::getInt64Ty(*Context), cast<ConstantInt>(I.getLength())->getZExtValue());
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can you just use I.getLength() instead of create a new ConstantInt?

auto Align = I.getAlignment();
auto Volatile = I.isVolatile();

auto Alloca = Builder.CreateAlloca(SrcTy->getPointerElementType(), SrcTy->isArrayTy() ? Builder.getInt32(SrcTy->getArrayNumElements()) : nullptr);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

this line may exceed 80 columns.
Please follow http://llvm.org/docs/CodingStandards.html

auto Src = I.getRawSource();
auto SrcTy = Src->getType();
ConstantInt *Length = nullptr;
if(isa<ConstantInt>(I.getLength())) //the length (in bytes) could be non-constant
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If it is not constant we can create a loop and each time copy one chunk of fixed length.

However, for now, maybe just report_fatal_error as "llvm.memmov of non-constant length not supported".

Length = ConstantInt::get(Type::getInt64Ty(*Context), cast<ConstantInt>(I.getLength())->getZExtValue());
if(isa<BitCastInst>(Src)) //the source could be bit-cast from another type, need the original type for the allocation
SrcTy = cast<BitCastInst>(Src)->getOperand(0)->getType();
auto Align = I.getAlignment();
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

need to check the size of the variable matches the length of memmov. If not report_fatal_error.

@@ -105,6 +125,7 @@ class SPIRVLowerMemmove: public ModulePass,
static char ID;
private:
LLVMContext *Context;
Module *Mod;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this member used?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

line 90 to query the size of the allocated type: Mod->getDataLayout()->getTypeSizeInBits(SrcTy->getPointerElementType()). I created the member, since I didn't know how else to get the Module instance.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sorry I missed that. Then it looks fine.

@yxsamliu
Copy link
Contributor

There are too many commits in this pull request. Can you squash them into one commit? Otherwise LGTM.

@doe300 doe300 force-pushed the khronos/spirv-3.6.1 branch from c549c2b to d60f924 Compare June 19, 2017 19:01
@yxsamliu yxsamliu merged commit 034d5dd into KhronosGroup:khronos/spirv-3.6.1 Jun 29, 2017
robclark pushed a commit to freedreno/SPIRV-LLVM that referenced this pull request Jan 26, 2018
AlexeySotkin pushed a commit to AlexeySotkin/SPIRV-LLVM that referenced this pull request Mar 12, 2018
AlexeySotkin pushed a commit to AlexeySotkin/SPIRV-LLVM that referenced this pull request Apr 3, 2018
Sign up for free to subscribe to this conversation on GitHub. Already have an account? Sign in.
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants