maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[OpenCL] Fix pipe size in TypeInfo. 

Pipes are now the size of pointers rather than the size
of the type that they contain.

Patch by Simon Perretta!

Differential Revision: https://reviews.llvm.org/D33597

llvm-svn: 304708
This commit is contained in:
Anastasia Stulova 2017-06-05 11:27:03 +00:00
parent f5d486f43d
commit b3398936ab
2 changed files with 18 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
clang/lib/AST/ASTContext.cpp
View File

@ -1939,9 +1939,8 @@ TypeInfo ASTContext::getTypeInfoImpl(const Type *T) const {
break;
case Type::Pipe: {
TypeInfo Info = getTypeInfo(cast<PipeType>(T)->getElementType());
Width = Info.Width;
Align = Info.Align;
Width = Target->getPointerWidth(getTargetAddressSpace(LangAS::opencl_global));
Align = Target->getPointerAlign(getTargetAddressSpace(LangAS::opencl_global));
}
}

16
clang/test/Index/pipe-size.cl Normal file
View File

@ -0,0 +1,16 @@
// RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple x86_64-unknown-linux-gnu %s -o - | FileCheck %s --check-prefix=X86
// RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR
// RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR64
// RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple amdgcn-amd-amdhsa-amdgizcl %s -o - | FileCheck %s --check-prefix=AMD
__kernel void testPipe( pipe int test )
{
int s = sizeof(test);
// X86: store %opencl.pipe_t* %test, %opencl.pipe_t** %test.addr, align 8
// X86: store i32 8, i32* %s, align 4
// SPIR: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)** %test.addr, align 4
// SPIR: store i32 4, i32* %s, align 4
// SPIR64: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)** %test.addr, align 8
// SPIR64: store i32 8, i32* %s, align 4
// AMD: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)* addrspace(5)* %test.addr, align 4
// AMD: store i32 8, i32 addrspace(5)* %s, align 4
}