[SYCL] Add DPCPP_ENABLE_OPAQUE_POINTERS option (#6378)

bader · web-flow · commit a364383a9ee0 · 2022-07-21T19:39:39.000+03:00
This option controls the default behavior of LLVM projects regarding
opaque pointers. The default value is OFF as most of the SYCL compiler
components are not ready to handle opaque pointers.

This change adds a new nightly job testing SYCL compiler with enabled
opaque pointers by default.
diff --git a/.github/workflows/sycl_nightly.yml b/.github/workflows/sycl_nightly.yml
@@ -19,6 +19,16 @@ jobs:
       build_configure_extra_args: ''
       lts_config: "ocl_gen9;ocl_x64"
 
+  ubuntu2004_opaque_pointers_build_test:
+    if: github.repository == 'intel/llvm'
+    uses: ./.github/workflows/sycl_linux_build_and_test.yml
+    with:
+      build_cache_root: "/__w/"
+      build_cache_suffix: opaque_pointers
+      build_artifact_suffix: opaque_pointers
+      build_configure_extra_args: "--hip --cuda --enable-esimd-emulator --cmake-opt=-DDPCPP_ENABLE_OPAQUE_POINTERS=TRUE"
+      lts_config: "ocl_gen9;ocl_x64"
+
   windows_default:
     name: Windows
     if: github.repository == 'intel/llvm'
diff --git a/clang/CMakeLists.txt b/clang/CMakeLists.txt
@@ -255,7 +255,11 @@ option(CLANG_DEFAULT_PIE_ON_LINUX "Default to -fPIE and -pie on linux-gnu" OFF)
 set(CLANG_ENABLE_OPAQUE_POINTERS "DEFAULT" CACHE STRING
     "Enable opaque pointers by default")
 if(CLANG_ENABLE_OPAQUE_POINTERS STREQUAL "DEFAULT")
-  set(CLANG_ENABLE_OPAQUE_POINTERS_INTERNAL ON)
+  if(DPCPP_ENABLE_OPAQUE_POINTERS)
+    set(CLANG_ENABLE_OPAQUE_POINTERS_INTERNAL ON)
+  else()
+    set(CLANG_ENABLE_OPAQUE_POINTERS_INTERNAL OFF)
+  endif()
 elseif(CLANG_ENABLE_OPAQUE_POINTERS)
   set(CLANG_ENABLE_OPAQUE_POINTERS_INTERNAL ON)
 else()
diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td
@@ -5855,9 +5855,7 @@ defm enable_noundef_analysis : BoolOption<"",
 defm opaque_pointers : BoolOption<"",
   "opaque-pointers",
   CodeGenOpts<"OpaquePointers">,
-  // FIXME: Upstream default is DefaultTrue, but llvm-spirv is broken after this
-  //        change. Override to DefaultFalse until fixed downstream.
-  DefaultFalse,
+  DefaultTrue,
   PosFlag<SetTrue, [], "Enable">,
   NegFlag<SetFalse, [], "Disable">,
   BothFlags<[], " opaque pointers">>;
diff --git a/clang/test/CodeGenOpenCL/sampled_image.cl b/clang/test/CodeGenOpenCL/sampled_image.cl
@@ -5,8 +5,8 @@ __attribute__((overloadable)) void my_read_image(__ocl_sampled_image1d_ro_t img)
 __attribute__((overloadable)) void my_read_image(__ocl_sampled_image2d_ro_t img);
 
 void test_read_image(__ocl_sampled_image1d_ro_t img_ro, __ocl_sampled_image2d_ro_t img_2d) {
-  // CHECK: call void @_Z13my_read_image32__spirv_SampledImage__image1d_ro(%spirv.SampledImage.image1d_ro_t* %{{[0-9]+}})
+  // CHECK: call void @_Z13my_read_image32__spirv_SampledImage__image1d_ro(ptr %{{[0-9]+}})
   my_read_image(img_ro);
-  // CHECK: call void @_Z13my_read_image32__spirv_SampledImage__image2d_ro(%spirv.SampledImage.image2d_ro_t* %{{[0-9]+}})
+  // CHECK: call void @_Z13my_read_image32__spirv_SampledImage__image2d_ro(ptr %{{[0-9]+}})
   my_read_image(img_2d);
 }
diff --git a/clang/test/CodeGenSPIRV/intel/is_valid_event.cl b/clang/test/CodeGenSPIRV/intel/is_valid_event.cl
diff --git a/clang/test/CodeGenSPIRV/intel/private-array-initialization.cl b/clang/test/CodeGenSPIRV/intel/private-array-initialization.cl
diff --git a/clang/test/CodeGenSYCL/const-wg-init.cpp b/clang/test/CodeGenSYCL/const-wg-init.cpp
@@ -14,8 +14,7 @@ int main() {
   kernel_parallel_for_work_group<class kernel>([=](cl::sycl::group<1> G) {
     const int WG_CONST = 10;
   });
-// CHECK:  store i32 10, i32 addrspace(4)* addrspacecast (i32 addrspace(3)* @{{.*}}WG_CONST{{.*}} to i32 addrspace(4)*)
-// CHECK:  %{{[0-9]+}} = call {}* @llvm.invariant.start.p4i8(i64 4, i8 addrspace(4)* addrspacecast (i8 addrspace(3)* bitcast (i32 addrspace(3)* @{{.*}}WG_CONST{{.*}} to i8 addrspace(3)*) to i8 addrspace(4)*))
-
+  // CHECK:  store i32 10, ptr addrspace(4) addrspacecast (ptr addrspace(3) @{{.*}}WG_CONST{{.*}} to ptr addrspace(4))
+  // CHECK:  %{{[0-9]+}} = call ptr @llvm.invariant.start.p4(i64 4, ptr addrspace(4) addrspacecast (ptr addrspace(3) @{{.*}}WG_CONST{{.*}} to ptr addrspace(4)))
   return 0;
 }
diff --git a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp
@@ -25,7 +25,7 @@ int main() {
 }
 
 // CHECK: define{{.*}} spir_kernel {{.*}}19use_kernel_for_test({{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{
-// CHECK: getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 0, !dbg [[LINE_A0:![0-9]+]]
+// CHECK: getelementptr inbounds %class.anon, {{.*}}, i32 0, i32 0, !dbg [[LINE_A0:![0-9]+]]
 // CHECK: call spir_func void {{.*}}6__init{{.*}} !dbg [[LINE_A0]]
 // CHECK: call spir_func void @_ZZ4mainENKUlvE_clEv{{.*}} !dbg [[LINE_B0:![0-9]+]]
 // CHECK: ret void, !dbg [[LINE_C0:![0-9]+]]
diff --git a/clang/test/CodeGenSYCL/kernel-annotation.cpp b/clang/test/CodeGenSYCL/kernel-annotation.cpp
@@ -25,12 +25,12 @@ class Functor {
 // CHECK-SPIR-NOT: annotations =
 
 // CHECK-NVPTX: nvvm.annotations = !{[[FIRST:![0-9]]], [[SECOND:![0-9]]]}
-// CHECK-NVPTX: [[FIRST]] = !{void ()* @_ZTS7Functor, !"kernel", i32 1}
-// CHECK-NVPTX: [[SECOND]] = !{void ()* @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_E5foo_2, !"kernel", i32 1}
+// CHECK-NVPTX: [[FIRST]] = !{ptr @_ZTS7Functor, !"kernel", i32 1}
+// CHECK-NVPTX: [[SECOND]] = !{ptr @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_E5foo_2, !"kernel", i32 1}
 
 // CHECK-AMDGCN: amdgcn.annotations = !{[[FIRST:![0-9]]], [[SECOND:![0-9]]]}
-// CHECK-AMDGCN: [[FIRST]] = !{void ()* @_ZTS7Functor, !"kernel", i32 1}
-// CHECK-AMDGCN: [[SECOND]] = !{void ()* @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_E5foo_2, !"kernel", i32 1}
+// CHECK-AMDGCN: [[FIRST]] = !{ptr @_ZTS7Functor, !"kernel", i32 1}
+// CHECK-AMDGCN: [[SECOND]] = !{ptr @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_E5foo_2, !"kernel", i32 1}
 
 int main() {
   sycl::queue q;
diff --git a/clang/test/Driver/clang-offload-deps.c b/clang/test/Driver/clang-offload-deps.c
@@ -1,5 +1,8 @@
 // REQUIRES: x86-registered-target
 
+// FIXME: enable opaque pointers support
+// UNSUPPORTED: enable-opaque-pointers
+
 //
 // Check help message.
 //
diff --git a/clang/test/Driver/clang-offload-wrapper.c b/clang/test/Driver/clang-offload-wrapper.c
@@ -1,5 +1,8 @@
 // REQUIRES: x86-registered-target
 
+// FIXME: enable opaque pointers support
+// UNSUPPORTED: enable-opaque-pointers
+
 //
 // Check help message.
 //
diff --git a/clang/test/Driver/hip-options.hip b/clang/test/Driver/hip-options.hip
@@ -79,7 +79,7 @@
 // HIPTHINLTO-NOT: "-cc1"{{.*}} "-triple" "x86_64-unknown-linux-gnu" {{.*}} "-flto-unit"
 // HIPTHINLTO: "-cc1"{{.*}} "-triple" "amdgcn-amd-amdhsa" {{.*}} "-flto=thin" "-flto-unit" {{.*}} "-fwhole-program-vtables"
 // HIPTHINLTO-NOT: "-cc1"{{.*}} "-triple" "x86_64-unknown-linux-gnu" {{.*}} "-flto-unit"
-// HIPTHINLTO: lld{{.*}}"-plugin-opt=mcpu=gfx906" "-plugin-opt=thinlto" "-plugin-opt=-force-import-all"
+// HIPTHINLTO: lld{{.*}}"-plugin-opt=mcpu=gfx906"{{.*}}"-plugin-opt=thinlto"{{.*}}"-plugin-opt=-force-import-all"
 
 // Check that -flto=thin is handled correctly, particularly with -fwhole-program-vtables.
 //
diff --git a/lld/ELF/Driver.cpp b/lld/ELF/Driver.cpp
@@ -1123,8 +1123,13 @@ static void readConfigs(opt::InputArgList &args) {
   config->nostdlib = args.hasArg(OPT_nostdlib);
   config->oFormatBinary = isOutputFormatBinary(args);
   config->omagic = args.hasFlag(OPT_omagic, OPT_no_omagic, false);
+#if ENABLE_OPAQUE_POINTERS
+  config->opaquePointers = args.hasFlag(
+      OPT_plugin_opt_opaque_pointers, OPT_plugin_opt_no_opaque_pointers, true);
+#else
   config->opaquePointers = args.hasFlag(
       OPT_plugin_opt_opaque_pointers, OPT_plugin_opt_no_opaque_pointers, false);
+#endif
   config->optRemarksFilename = args.getLastArgValue(OPT_opt_remarks_filename);
   config->optStatsFilename = args.getLastArgValue(OPT_plugin_opt_stats_file);
 
diff --git a/llvm-spirv/tools/llvm-spirv/llvm-spirv.cpp b/llvm-spirv/tools/llvm-spirv/llvm-spirv.cpp
@@ -153,8 +153,15 @@ static cl::opt<bool>
     SPIRVToolsDis("spirv-tools-dis", cl::init(false),
                   cl::desc("Emit textual assembly using SPIRV-Tools"));
 
+#if ENABLE_OPAQUE_POINTERS
+constexpr static bool SPIRVOpaquePointersDefault = true;
+#else
+constexpr static bool SPIRVOpaquePointersDefault = false;
+#endif
+
 static cl::opt<bool>
-    EmitOpaquePointers("emit-opaque-pointers", cl::init(false),
+    EmitOpaquePointers("emit-opaque-pointers",
+                       cl::init(SPIRVOpaquePointersDefault),
                        cl::desc("Emit opaque instead of typed LLVM pointers "
                                 "for the translation from SPIR-V."),
                        cl::Hidden);
diff --git a/llvm/CMakeLists.txt b/llvm/CMakeLists.txt
@@ -806,6 +806,12 @@ if(NOT LLVM_ENABLE_NEW_PASS_MANAGER)
                       " no longer supported.")
 endif()
 
+set(DPCPP_ENABLE_OPAQUE_POINTERS FALSE CACHE BOOL
+   "Enable opaque pointers by default.")
+if (DPCPP_ENABLE_OPAQUE_POINTERS)
+  add_definitions("-DENABLE_OPAQUE_POINTERS=1")
+endif(DPCPP_ENABLE_OPAQUE_POINTERS)
+
 include(HandleLLVMOptions)
 
 find_package(Python3 ${LLVM_MINIMUM_PYTHON_VERSION} REQUIRED
diff --git a/llvm/include/llvm/LTO/Config.h b/llvm/include/llvm/LTO/Config.h
@@ -179,7 +179,11 @@ struct Config {
 
   /// Use opaque pointer types. Used to call LLVMContext::setOpaquePointers
   /// unless already set by the `-opaque-pointers` commandline option.
+#if ENABLE_OPAQUE_POINTERS
+  bool OpaquePointers = true;
+#else
   bool OpaquePointers = false;
+#endif
 
   /// If this field is set, LTO will write input file paths and symbol
   /// resolutions here in llvm-lto2 command line flag format. This can be
diff --git a/llvm/lib/IR/LLVMContextImpl.cpp b/llvm/lib/IR/LLVMContextImpl.cpp
@@ -33,9 +33,15 @@
 
 using namespace llvm;
 
-static cl::opt<bool>
-    OpaquePointersCL("opaque-pointers", cl::desc("Use opaque pointers"),
-                     cl::init(false));
+#if ENABLE_OPAQUE_POINTERS
+static cl::opt<bool> OpaquePointersCL("opaque-pointers",
+                                      cl::desc("Use opaque pointers"),
+                                      cl::init(true));
+#else
+static cl::opt<bool> OpaquePointersCL("opaque-pointers",
+                                      cl::desc("Use opaque pointers"),
+                                      cl::init(false));
+#endif
 
 LLVMContextImpl::LLVMContextImpl(LLVMContext &C)
     : DiagHandler(std::make_unique<DiagnosticHandler>()),
diff --git a/llvm/test/tools/spirv-to-ir-wrapper/spirv-to-ir-wrapper.ll b/llvm/test/tools/spirv-to-ir-wrapper/spirv-to-ir-wrapper.ll
@@ -1,14 +1,16 @@
+; TODO: switch to opaque pointers once llvm-spirv tool is able to handle them.
+
 ; Check for passthrough abilities
-; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-as -opaque-pointers=0 %s -o %t.bc
 ; RUN: spirv-to-ir-wrapper %t.bc -o %t_1.bc -skip-unknown-input
-; RUN: llvm-dis %t_1.bc -o %t_1.ll
+; RUN: llvm-dis -opaque-pointers=0 %t_1.bc -o %t_1.ll
 ; RUN: FileCheck %s --input-file %t_1.ll
 
 ; Check for SPIR-V conversion
-; RUN: llvm-spirv %t.bc -o %t.spv
-; RUN: spirv-to-ir-wrapper %t.spv -o %t_2.bc
-; RUN: llvm-dis %t_2.bc -o %t_2.ll
-; RUN: FileCheck %s --input-file %t_2.ll
+; RUN: llvm-spirv -opaque-pointers=0 %t.bc -o %t.spv
+; RUN: spirv-to-ir-wrapper -llvm-spirv-opts "-emit-opaque-pointers=false" %t.spv -o %t_2.bc
+; RUN: llvm-dis -opaque-pointers=0 %t_2.bc -o %t_2.ll
+; RUNx: FileCheck %s --input-file %t_2.ll
 
 ; CHECK: target datalayout
 ; CHECK-NEXT: target triple = "spir-unknown-unknown"
diff --git a/llvm/tools/gold/gold-plugin.cpp b/llvm/tools/gold/gold-plugin.cpp
@@ -209,7 +209,11 @@ namespace options {
   // Asserts that LTO link has whole program visibility
   static bool whole_program_visibility = false;
   // Use opaque pointer types.
+#if ENABLE_OPAQUE_POINTERS
+  static bool opaque_pointers = true;
+#else
   static bool opaque_pointers = false;
+#endif
 
   // Optimization remarks filename, accepted passes and hotness options
   static std::string RemarksFilename;
diff --git a/llvm/tools/llvm-c-test/echo.cpp b/llvm/tools/llvm-c-test/echo.cpp
@@ -1383,7 +1383,11 @@ int llvm_echo(bool OpaquePointers) {
   size_t ModuleIdentLen;
   const char *ModuleName = LLVMGetModuleIdentifier(Src, &ModuleIdentLen);
   LLVMContextRef Ctx = LLVMContextCreate();
-  if (OpaquePointers)
+  // FIXME: Delete else branch once after the project is ready for opaque
+  // pointers. Original code assued that "default" value is "true".
+  if (!OpaquePointers)
+    LLVMContextSetOpaquePointers(Ctx, false);
+  else
     LLVMContextSetOpaquePointers(Ctx, true);
   LLVMModuleRef M = LLVMModuleCreateWithNameInContext(ModuleName, Ctx);
 
diff --git a/llvm/tools/llvm-lto2/llvm-lto2.cpp b/llvm/tools/llvm-lto2/llvm-lto2.cpp
@@ -168,9 +168,15 @@ static cl::opt<bool>
                  cl::desc("Run PGO context sensitive IR instrumentation"),
                  cl::Hidden);
 
+#if ENABLE_OPAQUE_POINTERS
+static cl::opt<bool> LtoOpaquePointers("lto-opaque-pointers",
+                                       cl::desc("Enable opaque pointer types"),
+                                       cl::init(true), cl::Hidden);
+#else
 static cl::opt<bool> LtoOpaquePointers("lto-opaque-pointers",
                                        cl::desc("Enable opaque pointer types"),
                                        cl::init(false), cl::Hidden);
+#endif
 
 static cl::opt<bool>
     DebugPassManager("debug-pass-manager", cl::Hidden,
diff --git a/sycl/test/check_device_code/id_queries_fit_int.cpp b/sycl/test/check_device_code/id_queries_fit_int.cpp
@@ -4,7 +4,7 @@
 
 using namespace sycl;
 
-// CHECK: define {{.*}}dso_local spir_func void @{{.*}}testItem{{.*}}(%"class.{{.*}}item"*{{.*}}%{{.*}})
+// CHECK: define {{.*}}dso_local spir_func void @{{.*}}testItem{{.*}}(%"class.{{.*}}item"{{.*}}%{{.*}})
 SYCL_EXTERNAL void testItem(item<1> TestItem) {
   // CHECK: call void @llvm.assume(i1 {{.*}})
   int Id = TestItem.get_id(0);
@@ -14,7 +14,7 @@ SYCL_EXTERNAL void testItem(item<1> TestItem) {
   int LinearId = TestItem.get_linear_id();
 }
 
-// CHECK: define {{.*}}dso_local spir_func void @{{.*}}testNDItem{{.*}}(%"class.{{.*}}nd_item"*{{.*}}%{{.*}})
+// CHECK: define {{.*}}dso_local spir_func void @{{.*}}testNDItem{{.*}}(%"class.{{.*}}nd_item"{{.*}}%{{.*}})
 SYCL_EXTERNAL void testNDItem(nd_item<1> TestNDItem) {
   // CHECK: call void @llvm.assume(i1 {{.*}})
   int GlobalId = TestNDItem.get_global_id(0);
diff --git a/sycl/test/check_device_code/kernel_arguments_as.cpp b/sycl/test/check_device_code/kernel_arguments_as.cpp
@@ -1,16 +1,16 @@
-// RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes
+// RUN: %clangxx -fsycl-device-only -Xclang -opaque-pointers -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes
 // RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE
-// RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__
+// RUN: %clangxx -fsycl-device-only -Xclang -opaque-pointers -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__
 // RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE
 //
 // Check the address space of the pointer in accessor class.
 //
 // CHECK: %struct.AccWrapper = type { %"class.cl::sycl::accessor[[NUMBER_SUFFIX:\.?[0-9]*]]" }
 // CHECK: %"class.cl::sycl::accessor[[NUMBER_SUFFIX]]" = type { %"class{{.*}}AccessorImplDevice", %[[UNION:.*]] }
-// CHECK-DISABLE: %[[UNION]] = type { i32 addrspace(1)* }
-// CHECK-ENABLE: %[[UNION]] = type { i32 addrspace(5)* }
+// CHECK-DISABLE: %[[UNION]] = type { ptr addrspace(1) }
+// CHECK-ENABLE: %[[UNION]] = type { ptr addrspace(5) }
 // CHECK: %struct.AccWrapper.{{[0-9]+}} = type { %"class.cl::sycl::accessor.[[NUM:[0-9]+]]" }
-// CHECK-NEXT: %"class.cl::sycl::accessor.[[NUM]]" = type { %"class{{.*}}LocalAccessorBaseDevice", i32 addrspace(3)* }
+// CHECK-NEXT: %"class.cl::sycl::accessor.[[NUM]]" = type { %"class{{.*}}LocalAccessorBaseDevice", ptr addrspace(3) }
 //
 // Check that kernel arguments doesn't have generic address space.
 //
diff --git a/sycl/test/check_device_code/no_offset.cpp b/sycl/test/check_device_code/no_offset.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl-device-only -fsycl-early-optimizations -fsycl-dead-args-optimization -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clangxx -fsycl-device-only -Xclang -opaque-pointers -fsycl-early-optimizations -fsycl-dead-args-optimization -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ -S -emit-llvm -o - %s | FileCheck %s
 
 #include <sycl/sycl.hpp>
 
@@ -15,7 +15,7 @@ int main() {
             sycl::ext::oneapi::accessor_property_list PL{sycl::ext::oneapi::no_offset, sycl::no_init};
             sycl::accessor acc_a(a, cgh, sycl::write_only, PL);
             sycl::accessor acc_b{b, cgh, sycl::read_only};
-            // CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_EUlT_E_(i32 addrspace(1)* {{.*}}, i32 addrspace(1)* noundef readonly {{.*}}, %"class.cl::sycl::id"* noundef byval(%"class.cl::sycl::id") align 8 {{.*}})
+            // CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_EUlT_E_(ptr addrspace(1) {{.*}}, ptr addrspace(1) noundef readonly {{.*}}, ptr noundef byval(%"class.cl::sycl::id") align 8 {{.*}})
             cgh.parallel_for(size, [=](auto i) {
                 acc_a[i] = acc_b[i];
             });
@@ -33,7 +33,7 @@ int main() {
         q.submit([&](sycl::handler &cgh) {
             sycl::accessor acc_a(a, cgh, sycl::write_only);
             sycl::accessor acc_b{b, cgh, sycl::read_only};
-            // CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_EUlT_E_(i32 addrspace(1)* {{.*}}, %"class.cl::sycl::id"* noundef byval(%"class.cl::sycl::id") align 8 {{.*}}, i32 addrspace(1)* noundef readonly {{.*}}, %"class.cl::sycl::id"* noundef byval(%"class.cl::sycl::id") align 8 {{.*}})
+            // CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_EUlT_E_(ptr addrspace(1) {{.*}}, ptr noundef byval(%"class.cl::sycl::id") align 8 {{.*}}, ptr addrspace(1) noundef readonly {{.*}}, ptr noundef byval(%"class.cl::sycl::id") align 8 {{.*}})
             cgh.parallel_for(size, [=](auto i) {
                 acc_a[i] = acc_b[i];
             });
diff --git a/sycl/test/extensions/sub_group_as.cpp b/sycl/test/extensions/sub_group_as.cpp

Original file line number	Diff line number	Diff line change
`@@ -25,7 +25,7 @@ int main() {`
`25`	`25`	`}`
`26`	`26`
`27`	`27`	`// CHECK: define{{.}} spir_kernel {{.}}19use_kernel_for_test({{.}}){{.}} !dbg [[KERNEL:![0-9]+]] {{.*}}{`
`28`		`-// CHECK: getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 0, !dbg [[LINE_A0:![0-9]+]]`
	`28`	`+// CHECK: getelementptr inbounds %class.anon, {{.*}}, i32 0, i32 0, !dbg [[LINE_A0:![0-9]+]]`
`29`	`29`	`// CHECK: call spir_func void {{.}}6__init{{.}} !dbg [[LINE_A0]]`
`30`	`30`	`// CHECK: call spir_func void @_ZZ4mainENKUlvE_clEv{{.*}} !dbg [[LINE_B0:![0-9]+]]`
`31`	`31`	`// CHECK: ret void, !dbg [[LINE_C0:![0-9]+]]`