[Comgr] Add initial SPIR-V translation support

In this commit, we add a new SPIR-V data type:

  AMD_COMGR_DATA_KIND_SPIRV

and translation action:

  AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC

With this data type and action, users can provide a set of SPIR-V (.spv)
files, and have Comgr translate them into a set of LLVM IR bitcode (.bc)
objects.

We also add a HIP and OpenCL test for the new translation action. The
HIP test is currently disabled until compiler support is finalized.

Change-Id: Ic05c880158b9c4fd49550a5b2e51bb26dfed0efc

Author: lamb-j

amd/comgr/CMakeLists.txt

@@ -405,6 +405,7 @@ else()
     Support
     Symbolize
     TargetParser
+    SPIRVAMDLib
     )
 endif()
 

amd/comgr/docs/ReleaseNotes.md

@@ -156,7 +156,11 @@ action, and Comgr will internally unbundle and link via the OffloadBundler and l
 - (Data Type) AMD\_COMGR\_DATA\_KIND\_OBJ\_BUNDLE
   - This data kind represents a clang-offload-bundle of object files, and can be
   passed when calling the AMD\_COMGR\_ACTION\_UNBUNDLE action
-
+- (Data Type) AMD\_COMGR\_DATA\_KIND\_SPIRV
+  - This data kind represents a SPIR-V binary file (.spv)
+- (Action) AMD\_COMGR\_ACTION\_TRANSLATE\_SPIRV\_TO\_BC
+  - This accepts a set of SPIR-V (.spv) inputs, and returns a set of translated
+  bitcode (.bc) outputs
 
 Deprecated Comgr Actions and Data Types
 ---------------------------------------

amd/comgr/include/amd_comgr.h.in

@@ -392,10 +392,14 @@ typedef enum amd_comgr_data_kind_s {
    * The data is an object file bundle.
    */
   AMD_COMGR_DATA_KIND_OBJ_BUNDLE = 0x14,
+  /**
+   * The data is SPIR-V IR
+   */
+  AMD_COMGR_DATA_KIND_SPIRV = 0x15,
   /**
    * Marker for last valid data kind.
    */
-  AMD_COMGR_DATA_KIND_LAST = AMD_COMGR_DATA_KIND_OBJ_BUNDLE
+  AMD_COMGR_DATA_KIND_LAST = AMD_COMGR_DATA_KIND_SPIRV
 } amd_comgr_data_kind_t;
 
 /**
@@ -1772,10 +1776,21 @@ typedef enum amd_comgr_action_kind_s {
    */
   AMD_COMGR_ACTION_UNBUNDLE = 0xF,
 
+   /**
+   * Translate each source SPIR-V object in @p input into LLVM IR Bitcode.
+   * For each successful translation, add a bc object to @p result   *
+   *
+   * Return @p AMD_COMGR_STATUS_ERROR if any translation fails
+   *
+   * Return @p AMD_COMGR_STATUS_ERROR_INVALID_ARGUMENT
+   * if any input is not SPIR-V.
+   */
+  AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC = 0x13,
+
   /**
    * Marker for last valid action kind.
    */
-  AMD_COMGR_ACTION_LAST = AMD_COMGR_ACTION_UNBUNDLE
+  AMD_COMGR_ACTION_LAST = AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC
 } amd_comgr_action_kind_t;
 
 /**

amd/comgr/src/comgr-compiler.cpp

@@ -82,9 +82,11 @@
 #include "llvm/Support/VirtualFileSystem.h"
 #include "llvm/TargetParser/Host.h"
 
+#include "LLVMSPIRVLib/LLVMSPIRVLib.h"
 #include "time-stat/ts-interface.h"
 
 #include <csignal>
+#include <sstream>
 
 LLD_HAS_DRIVER(elf)
 
@@ -1851,6 +1853,63 @@ amd_comgr_status_t AMDGPUCompiler::linkToExecutable() {
   return amd_comgr_data_set_add(OutSetT, OutputT);
 }
 
+amd_comgr_status_t AMDGPUCompiler::translateSpirvToBitcode() {
+  if (auto Status = createTmpDirs()) {
+    return Status;
+  }
+
+  LLVMContext Context;
+  Context.setDiagnosticHandler(
+    std::make_unique<AMDGPUCompilerDiagnosticHandler>(this), true);
+
+  for (auto *Input : InSet->DataObjects) {
+
+    if (Input->DataKind != AMD_COMGR_DATA_KIND_SPIRV) {
+      return AMD_COMGR_STATUS_ERROR_INVALID_ARGUMENT;
+    }
+
+    // TODO: With C++23, we should investigate replacing with spanstream
+    // to avoid memory copies:
+    //  https://en.cppreference.com/w/cpp/io/basic_ispanstream
+    std::istringstream ISS(std::string(Input->Data, Input->Size));
+
+    llvm::Module *M;
+    std::string Err;
+
+    if (!llvm::readSpirv(Context, ISS, M, Err)) {
+      LogS << "Failed to load SPIR-V as LLVM Module: " << Err << '\n';
+      return AMD_COMGR_STATUS_ERROR;
+    }
+
+    SmallString<0> OutBuf;
+    BitcodeWriter Writer(OutBuf);
+    Writer.writeModule(*M, false, nullptr, false, nullptr);
+    Writer.writeSymtab();
+    Writer.writeStrtab();
+
+    amd_comgr_data_t OutputT;
+    if (auto Status = amd_comgr_create_data(AMD_COMGR_DATA_KIND_BC, &OutputT)) {
+      return Status;
+    }
+
+    // OutputT can be released after addition to the data_set
+    ScopedDataObjectReleaser SDOR(OutputT);
+
+    DataObject *Output = DataObject::convert(OutputT);
+    Output->setName(std::string(Input->Name) + std::string(".bc"));
+    Output->setData(OutBuf);
+
+    if (auto Status = amd_comgr_data_set_add(OutSetT, OutputT)) {
+      return Status;
+    }
+
+    LogS << "SPIR-V Translation: amd-llvm-spirv -r " << Input->Name << " " <<
+      Output->Name << "\n";
+  }
+
+  return AMD_COMGR_STATUS_SUCCESS;
+}
+
 AMDGPUCompiler::AMDGPUCompiler(DataAction *ActionInfo, DataSet *InSet,
                                DataSet *OutSet, raw_ostream &LogS)
     : ActionInfo(ActionInfo), InSet(InSet), OutSetT(DataSet::convert(OutSet)),

amd/comgr/src/comgr-compiler.h

@@ -140,6 +140,7 @@ class AMDGPUCompiler {
   amd_comgr_status_t linkToRelocatable();
   amd_comgr_status_t linkToExecutable();
   amd_comgr_status_t compileToExecutable();
+  amd_comgr_status_t translateSpirvToBitcode();
 
   amd_comgr_language_t getLanguage() const { return ActionInfo->Language; }
 };

amd/comgr/src/comgr.cpp

@@ -187,6 +187,8 @@ amd_comgr_status_t dispatchCompilerAction(amd_comgr_action_kind_t ActionKind,
     return Compiler.compileToBitcode(true);
   case AMD_COMGR_ACTION_COMPILE_SOURCE_TO_EXECUTABLE:
     return Compiler.compileToExecutable();
+  case AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC:
+    return Compiler.translateSpirvToBitcode();
 
   default:
     return AMD_COMGR_STATUS_ERROR_INVALID_ARGUMENT;
@@ -289,6 +291,8 @@ StringRef getActionKindName(amd_comgr_action_kind_t ActionKind) {
     return "AMD_COMGR_ACTION_COMPILE_SOURCE_TO_EXECUTABLE";
   case AMD_COMGR_ACTION_UNBUNDLE:
     return "AMD_COMGR_ACTION_UNBUNDLE";
+  case AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC:
+    return "AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC";
   }
 
   llvm_unreachable("invalid action");
@@ -1368,6 +1372,7 @@ amd_comgr_status_t AMD_COMGR_API
     case AMD_COMGR_ACTION_COMPILE_SOURCE_TO_RELOCATABLE:
     case AMD_COMGR_ACTION_COMPILE_SOURCE_WITH_DEVICE_LIBS_TO_BC:
     case AMD_COMGR_ACTION_COMPILE_SOURCE_TO_EXECUTABLE:
+    case AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC:
       ActionStatus = dispatchCompilerAction(ActionKind, ActionInfoP, InputSetP,
                                             ResultSetP, *LogP);
       break;

amd/comgr/test-lit/CMakeLists.txt

@@ -26,5 +26,6 @@ endmacro()
 set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR})
 
 add_comgr_lit_binary(source-to-bc-with-dev-libs)
+add_comgr_lit_binary(spirv-translator)
 
 add_dependencies(check-comgr test-lit)

amd/comgr/test-lit/comgr-sources/spirv-translator.c

@@ -0,0 +1,60 @@
+#include "amd_comgr.h"
+#include "common.h"
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+// Tests the AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC action
+//     Accepts one or more .spv files, and returns one or more .bc files
+
+int main(int argc, char *argv[]) {
+  char *BufSpirv;
+  size_t SizeSpirv;
+  amd_comgr_data_t DataSpirv;
+  amd_comgr_data_set_t DataSetSpirv, DataSetBc;
+  amd_comgr_action_info_t DataAction;
+  amd_comgr_status_t Status;
+  size_t Count;
+
+  if (argc != 4) {
+    fprintf(stderr, "Usage: spirv-translator file.spv -o file.spv.bc\n");
+    exit(1);
+  }
+
+  SizeSpirv = setBuf(argv[1], &BufSpirv);
+
+  amd_comgr_(create_data_set(&DataSetSpirv));
+  amd_comgr_(create_data(AMD_COMGR_DATA_KIND_SPIRV, &DataSpirv));
+  amd_comgr_(set_data(DataSpirv, SizeSpirv, BufSpirv));
+  amd_comgr_(set_data_name(DataSpirv, "source.spv"));
+  amd_comgr_(data_set_add(DataSetSpirv, DataSpirv));
+
+  amd_comgr_(create_action_info(&DataAction));
+  amd_comgr_(create_data_set(&DataSetBc));
+
+  amd_comgr_(do_action(AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC,
+                      DataAction, DataSetSpirv, DataSetBc));
+
+  amd_comgr_(action_data_count(DataSetBc, AMD_COMGR_DATA_KIND_BC, &Count));
+
+  if (Count != 1) {
+    printf("AMD_COMGR_ACTION_COMPILE_SOURCE_TO_BC Failed: "
+           "produced %zu BC objects (expected 1)\n",
+           Count);
+    exit(1);
+  }
+
+  // Write bitcode to file
+  amd_comgr_data_t DataSpirvBc;
+
+  amd_comgr_(action_data_get_data(
+      DataSetBc, AMD_COMGR_DATA_KIND_BC, 0, &DataSpirvBc));
+
+  dumpData(DataSpirvBc, argv[3]);
+
+  amd_comgr_(release_data(DataSpirv));
+  amd_comgr_(destroy_data_set(DataSetSpirv));
+  amd_comgr_(destroy_data_set(DataSetBc));
+  amd_comgr_(destroy_action_info(DataAction));
+  free(BufSpirv);
+}

amd/comgr/test-lit/spirv-translator.cl

@@ -0,0 +1,22 @@
+// COM: Enable this test once changes from amdspirv docker land
+
+// COM: Generate a spirv-targeted LLVM IR file from an OpenCL kernel
+// RUN: clang -c -emit-llvm --target=spirv64 %s -o %t.bc
+
+// COM: Translate LLVM IR to SPIRV format
+// RUN: amd-llvm-spirv %t.bc -o %t.spv
+
+// COM: Run Comgr Translator to covert SPIRV back to LLVM IR
+// RUN: spirv-translator %t.spv -o %t.translated.bc
+
+// COM: Dissasemble LLVM IR bitcode to LLVM IR text
+// RUN: llvm-dis %t.translated.bc -o - | FileCheck %s
+
+// COM: Verify LLVM IR text
+// CHECK: target triple = "spir64-unknown-unknown"
+// CHECK: define spir_kernel void @source
+
+void kernel source(__global int *j) {
+  *j += 2;
+}
+

amd/comgr/test-lit/spirv-translator.hip

@@ -0,0 +1,54 @@
+// XFAIL: *
+// COM: Generate a SPIRV file from a HIP kernel
+// RUN: clang -x hip --offload-arch=amdgcnspirv -nogpulib -nogpuinc \
+// RUN:    --no-gpu-bundle-output --offload-device-only -O3 %s -o %t.spv
+
+// COM: Run Comgr Translator to covert SPIRV back to LLVM IR
+// RUN: spirv-translator %t.spv -o %t.translated.bc
+
+// COM: Dissasemble LLVM IR bitcode to LLVM IR text
+// RUN: llvm-dis %t.translated.bc -o - | FileCheck %s
+
+// COM: Verify LLVM IR text
+// CHECK: target triple = "amdgcn-amd-amdhsa"
+// CHECK: define void @_Z11clean_valuePf
+// CHECK: define amdgpu_kernel void @_Z9add_valuePfS_S_
+
+#include <cstdlib>
+
+#define __constant__ __attribute__((constant))
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+#define __shared__ __attribute__((shared))
+#define __managed__ __attribute__((managed))
+#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
+
+struct dim3 {
+  unsigned x, y, z;
+  __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
+};
+
+#ifdef __HIP__
+typedef struct hipStream *hipStream_t;
+typedef enum hipError {} hipError_t;
+int hipConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
+                     hipStream_t stream = 0);
+extern "C" hipError_t __hipPushCallConfiguration(dim3 gridSize, dim3 blockSize,
+                                                 size_t sharedSize = 0,
+                                                 hipStream_t stream = 0);
+extern "C" hipError_t hipLaunchKernel(const void *func, dim3 gridDim,
+                                      dim3 blockDim, void **args,
+                                      size_t sharedMem,
+                                      hipStream_t stream);
+#endif
+
+__attribute__((device))
+void clean_value(float* ptr) { *ptr = 0; }
+
+__attribute__((global))
+void add_value(float* a, float* b, float* res) {
+    *res = *a + *b;
+
+    clean_value(a);
+}