diff --git a/clang/include/clang/Basic/Targets/SPIR.h b/clang/include/clang/Basic/Targets/SPIR.h new file mode 100644 index 0000000000000..95ffda3b8f230 --- /dev/null +++ b/clang/include/clang/Basic/Targets/SPIR.h @@ -0,0 +1,19 @@ +//===---- SPIR.h - Declare SPIR and SPIR-V target interfaces ----*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +namespace clang { +namespace targets { + +// Used by both the SPIR and SPIR-V targets. Code of the generic address space +// for the target +constexpr unsigned SPIR_GENERIC_AS = 4u; + +} // namespace targets +} // namespace clang diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp index 14fcfeb9fdb54..fc4ed554401b2 100644 --- a/clang/lib/CodeGen/BackendUtil.cpp +++ b/clang/lib/CodeGen/BackendUtil.cpp @@ -11,6 +11,7 @@ #include "clang/Basic/Diagnostic.h" #include "clang/Basic/LangOptions.h" #include "clang/Basic/TargetOptions.h" +#include "clang/Basic/Targets/SPIR.h" #include "clang/Frontend/FrontendDiagnostic.h" #include "clang/Frontend/Utils.h" #include "clang/Lex/HeaderSearchOptions.h" @@ -122,6 +123,8 @@ class EmitAssemblyHelper { std::unique_ptr OS; + Triple TargetTriple; + TargetIRAnalysis getTargetIRAnalysis() const { if (TM) return TM->getTargetIRAnalysis(); @@ -174,7 +177,8 @@ class EmitAssemblyHelper { const LangOptions &LOpts, Module *M) : Diags(_Diags), HSOpts(HeaderSearchOpts), CodeGenOpts(CGOpts), TargetOpts(TOpts), LangOpts(LOpts), TheModule(M), - CodeGenerationTime("codegen", "Code Generation Time") {} + CodeGenerationTime("codegen", "Code Generation Time"), + TargetTriple(TheModule->getTargetTriple()) {} ~EmitAssemblyHelper() { if (CodeGenOpts.DisableFree) @@ -698,7 +702,6 @@ void EmitAssemblyHelper::CreatePasses(legacy::PassManager &MPM, // manually (and not via PMBuilder), since some passes (eg. InstrProfiling) // are inserted before PMBuilder ones - they'd get the default-constructed // TLI with an unknown target otherwise. - Triple TargetTriple(TheModule->getTargetTriple()); std::unique_ptr TLII( createTLII(TargetTriple, CodeGenOpts)); @@ -971,7 +974,6 @@ bool EmitAssemblyHelper::AddEmitPasses(legacy::PassManager &CodeGenPasses, raw_pwrite_stream &OS, raw_pwrite_stream *DwoOS) { // Add LibraryInfo. - llvm::Triple TargetTriple(TheModule->getTargetTriple()); std::unique_ptr TLII( createTLII(TargetTriple, CodeGenOpts)); CodeGenPasses.add(new TargetLibraryInfoWrapperPass(*TLII)); @@ -1050,7 +1052,7 @@ void EmitAssemblyHelper::EmitAssemblyWithLegacyPassManager( // -fsycl-instrument-device-code option was passed. This option can be // used only with spir triple. if (CodeGenOpts.SPIRITTAnnotations) { - assert(llvm::Triple(TheModule->getTargetTriple()).isSPIR() && + assert(TargetTriple.isSPIR() && "ITT annotations can only be added to a module with spir target"); PerModulePasses.add(createSPIRITTAnnotationsLegacyPass()); } @@ -1066,6 +1068,15 @@ void EmitAssemblyHelper::EmitAssemblyWithLegacyPassManager( PerModulePasses.add(createSYCLMutatePrintfAddrspaceLegacyPass()); } + if (!CodeGenOpts.DisableLLVMPasses) { + // Add the InferAddressSpaces pass for all the SPIR[V] targets + if (TargetTriple.isSPIR() || TargetTriple.isSPIRV()) { + // This function pass should run after inlining, so it is added to MPM + PerModulePasses.add( + createInferAddressSpacesPass(targets::SPIR_GENERIC_AS)); + } + } + switch (Action) { case Backend_EmitNothing: break; @@ -1086,10 +1097,8 @@ void EmitAssemblyHelper::EmitAssemblyWithLegacyPassManager( // Emit a module summary by default for Regular LTO except for ld64 // targets bool EmitLTOSummary = - (CodeGenOpts.PrepareForLTO && - !CodeGenOpts.DisableLLVMPasses && - llvm::Triple(TheModule->getTargetTriple()).getVendor() != - llvm::Triple::Apple); + (CodeGenOpts.PrepareForLTO && !CodeGenOpts.DisableLLVMPasses && + TargetTriple.getVendor() != llvm::Triple::Apple); if (EmitLTOSummary) { if (!TheModule->getModuleFlag("ThinLTO")) TheModule->addModuleFlag(Module::Error, "ThinLTO", uint32_t(0)); @@ -1370,7 +1379,6 @@ void EmitAssemblyHelper::RunOptimizationPipeline( // Register the target library analysis directly and give it a customized // preset TLI. - Triple TargetTriple(TheModule->getTargetTriple()); std::unique_ptr TLII( createTLII(TargetTriple, CodeGenOpts)); FAM.registerPass([&] { return TargetLibraryAnalysis(*TLII); }); @@ -1495,7 +1503,7 @@ void EmitAssemblyHelper::RunOptimizationPipeline( // -fsycl-instrument-device-code option was passed. This option can be used // only with spir triple. if (CodeGenOpts.SPIRITTAnnotations) { - assert(llvm::Triple(TheModule->getTargetTriple()).isSPIR() && + assert(TargetTriple.isSPIR() && "ITT annotations can only be added to a module with spir target"); MPM.addPass(SPIRITTAnnotationsPass()); } @@ -1535,8 +1543,7 @@ void EmitAssemblyHelper::RunOptimizationPipeline( // targets bool EmitLTOSummary = (CodeGenOpts.PrepareForLTO && !CodeGenOpts.DisableLLVMPasses && - llvm::Triple(TheModule->getTargetTriple()).getVendor() != - llvm::Triple::Apple); + TargetTriple.getVendor() != llvm::Triple::Apple); if (EmitLTOSummary) { if (!TheModule->getModuleFlag("ThinLTO")) TheModule->addModuleFlag(Module::Error, "ThinLTO", uint32_t(0)); diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 7200b51695d2d..670c492212af5 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -149,6 +149,7 @@ template struct id { template id(T... args) {} // fake constructor + int get(int) const { return 0; } // fake getter private: // Some fake field added to see using of id arguments in the // kernel wrapper @@ -206,6 +207,14 @@ class __attribute__((sycl_special_class)) accessor { void use(T... args) const {} _ImplT impl; + // Operator returns a reference to a temporary value but this is a fake + // operator for testings only. Operator is marked as 'const' to let us + // use it in kernels. + dataT &operator[](int) const { + const dataT Data{}; + return const_cast(Data); + } + private: void __init(__attribute__((opencl_global)) dataT *Ptr, range AccessRange, range MemRange, id Offset) {} diff --git a/clang/test/CodeGenSYCL/infer-address-spaces.cpp b/clang/test/CodeGenSYCL/infer-address-spaces.cpp new file mode 100644 index 0000000000000..0cb8de4d9b7bd --- /dev/null +++ b/clang/test/CodeGenSYCL/infer-address-spaces.cpp @@ -0,0 +1,32 @@ +// RUN: %clang_cc1 -O1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -emit-llvm %s -o - | FileCheck %s + +#include "sycl.hpp" + +#define BLOCK_SIZE 16 + +using namespace cl::sycl; + +int main() { + queue Q; + auto MatrixTemp = buffer{range<1>{BLOCK_SIZE * 512}}; + Q.submit([&](handler &cgh) { + auto temp_dst_acc = MatrixTemp.get_access(cgh); + auto temp_t = accessor(); + cgh.parallel_for(range<1>(BLOCK_SIZE), [=](id<1> id) { + int index = 64 * id.get(0); + temp_dst_acc[index] = temp_t[index]; + }); + }); + + return 0; +} + +// No addrspacecast before loading and storing values +// CHECK: %[[#VALUE_1:]] = getelementptr inbounds %"struct.cl::sycl::range", %"struct.cl::sycl::range"* %{{.*}}, i64 0, i32 0 +// CHECK-NOT: %{{.*}} = addrspacecast i32* %[[#VALUE_1]] to i32 addrspace(4)* +// CHECK: %[[#VALUE_2:]] = getelementptr inbounds %"struct.cl::sycl::range", %"struct.cl::sycl::range"* %{{.*}}, i64 0, i32 0 +// CHECK-NOT: %{{.*}} = addrspacecast i32* %[[#VALUE_2]] to i32 addrspace(4)* +// CHECK-NOT: %{{.*}} = load i32, i32 addrspace(4)* %[[#VALUE_1]], align 4, !tbaa !6 +// CHECK: %[[#VALUE_3:]] = load i32, i32* %[[#VALUE_1]], align 4, !tbaa !6 +// CHECK-NOT: store i32 %[[#VALUE_3]], i32 addrspace(4)* %[[#VALUE_2]], align 4, !tbaa !6 +// CHECK: store i32 %[[#VALUE_3]], i32* %[[#VALUE_2]], align 4, !tbaa !6