r304708 - [OpenCL] Fix pipe size in TypeInfo.

[Anastasia Stulova via cfe-commits]

Author: stulova
Date: Mon Jun  5 06:27:03 2017
New Revision: 304708

URL: http://llvm.org/viewvc/llvm-project?rev=304708&view=rev
Log:
[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


Added:
    cfe/trunk/test/Index/pipe-size.cl
Modified:
    cfe/trunk/lib/AST/ASTContext.cpp

Modified: cfe/trunk/lib/AST/ASTContext.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ASTContext.cpp?rev=304708&r1=304707&r2=304708&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ASTContext.cpp (original)
+++ cfe/trunk/lib/AST/ASTContext.cpp Mon Jun  5 06:27:03 2017
@@ -1939,9 +1939,8 @@ TypeInfo ASTContext::getTypeInfoImpl(con
   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));
   }
 
   }

Added: cfe/trunk/test/Index/pipe-size.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Index/pipe-size.cl?rev=304708&view=auto
==============================================================================
--- cfe/trunk/test/Index/pipe-size.cl (added)
+++ cfe/trunk/test/Index/pipe-size.cl Mon Jun  5 06:27:03 2017
@@ -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
+}