[libc-commits] [clang] [libc] [llvm] [OpenMP] [HIP] [CUDA] Support for Emissary APIs. (PR #218473)
Greg Rodgers via libc-commits
libc-commits at lists.llvm.org
Tue Aug 25 05:19:50 PDT 2026
https://github.com/gregrodgers updated https://github.com/llvm/llvm-project/pull/218473
>From cf6351e5c0470e01ed7a08cefe7d8b54d6a3798c Mon Sep 17 00:00:00 2001
From: gregrodgers <Gregory.Rodgers at amd.com>
Date: Mon, 24 Aug 2026 07:54:29 -0500
Subject: [PATCH 1/8] [OPENMP] Support for Emissary APIs
---
clang/lib/CodeGen/CGEmitEmissaryExec.cpp | 419 +++++++++++++
clang/lib/CodeGen/CGExpr.cpp | 8 +
clang/lib/CodeGen/CMakeLists.txt | 1 +
clang/lib/CodeGen/CodeGenFunction.h | 9 +
clang/lib/Headers/CMakeLists.txt | 1 +
clang/lib/Headers/EmissaryIds.h | 105 ++++
clang/lib/Headers/llvm_libc_wrappers/stdio.h | 3 +
...gcn_target_printf_unknown_size_arguments.c | 51 ++
libc/docs/gpu/emissary.rst | 43 ++
libc/docs/gpu/index.rst | 1 +
libc/shared/CMakeLists.txt | 1 +
libc/shared/emissary_rpc_server.h | 577 ++++++++++++++++++
libc/shared/rpc_util.h | 2 +-
libc/src/__support/RPC/CMakeLists.txt | 1 +
.../__support/RPC/emissary_device_utils.cpp | 104 ++++
libc/test/shared/CMakeLists.txt | 20 +
libc/test/shared/emissary_registry_test.cpp | 63 ++
offload/liboffload/exports | 1 +
offload/libomptarget/exports | 1 +
offload/plugins-nextgen/common/src/RPC.cpp | 4 +
20 files changed, 1414 insertions(+), 1 deletion(-)
create mode 100644 clang/lib/CodeGen/CGEmitEmissaryExec.cpp
create mode 100644 clang/lib/Headers/EmissaryIds.h
create mode 100644 clang/test/OpenMP/amdgcn_target_printf_unknown_size_arguments.c
create mode 100644 libc/docs/gpu/emissary.rst
create mode 100644 libc/shared/emissary_rpc_server.h
create mode 100644 libc/src/__support/RPC/emissary_device_utils.cpp
create mode 100644 libc/test/shared/emissary_registry_test.cpp
diff --git a/clang/lib/CodeGen/CGEmitEmissaryExec.cpp b/clang/lib/CodeGen/CGEmitEmissaryExec.cpp
new file mode 100644
index 0000000000000..8f269608e2c95
--- /dev/null
+++ b/clang/lib/CodeGen/CGEmitEmissaryExec.cpp
@@ -0,0 +1,419 @@
+//===- CGEmitEmissaryExec.cpp - Codegen for _emissary_exec ---------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// EmitEmissaryExec:
+//
+// When a device call to the variadic function _emissary_exec is encountered
+// (in CGExpr.cpp) EmitEmissaryExec does these steps:
+//
+// 1. If string lens are runtime dependent, Emit code to determine runtime len.
+// 2. Emits call to allocate memory __llvm_emissary_premalloc,
+// 3. Emit stores of each arg into arg buffer,
+// 4. Emits call to function __llvm_emissary_rpc or __llvm_emissary_rpc_dm
+//
+// The arg buffer is a struct that contains the length, number of args, an
+// array of 4-byte keys that represent the type of each arg, an array of
+// aligned "data" values for each arg, and finally the runtime string values.
+// If an arg is a string the data value is the runtime length of the string.
+// Each 4-byte key contains the llvm type ID and the number of bits for the
+// type. encoded by the macro PACK_TY_BITLEN(x,y) ((uint32_t)x << 16) |
+// ((uint32_t)y)
+//
+//===----------------------------------------------------------------------===//
+
+#include "../../../clang/lib/Headers/EmissaryIds.h"
+#include "CodeGenFunction.h"
+#include "clang/Basic/Builtins.h"
+#include "llvm/IR/DataLayout.h"
+#include "llvm/IR/Instruction.h"
+#include "llvm/Support/MathExtras.h"
+#include "llvm/Transforms/Utils/AMDGPUEmitPrintf.h"
+
+using namespace clang;
+using namespace CodeGen;
+
+// These static helper functions support EmitEmissaryExec.
+static llvm::Function *getOmpStrlenDeclaration(CodeGenModule &CGM) {
+ auto &M = CGM.getModule();
+ // Args are pointer to char and maxstringlen
+ llvm::Type *ArgTypes[] = {CGM.Int8PtrTy, CGM.Int32Ty};
+ llvm::FunctionType *OmpStrlenFTy =
+ llvm::FunctionType::get(CGM.Int32Ty, ArgTypes, false);
+ if (auto *F = M.getFunction("__strlen_max")) {
+ assert(F->getFunctionType() == OmpStrlenFTy);
+ return F;
+ }
+ llvm::Function *FN = llvm::Function::Create(
+ OmpStrlenFTy, llvm::GlobalVariable::ExternalLinkage, "__strlen_max", &M);
+ return FN;
+}
+
+// Determines if an expression is a string with variable length
+static bool isVarString(const clang::Expr *ArgX, const clang::Type *ArgXTy,
+ const llvm::Value *Arg) {
+ if ((ArgXTy->isPointerType() || ArgXTy->isConstantArrayType()) &&
+ ArgXTy->getPointeeOrArrayElementType()->isCharType() && !ArgX->isLValue())
+ return true;
+ // Ensure the VarDecl has an initializer
+ if (const auto *DRE = dyn_cast<DeclRefExpr>(ArgX))
+ if (const auto *VD = dyn_cast<VarDecl>(DRE->getDecl()))
+ if (!VD->getInit() ||
+ !llvm::isa<StringLiteral>(VD->getInit()->IgnoreImplicit()))
+ return true;
+ return false;
+}
+
+// Determines if an argument is a string
+static bool isString(const clang::Type *ArgXTy) {
+ if ((ArgXTy->isPointerType() || ArgXTy->isConstantArrayType()) &&
+ ArgXTy->getPointeeOrArrayElementType()->isCharType())
+ return true;
+ else
+ return false;
+}
+
+// Gets a string literal to write into the transfer buffer
+static const StringLiteral *getSL(const clang::Expr *ArgX,
+ const clang::Type *ArgXTy) {
+ // String in ArgX has known constant length
+ if (!ArgXTy->isConstantArrayType()) {
+ // Allow constant string to be a declared variable,
+ // But it must be constant and initialized.
+ const DeclRefExpr *DRE = cast<DeclRefExpr>(ArgX);
+ const VarDecl *VarD = cast<VarDecl>(DRE->getDecl());
+ ArgX = VarD->getInit()->IgnoreImplicit();
+ }
+ const StringLiteral *SL = cast<StringLiteral>(ArgX);
+ return SL;
+}
+
+// Returns a function pointer to __llvm_emissary_premalloc
+static llvm::Function *getEmissaryAllocDeclaration(CodeGenModule &CGM) {
+ auto &M = CGM.getModule();
+ const char *ExecuteName = "__llvm_emissary_premalloc";
+ llvm::Type *ArgTypes[] = {CGM.Int32Ty};
+ llvm::Function *FN;
+ // Maybe this should be pointer to char instead of pointer to void
+ llvm::FunctionType *VargsFnAllocFuncType = llvm::FunctionType::get(
+ CGM.getTypes().ConvertType(
+ CGM.getContext().getPointerType(CGM.getContext().VoidTy)),
+ ArgTypes, false);
+ if (!(FN = M.getFunction(ExecuteName)))
+ FN = llvm::Function::Create(VargsFnAllocFuncType,
+ llvm::GlobalVariable::ExternalLinkage,
+ ExecuteName, &M);
+ assert(FN->getFunctionType() == VargsFnAllocFuncType);
+ return FN;
+}
+
+// Returns a function pointer to __llvm_emissary_rpc
+static llvm::Function *getEmissaryExecDeclaration(CodeGenModule &CGM,
+ bool HasXfers) {
+ const char *ExecuteName =
+ HasXfers ? "__llvm_emissary_rpc_dm" : "__llvm_emissary_rpc";
+ auto &M = CGM.getModule();
+ llvm::Type *ArgTypes[] = {
+ CGM.Int32Ty, CGM.getTypes().ConvertType(CGM.getContext().getPointerType(
+ CGM.getContext().VoidTy))};
+ llvm::Function *FN;
+ llvm::FunctionType *VarfnFuncType =
+ llvm::FunctionType::get(CGM.Int64Ty, ArgTypes, false);
+ if (!(FN = M.getFunction(ExecuteName)))
+ FN = llvm::Function::Create(
+ VarfnFuncType, llvm::GlobalVariable::ExternalLinkage, ExecuteName, &M);
+ assert(FN->getFunctionType() == VarfnFuncType);
+ return FN;
+}
+
+// A macro to pack the llvm type ID and NumBits into 4-byte key
+#define PACK_TY_BITLEN(x, y) ((uint32_t)x << 16) | ((uint32_t)y)
+
+static EmisTyID getEmisTyID(llvm::Type::TypeID TyId) {
+ switch (TyId) {
+ case llvm::Type::HalfTyID: ///< 16-bit floating point type
+ case llvm::Type::X86_FP80TyID: ///< 80-bit floating point type (X87)
+ case llvm::Type::BFloatTyID: ///< 16-bit floating point type (7-bit
+ ///< significand)
+ return EmisInvalidTy;
+ case llvm::Type::FloatTyID: ///< 32-bit floating point type
+ case llvm::Type::DoubleTyID: ///< 64-bit floating point type
+ case llvm::Type::FP128TyID: ///< 128-bit floating point type (112-bit
+ ///< significand)
+ return EmisFloatTy;
+ case llvm::Type::PPC_FP128TyID: ///< 128-bit floating point type (two 64-bits,
+ ///< PowerPC)
+ case llvm::Type::VoidTyID: ///< type with no size
+ case llvm::Type::LabelTyID: ///< Labels
+ case llvm::Type::MetadataTyID: ///< Metadata
+ case llvm::Type::X86_AMXTyID: ///< AMX vectors (8192 bits, X86 specific)
+ case llvm::Type::TokenTyID: ///< Tokens
+ return EmisInvalidTy;
+ // Derived types... see DerivedTypes.h file.
+ case llvm::Type::IntegerTyID: ///< Arbitrary bit width integers
+ return EmisIntegerTy;
+ case llvm::Type::ByteTyID: ///< Arbitrary bit width bytes
+ case llvm::Type::FunctionTyID: ///< Functions
+ return EmisInvalidTy;
+ case llvm::Type::PointerTyID: ///< Pointers
+ return EmisPointerTy;
+ case llvm::Type::StructTyID: ///< Structures
+ case llvm::Type::ArrayTyID: ///< Arrays
+ case llvm::Type::FixedVectorTyID: ///< Fixed width SIMD vector type
+ case llvm::Type::ScalableVectorTyID: ///< Scalable SIMD vector type
+ case llvm::Type::TypedPointerTyID: ///< Typed pointer used by some GPU targets
+ case llvm::Type::TargetExtTyID: ///< Target extension type
+ return EmisInvalidTy;
+ default:
+ return EmisInvalidTy;
+ }
+}
+
+// ----- External function EmitEmissaryExec called from CGExpr.cpp -----
+RValue CodeGenFunction::EmitEmissaryExec(const CallExpr *E) {
+ assert(getTarget().getTriple().isAMDGCN() ||
+ getTarget().getTriple().isNVPTX());
+ assert(E->getNumArgs() >= 1); // _emissary_exec always has at least one arg.
+ const llvm::DataLayout &DL = CGM.getDataLayout();
+ CallArgList Args;
+
+ EmitCallArgs(Args,
+ E->getDirectCallee()->getType()->getAs<FunctionProtoType>(),
+ E->arguments(), E->getDirectCallee(),
+ /* ParamsToSkip = */ 0);
+
+ // We don't know how to emit non-scalar varargs.
+ if (std::any_of(Args.begin() + 1, Args.end(), [&](const CallArg &A) {
+ return !A.getRValue(*this).isScalar();
+ })) {
+ CGM.ErrorUnsupported(E, "non-scalar arg in GPU vargs function");
+ return RValue::get(llvm::ConstantInt::get(IntTy, 0));
+ }
+ // Arg 0 is the packed emisid supplied by the caller, so Args maps 1:1 onto
+ // E->arguments(). It has to be a compile-time constant because the buffer
+ // layout below depends on the transfer counts encoded in it. _PACK_EMIS_IDS()
+ // folds to a constant, so this only fires on a malformed hand-written call --
+ // diagnose it instead of crashing on the cast.
+ RValue EmisIdRV = Args[0].getKnownRValue();
+ if (!EmisIdRV.isScalar() ||
+ !llvm::isa<llvm::ConstantInt>(EmisIdRV.getScalarVal())) {
+ CGM.ErrorUnsupported(E, "non-constant emissary id in _emissary_exec call");
+ return RValue::get(llvm::ConstantInt::get(IntTy, 0));
+ }
+
+ unsigned NumArgs = (unsigned)Args.size();
+ llvm::SmallVector<llvm::Type *, 32> ArgTypes;
+ llvm::SmallVector<llvm::Value *, 32> VarStrLengths;
+ llvm::Value *TotalVarStrsLength = llvm::ConstantInt::get(Int32Ty, 0);
+ bool HasVarStrings = false;
+ ArgTypes.push_back(Int32Ty); // 1st field in struct is total DataLen
+ ArgTypes.push_back(Int32Ty); // 2nd field in struct will be num args
+ // An array of 4-byte keys that describe the arg type
+ for (unsigned I = 0; I < NumArgs; ++I)
+ ArgTypes.push_back(Int32Ty);
+
+ // Track the size of the numeric data length and string length
+ unsigned DataLenCT = (unsigned)(DL.getTypeAllocSize(Int32Ty)) * (NumArgs + 2);
+ unsigned AllStringsLenCT = 0;
+
+ // --- 1st Pass over Args to create ArgTypes and count size ---
+ size_t StructOffset = 4 * (NumArgs + 2);
+ for (unsigned I = 0; I < NumArgs; I++) {
+ llvm::Value *Arg = Args[I].getRValue(*this).getScalarVal();
+ llvm::Type *ArgType = Arg->getType();
+ // Skip string processing on arg0 which may not be in E->getArg(0)
+ if (I != 0) {
+ const Expr *ArgX = E->getArg(I)->IgnoreParenCasts();
+ auto *ArgXTy = ArgX->getType().getTypePtr();
+ if (isString(ArgXTy)) {
+ if (isVarString(ArgX, ArgXTy, Arg)) {
+ HasVarStrings = true;
+ if (auto *PtrTy = dyn_cast<llvm::PointerType>(ArgType))
+ if (PtrTy->getPointerAddressSpace()) {
+ Arg = Builder.CreateAddrSpaceCast(Arg, CGM.Int8PtrTy);
+ ArgType = Arg->getType();
+ }
+ llvm::Value *VarStrLen =
+ Builder.CreateCall(getOmpStrlenDeclaration(CGM),
+ {Arg, llvm::ConstantInt::get(Int32Ty, 1024)});
+ VarStrLengths.push_back(VarStrLen);
+ TotalVarStrsLength = Builder.CreateAdd(TotalVarStrsLength, VarStrLen,
+ "sum_of_var_strings_length");
+ ArgType = Int32Ty;
+ } else {
+ const StringLiteral *SL = getSL(ArgX, ArgXTy);
+ StringRef ArgString = SL->getString();
+ AllStringsLenCT += ((int)ArgString.size() + 1);
+ // change ArgType from char ptr to int to contain string length
+ ArgType = Int32Ty;
+ }
+ } // end of processing string argument
+ } // End of skip 1st arg
+ // if ArgTypeSize is >4 bytes we need to insert dummy align
+ // values in the struct so all stores can be aligned .
+ // These dummy fields must be inserted before the arg.
+ //
+ // In the pass below where the stores are generated careful
+ // tracking of the index into the struct is necessary.
+ size_t NeedsPadding = (StructOffset % (size_t)DL.getTypeAllocSize(ArgType));
+ if (NeedsPadding) {
+ DataLenCT += (unsigned)NeedsPadding;
+ StructOffset += NeedsPadding;
+ ArgTypes.push_back(Int32Ty); // could assert that NeedsPadding == 4 here
+ }
+
+ ArgTypes.push_back(ArgType);
+ DataLenCT += ((int)DL.getTypeAllocSize(ArgType));
+ StructOffset += (size_t)DL.getTypeAllocSize(ArgType);
+ }
+
+ // --- Generate call to __llvm_emissary_premalloc to get data pointer
+ if (HasVarStrings)
+ TotalVarStrsLength = Builder.CreateAdd(
+ TotalVarStrsLength,
+ llvm::ConstantInt::get(Int32Ty, AllStringsLenCT + DataLenCT),
+ "total_buffer_size");
+ llvm::Value *BufferLen =
+ HasVarStrings
+ ? TotalVarStrsLength
+ : llvm::ConstantInt::get(Int32Ty, AllStringsLenCT + DataLenCT);
+ llvm::Value *DataStructPtr =
+ Builder.CreateCall(getEmissaryAllocDeclaration(CGM), {BufferLen});
+
+ // --- Cast the generic return pointer to be a struct in device global memory
+ llvm::StructType *DataStructTy =
+ llvm::StructType::create(ArgTypes, "varfn_args_store");
+ unsigned AS = getContext().getTargetAddressSpace(LangAS::cuda_device);
+ llvm::Value *BufferPtr = Builder.CreatePointerCast(
+ DataStructPtr, llvm::PointerType::get(CGM.getLLVMContext(), AS),
+ "varfn_args_store_casted");
+ // --- Header of struct contains length and NumArgs ---
+ llvm::Value *DataLenField = llvm::ConstantInt::get(Int32Ty, DataLenCT);
+ llvm::Value *P = Builder.CreateStructGEP(DataStructTy, BufferPtr, 0);
+ Builder.CreateAlignedStore(DataLenField, P,
+ DL.getPrefTypeAlign(DataLenField->getType()));
+ llvm::Value *NumArgsField = llvm::ConstantInt::get(Int32Ty, NumArgs);
+ P = Builder.CreateStructGEP(DataStructTy, BufferPtr, 1);
+ Builder.CreateAlignedStore(NumArgsField, P,
+ DL.getPrefTypeAlign(NumArgsField->getType()));
+
+ // --- 2nd Pass: create array of 4-byte keys to describe each arg
+ for (unsigned I = 0; I < NumArgs; I++) {
+ llvm::Type *Ty = Args[I].getRValue(*this).getScalarVal()->getType();
+ llvm::Type::TypeID ArgTypeId =
+ Args[I].getRValue(*this).getScalarVal()->getType()->getTypeID();
+ EmisTyID EmisTypeId = getEmisTyID(ArgTypeId);
+
+ // Get type size in bits. Usually 64 or 32.
+ uint32_t NumBits = 0;
+ if (I > 0 &&
+ isString(E->getArg(I)->IgnoreParenCasts()->getType().getTypePtr()))
+ // The llvm typeID for string is pointer. Since pointer NumBits is 0,
+ // we set NumBits to 1 to distinguish pointer type ID as string pointer.
+ NumBits = 1;
+ else
+ NumBits = Ty->getScalarSizeInBits();
+ // Create a key that combines llvm typeID and size
+ llvm::Value *Key =
+ llvm::ConstantInt::get(Int32Ty, PACK_TY_BITLEN(EmisTypeId, NumBits));
+ P = Builder.CreateStructGEP(DataStructTy, BufferPtr, I + 2);
+ Builder.CreateAlignedStore(Key, P, DL.getPrefTypeAlign(Key->getType()));
+ }
+
+ // --- 3rd Pass: Store data values for each arg ---
+ unsigned VarStringIndex = 0;
+ unsigned StructIndex = 2 + NumArgs;
+ StructOffset = 4 * StructIndex;
+ bool HasXfers;
+ for (unsigned I = 0; I < NumArgs; I++) {
+ llvm::Value *Arg = nullptr;
+ if (I == 0) {
+ Arg = Args[I].getKnownRValue().getScalarVal();
+ uint64_t UInt64Value = llvm::cast<llvm::ConstantInt>(Arg)->getZExtValue();
+ uint32_t Lower32 = (uint32_t)(UInt64Value & 0xFFFFFFFF);
+ HasXfers = Lower32 ? true : false;
+ } else {
+ const Expr *ArgX = E->getArg(I)->IgnoreParenCasts();
+ auto *ArgXTy = ArgX->getType().getTypePtr();
+ if (isString(ArgXTy)) {
+ if (isVarString(ArgX, ArgXTy, Arg)) {
+ Arg = VarStrLengths[VarStringIndex];
+ VarStringIndex++;
+ } else {
+ const StringLiteral *SL = getSL(ArgX, ArgXTy);
+ StringRef ArgString = SL->getString();
+ int ArgStrLen = (int)ArgString.size() + 1;
+ // Change Arg from a char pointer to the integer string length
+ Arg = llvm::ConstantInt::get(Int32Ty, ArgStrLen);
+ }
+ } else {
+ Arg = Args[I].getKnownRValue().getScalarVal();
+ }
+ }
+ size_t StructElementSize = (size_t)DL.getTypeAllocSize(Arg->getType());
+ size_t NeedsPadding = (StructOffset % StructElementSize);
+ if (NeedsPadding) {
+ // Skip over dummy fields in struct to align
+ StructOffset += NeedsPadding; // should assert NeedsPadding == 4
+ StructIndex++;
+ }
+ P = Builder.CreateStructGEP(DataStructTy, BufferPtr, StructIndex);
+ Builder.CreateAlignedStore(Arg, P, DL.getPrefTypeAlign(Arg->getType()));
+ StructOffset += StructElementSize;
+ StructIndex++;
+ }
+
+ // --- 4th Pass: memcpy all strings after the data values ---
+ // bitcast the struct in device global memory as a char buffer
+ Address BufferPtrByteAddr =
+ Address(Builder.CreatePointerCast(
+ BufferPtr, llvm::PointerType::get(CGM.getLLVMContext(), AS),
+ "_casted"),
+ Int8Ty, CharUnits::fromQuantity(1));
+
+ // BufferPtrByteAddr is a pointer to where we want to write the next string
+ BufferPtrByteAddr = Builder.CreateConstInBoundsByteGEP(
+ BufferPtrByteAddr, CharUnits::fromQuantity(DataLenCT));
+ VarStringIndex = 0;
+ // Skip string processing on arg0 which may not be in E->getArg(0)
+ for (unsigned I = 1; I < NumArgs; ++I) {
+ llvm::Value *Arg = Args[I].getKnownRValue().getScalarVal();
+ const Expr *ArgX = E->getArg(I)->IgnoreParenCasts();
+ auto *ArgXTy = ArgX->getType().getTypePtr();
+ if (isString(ArgXTy)) {
+ if (isVarString(ArgX, ArgXTy, Arg)) {
+ llvm::Value *VarStrLength = VarStrLengths[VarStringIndex];
+ VarStringIndex++;
+ Address SrcAddr = Address(Arg, Int8Ty, CharUnits::fromQuantity(1));
+ Builder.CreateMemCpy(BufferPtrByteAddr, SrcAddr, VarStrLength);
+ // update BufferPtrByteAddr for next string memcpy
+ llvm::Value *PtrAsInt = BufferPtrByteAddr.emitRawPointer(*this);
+ BufferPtrByteAddr =
+ Address(Builder.CreateGEP(Int8Ty, PtrAsInt,
+ ArrayRef<llvm::Value *>(VarStrLength)),
+ Int8Ty, CharUnits::fromQuantity(1));
+ } else {
+ const StringLiteral *SL = getSL(ArgX, ArgXTy);
+ StringRef ArgString = SL->getString();
+ int ArgStrLen = (int)ArgString.size() + 1;
+ Address SrcAddr = CGM.GetAddrOfConstantStringFromLiteral(SL);
+ Builder.CreateMemCpy(BufferPtrByteAddr, SrcAddr, ArgStrLen);
+ // update BufferPtrByteAddr for next memcpy
+ BufferPtrByteAddr = Builder.CreateConstInBoundsByteGEP(
+ BufferPtrByteAddr, CharUnits::fromQuantity(ArgStrLen));
+ }
+ }
+ }
+ // --- Generate call to __llvm_emissary_rpc and return RValue
+ llvm::Value *EmisRc = Builder.CreateCall(
+ getEmissaryExecDeclaration(CGM, HasXfers), {BufferLen, DataStructPtr});
+ // truncate long long int to int for printf return value.
+ if ((E->getDirectCallee()->getNameAsString() == "fprintf") ||
+ (E->getDirectCallee()->getNameAsString() == "printf"))
+ EmisRc = Builder.CreateTrunc(EmisRc, CGM.Int32Ty, "emis_rc");
+ return RValue::get(EmisRc);
+}
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index eff6a7de320d7..470095dc3a9b1 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -7131,6 +7131,14 @@ RValue CodeGenFunction::EmitCall(QualType CalleeType,
StaticOperator = true;
}
+ // Call EmitEmissaryExec(E) on device pass calls to _emissary_exec.
+ if ((CGM.getTriple().isAMDGCN() || CGM.getTriple().isNVPTX()) && FnType &&
+ isa<FunctionProtoType>(FnType) &&
+ cast<FunctionProtoType>(FnType)->isVariadic() && E->getDirectCallee() &&
+ E->getDirectCallee()->getIdentifier() &&
+ E->getDirectCallee()->getIdentifier()->isStr("_emissary_exec"))
+ return EmitEmissaryExec(E);
+
auto Arguments = E->arguments();
if (StaticOperator) {
// If we're calling a static operator, we need to emit the object argument
diff --git a/clang/lib/CodeGen/CMakeLists.txt b/clang/lib/CodeGen/CMakeLists.txt
index 853f038c4186b..13ba89ec5a577 100644
--- a/clang/lib/CodeGen/CMakeLists.txt
+++ b/clang/lib/CodeGen/CMakeLists.txt
@@ -75,6 +75,7 @@ add_clang_library(clangCodeGen
CGDebugInfo.cpp
CGDecl.cpp
CGDeclCXX.cpp
+ CGEmitEmissaryExec.cpp
CGException.cpp
CGExpr.cpp
CGExprAgg.cpp
diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index dfb6f2ff65a7d..14cecbd55beba 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -4626,6 +4626,15 @@ class CodeGenFunction : public CodeGenTypeCache {
llvm::CallBase **CallOrInvoke = nullptr,
CGFunctionInfo const **ResolvedFnInfo = nullptr);
+ /// EmitEmissaryExec generates IR to allocate an arg buffer, fill buffer with
+ /// args, then generate a call to __llvm_emissary_rpc(sz,buf) when a call-site
+ /// to _emissary_exec(...) is encountered. _emissary_exec greatly simplifies
+ /// construction of device stub functions when creating an emissary API.
+ /// The LLVM RPC device utility __llvm_emissary_rpc triggers the rpc client
+ /// server exchange where the RPC host server executes the designated function
+ /// for each active lane in the GPU warp.
+ RValue EmitEmissaryExec(const CallExpr *E);
+
// If a Call or Invoke instruction was emitted for this CallExpr, this method
// writes the pointer to `CallOrInvoke` if it's not null.
RValue EmitCallExpr(const CallExpr *E,
diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt
index 21b6eb5e38052..867c7f8d3bbfb 100644
--- a/clang/lib/Headers/CMakeLists.txt
+++ b/clang/lib/Headers/CMakeLists.txt
@@ -42,6 +42,7 @@ set(core_files
tgmath.h
unwind.h
varargs.h
+ EmissaryIds.h
)
set(arm_common_files
diff --git a/clang/lib/Headers/EmissaryIds.h b/clang/lib/Headers/EmissaryIds.h
new file mode 100644
index 0000000000000..d7891f0e5aa1e
--- /dev/null
+++ b/clang/lib/Headers/EmissaryIds.h
@@ -0,0 +1,105 @@
+//===-- EmissaryIds.h - Emissary API identifiers ------------- C/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
+//
+//===----------------------------------------------------------------------===//
+//
+// Defines Emissary API identifiers. This header is used by both host
+// and device compilations.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef OFFLOAD_EMISSARY_IDS_H
+#define OFFLOAD_EMISSARY_IDS_H
+
+#define __DEVATTR__
+#if defined(__NVPTX__) || defined(__AMDGCN__)
+#if defined(__HIP__) || defined(__CUDA__)
+#if defined(__DEVATTR__)
+#undef __DEVATTR__
+#endif
+#define __DEVATTR__ __device__
+#endif
+#endif
+
+#if defined(__cplusplus)
+extern "C" {
+#endif
+
+__DEVATTR__ unsigned long long int _emissary_exec(const unsigned long long int,
+ ...);
+
+#if defined(__cplusplus)
+}
+#endif
+
+#define _PACK_EMIS_IDS(a, b, c, d) \
+ ((unsigned long long)a << 48) | ((unsigned long long)b << 32) | \
+ ((unsigned long long)c << 16) | ((unsigned long long)d)
+
+enum EmisTyID {
+ EmisInvalidTy = 0,
+ EmisFloatTy,
+ EmisIntegerTy,
+ EmisPointerTy,
+};
+
+/// These are the various Emissary APIs currently defined.
+/// MPI, HDF5, and, RESERVE are "external" Emissary APIs whose device stubs and
+/// host runtime support are provided by library maintainers typically in the
+/// form of a header such as "EmissaryMPI.h". The stubs call _emissary_exec.
+/// The host runtime support will call functions from the actual host library
+/// which are often platform specific and thus only linkable by an application.
+/// A small demo of an external Emissary API (EmissaryMPI.h) is found in docs.
+
+typedef enum {
+ EMIS_ID_INVALID,
+ EMIS_ID_FORTRT,
+ EMIS_ID_PRINT,
+ EMIS_ID_MPI,
+ EMIS_ID_HDF5,
+ EMIS_ID_RESERVE,
+} offload_emis_id_t;
+
+/// This structure is created by emisExtractArgBuf to get information
+/// from the data buffer passed by rpc.
+typedef struct {
+ unsigned int DataLen;
+ unsigned int NumArgs;
+ unsigned int emisid;
+ unsigned int emisfnid;
+ unsigned int NumSendXfers;
+ unsigned int NumRecvXfers;
+ unsigned long long data_not_used;
+ char *keyptr;
+ char *argptr;
+ char *strptr;
+} emisArgBuf_t;
+
+typedef unsigned long long EmissaryReturn_t;
+typedef unsigned long long emis_argptr_t;
+typedef EmissaryReturn_t emisfn_t(void *, ...);
+
+typedef enum service_rc {
+ _ERC_SUCCESS = 0,
+ _ERC_STATUS_ERROR = 1,
+ _ERC_DATA_USED_ERROR = 2,
+ _ERC_ADDINT_ERROR = 3,
+ _ERC_ADDFLOAT_ERROR = 4,
+ _ERC_ADDSTRING_ERROR = 5,
+ _ERC_UNSUPPORTED_ID_ERROR = 6,
+ _ERC_INVALID_ID_ERROR = 7,
+ _ERC_ERROR_INVALID_REQUEST = 8,
+} service_rc;
+
+#define LLVM_EMISSARY_BASE 'e'
+#define LLVM_EMISSARY_OPCODE(n) (LLVM_EMISSARY_BASE << 24 | n)
+
+typedef enum {
+ OFFLOAD_EMISSARY = LLVM_EMISSARY_OPCODE(1),
+ OFFLOAD_EMISSARY_DM = LLVM_EMISSARY_OPCODE(2),
+} offload_emissary_t;
+
+#endif // OFFLOAD_EMISSARY_IDS_H
diff --git a/clang/lib/Headers/llvm_libc_wrappers/stdio.h b/clang/lib/Headers/llvm_libc_wrappers/stdio.h
index 0c3e44823da70..9d7be5efeb33b 100644
--- a/clang/lib/Headers/llvm_libc_wrappers/stdio.h
+++ b/clang/lib/Headers/llvm_libc_wrappers/stdio.h
@@ -14,6 +14,9 @@
#endif
#include_next <stdio.h>
+#if __has_include("emissary_print.h")
+#include "emissary_print.h"
+#endif
#if defined(__HIP__) || defined(__CUDA__)
#define __LIBC_ATTRS __attribute__((device))
diff --git a/clang/test/OpenMP/amdgcn_target_printf_unknown_size_arguments.c b/clang/test/OpenMP/amdgcn_target_printf_unknown_size_arguments.c
new file mode 100644
index 0000000000000..09c6511a427c4
--- /dev/null
+++ b/clang/test/OpenMP/amdgcn_target_printf_unknown_size_arguments.c
@@ -0,0 +1,51 @@
+// REQUIRES: amdgpu-registered-target
+// REQUIRES: x86-registered-target
+
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple amdgcn-amd-amdhsa -fopenmp-is-device -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-host-ir-file-path %t-host.bc -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK
+// expected-no-diagnostics
+
+extern int printf(const char *, ...);
+
+int CheckMultipleArgs(int a) {
+ char *test = "testing";
+ char *t;
+#pragma omp target private(t)
+ {
+ t = test + a;
+ printf("%s %d %s", t, 21, test);
+// CHECK-LABEL: define weak_odr protected amdgpu_kernel void @{{.*}}CheckMultipleArgs
+// CHECK: entry:
+// CHECK: [[TEST_ADDR:%[a-zA-Z0-9_.]+]] = alloca ptr, align 8, addrspace(5)
+// CHECK: [[A_ADDR:%[a-zA-Z0-9_.]+]] = alloca i64, align 8, addrspace(5)
+// CHECK: [[DYN_PTR_ADDR:%[a-zA-Z0-9_.]+]] = alloca ptr, align 8, addrspace(5)
+// CHECK: [[T_ADDR:%[a-zA-Z0-9_.]+]] = alloca ptr, align 8, addrspace(5)
+// CHECK: [[TEST_CAST:%[a-zA-Z0-9_.]+]] = addrspacecast ptr addrspace(5) [[TEST_ADDR]] to ptr
+// CHECK: [[A_CAST:%[a-zA-Z0-9_.]+]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
+// CHECK: [[DYN_PTR_CAST:%[a-zA-Z0-9_.]+]] = addrspacecast ptr addrspace(5) [[DYN_PTR_ADDR]] to ptr
+// CHECK: [[T_CAST:%[a-zA-Z0-9_.]+]] = addrspacecast ptr addrspace(5) [[T_ADDR]] to ptr
+// CHECK: store ptr %test, ptr [[TEST_CAST]], align 8
+// CHECK: store i64 %a, ptr [[A_CAST]], align 8
+// CHECK: store ptr %dyn_ptr, ptr [[DYN_PTR_CAST]], align 8
+// CHECK: [[INIT_CALL:%[a-zA-Z0-9_.]+]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) {{.*}} to ptr), ptr %dyn_ptr)
+// CHECK: [[EXEC_USER_CODE:%[a-zA-Z0-9_.]+]] = icmp eq i32 [[INIT_CALL]], -1
+// CHECK: br i1 [[EXEC_USER_CODE]], label %[[USER_CODE_ENTRY:.+]], label %[[WORKER_EXIT:.+]]
+
+// CHECK: [[USER_CODE_ENTRY]]:
+// CHECK: [[LOAD_TEST:%[0-9]+]] = load ptr, ptr [[TEST_CAST]], align 8
+// CHECK: [[LOAD_A:%[0-9]+]] = load i32, ptr [[A_CAST]], align 4
+// CHECK: %idx.ext = sext i32 [[LOAD_A]] to i64
+// CHECK: %add.ptr = getelementptr inbounds i8, ptr [[LOAD_TEST]], i64 %idx.ext
+// CHECK: store ptr %add.ptr, ptr [[T_CAST]], align 8
+// CHECK: [[LOAD_T:%[0-9]+]] = load ptr, ptr [[T_CAST]], align 8
+// CHECK: [[LOAD_TEST_AGAIN:%[0-9]+]] = load ptr, ptr [[TEST_CAST]], align 8
+// CHECK: call i32 (ptr, ...) @printf(ptr noundef addrspacecast (ptr addrspace(4) @.str to ptr), ptr noundef [[LOAD_T]], i32 noundef 21, ptr noundef [[LOAD_TEST_AGAIN]])
+// CHECK: call void @__kmpc_target_deinit()
+// CHECK: ret void
+
+// CHECK: [[WORKER_EXIT]]:
+// CHECK: ret void
+ }
+
+ return 0;
+}
diff --git a/libc/docs/gpu/emissary.rst b/libc/docs/gpu/emissary.rst
new file mode 100644
index 0000000000000..f863a3f5171ab
--- /dev/null
+++ b/libc/docs/gpu/emissary.rst
@@ -0,0 +1,43 @@
+.. _libc_gpu_emissary:
+
+========
+Emissary
+========
+
+.. note:: This feature is experimental and may change in the future.
+
+Emissary lets GPU device code call **host library** functions -- MPI, HDF5,
+``printf``, or an application-supplied API -- by name. Clang lowers a call to
+the variadic entry point ``_emissary_exec`` into a packed argument buffer and
+a GPU RPC request; a host-side server unpacks the buffer and dispatches to a
+handler that invokes the real library.
+
+It builds on the RPC transport described in :ref:`libc_gpu_rpc`, adding two
+opcodes (``OFFLOAD_EMISSARY`` and ``OFFLOAD_EMISSARY_DM``) and a runtime
+handler registry so that adding support for a new host library is a library
+change rather than a compiler release.
+
+Components in this repository
+=============================
+
+.. list-table::
+ :header-rows: 1
+
+ * - File
+ - Role
+ * - ``clang/lib/Headers/EmissaryIds.h``
+ - Wire ABI: ``_emissary_exec``, ``_PACK_EMIS_IDS``, ``emisArgBuf_t``, API
+ id enum, RPC opcodes.
+ * - ``clang/lib/CodeGen/CGEmitEmissaryExec.cpp``
+ - Packs call-site arguments into the buffer and emits the RPC call.
+ Interception lives in ``CGExpr.cpp``.
+ * - ``libc/src/__support/RPC/emissary_device_utils.cpp``
+ - Device helpers: ``__llvm_emissary_premalloc``,
+ ``__llvm_emissary_rpc``, ``__llvm_emissary_rpc_dm``.
+ * - ``libc/shared/emissary_rpc_server.h``
+ - Host registry (``EmissaryRegister`` / ``EmissaryLookup``), buffer
+ unpack, ``EmissaryTop``, ``handleEmissaryOpcodes``.
+ * - ``offload/plugins-nextgen/common/src/RPC.cpp``
+ - Server thread and opcode routing.
+ * - ``libc/test/shared/emissary_registry_test.cpp``
+ - Registry unit tests.
diff --git a/libc/docs/gpu/index.rst b/libc/docs/gpu/index.rst
index 1fca67205acb4..9d2c3dc602031 100644
--- a/libc/docs/gpu/index.rst
+++ b/libc/docs/gpu/index.rst
@@ -18,3 +18,4 @@ learn more about this project.
rpc
testing
motivation
+ emissary
diff --git a/libc/shared/CMakeLists.txt b/libc/shared/CMakeLists.txt
index 1237254aeebfe..64bd0de13de8d 100644
--- a/libc/shared/CMakeLists.txt
+++ b/libc/shared/CMakeLists.txt
@@ -3,6 +3,7 @@ set(LLVM_LIBC_SHARED_RPC_EXPORT_HEADERS
"${CMAKE_CURRENT_SOURCE_DIR}/rpc_util.h"
"${CMAKE_CURRENT_SOURCE_DIR}/rpc_dispatch.h"
"${CMAKE_CURRENT_SOURCE_DIR}/rpc_server.h"
+ "${CMAKE_CURRENT_SOURCE_DIR}/emissary_rpc_server.h"
"${CMAKE_CURRENT_SOURCE_DIR}/rpc_opcodes.h"
)
diff --git a/libc/shared/emissary_rpc_server.h b/libc/shared/emissary_rpc_server.h
new file mode 100644
index 0000000000000..31d8aeb5a732b
--- /dev/null
+++ b/libc/shared/emissary_rpc_server.h
@@ -0,0 +1,577 @@
+//===-- Shared memory RPC server instantiation ------------------*- 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
+//
+//===----------------------------------------------------------------------===//
+//
+// This file is an extension of rpc_server.h
+//
+// Consumers must add the Clang resource header directory to the include path
+// when compiling translation units that include this header (directly or via
+// <shared/emissary_rpc_server.h>). EmissaryIds.h is installed there, not under
+// lib/llvm/include:
+//
+// -I$("$CXX" -print-resource-dir)/include
+//
+// Typical HIP/OpenMP demo builds also pass -I for lib/llvm/include (or
+// llvm/include) so that <shared/emissary_rpc_server.h> resolves.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIBC_SRC___SUPPORT_RPC_EMISSARY_RPC_SERVER_H
+#define LLVM_LIBC_SRC___SUPPORT_RPC_EMISSARY_RPC_SERVER_H
+
+#if __has_include("../clang/lib/Headers/EmissaryIds.h")
+#include "../clang/lib/Headers/EmissaryIds.h"
+#else
+#include "EmissaryIds.h"
+#endif
+
+#include "rpc.h"
+#include "rpc_opcodes.h"
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <unordered_map>
+
+//===----------------------------------------------------------------------===//
+// Emissary host handler registry
+//
+// Runtime registry that maps an Emissary API id to the host handler that
+// services it. It lets a client library register its dispatcher at load time
+// so the RPC server (EmissaryTop) can route requests without a compile-time
+// switch over every known client.
+//
+// The registry is a fixed-size table with no dynamic allocation.
+// EmissaryRegister and EmissaryLookup have C linkage so a client shared object
+// can register against the server without C++ name-mangling coupling. The
+// backing table is a C++17 inline variable, so all translation units in a
+// program share one instance; on ELF it also merges across shared objects
+// under default visibility, letting a client .so and the server share the same
+// table.
+//===----------------------------------------------------------------------===//
+
+/// Upper bound on the number of distinct Emissary API ids. The table is indexed
+/// directly by \c emisid, so this also bounds the largest id that can be
+/// registered. It comfortably exceeds the current \c offload_emis_id_t range
+/// and leaves room for reserved/dynamic ids.
+#define EMISSARY_MAX_REGISTERED_IDS 64
+
+/// Host handler for one Emissary API id. The signature matches the per-client
+/// dispatchers (\c EmissaryMPI, \c EmissaryHDF5, ...): it receives the RPC data
+/// buffer, the decoded argument descriptor, and the unpacked argument vector.
+using EmissaryHandler_t = EmissaryReturn_t (*)(char *data, emisArgBuf_t *ab,
+ emis_argptr_t *args[]);
+
+namespace emissary_registry_detail {
+/// Backing table, indexed directly by Emissary API id. As a C++17 inline
+/// variable it has exactly one instance across the whole program, but that
+/// only holds across shared objects if the symbol keeps default visibility
+/// *and* the linker is told to export it. Two things are needed to make that
+/// true even when a consuming DSO is linked with an explicit
+/// `--version-script` (as libomptarget/liboffload are) and/or
+/// -fvisibility-inlines-hidden (LLVM's default project-wide flag), either of
+/// which would otherwise localize this vague-linkage symbol into each DSO and
+/// silently defeat the cross-DSO sharing this registry depends on:
+/// 1. Explicit default visibility, so the compiler doesn't hide it.
+/// 2. A stable `asm` symbol name, so it can be listed by name in a
+/// `--version-script` `global:` clause without embedding an
+/// Itanium-mangled C++ symbol (`_ZN...`) in linker input -- mangled names
+/// are an ABI/compiler-version implementation detail, not something
+/// version scripts should hardcode.
+/// \internal
+__attribute__((visibility("default"))) inline EmissaryHandler_t
+ Table[EMISSARY_MAX_REGISTERED_IDS] asm("EmissaryRegistryTable") = {};
+} // namespace emissary_registry_detail
+
+extern "C" {
+
+/// Register a host handler for an Emissary API id.
+///
+/// \param emisid the Emissary API id (an \c offload_emis_id_t value or a
+/// reserved dynamic id) to associate with \p Handler.
+/// \param Handler the host dispatcher to invoke for \p emisid; must not be
+/// null.
+/// \returns \c true on success; \c false if \p emisid is out of range,
+/// \p Handler is null, or a *different* handler is already registered for
+/// \p emisid. Re-registering the identical handler is idempotent and
+/// succeeds.
+///
+/// Explicit default visibility (see \c Table above): this symbol must merge
+/// across DSOs even when the caller is built with -fvisibility-inlines-hidden.
+__attribute__((visibility("default"))) inline bool
+EmissaryRegister(unsigned int emisid, EmissaryHandler_t Handler) {
+ if (emisid >= EMISSARY_MAX_REGISTERED_IDS || Handler == nullptr)
+ return false;
+ EmissaryHandler_t &Slot = emissary_registry_detail::Table[emisid];
+ // Reject last-wins: two libraries claiming the same id is a configuration
+ // error, not something to silently overwrite.
+ if (Slot != nullptr && Slot != Handler)
+ return false;
+ Slot = Handler;
+ return true;
+}
+
+/// Look up the host handler registered for an Emissary API id.
+///
+/// \param emisid the Emissary API id to look up.
+/// \returns the registered handler, or null if \p emisid is out of range or no
+/// handler is registered for it.
+///
+/// Explicit default visibility (see \c Table above): this symbol must merge
+/// across DSOs even when the caller is built with -fvisibility-inlines-hidden.
+__attribute__((visibility("default"))) inline EmissaryHandler_t
+EmissaryLookup(unsigned int emisid) {
+ if (emisid >= EMISSARY_MAX_REGISTERED_IDS)
+ return nullptr;
+ return emissary_registry_detail::Table[emisid];
+}
+
+} // extern "C"
+
+// No Emissary API host handler is declared or called here by name.
+// Every client -- MPI, HDF5, PRINT, and RESERVE -- self-registers its host
+// dispatcher with the runtime registry defined above, so EmissaryTop routes all
+// of them through EmissaryLookup without a compile-time weak symbol or switch
+// case. PRINT (device printf/fprintf) is no longer built into this header
+// either: it is an ordinary client library (libemissary_print) that an
+// application links like libemissary_mpi. RESERVE is registered by the
+// user/reserved client library.
+extern "C" {
+/// Optional FORCE_OPT=1 SDMA path for device MPI Put/Get (libemissary_mpi).
+/// This is an internal optimization hook, not an Emissary API, so it keeps its
+/// weak-stub design: libLLVMOffload links without libemissary_mpi; the app
+/// overrides with a strong definition from libemissary_mpi when FORCE_OPT SDMA
+/// is used.
+__attribute__((weak)) int
+emissary_mpi_sdma_try_dm_buffer(char *rpc_buffer,
+ unsigned long long *out_result) {
+ (void)rpc_buffer;
+ (void)out_result;
+ return -1;
+}
+} // end extern "C"
+
+namespace rpc {
+namespace internal {
+
+// emisExtractArgBuf extract ArgBuf using protocol EmitEmissaryExec makes.
+static void emisExtractArgBuf(char *data, emisArgBuf_t *ab) {
+
+ uint32_t *Int32Data = (uint32_t *)data;
+ ab->DataLen = Int32Data[0];
+ ab->NumArgs = Int32Data[1];
+
+ // Note: while the data buffer contains all args including strings,
+ // ab->DataLen does not include strings. It only counts header, keys,
+ // and aligned numerics.
+
+ ab->keyptr = data + (2 * sizeof(int));
+ ab->argptr = ab->keyptr + (ab->NumArgs * sizeof(int));
+ ab->strptr = data + (size_t)ab->DataLen;
+ int AlignFill = 0;
+ if (((size_t)ab->argptr) % (size_t)8) {
+ ab->argptr += 4;
+ AlignFill = 4;
+ }
+
+ // Extract the two emissary identifiers and number of send
+ // and recv device data transfers. These are 4 16 bit values
+ // packed into a single 64-bit field.
+ uint64_t Arg1 = *(uint64_t *)ab->argptr;
+ ab->emisid = (unsigned int)((Arg1 >> 48) & 0xFFFF);
+ ab->emisfnid = (unsigned int)((Arg1 >> 32) & 0xFFFF);
+ ab->NumSendXfers = (unsigned int)((Arg1 >> 16) & 0xFFFF);
+ ab->NumRecvXfers = (unsigned int)((Arg1) & 0xFFFF);
+
+ // skip the uint64_t emissary id arg which is first arg in _emissary_exec.
+ ab->keyptr += sizeof(int);
+ ab->argptr += sizeof(uint64_t);
+ ab->NumArgs -= 1;
+
+ // data_not_used used for testing consistency.
+ ab->data_not_used =
+ (size_t)(ab->DataLen) - (((size_t)(3 + ab->NumArgs) * sizeof(int)) +
+ sizeof(uint64_t) + AlignFill);
+
+ // Ensure first arg after emissary id arg is aligned.
+ if (((size_t)ab->argptr) % (size_t)8) {
+ ab->argptr += 4;
+ ab->data_not_used -= 4;
+ }
+}
+
+/// Get uint32 value extended to uint64_t value from a char ptr
+static uint64_t getUInt32(char *val) {
+ uint32_t i32 = *(uint32_t *)val;
+ return (uint64_t)i32;
+}
+
+/// Get uint64_t value from a char ptr
+static uint64_t getUInt64(char *val) { return *(uint64_t *)val; }
+
+// build argument array to create call to variadic wrappers
+static uint32_t
+emissaryBuildVargs(int NumArgs, char *keyptr, char *dataptr, char *strptr,
+ unsigned long long *data_not_used, emis_argptr_t *a[],
+ std::unordered_map<void *, void *> *D2HAddrList) {
+ size_t NumBytes;
+ size_t BytesConsumed;
+ size_t StrSz;
+ size_t FillerNeeded;
+ uint32_t ArgCount = 0;
+ for (int ArgNum = 0; ArgNum < NumArgs; ArgNum++) {
+ NumBytes = 0;
+ StrSz = 0;
+ unsigned int Key = *(unsigned int *)keyptr;
+ unsigned int EmisId = Key >> 16;
+ unsigned int NumBits = (Key << 16) >> 16;
+
+ switch (EmisId) {
+ case EmisFloatTy:
+ NumBytes = NumBits / 8;
+ BytesConsumed = NumBytes;
+ FillerNeeded = ((size_t)dataptr) % NumBytes;
+ if (FillerNeeded) {
+ dataptr += FillerNeeded;
+ BytesConsumed += FillerNeeded;
+ }
+ if ((*data_not_used) < BytesConsumed)
+ return _ERC_DATA_USED_ERROR;
+
+ if (NumBytes == 4)
+ a[ArgCount] = (emis_argptr_t *)getUInt32(dataptr);
+ else {
+ double *value = (double *)dataptr;
+ a[ArgCount] = (emis_argptr_t *)(uint64_t)*value;
+ }
+ break;
+
+ case EmisIntegerTy:
+ NumBytes = NumBits / 8;
+ BytesConsumed = NumBytes;
+ FillerNeeded = ((size_t)dataptr) % NumBytes;
+ if (FillerNeeded) {
+ dataptr += FillerNeeded;
+ BytesConsumed += FillerNeeded;
+ }
+ if ((*data_not_used) < BytesConsumed)
+ return _ERC_DATA_USED_ERROR;
+
+ if (NumBytes == 4)
+ a[ArgCount] = (emis_argptr_t *)getUInt32(dataptr);
+ else
+ a[ArgCount] = (emis_argptr_t *)getUInt64(dataptr);
+ break;
+
+ case EmisPointerTy: {
+ if (NumBits == 1) { // This is a pointer to string
+ NumBytes = 4;
+ BytesConsumed = NumBytes;
+ StrSz = (size_t)*(unsigned int *)dataptr;
+ if ((*data_not_used) < BytesConsumed)
+ return _ERC_DATA_USED_ERROR;
+ a[ArgCount] = (emis_argptr_t *)((char *)strptr);
+ } else {
+ NumBytes = 8;
+ BytesConsumed = NumBytes;
+ FillerNeeded = ((size_t)dataptr) % NumBytes;
+ if (FillerNeeded) {
+ dataptr += FillerNeeded; // dataptr is now aligned
+ BytesConsumed += FillerNeeded;
+ }
+ if ((*data_not_used) < BytesConsumed)
+ return _ERC_DATA_USED_ERROR;
+ a[ArgCount] = (emis_argptr_t *)getUInt64(dataptr);
+ }
+ if (D2HAddrList) {
+ auto Found = D2HAddrList->find((void *)a[ArgCount]);
+ if (Found != D2HAddrList->end())
+ a[ArgCount] = (emis_argptr_t *)Found->second;
+ }
+ } break;
+
+ default:
+ return _ERC_INVALID_ID_ERROR;
+ }
+ // Move to next argument
+ dataptr += NumBytes;
+ strptr += StrSz;
+ *data_not_used -= BytesConsumed;
+ keyptr += 4;
+ ArgCount++;
+ }
+ return _ERC_SUCCESS;
+}
+
+// Utility to skip two args in the ArgBuf
+static void emisSkipXferArgSet(emisArgBuf_t *ab) {
+ // Skip the ptr and size of the Xfer
+ ab->NumArgs -= 2;
+ ab->keyptr += 2 * sizeof(uint32_t);
+ ab->argptr += 2 * sizeof(void *);
+ ab->data_not_used -= 2 * sizeof(void *);
+}
+
+static EmissaryReturn_t
+EmissaryTop(char *data, emisArgBuf_t *ab,
+ std::unordered_map<void *, void *> *D2HAddrList) {
+ // Registry-only dispatch (D1): every Emissary API -- MPI, HDF5, PRINT,
+ // RESERVE, and any out-of-tree client -- is serviced through the runtime
+ // registry. A client's handler is present because its library
+ // self-registered at load time. There is no per-client switch and no weak
+ // symbol fallback: an unregistered id (for example the reserved-but-unused
+ // EMIS_ID_FORTRT) is simply unsupported.
+ if (ab->emisid == EMIS_ID_INVALID) {
+ fprintf(stderr, "Emissary (host execution) got invalid EMIS_ID\n");
+ return (EmissaryReturn_t)0;
+ }
+
+ EmissaryHandler_t Handler = EmissaryLookup(ab->emisid);
+ if (Handler == nullptr) {
+ fprintf(stderr,
+ "Emissary (host execution) EMIS_ID:%d fnid:%d not supported\n",
+ ab->emisid, ab->emisfnid);
+ return (EmissaryReturn_t)0;
+ }
+
+ emis_argptr_t **args = (emis_argptr_t **)aligned_alloc(
+ sizeof(emis_argptr_t), ab->NumArgs * sizeof(emis_argptr_t *));
+
+ // Build the unpacked argument vector against a scratch copy of data_not_used
+ // so the buffer descriptor (ab) is left pristine for the handler. PRINT walks
+ // the raw buffer itself and relies on ab->data_not_used being intact; the
+ // other handlers use only the argument vector, so this is safe for all of
+ // them and keeps a single uniform dispatch path.
+ unsigned long long data_not_used = ab->data_not_used;
+ if (emissaryBuildVargs(ab->NumArgs, ab->keyptr, ab->argptr, ab->strptr,
+ &data_not_used, &args[0],
+ D2HAddrList) != _ERC_SUCCESS) {
+ free(args);
+ return (EmissaryReturn_t)0;
+ }
+
+ EmissaryReturn_t result = Handler(data, ab, args);
+ free(args);
+ return result;
+}
+
+// -----------------------------------------------------------------
+// -- Handle OFFLOAD_EMISSARY and OFFLOAD_EMISSARY_DM opcodes --
+// -- handleEmissaryImpl calls EmissaryTop for each active lane --
+// -----------------------------------------------------------------
+template <uint32_t NumLanes>
+inline RPCStatus handleEmissaryImpl(Server::Port &port) {
+
+ switch (port.get_opcode()) {
+
+ // This case handles the device function __llvm_emissary_rpc for emissary
+ // APIs that require no d2h or h2d memory transfer.
+ case OFFLOAD_EMISSARY: {
+ uint64_t Sizes[NumLanes] = {0};
+ unsigned long long Results[NumLanes] = {0};
+ void *BufPtrs[NumLanes] = {nullptr};
+ port.recv_n(BufPtrs, Sizes, [&](uint64_t Size) { return new char[Size]; });
+ uint32_t id = 0;
+ for (void *BufferPtr : BufPtrs) {
+ if (BufferPtr) {
+ emisArgBuf_t ab;
+ emisExtractArgBuf((char *)BufferPtr, &ab);
+ Results[id++] = EmissaryTop((char *)BufferPtr, &ab, nullptr);
+ }
+ }
+ port.send([&](::rpc::Buffer *Buffer, uint32_t ID) {
+ Buffer->data[0] = static_cast<uint64_t>(Results[ID]);
+ });
+ for (void *BufferPtr : BufPtrs) {
+ if (BufferPtr) {
+ delete[] reinterpret_cast<char *>(BufferPtr);
+ }
+ }
+ break;
+ }
+
+ // This case handles the device function __llvm_emissary_rpc_dm for emissary
+ // APIs require D2H or H2D transfer vectors to be processed through the port.
+ // FIXME: test with multiple transfer vectors of the same type.
+ case OFFLOAD_EMISSARY_DM: {
+ uint64_t Sizes[NumLanes] = {0};
+ unsigned long long Results[NumLanes] = {0};
+ void *BufPtrs[NumLanes] = {nullptr};
+ port.recv_n(BufPtrs, Sizes, [&](uint64_t Size) { return new char[Size]; });
+
+ uint32_t id = 0;
+ emisArgBuf_t AB[NumLanes];
+ std::unordered_map<void *, void *> D2HAddrList;
+ void *Xfers[NumLanes] = {nullptr};
+ void *DevXfers[NumLanes] = {nullptr};
+ uint64_t XferSzs[NumLanes] = {0};
+ bool SdmaHandled[NumLanes] = {false};
+ uint32_t TotalSendXfers = 0;
+ id = 0;
+
+ for (void *BufferPtr : BufPtrs) {
+ if (BufferPtr) {
+
+ emisArgBuf_t *ab = &AB[id];
+ emisExtractArgBuf((char *)BufferPtr, ab);
+ unsigned long long SdmaResult = 0;
+ if (emissary_mpi_sdma_try_dm_buffer((char *)BufferPtr, &SdmaResult) ==
+ 0) {
+ Results[id] = SdmaResult;
+ SdmaHandled[id] = true;
+ id++;
+ continue;
+ }
+ for (uint32_t idx = 0; idx < ab->NumSendXfers; idx++) {
+ TotalSendXfers++;
+ DevXfers[id] = (void *)*((uint64_t *)ab->argptr);
+ XferSzs[id] = (size_t)*((size_t *)(ab->argptr + sizeof(void *)));
+ emisSkipXferArgSet(ab);
+ }
+ // Allocate the host space for the receive Xfers
+ for (uint32_t idx = 0; idx < ab->NumRecvXfers; idx++) {
+ void *DevAddr = (void *)*((uint64_t *)ab->argptr);
+ size_t DevSz = (((size_t)*((size_t *)(ab->argptr + sizeof(void *)))) &
+ 0x00000000FFFFFFFF);
+ void *HostAddr = new char[DevSz];
+ D2HAddrList.insert(std::pair<void *, void *>(DevAddr, HostAddr));
+ emisSkipXferArgSet(ab);
+ }
+ id++;
+ }
+ }
+
+ // recv_n for device send_n into new host-allocated Xfers
+ if (TotalSendXfers)
+ port.recv_n(Xfers, XferSzs,
+ [&](uint64_t Size) { return new char[Size]; });
+
+ // Xfers now contains just allocated host addrs for sends and
+ // DevXfers contains corresponding DevAddr for those sends
+ // Build map to pass to Emissary
+ id = 0;
+ for (void *Xfer : Xfers) {
+ if (Xfer) {
+ D2HAddrList.insert(std::pair<void *, void *>(DevXfers[id], Xfer));
+ id++;
+ }
+ }
+
+ // Call EmissaryTop for each active lane
+ id = 0;
+ for (void *BufferPtr : BufPtrs) {
+ if (BufferPtr) {
+ if (SdmaHandled[id]) {
+ id++;
+ continue;
+ }
+ emisArgBuf_t *ab = &AB[id];
+ emisExtractArgBuf((char *)BufferPtr, ab);
+ for (uint32_t idx = 0; idx < ab->NumSendXfers; idx++)
+ emisSkipXferArgSet(ab);
+ for (uint32_t idx = 0; idx < ab->NumRecvXfers; idx++)
+ emisSkipXferArgSet(ab);
+ Results[id] = EmissaryTop((char *)BufferPtr, ab, &D2HAddrList);
+ id++;
+ }
+ }
+
+ // Process send_n for the H2D Xfers.
+ void *recvXfers[NumLanes] = {nullptr};
+ uint64_t recvXferSzs[NumLanes] = {0};
+ id = 0;
+ uint32_t TotalRecvXfers = 0;
+ for (void *BufferPtr : BufPtrs) {
+ if (BufferPtr) {
+ if (SdmaHandled[id]) {
+ id++;
+ continue;
+ }
+ emisArgBuf_t *ab = &AB[id];
+ // Reset ArgBuf tracker
+ emisExtractArgBuf((char *)BufferPtr, ab);
+ for (uint32_t idx = 0; idx < ab->NumSendXfers; idx++)
+ emisSkipXferArgSet(ab);
+ for (uint32_t idx = 0; idx < ab->NumRecvXfers; idx++) {
+ TotalRecvXfers++;
+ void *DevAddr = (void *)*((uint64_t *)ab->argptr);
+ recvXfers[id] = D2HAddrList[DevAddr];
+ recvXferSzs[id] =
+ (((uint64_t)*((size_t *)(ab->argptr + sizeof(void *)))) &
+ 0x00000000FFFFFFFF);
+ emisSkipXferArgSet(ab);
+ }
+ id++;
+ }
+ }
+ if (TotalRecvXfers)
+ port.send_n(recvXfers, recvXferSzs);
+
+ // Cleanup all host allocated transfer buffers
+ id = 0;
+ for (void *BufferPtr : BufPtrs) {
+ if (BufferPtr) {
+ if (SdmaHandled[id]) {
+ id++;
+ continue;
+ }
+ emisArgBuf_t *ab = &AB[id];
+ // Reset the ArgBuf tracker ab
+ emisExtractArgBuf((char *)BufferPtr, ab);
+ // Cleanup host allocated send Xfers
+ for (uint32_t idx = 0; idx < ab->NumSendXfers; idx++) {
+ void *DevAddr = (void *)*((uint64_t *)ab->argptr);
+ void *HostAddr = D2HAddrList[DevAddr];
+ delete[] reinterpret_cast<char *>(HostAddr);
+ emisSkipXferArgSet(ab);
+ }
+ // Cleanup host allocated bufs
+ for (uint32_t idx = 0; idx < ab->NumRecvXfers; idx++) {
+ void *DevAddr = (void *)*((uint64_t *)ab->argptr);
+ void *HostAddr = D2HAddrList[DevAddr];
+ delete[] reinterpret_cast<char *>(HostAddr);
+ emisSkipXferArgSet(ab);
+ }
+ id++;
+ }
+ }
+
+ port.send([&](::rpc::Buffer *Buffer, uint32_t ID) {
+ Buffer->data[0] = static_cast<uint64_t>(Results[ID]);
+ delete[] reinterpret_cast<char *>(BufPtrs[ID]);
+ });
+
+ break;
+ } // END CASE OFFLOAD_EMISSARY_DM
+
+ default: {
+ return ::rpc::RPC_UNHANDLED_OPCODE;
+ break;
+ }
+ }
+ return ::rpc::RPC_SUCCESS;
+} // end handleEmissaryImpl
+
+} // namespace internal
+
+// Handles any opcode generated from emissary client code.
+inline RPCStatus handleEmissaryOpcodes(Server::Port &port, uint32_t num_lanes) {
+ switch (num_lanes) {
+ case 1:
+ return internal::handleEmissaryImpl<1>(port);
+ case 32:
+ return internal::handleEmissaryImpl<32>(port);
+ case 64:
+ return internal::handleEmissaryImpl<64>(port);
+ default:
+ return RPC_ERROR;
+ }
+}
+
+} // namespace rpc
+
+#endif // LLVM_LIBC_SRC___SUPPORT_RPC_EMISSARY_RPC_SERVER_H
diff --git a/libc/shared/rpc_util.h b/libc/shared/rpc_util.h
index 5fabeb8069f5b..0f5fc4737983d 100644
--- a/libc/shared/rpc_util.h
+++ b/libc/shared/rpc_util.h
@@ -343,7 +343,7 @@ RPC_ATTRS void sleep_briefly() {
#if __has_builtin(__nvvm_reflect)
if (__nvvm_reflect("__CUDA_ARCH") >= 700)
asm("nanosleep.u32 64;" ::: "memory");
-#elif __has_builtin(__builtin_amdgcn_s_sleep)
+#elif __has_builtin(__builtin_amdgcn_s_sleep) && defined(RPC_TARGET_IS_GPU)
__builtin_amdgcn_s_sleep(2);
#elif __has_builtin(__builtin_ia32_pause)
__builtin_ia32_pause();
diff --git a/libc/src/__support/RPC/CMakeLists.txt b/libc/src/__support/RPC/CMakeLists.txt
index cac9c4e05e369..876bd221dab03 100644
--- a/libc/src/__support/RPC/CMakeLists.txt
+++ b/libc/src/__support/RPC/CMakeLists.txt
@@ -6,6 +6,7 @@ add_object_library(
rpc_client
SRCS
rpc_client.cpp
+ emissary_device_utils.cpp
HDRS
rpc_client.h
DEPENDS
diff --git a/libc/src/__support/RPC/emissary_device_utils.cpp b/libc/src/__support/RPC/emissary_device_utils.cpp
new file mode 100644
index 0000000000000..0b1c0aa15ebf7
--- /dev/null
+++ b/libc/src/__support/RPC/emissary_device_utils.cpp
@@ -0,0 +1,104 @@
+//===- emissary_device_utils.cpp - utils for Emissary APIs ------*- 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Device functions emitted by clang/lib/CodeGen/CGEmitEmissaryExec.cpp
+//
+//===----------------------------------------------------------------------===//
+
+#include "EmissaryIds.h"
+#include "rpc_client.h"
+#include "shared/rpc.h"
+#include "src/__support/macros/config.h"
+#include "src/stdlib/free.h"
+#include "src/stdlib/malloc.h"
+
+extern "C" {
+
+#ifdef __NVPTX__
+[[gnu::leaf]] void *malloc(size_t Size);
+[[gnu::leaf]] void free(void *Ptr);
+#endif
+
+// The clang compiler will generate calls to __strlen_max when string length
+// is not compile time constant.
+uint32_t __strlen_max(const char *InStr, uint32_t MaxStrLen) {
+ if (InStr == 0) // encountered a null pointer to string
+ return 0;
+ for (uint32_t I = 0; I < MaxStrLen; I++)
+ if (InStr[I] == (char)0)
+ return (uint32_t)(I + 1);
+ return MaxStrLen;
+}
+
+void *__llvm_emissary_premalloc(uint32_t Sz) {
+#ifdef __NVPTX__
+ return malloc((size_t)Sz);
+#else
+ return LIBC_NAMESPACE::malloc((size_t)Sz);
+#endif
+}
+unsigned long long __llvm_emissary_rpc(uint32_t Sz32, void *BufData) {
+ rpc::Client::Port Port = LIBC_NAMESPACE::rpc::client.open<OFFLOAD_EMISSARY>();
+ Port.send_n(BufData, (size_t)Sz32);
+ unsigned long long Ret;
+ Port.recv([&](rpc::Buffer *Buffer, uint32_t) {
+ Ret = static_cast<unsigned long long>(Buffer->data[0]);
+ });
+#ifdef __NVPTX__
+ free(BufData);
+#else
+ LIBC_NAMESPACE::free(BufData);
+#endif
+ return Ret;
+}
+
+// This is for emissary APIs that require d2h or h2d memory transfers.
+unsigned long long __llvm_emissary_rpc_dm(uint32_t Sz32, void *BufData) {
+ rpc::Client::Port Port =
+ LIBC_NAMESPACE::rpc::client.open<OFFLOAD_EMISSARY_DM>();
+ Port.send_n(BufData, (size_t)Sz32);
+ char *Data = (char *)BufData;
+ uint32_t *Int32Data = (uint32_t *)Data;
+ uint32_t NumArgs = Int32Data[1];
+ char *KeyPtr = Data + (2 * sizeof(int));
+ char *ArgPtr = KeyPtr + (NumArgs * sizeof(int));
+ if (((size_t)ArgPtr) % (size_t)8)
+ ArgPtr += 4; // ArgPtr must be aligned
+ uint64_t Arg1 = *(uint64_t *)ArgPtr;
+ uint32_t NumSendXfers = (unsigned int)((Arg1 >> 16) & 0xFFFF);
+ uint32_t NumRecvXfers = (unsigned int)((Arg1) & 0xFFFF);
+ // Skip by Arg1 and process Send and Recv Xfers if any
+ ArgPtr += sizeof(uint64_t);
+ for (uint32_t idx = 0; idx < NumSendXfers; idx++) {
+ void *D2HData = (void *)*((uint64_t *)ArgPtr);
+ ArgPtr += sizeof(void *);
+ size_t D2HSize = ((size_t)*((size_t *)ArgPtr) & 0x00000000FFFFFFFF);
+ ArgPtr += sizeof(size_t);
+ Port.send_n(D2HData, D2HSize);
+ }
+ for (uint32_t idx = 0; idx < NumRecvXfers; idx++) {
+ void *H2DData = (void *)*((uint64_t *)ArgPtr);
+ ArgPtr += sizeof(void *);
+ ArgPtr += sizeof(size_t);
+ uint64_t RecvSize;
+ void *Buf = nullptr;
+ Port.recv_n(&Buf, &RecvSize,
+ [&](uint64_t) { return reinterpret_cast<void *>(H2DData); });
+ }
+ unsigned long long Ret;
+ Port.recv([&](rpc::Buffer *Buffer, uint32_t) {
+ Ret = static_cast<unsigned long long>(Buffer->data[0]);
+ });
+#ifdef __NVPTX__
+ free(BufData);
+#else
+ LIBC_NAMESPACE::free(BufData);
+#endif
+ return Ret;
+}
+} // end extern "C"
diff --git a/libc/test/shared/CMakeLists.txt b/libc/test/shared/CMakeLists.txt
index 74eda7cdaf065..8832b7738d749 100644
--- a/libc/test/shared/CMakeLists.txt
+++ b/libc/test/shared/CMakeLists.txt
@@ -917,3 +917,23 @@ add_fp_unittest(
libc.src.__support.CPP.array
libc.src.__support.FPUtil.fp_bits
)
+
+# Emissary runtime handler registry test.
+# include <EmissaryIds.h> from the Clang resource headers source tree, so add
+# that directory to the include path. Skip if it is not present (e.g. a libc
+# checkout without the clang component).
+set(_emissary_ids_dir "${LIBC_SOURCE_DIR}/../clang/lib/Headers")
+if(EXISTS "${_emissary_ids_dir}/EmissaryIds.h")
+ add_libc_test(
+ emissary_registry_test
+ UNIT_TEST_ONLY
+ SUITE
+ libc-shared-tests
+ SRCS
+ emissary_registry_test.cpp
+ COMPILE_OPTIONS
+ -I${_emissary_ids_dir}
+ DEPENDS
+ libc.src.__support.CPP.array
+ )
+endif()
diff --git a/libc/test/shared/emissary_registry_test.cpp b/libc/test/shared/emissary_registry_test.cpp
new file mode 100644
index 0000000000000..66cf82d007822
--- /dev/null
+++ b/libc/test/shared/emissary_registry_test.cpp
@@ -0,0 +1,63 @@
+//===-- Unittests for the Emissary host handler registry ------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "shared/emissary_rpc_server.h"
+#include "test/UnitTest/Test.h"
+
+namespace {
+
+// Distinct handler bodies so lookups can be told apart by return value. The
+// bodies are never actually executed against a real RPC buffer here; the tests
+// only compare function pointers and invoke them with a stub descriptor.
+EmissaryReturn_t handlerA(char *, emisArgBuf_t *, emis_argptr_t *[]) {
+ return 0xA;
+}
+EmissaryReturn_t handlerB(char *, emisArgBuf_t *, emis_argptr_t *[]) {
+ return 0xB;
+}
+
+// Use ids near the top of the range to avoid colliding with real
+// offload_emis_id_t values a co-linked client might register.
+constexpr unsigned int kIdA = EMISSARY_MAX_REGISTERED_IDS - 2;
+constexpr unsigned int kIdB = EMISSARY_MAX_REGISTERED_IDS - 3;
+
+} // namespace
+
+TEST(LlvmLibcEmissaryRegistryTest, LookupUnregisteredIsNull) {
+ EXPECT_EQ(EmissaryLookup(kIdA), static_cast<EmissaryHandler_t>(nullptr));
+}
+
+TEST(LlvmLibcEmissaryRegistryTest, RegisterThenLookup) {
+ ASSERT_TRUE(EmissaryRegister(kIdA, &handlerA));
+ EmissaryHandler_t Got = EmissaryLookup(kIdA);
+ ASSERT_TRUE(Got == &handlerA);
+
+ // The looked-up handler is really callable and is the one we stored.
+ emisArgBuf_t Ab = {};
+ EXPECT_EQ(Got(nullptr, &Ab, nullptr), static_cast<EmissaryReturn_t>(0xA));
+}
+
+TEST(LlvmLibcEmissaryRegistryTest, IdempotentReregisterSucceeds) {
+ ASSERT_TRUE(EmissaryRegister(kIdB, &handlerB));
+ // Same id, same handler: allowed.
+ EXPECT_TRUE(EmissaryRegister(kIdB, &handlerB));
+ // Same id, different handler: rejected, original preserved.
+ EXPECT_FALSE(EmissaryRegister(kIdB, &handlerA));
+ EXPECT_TRUE(EmissaryLookup(kIdB) == &handlerB);
+}
+
+TEST(LlvmLibcEmissaryRegistryTest, RejectsNullHandler) {
+ EXPECT_FALSE(EmissaryRegister(EMISSARY_MAX_REGISTERED_IDS - 4, nullptr));
+}
+
+TEST(LlvmLibcEmissaryRegistryTest, RejectsOutOfRangeId) {
+ EXPECT_FALSE(EmissaryRegister(EMISSARY_MAX_REGISTERED_IDS, &handlerA));
+ EXPECT_FALSE(EmissaryRegister(EMISSARY_MAX_REGISTERED_IDS + 100, &handlerA));
+ EXPECT_EQ(EmissaryLookup(EMISSARY_MAX_REGISTERED_IDS),
+ static_cast<EmissaryHandler_t>(nullptr));
+}
diff --git a/offload/liboffload/exports b/offload/liboffload/exports
index 168341aa7d938..e029bbe15cde9 100644
--- a/offload/liboffload/exports
+++ b/offload/liboffload/exports
@@ -1,6 +1,7 @@
VERS1.0 {
global:
ol*;
+ EmissaryRegistryTable;
local:
*;
};
diff --git a/offload/libomptarget/exports b/offload/libomptarget/exports
index 1831c43cc5f29..64dbc3479613c 100644
--- a/offload/libomptarget/exports
+++ b/offload/libomptarget/exports
@@ -83,6 +83,7 @@ VERS1.0 {
__llvmPushCallConfiguration;
__llvmPopCallConfiguration;
llvmLaunchKernel;
+ EmissaryRegistryTable;
local:
*;
};
diff --git a/offload/plugins-nextgen/common/src/RPC.cpp b/offload/plugins-nextgen/common/src/RPC.cpp
index 7c03c916058fb..2d896b6231cc3 100644
--- a/offload/plugins-nextgen/common/src/RPC.cpp
+++ b/offload/plugins-nextgen/common/src/RPC.cpp
@@ -13,6 +13,7 @@
#include "PluginInterface.h"
+#include "shared/emissary_rpc_server.h"
#include "shared/rpc.h"
#include "shared/rpc_opcodes.h"
#include "shared/rpc_server.h"
@@ -111,6 +112,9 @@ runServer(plugin::GenericDeviceTy &Device, void *Buffer,
if (Status == rpc::RPC_UNHANDLED_OPCODE)
Status = rpc::handle_libc_opcodes(*Port, NumLanes);
+ if (Status == rpc::RPC_UNHANDLED_OPCODE)
+ Status = rpc::handleEmissaryOpcodes(*Port, NumLanes);
+
return Status;
}
>From 1f078c1c359553d7fabf63d63f9b06ff1a94d3cf Mon Sep 17 00:00:00 2001
From: gregrodgers <Gregory.Rodgers at amd.com>
Date: Mon, 24 Aug 2026 10:58:00 -0500
Subject: [PATCH 2/8] [OpenMP] test case demo of an Emissary API.
---
offload/test/offloading/emissary.cpp | 112 +++++++++++++++++++++++++++
1 file changed, 112 insertions(+)
create mode 100644 offload/test/offloading/emissary.cpp
diff --git a/offload/test/offloading/emissary.cpp b/offload/test/offloading/emissary.cpp
new file mode 100644
index 0000000000000..a2c081712b82c
--- /dev/null
+++ b/offload/test/offloading/emissary.cpp
@@ -0,0 +1,112 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic -I %S/../../../libc
+// RUN: env LIBOMPTARGET_INFO=16 \
+// RUN: %libomptarget-run-generic 2>&1
+
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// REQUIRES: gpu
+// XFAIL: intelgpu
+
+// ---------- These are the functions in our host library ----------
+#if (!defined(__NVPTX__) && !defined(__AMDGCN__))
+extern "C" int foo(int x, double y, int *iarray) { return (int)y * iarray[0]; }
+extern "C" double bar(int x, double y, double *darray) { return y * darray[0]; }
+#endif
+
+// ---------- Emissary API definition for foobar_openmp library ----------
+// This 4-part definition file is typically in it's own file:
+//
+// 1. === Includes, always include EmissaryIds.h
+#include <EmissaryIds.h>
+#include <stdarg.h>
+
+// 2.=== Enum with index for each function provided by Emissary API.
+typedef enum {
+ _RESERVE_INVALID, // recommend 0 is INVALID
+ _RESERVE_foo_idx,
+ _RESERVE_bar_idx,
+} offload_emis_rsrv_t;
+
+// 3. === Device Stubs for each function in the API
+// This section ONLY for device compilation
+#if (defined(__NVPTX__) || defined(__AMDGCN__))
+extern "C" int foo(int x, double y, int *iarray) {
+ return (int)_emissary_exec(
+ _PACK_EMIS_IDS(EMIS_ID_RESERVE, _RESERVE_foo_idx, 0, 0), x, y, iarray); }
+extern "C" double bar(int x, double y, double *darray) {
+ return (double)_emissary_exec(
+ _PACK_EMIS_IDS(EMIS_ID_RESERVE, _RESERVE_bar_idx, 0, 0), x, y, darray); }
+
+#else // end device stub definitions
+
+// 4. === Define host selector function for Emissary API reserve
+// Section 4 is only compiled on host pass
+#include <cstdint>
+#include <shared/emissary_rpc_server.h>
+#define _PTR_TO_64BIT_ (unsigned long long int)
+// This is the EmissaryReserve selector function. It is called when the emissary
+// runtime sees EMIS_ID_RESERVE as the API identifier. It is dispatched through
+// the runtime registry (see the self-registration constructor below)
+// This function invokes host function based on function index (emisfnid).
+extern "C" EmissaryReturn_t EmissaryReserve(char *data, emisArgBuf_t *ab,
+ emis_argptr_t *a[]) {
+ switch (ab->emisfnid) {
+
+ case _RESERVE_foo_idx: {
+ return (EmissaryReturn_t) foo (
+ (int)(_PTR_TO_64BIT_ a[0]),
+ (double)(_PTR_TO_64BIT_ a[1]),
+ (int *)(_PTR_TO_64BIT_ a[2]));
+ }
+
+ case _RESERVE_bar_idx: {
+ return (EmissaryReturn_t) bar (
+ (int)(_PTR_TO_64BIT_ a[0]),
+ (double)(_PTR_TO_64BIT_ a[1]),
+ (double *)(_PTR_TO_64BIT_ a[2]));
+ }
+
+ } // end switch statement
+ return (EmissaryReturn_t)0;
+} // end EmissaryReserve function selector
+
+// Self-register the host selector fn at load time so the RPC server
+// dispatches EMIS_ID_RESERVE through the runtime registry
+extern "C" __attribute__((constructor)) void
+emissary_reserve_self_register(void) {
+ EmissaryRegister(EMIS_ID_RESERVE, &EmissaryReserve);
+}
+
+#undef _PTR_TO_64BIT_
+#endif
+//== End section 4 and end of Emissary API definition for foobar_openmp library
+
+// ---------- Demo app using foobar_openmp lib on host AND device ----------
+#define VSIZE 10
+#include <stdio.h>
+int main(int argc, char *argv[]) {
+ double yfoo = 2.0;
+ int iarray[2] = {4,42};
+ int foo_rc = foo(-1, yfoo, iarray);
+
+ double ybar = 3.0;
+ double darray[2] = {4.0 , 42.0};;
+ double bar_rc = bar(-2, ybar, darray);
+ printf("MAIN foo_rc:%d bar_rc:%f\n",foo_rc, bar_rc);
+ foo_rc = 1; bar_rc=1;
+
+ printf("PREREGION foo_rc:%d bar_rc:%f yfoo:%f \n",foo_rc, bar_rc, yfoo);
+#pragma omp target teams distribute parallel for map(to:yfoo,ybar) map(from: foo_rc,bar_rc) is_device_ptr(iarray, darray)
+ for (int i = 0; i < VSIZE; i++) {
+ foo_rc = foo(i, yfoo, iarray);
+ bar_rc = bar(i, ybar, darray);
+ }
+ printf("POSTREGION foo_rc:%d bar_rc:%f yfoo:%f \n",foo_rc, bar_rc, yfoo);
+ int rc = 0;
+ if (foo_rc != 8 )
+ rc = 1;
+ if (bar_rc != 12.0 )
+ rc = 2;
+ return rc;
+}
>From 4bb1f17a536e21a12d2417387c0928fbaacaaf89 Mon Sep 17 00:00:00 2001
From: gregrodgers <Gregory.Rodgers at amd.com>
Date: Mon, 24 Aug 2026 13:53:42 -0500
Subject: [PATCH 3/8] [OPENMP] add rpc_server.h when including
emissary_rpc_server.h
---
libc/test/shared/emissary_registry_test.cpp | 1 +
offload/test/offloading/emissary.cpp | 1 +
2 files changed, 2 insertions(+)
diff --git a/libc/test/shared/emissary_registry_test.cpp b/libc/test/shared/emissary_registry_test.cpp
index 66cf82d007822..8f15b292ae9e1 100644
--- a/libc/test/shared/emissary_registry_test.cpp
+++ b/libc/test/shared/emissary_registry_test.cpp
@@ -7,6 +7,7 @@
//===----------------------------------------------------------------------===//
#include "shared/emissary_rpc_server.h"
+#include "shared/rpc_server.h"
#include "test/UnitTest/Test.h"
namespace {
diff --git a/offload/test/offloading/emissary.cpp b/offload/test/offloading/emissary.cpp
index a2c081712b82c..dd1579f5299ee 100644
--- a/offload/test/offloading/emissary.cpp
+++ b/offload/test/offloading/emissary.cpp
@@ -43,6 +43,7 @@ extern "C" double bar(int x, double y, double *darray) {
// 4. === Define host selector function for Emissary API reserve
// Section 4 is only compiled on host pass
#include <cstdint>
+#include <shared/rpc_server.h>
#include <shared/emissary_rpc_server.h>
#define _PTR_TO_64BIT_ (unsigned long long int)
// This is the EmissaryReserve selector function. It is called when the emissary
>From 94dc95b3eb2f1a17dade2e6944db4eabf1febc9f Mon Sep 17 00:00:00 2001
From: gregrodgers <Gregory.Rodgers at amd.com>
Date: Mon, 24 Aug 2026 14:42:07 -0500
Subject: [PATCH 4/8] [Emissary] remove emissary_registry_test because it is
malformed and also covered in offloading test
---
libc/test/shared/CMakeLists.txt | 20 -------
libc/test/shared/emissary_registry_test.cpp | 64 ---------------------
2 files changed, 84 deletions(-)
delete mode 100644 libc/test/shared/emissary_registry_test.cpp
diff --git a/libc/test/shared/CMakeLists.txt b/libc/test/shared/CMakeLists.txt
index 8832b7738d749..74eda7cdaf065 100644
--- a/libc/test/shared/CMakeLists.txt
+++ b/libc/test/shared/CMakeLists.txt
@@ -917,23 +917,3 @@ add_fp_unittest(
libc.src.__support.CPP.array
libc.src.__support.FPUtil.fp_bits
)
-
-# Emissary runtime handler registry test.
-# include <EmissaryIds.h> from the Clang resource headers source tree, so add
-# that directory to the include path. Skip if it is not present (e.g. a libc
-# checkout without the clang component).
-set(_emissary_ids_dir "${LIBC_SOURCE_DIR}/../clang/lib/Headers")
-if(EXISTS "${_emissary_ids_dir}/EmissaryIds.h")
- add_libc_test(
- emissary_registry_test
- UNIT_TEST_ONLY
- SUITE
- libc-shared-tests
- SRCS
- emissary_registry_test.cpp
- COMPILE_OPTIONS
- -I${_emissary_ids_dir}
- DEPENDS
- libc.src.__support.CPP.array
- )
-endif()
diff --git a/libc/test/shared/emissary_registry_test.cpp b/libc/test/shared/emissary_registry_test.cpp
deleted file mode 100644
index 8f15b292ae9e1..0000000000000
--- a/libc/test/shared/emissary_registry_test.cpp
+++ /dev/null
@@ -1,64 +0,0 @@
-//===-- Unittests for the Emissary host handler registry ------------------===//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-
-#include "shared/emissary_rpc_server.h"
-#include "shared/rpc_server.h"
-#include "test/UnitTest/Test.h"
-
-namespace {
-
-// Distinct handler bodies so lookups can be told apart by return value. The
-// bodies are never actually executed against a real RPC buffer here; the tests
-// only compare function pointers and invoke them with a stub descriptor.
-EmissaryReturn_t handlerA(char *, emisArgBuf_t *, emis_argptr_t *[]) {
- return 0xA;
-}
-EmissaryReturn_t handlerB(char *, emisArgBuf_t *, emis_argptr_t *[]) {
- return 0xB;
-}
-
-// Use ids near the top of the range to avoid colliding with real
-// offload_emis_id_t values a co-linked client might register.
-constexpr unsigned int kIdA = EMISSARY_MAX_REGISTERED_IDS - 2;
-constexpr unsigned int kIdB = EMISSARY_MAX_REGISTERED_IDS - 3;
-
-} // namespace
-
-TEST(LlvmLibcEmissaryRegistryTest, LookupUnregisteredIsNull) {
- EXPECT_EQ(EmissaryLookup(kIdA), static_cast<EmissaryHandler_t>(nullptr));
-}
-
-TEST(LlvmLibcEmissaryRegistryTest, RegisterThenLookup) {
- ASSERT_TRUE(EmissaryRegister(kIdA, &handlerA));
- EmissaryHandler_t Got = EmissaryLookup(kIdA);
- ASSERT_TRUE(Got == &handlerA);
-
- // The looked-up handler is really callable and is the one we stored.
- emisArgBuf_t Ab = {};
- EXPECT_EQ(Got(nullptr, &Ab, nullptr), static_cast<EmissaryReturn_t>(0xA));
-}
-
-TEST(LlvmLibcEmissaryRegistryTest, IdempotentReregisterSucceeds) {
- ASSERT_TRUE(EmissaryRegister(kIdB, &handlerB));
- // Same id, same handler: allowed.
- EXPECT_TRUE(EmissaryRegister(kIdB, &handlerB));
- // Same id, different handler: rejected, original preserved.
- EXPECT_FALSE(EmissaryRegister(kIdB, &handlerA));
- EXPECT_TRUE(EmissaryLookup(kIdB) == &handlerB);
-}
-
-TEST(LlvmLibcEmissaryRegistryTest, RejectsNullHandler) {
- EXPECT_FALSE(EmissaryRegister(EMISSARY_MAX_REGISTERED_IDS - 4, nullptr));
-}
-
-TEST(LlvmLibcEmissaryRegistryTest, RejectsOutOfRangeId) {
- EXPECT_FALSE(EmissaryRegister(EMISSARY_MAX_REGISTERED_IDS, &handlerA));
- EXPECT_FALSE(EmissaryRegister(EMISSARY_MAX_REGISTERED_IDS + 100, &handlerA));
- EXPECT_EQ(EmissaryLookup(EMISSARY_MAX_REGISTERED_IDS),
- static_cast<EmissaryHandler_t>(nullptr));
-}
>From 6bfc7ad0c20827061b7d5e68bcfbf5715f1c36f5 Mon Sep 17 00:00:00 2001
From: gregrodgers <Gregory.Rodgers at amd.com>
Date: Mon, 24 Aug 2026 15:05:47 -0500
Subject: [PATCH 5/8] [EMISSARY] make include of EmissaryIds.h work before
compiler is installed
---
libc/src/__support/RPC/emissary_device_utils.cpp | 5 +++++
offload/test/offloading/emissary.cpp | 6 +++++-
2 files changed, 10 insertions(+), 1 deletion(-)
diff --git a/libc/src/__support/RPC/emissary_device_utils.cpp b/libc/src/__support/RPC/emissary_device_utils.cpp
index 0b1c0aa15ebf7..4584020d5a429 100644
--- a/libc/src/__support/RPC/emissary_device_utils.cpp
+++ b/libc/src/__support/RPC/emissary_device_utils.cpp
@@ -10,7 +10,12 @@
//
//===----------------------------------------------------------------------===//
+#if __has_include("../../../clang/lib/Headers/EmissaryIds.h")
+#include "../../../clang/lib/Headers/EmissaryIds.h"
+#else
#include "EmissaryIds.h"
+#endif
+
#include "rpc_client.h"
#include "shared/rpc.h"
#include "src/__support/macros/config.h"
diff --git a/offload/test/offloading/emissary.cpp b/offload/test/offloading/emissary.cpp
index dd1579f5299ee..37d0e8c487c06 100644
--- a/offload/test/offloading/emissary.cpp
+++ b/offload/test/offloading/emissary.cpp
@@ -18,7 +18,11 @@ extern "C" double bar(int x, double y, double *darray) { return y * darray[0]; }
// This 4-part definition file is typically in it's own file:
//
// 1. === Includes, always include EmissaryIds.h
-#include <EmissaryIds.h>
+#if __has_include("../../../clang/lib/Headers/EmissaryIds.h")
+#include "../clang/lib/Headers/EmissaryIds.h"
+#else
+#include "EmissaryIds.h"
+#endif
#include <stdarg.h>
// 2.=== Enum with index for each function provided by Emissary API.
>From f9cd8991142be193d135787a708fac648d133dc3 Mon Sep 17 00:00:00 2001
From: gregrodgers <Gregory.Rodgers at amd.com>
Date: Mon, 24 Aug 2026 15:55:50 -0500
Subject: [PATCH 6/8] [EMISSARY] cleanup CI warning on exhaustive switch
---
clang/lib/CodeGen/CGEmitEmissaryExec.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CodeGen/CGEmitEmissaryExec.cpp b/clang/lib/CodeGen/CGEmitEmissaryExec.cpp
index 8f269608e2c95..6a22e05a55318 100644
--- a/clang/lib/CodeGen/CGEmitEmissaryExec.cpp
+++ b/clang/lib/CodeGen/CGEmitEmissaryExec.cpp
@@ -168,9 +168,9 @@ static EmisTyID getEmisTyID(llvm::Type::TypeID TyId) {
case llvm::Type::TypedPointerTyID: ///< Typed pointer used by some GPU targets
case llvm::Type::TargetExtTyID: ///< Target extension type
return EmisInvalidTy;
- default:
- return EmisInvalidTy;
+ break;
}
+ return EmisInvalidTy;
}
// ----- External function EmitEmissaryExec called from CGExpr.cpp -----
>From bb95347401838fbfe4ee9e095ee4a38e97ef6620 Mon Sep 17 00:00:00 2001
From: gregrodgers <Gregory.Rodgers at amd.com>
Date: Mon, 24 Aug 2026 16:00:44 -0500
Subject: [PATCH 7/8] [EMISSARY] clang format previous change
---
clang/lib/CodeGen/CGEmitEmissaryExec.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/CodeGen/CGEmitEmissaryExec.cpp b/clang/lib/CodeGen/CGEmitEmissaryExec.cpp
index 6a22e05a55318..0ffefbce6bc39 100644
--- a/clang/lib/CodeGen/CGEmitEmissaryExec.cpp
+++ b/clang/lib/CodeGen/CGEmitEmissaryExec.cpp
@@ -168,7 +168,7 @@ static EmisTyID getEmisTyID(llvm::Type::TypeID TyId) {
case llvm::Type::TypedPointerTyID: ///< Typed pointer used by some GPU targets
case llvm::Type::TargetExtTyID: ///< Target extension type
return EmisInvalidTy;
- break;
+ break;
}
return EmisInvalidTy;
}
>From 430e0397eb48edb66b79e3fdefe97a6b7d2a9afd Mon Sep 17 00:00:00 2001
From: gregrodgers <Gregory.Rodgers at amd.com>
Date: Tue, 25 Aug 2026 07:18:39 -0500
Subject: [PATCH 8/8] [EMISSARY] Removed printf test. This was a new file but
the test has nothing to do with Emissary API support
---
...gcn_target_printf_unknown_size_arguments.c | 51 -------------------
1 file changed, 51 deletions(-)
delete mode 100644 clang/test/OpenMP/amdgcn_target_printf_unknown_size_arguments.c
diff --git a/clang/test/OpenMP/amdgcn_target_printf_unknown_size_arguments.c b/clang/test/OpenMP/amdgcn_target_printf_unknown_size_arguments.c
deleted file mode 100644
index 09c6511a427c4..0000000000000
--- a/clang/test/OpenMP/amdgcn_target_printf_unknown_size_arguments.c
+++ /dev/null
@@ -1,51 +0,0 @@
-// REQUIRES: amdgpu-registered-target
-// REQUIRES: x86-registered-target
-
-// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
-// RUN: %clang_cc1 -verify -fopenmp -x c -triple amdgcn-amd-amdhsa -fopenmp-is-device -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-host-ir-file-path %t-host.bc -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK
-// expected-no-diagnostics
-
-extern int printf(const char *, ...);
-
-int CheckMultipleArgs(int a) {
- char *test = "testing";
- char *t;
-#pragma omp target private(t)
- {
- t = test + a;
- printf("%s %d %s", t, 21, test);
-// CHECK-LABEL: define weak_odr protected amdgpu_kernel void @{{.*}}CheckMultipleArgs
-// CHECK: entry:
-// CHECK: [[TEST_ADDR:%[a-zA-Z0-9_.]+]] = alloca ptr, align 8, addrspace(5)
-// CHECK: [[A_ADDR:%[a-zA-Z0-9_.]+]] = alloca i64, align 8, addrspace(5)
-// CHECK: [[DYN_PTR_ADDR:%[a-zA-Z0-9_.]+]] = alloca ptr, align 8, addrspace(5)
-// CHECK: [[T_ADDR:%[a-zA-Z0-9_.]+]] = alloca ptr, align 8, addrspace(5)
-// CHECK: [[TEST_CAST:%[a-zA-Z0-9_.]+]] = addrspacecast ptr addrspace(5) [[TEST_ADDR]] to ptr
-// CHECK: [[A_CAST:%[a-zA-Z0-9_.]+]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
-// CHECK: [[DYN_PTR_CAST:%[a-zA-Z0-9_.]+]] = addrspacecast ptr addrspace(5) [[DYN_PTR_ADDR]] to ptr
-// CHECK: [[T_CAST:%[a-zA-Z0-9_.]+]] = addrspacecast ptr addrspace(5) [[T_ADDR]] to ptr
-// CHECK: store ptr %test, ptr [[TEST_CAST]], align 8
-// CHECK: store i64 %a, ptr [[A_CAST]], align 8
-// CHECK: store ptr %dyn_ptr, ptr [[DYN_PTR_CAST]], align 8
-// CHECK: [[INIT_CALL:%[a-zA-Z0-9_.]+]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) {{.*}} to ptr), ptr %dyn_ptr)
-// CHECK: [[EXEC_USER_CODE:%[a-zA-Z0-9_.]+]] = icmp eq i32 [[INIT_CALL]], -1
-// CHECK: br i1 [[EXEC_USER_CODE]], label %[[USER_CODE_ENTRY:.+]], label %[[WORKER_EXIT:.+]]
-
-// CHECK: [[USER_CODE_ENTRY]]:
-// CHECK: [[LOAD_TEST:%[0-9]+]] = load ptr, ptr [[TEST_CAST]], align 8
-// CHECK: [[LOAD_A:%[0-9]+]] = load i32, ptr [[A_CAST]], align 4
-// CHECK: %idx.ext = sext i32 [[LOAD_A]] to i64
-// CHECK: %add.ptr = getelementptr inbounds i8, ptr [[LOAD_TEST]], i64 %idx.ext
-// CHECK: store ptr %add.ptr, ptr [[T_CAST]], align 8
-// CHECK: [[LOAD_T:%[0-9]+]] = load ptr, ptr [[T_CAST]], align 8
-// CHECK: [[LOAD_TEST_AGAIN:%[0-9]+]] = load ptr, ptr [[TEST_CAST]], align 8
-// CHECK: call i32 (ptr, ...) @printf(ptr noundef addrspacecast (ptr addrspace(4) @.str to ptr), ptr noundef [[LOAD_T]], i32 noundef 21, ptr noundef [[LOAD_TEST_AGAIN]])
-// CHECK: call void @__kmpc_target_deinit()
-// CHECK: ret void
-
-// CHECK: [[WORKER_EXIT]]:
-// CHECK: ret void
- }
-
- return 0;
-}
More information about the libc-commits
mailing list