[Backport to 15] [DebugInfo] Round up #elts of TypeVector when calcul…

…ating memory size (#2648) * [DebugInfo] Round up #elts of TypeVector when calculating memory size (#2504) Round up number of elements in a Vector to a power of 2 when calculating memory size. Memory size will be calculated as BaseType * bit_ceil(ComponentCount). The previous calculation already rounded 3 elements to 4 elements. Signed-off-by: Lu, John <[email protected]> * Avoid requiring C++20 (#2600) bit_ceil requires C++20. Use llvm::bit_ceil to avoid this requirement. Signed-off-by: Lu, John <[email protected]> * Remove unecessary include that require c++20 (#2628) * Add own bit_ceil implementation. Signed-off-by: Marcos Maronas <[email protected]> Signed-off-by: Lu, John <[email protected]>
KhronosGroup · Jul 18, 2024 · a091a82 · a091a82
1 parent 0303992
commit a091a82
Show file tree

Hide file tree

Showing 4 changed files with 71 additions and 9 deletions.
diff --git a/lib/SPIRV/SPIRVInternal.h b/lib/SPIRV/SPIRVInternal.h
@@ -1098,6 +1098,15 @@ bool postProcessBuiltinsWithArrayArguments(Module *M, bool IsCpp = false);
 
 template <typename T>
 MetadataAsValue *map2MDString(LLVMContext &C, SPIRVValue *V);
+
+/// Returns the smallest integral power of two no smaller than Value if Value is
+/// nonzero.  Returns 1 otherwise.
+///
+/// Ex. bitCeil(5) == 8.
+///
+/// The return value is undefined if the input is larger than the largest power
+/// of two representable in SPIRVWord.
+[[nodiscard]] SPIRVWord bitCeil(SPIRVWord Value);
 } // namespace SPIRV
 
 #endif // SPIRV_SPIRVINTERNAL_H
diff --git a/lib/SPIRV/SPIRVToLLVMDbgTran.cpp b/lib/SPIRV/SPIRVToLLVMDbgTran.cpp
@@ -495,15 +495,13 @@ SPIRVToLLVMDbgTran::transTypeVector(const SPIRVExtInst *DebugInst) {
       transNonNullDebugType(BM->get<SPIRVExtInst>(Ops[BaseTypeIdx]));
   SPIRVWord Count = getConstantValueOrLiteral(Ops, ComponentCountIdx,
                                               DebugInst->getExtSetKind());
-  // FIXME: The current design of SPIR-V Debug Info doesn't provide a field
-  // for the derived memory size. Meanwhile, OpenCL/SYCL 3-element vectors
-  // occupy the same amount of memory as 4-element vectors, hence the simple
-  // elem_count * elem_size formula fails in this edge case.
-  // Once the specification is updated to reflect the whole memory block's
-  // size in SPIR-V, the calculations below must be replaced with a simple
-  // translation of the known size.
-  SPIRVWord SizeCount = (Count == 3) ? 4 : Count;
-  uint64_t Size = getDerivedSizeInBits(BaseTy) * SizeCount;
+  // Round up to a power of two.
+  // OpenCL/SYCL 3-element vectors
+  // occupy the same amount of memory as 4-element vectors
+  // Clang rounds up the memory size of vectors to a power of 2.
+  // Vulkan allows vec3 to have a memory size of 12, but in RenderDoc memory
+  // size is not derived from debug info.
+  const uint64_t Size = getDerivedSizeInBits(BaseTy) * bitCeil(Count);
 
   SmallVector<llvm::Metadata *, 8> Subscripts;
   Subscripts.push_back(getDIBuilder(DebugInst).getOrCreateSubrange(0, Count));

diff --git a/lib/SPIRV/SPIRVUtil.cpp b/lib/SPIRV/SPIRVUtil.cpp
@@ -2492,4 +2492,20 @@ template MetadataAsValue *
 map2MDString<internal::InternalJointMatrixLayout>(LLVMContext &, SPIRVValue *);
 template MetadataAsValue *map2MDString<spv::Scope>(LLVMContext &, SPIRVValue *);
 
+[[nodiscard]] SPIRVWord bitCeil(SPIRVWord Value) {
+  if (Value < 2)
+    return 1;
+
+  // If Value is already a power of 2, just return it.
+  if ((Value & (Value - 1)) == 0)
+    return Value;
+
+  Value--;
+  for (SPIRVWord Shift = std::numeric_limits<SPIRVWord>::digits >> 1; Shift;
+       Shift >>= 1) {
+    Value |= Value >> Shift;
+  }
+  return ++Value;
+}
+
 } // namespace SPIRV
diff --git a/test/DebugInfo/DebugInfoVector.ll b/test/DebugInfo/DebugInfoVector.ll
@@ -0,0 +1,39 @@
+; Ensure that a vector type's memory size is calculated as bit_ceil(# elements) * element size
+; even if the (# elements) is not 3.
+;
+; This test was derived from DebugInfo/X86/sycl-vec-3.ll.
+
+; RUN: llvm-as < %s -o %t.bc
+
+; RUN: llvm-spirv %t.bc -o %t.spv -spirv-ext=+SPV_INTEL_vector_compute
+; RUN: llvm-spirv -r %t.spv -o %t.bc
+; RUN: llvm-dis %t.bc -o - | FileCheck %s --check-prefixes=CHECK
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+%"class.cl::sycl::vec" = type { <832 x i32> }
+@vector = dso_local addrspace(1) global %"class.cl::sycl::vec" zeroinitializer, align 16, !dbg !0
+
+!llvm.dbg.cu = !{!9}
+!llvm.module.flags = !{!10, !11, !12, !13, !14}
+
+!0 = !DIGlobalVariableExpression(var: !1, expr: !DIExpression())
+!1 = distinct !DIGlobalVariable(name: "vector", scope: null, file: !2, line: 3, type: !3, isLocal: false, isDefinition: true)
+!2 = !DIFile(filename: "sycl-vec-3.cpp", directory: "/tmp")
+; CHECK: !DICompositeType(tag: DW_TAG_array_type, baseType: ![[BASE_TY:[0-9]+]],{{.*}} size: 32768, flags: DIFlagVector, elements: ![[ELEMS:[0-9]+]])
+!3 = distinct !DICompositeType(tag: DW_TAG_array_type, baseType: !6, file: !2, line: 3, size: 32768, flags: DIFlagVector, elements: !4, identifier: "_ZTSN2cl4sycl3vecIiLi3EEE")
+; CHECK-DAG: ![[ELEMS]] = !{![[ELEMS_RANGE:[0-9]+]]}
+!4 = !{!5}
+; CHECK-DAG: ![[ELEMS_RANGE]] = !DISubrange(count: 832{{.*}})
+!5 = !DISubrange(count: 832)
+; CHECK-DAG: ![[BASE_TY]] = !DIBasicType(name: "int", size: 32,{{.*}} encoding: DW_ATE_signed)
+!6 = !DIBasicType(name: "int", size: 32, align: 32, encoding: DW_ATE_signed)
+!7 = !{}
+!8 = !{!0}
+!9 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus, file: !2, producer: "clang version 13.0.0 (https://github.com/intel/llvm.git)", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, enums: !7, retainedTypes: !7, globals: !8, imports: !7)
+!10 = !{i32 7, !"Dwarf Version", i32 4}
+!11 = !{i32 2, !"Debug Info Version", i32 3}
+!12 = !{i32 1, !"wchar_size", i32 4}
+!13 = !{i32 7, !"uwtable", i32 1}
+!14 = !{i32 7, !"frame-pointer", i32 2}