[PATCH] D33597: [OpenCL] Fix pipe size in TypeInfo.

[Anastasia Stulova via Phabricator via cfe-commits]

Anastasia updated this revision to Diff 101031.
Anastasia edited reviewers, added: yaxunl; removed: echuraev.
Anastasia removed a subscriber: yaxunl.
Anastasia added a comment.

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?


https://reviews.llvm.org/D33597

Files:
  lib/AST/ASTContext.cpp
  test/Index/pipe-size.cl


Index: test/Index/pipe-size.cl
===================================================================
--- /dev/null
+++ test/Index/pipe-size.cl
@@ -0,0 +1,13 @@
+// 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
+__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
+}
Index: lib/AST/ASTContext.cpp
===================================================================
--- lib/AST/ASTContext.cpp
+++ lib/AST/ASTContext.cpp
@@ -1939,9 +1939,8 @@
   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));
   }
 
   }



Index: test/Index/pipe-size.cl
===================================================================
--- /dev/null
+++ test/Index/pipe-size.cl
@@ -0,0 +1,13 @@
+// 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
+__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
+}
Index: lib/AST/ASTContext.cpp
===================================================================
--- lib/AST/ASTContext.cpp
+++ lib/AST/ASTContext.cpp
@@ -1939,9 +1939,8 @@
   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));
   }
 
   }

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits