Add DPCPP_ENABLE_OPAQUE_POINTERS option

.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'

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()

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">>;

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);
 }

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;
 }

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]+]]

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;

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.
 //

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.
 //

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.
 //

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);
 

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);

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

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

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>()),

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"

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;

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);
 

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,

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);

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.
 //

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];
             });