[SYCL] Prepare FE for upcoming opaque ptr changes

This change adds -opaque-ptr explicitly to CodeGenSYCL tests which
were failing when it was enabled by default.   It then modifies
these tests to handle the equivalent opaque pointer changes in
the emitted IR (ptr instead of i8*, removal of bitcasts etc.)

This change also creates no_opaque_* equivalents for each of these
tests where -no-opaque-ptr is explicitly added to the tests; this
is to continue testing the current behavior.   When opaque pointer
becomes the default and is enabled, the no_opaque_* tests will be
deleted.

Author: premanandrao

clang/lib/CodeGen/CGExprScalar.cpp

@@ -1631,8 +1631,9 @@ ScalarExprEmitter::VisitSYCLUniqueStableIdExpr(SYCLUniqueStableIdExpr *E) {
   if (GlobalConstStr->getType()->getPointerAddressSpace() == ExprAS)
     return GlobalConstStr;
 
-  llvm::Type *EltTy = GlobalConstStr->getType()->getPointerElementType();
-  llvm::PointerType *NewPtrTy = llvm::PointerType::get(EltTy, ExprAS);
+  llvm::PointerType *PtrTy = cast<llvm::PointerType>(GlobalConstStr->getType());
+  llvm::PointerType *NewPtrTy =
+      llvm::PointerType::getWithSamePointeeType(PtrTy, ExprAS);
   return Builder.CreateAddrSpaceCast(GlobalConstStr, NewPtrTy,
                                      "usid_addr_cast");
 }

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -2702,7 +2702,7 @@ Address CodeGenFunction::EmitFieldAnnotations(const FieldDecl *D,
 
   // llvm.ptr.annotation intrinsic accepts a pointer to integer of any width -
   // don't perform bitcasts if value is integer
-  if (VTy->getPointerElementType()->isIntegerTy()) {
+  if (Addr.getElementType()->isIntegerTy()) {
     llvm::Function *F = CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, VTy);
 
     for (const auto *I : D->specific_attrs<AnnotateAttr>())
@@ -2744,7 +2744,7 @@ Address CodeGenFunction::EmitFieldSYCLAnnotations(const FieldDecl *D,
   auto *PTy = dyn_cast<llvm::PointerType>(VTy);
   unsigned AS = PTy ? PTy->getAddressSpace() : 0;
   llvm::Type *IntrType = VTy;
-  if (!VTy->getPointerElementType()->isIntegerTy())
+  if (!Addr.getElementType()->isIntegerTy())
     IntrType = llvm::PointerType::getWithSamePointeeType(CGM.Int8PtrTy, AS);
   llvm::Function *F =
       CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, IntrType);
@@ -2770,7 +2770,7 @@ Address CodeGenFunction::EmitIntelFPGAFieldAnnotations(SourceLocation Location,
   llvm::Type *VTy = V->getType();
   // llvm.ptr.annotation intrinsic accepts a pointer to integer of any width -
   // don't perform bitcasts if value is integer
-  if (VTy->getPointerElementType()->isIntegerTy()) {
+  if (Addr.getElementType()->isIntegerTy()) {
     llvm::Function *F =
         CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, VTy);
     V = EmitAnnotationCall(F, V, AnnotStr, Location);

clang/test/CodeGenSYCL/accessor_inheritance.cpp

@@ -1,4 +1,4 @@
-// RUN:  %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN:  %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -opaque-pointers -emit-llvm %s -o - | FileCheck %s
 #include "Inputs/sycl.hpp"
 
 struct Base {
@@ -26,72 +26,67 @@ int main() {
 // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE6kernel
 // CHECK-SAME: i32 noundef [[ARG_A:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: i32 noundef [[ARG_B:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: i8 addrspace(1)* noundef readonly align 1 [[ACC1_DATA:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: %[[RANGE_TYPE]]* noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE1:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: %[[RANGE_TYPE]]* noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE2:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: %[[ID_TYPE]]* noundef byval(%[[ID_TYPE]]) align 4 [[ACC1_ID:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: i8 addrspace(1)* noundef readonly align 1 [[ACC2_DATA:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: %[[RANGE_TYPE]]* noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE1:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: %[[RANGE_TYPE]]* noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE2:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: %[[ID_TYPE]]* noundef byval(%[[ID_TYPE]]) align 4 [[ACC2_ID:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef readonly align 1 [[ACC1_DATA:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE2:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr noundef byval(%[[ID_TYPE]]) align 4 [[ACC1_ID:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef readonly align 1 [[ACC2_DATA:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE2:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr noundef byval(%[[ID_TYPE]]) align 4 [[ACC2_ID:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: i32 noundef [[ARG_C:%[a-zA-Z0-9_]+]])
 
 // Allocas and addrspacecasts for kernel parameters
 // CHECK: [[ARG_A]].addr = alloca i32
 // CHECK: [[ARG_B]].addr = alloca i32
-// CHECK: [[ACC1_DATA]].addr = alloca i8 addrspace(1)
-// CHECK: [[ACC2_DATA]].addr = alloca i8 addrspace(1)*
+// CHECK: [[ACC1_DATA]].addr = alloca ptr addrspace(1)
+// CHECK: [[ACC2_DATA]].addr = alloca ptr addrspace(1)
 // CHECK: [[ARG_C]].addr = alloca i32
 // CHECK: [[KERNEL:%[a-zA-Z0-9_]+]] = alloca %class{{.*}}.anon
-// CHECK: [[ARG_A]].addr.ascast = addrspacecast i32* [[ARG_A]].addr to i32 addrspace(4)*
-// CHECK: [[ARG_B]].addr.ascast = addrspacecast i32* [[ARG_B]].addr to i32 addrspace(4)*
-// CHECK: [[ACC1_DATA]].addr.ascast = addrspacecast i8 addrspace(1)** [[ACC1_DATA]].addr to i8 addrspace(1)* addrspace(4)*
-// CHECK: [[ACC2_DATA]].addr.ascast = addrspacecast i8 addrspace(1)** [[ACC2_DATA]].addr to i8 addrspace(1)* addrspace(4)*
-// CHECK: [[ARG_C]].addr.ascast = addrspacecast i32* [[ARG_C]].addr to i32 addrspace(4)*
+// CHECK: [[ARG_A]].addr.ascast = addrspacecast ptr [[ARG_A]].addr to ptr addrspace(4)
+// CHECK: [[ARG_B]].addr.ascast = addrspacecast ptr [[ARG_B]].addr to ptr addrspace(4)
+// CHECK: [[ACC1_DATA]].addr.ascast = addrspacecast ptr [[ACC1_DATA]].addr to ptr addrspace(4)
+// CHECK: [[ACC2_DATA]].addr.ascast = addrspacecast ptr [[ACC2_DATA]].addr to ptr addrspace(4)
+// CHECK: [[ARG_C]].addr.ascast = addrspacecast ptr [[ARG_C]].addr to ptr addrspace(4)
 //
 // Lambda object alloca
-// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = addrspacecast %class{{.*}}.anon* [[KERNEL]] to %class{{.*}}.anon addrspace(4)*
+// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = addrspacecast ptr [[KERNEL]] to ptr addrspace(4)
 //
 // Kernel argument stores
-// CHECK: store i32 [[ARG_A]], i32 addrspace(4)* [[ARG_A]].addr.ascast
-// CHECK: store i32 [[ARG_B]], i32 addrspace(4)* [[ARG_B]].addr.ascast
-// CHECK: store i8 addrspace(1)* [[ACC1_DATA]], i8 addrspace(1)* addrspace(4)* [[ACC1_DATA]].addr.ascast
-// CHECK: store i8 addrspace(1)* [[ACC2_DATA]], i8 addrspace(1)* addrspace(4)* [[ACC2_DATA]].addr.ascast
-// CHECK: store i32 [[ARG_C]], i32 addrspace(4)* [[ARG_C]].addr.ascast
+// CHECK: store i32 [[ARG_A]], ptr addrspace(4) [[ARG_A]].addr.ascast
+// CHECK: store i32 [[ARG_B]], ptr addrspace(4) [[ARG_B]].addr.ascast
+// CHECK: store ptr addrspace(1) [[ACC1_DATA]], ptr addrspace(4) [[ACC1_DATA]].addr.ascast
+// CHECK: store ptr addrspace(1) [[ACC2_DATA]], ptr addrspace(4) [[ACC2_DATA]].addr.ascast
+// CHECK: store i32 [[ARG_C]], ptr addrspace(4) [[ARG_C]].addr.ascast
 //
 // Check A and B scalar fields initialization
-// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
-// CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to %struct{{.*}}Base addrspace(4)*
-// CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 0
-// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast
-// CHECK: store i32 [[ARG_A_LOAD]], i32 addrspace(4)* [[FIELD_A]]
-// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 1
-// CHECK: [[ARG_B_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_B]].addr.ascast
-// CHECK: store i32 [[ARG_B_LOAD]], i32 addrspace(4)* [[FIELD_B]]
+// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP]], i32 0, i32 0
+// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, ptr addrspace(4) [[ARG_A]].addr.ascast
+// CHECK: store i32 [[ARG_A_LOAD]], ptr addrspace(4) [[FIELD_A]]
+// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP]], i32 0, i32 1
+// CHECK: [[ARG_B_LOAD:%[a-zA-Z0-9_]+]] = load i32, ptr addrspace(4) [[ARG_B]].addr.ascast
+// CHECK: store i32 [[ARG_B_LOAD]], ptr addrspace(4) [[FIELD_B]]
 //
 // Check accessors initialization
-// CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 2
+// CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP]], i32 0, i32 2
 // Default constructor call
-// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC_FIELD]])
-// CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to i8 addrspace(4)*
-// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8 addrspace(4)* [[BITCAST1]], i64 20
-// CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8 addrspace(4)* [[GEP1]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
+// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) {{[^,]*}} [[ACC_FIELD]])
+// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, ptr addrspace(4) [[GEP]], i64 20
 // Default constructor call
-// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[BITCAST2]])
+// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev(ptr addrspace(4) {{[^,]*}} [[GEP1]])
 
 // CHECK C field initialization
-// CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured addrspace(4)* [[GEP]], i32 0, i32 2
-// CHECK: [[ARG_C_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_C]].addr.ascast
-// CHECK: store i32 [[ARG_C_LOAD]], i32 addrspace(4)* [[FIELD_C]]
+// CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, ptr addrspace(4) [[GEP]], i32 0, i32 2
+// CHECK: [[ARG_C_LOAD:%[a-zA-Z0-9_]+]] = load i32, ptr addrspace(4) [[ARG_C]].addr.ascast
+// CHECK: store i32 [[ARG_C_LOAD]], ptr addrspace(4) [[FIELD_C]]
 //
 // Check __init method calls
-// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
-// CHECK: [[BITCAST3:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP2]] to %struct{{.*}}Base addrspace(4)*
-// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST3]], i32 0, i32 2
-// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC1_DATA]].addr.ascast
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC1_FIELD]], i8 addrspace(1)* noundef [[ACC1_DATA_LOAD]]
+// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP2]], i32 0, i32 2
+// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[ACC1_DATA]].addr.ascast
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}} [[ACC1_FIELD]], ptr addrspace(1) noundef [[ACC1_DATA_LOAD]]
 //
-// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
-// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC2_DATA]].addr.ascast
-// CHECK: [[BITCAST4:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP3]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[BITCAST4]], i8 addrspace(1)* noundef [[ACC2_DATA_LOAD]]
+// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[ACC2_DATA]].addr.ascast
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}}, ptr addrspace(1) noundef [[ACC2_DATA_LOAD]]