From d5e54601f1177956d8206c08b1c8d7d5fd8e2407 Mon Sep 17 00:00:00 2001 From: Anastasia Stulova Date: Mon, 5 Jun 2017 11:27:03 +0000 Subject: [PATCH] [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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@304708 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/AST/ASTContext.cpp | 5 ++--- test/Index/pipe-size.cl | 16 ++++++++++++++++ 2 files changed, 18 insertions(+), 3 deletions(-) create mode 100644 test/Index/pipe-size.cl diff --git a/lib/AST/ASTContext.cpp b/lib/AST/ASTContext.cpp index 34c4d2617e..62b19685c6 100644 --- a/lib/AST/ASTContext.cpp +++ b/lib/AST/ASTContext.cpp @@ -1939,9 +1939,8 @@ TypeInfo ASTContext::getTypeInfoImpl(const Type *T) const { break; case Type::Pipe: { - TypeInfo Info = getTypeInfo(cast(T)->getElementType()); - Width = Info.Width; - Align = Info.Align; + Width = Target->getPointerWidth(getTargetAddressSpace(LangAS::opencl_global)); + Align = Target->getPointerAlign(getTargetAddressSpace(LangAS::opencl_global)); } } diff --git a/test/Index/pipe-size.cl b/test/Index/pipe-size.cl new file mode 100644 index 0000000000..d07d7067da --- /dev/null +++ b/test/Index/pipe-size.cl @@ -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 +} -- 2.40.0