[CBS] Fix runtime issues with opaque pointers #765

fodinabor · 2022-07-09T13:28:34Z

As a follow-up for #763 and should fix #764

This fixes a bunch of (SYCL compilation time) issues that are introduced with LLVM 15, due to defaulting to opaque pointers now.

The issues range from getPointerElementType being illegal on opaque ptrs to the annotations no longer using bitcast or getelementptr inline to specify their targets.

Should make hipSYCL ready for LLVM 15 and hopefully LLVM 16 as well, where according to https://llvm.org/docs/OpaquePointers.html#version-support, the non-opaque ptrs shall be fully removed.

cc @aaronmondal

This commit adds two missing headers whose contents have been moved in upstream Clang/LLVM. It also changes a call from `PointerType::getElementType()` to `Type::getPointerElementType()` as a temporary fix until opaque pointers are introduced. The former was already deprecated in Clang/LLVM 14 and has been removed in upstream.

aaronmondal · 2022-07-09T14:15:58Z

@fodinabor Cool ill try this out 😊

src/compiler/cbs/VectorShapeTransformer.cpp

src/compiler/cbs/SubCfgFormation.cpp

aaronmondal

The compile-time issues for the plugin itself, as well as the warnings appear to be gone now.

The plugin still does not want to run though. It seems like something changed with the way the passes are registered. I think this is somehow related to the Analysis Passes. Even with -mllvm -opt-bisect-limit=0 Clang segfaults with

llvm::PassInstrumentationAnalysis::Result&
    llvm::AnalysisManager<llvm::Module>::getResult<llvm::PassInstrumentationAnalysis (llvm::Module&)
        .../llvm/IR/PassManager.h:784:23

and

clang++: .../llvm/include/llvm/IR/PassManager.h:782:
  typename PassT::Result &llvm::AnalysisManager<llvm::Module>::getResult(IRUnitT &, ExtraArgTs...)
[IRUnitT = llvm::Module, ExtraArgTs = <>, PassT = llvm::PassInstrumentationAnalysis]:
  Assertion `AnalysisPasses.count(PassT::ID()) && "This analysis pass was not registered prior to being queried"' failed.

I guess that should be handled in a different issue though.

aaronmondal · 2022-07-12T18:11:36Z

@fodinabor My mistake. Warnings are not gone yet. Deprecation warnings persist at

VectorShapeTransformer.cpp:72:19
VectorShapeTransformer.cpp:504:26
VectorShapeTransformer.cpp:526:36
VectorShapeTransformer.cpp:546:34

all with "getPointerElementType is deprecated".

fodinabor · 2022-07-12T18:16:50Z

Jup, but I guess should keep those around for LLVM <= 15 in non-opaque pointer mode, and maybe the if constexpr should save us for LLVM 16, when it's removed? (I always forget how illegal if constexpr unreachable code may be for the program to still compile..)

fodinabor · 2022-07-13T11:35:40Z

Alright, I checked back, the if constexpr is only enough if getPointerElementTy was resolved in a delayed fashion due to depending on a template parameter. Since that's not the case, I decided to add some more macros proactively...

aaronmondal · 2022-07-13T15:00:59Z

Looking back at the opaqueptr migration instructions I found

As of April 2022 both LLVM and Clang have complete support for opaque pointers, and opaque pointers are enabled by default in LLVM and Clang.

I wonder whether it makes sense to completely drop the isOpaquePointerTy check. Maybe something like below is actually enough. Not sure whether this breaks older builds though (diffed against e72f3a6)

diff --git a/src/compiler/cbs/VectorShapeTransformer.cpp b/src/compiler/cbs/VectorShapeTransformer.cpp
index 0eaaca5..667ca6f 100644
--- a/src/compiler/cbs/VectorShapeTransformer.cpp
+++ b/src/compiler/cbs/VectorShapeTransformer.cpp
@@ -25,12 +25,8 @@
 using namespace hipsycl::compiler;
 using namespace llvm;
 
-#if LLVM_VERSION_MAJOR < 13
-#define IS_OPAQUE(pointer) constexpr(false)
-#elif LLVM_VERSION_MAJOR < 16
-#define IS_OPAQUE(pointer) (pointer->isOpaquePointerTy())
-#else
-#define IS_OPAQUE(pointer) constexpr(true)
+#if LLVM_VERSION_MAJOR < 14
+#define HIPSYCL_NO_OPAQUE_PTR
 #endif
 
 hipsycl::compiler::VectorShape GenericTransfer(hipsycl::compiler::VectorShape a) {
@@ -67,9 +63,11 @@ static Type *getElementType(Type *Ty) {
     return VecTy->getElementType();
   }
   if (auto PtrTy = dyn_cast<PointerType>(Ty)) {
-    if IS_OPAQUE (PtrTy)
-      return nullptr;
+#ifdef HIPSYCL_NO_OPAQUE_PTR
     return PtrTy->getPointerElementType();
+#else
+    return nullptr;
+#endif
   }
   if (auto ArrTy = dyn_cast<ArrayType>(Ty)) {
     return ArrTy->getElementType();
@@ -209,10 +207,11 @@ VectorShape VectorShapeTransformer::computeIdealShapeForInst(const Instruction &
         result.setAlignment(newalignment);
       } else {
         // NOTE: If indexShape is varying, this still reasons about alignment
-        if IS_OPAQUE (subT)
-          subT = gep.getSourceElementType();
-        else
-          subT = getElementType(subT);
+#ifdef HIPSYCL_NO_OPAQUE_PTR
+        subT = getElementType(subT);
+#else
+        subT = gep.getSourceElementType();
+#endif
         assert(subT && "Unknown LLVM element type .. IR type system change?");
 
         const size_t typeSizeInBytes = (size_t)layout.getTypeStoreSize(subT);
@@ -498,10 +497,11 @@ bool returnsVoidPtr(const Instruction &inst) {
     return false;
   if (!inst.getType()->isPointerTy())
     return false;
-  if IS_OPAQUE (inst.getType())
-    return true;
-
+#ifdef HIPSYCL_NO_OPAQUE_PTR
   return inst.getType()->getPointerElementType()->isIntegerTy(8);
+#else
+  return true;
+#endif
 }
 
 VectorShape VectorShapeTransformer::computeShapeForCastInst(const CastInst &castI) const {
@@ -520,9 +520,7 @@ VectorShape VectorShapeTransformer::computeShapeForCastInst(const CastInst &cast
   switch (castI.getOpcode()) {
   case Instruction::IntToPtr: {
     PointerType *DestType = cast<PointerType>(castI.getDestTy());
-    if IS_OPAQUE (DestType)
-      return VectorShape::strided(castOpStride, 1);
-
+#ifdef HIPSYCL_NO_UNIQUE_PTR
     Type *DestPointsTo = DestType->getPointerElementType();
 
     // FIXME: void pointers are char pointers (i8*), but what
@@ -536,18 +534,22 @@ VectorShape VectorShapeTransformer::computeShapeForCastInst(const CastInst &cast
       return VectorShape::varying();
 
     return VectorShape::strided(castOpStride / typeSize, 1);
+#else
+    return VectorShape::strided(castOpStride, 1);
+#endif
   }
 
   case Instruction::PtrToInt: {
     Type *SrcType = castI.getSrcTy();
-    if IS_OPAQUE (SrcType)
-      return VectorShape::strided(castOpStride, aligned);
-
+#ifdef HIPSYCL_NO_UNIQUE_PTR
     Type *SrcElemType = SrcType->getPointerElementType();
 
     unsigned typeSize = (unsigned)layout.getTypeStoreSize(SrcElemType);
 
     return VectorShape::strided(typeSize * castOpStride, aligned);
+#else
+    return VectorShape::strided(castOpStride, aligned);
+#endif
   }
 
     // Truncation reinterprets the stride modulo the target type width

fodinabor · 2022-07-13T15:07:36Z

I don't think we gain much by dropping that check and since in LLVM 15 there still is the flag -Xclang -no-opaque-pointers that can be used to switch Clang to continue emitting typed pointers, I'd prefer our LLVM 13-15 support to be able to deal with both situations.

aaronmondal · 2022-07-13T15:42:54Z

Yeah I agree. I guess as long as -no-opaque-pointers is around it needs to be handled. The current patches still trigger deprecation warnings though. I think it is desirable to have these gone.

aaronmondal · 2022-07-13T16:59:45Z

I don't think that this should be disabled by ignoring compiler flags. If upstream removes pointer types at an unexpected version number builds will break with missing symbol errors.

fodinabor · 2022-07-13T17:27:11Z

Since the plan for typed pointers removal from LLVM is quite well cut-out (LLVM 15 will switch to opaque ptrs by default but retain best-effort support for typed pointers, LLVM 16 will remove typed pointers) and we only use typed pointer functionality for LLVM_VERSION_MAJOR < 16, I believe we did our best to keep maximum compatibility with all LLVM versions. Deprecation warnings are so that one can take early enough action before something is removed. Since we proactively handled the scheduled removal of the function, I added the warning silencing.
This also should not unnecessarily break compatibility any further with downstream targets like ROCm's LLVM that release some random commit, since LLVM_VERSION_MAJOR will be set to 16 as soon as release/15.x is branched off, therefore disabling our usage of typed ptrs. Hence ROCm clang may release anytime, and we should support at least that change. For release/15.x the typed-ptr functions should in theory stay available.

Since LLVM is generally a fast-moving target, hipSYCL's policy generally did not include support for top-of-tree versions anyways. (which is e.g. why we don't have an LLVM main CI run configured, it would break far too often, partially due to changes that will be reverted until the next release anyway)

If you have any suggestions on how we can handle API differences (e.g. removal of functions) even more gracefully, they are very welcome, since LLVM downstream targets are a mess, and having a solution for that would be great (see also #709).

illuhad

I think I agree with @fodinabor. To me, this looks good. Obviously, handling fast-moving targets like LLVM is always difficult, especially in periods of technical transitions as currently. Given the plethora of different LLVM versions that people use hipSYCL with, there might not be a perfect solution.
Of course, suggestions for improvements are always welcome :)

aaronmondal · 2022-07-13T18:15:58Z

Ah I didn't know that opaque pointer removal points were already known. Then I'm fine with the current state of the PR 😊

al42and · 2022-08-03T12:36:40Z

src/compiler/cbs/VectorShapeTransformer.cpp

@@ -23,6 +25,17 @@
 using namespace hipsycl::compiler;
 using namespace llvm;

+#if LLVM_VERSION_MAJOR < 13


FYI, ROCm's clang strikes again: LLVM bundled with ROCm 4.5.2 lacks isOpaquePointerTy.

Aaaaaaa :/ Thanks for the info. Guess we need add a macro to force same treatment as clang < 13 for ROCm :/ FWIW, I believe hipSYCL works out-of-the-box with all functionality with the clang from ROCm 5.2.

FWIW, I believe hipSYCL works out-of-the-box with all functionality with the clang from ROCm 5.2.

Yes, can confirm that. Unfortunately, on our local test machine, 4.5 is the most stable.

Guess we need add a macro to force same treatment as clang < 13 for ROCm :/

I just used || defined(HIPSYCL_NO_DEVICE_MANGLER) here as a workaround, but perhaps it would be easier for users if there was a single USE_ROCM_CLANG flag that automatically enables the required workarounds for the current version (HIPSYCL_NO_DEVICE_MANGLER, #709, this one here, etc).

Unfortunately, ROCm/ROCm#1680 does not see any action 😒

aaronmondal and others added 2 commits July 8, 2022 18:01

[CBS] Fix runtime compat with opaque ptrs.

c6e30ff

aaronmondal reviewed Jul 9, 2022

View reviewed changes

src/compiler/cbs/VectorShapeTransformer.cpp Show resolved Hide resolved

aaronmondal reviewed Jul 9, 2022

View reviewed changes

src/compiler/cbs/SubCfgFormation.cpp Outdated Show resolved Hide resolved

Make IS_OPAQUE usage more explicitly an if & braces.

e72f3a6

aaronmondal approved these changes Jul 11, 2022

View reviewed changes

aaronmondal mentioned this pull request Jul 11, 2022

Clang plugin segfaulting with unregistered passes #766

Open

[CBS] Improve guarding against typed ptr removal.

65b5f09

[CBS] Silence deprecation warnings.

8d7787b

illuhad approved these changes Jul 13, 2022

View reviewed changes

illuhad merged commit 2d93799 into AdaptiveCpp:develop Jul 13, 2022

aaronmondal mentioned this pull request Jul 13, 2022

Make OpenSYCL live at HEAD #769

Open

al42and reviewed Aug 3, 2022

View reviewed changes

fodinabor deleted the feature/opaque-ptr branch August 4, 2022 06:43

illuhad mentioned this pull request Aug 7, 2022

cmake options to select the ROCm LLVM #797

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[CBS] Fix runtime issues with opaque pointers #765

[CBS] Fix runtime issues with opaque pointers #765

fodinabor commented Jul 9, 2022

aaronmondal commented Jul 9, 2022

aaronmondal left a comment •

edited

aaronmondal commented Jul 12, 2022

fodinabor commented Jul 12, 2022

fodinabor commented Jul 13, 2022

aaronmondal commented Jul 13, 2022

fodinabor commented Jul 13, 2022

aaronmondal commented Jul 13, 2022

aaronmondal commented Jul 13, 2022

fodinabor commented Jul 13, 2022

illuhad left a comment

aaronmondal commented Jul 13, 2022

al42and Aug 3, 2022

illuhad Aug 3, 2022

al42and Aug 3, 2022

[CBS] Fix runtime issues with opaque pointers #765

[CBS] Fix runtime issues with opaque pointers #765

Conversation

fodinabor commented Jul 9, 2022

aaronmondal commented Jul 9, 2022

aaronmondal left a comment • edited

Choose a reason for hiding this comment

aaronmondal commented Jul 12, 2022

fodinabor commented Jul 12, 2022

fodinabor commented Jul 13, 2022

aaronmondal commented Jul 13, 2022

fodinabor commented Jul 13, 2022

aaronmondal commented Jul 13, 2022

aaronmondal commented Jul 13, 2022

fodinabor commented Jul 13, 2022

illuhad left a comment

Choose a reason for hiding this comment

aaronmondal commented Jul 13, 2022

al42and Aug 3, 2022

Choose a reason for hiding this comment

illuhad Aug 3, 2022

Choose a reason for hiding this comment

al42and Aug 3, 2022

Choose a reason for hiding this comment

aaronmondal left a comment •

edited