From a4c5deddc3a9a888ca806639b3ef6a7c2c2124a3 Mon Sep 17 00:00:00 2001 From: Jacob Lambert Date: Thu, 6 Nov 2025 20:35:29 -0800 Subject: [PATCH 1/2] [Comgr] Add AMD_COMGR_ACTION_COMPILE_SOURCE_TO_SPIRV action - Add new action to compile HIP source directly to SPIRV - Implement compileSourceToSpirv() using --offload-arch=amdgcnspirv - Add source-to-spirv test driver binary - Add LIT tests for source-to-spirv compilation --- amd/comgr/include/amd_comgr.h.in | 17 +++++- amd/comgr/src/comgr-compiler.cpp | 36 ++++++++++++ amd/comgr/src/comgr-compiler.h | 1 + amd/comgr/src/comgr.cpp | 5 ++ amd/comgr/test-lit/CMakeLists.txt | 1 + .../test-lit/comgr-sources/source-to-spirv.c | 55 +++++++++++++++++++ .../source-to-spirv-with-options.hip | 49 +++++++++++++++++ .../test-lit/spirv-tests/source-to-spirv.hip | 39 +++++++++++++ 8 files changed, 202 insertions(+), 1 deletion(-) create mode 100644 amd/comgr/test-lit/comgr-sources/source-to-spirv.c create mode 100644 amd/comgr/test-lit/spirv-tests/source-to-spirv-with-options.hip create mode 100644 amd/comgr/test-lit/spirv-tests/source-to-spirv.hip diff --git a/amd/comgr/include/amd_comgr.h.in b/amd/comgr/include/amd_comgr.h.in index 6a91b38e39e98..e4d66f3491403 100644 --- a/amd/comgr/include/amd_comgr.h.in +++ b/amd/comgr/include/amd_comgr.h.in @@ -1823,10 +1823,25 @@ typedef enum amd_comgr_action_kind_s { */ AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC = 0x13, + /** + * Compile each HIP source data object in @p input in order. For each + * successful compilation add a SPIR-V data object to @p result. Resolve + * any include source names using the names of include data objects in + * @p input. Resolve any include relative path names using the working + * directory path in @p info. Compile the source for the language in @p + * info. + * + * Return @p AMD_COMGR_STATUS_ERROR if any compilation fails. + * + * Return @p AMD_COMGR_STATUS_ERROR_INVALID_ARGUMENT if language is not + * HIP in @p info. + */ + AMD_COMGR_ACTION_COMPILE_SOURCE_TO_SPIRV = 0x14, + /** * Marker for last valid action kind. */ - AMD_COMGR_ACTION_LAST = AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC + AMD_COMGR_ACTION_LAST = AMD_COMGR_ACTION_COMPILE_SOURCE_TO_SPIRV } amd_comgr_action_kind_t; /** diff --git a/amd/comgr/src/comgr-compiler.cpp b/amd/comgr/src/comgr-compiler.cpp index 2f9126d2f030e..fc73df6fa8e50 100644 --- a/amd/comgr/src/comgr-compiler.cpp +++ b/amd/comgr/src/comgr-compiler.cpp @@ -2156,6 +2156,42 @@ amd_comgr_status_t AMDGPUCompiler::compileSpirvToRelocatable() { return processFiles(AMD_COMGR_DATA_KIND_RELOCATABLE, ".o", TranslatedSpirv); } +amd_comgr_status_t AMDGPUCompiler::compileSourceToSpirv() { + if (auto Status = createTmpDirs()) { + return Status; + } + + if (ActionInfo->Language != AMD_COMGR_LANGUAGE_HIP) { + return AMD_COMGR_STATUS_ERROR_INVALID_ARGUMENT; + } + + if (auto Status = addIncludeFlags()) { + return Status; + } + + if (auto Status = addCompilationFlags()) { + return Status; + } + + // Add SPIRV-specific compilation flags + Args.push_back("--offload-arch=amdgcnspirv"); + Args.push_back("--no-gpu-bundle-output"); + Args.push_back("-c"); + + +#if _WIN32 + Args.push_back("-fshort-wchar"); +#endif + + if (ActionInfo->ShouldLinkDeviceLibs) { + if (auto Status = addDeviceLibraries()) { + return Status; + } + } + + return processFiles(AMD_COMGR_DATA_KIND_SPIRV, ".spv"); +} + AMDGPUCompiler::AMDGPUCompiler(DataAction *ActionInfo, DataSet *InSet, DataSet *OutSet, raw_ostream &LogS) : ActionInfo(ActionInfo), InSet(InSet), OutSetT(DataSet::convert(OutSet)), diff --git a/amd/comgr/src/comgr-compiler.h b/amd/comgr/src/comgr-compiler.h index 455a179c85050..a950a130ecc41 100644 --- a/amd/comgr/src/comgr-compiler.h +++ b/amd/comgr/src/comgr-compiler.h @@ -81,6 +81,7 @@ class AMDGPUCompiler { amd_comgr_status_t compileToExecutable(); amd_comgr_status_t compileSpirvToRelocatable(); amd_comgr_status_t translateSpirvToBitcode(); + amd_comgr_status_t compileSourceToSpirv(); amd_comgr_language_t getLanguage() const { return ActionInfo->Language; } }; diff --git a/amd/comgr/src/comgr.cpp b/amd/comgr/src/comgr.cpp index f0ace22833952..53ea784db97ab 100644 --- a/amd/comgr/src/comgr.cpp +++ b/amd/comgr/src/comgr.cpp @@ -102,6 +102,8 @@ amd_comgr_status_t dispatchCompilerAction(amd_comgr_action_kind_t ActionKind, return Compiler.compileSpirvToRelocatable(); case AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC: return Compiler.translateSpirvToBitcode(); + case AMD_COMGR_ACTION_COMPILE_SOURCE_TO_SPIRV: + return Compiler.compileSourceToSpirv(); default: return AMD_COMGR_STATUS_ERROR_INVALID_ARGUMENT; @@ -193,6 +195,8 @@ StringRef getActionKindName(amd_comgr_action_kind_t ActionKind) { return "AMD_COMGR_ACTION_COMPILE_SPIRV_TO_RELOCATABLE"; case AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC: return "AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC"; + case AMD_COMGR_ACTION_COMPILE_SOURCE_TO_SPIRV: + return "AMD_COMGR_ACTION_COMPILE_SOURCE_TO_SPIRV"; } llvm_unreachable("invalid action"); @@ -1302,6 +1306,7 @@ amd_comgr_status_t AMD_COMGR_API case AMD_COMGR_ACTION_COMPILE_SOURCE_TO_EXECUTABLE: case AMD_COMGR_ACTION_COMPILE_SPIRV_TO_RELOCATABLE: case AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC: + case AMD_COMGR_ACTION_COMPILE_SOURCE_TO_SPIRV: ActionStatus = dispatchCompilerAction(ActionKind, ActionInfoP, InputSetP, ResultSetP, *LogP); break; diff --git a/amd/comgr/test-lit/CMakeLists.txt b/amd/comgr/test-lit/CMakeLists.txt index b3e478e4002a8..bfc765e4a743f 100644 --- a/amd/comgr/test-lit/CMakeLists.txt +++ b/amd/comgr/test-lit/CMakeLists.txt @@ -42,6 +42,7 @@ add_comgr_lit_binary(source-to-bc-with-dev-libs c) add_comgr_lit_binary(spirv-translator c) add_comgr_lit_binary(compile-opencl-minimal c) add_comgr_lit_binary(spirv-to-reloc c) +add_comgr_lit_binary(source-to-spirv c) add_comgr_lit_binary(unbundle c) add_comgr_lit_binary(get-version c) add_comgr_lit_binary(status-string c) diff --git a/amd/comgr/test-lit/comgr-sources/source-to-spirv.c b/amd/comgr/test-lit/comgr-sources/source-to-spirv.c new file mode 100644 index 0000000000000..12d9fd98d2c64 --- /dev/null +++ b/amd/comgr/test-lit/comgr-sources/source-to-spirv.c @@ -0,0 +1,55 @@ +//===- source-to-spirv.c --------------------------------------------------===// +// +// Part of Comgr, under the Apache License v2.0 with LLVM Exceptions. See +// amd/comgr/LICENSE.TXT in this repository for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "amd_comgr.h" +#include "common.h" +#include +#include +#include + +int main(int argc, char *argv[]) { + char *BufSource; + size_t SizeSource; + amd_comgr_data_t DataSource; + amd_comgr_data_set_t DataSetSource, DataSetSpirv; + amd_comgr_action_info_t DataAction; + size_t Count; + + if (argc != 3) { + fprintf(stderr, "Usage: source-to-spirv file.hip file.spv\n"); + exit(1); + } + + SizeSource = setBuf(argv[1], &BufSource); + + amd_comgr_(create_data(AMD_COMGR_DATA_KIND_SOURCE, &DataSource)); + amd_comgr_(set_data(DataSource, SizeSource, BufSource)); + amd_comgr_(set_data_name(DataSource, "file.hip")); + + amd_comgr_(create_data_set(&DataSetSource)); + amd_comgr_(data_set_add(DataSetSource, DataSource)); + + amd_comgr_(create_action_info(&DataAction)); + amd_comgr_(action_info_set_language(DataAction, AMD_COMGR_LANGUAGE_HIP)); + + amd_comgr_(create_data_set(&DataSetSpirv)); + amd_comgr_(do_action(AMD_COMGR_ACTION_COMPILE_SOURCE_TO_SPIRV, + DataAction, DataSetSource, DataSetSpirv)); + + amd_comgr_data_t DataSpirv; + amd_comgr_(action_data_get_data(DataSetSpirv, AMD_COMGR_DATA_KIND_SPIRV, + 0, &DataSpirv)); + dumpData(DataSpirv, argv[2]); + + amd_comgr_(release_data(DataSource)); + amd_comgr_(destroy_data_set(DataSetSource)); + amd_comgr_(destroy_data_set(DataSetSpirv)); + amd_comgr_(destroy_action_info(DataAction)); + free(BufSource); +} + diff --git a/amd/comgr/test-lit/spirv-tests/source-to-spirv-with-options.hip b/amd/comgr/test-lit/spirv-tests/source-to-spirv-with-options.hip new file mode 100644 index 0000000000000..174d6c629ab89 --- /dev/null +++ b/amd/comgr/test-lit/spirv-tests/source-to-spirv-with-options.hip @@ -0,0 +1,49 @@ +// REQUIRES: comgr-has-spirv + +// COM: Compile HIP source to SPIR-V using Comgr with additional options +// RUN: AMD_COMGR_EMIT_VERBOSE_LOGS=1 AMD_COMGR_REDIRECT_LOGS=source-to-spirv-logs.txt \ +// RUN: source-to-spirv %s %t.spv + +// COM: Verify verbose logs show the SPIR-V target architecture +// RUN: grep 'offload-arch=amdgcnspirv' source-to-spirv-logs.txt + +// COM: Verify the SPIR-V file was created and is non-empty +// RUN: test -s %t.spv + +// COM: Translate SPIR-V back to LLVM IR bitcode +// RUN: spirv-translator %t.spv -o %t.bc + +// COM: Disassemble LLVM IR bitcode to text and verify content +// RUN: llvm-dis %t.bc -o - | FileCheck %s + +// COM: Verify LLVM IR contains expected functions, attributes, and target triple +// CHECK: target triple = "amdgcn-amd-amdhsa" +// CHECK: define void @_Z11clean_valuePf +// CHECK: define amdgpu_kernel void @_Z12complex_mathPfS_S_ +// CHECK: call {{.*}} @llvm.sqrt +// CHECK: call {{.*}} @llvm.sin + +// RUN: rm source-to-spirv-logs.txt + +#include +#include + +#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__))) + +__attribute__((device)) +void clean_value(float* ptr) { *ptr = 0; } + +__attribute__((global)) +void complex_math(float* a, float* b, float* res) { + float temp = sqrt(*a) + sin(*b); + *res = temp * 2.0f; + + clean_value(a); +} + diff --git a/amd/comgr/test-lit/spirv-tests/source-to-spirv.hip b/amd/comgr/test-lit/spirv-tests/source-to-spirv.hip new file mode 100644 index 0000000000000..df87ac5f4c90e --- /dev/null +++ b/amd/comgr/test-lit/spirv-tests/source-to-spirv.hip @@ -0,0 +1,39 @@ +// REQUIRES: comgr-has-spirv + +// COM: Compile HIP source to SPIR-V using Comgr +// RUN: source-to-spirv %s %t.spv + +// COM: Verify the SPIR-V file was created and is non-empty +// RUN: test -s %t.spv + +// COM: Translate SPIR-V back to LLVM IR bitcode +// RUN: spirv-translator %t.spv -o %t.bc + +// COM: Disassemble LLVM IR bitcode to text +// RUN: llvm-dis %t.bc -o - | FileCheck %s + +// COM: Verify LLVM IR contains expected functions and target triple +// CHECK: target triple = "amdgcn-amd-amdhsa" +// CHECK: define void @_Z11clean_valuePf +// CHECK: define amdgpu_kernel void @_Z9add_valuePfS_S_ + +#include + +#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__))) + +__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); +} + From faa3ea59f596d5b643bcdcb8eefe3cc91e9dd428 Mon Sep 17 00:00:00 2001 From: Jacob Lambert Date: Thu, 6 Nov 2025 21:17:10 -0800 Subject: [PATCH 2/2] remove LIT tests for now --- .../source-to-spirv-with-options.hip | 49 ------------------- .../test-lit/spirv-tests/source-to-spirv.hip | 39 --------------- 2 files changed, 88 deletions(-) delete mode 100644 amd/comgr/test-lit/spirv-tests/source-to-spirv-with-options.hip delete mode 100644 amd/comgr/test-lit/spirv-tests/source-to-spirv.hip diff --git a/amd/comgr/test-lit/spirv-tests/source-to-spirv-with-options.hip b/amd/comgr/test-lit/spirv-tests/source-to-spirv-with-options.hip deleted file mode 100644 index 174d6c629ab89..0000000000000 --- a/amd/comgr/test-lit/spirv-tests/source-to-spirv-with-options.hip +++ /dev/null @@ -1,49 +0,0 @@ -// REQUIRES: comgr-has-spirv - -// COM: Compile HIP source to SPIR-V using Comgr with additional options -// RUN: AMD_COMGR_EMIT_VERBOSE_LOGS=1 AMD_COMGR_REDIRECT_LOGS=source-to-spirv-logs.txt \ -// RUN: source-to-spirv %s %t.spv - -// COM: Verify verbose logs show the SPIR-V target architecture -// RUN: grep 'offload-arch=amdgcnspirv' source-to-spirv-logs.txt - -// COM: Verify the SPIR-V file was created and is non-empty -// RUN: test -s %t.spv - -// COM: Translate SPIR-V back to LLVM IR bitcode -// RUN: spirv-translator %t.spv -o %t.bc - -// COM: Disassemble LLVM IR bitcode to text and verify content -// RUN: llvm-dis %t.bc -o - | FileCheck %s - -// COM: Verify LLVM IR contains expected functions, attributes, and target triple -// CHECK: target triple = "amdgcn-amd-amdhsa" -// CHECK: define void @_Z11clean_valuePf -// CHECK: define amdgpu_kernel void @_Z12complex_mathPfS_S_ -// CHECK: call {{.*}} @llvm.sqrt -// CHECK: call {{.*}} @llvm.sin - -// RUN: rm source-to-spirv-logs.txt - -#include -#include - -#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__))) - -__attribute__((device)) -void clean_value(float* ptr) { *ptr = 0; } - -__attribute__((global)) -void complex_math(float* a, float* b, float* res) { - float temp = sqrt(*a) + sin(*b); - *res = temp * 2.0f; - - clean_value(a); -} - diff --git a/amd/comgr/test-lit/spirv-tests/source-to-spirv.hip b/amd/comgr/test-lit/spirv-tests/source-to-spirv.hip deleted file mode 100644 index df87ac5f4c90e..0000000000000 --- a/amd/comgr/test-lit/spirv-tests/source-to-spirv.hip +++ /dev/null @@ -1,39 +0,0 @@ -// REQUIRES: comgr-has-spirv - -// COM: Compile HIP source to SPIR-V using Comgr -// RUN: source-to-spirv %s %t.spv - -// COM: Verify the SPIR-V file was created and is non-empty -// RUN: test -s %t.spv - -// COM: Translate SPIR-V back to LLVM IR bitcode -// RUN: spirv-translator %t.spv -o %t.bc - -// COM: Disassemble LLVM IR bitcode to text -// RUN: llvm-dis %t.bc -o - | FileCheck %s - -// COM: Verify LLVM IR contains expected functions and target triple -// CHECK: target triple = "amdgcn-amd-amdhsa" -// CHECK: define void @_Z11clean_valuePf -// CHECK: define amdgpu_kernel void @_Z9add_valuePfS_S_ - -#include - -#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__))) - -__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); -} -