https://github.com/gregrodgers updated https://github.com/llvm/llvm-project/pull/218473
>From cf6351e5c0470e01ed7a08cefe7d8b54d6a3798c Mon Sep 17 00:00:00 2001 From: gregrodgers <[email protected]> Date: Mon, 24 Aug 2026 07:54:29 -0500 Subject: [PATCH 1/3] [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 <[email protected]> Date: Mon, 24 Aug 2026 10:58:00 -0500 Subject: [PATCH 2/3] [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 <[email protected]> Date: Mon, 24 Aug 2026 13:53:42 -0500 Subject: [PATCH 3/3] [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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
