Merge from 'sycl' to 'sycl-web' (3 commits)

iclsrc · iclsrc · commit 2cfa1a66042b · 2021-10-07T03:32:25.000-07:00
diff --git a/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_simd_codegen.cpp b/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_simd_codegen.cpp
diff --git a/libclc/utils/libclc-remangler/LibclcRemangler.cpp b/libclc/utils/libclc-remangler/LibclcRemangler.cpp
@@ -367,6 +367,20 @@ class Remangler {
       addSub(typeNode);
       break;
     }
+    case Node::Kind::KBinaryFPType: {
+      if (remangleSub(typeNode, S))
+        return;
+
+      const BinaryFPType *BFPType = static_cast<const BinaryFPType *>(typeNode);
+      assert(BFPType->getDimension()->getKind() == Node::Kind::KNameType);
+      const NameType *dims =
+          static_cast<const NameType *>(BFPType->getDimension());
+
+      S << "DF";
+      S << dims->getName();
+      S << '_';
+      break;
+    }
     case Node::Kind::KVendorExtQualType: {
       if (remangleSub(typeNode, S))
         return;
diff --git a/llvm-spirv/lib/SPIRV/SPIRVToOCL.h b/llvm-spirv/lib/SPIRV/SPIRVToOCL.h
@@ -187,7 +187,7 @@ class SPIRVToOCLBase : public InstVisitor<SPIRVToOCLBase> {
 
   /// Transform __spirv_OpAtomicCompareExchange and
   /// __spirv_OpAtomicCompareExchangeWeak
-  virtual Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI, Op OC) = 0;
+  virtual Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI) = 0;
 
   /// Transform __spirv_OpAtomicIIncrement/OpAtomicIDecrement to:
   /// - OCL2.0: atomic_fetch_add_explicit/atomic_fetch_sub_explicit
@@ -293,10 +293,10 @@ class SPIRVToOCL12Base : public SPIRVToOCLBase {
   /// (bool)atomic_xchg(*ptr, 1)
   Instruction *visitCallSPIRVAtomicFlagTestAndSet(CallInst *CI);
 
-  /// Transform __spirv_OpAtomicCompareExchange and
-  /// __spirv_OpAtomicCompareExchangeWeak into atomic_cmpxchg. There is no
-  /// weak version of function in OpenCL 1.2
-  Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI, Op OC) override;
+  /// Transform __spirv_OpAtomicCompareExchange/Weak into atomic_cmpxchg
+  /// OpAtomicCompareExchangeWeak is not "weak" at all, but instead has
+  /// the same semantics as OpAtomicCompareExchange.
+  Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI) override;
 
   /// Conduct generic mutations for all atomic builtins
   CallInst *mutateCommonAtomicArguments(CallInst *CI, Op OC) override;
@@ -371,8 +371,10 @@ class SPIRVToOCL20Base : public SPIRVToOCLBase {
   std::string mapFPAtomicName(Op OC) override;
 
   /// Transform __spirv_OpAtomicCompareExchange/Weak into
-  /// compare_exchange_strong/weak_explicit
-  Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI, Op OC) override;
+  /// atomic_compare_exchange_strong_explicit
+  /// OpAtomicCompareExchangeWeak is not "weak" at all, but instead has
+  /// the same semantics as OpAtomicCompareExchange.
+  Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI) override;
 };
 
 class SPIRVToOCL20Pass : public llvm::PassInfoMixin<SPIRVToOCL20Pass>,
diff --git a/llvm-spirv/lib/SPIRV/SPIRVToOCL12.cpp b/llvm-spirv/lib/SPIRV/SPIRVToOCL12.cpp
@@ -205,8 +205,7 @@ SPIRVToOCL12Base::visitCallSPIRVAtomicFlagTestAndSet(CallInst *CI) {
       &Attrs);
 }
 
-Instruction *SPIRVToOCL12Base::visitCallSPIRVAtomicCmpExchg(CallInst *CI,
-                                                            Op OC) {
+Instruction *SPIRVToOCL12Base::visitCallSPIRVAtomicCmpExchg(CallInst *CI) {
   AttributeList Attrs = CI->getCalledFunction()->getAttributes();
   return mutateCallInstOCL(
       M, CI,
@@ -247,7 +246,7 @@ Instruction *SPIRVToOCL12Base::visitCallSPIRVAtomicBuiltin(CallInst *CI,
     break;
   case OpAtomicCompareExchange:
   case OpAtomicCompareExchangeWeak:
-    NewCI = visitCallSPIRVAtomicCmpExchg(CI, OC);
+    NewCI = visitCallSPIRVAtomicCmpExchg(CI);
     break;
   default:
     NewCI = mutateCommonAtomicArguments(CI, OC);
diff --git a/llvm-spirv/lib/SPIRV/SPIRVToOCL20.cpp b/llvm-spirv/lib/SPIRV/SPIRVToOCL20.cpp
@@ -159,7 +159,7 @@ Instruction *SPIRVToOCL20Base::visitCallSPIRVAtomicBuiltin(CallInst *CI,
     break;
   case OpAtomicCompareExchange:
   case OpAtomicCompareExchangeWeak:
-    NewCI = visitCallSPIRVAtomicCmpExchg(CIG, OC);
+    NewCI = visitCallSPIRVAtomicCmpExchg(CIG);
     break;
   default:
     NewCI = mutateAtomicName(CIG, OC);
@@ -232,8 +232,7 @@ CallInst *SPIRVToOCL20Base::mutateCommonAtomicArguments(CallInst *CI, Op OC) {
       &Attrs);
 }
 
-Instruction *SPIRVToOCL20Base::visitCallSPIRVAtomicCmpExchg(CallInst *CI,
-                                                            Op OC) {
+Instruction *SPIRVToOCL20Base::visitCallSPIRVAtomicCmpExchg(CallInst *CI) {
   assert(CI->getCalledFunction() && "Unexpected indirect call");
   AttributeList Attrs = CI->getCalledFunction()->getAttributes();
   Instruction *PInsertBefore = CI;
@@ -242,7 +241,7 @@ Instruction *SPIRVToOCL20Base::visitCallSPIRVAtomicCmpExchg(CallInst *CI,
       M, CI,
       [=](CallInst *, std::vector<Value *> &Args, Type *&RetTy) {
         // OpAtomicCompareExchange[Weak] semantics is different from
-        // atomic_compare_exchange_[strong|weak] semantics as well as
+        // atomic_compare_exchange_strong semantics as well as
         // arguments order.
         // OCL built-ins returns boolean value and stores a new/original
         // value by pointer passed as 2nd argument (aka expected) while SPIR-V
@@ -263,7 +262,9 @@ Instruction *SPIRVToOCL20Base::visitCallSPIRVAtomicCmpExchg(CallInst *CI,
         std::swap(Args[3], Args[4]);
         std::swap(Args[2], Args[3]);
         RetTy = Type::getInt1Ty(*Ctx);
-        return OCLSPIRVBuiltinMap::rmap(OC);
+        // OpAtomicCompareExchangeWeak is not "weak" at all, but instead has
+        // the same semantics as OpAtomicCompareExchange.
+        return "atomic_compare_exchange_strong_explicit";
       },
       [=](CallInst *CI) -> Instruction * {
         // OCL built-ins atomic_compare_exchange_[strong|weak] return boolean
diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
@@ -2332,11 +2332,13 @@ AnnotationDecorations tryParseAnnotationString(SPIRVModule *BM,
   RegexIterT DecorationsIt(AnnotatedCode.begin(), AnnotatedCode.end(),
                            DecorationRegex);
   RegexIterT DecorationsEnd;
-  // If we didn't find any FPGA specific annotations then add a UserSemantic
-  // decoration
+  // If we didn't find any FPGA specific annotations that are seprated as
+  // described above, then add a UserSemantic decoration
   if (DecorationsIt == DecorationsEnd)
     Decorates.MemoryAttributesVec.emplace_back(DecorationUserSemantic,
                                                AnnotatedCode.str());
+  bool IntelFPGADecorationFound = false;
+  DecorationsInfoVec IntelFPGADecorationsVec;
   for (; DecorationsIt != DecorationsEnd; ++DecorationsIt) {
     // Drop the braces surrounding the actual decoration
     const StringRef AnnotatedDecoration = AnnotatedCode.substr(
@@ -2346,12 +2348,14 @@ AnnotationDecorations tryParseAnnotationString(SPIRVModule *BM,
     StringRef Name = Split.first, ValueStr = Split.second;
     if (AllowFPGAMemAccesses) {
       if (Name == "params") {
+        IntelFPGADecorationFound = true;
         unsigned ParamsBitMask = 0;
         bool Failure = ValueStr.getAsInteger(10, ParamsBitMask);
         assert(!Failure && "Non-integer LSU controls value");
         (void)Failure;
         LSUControls.setWithBitMask(ParamsBitMask);
       } else if (Name == "cache-size") {
+        IntelFPGADecorationFound = true;
         if (!LSUControls.CacheSizeInfo.hasValue())
           continue;
         unsigned CacheSizeValue = 0;
@@ -2365,12 +2369,15 @@ AnnotationDecorations tryParseAnnotationString(SPIRVModule *BM,
       StringRef Annotation;
       Decoration Dec;
       if (Name == "pump") {
+        IntelFPGADecorationFound = true;
         Dec = llvm::StringSwitch<Decoration>(ValueStr)
                   .Case("1", DecorationSinglepumpINTEL)
                   .Case("2", DecorationDoublepumpINTEL);
       } else if (Name == "register") {
+        IntelFPGADecorationFound = true;
         Dec = DecorationRegisterINTEL;
       } else if (Name == "simple_dual_port") {
+        IntelFPGADecorationFound = true;
         Dec = DecorationSimpleDualPortINTEL;
       } else {
         Dec = llvm::StringSwitch<Decoration>(Name)
@@ -2384,14 +2391,25 @@ AnnotationDecorations tryParseAnnotationString(SPIRVModule *BM,
                   .Case("force_pow2_depth", DecorationForcePow2DepthINTEL)
                   .Default(DecorationUserSemantic);
         if (Dec == DecorationUserSemantic)
-          Annotation = AnnotatedCode;
-        else
+          Annotation = AnnotatedDecoration;
+        else {
+          IntelFPGADecorationFound = true;
           Annotation = ValueStr;
+        }
       }
-      Decorates.MemoryAttributesVec.emplace_back(Dec, Annotation.str());
+      IntelFPGADecorationsVec.emplace_back(Dec, Annotation.str());
     }
   }
+  // Even if there is an annotation string that is split in blocks like Intel
+  // FPGA annotation, it's not necessarily an FPGA annotation. Translate the
+  // whole string as UserSemantic decoration in this case.
+  if (IntelFPGADecorationFound)
+    Decorates.MemoryAttributesVec = IntelFPGADecorationsVec;
+  else
+    Decorates.MemoryAttributesVec.emplace_back(DecorationUserSemantic,
+                                               AnnotatedCode.str());
   Decorates.MemoryAccessesVec = LSUControls.getDecorationsFromCurrentState();
+
   return Decorates;
 }
 
diff --git a/llvm-spirv/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl b/llvm-spirv/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl
@@ -47,5 +47,5 @@ __kernel void testAtomicCompareExchangeExplicit_cl20(
 
 //CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected1.as, i32 %desired, i32 3, i32 0, i32 2)
 //CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected3.as, i32 %desired, i32 4, i32 0, i32 1)
-//CHECK-LLVM: call spir_func i1 @_Z37atomic_compare_exchange_weak_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected5.as, i32 %desired, i32 3, i32 0, i32 2)
-//CHECK-LLVM: call spir_func i1 @_Z37atomic_compare_exchange_weak_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected7.as, i32 %desired, i32 4, i32 0, i32 1)
+//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected5.as, i32 %desired, i32 3, i32 0, i32 2)
+//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected7.as, i32 %desired, i32 4, i32 0, i32 1)
diff --git a/llvm-spirv/test/transcoding/AtomicCompareExchange_cl20.ll b/llvm-spirv/test/transcoding/AtomicCompareExchange_cl20.ll
@@ -9,6 +9,13 @@ target triple = "spir-unknown-unknown"
 
 ; Check 'LLVM ==> SPIR-V ==> LLVM' conversion of atomic_compare_exchange_strong and atomic_compare_exchange_weak.
 
+; SPIR-V does not include an equivalent of atomic_compare_exchange_weak
+; (OpAtomicCompareExchangeWeak is identical to OpAtomicCompareExchange and
+; is deprecated, and removed in SPIR-V 1.4.)
+; This breaks the round trip for atomic_compare_exchange_weak, which must be
+; translated back to LLVM IR as atomic_compare_exchange_strong, regardless
+; of whether OpAtomicCompareExchange or OpAtomicCompareExchangeWeak is used.
+
 ; Function Attrs: nounwind
 
 ; CHECK-LABEL:   define spir_func void @test_strong
@@ -24,7 +31,7 @@ target triple = "spir-unknown-unknown"
 ; CHECK:         [[PTR_WEAK:%expected[0-9]*]] = alloca i32, align 4
 ; CHECK:         store i32 {{.*}}, i32* [[PTR_WEAK]]
 ; CHECK:         [[PTR_WEAK]].as = addrspacecast i32* [[PTR_WEAK]] to i32 addrspace(4)*
-; CHECK:         call spir_func i1 @_Z37atomic_compare_exchange_weak_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope{{.*}}(i32 {{.*}}* %object, i32 {{.*}}* [[PTR_WEAK]].as, i32 %desired, i32 5, i32 5, i32 2)
+; CHECK:         call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope{{.*}}(i32 {{.*}}* %object, i32 {{.*}}* [[PTR_WEAK]].as, i32 %desired, i32 5, i32 5, i32 2)
 ; CHECK:         load i32, i32 addrspace(4)* [[PTR_WEAK]].as
 
 ; Check that alloca for atomic_compare_exchange is being created in the entry block.
diff --git a/llvm-spirv/test/transcoding/IntelFPGAMemoryAttributes.ll b/llvm-spirv/test/transcoding/IntelFPGAMemoryAttributes.ll
@@ -249,6 +249,8 @@
 ; CHECK-SPIRV: Decorate {{[0-9]+}} ForcePow2DepthINTEL 0
 ; CHECK-SPIRV: Decorate {{[0-9]+}} ForcePow2DepthINTEL 1
 
+; CHECK-SPIRV-NOT: Decorate [[#]] UserSemantic "{memory:MLAB}{sizeinfo:4,500}"
+
 target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
 target triple = "spir"
 
diff --git a/llvm-spirv/test/transcoding/atomics.spt b/llvm-spirv/test/transcoding/atomics.spt
@@ -1,10 +1,10 @@
-119734787 65536 393230 32 0 
+119734787 65536 393230 33 0 
 2 Capability Addresses 
 2 Capability Kernel 
 5 ExtInstImport 1 "OpenCL.std"
 3 MemoryModel 2 2 
 8 EntryPoint 6 6 "test_atomic_global"
-13 String 31 "kernel_arg_type.test_atomic_global.int*,"
+13 String 32 "kernel_arg_type.test_atomic_global.int*,"
 3 Source 3 102000 
 3 Name 7 "dst"
 4 Name 8 "object"
@@ -18,7 +18,7 @@
 2 TypeVoid 2 
 4 TypePointer 4 5 3 
 7 TypeFunction 5 2 4 4 4 3 
-2 TypeBool 29 
+2 TypeBool 30 
 
 
 5 Function 2 6 0 5 
@@ -40,10 +40,11 @@
 7 AtomicXor 3 24 7 13 14 13 
 7 AtomicAnd 3 25 7 13 14 13 
 9 AtomicCompareExchange 3 26 7 13 14 14 13 17 
-7 AtomicExchange 3 27 7 13 14 13 
-6 AtomicLoad 3 28 8 13 14 
+9 AtomicCompareExchangeWeak 3 27 7 13 14 14 13 17 
+7 AtomicExchange 3 28 7 13 14 13 
+6 AtomicLoad 3 29 8 13 14 
 5 AtomicStore 8 13 14 10 
-6 AtomicFlagTestAndSet 29 30 8 13 14 
+6 AtomicFlagTestAndSet 30 31 8 13 14 
 4 AtomicFlagClear 8 13 14 
 1 Return 
 
@@ -66,7 +67,7 @@
 ; CHECK-LLVM-12: call spir_func i32 @_Z9atomic_orPU3AS1Vii(i32 addrspace(1)* %dst, i32 1) [[attr]]
 ; CHECK-LLVM-12: call spir_func i32 @_Z10atomic_xorPU3AS1Vii(i32 addrspace(1)* %dst, i32 1) [[attr]]
 ; CHECK-LLVM-12: call spir_func i32 @_Z10atomic_andPU3AS1Vii(i32 addrspace(1)* %dst, i32 1) [[attr]]
-; CHECK-LLVM-12: call spir_func i32 @_Z14atomic_cmpxchgPU3AS1Viii(i32 addrspace(1)* %dst, i32 0, i32 1) [[attr]]
+; CHECK-LLVM-12-COUNT-2: call spir_func i32 @_Z14atomic_cmpxchgPU3AS1Viii(i32 addrspace(1)* %dst, i32 0, i32 1) [[attr]]
 ; CHECK-LLVM-12: call spir_func i32 @_Z11atomic_xchgPU3AS1Vii(i32 addrspace(1)* %dst, i32 1) [[attr]]
 ; CHECK-LLVM-12: call spir_func i32 @_Z10atomic_addPU3AS1Vii(i32 addrspace(1)* %object, i32 0) [[attr]]
 ; CHECK-LLVM-12: call spir_func i32 @_Z11atomic_xchgPU3AS1Vii(i32 addrspace(1)* %object, i32 %desired) [[attr]]
@@ -89,11 +90,12 @@
 ; CHECK-LLVM-20: call spir_func i32 @_Z25atomic_fetch_xor_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(i32 addrspace(4)* %dst.as9, i32 1, i32 5, i32 2) [[attr]]
 ; CHECK-LLVM-20: call spir_func i32 @_Z25atomic_fetch_and_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(i32 addrspace(4)* %dst.as10, i32 1, i32 5, i32 2) [[attr]]
 ; CHECK-LLVM-20: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %dst.as11, i32 addrspace(4)* %expected12.as, i32 1, i32 5, i32 5, i32 2) [[attr]]
-; CHECK-LLVM-20: call spir_func i32 @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(i32 addrspace(4)* %dst.as13, i32 1, i32 5, i32 2) [[attr]]
+; CHECK-LLVM-20: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %dst.as13, i32 addrspace(4)* %expected14.as, i32 1, i32 5, i32 5, i32 2) [[attr]]
+; CHECK-LLVM-20: call spir_func i32 @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(i32 addrspace(4)* %dst.as15, i32 1, i32 5, i32 2) [[attr]]
 ; CHECK-LLVM-20: call spir_func i32 @_Z20atomic_load_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(i32 addrspace(4)* %object.as, i32 5, i32 2) [[attr]]
-; CHECK-LLVM-20: call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(i32 addrspace(4)* %object.as14, i32 %desired, i32 5, i32 2) [[attr]]
-; CHECK-LLVM-20: call spir_func i1 @_Z33atomic_flag_test_and_set_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(i32 addrspace(4)* %object.as15, i32 5, i32 2) [[attr]]
-; CHECK-LLVM-20: call spir_func void @_Z26atomic_flag_clear_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(i32 addrspace(4)* %object.as16, i32 5, i32 2) [[attr]]
+; CHECK-LLVM-20: call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(i32 addrspace(4)* %object.as16, i32 %desired, i32 5, i32 2) [[attr]]
+; CHECK-LLVM-20: call spir_func i1 @_Z33atomic_flag_test_and_set_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(i32 addrspace(4)* %object.as17, i32 5, i32 2) [[attr]]
+; CHECK-LLVM-20: call spir_func void @_Z26atomic_flag_clear_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(i32 addrspace(4)* %object.as18, i32 5, i32 2) [[attr]]
 
 ; RUN: llvm-spirv -r %t1.spv -o %t2.bc --spirv-target-env="SPV-IR"
 ; RUN: llvm-dis < %t2.bc | FileCheck %s --check-prefix=CHECK-LLVM-SPV-IR
diff --git a/llvm/include/llvm/Demangle/ItaniumDemangle.h b/llvm/include/llvm/Demangle/ItaniumDemangle.h
@@ -1104,6 +1104,8 @@ class BinaryFPType final : public Node {
 
   template<typename Fn> void match(Fn F) const { F(Dimension); }
 
+  const Node *getDimension() const { return Dimension; }
+
   void printLeft(OutputStream &S) const override {
     S += "_Float";
     Dimension->print(S);
diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp
@@ -1428,7 +1428,16 @@ pi_result hip_piDeviceGetInfo(pi_device device, pi_device_info param_name,
     return getInfo(param_value_size, param_value, param_value_size_ret, "");
   }
   case PI_DEVICE_INFO_EXTENSIONS: {
-    return getInfo(param_value_size, param_value, param_value_size_ret, "");
+    // TODO: Remove comment when HIP support native asserts.
+    // DEVICELIB_ASSERT extension is set so fallback assert
+    // postprocessing is NOP. HIP 4.3 docs indicate support for
+    // native asserts are in progress
+    std::string SupportedExtensions = "";
+    SupportedExtensions += PI_DEVICE_INFO_EXTENSION_DEVICELIB_ASSERT;
+    SupportedExtensions += " ";
+
+    return getInfo(param_value_size, param_value, param_value_size_ret,
+                   SupportedExtensions.c_str());
   }
   case PI_DEVICE_INFO_PRINTF_BUFFER_SIZE: {
     // The minimum value for the FULL profile is 1 MB.
diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp
@@ -419,6 +419,23 @@ class queue_impl {
   }
 
 private:
+  void finalizeHandler(handler &Handler, bool NeedSeparateDependencyMgmt,
+                       event &EventRet) {
+    if (MIsInorder) {
+      // Accessing and changing of an event isn't atomic operation.
+      // Hence, here is the lock for thread-safety.
+      std::lock_guard<std::mutex> Lock{MLastEventMtx};
+
+      if (NeedSeparateDependencyMgmt)
+        Handler.depends_on(MLastEvent);
+
+      EventRet = Handler.finalize();
+
+      MLastEvent = EventRet;
+    } else
+      EventRet = Handler.finalize();
+  }
+
   /// Performs command group submission to the queue.
   ///
   /// \param CGF is a function object containing command group.
@@ -437,9 +454,9 @@ class queue_impl {
     // Host and interop tasks, however, are not submitted to low-level runtimes
     // and require separate dependency management.
     const CG::CGTYPE Type = Handler.getType();
-    if (MIsInorder && (Type == CG::CGTYPE::CodeplayHostTask ||
-                       Type == CG::CGTYPE::CodeplayInteropTask))
-      Handler.depends_on(MLastEvent);
+    bool NeedSeparateDependencyMgmt =
+        MIsInorder && (Type == CG::CGTYPE::CodeplayHostTask ||
+                       Type == CG::CGTYPE::CodeplayInteropTask);
 
     event Event;
 
@@ -452,14 +469,11 @@ class queue_impl {
                             : ProgramManager::getInstance().kernelUsesAssert(
                                   Handler.MOSModuleHandle, Handler.MKernelName);
 
-      Event = Handler.finalize();
+      finalizeHandler(Handler, NeedSeparateDependencyMgmt, Event);
 
       (*PostProcess)(IsKernel, KernelUsesAssert, Event);
     } else
-      Event = Handler.finalize();
-
-    if (MIsInorder)
-      MLastEvent = Event;
+      finalizeHandler(Handler, NeedSeparateDependencyMgmt, Event);
 
     addEvent(Event);
     return Event;
@@ -521,7 +535,11 @@ class queue_impl {
   // Buffer to store assert failure descriptor
   buffer<AssertHappened, 1> MAssertHappenedBuffer;
 
+  // This event is employed for enhanced dependency tracking with in-order queue
+  // Access to the event should be guarded with MLastEventMtx
   event MLastEvent;
+  std::mutex MLastEventMtx;
+
   const bool MIsInorder;
 };
 
