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

Original commit:
KhronosGroup/SPIRV-LLVM-Translator@a6398f45674bb5d

Author: LU-JOHN

llvm-spirv/lib/SPIRV/SPIRVToLLVMDbgTran.cpp

@@ -50,6 +50,7 @@
 #include "llvm/IR/DebugProgramInstruction.h"
 #include "llvm/IR/IntrinsicInst.h"
 #include "llvm/IR/Module.h"
+#include <bit>
 
 using namespace std;
 using namespace SPIRVDebug::Operand;
@@ -484,15 +485,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.
+  uint64_t Size = getDerivedSizeInBits(BaseTy) * bit_ceil(Count);
 
   SmallVector<llvm::Metadata *, 8> Subscripts;
   Subscripts.push_back(getDIBuilder(DebugInst).getOrCreateSubrange(0, Count));

llvm-spirv/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}