[InferAddressSpaces] Add InferAddressSpaces pass to pipeline for SPIR

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

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<raw_pwrite_stream> 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<TargetLibraryInfoImpl> 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<TargetLibraryInfoImpl> 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<TargetLibraryInfoImpl> 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));

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -149,6 +149,7 @@ template <int dim>
 struct id {
   template <typename... T>
   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<dimensions> 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<dataT &>(Data);
+  }
+
 private:
   void __init(__attribute__((opencl_global)) dataT *Ptr, range<dimensions> AccessRange,
               range<dimensions> MemRange, id<dimensions> Offset) {}

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<float, 1>{range<1>{BLOCK_SIZE * 512}};
+  Q.submit([&](handler &cgh) {
+    auto temp_dst_acc = MatrixTemp.get_access<access::mode::write>(cgh);
+    auto temp_t = accessor<float, 1, access::mode::read_write, access::target::local>();
+    cgh.parallel_for<class test>(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