[SYCL] Generalize GlobalOffset and enable it for AMDGPU (#5855)

The purpose of this patch is to generalize SYCL global offset pass and enable it for AMDGPU.

* enable global offset in AMD's HIP
* decorate SYCL kernel with dedicated MDNode: This removes the need for command line options added by the SYCL driver, discussed here: [SYCL] Generalize local accessor to shared mem pass #5149 (comment)
* extract common helpers for local accessor and global offset passes
* generalize the pass
* introduce builtin_amdgcn_implicit_offset and enable the pass for ADMGPU
* implement spirv_GlobalOffset_[x,y,z]
* update the docs

The main deviation from the NVPTX is the need for supporting address spaces. For AMD kernel arguments reside in constant address space, which for the case with offset forces a copy to private AS, in order to keep the call-graph interface coherent (we can't allocate const address space for the case without offset).

Corresponding test-suit PR: intel/llvm-test-suite#941

Author: jchlanda

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -316,5 +316,10 @@ TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x16_bf16, "V16fV4sV8sV16fiIiIi",
 TARGET_BUILTIN(__builtin_amdgcn_smfmac_i32_16x16x64_i8, "V4iV2iV4iV4iiIiIi", "nc", "mai-insts")
 TARGET_BUILTIN(__builtin_amdgcn_smfmac_i32_32x32x32_i8, "V16iV2iV4iV16iiIiIi", "nc", "mai-insts")
 
+//===----------------------------------------------------------------------===//
+// SYCL builtin.
+//===----------------------------------------------------------------------===//
+BUILTIN(__builtin_amdgcn_implicit_offset, "Ui*5", "nc")
+
 #undef BUILTIN
 #undef TARGET_BUILTIN

clang/lib/CodeGen/TargetInfo.cpp

@@ -151,6 +151,26 @@ static bool occupiesMoreThan(CodeGenTypes &cgt,
   return (intCount + fpCount > maxAllRegisters);
 }
 
+/// Helper function for AMDGCN and NVVM targets, adds a NamedMDNode with GV,
+/// Name, and Operand as operands, and adds the resulting MDNode to the
+/// AnnotationName MDNode.
+static void addAMDGCOrNVVMMetadata(const char *AnnotationName,
+                                   llvm::GlobalValue *GV, StringRef Name,
+                                   int Operand) {
+  llvm::Module *M = GV->getParent();
+  llvm::LLVMContext &Ctx = M->getContext();
+
+  // Get annotations metadata node.
+  llvm::NamedMDNode *MD = M->getOrInsertNamedMetadata(AnnotationName);
+
+  llvm::Metadata *MDVals[] = {
+      llvm::ConstantAsMetadata::get(GV), llvm::MDString::get(Ctx, Name),
+      llvm::ConstantAsMetadata::get(
+          llvm::ConstantInt::get(llvm::Type::getInt32Ty(Ctx), Operand))};
+  // Append metadata to annotations node.
+  MD->addOperand(llvm::MDNode::get(Ctx, MDVals));
+}
+
 bool SwiftABIInfo::isLegalVectorTypeForSwift(CharUnits vectorSize,
                                              llvm::Type *eltTy,
                                              unsigned numElts) const {
@@ -7327,18 +7347,7 @@ void NVPTXTargetCodeGenInfo::setTargetAttributes(
 
 void NVPTXTargetCodeGenInfo::addNVVMMetadata(llvm::GlobalValue *GV,
                                              StringRef Name, int Operand) {
-  llvm::Module *M = GV->getParent();
-  llvm::LLVMContext &Ctx = M->getContext();
-
-  // Get "nvvm.annotations" metadata node
-  llvm::NamedMDNode *MD = M->getOrInsertNamedMetadata("nvvm.annotations");
-
-  llvm::Metadata *MDVals[] = {
-      llvm::ConstantAsMetadata::get(GV), llvm::MDString::get(Ctx, Name),
-      llvm::ConstantAsMetadata::get(
-          llvm::ConstantInt::get(llvm::Type::getInt32Ty(Ctx), Operand))};
-  // Append metadata to nvvm.annotations
-  MD->addOperand(llvm::MDNode::get(Ctx, MDVals));
+  addAMDGCOrNVVMMetadata("nvvm.annotations", GV, Name, Operand);
 }
 
 bool NVPTXTargetCodeGenInfo::shouldEmitStaticExternCAliases() const {
@@ -9308,6 +9317,12 @@ class AMDGPUTargetCodeGenInfo : public TargetCodeGenInfo {
                             llvm::Type *BlockTy) const override;
   bool shouldEmitStaticExternCAliases() const override;
   void setCUDAKernelCallingConvention(const FunctionType *&FT) const override;
+
+private:
+  // Adds a NamedMDNode with GV, Name, and Operand as operands, and adds the
+  // resulting MDNode to the amdgcn.annotations MDNode.
+  static void addAMDGCNMetadata(llvm::GlobalValue *GV, StringRef Name,
+                                int Operand);
 };
 }
 
@@ -9324,6 +9339,11 @@ static bool requiresAMDGPUProtectedVisibility(const Decl *D,
            cast<VarDecl>(D)->getType()->isCUDADeviceBuiltinTextureType()));
 }
 
+void AMDGPUTargetCodeGenInfo::addAMDGCNMetadata(llvm::GlobalValue *GV,
+                                                StringRef Name, int Operand) {
+  addAMDGCOrNVVMMetadata("amdgcn.annotations", GV, Name, Operand);
+}
+
 void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes(
     const FunctionDecl *FD, llvm::Function *F, CodeGenModule &M) const {
   const auto *ReqdWGS =
@@ -9425,10 +9445,15 @@ void AMDGPUTargetCodeGenInfo::setTargetAttributes(
 
   const bool IsHIPKernel =
       M.getLangOpts().HIP && FD && FD->hasAttr<CUDAGlobalAttr>();
-
   if (IsHIPKernel)
     F->addFnAttr("uniform-work-group-size", "true");
 
+  // Create !{<func-ref>, metadata !"kernel", i32 1} node for SYCL kernels.
+  const bool IsSYCLKernel =
+      FD && M.getLangOpts().SYCLIsDevice && FD->hasAttr<SYCLKernelAttr>();
+  if (IsSYCLKernel)
+    addAMDGCNMetadata(F, "kernel", 1);
+
   if (M.getContext().getTargetInfo().allowAMDGPUUnsafeFPAtomics())
     F->addFnAttr("amdgpu-unsafe-fp-atomics", "true");
 

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5842,12 +5842,6 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     CmdArgs.push_back("-treat-scalable-fixed-error-as-warning");
   }
 
-  // Enable local accessor to shared memory pass for SYCL.
-  if (isa<BackendJobAction>(JA) && IsSYCLOffloadDevice &&
-      (Triple.isNVPTX() || Triple.isAMDGCN())) {
-    CmdArgs.push_back("-mllvm");
-    CmdArgs.push_back("-sycl-enable-local-accessor");
-  }
   // These two are potentially updated by AddClangCLArgs.
   codegenoptions::DebugInfoKind DebugInfoKind = codegenoptions::NoDebugInfo;
   bool EmitCodeView = false;

clang/lib/Driver/ToolChains/HIPAMD.cpp

@@ -78,12 +78,8 @@ void AMDGCN::Linker::constructLldCommand(Compilation &C, const JobAction &JA,
                                          const llvm::opt::ArgList &Args) const {
   // Construct lld command.
   // The output from ld.lld is an HSA code object file.
-  ArgStringList LldArgs{"-flavor",
-                        "gnu",
-                        "--no-undefined",
-                        "-shared",
-                        "-plugin-opt=-amdgpu-internalize-symbols",
-                        "-plugin-opt=-sycl-enable-local-accessor"};
+  ArgStringList LldArgs{"-flavor", "gnu", "--no-undefined", "-shared",
+                        "-plugin-opt=-amdgpu-internalize-symbols"};
 
   auto &TC = getToolChain();
   auto &D = TC.getDriver();

clang/test/CodeGenSYCL/kernel-annotation.cpp

@@ -0,0 +1,48 @@
+// RUN: %clang_cc1 -triple spir64-unknown-unknown -fsycl-is-device \
+// RUN: -S -emit-llvm %s -o %temp.ll
+// RUN: FileCheck -check-prefix=CHECK-SPIR --input-file %temp.ll %s
+
+// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fsycl-is-device \
+// RUN: -S -emit-llvm %s -o %temp.ll
+// RUN: FileCheck -check-prefix=CHECK-NVPTX --input-file %temp.ll %s
+
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -fsycl-is-device \
+// RUN: -S -emit-llvm %s -o %temp.ll
+// RUN: FileCheck -check-prefix=CHECK-AMDGCN --input-file %temp.ll %s
+
+// The test makes sure that `[nnvm|amdgcn].annotations are correctly generated
+// only for their respective targets.
+
+#include "Inputs/sycl.hpp"
+
+sycl::handler H;
+
+class Functor {
+public:
+  void operator()() const {}
+};
+
+// CHECK-SPIR-NOT: annotations =
+
+// CHECK-NVPTX: nvvm.annotations = !{[[FIRST:![0-9]]], [[SECOND:![0-9]]]}
+// CHECK-NVPTX: [[FIRST]] = !{void ()* @_ZTS7Functor, !"kernel", i32 1}
+// CHECK-NVPTX: [[SECOND]] = !{void ()* @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_E5foo_2, !"kernel", i32 1}
+
+// CHECK-AMDGCN: amdgcn.annotations = !{[[FIRST:![0-9]]], [[SECOND:![0-9]]]}
+// CHECK-AMDGCN: [[FIRST]] = !{void ()* @_ZTS7Functor, !"kernel", i32 1}
+// CHECK-AMDGCN: [[SECOND]] = !{void ()* @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_E5foo_2, !"kernel", i32 1}
+
+int main() {
+  sycl::queue q;
+  q.submit([&](sycl::handler &cgh) {
+    Functor foo{};
+    cgh.single_task(foo);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.parallel_for<class foo_2>(cl::sycl::range<1>(1),
+                                  [=](cl::sycl::item<1> item) {
+                                  });
+  });
+  return 0;
+}

clang/test/Driver/sycl-local-accessor-opt.cpp

@@ -1,4 +1,5 @@
 
+workitem/get_global_offset.ll
 group/collectives.cl
 group/collectives_helpers.ll
 atomic/loadstore_helpers.ll

libclc/amdgcn-amdhsa/libspirv/SOURCES

@@ -0,0 +1,42 @@
+;;===----------------------------------------------------------------------===//
+;;
+;; 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
+;;
+;;===----------------------------------------------------------------------===//
+
+#if __clang_major__ >= 7
+target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7"
+#else
+target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7"
+#endif
+
+; Function Attrs: nounwind readnone speculatable
+declare i32 addrspace(5)* @llvm.amdgcn.implicit.offset()
+
+define hidden i64 @_Z22__spirv_GlobalOffset_xv() nounwind alwaysinline {
+entry:
+  %0 = tail call i32 addrspace(5)* @llvm.amdgcn.implicit.offset()
+  %1 = load i32, i32 addrspace(5)* %0, align 4
+  %zext = zext i32 %1 to i64
+  ret i64 %zext
+}
+
+define hidden i64 @_Z22__spirv_GlobalOffset_yv() nounwind alwaysinline {
+entry:
+  %0 = tail call i32 addrspace(5)* @llvm.amdgcn.implicit.offset()
+  %arrayidx = getelementptr inbounds i32, i32 addrspace(5)* %0, i64 1
+  %1 = load i32, i32 addrspace(5)* %arrayidx, align 4
+  %zext = zext i32 %1 to i64
+  ret i64 %zext
+}
+
+define hidden i64 @_Z22__spirv_GlobalOffset_zv() nounwind alwaysinline {
+entry:
+  %0 = tail call i32 addrspace(5)* @llvm.amdgcn.implicit.offset()
+  %arrayidx = getelementptr inbounds i32, i32 addrspace(5)* %0, i64 2
+  %1 = load i32, i32 addrspace(5)* %arrayidx, align 4
+  %zext = zext i32 %1 to i64
+  ret i64 %zext
+}

libclc/amdgcn-amdhsa/libspirv/workitem/get_global_offset.ll

@@ -1,4 +1,3 @@
-workitem/get_global_offset.cl
 workitem/get_group_id.cl
 workitem/get_global_size.cl
 workitem/get_local_id.cl