[OpenCL] Fix pipe size in TypeInfo.

[Anastasia]

Summary

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

Patch by Simon Perretta!

[yaxunl]

lib/AST/ASTContext.cpp

1942 ↗

(On Diff #100403)

pipe is in global address space. shouldn't it be

Width = Target->getPointerWidth(getTargetAddressSpace(LangAS::opencl_global));

the same with align.

[Anastasia]

Use global AS pointer for pipe size.

@sam, I am moving you to the reviewer to finish this change. Do you think it makes sense to add RUN line with some AMD GPU in triple to the test?

[yaxunl]

In D33597#770137, @Anastasia wrote:

Use global AS pointer for pipe size.

@sam, I am moving you to the reviewer to finish this change. Do you think it makes sense to add RUN line with some AMD GPU in triple to the test?

I would appreciate that if you can add one RUN line for triple amdgcn-amd-amdhsa-amdgizcl. Thanks.

[Anastasia]

Added RUN line for AMD

[yaxunl]

LGTM. Thanks!

cfe/trunk/lib/AST/ASTContext.cpp

if (Width != 0 && Width <= Target->getMaxAtomicPromoteWidth()) {
 
       // Set the alignment equal to the size.
       Align = static_cast<unsigned>(Width);
     }
   }
   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));
   }
 
   }
 
   assert(llvm::isPowerOf2_32(Align) && "Alignment must be power of 2");
   return TypeInfo(Width, Align, AlignIsRequired);
 }
 

cfe/trunk/test/Index/pipe-size.cl

+// 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
+}