diff --git a/lib/SPIRV/SPIRVInternal.h b/lib/SPIRV/SPIRVInternal.h index 2db4658d31..d5448f5cd3 100644 --- a/lib/SPIRV/SPIRVInternal.h +++ b/lib/SPIRV/SPIRVInternal.h @@ -1128,6 +1128,15 @@ bool postProcessBuiltinsWithArrayArguments(Module *M, bool IsCpp = false); template 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 index d99b220f14..033ecc7a5b 100644 --- a/lib/SPIRV/SPIRVToLLVMDbgTran.cpp +++ b/lib/SPIRV/SPIRVToLLVMDbgTran.cpp @@ -489,15 +489,13 @@ SPIRVToLLVMDbgTran::transTypeVector(const SPIRVExtInst *DebugInst) { transNonNullDebugType(BM->get(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) * bitCeil(Count); SmallVector Subscripts; Subscripts.push_back(getDIBuilder(DebugInst).getOrCreateSubrange(0, Count)); diff --git a/lib/SPIRV/SPIRVUtil.cpp b/lib/SPIRV/SPIRVUtil.cpp index a893c68d73..fa62763802 100644 --- a/lib/SPIRV/SPIRVUtil.cpp +++ b/lib/SPIRV/SPIRVUtil.cpp @@ -2446,4 +2446,20 @@ template MetadataAsValue * map2MDString(LLVMContext &, SPIRVValue *); template MetadataAsValue *map2MDString(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::digits >> 1; Shift; + Shift >>= 1) { + Value |= Value >> Shift; + } + return ++Value; +} + } // namespace SPIRV diff --git a/test/DebugInfo/DebugInfoVector.ll b/test/DebugInfo/DebugInfoVector.ll new file mode 100644 index 0000000000..07f642d75e --- /dev/null +++ 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}