From 85ab044fff82a0c588dbed83b96a89e66a4c127c Mon Sep 17 00:00:00 2001 From: Omar Ahmed Date: Thu, 8 Dec 2022 17:08:26 +0000 Subject: [PATCH 1/3] Add TransformGenericVolatileMemoryAccess pass --- include/clspv/FeatureMacro.h | 9 ++- include/clspv/Option.h | 3 +- lib/CMakeLists.txt | 1 + lib/Compiler.cpp | 15 ++++ lib/FeatureMacro.cpp | 1 - lib/PassRegistry.def | 1 + lib/Passes.h | 1 + lib/TransformGenericVolatileMemoryAccess.cpp | 73 ++++++++++++++++++++ lib/TransformGenericVolatileMemoryAccess.h | 30 ++++++++ test/generic_volatile_memory_access.cl | 36 ++++++++++ 10 files changed, 163 insertions(+), 7 deletions(-) create mode 100644 lib/TransformGenericVolatileMemoryAccess.cpp create mode 100644 lib/TransformGenericVolatileMemoryAccess.h create mode 100644 test/generic_volatile_memory_access.cl diff --git a/include/clspv/FeatureMacro.h b/include/clspv/FeatureMacro.h index 2b40f3a9d..cab8b7a46 100644 --- a/include/clspv/FeatureMacro.h +++ b/include/clspv/FeatureMacro.h @@ -30,7 +30,6 @@ enum class FeatureMacro { __opencl_c_subgroups, // following items are not supported __opencl_c_device_enqueue, - __opencl_c_generic_address_space, __opencl_c_pipes, __opencl_c_program_scope_global_variables, // following items are always enabled, but no point in complaining if they are @@ -44,7 +43,8 @@ enum class FeatureMacro { __opencl_c_read_write_images, __opencl_c_atomic_scope_device, __opencl_c_atomic_scope_all_devices, - __opencl_c_work_group_collective_functions + __opencl_c_work_group_collective_functions, + __opencl_c_generic_address_space, }; #define FeatureStr(f) std::make_pair(FeatureMacro::f, #f) @@ -53,6 +53,7 @@ constexpr std::array, 15> FeatureStr(__opencl_c_3d_image_writes), FeatureStr(__opencl_c_atomic_order_acq_rel), FeatureStr(__opencl_c_fp64), FeatureStr(__opencl_c_images), + FeatureStr(__opencl_c_generic_address_space), FeatureStr(__opencl_c_subgroups), // following items are always enabled by clang FeatureStr(__opencl_c_int64), @@ -62,9 +63,7 @@ constexpr std::array, 15> FeatureStr(__opencl_c_atomic_scope_all_devices), FeatureStr(__opencl_c_work_group_collective_functions), // following items cannot be enabled so are automatically disabled - FeatureStr(__opencl_c_device_enqueue), - FeatureStr(__opencl_c_generic_address_space), - FeatureStr(__opencl_c_pipes), + FeatureStr(__opencl_c_device_enqueue), FeatureStr(__opencl_c_pipes), FeatureStr(__opencl_c_program_scope_global_variables)}; #undef FeatureStr diff --git a/include/clspv/Option.h b/include/clspv/Option.h index ad355bcf2..75cc5dd79 100644 --- a/include/clspv/Option.h +++ b/include/clspv/Option.h @@ -168,7 +168,8 @@ SourceLanguage Language(); // Returns true when the source language makes use of the generic address space. inline bool LanguageUsesGenericAddressSpace() { return (Language() == SourceLanguage::OpenCL_CPP) || - ((Language() == SourceLanguage::OpenCL_C_20)); + (Language() == SourceLanguage::OpenCL_C_20) || + (Language() == SourceLanguage::OpenCL_C_30); } // Return the SPIR-V binary version diff --git a/lib/CMakeLists.txt b/lib/CMakeLists.txt index b68a3d01b..ea2232aac 100644 --- a/lib/CMakeLists.txt +++ b/lib/CMakeLists.txt @@ -65,6 +65,7 @@ add_library(clspv_passes OBJECT ${CMAKE_CURRENT_SOURCE_DIR}/SPIRVOp.cpp ${CMAKE_CURRENT_SOURCE_DIR}/SPIRVProducerPass.cpp ${CMAKE_CURRENT_SOURCE_DIR}/RemoveUnusedArguments.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/TransformGenericVolatileMemoryAccess.cpp ${CMAKE_CURRENT_SOURCE_DIR}/ReorderBasicBlocksPass.cpp ${CMAKE_CURRENT_SOURCE_DIR}/ReplaceLLVMIntrinsicsPass.cpp ${CMAKE_CURRENT_SOURCE_DIR}/ReplaceOpenCLBuiltinPass.cpp diff --git a/lib/Compiler.cpp b/lib/Compiler.cpp index 481198e93..363294a99 100644 --- a/lib/Compiler.cpp +++ b/lib/Compiler.cpp @@ -537,6 +537,21 @@ int RunPassPipeline(llvm::Module &M, llvm::raw_svector_ostream *binaryStream) { pm.addPass(clspv::FixupBuiltinsPass()); pm.addPass(clspv::ThreeElementVectorLoweringPass()); + // Lower longer vectors when requested. Note that this pass depends on + // ReplaceOpenCLBuiltinPass and expects DeadCodeEliminationPass to be run + // afterwards. + if (clspv::Option::LongVectorSupport()) { + pm.addPass(clspv::LongVectorLoweringPass()); + } + // Volatile information on loads and stores are not used inside + // SPIRVProducer pass, so it doesn't have any effect on the generated code, + // but they stop the mem2reg pass from optimizing them. CLSPV try to get rid + // of generic address spaces by inferring them and optimizing them through + // "InferAddressSpacePass" and "mem2reg" pass. However, volatile loads and + // stores will stop mem2reg. So, we remove volatile info on loads and stores + // so that we could use mem2reg optimization on them and remove generic + // address spaces. + pm.addPass(clspv::TransformGenericVolatileMemoryAccess()); // We need to run mem2reg and inst combine early because our // createInlineFuncWithPointerBitCastArgPass pass cannot handle the // pattern diff --git a/lib/FeatureMacro.cpp b/lib/FeatureMacro.cpp index 2750859db..60e604b85 100644 --- a/lib/FeatureMacro.cpp +++ b/lib/FeatureMacro.cpp @@ -23,7 +23,6 @@ namespace clspv { FeatureMacro FeatureMacroLookup(const std::string &name) { constexpr std::array NotSuppported{ FeatureMacro::__opencl_c_pipes, - FeatureMacro::__opencl_c_generic_address_space, FeatureMacro::__opencl_c_device_enqueue, FeatureMacro::__opencl_c_program_scope_global_variables}; diff --git a/lib/PassRegistry.def b/lib/PassRegistry.def index ef5adb2da..2132bc4d2 100644 --- a/lib/PassRegistry.def +++ b/lib/PassRegistry.def @@ -40,6 +40,7 @@ MODULE_PASS("native-math", clspv::NativeMathPass) MODULE_PASS("opencl-inliner", clspv::OpenCLInlinerPass) MODULE_PASS("physical-pointer-args", clspv::PhysicalPointerArgsPass) MODULE_PASS("remove-unused-arguments", clspv::RemoveUnusedArguments) +MODULE_PASS("transform-generic-volatile-memory-access", clspv::TransformGenericVolatileMemoryAccess) MODULE_PASS("replace-llvm-intrinsics", clspv::ReplaceLLVMIntrinsicsPass) MODULE_PASS("replace-opencl-builtin", clspv::ReplaceOpenCLBuiltinPass) MODULE_PASS("replace-pointer-bitcast", clspv::ReplacePointerBitcastPass) diff --git a/lib/Passes.h b/lib/Passes.h index 5ec07579f..cf5b46e02 100644 --- a/lib/Passes.h +++ b/lib/Passes.h @@ -55,6 +55,7 @@ #include "SplatSelectCondition.h" #include "StripFreezePass.h" #include "ThreeElementVectorLoweringPass.h" +#include "TransformGenericVolatileMemoryAccess.h" #include "UBOTypeTransformPass.h" #include "UndoBoolPass.h" #include "UndoByvalPass.h" diff --git a/lib/TransformGenericVolatileMemoryAccess.cpp b/lib/TransformGenericVolatileMemoryAccess.cpp new file mode 100644 index 000000000..a866546f5 --- /dev/null +++ b/lib/TransformGenericVolatileMemoryAccess.cpp @@ -0,0 +1,73 @@ +// Copyright 2019 The Clspv Authors. All rights reserved. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "llvm/IR/Function.h" +#include "llvm/IR/IRBuilder.h" +#include "llvm/IR/Instructions.h" +#include "llvm/IR/Module.h" +#include "llvm/Pass.h" + +#include "clspv/AddressSpace.h" + +#include "TransformGenericVolatileMemoryAccess.h" + +using namespace llvm; + +PreservedAnalyses +clspv::TransformGenericVolatileMemoryAccess::run(Module &M, + ModuleAnalysisManager &) { + PreservedAnalyses PA; + + SmallVector DeadInsts; + + for (auto &F : M.functions()) { + for (auto &BB : F) { + for (auto &I : BB) { + IRBuilder<> B(&I); + + if (auto *load = dyn_cast(&I)) { + if (load->isVolatile() && + getPointerAddressSpace(load->getPointerOperandType()) == + clspv::AddressSpace::Generic) { + auto NonVolatileLoad = + B.CreateLoad(load->getType(), load->getPointerOperand()); + load->replaceAllUsesWith(NonVolatileLoad); + DeadInsts.push_back(load); + } + } else if (auto *store = dyn_cast(&I)) { + if (store->isVolatile() && + getPointerAddressSpace(store->getPointerOperandType()) == + clspv::AddressSpace::Generic) { + B.CreateStore(store->getValueOperand(), store->getPointerOperand()); + DeadInsts.push_back(store); + } + } + } + } + } + + for (auto Inst : DeadInsts) { + Inst->eraseFromParent(); + } + + return PA; +} + +unsigned clspv::TransformGenericVolatileMemoryAccess::getPointerAddressSpace( + Type *PtrTy) const { + if (PtrTy->getNonOpaquePointerElementType()->isPointerTy()) { + return getPointerAddressSpace(PtrTy->getNonOpaquePointerElementType()); + } + return PtrTy->getPointerAddressSpace(); +} diff --git a/lib/TransformGenericVolatileMemoryAccess.h b/lib/TransformGenericVolatileMemoryAccess.h new file mode 100644 index 000000000..2f73ca3ac --- /dev/null +++ b/lib/TransformGenericVolatileMemoryAccess.h @@ -0,0 +1,30 @@ +// Copyright 2022 The Clspv Authors. All rights reserved. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "llvm/IR/Module.h" +#include "llvm/IR/PassManager.h" + +#ifndef _CLSPV_LIB_TRANSFORM_GENERIC_VOLATILE_MEMORY_ACCESS_PASS_H +#define _CLSPV_LIB_TRANSFORM_GENERIC_VOLATILE_MEMORY_ACCESS_PASS_H + +namespace clspv { +struct TransformGenericVolatileMemoryAccess + : llvm::PassInfoMixin { + llvm::PreservedAnalyses run(llvm::Module &M, llvm::ModuleAnalysisManager &); + + unsigned getPointerAddressSpace(llvm::Type *ptr) const; +}; +} // namespace clspv + +#endif // _CLSPV_LIB_TRANSFORM_GENERIC_VOLATILE_MEMORY_ACCESS_PASS_H diff --git a/test/generic_volatile_memory_access.cl b/test/generic_volatile_memory_access.cl new file mode 100644 index 000000000..53d695a74 --- /dev/null +++ b/test/generic_volatile_memory_access.cl @@ -0,0 +1,36 @@ +// RUN: clspv %target %s -o %t.spv -cl-std=CL2.0 -inline-entry-points +// RUN: spirv-dis -o %t2.spvasm %t.spv +// RUN: FileCheck %s < %t2.spvasm +// RUN: spirv-val --target-env vulkan1.0 %t.spv + +bool isFenceValid(cl_mem_fence_flags fence) { + if ((fence == 0) || (fence == CLK_GLOBAL_MEM_FENCE) || (fence == CLK_LOCAL_MEM_FENCE) + || (fence == (CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE))) + return true; + else + return false; +} + + +bool helperFunction(float *floatp, float val) { + if (!isFenceValid(get_fence(floatp))) + return false; + + if (*floatp != val) + return false; + + return true; +} + +__kernel void testKernel(__global uint *results) { + uint tid = get_global_id(0); + + __private float val; + val = 0.1f; + float * volatile ptr = &val; + + results[tid] = helperFunction(ptr, val); +} + +// CHECK-NOT: store volatile float addrspace(4)* {{.*}}, float addrspace(4)** {{.*}}, align 4 +// CHECK-NOT: {{.*}} = load volatile float addrspace(4)*, float addrspace(4)** {{.*}}, align 4 From e447ac8578ca5aab38fd56b900cb9a3e408f0641 Mon Sep 17 00:00:00 2001 From: Omar Ahmed Date: Mon, 12 Dec 2022 16:34:31 +0000 Subject: [PATCH 2/3] Add support for addrSpaceCast to longVectorLowering --- include/clspv/FeatureMacro.h | 4 ++-- lib/BuiltinsEnum.h | 3 +++ lib/BuiltinsMap.inc | 3 +++ lib/FeatureMacro.cpp | 3 +-- lib/LongVectorLoweringPass.cpp | 6 +++++- lib/ReplaceOpenCLBuiltinPass.cpp | 20 ++++++++++++++++++++ lib/ReplaceOpenCLBuiltinPass.h | 1 + lib/SPIRVProducerPass.cpp | 2 +- 8 files changed, 36 insertions(+), 6 deletions(-) diff --git a/include/clspv/FeatureMacro.h b/include/clspv/FeatureMacro.h index cab8b7a46..412f2e4ad 100644 --- a/include/clspv/FeatureMacro.h +++ b/include/clspv/FeatureMacro.h @@ -54,6 +54,7 @@ constexpr std::array, 15> FeatureStr(__opencl_c_atomic_order_acq_rel), FeatureStr(__opencl_c_fp64), FeatureStr(__opencl_c_images), FeatureStr(__opencl_c_generic_address_space), + FeatureStr(__opencl_c_program_scope_global_variables), FeatureStr(__opencl_c_subgroups), // following items are always enabled by clang FeatureStr(__opencl_c_int64), @@ -63,8 +64,7 @@ constexpr std::array, 15> FeatureStr(__opencl_c_atomic_scope_all_devices), FeatureStr(__opencl_c_work_group_collective_functions), // following items cannot be enabled so are automatically disabled - FeatureStr(__opencl_c_device_enqueue), FeatureStr(__opencl_c_pipes), - FeatureStr(__opencl_c_program_scope_global_variables)}; + FeatureStr(__opencl_c_device_enqueue), FeatureStr(__opencl_c_pipes)}; #undef FeatureStr FeatureMacro FeatureMacroLookup(const std::string &name); diff --git a/lib/BuiltinsEnum.h b/lib/BuiltinsEnum.h index 7c602480f..f51951f60 100644 --- a/lib/BuiltinsEnum.h +++ b/lib/BuiltinsEnum.h @@ -107,6 +107,9 @@ enum BuiltinType : unsigned int { kType_MemoryFence_Start, kGetFence, + kToGlobal, + kToLocal, + kToPrivate, kMemFence, kReadMemFence, kWriteMemFence, diff --git a/lib/BuiltinsMap.inc b/lib/BuiltinsMap.inc index e9e5637e8..bab183e0f 100644 --- a/lib/BuiltinsMap.inc +++ b/lib/BuiltinsMap.inc @@ -828,6 +828,9 @@ static std::unordered_map NotSuppported{ FeatureMacro::__opencl_c_pipes, - FeatureMacro::__opencl_c_device_enqueue, - FeatureMacro::__opencl_c_program_scope_global_variables}; + FeatureMacro::__opencl_c_device_enqueue}; const auto macro_itr = std::find_if( FeatureMacroList.begin(), FeatureMacroList.end(), diff --git a/lib/LongVectorLoweringPass.cpp b/lib/LongVectorLoweringPass.cpp index 1142d5de3..14d63be1a 100644 --- a/lib/LongVectorLoweringPass.cpp +++ b/lib/LongVectorLoweringPass.cpp @@ -949,7 +949,11 @@ Value *clspv::LongVectorLoweringPass::visitCastInst(CastInst &I) { V = B.CreateIntToPtr(EquivalentValue, EquivalentDestTy, I.getName()); break; } - + case Instruction::AddrSpaceCast: { + IRBuilder<> B(&I); + V = B.CreateAddrSpaceCast(EquivalentValue, EquivalentDestTy, I.getName()); + break; + } default: llvm_unreachable("Cast unsupported."); break; diff --git a/lib/ReplaceOpenCLBuiltinPass.cpp b/lib/ReplaceOpenCLBuiltinPass.cpp index ae893d452..8272932a8 100644 --- a/lib/ReplaceOpenCLBuiltinPass.cpp +++ b/lib/ReplaceOpenCLBuiltinPass.cpp @@ -76,6 +76,9 @@ std::set ReplaceOpenCLBuiltinPass::ReplaceableBuiltins = Builtins::kSubGroupBarrier, Builtins::kAtomicWorkItemFence, Builtins::kGetFence, + Builtins::kToGlobal, + Builtins::kToLocal, + Builtins::kToPrivate, Builtins::kMemFence, Builtins::kReadMemFence, Builtins::kWriteMemFence, @@ -528,6 +531,12 @@ bool ReplaceOpenCLBuiltinPass::runOnFunction(Function &F) { return replaceAtomicLoad(F); case Builtins::kGetFence: return replaceGetFence(F); + case Builtins::kToGlobal: + return replaceAddressSpaceQualifiers(F, AddressSpace::Global); + case Builtins::kToLocal: + return replaceAddressSpaceQualifiers(F, AddressSpace::Local); + case Builtins::kToPrivate: + return replaceAddressSpaceQualifiers(F, AddressSpace::Private); case Builtins::kAtomicInit: case Builtins::kAtomicStore: case Builtins::kAtomicStoreExplicit: @@ -3783,6 +3792,17 @@ bool ReplaceOpenCLBuiltinPass::replaceGetFence(Function &F) { }); } +bool ReplaceOpenCLBuiltinPass::replaceAddressSpaceQualifiers(Function &F, unsigned ToAddressSpace) { + return replaceCallsWithValue(F, [=](CallInst *Call) { + auto pointer = Call->getArgOperand(0); + // Clang emits an address space cast to the generic address space. Skip the + // cast and use the input directly. + + IRBuilder<> builder(Call); + return builder.CreateAddrSpaceCast(pointer, PointerType::get(pointer->getType()->getNonOpaquePointerElementType(), ToAddressSpace)); + }); +} + bool ReplaceOpenCLBuiltinPass::replaceExplicitAtomics( Function &F, spv::Op Op, spv::MemorySemanticsMask semantics) { return replaceCallsWithValue(F, [Op, semantics](CallInst *Call) { diff --git a/lib/ReplaceOpenCLBuiltinPass.h b/lib/ReplaceOpenCLBuiltinPass.h index 0574cbb6e..b1f9229e4 100644 --- a/lib/ReplaceOpenCLBuiltinPass.h +++ b/lib/ReplaceOpenCLBuiltinPass.h @@ -86,6 +86,7 @@ struct ReplaceOpenCLBuiltinPass bool replaceAtomics(llvm::Function &F, llvm::AtomicRMWInst::BinOp Op); bool replaceAtomicLoad(llvm::Function &F); bool replaceGetFence(llvm::Function &F); + bool replaceAddressSpaceQualifiers(llvm::Function &F, unsigned ToAddressSpace); bool replaceExplicitAtomics(llvm::Function &F, spv::Op Op, spv::MemorySemanticsMask semantics = spv::MemorySemanticsAcquireReleaseMask); diff --git a/lib/SPIRVProducerPass.cpp b/lib/SPIRVProducerPass.cpp index d7abb6840..159b1483c 100644 --- a/lib/SPIRVProducerPass.cpp +++ b/lib/SPIRVProducerPass.cpp @@ -2312,7 +2312,7 @@ SPIRVID SPIRVProducerPassImpl::getSPIRVConstant(Constant *C) { llvm_unreachable("Unhandled function declaration/definition"); } else if (auto *ConstExpr = dyn_cast(Cst)) { // If there is exactly one use we know where to insert the instruction - if (ConstExpr->getNumUses() == 1) { + if (ConstExpr->getNumUses() <= 2) { auto *User = *ConstExpr->user_begin(); auto *EquivInstr = ConstExpr->getAsInstruction(dyn_cast(User)); From 804332c4b2b2c4f96e8b99fa7c0b4b582b2dbb95 Mon Sep 17 00:00:00 2001 From: Omar Ahmed Date: Fri, 30 Dec 2022 17:41:52 +0000 Subject: [PATCH 3/3] Add global variables support --- lib/ClusterConstants.cpp | 21 +++++++++++++++++++-- lib/NormalizeGlobalVariable.cpp | 2 +- lib/SPIRVProducerPass.cpp | 6 +++--- 3 files changed, 23 insertions(+), 6 deletions(-) diff --git a/lib/ClusterConstants.cpp b/lib/ClusterConstants.cpp index 098a1859b..a1367d4a8 100644 --- a/lib/ClusterConstants.cpp +++ b/lib/ClusterConstants.cpp @@ -65,6 +65,20 @@ clspv::ClusterModuleScopeConstantVars::run(Module &M, ModuleAnalysisManager &) { initializers_alignment[GV.getInitializer()] = GV.getAlignment(); } } + else if (GV.getType()->getPointerAddressSpace() == clspv::AddressSpace::Global) { + if (GV.use_empty()) { + dead_global_constants.push_back(&GV); + } else { + global_constants.push_back(&GV); + if (GV.hasInitializer()) { + initializers.insert(GV.getInitializer()); + initializers_alignment[GV.getInitializer()] = GV.getAlignment(); + } else { + initializers.insert(Constant::getNullValue(GV.getType())); + initializers_alignment[Constant::getNullValue(GV.getType())] = GV.getAlignment(); + } + } + } } for (GlobalVariable *GV : dead_global_constants) { @@ -127,10 +141,10 @@ clspv::ClusterModuleScopeConstantVars::run(Module &M, ModuleAnalysisManager &) { Constant *clustered_initializer = ConstantStruct::get(type, initializers_as_vec); GlobalVariable *clustered_gv = new GlobalVariable( - M, type, true, GlobalValue::InternalLinkage, clustered_initializer, + M, type, false, GlobalValue::InternalLinkage, clustered_initializer, clspv::ClusteredConstantsVariableName(), nullptr, GlobalValue::ThreadLocalMode::NotThreadLocal, - clspv::AddressSpace::Constant); + clspv::AddressSpace::Global); assert(clustered_gv); clustered_gv->setAlignment(MaybeAlign(max_alignment)); @@ -148,7 +162,10 @@ clspv::ClusterModuleScopeConstantVars::run(Module &M, ModuleAnalysisManager &) { Instruction *gep = GetElementPtrInst::CreateInBounds( clustered_gv->getValueType(), clustered_gv, {zero, Builder.getInt32(index)}, "", inst); + // TODO: Handle cases for module constants where we have the user is a gep and a load user for the gep. This case happens when we have a module constant and a global variable in the same opencl module. user->replaceUsesOfWith(GV, gep); + } else if (dyn_cast(user)) { + // Will be handled with instructions. } else { errs() << "Don't know how to handle updating user of __constant: " << *user << "\n"; diff --git a/lib/NormalizeGlobalVariable.cpp b/lib/NormalizeGlobalVariable.cpp index 2dec2df01..74e3993eb 100644 --- a/lib/NormalizeGlobalVariable.cpp +++ b/lib/NormalizeGlobalVariable.cpp @@ -231,7 +231,7 @@ void NormalizeGlobalVariables(Module &M) { SmallVector globals; for (auto &GV : M.globals()) { if (GV.hasInitializer() && GV.getType()->getPointerAddressSpace() == - clspv::AddressSpace::Constant) { + clspv::AddressSpace::Global) { globals.push_back(&GV); } } diff --git a/lib/SPIRVProducerPass.cpp b/lib/SPIRVProducerPass.cpp index 159b1483c..99d26ce4d 100644 --- a/lib/SPIRVProducerPass.cpp +++ b/lib/SPIRVProducerPass.cpp @@ -991,7 +991,7 @@ void SPIRVProducerPassImpl::FindGlobalConstVars() { SmallVector GVList; SmallVector DeadGVList; for (GlobalVariable &GV : module->globals()) { - if (GV.getType()->getAddressSpace() == AddressSpace::Constant) { + if (GV.getType()->getAddressSpace() == AddressSpace::Global) { if (GV.use_empty()) { DeadGVList.push_back(&GV); } else { @@ -1260,7 +1260,7 @@ void SPIRVProducerPassImpl::FindTypesForResourceVars() { PointerType *PTy = cast(GV.getType()); const auto AS = PTy->getAddressSpace(); const bool module_scope_constant_external_init = - (AS == AddressSpace::Constant) && GV.hasInitializer(); + (AS == AddressSpace::Global || AS == AddressSpace::Constant) && GV.hasInitializer(); const spv::BuiltIn BuiltinType = GetBuiltin(GV.getName()); if (module_scope_constant_external_init && spv::BuiltInMax == BuiltinType) { @@ -2803,7 +2803,7 @@ void SPIRVProducerPassImpl::GenerateGlobalVar(GlobalVariable &GV) { const auto spvSC = GetStorageClass(AS); const bool module_scope_constant_external_init = - (AS == AddressSpace::Constant) && GV.hasInitializer() && + (AS == AddressSpace::Global || AS == AddressSpace::Constant) && GV.hasInitializer() && clspv::Option::ModuleConstantsInStorageBuffer(); if (GV.hasInitializer()) {