diff --git a/flang/include/flang/Common/real.h b/flang/include/flang/Common/real.h index 50aab7d89a597e..9ca58bed2dd7c2 100644 --- a/flang/include/flang/Common/real.h +++ b/flang/include/flang/Common/real.h @@ -13,6 +13,7 @@ // The various representations are distinguished by their binary precisions // (number of explicit significand bits and any implicit MSB in the fraction). +#include "flang/Runtime/api-attrs.h" #include namespace Fortran::common { @@ -119,6 +120,7 @@ template class RealDetails { } public: + RT_OFFLOAD_VAR_GROUP_BEGIN static constexpr int binaryPrecision{BINARY_PRECISION}; static constexpr int bits{BitsForBinaryPrecision(binaryPrecision)}; static constexpr bool isImplicitMSB{binaryPrecision != 64 /*x87*/}; @@ -138,6 +140,7 @@ template class RealDetails { static constexpr int maxHexadecimalConversionDigits{ MaxHexadecimalConversionDigits(binaryPrecision)}; + RT_OFFLOAD_VAR_GROUP_END static_assert(binaryPrecision > 0); static_assert(exponentBits > 1); diff --git a/flang/include/flang/Common/uint128.h b/flang/include/flang/Common/uint128.h index 03e44eb6997d5b..55841c0d9b9028 100644 --- a/flang/include/flang/Common/uint128.h +++ b/flang/include/flang/Common/uint128.h @@ -20,6 +20,7 @@ #endif #include "leading-zero-bit-count.h" +#include "flang/Runtime/api-attrs.h" #include #include @@ -260,7 +261,9 @@ template class Int128 { return LeadingZeroBitCount(high_); } } + RT_VAR_GROUP_BEGIN static constexpr std::uint64_t topBit{std::uint64_t{1} << 63}; + RT_VAR_GROUP_END #if FLANG_LITTLE_ENDIAN std::uint64_t low_{0}, high_{0}; #elif FLANG_BIG_ENDIAN diff --git a/flang/include/flang/Decimal/binary-floating-point.h b/flang/include/flang/Decimal/binary-floating-point.h index d1992819f85aa6..1c8829550043de 100644 --- a/flang/include/flang/Decimal/binary-floating-point.h +++ b/flang/include/flang/Decimal/binary-floating-point.h @@ -14,6 +14,7 @@ #include "flang/Common/real.h" #include "flang/Common/uint128.h" +#include "flang/Runtime/api-attrs.h" #include #include #include @@ -47,9 +48,11 @@ class BinaryFloatingPointNumber : public common::RealDetails { using RawType = common::HostUnsignedIntType; static_assert(CHAR_BIT * sizeof(RawType) >= bits); + RT_OFFLOAD_VAR_GROUP_BEGIN static constexpr RawType significandMask{(RawType{1} << significandBits) - 1}; - constexpr BinaryFloatingPointNumber() {} // zero + constexpr RT_API_ATTRS BinaryFloatingPointNumber() {} // zero + RT_OFFLOAD_VAR_GROUP_END constexpr BinaryFloatingPointNumber( const BinaryFloatingPointNumber &that) = default; constexpr BinaryFloatingPointNumber( @@ -58,26 +61,30 @@ class BinaryFloatingPointNumber : public common::RealDetails { const BinaryFloatingPointNumber &that) = default; constexpr BinaryFloatingPointNumber &operator=( BinaryFloatingPointNumber &&that) = default; - constexpr explicit BinaryFloatingPointNumber(RawType raw) : raw_{raw} {} + constexpr explicit RT_API_ATTRS BinaryFloatingPointNumber(RawType raw) + : raw_{raw} {} - RawType raw() const { return raw_; } + RT_API_ATTRS RawType raw() const { return raw_; } - template explicit constexpr BinaryFloatingPointNumber(A x) { + template + explicit constexpr RT_API_ATTRS BinaryFloatingPointNumber(A x) { static_assert(sizeof raw_ <= sizeof x); std::memcpy(reinterpret_cast(&raw_), reinterpret_cast(&x), sizeof raw_); } - constexpr int BiasedExponent() const { + constexpr RT_API_ATTRS int BiasedExponent() const { return static_cast( (raw_ >> significandBits) & ((1 << exponentBits) - 1)); } - constexpr int UnbiasedExponent() const { + constexpr RT_API_ATTRS int UnbiasedExponent() const { int biased{BiasedExponent()}; return biased - exponentBias + (biased == 0); } - constexpr RawType Significand() const { return raw_ & significandMask; } - constexpr RawType Fraction() const { + constexpr RT_API_ATTRS RawType Significand() const { + return raw_ & significandMask; + } + constexpr RT_API_ATTRS RawType Fraction() const { RawType sig{Significand()}; if (isImplicitMSB && BiasedExponent() > 0) { sig |= RawType{1} << significandBits; @@ -85,10 +92,10 @@ class BinaryFloatingPointNumber : public common::RealDetails { return sig; } - constexpr bool IsZero() const { + constexpr RT_API_ATTRS bool IsZero() const { return (raw_ & ((RawType{1} << (bits - 1)) - 1)) == 0; } - constexpr bool IsNaN() const { + constexpr RT_API_ATTRS bool IsNaN() const { auto expo{BiasedExponent()}; auto sig{Significand()}; if constexpr (bits == 80) { // x87 @@ -102,7 +109,7 @@ class BinaryFloatingPointNumber : public common::RealDetails { return expo == maxExponent && sig != 0; } } - constexpr bool IsInfinite() const { + constexpr RT_API_ATTRS bool IsInfinite() const { if constexpr (bits == 80) { // x87 return BiasedExponent() == maxExponent && Significand() == ((significandMask >> 1) + 1); @@ -110,27 +117,30 @@ class BinaryFloatingPointNumber : public common::RealDetails { return BiasedExponent() == maxExponent && Significand() == 0; } } - constexpr bool IsMaximalFiniteMagnitude() const { + constexpr RT_API_ATTRS bool IsMaximalFiniteMagnitude() const { return BiasedExponent() == maxExponent - 1 && Significand() == significandMask; } - constexpr bool IsNegative() const { return ((raw_ >> (bits - 1)) & 1) != 0; } + constexpr RT_API_ATTRS bool IsNegative() const { + return ((raw_ >> (bits - 1)) & 1) != 0; + } - constexpr void Negate() { raw_ ^= RawType{1} << (bits - 1); } + constexpr RT_API_ATTRS void Negate() { raw_ ^= RawType{1} << (bits - 1); } // For calculating the nearest neighbors of a floating-point value - constexpr void Previous() { + constexpr RT_API_ATTRS void Previous() { RemoveExplicitMSB(); --raw_; InsertExplicitMSB(); } - constexpr void Next() { + constexpr RT_API_ATTRS void Next() { RemoveExplicitMSB(); ++raw_; InsertExplicitMSB(); } - static constexpr BinaryFloatingPointNumber Infinity(bool isNegative) { + static constexpr RT_API_ATTRS BinaryFloatingPointNumber Infinity( + bool isNegative) { RawType result{RawType{maxExponent} << significandBits}; if (isNegative) { result |= RawType{1} << (bits - 1); @@ -139,7 +149,8 @@ class BinaryFloatingPointNumber : public common::RealDetails { } // Returns true when the result is exact - constexpr bool RoundToBits(int keepBits, enum FortranRounding mode) { + constexpr RT_API_ATTRS bool RoundToBits( + int keepBits, enum FortranRounding mode) { if (IsNaN() || IsInfinite() || keepBits >= binaryPrecision) { return true; } @@ -180,12 +191,12 @@ class BinaryFloatingPointNumber : public common::RealDetails { } private: - constexpr void RemoveExplicitMSB() { + constexpr RT_API_ATTRS void RemoveExplicitMSB() { if constexpr (!isImplicitMSB) { raw_ = (raw_ & (significandMask >> 1)) | ((raw_ & ~significandMask) >> 1); } } - constexpr void InsertExplicitMSB() { + constexpr RT_API_ATTRS void InsertExplicitMSB() { if constexpr (!isImplicitMSB) { constexpr RawType mask{significandMask >> 1}; raw_ = (raw_ & mask) | ((raw_ & ~mask) << 1); diff --git a/flang/include/flang/Decimal/decimal.h b/flang/include/flang/Decimal/decimal.h index f0997fb63df018..aeda01c44fa6f6 100644 --- a/flang/include/flang/Decimal/decimal.h +++ b/flang/include/flang/Decimal/decimal.h @@ -12,6 +12,7 @@ #ifndef FORTRAN_DECIMAL_DECIMAL_H_ #define FORTRAN_DECIMAL_DECIMAL_H_ +#include "flang/Runtime/api-attrs.h" #include #ifdef __cplusplus @@ -65,27 +66,27 @@ enum DecimalConversionFlags { #ifdef __cplusplus template -ConversionToDecimalResult ConvertToDecimal(char *, size_t, +RT_API_ATTRS ConversionToDecimalResult ConvertToDecimal(char *, size_t, DecimalConversionFlags, int digits, enum FortranRounding rounding, BinaryFloatingPointNumber x); -extern template ConversionToDecimalResult ConvertToDecimal<8>(char *, size_t, - enum DecimalConversionFlags, int, enum FortranRounding, +extern template RT_API_ATTRS ConversionToDecimalResult ConvertToDecimal<8>( + char *, size_t, enum DecimalConversionFlags, int, enum FortranRounding, BinaryFloatingPointNumber<8>); -extern template ConversionToDecimalResult ConvertToDecimal<11>(char *, size_t, - enum DecimalConversionFlags, int, enum FortranRounding, +extern template RT_API_ATTRS ConversionToDecimalResult ConvertToDecimal<11>( + char *, size_t, enum DecimalConversionFlags, int, enum FortranRounding, BinaryFloatingPointNumber<11>); -extern template ConversionToDecimalResult ConvertToDecimal<24>(char *, size_t, - enum DecimalConversionFlags, int, enum FortranRounding, +extern template RT_API_ATTRS ConversionToDecimalResult ConvertToDecimal<24>( + char *, size_t, enum DecimalConversionFlags, int, enum FortranRounding, BinaryFloatingPointNumber<24>); -extern template ConversionToDecimalResult ConvertToDecimal<53>(char *, size_t, - enum DecimalConversionFlags, int, enum FortranRounding, +extern template RT_API_ATTRS ConversionToDecimalResult ConvertToDecimal<53>( + char *, size_t, enum DecimalConversionFlags, int, enum FortranRounding, BinaryFloatingPointNumber<53>); -extern template ConversionToDecimalResult ConvertToDecimal<64>(char *, size_t, - enum DecimalConversionFlags, int, enum FortranRounding, +extern template RT_API_ATTRS ConversionToDecimalResult ConvertToDecimal<64>( + char *, size_t, enum DecimalConversionFlags, int, enum FortranRounding, BinaryFloatingPointNumber<64>); -extern template ConversionToDecimalResult ConvertToDecimal<113>(char *, size_t, - enum DecimalConversionFlags, int, enum FortranRounding, +extern template RT_API_ATTRS ConversionToDecimalResult ConvertToDecimal<113>( + char *, size_t, enum DecimalConversionFlags, int, enum FortranRounding, BinaryFloatingPointNumber<113>); template struct ConversionToBinaryResult { @@ -94,20 +95,20 @@ template struct ConversionToBinaryResult { }; template -ConversionToBinaryResult ConvertToBinary(const char *&, +RT_API_ATTRS ConversionToBinaryResult ConvertToBinary(const char *&, enum FortranRounding = RoundNearest, const char *end = nullptr); -extern template ConversionToBinaryResult<8> ConvertToBinary<8>( +extern template RT_API_ATTRS ConversionToBinaryResult<8> ConvertToBinary<8>( const char *&, enum FortranRounding, const char *end); -extern template ConversionToBinaryResult<11> ConvertToBinary<11>( +extern template RT_API_ATTRS ConversionToBinaryResult<11> ConvertToBinary<11>( const char *&, enum FortranRounding, const char *end); -extern template ConversionToBinaryResult<24> ConvertToBinary<24>( +extern template RT_API_ATTRS ConversionToBinaryResult<24> ConvertToBinary<24>( const char *&, enum FortranRounding, const char *end); -extern template ConversionToBinaryResult<53> ConvertToBinary<53>( +extern template RT_API_ATTRS ConversionToBinaryResult<53> ConvertToBinary<53>( const char *&, enum FortranRounding, const char *end); -extern template ConversionToBinaryResult<64> ConvertToBinary<64>( +extern template RT_API_ATTRS ConversionToBinaryResult<64> ConvertToBinary<64>( const char *&, enum FortranRounding, const char *end); -extern template ConversionToBinaryResult<113> ConvertToBinary<113>( +extern template RT_API_ATTRS ConversionToBinaryResult<113> ConvertToBinary<113>( const char *&, enum FortranRounding, const char *end); } // namespace Fortran::decimal extern "C" { @@ -116,21 +117,21 @@ extern "C" { #define NS(x) x #endif /* C++ */ -struct NS(ConversionToDecimalResult) +RT_API_ATTRS struct NS(ConversionToDecimalResult) ConvertFloatToDecimal(char *, size_t, enum NS(DecimalConversionFlags), int digits, enum NS(FortranRounding), float); -struct NS(ConversionToDecimalResult) +RT_API_ATTRS struct NS(ConversionToDecimalResult) ConvertDoubleToDecimal(char *, size_t, enum NS(DecimalConversionFlags), int digits, enum NS(FortranRounding), double); -struct NS(ConversionToDecimalResult) +RT_API_ATTRS struct NS(ConversionToDecimalResult) ConvertLongDoubleToDecimal(char *, size_t, enum NS(DecimalConversionFlags), int digits, enum NS(FortranRounding), long double); -enum NS(ConversionResultFlags) +RT_API_ATTRS enum NS(ConversionResultFlags) ConvertDecimalToFloat(const char **, float *, enum NS(FortranRounding)); -enum NS(ConversionResultFlags) +RT_API_ATTRS enum NS(ConversionResultFlags) ConvertDecimalToDouble(const char **, double *, enum NS(FortranRounding)); -enum NS(ConversionResultFlags) ConvertDecimalToLongDouble( +RT_API_ATTRS enum NS(ConversionResultFlags) ConvertDecimalToLongDouble( const char **, long double *, enum NS(FortranRounding)); #undef NS #ifdef __cplusplus diff --git a/flang/include/flang/Runtime/api-attrs.h b/flang/include/flang/Runtime/api-attrs.h index fc3eb42e1b73f5..050d2366b8e165 100644 --- a/flang/include/flang/Runtime/api-attrs.h +++ b/flang/include/flang/Runtime/api-attrs.h @@ -102,7 +102,7 @@ * to appear as part of a C++ decl-specifier. */ #ifndef RT_CONST_VAR_ATTRS -#if defined(__CUDACC__) || defined(__CUDA__) +#if (defined(__CUDACC__) || defined(__CUDA__)) && defined(__CUDA_ARCH__) #define RT_CONST_VAR_ATTRS __constant__ #else #define RT_CONST_VAR_ATTRS diff --git a/flang/include/flang/Runtime/io-api.h b/flang/include/flang/Runtime/io-api.h index 556cc20c5a121e..fae4e83e1b35db 100644 --- a/flang/include/flang/Runtime/io-api.h +++ b/flang/include/flang/Runtime/io-api.h @@ -58,6 +58,14 @@ extern "C" { #define IONAME(name) RTNAME(io##name) +#ifndef IODECL +#define IODECL(name) RT_API_ATTRS IONAME(name) +#endif + +#ifndef IODEF +#define IODEF(name) RT_API_ATTRS IONAME(name) +#endif + // These functions initiate data transfer statements (READ, WRITE, PRINT). // Example: PRINT *, 666 is implemented as the series of calls: // Cookie cookie{BeginExternalListOutput(DefaultOutputUnit, @@ -139,7 +147,7 @@ enum Iostat IONAME(CheckUnitNumberInRange128)(common::int128_t unit, const char *sourceFile = nullptr, int sourceLine = 0); // External synchronous I/O initiation -Cookie IONAME(BeginExternalListOutput)(ExternalUnit = DefaultOutputUnit, +Cookie IODECL(BeginExternalListOutput)(ExternalUnit = DefaultOutputUnit, const char *sourceFile = nullptr, int sourceLine = 0); Cookie IONAME(BeginExternalListInput)(ExternalUnit = DefaultInputUnit, const char *sourceFile = nullptr, int sourceLine = 0); @@ -253,7 +261,7 @@ bool IONAME(InputDescriptor)(Cookie, const Descriptor &); // Formatted (including list directed) I/O data items bool IONAME(OutputInteger8)(Cookie, std::int8_t); bool IONAME(OutputInteger16)(Cookie, std::int16_t); -bool IONAME(OutputInteger32)(Cookie, std::int32_t); +bool IODECL(OutputInteger32)(Cookie, std::int32_t); bool IONAME(OutputInteger64)(Cookie, std::int64_t); bool IONAME(OutputInteger128)(Cookie, common::int128_t); bool IONAME(InputInteger)(Cookie, std::int64_t &, int kind = 8); @@ -357,7 +365,7 @@ bool IONAME(InquireInteger64)( // returned is guaranteed to only be one of the problems that the // EnableHandlers() call has indicated should be handled in compiled code // rather than by terminating the image. -enum Iostat IONAME(EndIoStatement)(Cookie); +enum Iostat IODECL(EndIoStatement)(Cookie); } // extern "C" } // namespace Fortran::runtime::io diff --git a/flang/include/flang/Runtime/iostat.h b/flang/include/flang/Runtime/iostat.h index afce509cf1f564..0c947f6cb552c1 100644 --- a/flang/include/flang/Runtime/iostat.h +++ b/flang/include/flang/Runtime/iostat.h @@ -11,6 +11,7 @@ #ifndef FORTRAN_RUNTIME_IOSTAT_H_ #define FORTRAN_RUNTIME_IOSTAT_H_ +#include "flang/Runtime/api-attrs.h" #include "flang/Runtime/magic-numbers.h" namespace Fortran::runtime::io { @@ -88,7 +89,7 @@ enum Iostat { IostatNonExternalDefinedUnformattedIo, }; -const char *IostatErrorString(int); +const RT_API_ATTRS char *IostatErrorString(int); } // namespace Fortran::runtime::io #endif // FORTRAN_RUNTIME_IOSTAT_H_ diff --git a/flang/include/flang/Runtime/memory.h b/flang/include/flang/Runtime/memory.h index e24c509f4e90cb..5a1d352963400b 100644 --- a/flang/include/flang/Runtime/memory.h +++ b/flang/include/flang/Runtime/memory.h @@ -128,12 +128,21 @@ inline RT_API_ATTRS bool operator!=(std::nullptr_t, const OwningPtr &x) { template class SizedNew { public: - explicit SizedNew(const Terminator &terminator) : terminator_{terminator} {} + explicit RT_API_ATTRS SizedNew(const Terminator &terminator) + : terminator_{terminator} {} + + // Disable warnings about calling constructors for host-only + // classes (those thare are not supposed to be constructed) + // on the device. + RT_DIAG_PUSH + RT_DIAG_DISABLE_CALL_HOST_FROM_DEVICE_WARN template - [[nodiscard]] OwningPtr operator()(std::size_t bytes, X &&...x) { + [[nodiscard]] RT_API_ATTRS OwningPtr operator()( + std::size_t bytes, X &&...x) { return OwningPtr{new (AllocateMemoryOrCrash(terminator_, bytes)) A{std::forward(x)...}}; } + RT_DIAG_POP private: const Terminator &terminator_; @@ -141,7 +150,8 @@ template class SizedNew { template struct New : public SizedNew { using SizedNew::SizedNew; - template [[nodiscard]] OwningPtr operator()(X &&...x) { + template + [[nodiscard]] RT_API_ATTRS OwningPtr operator()(X &&...x) { return SizedNew::operator()(sizeof(A), std::forward(x)...); } }; diff --git a/flang/include/flang/Runtime/type-code.h b/flang/include/flang/Runtime/type-code.h index f7419249c2ba9c..8e7314e0af1efc 100644 --- a/flang/include/flang/Runtime/type-code.h +++ b/flang/include/flang/Runtime/type-code.h @@ -12,7 +12,6 @@ #include "flang/Common/Fortran.h" #include "flang/Common/optional.h" #include "flang/ISO_Fortran_binding_wrapper.h" -#include #include namespace Fortran::runtime { diff --git a/flang/runtime/CMakeLists.txt b/flang/runtime/CMakeLists.txt index 7dd60b5edcd5fb..984e7b97537752 100644 --- a/flang/runtime/CMakeLists.txt +++ b/flang/runtime/CMakeLists.txt @@ -179,18 +179,29 @@ set(supported_files array-constructor.cpp assign.cpp character.cpp + connection.cpp copy.cpp derived-api.cpp derived.cpp descriptor.cpp dot-product.cpp + edit-input.cpp + edit-output.cpp + environment.cpp extrema.cpp findloc.cpp + format.cpp inquiry.cpp + internal-unit.cpp + io-api.cpp + io-error.cpp + io-stmt.cpp + iostat.cpp matmul-transpose.cpp matmul.cpp memory.cpp misc-intrinsic.cpp + namelist.cpp numeric.cpp pointer.cpp product.cpp @@ -203,6 +214,7 @@ set(supported_files transformational.cpp type-code.cpp type-info.cpp + utf.cpp ) if (FLANG_EXPERIMENTAL_CUDA_RUNTIME) diff --git a/flang/runtime/connection.cpp b/flang/runtime/connection.cpp index 91ac9a0e14e47b..94a16222fc1eda 100644 --- a/flang/runtime/connection.cpp +++ b/flang/runtime/connection.cpp @@ -13,31 +13,33 @@ namespace Fortran::runtime::io { -std::size_t ConnectionState::RemainingSpaceInRecord() const { +RT_OFFLOAD_API_GROUP_BEGIN +RT_API_ATTRS std::size_t ConnectionState::RemainingSpaceInRecord() const { auto recl{recordLength.value_or(openRecl.value_or( executionEnvironment.listDirectedOutputLineLengthLimit))}; return positionInRecord >= recl ? 0 : recl - positionInRecord; } -bool ConnectionState::NeedAdvance(std::size_t width) const { +RT_API_ATTRS bool ConnectionState::NeedAdvance(std::size_t width) const { return positionInRecord > 0 && width > RemainingSpaceInRecord(); } -bool ConnectionState::IsAtEOF() const { +RT_API_ATTRS bool ConnectionState::IsAtEOF() const { return endfileRecordNumber && currentRecordNumber >= *endfileRecordNumber; } -bool ConnectionState::IsAfterEndfile() const { +RT_API_ATTRS bool ConnectionState::IsAfterEndfile() const { return endfileRecordNumber && currentRecordNumber > *endfileRecordNumber; } -void ConnectionState::HandleAbsolutePosition(std::int64_t n) { +RT_API_ATTRS void ConnectionState::HandleAbsolutePosition(std::int64_t n) { positionInRecord = std::max(n, std::int64_t{0}) + leftTabLimit.value_or(0); } -void ConnectionState::HandleRelativePosition(std::int64_t n) { +RT_API_ATTRS void ConnectionState::HandleRelativePosition(std::int64_t n) { positionInRecord = std::max(leftTabLimit.value_or(0), positionInRecord + n); } +RT_OFFLOAD_API_GROUP_END SavedPosition::SavedPosition(IoStatementState &io) : io_{io} { ConnectionState &conn{io_.GetConnectionState()}; diff --git a/flang/runtime/connection.h b/flang/runtime/connection.h index c41970d47e7b09..6f1ea90a160e5e 100644 --- a/flang/runtime/connection.h +++ b/flang/runtime/connection.h @@ -31,12 +31,12 @@ struct ConnectionAttributes { unsigned char internalIoCharKind{0}; // 0->external, 1/2/4->internal Fortran::common::optional openRecl; // RECL= on OPEN - bool IsRecordFile() const { + RT_API_ATTRS bool IsRecordFile() const { // Formatted stream files are viewed as having records, at least on input return access != Access::Stream || !isUnformatted.value_or(true); } - template constexpr bool useUTF8() const { + template constexpr RT_API_ATTRS bool useUTF8() const { // For wide CHARACTER kinds, always use UTF-8 for formatted I/O. // For single-byte CHARACTER, encode characters >= 0x80 with // UTF-8 iff the mode is set. @@ -45,25 +45,28 @@ struct ConnectionAttributes { }; struct ConnectionState : public ConnectionAttributes { - bool IsAtEOF() const; // true when read has hit EOF or endfile record - bool IsAfterEndfile() const; // true after ENDFILE until repositioned + RT_API_ATTRS bool + IsAtEOF() const; // true when read has hit EOF or endfile record + RT_API_ATTRS bool + IsAfterEndfile() const; // true after ENDFILE until repositioned // All positions and measurements are always in units of bytes, // not characters. Multi-byte character encodings are possible in // both internal I/O (when the character kind of the variable is 2 or 4) // and external formatted I/O (when the encoding is UTF-8). - std::size_t RemainingSpaceInRecord() const; - bool NeedAdvance(std::size_t) const; - void HandleAbsolutePosition(std::int64_t); - void HandleRelativePosition(std::int64_t); + RT_API_ATTRS std::size_t RemainingSpaceInRecord() const; + RT_API_ATTRS bool NeedAdvance(std::size_t) const; + RT_API_ATTRS void HandleAbsolutePosition(std::int64_t); + RT_API_ATTRS void HandleRelativePosition(std::int64_t); - void BeginRecord() { + RT_API_ATTRS void BeginRecord() { positionInRecord = 0; furthestPositionInRecord = 0; unterminatedRecord = false; } - Fortran::common::optional EffectiveRecordLength() const { + RT_API_ATTRS Fortran::common::optional + EffectiveRecordLength() const { // When an input record is longer than an explicit RECL= from OPEN // it is effectively truncated on input. return openRecl && recordLength && *openRecl < *recordLength ? openRecl @@ -110,9 +113,9 @@ struct ConnectionState : public ConnectionAttributes { // Utility class for capturing and restoring a position in an input stream. class SavedPosition { public: - explicit SavedPosition(IoStatementState &); - ~SavedPosition(); - void Cancel() { cancelled_ = true; } + explicit RT_API_ATTRS SavedPosition(IoStatementState &); + RT_API_ATTRS ~SavedPosition(); + RT_API_ATTRS void Cancel() { cancelled_ = true; } private: IoStatementState &io_; diff --git a/flang/runtime/descriptor-io.h b/flang/runtime/descriptor-io.h index b6b0fefcff870b..63995961df5797 100644 --- a/flang/runtime/descriptor-io.h +++ b/flang/runtime/descriptor-io.h @@ -28,8 +28,8 @@ namespace Fortran::runtime::io::descr { template -inline A &ExtractElement(IoStatementState &io, const Descriptor &descriptor, - const SubscriptValue subscripts[]) { +inline RT_API_ATTRS A &ExtractElement(IoStatementState &io, + const Descriptor &descriptor, const SubscriptValue subscripts[]) { A *p{descriptor.Element(subscripts)}; if (!p) { io.GetIoErrorHandler().Crash("Bad address for I/O item -- null base " @@ -45,7 +45,7 @@ inline A &ExtractElement(IoStatementState &io, const Descriptor &descriptor, // NAMELIST array output. template -inline bool FormattedIntegerIO( +inline RT_API_ATTRS bool FormattedIntegerIO( IoStatementState &io, const Descriptor &descriptor) { std::size_t numElements{descriptor.Elements()}; SubscriptValue subscripts[maxRank]; @@ -78,7 +78,7 @@ inline bool FormattedIntegerIO( } template -inline bool FormattedRealIO( +inline RT_API_ATTRS bool FormattedRealIO( IoStatementState &io, const Descriptor &descriptor) { std::size_t numElements{descriptor.Elements()}; SubscriptValue subscripts[maxRank]; @@ -111,7 +111,7 @@ inline bool FormattedRealIO( } template -inline bool FormattedComplexIO( +inline RT_API_ATTRS bool FormattedComplexIO( IoStatementState &io, const Descriptor &descriptor) { std::size_t numElements{descriptor.Elements()}; SubscriptValue subscripts[maxRank]; @@ -159,7 +159,7 @@ inline bool FormattedComplexIO( } template -inline bool FormattedCharacterIO( +inline RT_API_ATTRS bool FormattedCharacterIO( IoStatementState &io, const Descriptor &descriptor) { std::size_t numElements{descriptor.Elements()}; SubscriptValue subscripts[maxRank]; @@ -199,7 +199,7 @@ inline bool FormattedCharacterIO( } template -inline bool FormattedLogicalIO( +inline RT_API_ATTRS bool FormattedLogicalIO( IoStatementState &io, const Descriptor &descriptor) { std::size_t numElements{descriptor.Elements()}; SubscriptValue subscripts[maxRank]; @@ -241,12 +241,12 @@ inline bool FormattedLogicalIO( } template -static bool DescriptorIO(IoStatementState &, const Descriptor &, +static RT_API_ATTRS bool DescriptorIO(IoStatementState &, const Descriptor &, const NonTbpDefinedIoTable * = nullptr); // For intrinsic (not defined) derived type I/O, formatted & unformatted template -static bool DefaultComponentIO(IoStatementState &io, +static RT_API_ATTRS bool DefaultComponentIO(IoStatementState &io, const typeInfo::Component &component, const Descriptor &origDescriptor, const SubscriptValue origSubscripts[], Terminator &terminator, const NonTbpDefinedIoTable *table) { @@ -269,7 +269,7 @@ static bool DefaultComponentIO(IoStatementState &io, } template -static bool DefaultComponentwiseFormattedIO(IoStatementState &io, +static RT_API_ATTRS bool DefaultComponentwiseFormattedIO(IoStatementState &io, const Descriptor &descriptor, const typeInfo::DerivedType &type, const NonTbpDefinedIoTable *table, const SubscriptValue subscripts[]) { IoErrorHandler &handler{io.GetIoErrorHandler()}; @@ -368,7 +368,12 @@ static bool FormattedDerivedTypeIO(IoStatementState &io, ++j, descriptor.IncrementSubscripts(subscripts)) { Fortran::common::optional result; if (special) { +#if !defined(RT_DEVICE_COMPILATION) result = DefinedFormattedIo(io, descriptor, *type, *special, subscripts); +#else + io.GetIoErrorHandler().Crash("not implemented yet: defined formatted IO"); + return false; +#endif } if (!result) { result = DefaultComponentwiseFormattedIO( @@ -488,8 +493,8 @@ static bool UnformattedDescriptorIO(IoStatementState &io, } template -static bool DescriptorIO(IoStatementState &io, const Descriptor &descriptor, - const NonTbpDefinedIoTable *table) { +static RT_API_ATTRS bool DescriptorIO(IoStatementState &io, + const Descriptor &descriptor, const NonTbpDefinedIoTable *table) { IoErrorHandler &handler{io.GetIoErrorHandler()}; if (handler.InError()) { return false; @@ -505,7 +510,12 @@ static bool DescriptorIO(IoStatementState &io, const Descriptor &descriptor, } } if (!io.get_if>()) { +#if !defined(RT_DEVICE_COMPILATION) return UnformattedDescriptorIO(io, descriptor, table); +#else + io.GetIoErrorHandler().Crash("not implemented yet: unformatted IO"); + return false; +#endif } if (auto catAndKind{descriptor.type().GetCategoryAndKind()}) { TypeCategory cat{catAndKind->first}; @@ -597,7 +607,12 @@ static bool DescriptorIO(IoStatementState &io, const Descriptor &descriptor, return false; } case TypeCategory::Derived: +#if !defined(RT_DEVICE_COMPILATION) return FormattedDerivedTypeIO(io, descriptor, table); +#else + io.GetIoErrorHandler().Crash("not implemented yet: derived type IO"); + return false; +#endif } } handler.Crash("DescriptorIO: bad type code (%d) in descriptor", diff --git a/flang/runtime/edit-input.cpp b/flang/runtime/edit-input.cpp index f7cbbc21e5956e..1c862f3792b9c5 100644 --- a/flang/runtime/edit-input.cpp +++ b/flang/runtime/edit-input.cpp @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include "edit-input.h" +#include "freestanding-tools.h" #include "namelist.h" #include "utf.h" #include "flang/Common/optional.h" @@ -17,16 +18,19 @@ namespace Fortran::runtime::io { +RT_OFFLOAD_API_GROUP_BEGIN + // Checks that a list-directed input value has been entirely consumed and // doesn't contain unparsed characters before the next value separator. -static inline bool IsCharValueSeparator(const DataEdit &edit, char32_t ch) { +static inline RT_API_ATTRS bool IsCharValueSeparator( + const DataEdit &edit, char32_t ch) { char32_t comma{ edit.modes.editingFlags & decimalComma ? char32_t{';'} : char32_t{','}}; return ch == ' ' || ch == '\t' || ch == comma || ch == '/' || (edit.IsNamelist() && (ch == '&' || ch == '$')); } -static bool CheckCompleteListDirectedField( +static RT_API_ATTRS bool CheckCompleteListDirectedField( IoStatementState &io, const DataEdit &edit) { if (edit.IsListDirected()) { std::size_t byteCount; @@ -52,7 +56,7 @@ static bool CheckCompleteListDirectedField( } template -static bool EditBOZInput( +static RT_API_ATTRS bool EditBOZInput( IoStatementState &io, const DataEdit &edit, void *n, std::size_t bytes) { // Skip leading white space & zeroes Fortran::common::optional remaining{io.CueUpInput(edit)}; @@ -151,13 +155,13 @@ static bool EditBOZInput( return CheckCompleteListDirectedField(io, edit); } -static inline char32_t GetRadixPointChar(const DataEdit &edit) { +static inline RT_API_ATTRS char32_t GetRadixPointChar(const DataEdit &edit) { return edit.modes.editingFlags & decimalComma ? char32_t{','} : char32_t{'.'}; } // Prepares input from a field, and returns the sign, if any, else '\0'. -static char ScanNumericPrefix(IoStatementState &io, const DataEdit &edit, - Fortran::common::optional &next, +static RT_API_ATTRS char ScanNumericPrefix(IoStatementState &io, + const DataEdit &edit, Fortran::common::optional &next, Fortran::common::optional &remaining) { remaining = io.CueUpInput(edit); next = io.NextInField(remaining, edit); @@ -174,7 +178,7 @@ static char ScanNumericPrefix(IoStatementState &io, const DataEdit &edit, return sign; } -bool EditIntegerInput( +RT_API_ATTRS bool EditIntegerInput( IoStatementState &io, const DataEdit &edit, void *n, int kind) { RUNTIME_CHECK(io.GetIoErrorHandler(), kind >= 1 && !(kind & (kind - 1))); switch (edit.descriptor) { @@ -279,18 +283,20 @@ struct ScannedRealInput { int exponent{0}; // adjusted as necessary; binary if isHexadecimal bool isHexadecimal{false}; // 0X... }; -static ScannedRealInput ScanRealInput( +static RT_API_ATTRS ScannedRealInput ScanRealInput( char *buffer, int bufferSize, IoStatementState &io, const DataEdit &edit) { Fortran::common::optional remaining; Fortran::common::optional next; int got{0}; Fortran::common::optional radixPointOffset; - auto Put{[&](char ch) -> void { + // The following lambda definition violates the conding style, + // but cuda-11.8 nvcc hits an internal error with the brace initialization. + auto Put = [&](char ch) -> void { if (got < bufferSize) { buffer[got] = ch; } ++got; - }}; + }; char sign{ScanNumericPrefix(io, edit, next, remaining)}; if (sign == '-') { Put('-'); @@ -487,13 +493,24 @@ static ScannedRealInput ScanRealInput( return {got, exponent, isHexadecimal}; } -static void RaiseFPExceptions(decimal::ConversionResultFlags flags) { +static RT_API_ATTRS void RaiseFPExceptions( + decimal::ConversionResultFlags flags) { +#if defined(RT_DEVICE_COMPILATION) + Terminator terminator(__FILE__, __LINE__); +#endif + #undef RAISE +#if defined(RT_DEVICE_COMPILATION) +#define RAISE(e) \ + terminator.Crash( \ + "not implemented yet: raising FP exception in device code: %s", #e); +#else // !defined(RT_DEVICE_COMPILATION) #ifdef feraisexcept // a macro in some environments; omit std:: #define RAISE feraiseexcept #else #define RAISE std::feraiseexcept #endif +#endif // !defined(RT_DEVICE_COMPILATION) if (flags & decimal::ConversionResultFlags::Overflow) { RAISE(FE_OVERFLOW); } @@ -514,7 +531,7 @@ static void RaiseFPExceptions(decimal::ConversionResultFlags flags) { // converter without modification, this fast path for real input // saves time by avoiding memory copies and reformatting of the exponent. template -static bool TryFastPathRealDecimalInput( +static RT_API_ATTRS bool TryFastPathRealDecimalInput( IoStatementState &io, const DataEdit &edit, void *n) { if (edit.modes.editingFlags & (blankZero | decimalComma)) { return false; @@ -586,7 +603,8 @@ static bool TryFastPathRealDecimalInput( } template -decimal::ConversionToBinaryResult ConvertHexadecimal( +RT_API_ATTRS decimal::ConversionToBinaryResult +ConvertHexadecimal( const char *&p, enum decimal::FortranRounding rounding, int expo) { using RealType = decimal::BinaryFloatingPointNumber; using RawType = typename RealType::RawType; @@ -689,7 +707,8 @@ decimal::ConversionToBinaryResult ConvertHexadecimal( } template -bool EditCommonRealInput(IoStatementState &io, const DataEdit &edit, void *n) { +RT_API_ATTRS bool EditCommonRealInput( + IoStatementState &io, const DataEdit &edit, void *n) { constexpr int binaryPrecision{common::PrecisionOfRealKind(KIND)}; if (TryFastPathRealDecimalInput(io, edit, n)) { return CheckCompleteListDirectedField(io, edit); @@ -785,7 +804,8 @@ bool EditCommonRealInput(IoStatementState &io, const DataEdit &edit, void *n) { } template -bool EditRealInput(IoStatementState &io, const DataEdit &edit, void *n) { +RT_API_ATTRS bool EditRealInput( + IoStatementState &io, const DataEdit &edit, void *n) { switch (edit.descriptor) { case DataEdit::ListDirected: if (IsNamelistNameOrSlash(io)) { @@ -819,7 +839,8 @@ bool EditRealInput(IoStatementState &io, const DataEdit &edit, void *n) { } // 13.7.3 in Fortran 2018 -bool EditLogicalInput(IoStatementState &io, const DataEdit &edit, bool &x) { +RT_API_ATTRS bool EditLogicalInput( + IoStatementState &io, const DataEdit &edit, bool &x) { switch (edit.descriptor) { case DataEdit::ListDirected: if (IsNamelistNameOrSlash(io)) { @@ -869,7 +890,7 @@ bool EditLogicalInput(IoStatementState &io, const DataEdit &edit, bool &x) { // See 13.10.3.1 paragraphs 7-9 in Fortran 2018 template -static bool EditDelimitedCharacterInput( +static RT_API_ATTRS bool EditDelimitedCharacterInput( IoStatementState &io, CHAR *x, std::size_t length, char32_t delimiter) { bool result{true}; while (true) { @@ -898,12 +919,12 @@ static bool EditDelimitedCharacterInput( --length; } } - std::fill_n(x, length, ' '); + Fortran::runtime::fill_n(x, length, ' '); return result; } template -static bool EditListDirectedCharacterInput( +static RT_API_ATTRS bool EditListDirectedCharacterInput( IoStatementState &io, CHAR *x, std::size_t length, const DataEdit &edit) { std::size_t byteCount{0}; auto ch{io.GetCurrentChar(byteCount)}; @@ -948,13 +969,13 @@ static bool EditListDirectedCharacterInput( remaining = --length > 0 ? maxUTF8Bytes : 0; } } - std::fill_n(x, length, ' '); + Fortran::runtime::fill_n(x, length, ' '); return true; } template -bool EditCharacterInput(IoStatementState &io, const DataEdit &edit, CHAR *x, - std::size_t lengthChars) { +RT_API_ATTRS bool EditCharacterInput(IoStatementState &io, const DataEdit &edit, + CHAR *x, std::size_t lengthChars) { switch (edit.descriptor) { case DataEdit::ListDirected: return EditListDirectedCharacterInput(io, x, lengthChars, edit); @@ -998,7 +1019,7 @@ bool EditCharacterInput(IoStatementState &io, const DataEdit &edit, CHAR *x, if (io.CheckForEndOfRecord(readyBytes)) { if (readyBytes == 0) { // PAD='YES' and no more data - std::fill_n(x, lengthChars, ' '); + Fortran::runtime::fill_n(x, lengthChars, ' '); return !io.GetIoErrorHandler().InError(); } else { // Do partial read(s) then pad on last iteration @@ -1075,23 +1096,31 @@ bool EditCharacterInput(IoStatementState &io, const DataEdit &edit, CHAR *x, readyBytes -= chunkBytes; } // Pad the remainder of the input variable, if any. - std::fill_n(x, lengthChars, ' '); + Fortran::runtime::fill_n(x, lengthChars, ' '); return CheckCompleteListDirectedField(io, edit); } -template bool EditRealInput<2>(IoStatementState &, const DataEdit &, void *); -template bool EditRealInput<3>(IoStatementState &, const DataEdit &, void *); -template bool EditRealInput<4>(IoStatementState &, const DataEdit &, void *); -template bool EditRealInput<8>(IoStatementState &, const DataEdit &, void *); -template bool EditRealInput<10>(IoStatementState &, const DataEdit &, void *); +template RT_API_ATTRS bool EditRealInput<2>( + IoStatementState &, const DataEdit &, void *); +template RT_API_ATTRS bool EditRealInput<3>( + IoStatementState &, const DataEdit &, void *); +template RT_API_ATTRS bool EditRealInput<4>( + IoStatementState &, const DataEdit &, void *); +template RT_API_ATTRS bool EditRealInput<8>( + IoStatementState &, const DataEdit &, void *); +template RT_API_ATTRS bool EditRealInput<10>( + IoStatementState &, const DataEdit &, void *); // TODO: double/double -template bool EditRealInput<16>(IoStatementState &, const DataEdit &, void *); +template RT_API_ATTRS bool EditRealInput<16>( + IoStatementState &, const DataEdit &, void *); -template bool EditCharacterInput( +template RT_API_ATTRS bool EditCharacterInput( IoStatementState &, const DataEdit &, char *, std::size_t); -template bool EditCharacterInput( +template RT_API_ATTRS bool EditCharacterInput( IoStatementState &, const DataEdit &, char16_t *, std::size_t); -template bool EditCharacterInput( +template RT_API_ATTRS bool EditCharacterInput( IoStatementState &, const DataEdit &, char32_t *, std::size_t); +RT_OFFLOAD_API_GROUP_END + } // namespace Fortran::runtime::io diff --git a/flang/runtime/edit-input.h b/flang/runtime/edit-input.h index 61844a1199a748..a90180b8ee2ebd 100644 --- a/flang/runtime/edit-input.h +++ b/flang/runtime/edit-input.h @@ -15,36 +15,38 @@ namespace Fortran::runtime::io { -bool EditIntegerInput(IoStatementState &, const DataEdit &, void *, int kind); +RT_API_ATTRS bool EditIntegerInput( + IoStatementState &, const DataEdit &, void *, int kind); template -bool EditRealInput(IoStatementState &, const DataEdit &, void *); +RT_API_ATTRS bool EditRealInput(IoStatementState &, const DataEdit &, void *); -bool EditLogicalInput(IoStatementState &, const DataEdit &, bool &); +RT_API_ATTRS bool EditLogicalInput( + IoStatementState &, const DataEdit &, bool &); template -bool EditCharacterInput( +RT_API_ATTRS bool EditCharacterInput( IoStatementState &, const DataEdit &, CHAR *, std::size_t); -extern template bool EditRealInput<2>( +extern template RT_API_ATTRS bool EditRealInput<2>( IoStatementState &, const DataEdit &, void *); -extern template bool EditRealInput<3>( +extern template RT_API_ATTRS bool EditRealInput<3>( IoStatementState &, const DataEdit &, void *); -extern template bool EditRealInput<4>( +extern template RT_API_ATTRS bool EditRealInput<4>( IoStatementState &, const DataEdit &, void *); -extern template bool EditRealInput<8>( +extern template RT_API_ATTRS bool EditRealInput<8>( IoStatementState &, const DataEdit &, void *); -extern template bool EditRealInput<10>( +extern template RT_API_ATTRS bool EditRealInput<10>( IoStatementState &, const DataEdit &, void *); // TODO: double/double -extern template bool EditRealInput<16>( +extern template RT_API_ATTRS bool EditRealInput<16>( IoStatementState &, const DataEdit &, void *); -extern template bool EditCharacterInput( +extern template RT_API_ATTRS bool EditCharacterInput( IoStatementState &, const DataEdit &, char *, std::size_t); -extern template bool EditCharacterInput( +extern template RT_API_ATTRS bool EditCharacterInput( IoStatementState &, const DataEdit &, char16_t *, std::size_t); -extern template bool EditCharacterInput( +extern template RT_API_ATTRS bool EditCharacterInput( IoStatementState &, const DataEdit &, char32_t *, std::size_t); } // namespace Fortran::runtime::io diff --git a/flang/runtime/edit-output.cpp b/flang/runtime/edit-output.cpp index 7267540370fc07..abb62bb48a2a90 100644 --- a/flang/runtime/edit-output.cpp +++ b/flang/runtime/edit-output.cpp @@ -15,8 +15,10 @@ namespace Fortran::runtime::io { +RT_OFFLOAD_API_GROUP_BEGIN + // In output statement, add a space between numbers and characters. -static void addSpaceBeforeCharacter(IoStatementState &io) { +static RT_API_ATTRS void addSpaceBeforeCharacter(IoStatementState &io) { if (auto *list{io.get_if>()}) { list->set_lastWasUndelimitedCharacter(false); } @@ -26,8 +28,8 @@ static void addSpaceBeforeCharacter(IoStatementState &io) { // representation of what is interpreted to be a single unsigned integer value. // When used with character data, endianness is exposed. template -static bool EditBOZOutput(IoStatementState &io, const DataEdit &edit, - const unsigned char *data0, std::size_t bytes) { +static RT_API_ATTRS bool EditBOZOutput(IoStatementState &io, + const DataEdit &edit, const unsigned char *data0, std::size_t bytes) { addSpaceBeforeCharacter(io); int digits{static_cast((bytes * 8) / LOG2_BASE)}; int get{static_cast(bytes * 8) - digits * LOG2_BASE}; @@ -107,7 +109,7 @@ static bool EditBOZOutput(IoStatementState &io, const DataEdit &edit, } template -bool EditIntegerOutput(IoStatementState &io, const DataEdit &edit, +bool RT_API_ATTRS EditIntegerOutput(IoStatementState &io, const DataEdit &edit, common::HostSignedIntType<8 * KIND> n) { addSpaceBeforeCharacter(io); char buffer[130], *end{&buffer[sizeof buffer]}, *p{end}; @@ -187,7 +189,7 @@ bool EditIntegerOutput(IoStatementState &io, const DataEdit &edit, } // Formats the exponent (see table 13.1 for all the cases) -const char *RealOutputEditingBase::FormatExponent( +RT_API_ATTRS const char *RealOutputEditingBase::FormatExponent( int expo, const DataEdit &edit, int &length) { char *eEnd{&exponent_[sizeof exponent_]}; char *exponent{eEnd}; @@ -226,7 +228,7 @@ const char *RealOutputEditingBase::FormatExponent( return overflow ? nullptr : exponent; } -bool RealOutputEditingBase::EmitPrefix( +RT_API_ATTRS bool RealOutputEditingBase::EmitPrefix( const DataEdit &edit, std::size_t length, std::size_t width) { if (edit.IsListDirected()) { int prefixLength{edit.descriptor == DataEdit::ListDirectedRealPart ? 2 @@ -247,7 +249,7 @@ bool RealOutputEditingBase::EmitPrefix( } } -bool RealOutputEditingBase::EmitSuffix(const DataEdit &edit) { +RT_API_ATTRS bool RealOutputEditingBase::EmitSuffix(const DataEdit &edit) { if (edit.descriptor == DataEdit::ListDirectedRealPart) { return EmitAscii( io_, edit.modes.editingFlags & decimalComma ? ";" : ",", 1); @@ -259,8 +261,10 @@ bool RealOutputEditingBase::EmitSuffix(const DataEdit &edit) { } template -decimal::ConversionToDecimalResult RealOutputEditing::ConvertToDecimal( +RT_API_ATTRS decimal::ConversionToDecimalResult +RealOutputEditing::ConvertToDecimal( int significantDigits, enum decimal::FortranRounding rounding, int flags) { +#if !defined(RT_DEVICE_COMPILATION) auto converted{decimal::ConvertToDecimal(buffer_, sizeof buffer_, static_cast(flags), significantDigits, rounding, x_)}; @@ -270,9 +274,13 @@ decimal::ConversionToDecimalResult RealOutputEditing::ConvertToDecimal( sizeof buffer_); } return converted; +#else // defined(RT_DEVICE_COMPILATION) + // TODO: enable Decimal library build for the device. + io_.GetIoErrorHandler().Crash("not implemented yet: decimal conversion"); +#endif // defined(RT_DEVICE_COMPILATION) } -static bool IsInfOrNaN(const char *p, int length) { +static RT_API_ATTRS bool IsInfOrNaN(const char *p, int length) { if (!p || length < 1) { return false; } @@ -287,7 +295,8 @@ static bool IsInfOrNaN(const char *p, int length) { // 13.7.2.3.3 in F'2018 template -bool RealOutputEditing::EditEorDOutput(const DataEdit &edit) { +RT_API_ATTRS bool RealOutputEditing::EditEorDOutput( + const DataEdit &edit) { addSpaceBeforeCharacter(io_); int editDigits{edit.digits.value_or(0)}; // 'd' field int editWidth{edit.width.value_or(0)}; // 'w' field @@ -423,7 +432,7 @@ bool RealOutputEditing::EditEorDOutput(const DataEdit &edit) { // 13.7.2.3.2 in F'2018 template -bool RealOutputEditing::EditFOutput(const DataEdit &edit) { +RT_API_ATTRS bool RealOutputEditing::EditFOutput(const DataEdit &edit) { addSpaceBeforeCharacter(io_); int fracDigits{edit.digits.value_or(0)}; // 'd' field const int editWidth{edit.width.value_or(0)}; // 'w' field @@ -553,12 +562,12 @@ bool RealOutputEditing::EditFOutput(const DataEdit &edit) { // 13.7.5.2.3 in F'2018 template -DataEdit RealOutputEditing::EditForGOutput(DataEdit edit) { +RT_API_ATTRS DataEdit RealOutputEditing::EditForGOutput(DataEdit edit) { edit.descriptor = 'E'; edit.variation = 'G'; // to suppress error for Ew.0 int editWidth{edit.width.value_or(0)}; - int significantDigits{ - edit.digits.value_or(BinaryFloatingPoint::decimalPrecision)}; // 'd' + int significantDigits{edit.digits.value_or( + static_cast(BinaryFloatingPoint::decimalPrecision))}; // 'd' if (editWidth > 0 && significantDigits == 0) { return edit; // Gw.0Ee -> Ew.0Ee for w > 0 } @@ -597,7 +606,8 @@ DataEdit RealOutputEditing::EditForGOutput(DataEdit edit) { // 13.10.4 in F'2018 template -bool RealOutputEditing::EditListDirectedOutput(const DataEdit &edit) { +RT_API_ATTRS bool RealOutputEditing::EditListDirectedOutput( + const DataEdit &edit) { decimal::ConversionToDecimalResult converted{ ConvertToDecimal(1, edit.modes.round)}; if (IsInfOrNaN(converted.str, static_cast(converted.length))) { @@ -631,7 +641,7 @@ bool RealOutputEditing::EditListDirectedOutput(const DataEdit &edit) { // E.g., 2. is edited into 0X8.0P-2 rather than 0X2.0P0. This implementation // follows that precedent so as to avoid a gratuitous incompatibility. template -auto RealOutputEditing::ConvertToHexadecimal( +RT_API_ATTRS auto RealOutputEditing::ConvertToHexadecimal( int significantDigits, enum decimal::FortranRounding rounding, int flags) -> ConvertToHexadecimalResult { if (x_.IsNaN() || x_.IsInfinite()) { @@ -689,7 +699,7 @@ auto RealOutputEditing::ConvertToHexadecimal( } template -bool RealOutputEditing::EditEXOutput(const DataEdit &edit) { +RT_API_ATTRS bool RealOutputEditing::EditEXOutput(const DataEdit &edit) { addSpaceBeforeCharacter(io_); int editDigits{edit.digits.value_or(0)}; // 'd' field int significantDigits{editDigits + 1}; @@ -740,7 +750,8 @@ bool RealOutputEditing::EditEXOutput(const DataEdit &edit) { EmitAscii(io_, exponent, expoLength); } -template bool RealOutputEditing::Edit(const DataEdit &edit) { +template +RT_API_ATTRS bool RealOutputEditing::Edit(const DataEdit &edit) { switch (edit.descriptor) { case 'D': return EditEorDOutput(edit); @@ -783,13 +794,14 @@ template bool RealOutputEditing::Edit(const DataEdit &edit) { return false; } -bool ListDirectedLogicalOutput(IoStatementState &io, +RT_API_ATTRS bool ListDirectedLogicalOutput(IoStatementState &io, ListDirectedStatementState &list, bool truth) { return list.EmitLeadingSpaceOrAdvance(io) && EmitAscii(io, truth ? "T" : "F", 1); } -bool EditLogicalOutput(IoStatementState &io, const DataEdit &edit, bool truth) { +RT_API_ATTRS bool EditLogicalOutput( + IoStatementState &io, const DataEdit &edit, bool truth) { switch (edit.descriptor) { case 'L': case 'G': @@ -813,7 +825,7 @@ bool EditLogicalOutput(IoStatementState &io, const DataEdit &edit, bool truth) { } template -bool ListDirectedCharacterOutput(IoStatementState &io, +RT_API_ATTRS bool ListDirectedCharacterOutput(IoStatementState &io, ListDirectedStatementState &list, const CHAR *x, std::size_t length) { bool ok{true}; @@ -870,8 +882,8 @@ bool ListDirectedCharacterOutput(IoStatementState &io, } template -bool EditCharacterOutput(IoStatementState &io, const DataEdit &edit, - const CHAR *x, std::size_t length) { +RT_API_ATTRS bool EditCharacterOutput(IoStatementState &io, + const DataEdit &edit, const CHAR *x, std::size_t length) { int len{static_cast(length)}; int width{edit.width.value_or(len)}; switch (edit.descriptor) { @@ -903,15 +915,15 @@ bool EditCharacterOutput(IoStatementState &io, const DataEdit &edit, EmitEncoded(io, x, std::min(width, len)); } -template bool EditIntegerOutput<1>( +template RT_API_ATTRS bool EditIntegerOutput<1>( IoStatementState &, const DataEdit &, std::int8_t); -template bool EditIntegerOutput<2>( +template RT_API_ATTRS bool EditIntegerOutput<2>( IoStatementState &, const DataEdit &, std::int16_t); -template bool EditIntegerOutput<4>( +template RT_API_ATTRS bool EditIntegerOutput<4>( IoStatementState &, const DataEdit &, std::int32_t); -template bool EditIntegerOutput<8>( +template RT_API_ATTRS bool EditIntegerOutput<8>( IoStatementState &, const DataEdit &, std::int64_t); -template bool EditIntegerOutput<16>( +template RT_API_ATTRS bool EditIntegerOutput<16>( IoStatementState &, const DataEdit &, common::int128_t); template class RealOutputEditing<2>; @@ -922,21 +934,23 @@ template class RealOutputEditing<10>; // TODO: double/double template class RealOutputEditing<16>; -template bool ListDirectedCharacterOutput(IoStatementState &, +template RT_API_ATTRS bool ListDirectedCharacterOutput(IoStatementState &, ListDirectedStatementState &, const char *, std::size_t chars); -template bool ListDirectedCharacterOutput(IoStatementState &, +template RT_API_ATTRS bool ListDirectedCharacterOutput(IoStatementState &, ListDirectedStatementState &, const char16_t *, std::size_t chars); -template bool ListDirectedCharacterOutput(IoStatementState &, +template RT_API_ATTRS bool ListDirectedCharacterOutput(IoStatementState &, ListDirectedStatementState &, const char32_t *, std::size_t chars); -template bool EditCharacterOutput( +template RT_API_ATTRS bool EditCharacterOutput( IoStatementState &, const DataEdit &, const char *, std::size_t chars); -template bool EditCharacterOutput( +template RT_API_ATTRS bool EditCharacterOutput( IoStatementState &, const DataEdit &, const char16_t *, std::size_t chars); -template bool EditCharacterOutput( +template RT_API_ATTRS bool EditCharacterOutput( IoStatementState &, const DataEdit &, const char32_t *, std::size_t chars); +RT_OFFLOAD_API_GROUP_END + } // namespace Fortran::runtime::io diff --git a/flang/runtime/edit-output.h b/flang/runtime/edit-output.h index 4e6d6b25b4dd2d..365bc2e2a4d10b 100644 --- a/flang/runtime/edit-output.h +++ b/flang/runtime/edit-output.h @@ -30,18 +30,20 @@ namespace Fortran::runtime::io { // one edit descriptor with a repeat factor may safely serve to edit // multiple elements of an array. template -bool EditIntegerOutput( +RT_API_ATTRS bool EditIntegerOutput( IoStatementState &, const DataEdit &, common::HostSignedIntType<8 * KIND>); // Encapsulates the state of a REAL output conversion. class RealOutputEditingBase { protected: - explicit RealOutputEditingBase(IoStatementState &io) : io_{io} {} + explicit RT_API_ATTRS RealOutputEditingBase(IoStatementState &io) : io_{io} {} // Returns null when the exponent overflows a fixed-size output field. - const char *FormatExponent(int, const DataEdit &edit, int &length); - bool EmitPrefix(const DataEdit &, std::size_t length, std::size_t width); - bool EmitSuffix(const DataEdit &); + RT_API_ATTRS const char *FormatExponent( + int, const DataEdit &edit, int &length); + RT_API_ATTRS bool EmitPrefix( + const DataEdit &, std::size_t length, std::size_t width); + RT_API_ATTRS bool EmitSuffix(const DataEdit &); IoStatementState &io_; int trailingBlanks_{0}; // created when Gw editing maps to Fw @@ -50,27 +52,29 @@ class RealOutputEditingBase { template class RealOutputEditing : public RealOutputEditingBase { public: + RT_VAR_GROUP_BEGIN static constexpr int binaryPrecision{common::PrecisionOfRealKind(KIND)}; + RT_VAR_GROUP_END using BinaryFloatingPoint = decimal::BinaryFloatingPointNumber; template - RealOutputEditing(IoStatementState &io, A x) + RT_API_ATTRS RealOutputEditing(IoStatementState &io, A x) : RealOutputEditingBase{io}, x_{x} {} - bool Edit(const DataEdit &); + RT_API_ATTRS bool Edit(const DataEdit &); private: // The DataEdit arguments here are const references or copies so that // the original DataEdit can safely serve multiple array elements when // it has a repeat count. - bool EditEorDOutput(const DataEdit &); - bool EditFOutput(const DataEdit &); - DataEdit EditForGOutput(DataEdit); // returns an E or F edit - bool EditEXOutput(const DataEdit &); - bool EditListDirectedOutput(const DataEdit &); + RT_API_ATTRS bool EditEorDOutput(const DataEdit &); + RT_API_ATTRS bool EditFOutput(const DataEdit &); + RT_API_ATTRS DataEdit EditForGOutput(DataEdit); // returns an E or F edit + RT_API_ATTRS bool EditEXOutput(const DataEdit &); + RT_API_ATTRS bool EditListDirectedOutput(const DataEdit &); - bool IsZero() const { return x_.IsZero(); } + RT_API_ATTRS bool IsZero() const { return x_.IsZero(); } - decimal::ConversionToDecimalResult ConvertToDecimal( + RT_API_ATTRS decimal::ConversionToDecimalResult ConvertToDecimal( int significantDigits, enum decimal::FortranRounding, int flags = 0); struct ConvertToHexadecimalResult { @@ -78,7 +82,7 @@ template class RealOutputEditing : public RealOutputEditingBase { int length; int exponent; }; - ConvertToHexadecimalResult ConvertToHexadecimal( + RT_API_ATTRS ConvertToHexadecimalResult ConvertToHexadecimal( int significantDigits, enum decimal::FortranRounding, int flags = 0); BinaryFloatingPoint x_; @@ -86,43 +90,43 @@ template class RealOutputEditing : public RealOutputEditingBase { EXTRA_DECIMAL_CONVERSION_SPACE]; }; -bool ListDirectedLogicalOutput( +RT_API_ATTRS bool ListDirectedLogicalOutput( IoStatementState &, ListDirectedStatementState &, bool); -bool EditLogicalOutput(IoStatementState &, const DataEdit &, bool); +RT_API_ATTRS bool EditLogicalOutput(IoStatementState &, const DataEdit &, bool); template -bool ListDirectedCharacterOutput(IoStatementState &, +RT_API_ATTRS bool ListDirectedCharacterOutput(IoStatementState &, ListDirectedStatementState &, const CHAR *, std::size_t chars); -extern template bool ListDirectedCharacterOutput(IoStatementState &, - ListDirectedStatementState &, const char *, - std::size_t chars); -extern template bool ListDirectedCharacterOutput(IoStatementState &, - ListDirectedStatementState &, const char16_t *, - std::size_t chars); -extern template bool ListDirectedCharacterOutput(IoStatementState &, - ListDirectedStatementState &, const char32_t *, - std::size_t chars); +extern template RT_API_ATTRS bool ListDirectedCharacterOutput( + IoStatementState &, ListDirectedStatementState &, + const char *, std::size_t chars); +extern template RT_API_ATTRS bool ListDirectedCharacterOutput( + IoStatementState &, ListDirectedStatementState &, + const char16_t *, std::size_t chars); +extern template RT_API_ATTRS bool ListDirectedCharacterOutput( + IoStatementState &, ListDirectedStatementState &, + const char32_t *, std::size_t chars); template -bool EditCharacterOutput( +RT_API_ATTRS bool EditCharacterOutput( IoStatementState &, const DataEdit &, const CHAR *, std::size_t chars); -extern template bool EditCharacterOutput( +extern template RT_API_ATTRS bool EditCharacterOutput( IoStatementState &, const DataEdit &, const char *, std::size_t chars); -extern template bool EditCharacterOutput( +extern template RT_API_ATTRS bool EditCharacterOutput( IoStatementState &, const DataEdit &, const char16_t *, std::size_t chars); -extern template bool EditCharacterOutput( +extern template RT_API_ATTRS bool EditCharacterOutput( IoStatementState &, const DataEdit &, const char32_t *, std::size_t chars); -extern template bool EditIntegerOutput<1>( +extern template RT_API_ATTRS bool EditIntegerOutput<1>( IoStatementState &, const DataEdit &, std::int8_t); -extern template bool EditIntegerOutput<2>( +extern template RT_API_ATTRS bool EditIntegerOutput<2>( IoStatementState &, const DataEdit &, std::int16_t); -extern template bool EditIntegerOutput<4>( +extern template RT_API_ATTRS bool EditIntegerOutput<4>( IoStatementState &, const DataEdit &, std::int32_t); -extern template bool EditIntegerOutput<8>( +extern template RT_API_ATTRS bool EditIntegerOutput<8>( IoStatementState &, const DataEdit &, std::int64_t); -extern template bool EditIntegerOutput<16>( +extern template RT_API_ATTRS bool EditIntegerOutput<16>( IoStatementState &, const DataEdit &, common::int128_t); extern template class RealOutputEditing<2>; diff --git a/flang/runtime/emit-encoded.h b/flang/runtime/emit-encoded.h index 864848c3b19c67..ac8c7d758a0d00 100644 --- a/flang/runtime/emit-encoded.h +++ b/flang/runtime/emit-encoded.h @@ -19,7 +19,8 @@ namespace Fortran::runtime::io { template -bool EmitEncoded(CONTEXT &to, const CHAR *data, std::size_t chars) { +RT_API_ATTRS bool EmitEncoded( + CONTEXT &to, const CHAR *data, std::size_t chars) { ConnectionState &connection{to.GetConnectionState()}; if (connection.access == Access::Stream && connection.internalIoCharKind == 0) { @@ -74,7 +75,7 @@ bool EmitEncoded(CONTEXT &to, const CHAR *data, std::size_t chars) { } template -bool EmitAscii(CONTEXT &to, const char *data, std::size_t chars) { +RT_API_ATTRS bool EmitAscii(CONTEXT &to, const char *data, std::size_t chars) { ConnectionState &connection{to.GetConnectionState()}; if (connection.internalIoCharKind <= 1 && connection.access != Access::Stream) { @@ -85,7 +86,7 @@ bool EmitAscii(CONTEXT &to, const char *data, std::size_t chars) { } template -bool EmitRepeated(CONTEXT &to, char ch, std::size_t n) { +RT_API_ATTRS bool EmitRepeated(CONTEXT &to, char ch, std::size_t n) { if (n <= 0) { return true; } diff --git a/flang/runtime/environment.h b/flang/runtime/environment.h index 9bc1158509615f..6c56993fb1d6ec 100644 --- a/flang/runtime/environment.h +++ b/flang/runtime/environment.h @@ -18,6 +18,7 @@ namespace Fortran::runtime { class Terminator; +RT_OFFLOAD_VAR_GROUP_BEGIN #if FLANG_BIG_ENDIAN constexpr bool isHostLittleEndian{false}; #elif FLANG_LITTLE_ENDIAN @@ -25,6 +26,7 @@ constexpr bool isHostLittleEndian{true}; #else #error host endianness is not known #endif +RT_OFFLOAD_VAR_GROUP_END // External unformatted I/O data conversions enum class Convert { Unknown, Native, LittleEndian, BigEndian, Swap }; diff --git a/flang/runtime/format-implementation.h b/flang/runtime/format-implementation.h index b84e3208271b75..45d4bd641f6f66 100644 --- a/flang/runtime/format-implementation.h +++ b/flang/runtime/format-implementation.h @@ -25,7 +25,7 @@ namespace Fortran::runtime::io { template -FormatControl::FormatControl(const Terminator &terminator, +RT_API_ATTRS FormatControl::FormatControl(const Terminator &terminator, const CharType *format, std::size_t formatLength, const Descriptor *formatDescriptor, int maxHeight) : maxHeight_{static_cast(maxHeight)}, format_{format}, @@ -63,7 +63,7 @@ FormatControl::FormatControl(const Terminator &terminator, } template -int FormatControl::GetIntField( +RT_API_ATTRS int FormatControl::GetIntField( IoErrorHandler &handler, CharType firstCh, bool *hadError) { CharType ch{firstCh ? firstCh : PeekNext()}; bool negate{ch == '-'}; @@ -114,7 +114,8 @@ int FormatControl::GetIntField( } template -static void HandleControl(CONTEXT &context, char ch, char next, int n) { +static RT_API_ATTRS void HandleControl( + CONTEXT &context, char ch, char next, int n) { MutableModes &modes{context.mutableModes()}; switch (ch) { case 'B': @@ -221,7 +222,8 @@ static void HandleControl(CONTEXT &context, char ch, char next, int n) { // Generally assumes that the format string has survived the common // format validator gauntlet. template -int FormatControl::CueUpNextDataEdit(Context &context, bool stop) { +RT_API_ATTRS int FormatControl::CueUpNextDataEdit( + Context &context, bool stop) { bool hitUnlimitedLoopEnd{false}; // Do repetitions remain on an unparenthesized data edit? while (height_ > 1 && format_[stack_[height_ - 1].start] != '(') { @@ -419,8 +421,8 @@ int FormatControl::CueUpNextDataEdit(Context &context, bool stop) { // Returns the next data edit descriptor template -Fortran::common::optional FormatControl::GetNextDataEdit( - Context &context, int maxRepeat) { +RT_API_ATTRS Fortran::common::optional +FormatControl::GetNextDataEdit(Context &context, int maxRepeat) { int repeat{CueUpNextDataEdit(context)}; auto start{offset_}; DataEdit edit; @@ -524,7 +526,7 @@ Fortran::common::optional FormatControl::GetNextDataEdit( } template -void FormatControl::Finish(Context &context) { +RT_API_ATTRS void FormatControl::Finish(Context &context) { CueUpNextDataEdit(context, true /* stop at colon or end of FORMAT */); if (freeFormat_) { FreeMemory(const_cast(format_)); diff --git a/flang/runtime/format.cpp b/flang/runtime/format.cpp index f219c29aaed142..6747253320f061 100644 --- a/flang/runtime/format.cpp +++ b/flang/runtime/format.cpp @@ -9,14 +9,18 @@ #include "format-implementation.h" namespace Fortran::runtime::io { +RT_OFFLOAD_API_GROUP_BEGIN template class FormatControl< InternalFormattedIoStatementState>; template class FormatControl< InternalFormattedIoStatementState>; +#if !defined(RT_DEVICE_COMPILATION) template class FormatControl< ExternalFormattedIoStatementState>; template class FormatControl< ExternalFormattedIoStatementState>; template class FormatControl>; template class FormatControl>; +#endif // !defined(RT_DEVICE_COMPILATION) +RT_OFFLOAD_API_GROUP_END } // namespace Fortran::runtime::io diff --git a/flang/runtime/format.h b/flang/runtime/format.h index e7d94559964041..f57cf920448712 100644 --- a/flang/runtime/format.h +++ b/flang/runtime/format.h @@ -12,6 +12,7 @@ #define FORTRAN_RUNTIME_FORMAT_H_ #include "environment.h" +#include "freestanding-tools.h" #include "io-error.h" #include "flang/Common/Fortran.h" #include "flang/Common/optional.h" @@ -49,20 +50,21 @@ struct DataEdit { char descriptor; // capitalized: one of A, I, B, O, Z, F, E(N/S/X), D, G // Special internal data edit descriptors for list-directed & NAMELIST I/O + RT_OFFLOAD_VAR_GROUP_BEGIN static constexpr char ListDirected{'g'}; // non-COMPLEX list-directed static constexpr char ListDirectedRealPart{'r'}; // emit "(r," or "(r;" static constexpr char ListDirectedImaginaryPart{'z'}; // emit "z)" static constexpr char ListDirectedNullValue{'n'}; // see 13.10.3.2 - constexpr bool IsListDirected() const { + static constexpr char DefinedDerivedType{'d'}; // DT defined I/O + RT_OFFLOAD_VAR_GROUP_END + constexpr RT_API_ATTRS bool IsListDirected() const { return descriptor == ListDirected || descriptor == ListDirectedRealPart || descriptor == ListDirectedImaginaryPart; } - constexpr bool IsNamelist() const { + constexpr RT_API_ATTRS bool IsNamelist() const { return IsListDirected() && modes.inNamelist; } - static constexpr char DefinedDerivedType{'d'}; // DT defined I/O - char variation{'\0'}; // N, S, or X for EN, ES, EX; G/l for original G/list Fortran::common::optional width; // the 'w' field; optional for A Fortran::common::optional digits; // the 'm' or 'd' field @@ -72,8 +74,10 @@ struct DataEdit { // "iotype" &/or "v_list" values for a DT'iotype'(v_list) // defined I/O data edit descriptor + RT_OFFLOAD_VAR_GROUP_BEGIN static constexpr std::size_t maxIoTypeChars{32}; static constexpr std::size_t maxVListEntries{4}; + RT_OFFLOAD_VAR_GROUP_END std::uint8_t ioTypeChars{0}; std::uint8_t vListEntries{0}; char ioType[maxIoTypeChars]; @@ -88,13 +92,13 @@ template class FormatControl { using Context = CONTEXT; using CharType = char; // formats are always default kind CHARACTER - FormatControl() {} - FormatControl(const Terminator &, const CharType *format, + RT_API_ATTRS FormatControl() {} + RT_API_ATTRS FormatControl(const Terminator &, const CharType *format, std::size_t formatLength, const Descriptor *formatDescriptor = nullptr, int maxHeight = maxMaxHeight); // For attempting to allocate in a user-supplied stack area - static std::size_t GetNeededSize(int maxHeight) { + static RT_API_ATTRS std::size_t GetNeededSize(int maxHeight) { return sizeof(FormatControl) - sizeof(Iteration) * (maxMaxHeight - maxHeight); } @@ -102,14 +106,15 @@ template class FormatControl { // Extracts the next data edit descriptor, handling control edit descriptors // along the way. If maxRepeat==0, this is a peek at the next data edit // descriptor. - Fortran::common::optional GetNextDataEdit( + RT_API_ATTRS Fortran::common::optional GetNextDataEdit( Context &, int maxRepeat = 1); // Emit any remaining character literals after the last data item (on output) // and perform remaining record positioning actions. - void Finish(Context &); + RT_API_ATTRS void Finish(Context &); private: + RT_OFFLOAD_VAR_GROUP_BEGIN static constexpr std::uint8_t maxMaxHeight{100}; struct Iteration { @@ -117,19 +122,20 @@ template class FormatControl { int start{0}; // offset in format_ of '(' or a repeated edit descriptor int remaining{0}; // while >0, decrement and iterate }; + RT_OFFLOAD_VAR_GROUP_END - void SkipBlanks() { + RT_API_ATTRS void SkipBlanks() { while (offset_ < formatLength_ && (format_[offset_] == ' ' || format_[offset_] == '\t' || format_[offset_] == '\v')) { ++offset_; } } - CharType PeekNext() { + RT_API_ATTRS CharType PeekNext() { SkipBlanks(); return offset_ < formatLength_ ? format_[offset_] : '\0'; } - CharType GetNextChar(IoErrorHandler &handler) { + RT_API_ATTRS CharType GetNextChar(IoErrorHandler &handler) { SkipBlanks(); if (offset_ >= formatLength_) { if (formatLength_ == 0) { @@ -143,7 +149,7 @@ template class FormatControl { } return format_[offset_++]; } - int GetIntField( + RT_API_ATTRS int GetIntField( IoErrorHandler &, CharType firstCh = '\0', bool *hadError = nullptr); // Advances through the FORMAT until the next data edit @@ -151,13 +157,14 @@ template class FormatControl { // along the way. Returns the repeat count that appeared // before the descriptor (defaulting to 1) and leaves offset_ // pointing to the data edit. - int CueUpNextDataEdit(Context &, bool stop = false); + RT_API_ATTRS int CueUpNextDataEdit(Context &, bool stop = false); - static constexpr CharType Capitalize(CharType ch) { + static constexpr RT_API_ATTRS CharType Capitalize(CharType ch) { return ch >= 'a' && ch <= 'z' ? ch + 'A' - 'a' : ch; } - void ReportBadFormat(Context &context, const char *msg, int offset) const { + RT_API_ATTRS void ReportBadFormat( + Context &context, const char *msg, int offset) const { if constexpr (std::is_same_v) { // Echo the bad format in the error message, but trim any leading or // trailing spaces. diff --git a/flang/runtime/freestanding-tools.h b/flang/runtime/freestanding-tools.h index bdc11ae93ac909..2bcbfee54cd21f 100644 --- a/flang/runtime/freestanding-tools.h +++ b/flang/runtime/freestanding-tools.h @@ -42,14 +42,19 @@ #define STD_REALLOC_UNSUPPORTED 1 #endif +#if !defined(STD_MEMCHR_UNSUPPORTED) && \ + (defined(__CUDACC__) || defined(__CUDA__)) && defined(__CUDA_ARCH__) +#define STD_MEMCHR_UNSUPPORTED 1 +#endif + namespace Fortran::runtime { #if STD_FILL_N_UNSUPPORTED // Provides alternative implementation for std::fill_n(), if // it is not supported. -template -static inline RT_API_ATTRS void fill_n( - A *start, std::size_t count, const A &value) { +template +static inline RT_API_ATTRS std::enable_if_t, void> +fill_n(A *start, std::size_t count, const B &value) { for (std::size_t j{0}; j < count; ++j) { start[j] = value; } @@ -134,5 +139,23 @@ static inline RT_API_ATTRS void *realloc(void *ptr, std::size_t newByteSize) { using std::realloc; #endif // !STD_REALLOC_UNSUPPORTED +#if STD_MEMCHR_UNSUPPORTED +// Provides alternative implementation for std::memchr(), if +// it is not supported. +static inline RT_API_ATTRS const void *memchr( + const void *ptr, int ch, std::size_t count) { + auto buf{reinterpret_cast(ptr)}; + auto c{static_cast(ch)}; + for (; count--; ++buf) { + if (*buf == c) { + return buf; + } + } + return nullptr; +} +#else // !STD_MEMCMP_UNSUPPORTED +using std::memchr; +#endif // !STD_MEMCMP_UNSUPPORTED + } // namespace Fortran::runtime #endif // FORTRAN_RUNTIME_FREESTANDING_TOOLS_H_ diff --git a/flang/runtime/internal-unit.cpp b/flang/runtime/internal-unit.cpp index 39fefd8c58cafb..93d046c48c814f 100644 --- a/flang/runtime/internal-unit.cpp +++ b/flang/runtime/internal-unit.cpp @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include "internal-unit.h" +#include "freestanding-tools.h" #include "io-error.h" #include "flang/Runtime/descriptor.h" #include @@ -14,8 +15,10 @@ namespace Fortran::runtime::io { +RT_OFFLOAD_API_GROUP_BEGIN + template -InternalDescriptorUnit::InternalDescriptorUnit(Scalar scalar, +RT_API_ATTRS InternalDescriptorUnit::InternalDescriptorUnit(Scalar scalar, std::size_t length, int kind, const Terminator &terminator, bool allocateOwnOutput) { internalIoCharKind = kind; @@ -37,7 +40,7 @@ InternalDescriptorUnit::InternalDescriptorUnit(Scalar scalar, } template -InternalDescriptorUnit::InternalDescriptorUnit( +RT_API_ATTRS InternalDescriptorUnit::InternalDescriptorUnit( const Descriptor &that, const Terminator &terminator) { auto thatType{that.type().GetCategoryAndKind()}; RUNTIME_CHECK(terminator, thatType.has_value()); @@ -52,7 +55,8 @@ InternalDescriptorUnit::InternalDescriptorUnit( endfileRecordNumber = d.Elements() + 1; } -template void InternalDescriptorUnit::EndIoStatement() { +template +RT_API_ATTRS void InternalDescriptorUnit::EndIoStatement() { if constexpr (DIR == Direction::Output) { if (usesOwnBuffer) { // Null terminate the buffer that contains just a single record. @@ -73,7 +77,7 @@ template void InternalDescriptorUnit::EndIoStatement() { } template -bool InternalDescriptorUnit::Emit( +RT_API_ATTRS bool InternalDescriptorUnit::Emit( const char *data, std::size_t bytes, IoErrorHandler &handler) { if constexpr (DIR == Direction::Input) { handler.Crash("InternalDescriptorUnit::Emit() called"); @@ -107,7 +111,7 @@ bool InternalDescriptorUnit::Emit( } template -std::size_t InternalDescriptorUnit::GetNextInputBytes( +RT_API_ATTRS std::size_t InternalDescriptorUnit::GetNextInputBytes( const char *&p, IoErrorHandler &handler) { if constexpr (DIR == Direction::Output) { handler.Crash("InternalDescriptorUnit::" @@ -128,7 +132,8 @@ std::size_t InternalDescriptorUnit::GetNextInputBytes( } template -bool InternalDescriptorUnit::AdvanceRecord(IoErrorHandler &handler) { +RT_API_ATTRS bool InternalDescriptorUnit::AdvanceRecord( + IoErrorHandler &handler) { if (currentRecordNumber >= endfileRecordNumber.value_or(0)) { if constexpr (DIR == Direction::Input) { handler.SignalEnd(); @@ -146,24 +151,25 @@ bool InternalDescriptorUnit::AdvanceRecord(IoErrorHandler &handler) { } template -void InternalDescriptorUnit::BlankFill(char *at, std::size_t bytes) { +RT_API_ATTRS void InternalDescriptorUnit::BlankFill( + char *at, std::size_t bytes) { switch (internalIoCharKind) { case 2: - std::fill_n(reinterpret_cast(at), bytes / 2, + Fortran::runtime::fill_n(reinterpret_cast(at), bytes / 2, static_cast(' ')); break; case 4: - std::fill_n(reinterpret_cast(at), bytes / 4, + Fortran::runtime::fill_n(reinterpret_cast(at), bytes / 4, static_cast(' ')); break; default: - std::fill_n(at, bytes, ' '); + Fortran::runtime::fill_n(at, bytes, ' '); break; } } template -void InternalDescriptorUnit::BlankFillOutputRecord() { +RT_API_ATTRS void InternalDescriptorUnit::BlankFillOutputRecord() { if constexpr (DIR == Direction::Output) { if (furthestPositionInRecord < recordLength.value_or(furthestPositionInRecord)) { @@ -174,18 +180,22 @@ void InternalDescriptorUnit::BlankFillOutputRecord() { } template -void InternalDescriptorUnit::BackspaceRecord(IoErrorHandler &handler) { +RT_API_ATTRS void InternalDescriptorUnit::BackspaceRecord( + IoErrorHandler &handler) { RUNTIME_CHECK(handler, currentRecordNumber > 1); --currentRecordNumber; BeginRecord(); } template -std::int64_t InternalDescriptorUnit::InquirePos() { +RT_API_ATTRS std::int64_t InternalDescriptorUnit::InquirePos() { return (currentRecordNumber - 1) * recordLength.value_or(0) + positionInRecord + 1; } template class InternalDescriptorUnit; template class InternalDescriptorUnit; + +RT_OFFLOAD_API_GROUP_END + } // namespace Fortran::runtime::io diff --git a/flang/runtime/internal-unit.h b/flang/runtime/internal-unit.h index fdf6e498eabf7e..92524d6a03a492 100644 --- a/flang/runtime/internal-unit.h +++ b/flang/runtime/internal-unit.h @@ -31,28 +31,30 @@ template class InternalDescriptorUnit : public ConnectionState { public: using Scalar = std::conditional_t; - InternalDescriptorUnit(Scalar, std::size_t chars, int kind, + RT_API_ATTRS InternalDescriptorUnit(Scalar, std::size_t chars, int kind, const Terminator &terminator, bool allocateOwnOutput = false); - InternalDescriptorUnit(const Descriptor &, const Terminator &); - void EndIoStatement(); + RT_API_ATTRS InternalDescriptorUnit(const Descriptor &, const Terminator &); + RT_API_ATTRS void EndIoStatement(); - bool Emit(const char *, std::size_t, IoErrorHandler &); - std::size_t GetNextInputBytes(const char *&, IoErrorHandler &); - bool AdvanceRecord(IoErrorHandler &); - void BackspaceRecord(IoErrorHandler &); - std::int64_t InquirePos(); + RT_API_ATTRS bool Emit(const char *, std::size_t, IoErrorHandler &); + RT_API_ATTRS std::size_t GetNextInputBytes(const char *&, IoErrorHandler &); + RT_API_ATTRS bool AdvanceRecord(IoErrorHandler &); + RT_API_ATTRS void BackspaceRecord(IoErrorHandler &); + RT_API_ATTRS std::int64_t InquirePos(); private: - Descriptor &descriptor() { return staticDescriptor_.descriptor(); } - const Descriptor &descriptor() const { + RT_API_ATTRS Descriptor &descriptor() { return staticDescriptor_.descriptor(); } - Scalar CurrentRecord() const { + RT_API_ATTRS const Descriptor &descriptor() const { + return staticDescriptor_.descriptor(); + } + RT_API_ATTRS Scalar CurrentRecord() const { return descriptor().template ZeroBasedIndexedElement( currentRecordNumber - 1); } - void BlankFill(char *, std::size_t); - void BlankFillOutputRecord(); + RT_API_ATTRS void BlankFill(char *, std::size_t); + RT_API_ATTRS void BlankFillOutputRecord(); StaticDescriptor staticDescriptor_; RT_OFFLOAD_VAR_GROUP_BEGIN diff --git a/flang/runtime/io-api.cpp b/flang/runtime/io-api.cpp index 094db5572f15c2..9cdacd13e0d3d5 100644 --- a/flang/runtime/io-api.cpp +++ b/flang/runtime/io-api.cpp @@ -99,13 +99,14 @@ Cookie IONAME(BeginInternalArrayFormattedInput)(const Descriptor &descriptor, } template -Cookie BeginInternalListIO( +RT_API_ATTRS Cookie BeginInternalListIO( std::conditional_t *internal, std::size_t internalLength, void ** /*scratchArea*/, - std::size_t /*scratchBytes*/, const char *sourceFile, int sourceLine) { + std::size_t /*scratchBytes*/, const char *sourceFile, int sourceLine, + bool allocateOwnOutput = false) { Terminator oom{sourceFile, sourceLine}; return &New>{oom}( - internal, internalLength, sourceFile, sourceLine) + internal, internalLength, sourceFile, sourceLine, allocateOwnOutput) .release() ->ioStatementState(); } @@ -227,11 +228,21 @@ Cookie BeginExternalListIO( } } +RT_EXT_API_GROUP_BEGIN +#if defined(RT_DEVICE_COMPILATION) +Cookie IODEF(BeginExternalListOutput)( + ExternalUnit unitNumber, const char *sourceFile, int sourceLine) { + return BeginInternalListIO(nullptr, 0, nullptr, 0, + sourceFile, sourceLine, /*allocateOwnOutput=*/true); +} +#else Cookie IONAME(BeginExternalListOutput)( ExternalUnit unitNumber, const char *sourceFile, int sourceLine) { return BeginExternalListIO( unitNumber, sourceFile, sourceLine); } +#endif +RT_EXT_API_GROUP_END Cookie IONAME(BeginExternalListInput)( ExternalUnit unitNumber, const char *sourceFile, int sourceLine) { @@ -1163,7 +1174,8 @@ bool IONAME(OutputInteger16)(Cookie cookie, std::int16_t n) { return descr::DescriptorIO(*cookie, descriptor); } -bool IONAME(OutputInteger32)(Cookie cookie, std::int32_t n) { +RT_EXT_API_GROUP_BEGIN +bool IODEF(OutputInteger32)(Cookie cookie, std::int32_t n) { if (!cookie->CheckFormattedStmtType("OutputInteger32")) { return false; } @@ -1173,6 +1185,7 @@ bool IONAME(OutputInteger32)(Cookie cookie, std::int32_t n) { TypeCategory::Integer, 4, reinterpret_cast(&n), 0); return descr::DescriptorIO(*cookie, descriptor); } +RT_EXT_API_GROUP_END bool IONAME(OutputInteger64)(Cookie cookie, std::int64_t n) { if (!cookie->CheckFormattedStmtType("OutputInteger64")) { @@ -1448,10 +1461,12 @@ bool IONAME(InquireInteger64)( return false; } -enum Iostat IONAME(EndIoStatement)(Cookie cookie) { +RT_EXT_API_GROUP_BEGIN +enum Iostat IODEF(EndIoStatement)(Cookie cookie) { IoStatementState &io{*cookie}; return static_cast(io.EndIoStatement()); } +RT_EXT_API_GROUP_END template static enum Iostat CheckUnitNumberInRangeImpl(INT unit, bool handleError, diff --git a/flang/runtime/io-error.cpp b/flang/runtime/io-error.cpp index c8f6675c60a6c8..18f353eeb1bce4 100644 --- a/flang/runtime/io-error.cpp +++ b/flang/runtime/io-error.cpp @@ -17,7 +17,10 @@ namespace Fortran::runtime::io { -void IoErrorHandler::SignalError(int iostatOrErrno, const char *msg, ...) { +RT_OFFLOAD_API_GROUP_BEGIN + +void RT_API_ATTRS IoErrorHandler::SignalError( + int iostatOrErrno, const char *msg, ...) { // Note that IOMSG= alone without IOSTAT=/END=/EOR=/ERR= does not suffice // for error recovery (see F'2018 subclause 12.11). switch (iostatOrErrno) { @@ -44,12 +47,20 @@ void IoErrorHandler::SignalError(int iostatOrErrno, const char *msg, ...) { if (ioStat_ <= 0) { ioStat_ = iostatOrErrno; // priority over END=/EOR= if (msg && (flags_ & hasIoMsg)) { +#if !defined(RT_DEVICE_COMPILATION) char buffer[256]; va_list ap; va_start(ap, msg); std::vsnprintf(buffer, sizeof buffer, msg, ap); - ioMsg_ = SaveDefaultCharacter(buffer, std::strlen(buffer) + 1, *this); va_end(ap); +#else + const char *buffer = "not implemented yet: IOSTAT with varargs"; +#endif + ioMsg_ = SaveDefaultCharacter( + buffer, Fortran::runtime::strlen(buffer) + 1, *this); +#if !defined(RT_DEVICE_COMPILATION) + va_end(ap); +#endif } } return; @@ -58,23 +69,31 @@ void IoErrorHandler::SignalError(int iostatOrErrno, const char *msg, ...) { } // I/O error not caught! if (msg) { +#if !defined(RT_DEVICE_COMPILATION) va_list ap; va_start(ap, msg); CrashArgs(msg, ap); va_end(ap); +#else + Crash("not implemented yet: IOSTAT with varargs"); +#endif } else if (const char *errstr{IostatErrorString(iostatOrErrno)}) { Crash(errstr); } else { +#if !defined(RT_DEVICE_COMPILATION) Crash("I/O error (errno=%d): %s", iostatOrErrno, std::strerror(iostatOrErrno)); +#else + Crash("I/O error (errno=%d)", iostatOrErrno); +#endif } } -void IoErrorHandler::SignalError(int iostatOrErrno) { +RT_API_ATTRS void IoErrorHandler::SignalError(int iostatOrErrno) { SignalError(iostatOrErrno, nullptr); } -void IoErrorHandler::Forward( +RT_API_ATTRS void IoErrorHandler::Forward( int ioStatOrErrno, const char *msg, std::size_t length) { if (ioStatOrErrno != IostatOk) { if (msg) { @@ -85,18 +104,20 @@ void IoErrorHandler::Forward( } } -void IoErrorHandler::SignalErrno() { SignalError(errno); } - -void IoErrorHandler::SignalEnd() { SignalError(IostatEnd); } +RT_API_ATTRS void IoErrorHandler::SignalEnd() { SignalError(IostatEnd); } -void IoErrorHandler::SignalEor() { SignalError(IostatEor); } +RT_API_ATTRS void IoErrorHandler::SignalEor() { SignalError(IostatEor); } -void IoErrorHandler::SignalPendingError() { +RT_API_ATTRS void IoErrorHandler::SignalPendingError() { int error{pendingError_}; pendingError_ = IostatOk; SignalError(error); } +RT_OFFLOAD_API_GROUP_END + +void IoErrorHandler::SignalErrno() { SignalError(errno); } + bool IoErrorHandler::GetIoMsg(char *buffer, std::size_t bufferLength) { const char *msg{ioMsg_.get()}; if (!msg) { @@ -132,7 +153,7 @@ bool IoErrorHandler::GetIoMsg(char *buffer, std::size_t bufferLength) { ToFortranDefaultCharacter(buffer, bufferLength, msg); return true; } else if (ok) { - std::size_t copied{std::strlen(buffer)}; + std::size_t copied{Fortran::runtime::strlen(buffer)}; if (copied < bufferLength) { std::memset(buffer + copied, ' ', bufferLength - copied); } diff --git a/flang/runtime/io-error.h b/flang/runtime/io-error.h index 565e7153351e7e..0fe11c9185c0a9 100644 --- a/flang/runtime/io-error.h +++ b/flang/runtime/io-error.h @@ -26,14 +26,15 @@ namespace Fortran::runtime::io { class IoErrorHandler : public Terminator { public: using Terminator::Terminator; - explicit IoErrorHandler(const Terminator &that) : Terminator{that} {} - void HasIoStat() { flags_ |= hasIoStat; } - void HasErrLabel() { flags_ |= hasErr; } - void HasEndLabel() { flags_ |= hasEnd; } - void HasEorLabel() { flags_ |= hasEor; } - void HasIoMsg() { flags_ |= hasIoMsg; } + explicit RT_API_ATTRS IoErrorHandler(const Terminator &that) + : Terminator{that} {} + RT_API_ATTRS void HasIoStat() { flags_ |= hasIoStat; } + RT_API_ATTRS void HasErrLabel() { flags_ |= hasErr; } + RT_API_ATTRS void HasEndLabel() { flags_ |= hasEnd; } + RT_API_ATTRS void HasEorLabel() { flags_ |= hasEor; } + RT_API_ATTRS void HasIoMsg() { flags_ |= hasIoMsg; } - bool InError() const { + RT_API_ATTRS bool InError() const { return ioStat_ != IostatOk || pendingError_ != IostatOk; } @@ -41,22 +42,25 @@ class IoErrorHandler : public Terminator { // Begin...() API routines before it is known whether they // have error handling control list items. Such statements // have an ErroneousIoStatementState with a pending error. - void SetPendingError(int iostat) { pendingError_ = iostat; } + RT_API_ATTRS void SetPendingError(int iostat) { pendingError_ = iostat; } - void SignalError(int iostatOrErrno, const char *msg, ...); - void SignalError(int iostatOrErrno); - template void SignalError(const char *msg, X &&...xs) { + RT_API_ATTRS void SignalError(int iostatOrErrno, const char *msg, ...); + RT_API_ATTRS void SignalError(int iostatOrErrno); + template + RT_API_ATTRS void SignalError(const char *msg, X &&...xs) { SignalError(IostatGenericError, msg, std::forward(xs)...); } - void Forward(int iostatOrErrno, const char *, std::size_t); + RT_API_ATTRS void Forward(int iostatOrErrno, const char *, std::size_t); void SignalErrno(); // SignalError(errno) - void SignalEnd(); // input only; EOF on internal write is an error - void SignalEor(); // non-advancing input only; EOR on write is an error - void SignalPendingError(); + RT_API_ATTRS void + SignalEnd(); // input only; EOF on internal write is an error + RT_API_ATTRS void + SignalEor(); // non-advancing input only; EOR on write is an error + RT_API_ATTRS void SignalPendingError(); - int GetIoStat() const { return ioStat_; } + RT_API_ATTRS int GetIoStat() const { return ioStat_; } bool GetIoMsg(char *, std::size_t); private: diff --git a/flang/runtime/io-stmt.cpp b/flang/runtime/io-stmt.cpp index 31d08d069dd4ce..093538e8ec812c 100644 --- a/flang/runtime/io-stmt.cpp +++ b/flang/runtime/io-stmt.cpp @@ -22,39 +22,46 @@ namespace Fortran::runtime::io { -bool IoStatementBase::Emit(const char *, std::size_t, std::size_t) { +RT_OFFLOAD_API_GROUP_BEGIN + +RT_API_ATTRS bool IoStatementBase::Emit( + const char *, std::size_t, std::size_t) { return false; } -std::size_t IoStatementBase::GetNextInputBytes(const char *&p) { +RT_API_ATTRS std::size_t IoStatementBase::GetNextInputBytes(const char *&p) { p = nullptr; return 0; } -bool IoStatementBase::AdvanceRecord(int) { return false; } +RT_API_ATTRS bool IoStatementBase::AdvanceRecord(int) { return false; } -void IoStatementBase::BackspaceRecord() {} +RT_API_ATTRS void IoStatementBase::BackspaceRecord() {} -bool IoStatementBase::Receive(char *, std::size_t, std::size_t) { +RT_API_ATTRS bool IoStatementBase::Receive(char *, std::size_t, std::size_t) { return false; } -Fortran::common::optional IoStatementBase::GetNextDataEdit( - IoStatementState &, int) { +RT_API_ATTRS Fortran::common::optional +IoStatementBase::GetNextDataEdit(IoStatementState &, int) { return Fortran::common::nullopt; } -ExternalFileUnit *IoStatementBase::GetExternalFileUnit() const { - return nullptr; -} +RT_API_ATTRS bool IoStatementBase::BeginReadingRecord() { return true; } + +RT_API_ATTRS void IoStatementBase::FinishReadingRecord() {} -bool IoStatementBase::BeginReadingRecord() { return true; } +RT_API_ATTRS void IoStatementBase::HandleAbsolutePosition(std::int64_t) {} -void IoStatementBase::FinishReadingRecord() {} +RT_API_ATTRS void IoStatementBase::HandleRelativePosition(std::int64_t) {} -void IoStatementBase::HandleAbsolutePosition(std::int64_t) {} +RT_API_ATTRS std::int64_t IoStatementBase::InquirePos() { return 0; } -void IoStatementBase::HandleRelativePosition(std::int64_t) {} +RT_OFFLOAD_API_GROUP_END + +ExternalFileUnit *IoStatementBase::GetExternalFileUnit() const { + return nullptr; +} bool IoStatementBase::Inquire(InquiryKeywordHash, char *, std::size_t) { return false; @@ -70,8 +77,6 @@ bool IoStatementBase::Inquire(InquiryKeywordHash, std::int64_t &) { return false; } -std::int64_t IoStatementBase::InquirePos() { return 0; } - void IoStatementBase::BadInquiryKeywordHashCrash(InquiryKeywordHash inquiry) { char buffer[16]; const char *decode{InquiryKeywordHashDecode(buffer, sizeof buffer, inquiry)}; @@ -79,19 +84,22 @@ void IoStatementBase::BadInquiryKeywordHashCrash(InquiryKeywordHash inquiry) { decode ? decode : "(cannot decode)"); } +RT_OFFLOAD_API_GROUP_BEGIN + template -InternalIoStatementState::InternalIoStatementState( - Buffer scalar, std::size_t length, const char *sourceFile, int sourceLine) +RT_API_ATTRS InternalIoStatementState::InternalIoStatementState( + Buffer scalar, std::size_t length, const char *sourceFile, int sourceLine, + bool allocateOwnOutput) : IoStatementBase{sourceFile, sourceLine}, - unit_{scalar, length, /*kind=*/1, *this} {} + unit_{scalar, length, /*kind=*/1, *this, allocateOwnOutput} {} template -InternalIoStatementState::InternalIoStatementState( +RT_API_ATTRS InternalIoStatementState::InternalIoStatementState( const Descriptor &d, const char *sourceFile, int sourceLine) : IoStatementBase{sourceFile, sourceLine}, unit_{d, *this} {} template -bool InternalIoStatementState::Emit( +RT_API_ATTRS bool InternalIoStatementState::Emit( const char *data, std::size_t bytes, std::size_t /*elementBytes*/) { if constexpr (DIR == Direction::Input) { Crash("InternalIoStatementState::Emit() called"); @@ -101,12 +109,13 @@ bool InternalIoStatementState::Emit( } template -std::size_t InternalIoStatementState::GetNextInputBytes(const char *&p) { +RT_API_ATTRS std::size_t InternalIoStatementState::GetNextInputBytes( + const char *&p) { return unit_.GetNextInputBytes(p, *this); } template -bool InternalIoStatementState::AdvanceRecord(int n) { +RT_API_ATTRS bool InternalIoStatementState::AdvanceRecord(int n) { while (n-- > 0) { if (!unit_.AdvanceRecord(*this)) { return false; @@ -115,11 +124,13 @@ bool InternalIoStatementState::AdvanceRecord(int n) { return true; } -template void InternalIoStatementState::BackspaceRecord() { +template +RT_API_ATTRS void InternalIoStatementState::BackspaceRecord() { unit_.BackspaceRecord(*this); } -template int InternalIoStatementState::EndIoStatement() { +template +RT_API_ATTRS int InternalIoStatementState::EndIoStatement() { if constexpr (DIR == Direction::Output) { unit_.EndIoStatement(); } @@ -131,39 +142,45 @@ template int InternalIoStatementState::EndIoStatement() { } template -void InternalIoStatementState::HandleAbsolutePosition(std::int64_t n) { +RT_API_ATTRS void InternalIoStatementState::HandleAbsolutePosition( + std::int64_t n) { return unit_.HandleAbsolutePosition(n); } template -void InternalIoStatementState::HandleRelativePosition(std::int64_t n) { +RT_API_ATTRS void InternalIoStatementState::HandleRelativePosition( + std::int64_t n) { return unit_.HandleRelativePosition(n); } template -std::int64_t InternalIoStatementState::InquirePos() { +RT_API_ATTRS std::int64_t InternalIoStatementState::InquirePos() { return unit_.InquirePos(); } template +RT_API_ATTRS InternalFormattedIoStatementState::InternalFormattedIoStatementState( Buffer buffer, std::size_t length, const CharType *format, std::size_t formatLength, const Descriptor *formatDescriptor, - const char *sourceFile, int sourceLine) - : InternalIoStatementState{buffer, length, sourceFile, sourceLine}, - ioStatementState_{*this}, format_{*this, format, formatLength, - formatDescriptor} {} + const char *sourceFile, int sourceLine, bool allocateOwnOutput) + : InternalIoStatementState{buffer, length, sourceFile, sourceLine, + allocateOwnOutput}, + ioStatementState_{*this}, + format_{*this, format, formatLength, formatDescriptor} {} template +RT_API_ATTRS InternalFormattedIoStatementState::InternalFormattedIoStatementState( const Descriptor &d, const CharType *format, std::size_t formatLength, const Descriptor *formatDescriptor, const char *sourceFile, int sourceLine) : InternalIoStatementState{d, sourceFile, sourceLine}, - ioStatementState_{*this}, format_{*this, format, formatLength, - formatDescriptor} {} + ioStatementState_{*this}, + format_{*this, format, formatLength, formatDescriptor} {} template -void InternalFormattedIoStatementState::CompleteOperation() { +RT_API_ATTRS void +InternalFormattedIoStatementState::CompleteOperation() { if (!this->completedOperation()) { if constexpr (DIR == Direction::Output) { format_.Finish(*this); @@ -174,25 +191,28 @@ void InternalFormattedIoStatementState::CompleteOperation() { } template -int InternalFormattedIoStatementState::EndIoStatement() { +RT_API_ATTRS int +InternalFormattedIoStatementState::EndIoStatement() { CompleteOperation(); return InternalIoStatementState::EndIoStatement(); } template -InternalListIoStatementState::InternalListIoStatementState( - Buffer buffer, std::size_t length, const char *sourceFile, int sourceLine) - : InternalIoStatementState{buffer, length, sourceFile, sourceLine}, +RT_API_ATTRS InternalListIoStatementState::InternalListIoStatementState( + Buffer buffer, std::size_t length, const char *sourceFile, int sourceLine, + bool allocateOwnOutput) + : InternalIoStatementState{buffer, length, sourceFile, sourceLine, + allocateOwnOutput}, ioStatementState_{*this} {} template -InternalListIoStatementState::InternalListIoStatementState( +RT_API_ATTRS InternalListIoStatementState::InternalListIoStatementState( const Descriptor &d, const char *sourceFile, int sourceLine) : InternalIoStatementState{d, sourceFile, sourceLine}, ioStatementState_{*this} {} template -void InternalListIoStatementState::CompleteOperation() { +RT_API_ATTRS void InternalListIoStatementState::CompleteOperation() { if (!this->completedOperation()) { if constexpr (DIR == Direction::Output) { if (unit_.furthestPositionInRecord > 0) { @@ -204,7 +224,7 @@ void InternalListIoStatementState::CompleteOperation() { } template -int InternalListIoStatementState::EndIoStatement() { +RT_API_ATTRS int InternalListIoStatementState::EndIoStatement() { CompleteOperation(); if constexpr (DIR == Direction::Input) { if (int status{ListDirectedStatementState::EndIoStatement()}; @@ -215,6 +235,8 @@ int InternalListIoStatementState::EndIoStatement() { return InternalIoStatementState::EndIoStatement(); } +RT_OFFLOAD_API_GROUP_END + ExternalIoStatementBase::ExternalIoStatementBase( ExternalFileUnit &unit, const char *sourceFile, int sourceLine) : IoStatementBase{sourceFile, sourceLine}, unit_{unit} {} diff --git a/flang/runtime/io-stmt.h b/flang/runtime/io-stmt.h index 4e17cee2becf87..7737302042b193 100644 --- a/flang/runtime/io-stmt.h +++ b/flang/runtime/io-stmt.h @@ -73,8 +73,8 @@ template <> class FormattedIoStatementState { template <> class FormattedIoStatementState { public: using AvailableOnDevice = std::true_type; - std::size_t GetEditDescriptorChars() const; - void GotChar(int); + RT_API_ATTRS std::size_t GetEditDescriptorChars() const; + RT_API_ATTRS void GotChar(int); private: // Account of characters read for edit descriptors (i.e., formatted I/O @@ -85,7 +85,7 @@ template <> class FormattedIoStatementState { // The Cookie type in the I/O API is a pointer (for C) to this class. class IoStatementState { public: - template explicit IoStatementState(A &x) : u_{x} {} + template explicit RT_API_ATTRS IoStatementState(A &x) : u_{x} {} // These member functions each project themselves into the active alternative. // They're used by per-data-item routines in the I/O API (e.g., OutputReal64) @@ -97,34 +97,37 @@ class IoStatementState { // It is called by EndIoStatement(), but it can be invoked earlier to // catch errors for (e.g.) GetIoMsg() and GetNewUnit(). If called // more than once, it is a no-op. - void CompleteOperation(); + RT_API_ATTRS void CompleteOperation(); // Completes an I/O statement and reclaims storage. - int EndIoStatement(); - - bool Emit(const char *, std::size_t bytes, std::size_t elementBytes = 0); - bool Receive(char *, std::size_t, std::size_t elementBytes = 0); - std::size_t GetNextInputBytes(const char *&); - bool AdvanceRecord(int = 1); - void BackspaceRecord(); - void HandleRelativePosition(std::int64_t byteOffset); - void HandleAbsolutePosition(std::int64_t byteOffset); // for r* in list I/O - Fortran::common::optional GetNextDataEdit(int maxRepeat = 1); + RT_API_ATTRS int EndIoStatement(); + + RT_API_ATTRS bool Emit( + const char *, std::size_t bytes, std::size_t elementBytes = 0); + RT_API_ATTRS bool Receive(char *, std::size_t, std::size_t elementBytes = 0); + RT_API_ATTRS std::size_t GetNextInputBytes(const char *&); + RT_API_ATTRS bool AdvanceRecord(int = 1); + RT_API_ATTRS void BackspaceRecord(); + RT_API_ATTRS void HandleRelativePosition(std::int64_t byteOffset); + RT_API_ATTRS void HandleAbsolutePosition( + std::int64_t byteOffset); // for r* in list I/O + RT_API_ATTRS Fortran::common::optional GetNextDataEdit( + int maxRepeat = 1); ExternalFileUnit *GetExternalFileUnit() const; // null if internal unit - bool BeginReadingRecord(); - void FinishReadingRecord(); + RT_API_ATTRS bool BeginReadingRecord(); + RT_API_ATTRS void FinishReadingRecord(); bool Inquire(InquiryKeywordHash, char *, std::size_t); bool Inquire(InquiryKeywordHash, bool &); bool Inquire(InquiryKeywordHash, std::int64_t, bool &); // PENDING= bool Inquire(InquiryKeywordHash, std::int64_t &); - std::int64_t InquirePos(); - void GotChar(signed int = 1); // for READ(SIZE=); can be <0 + RT_API_ATTRS std::int64_t InquirePos(); + RT_API_ATTRS void GotChar(signed int = 1); // for READ(SIZE=); can be <0 - MutableModes &mutableModes(); - ConnectionState &GetConnectionState(); - IoErrorHandler &GetIoErrorHandler() const; + RT_API_ATTRS MutableModes &mutableModes(); + RT_API_ATTRS ConnectionState &GetConnectionState(); + RT_API_ATTRS IoErrorHandler &GetIoErrorHandler() const; // N.B.: this also works with base classes - template A *get_if() const { + template RT_API_ATTRS A *get_if() const { [[maybe_unused]] std::size_t index{u_.index()}; return Fortran::common::visit( [=](auto &x) -> A * { @@ -145,7 +148,8 @@ class IoStatementState { } // Vacant after the end of the current record - Fortran::common::optional GetCurrentChar(std::size_t &byteCount); + RT_API_ATTRS Fortran::common::optional GetCurrentChar( + std::size_t &byteCount); // The "remaining" arguments to CueUpInput(), SkipSpaces(), & NextInField() // are always in units of bytes, not characters; the distinction matters @@ -153,7 +157,7 @@ class IoStatementState { // For fixed-width fields, return the number of remaining bytes. // Skip over leading blanks. - Fortran::common::optional CueUpInput(const DataEdit &edit) { + RT_API_ATTRS Fortran::common::optional CueUpInput(const DataEdit &edit) { Fortran::common::optional remaining; if (edit.IsListDirected()) { std::size_t byteCount{0}; @@ -171,7 +175,7 @@ class IoStatementState { return remaining; } - Fortran::common::optional SkipSpaces( + RT_API_ATTRS Fortran::common::optional SkipSpaces( Fortran::common::optional &remaining) { while (!remaining || *remaining > 0) { std::size_t byteCount{0}; @@ -196,15 +200,16 @@ class IoStatementState { // Acquires the next input character, respecting any applicable field width // or separator character. - Fortran::common::optional NextInField( + RT_API_ATTRS Fortran::common::optional NextInField( Fortran::common::optional &remaining, const DataEdit &); // Detect and signal any end-of-record condition after input. // Returns true if at EOR and remaining input should be padded with blanks. - bool CheckForEndOfRecord(std::size_t afterReading); + RT_API_ATTRS bool CheckForEndOfRecord(std::size_t afterReading); // Skips spaces, advances records, and ignores NAMELIST comments - Fortran::common::optional GetNextNonBlank(std::size_t &byteCount) { + RT_API_ATTRS Fortran::common::optional GetNextNonBlank( + std::size_t &byteCount) { auto ch{GetCurrentChar(byteCount)}; bool inNamelist{mutableModes().inNamelist}; while (!ch || *ch == ' ' || *ch == '\t' || (inNamelist && *ch == '!')) { @@ -218,7 +223,8 @@ class IoStatementState { return ch; } - template bool CheckFormattedStmtType(const char *name) { + template + RT_API_ATTRS bool CheckFormattedStmtType(const char *name) { if (get_if>()) { return true; } else { @@ -315,29 +321,31 @@ class IoStatementBase : public IoErrorHandler { public: using IoErrorHandler::IoErrorHandler; - bool completedOperation() const { return completedOperation_; } + RT_API_ATTRS bool completedOperation() const { return completedOperation_; } - void CompleteOperation() { completedOperation_ = true; } - int EndIoStatement() { return GetIoStat(); } + RT_API_ATTRS void CompleteOperation() { completedOperation_ = true; } + RT_API_ATTRS int EndIoStatement() { return GetIoStat(); } // These are default no-op backstops that can be overridden by descendants. - bool Emit(const char *, std::size_t bytes, std::size_t elementBytes = 0); - bool Receive(char *, std::size_t bytes, std::size_t elementBytes = 0); - std::size_t GetNextInputBytes(const char *&); - bool AdvanceRecord(int); - void BackspaceRecord(); - void HandleRelativePosition(std::int64_t); - void HandleAbsolutePosition(std::int64_t); - Fortran::common::optional GetNextDataEdit( + RT_API_ATTRS bool Emit( + const char *, std::size_t bytes, std::size_t elementBytes = 0); + RT_API_ATTRS bool Receive( + char *, std::size_t bytes, std::size_t elementBytes = 0); + RT_API_ATTRS std::size_t GetNextInputBytes(const char *&); + RT_API_ATTRS bool AdvanceRecord(int); + RT_API_ATTRS void BackspaceRecord(); + RT_API_ATTRS void HandleRelativePosition(std::int64_t); + RT_API_ATTRS void HandleAbsolutePosition(std::int64_t); + RT_API_ATTRS Fortran::common::optional GetNextDataEdit( IoStatementState &, int maxRepeat = 1); ExternalFileUnit *GetExternalFileUnit() const; - bool BeginReadingRecord(); - void FinishReadingRecord(); + RT_API_ATTRS bool BeginReadingRecord(); + RT_API_ATTRS void FinishReadingRecord(); bool Inquire(InquiryKeywordHash, char *, std::size_t); bool Inquire(InquiryKeywordHash, bool &); bool Inquire(InquiryKeywordHash, std::int64_t, bool &); bool Inquire(InquiryKeywordHash, std::int64_t &); - std::int64_t InquirePos(); + RT_API_ATTRS std::int64_t InquirePos(); void BadInquiryKeywordHashCrash(InquiryKeywordHash); @@ -352,14 +360,14 @@ class ListDirectedStatementState : public FormattedIoStatementState { public: using AvailableOnDevice = std::true_type; - bool EmitLeadingSpaceOrAdvance( + RT_API_ATTRS bool EmitLeadingSpaceOrAdvance( IoStatementState &, std::size_t = 1, bool isCharacter = false); - Fortran::common::optional GetNextDataEdit( + RT_API_ATTRS Fortran::common::optional GetNextDataEdit( IoStatementState &, int maxRepeat = 1); - bool lastWasUndelimitedCharacter() const { + RT_API_ATTRS bool lastWasUndelimitedCharacter() const { return lastWasUndelimitedCharacter_; } - void set_lastWasUndelimitedCharacter(bool yes = true) { + RT_API_ATTRS void set_lastWasUndelimitedCharacter(bool yes = true) { lastWasUndelimitedCharacter_ = yes; } @@ -371,20 +379,20 @@ class ListDirectedStatementState : public FormattedIoStatementState { public: using AvailableOnDevice = std::false_type; - bool inNamelistSequence() const { return inNamelistSequence_; } - int EndIoStatement(); + RT_API_ATTRS bool inNamelistSequence() const { return inNamelistSequence_; } + RT_API_ATTRS int EndIoStatement(); // Skips value separators, handles repetition and null values. // Vacant when '/' appears; present with descriptor == ListDirectedNullValue // when a null value appears. - Fortran::common::optional GetNextDataEdit( + RT_API_ATTRS Fortran::common::optional GetNextDataEdit( IoStatementState &, int maxRepeat = 1); // Each NAMELIST input item is treated like a distinct list-directed // input statement. This member function resets some state so that // repetition and null values work correctly for each successive // NAMELIST input item. - void ResetForNextNamelistItem(bool inNamelistSequence) { + RT_API_ATTRS void ResetForNextNamelistItem(bool inNamelistSequence) { remaining_ = 0; if (repeatPosition_) { repeatPosition_->Cancel(); @@ -412,21 +420,23 @@ class InternalIoStatementState : public IoStatementBase, std::true_type, std::false_type>; using Buffer = std::conditional_t; - InternalIoStatementState(Buffer, std::size_t, - const char *sourceFile = nullptr, int sourceLine = 0); - InternalIoStatementState( + RT_API_ATTRS InternalIoStatementState(Buffer, std::size_t, + const char *sourceFile = nullptr, int sourceLine = 0, + bool allocateOwnOutput = false); + RT_API_ATTRS InternalIoStatementState( const Descriptor &, const char *sourceFile = nullptr, int sourceLine = 0); - int EndIoStatement(); - - bool Emit(const char *data, std::size_t bytes, std::size_t elementBytes = 0); - std::size_t GetNextInputBytes(const char *&); - bool AdvanceRecord(int = 1); - void BackspaceRecord(); - ConnectionState &GetConnectionState() { return unit_; } - MutableModes &mutableModes() { return unit_.modes; } - void HandleRelativePosition(std::int64_t); - void HandleAbsolutePosition(std::int64_t); - std::int64_t InquirePos(); + RT_API_ATTRS int EndIoStatement(); + + RT_API_ATTRS bool Emit( + const char *data, std::size_t bytes, std::size_t elementBytes = 0); + RT_API_ATTRS std::size_t GetNextInputBytes(const char *&); + RT_API_ATTRS bool AdvanceRecord(int = 1); + RT_API_ATTRS void BackspaceRecord(); + RT_API_ATTRS ConnectionState &GetConnectionState() { return unit_; } + RT_API_ATTRS MutableModes &mutableModes() { return unit_.modes; } + RT_API_ATTRS void HandleRelativePosition(std::int64_t); + RT_API_ATTRS void HandleAbsolutePosition(std::int64_t); + RT_API_ATTRS std::int64_t InquirePos(); protected: bool free_{true}; @@ -442,17 +452,21 @@ class InternalFormattedIoStatementState std::true_type, std::false_type>; using CharType = CHAR; using typename InternalIoStatementState::Buffer; - InternalFormattedIoStatementState(Buffer internal, std::size_t internalLength, + RT_API_ATTRS InternalFormattedIoStatementState(Buffer internal, + std::size_t internalLength, const CharType *format, + std::size_t formatLength, const Descriptor *formatDescriptor = nullptr, + const char *sourceFile = nullptr, int sourceLine = 0, + bool allocateOwnOutput = false); + RT_API_ATTRS InternalFormattedIoStatementState(const Descriptor &, const CharType *format, std::size_t formatLength, const Descriptor *formatDescriptor = nullptr, const char *sourceFile = nullptr, int sourceLine = 0); - InternalFormattedIoStatementState(const Descriptor &, const CharType *format, - std::size_t formatLength, const Descriptor *formatDescriptor = nullptr, - const char *sourceFile = nullptr, int sourceLine = 0); - IoStatementState &ioStatementState() { return ioStatementState_; } - void CompleteOperation(); - int EndIoStatement(); - Fortran::common::optional GetNextDataEdit( + RT_API_ATTRS IoStatementState &ioStatementState() { + return ioStatementState_; + } + RT_API_ATTRS void CompleteOperation(); + RT_API_ATTRS int EndIoStatement(); + RT_API_ATTRS Fortran::common::optional GetNextDataEdit( IoStatementState &, int maxRepeat = 1) { return format_.GetNextDataEdit(*this, maxRepeat); } @@ -471,14 +485,17 @@ class InternalListIoStatementState : public InternalIoStatementState, using AvailableOnDevice = std::conditional_t; using typename InternalIoStatementState::Buffer; - InternalListIoStatementState(Buffer internal, std::size_t internalLength, - const char *sourceFile = nullptr, int sourceLine = 0); - InternalListIoStatementState( + RT_API_ATTRS InternalListIoStatementState(Buffer internal, + std::size_t internalLength, const char *sourceFile = nullptr, + int sourceLine = 0, bool allocateOwnOutput = false); + RT_API_ATTRS InternalListIoStatementState( const Descriptor &, const char *sourceFile = nullptr, int sourceLine = 0); - IoStatementState &ioStatementState() { return ioStatementState_; } + RT_API_ATTRS IoStatementState &ioStatementState() { + return ioStatementState_; + } using ListDirectedStatementState::GetNextDataEdit; - void CompleteOperation(); - int EndIoStatement(); + RT_API_ATTRS void CompleteOperation(); + RT_API_ATTRS int EndIoStatement(); private: IoStatementState ioStatementState_; // points to *this diff --git a/flang/runtime/iostat.cpp b/flang/runtime/iostat.cpp index c993b778e9e1f8..33d7b18cc79b72 100644 --- a/flang/runtime/iostat.cpp +++ b/flang/runtime/iostat.cpp @@ -9,7 +9,9 @@ #include "flang/Runtime/iostat.h" namespace Fortran::runtime::io { -const char *IostatErrorString(int iostat) { +RT_OFFLOAD_API_GROUP_BEGIN + +RT_API_ATTRS const char *IostatErrorString(int iostat) { switch (iostat) { case IostatOk: return "No error"; @@ -122,4 +124,6 @@ const char *IostatErrorString(int iostat) { } } +RT_OFFLOAD_API_GROUP_END + } // namespace Fortran::runtime::io diff --git a/flang/runtime/memory.cpp b/flang/runtime/memory.cpp index aa6ff9723d1a80..f3f0a6b047c243 100644 --- a/flang/runtime/memory.cpp +++ b/flang/runtime/memory.cpp @@ -7,12 +7,13 @@ //===----------------------------------------------------------------------===// #include "flang/Runtime/memory.h" +#include "freestanding-tools.h" #include "terminator.h" #include "tools.h" #include namespace Fortran::runtime { -RT_OFFLOAD_VAR_GROUP_BEGIN +RT_OFFLOAD_API_GROUP_BEGIN RT_API_ATTRS void *AllocateMemoryOrCrash( const Terminator &terminator, std::size_t bytes) { @@ -42,5 +43,5 @@ RT_API_ATTRS void *ReallocateMemoryOrCrash( RT_API_ATTRS void FreeMemory(void *p) { std::free(p); } -RT_OFFLOAD_VAR_GROUP_END +RT_OFFLOAD_API_GROUP_END } // namespace Fortran::runtime diff --git a/flang/runtime/namelist.cpp b/flang/runtime/namelist.cpp index ac9234f4af832b..4b48b8a788e65b 100644 --- a/flang/runtime/namelist.cpp +++ b/flang/runtime/namelist.cpp @@ -31,9 +31,12 @@ bool IONAME(OutputNamelist)(Cookie cookie, const NamelistGroup &group) { io.CheckFormattedStmtType("OutputNamelist"); io.mutableModes().inNamelist = true; ConnectionState &connection{io.GetConnectionState()}; + // The following lambda definition violates the conding style, + // but cuda-11.8 nvcc hits an internal error with the brace initialization. + // Internal function to advance records and convert case - const auto EmitUpperCase{[&](const char *prefix, std::size_t prefixLen, - const char *str, char suffix) -> bool { + const auto EmitUpperCase = [&](const char *prefix, std::size_t prefixLen, + const char *str, char suffix) -> bool { if ((connection.NeedAdvance(prefixLen) && !(io.AdvanceRecord() && EmitAscii(io, " ", 1))) || !EmitAscii(io, prefix, prefixLen) || @@ -49,7 +52,7 @@ bool IONAME(OutputNamelist)(Cookie cookie, const NamelistGroup &group) { } } return suffix == ' ' || EmitAscii(io, &suffix, 1); - }}; + }; // &GROUP if (!EmitUpperCase(" &", 2, group.groupName, ' ')) { return false; @@ -294,7 +297,7 @@ static bool HandleSubstring( ch = io.GetNextNonBlank(byteCount); } } - if (ch && ch == ':') { + if (ch && *ch == ':') { io.HandleRelativePosition(byteCount); ch = io.GetNextNonBlank(byteCount); if (ch) { @@ -587,7 +590,9 @@ bool IONAME(InputNamelist)(Cookie cookie, const NamelistGroup &group) { return true; } -bool IsNamelistNameOrSlash(IoStatementState &io) { +RT_OFFLOAD_API_GROUP_BEGIN + +RT_API_ATTRS bool IsNamelistNameOrSlash(IoStatementState &io) { if (auto *listInput{ io.get_if>()}) { if (listInput->inNamelistSequence()) { @@ -611,4 +616,6 @@ bool IsNamelistNameOrSlash(IoStatementState &io) { return false; } +RT_OFFLOAD_API_GROUP_END + } // namespace Fortran::runtime::io diff --git a/flang/runtime/namelist.h b/flang/runtime/namelist.h index 9a5da33a907e44..1fdc0eb4076eef 100644 --- a/flang/runtime/namelist.h +++ b/flang/runtime/namelist.h @@ -12,6 +12,7 @@ #define FORTRAN_RUNTIME_NAMELIST_H_ #include "non-tbp-dio.h" +#include "flang/Runtime/api-attrs.h" #include @@ -47,7 +48,7 @@ class NamelistGroup { // character; for use in disambiguating a name-like value (e.g. F or T) from a // NAMELIST group item name and for coping with short arrays. Always false // when not reading a NAMELIST. -bool IsNamelistNameOrSlash(IoStatementState &); +RT_API_ATTRS bool IsNamelistNameOrSlash(IoStatementState &); } // namespace Fortran::runtime::io #endif // FORTRAN_RUNTIME_NAMELIST_H_ diff --git a/flang/runtime/numeric-templates.h b/flang/runtime/numeric-templates.h index ecc3b2654d9652..cf4f7b6b7cd75a 100644 --- a/flang/runtime/numeric-templates.h +++ b/flang/runtime/numeric-templates.h @@ -122,12 +122,19 @@ template struct ABSTy { static constexpr RT_API_ATTRS T compute(T x) { return std::abs(x); } }; +// Suppress the warnings about calling __host__-only std::frexp, +// defined in C++ STD header files, from __device__ code. +RT_DIAG_PUSH +RT_DIAG_DISABLE_CALL_HOST_FROM_DEVICE_WARN + template struct FREXPTy { static constexpr RT_API_ATTRS T compute(T x, int *e) { return std::frexp(x, e); } }; +RT_DIAG_POP + template struct ILOGBTy { static constexpr RT_API_ATTRS int compute(T x) { return std::ilogb(x); } }; @@ -186,11 +193,6 @@ inline RT_API_ATTRS RESULT Exponent(ARG x) { } } -// Suppress the warnings about calling __host__-only std::frexp, -// defined in C++ STD header files, from __device__ code. -RT_DIAG_PUSH -RT_DIAG_DISABLE_CALL_HOST_FROM_DEVICE_WARN - // FRACTION (16.9.80) template inline RT_API_ATTRS T Fraction(T x) { if (ISNANTy::compute(x)) { @@ -205,8 +207,6 @@ template inline RT_API_ATTRS T Fraction(T x) { } } -RT_DIAG_POP - // SET_EXPONENT (16.9.171) template inline RT_API_ATTRS T SetExponent(T x, std::int64_t p) { if (ISNANTy::compute(x)) { diff --git a/flang/runtime/terminator.h b/flang/runtime/terminator.h index 444c68d109eedf..167574c7821b24 100644 --- a/flang/runtime/terminator.h +++ b/flang/runtime/terminator.h @@ -67,7 +67,7 @@ class Terminator { template RT_API_ATTRS void PrintCrashArgs(const char *message, Args... args) const { -#if RT_DEVICE_COMPILATION +#if defined(RT_DEVICE_COMPILATION) std::printf(message, args...); #else std::fprintf(stderr, message, args...); diff --git a/flang/runtime/tools.cpp b/flang/runtime/tools.cpp index e653323ed1de03..3d3fbaf70e2d00 100644 --- a/flang/runtime/tools.cpp +++ b/flang/runtime/tools.cpp @@ -175,7 +175,7 @@ RT_API_ATTRS void ShallowCopy(const Descriptor &to, const Descriptor &from) { RT_API_ATTRS char *EnsureNullTerminated( char *str, std::size_t length, Terminator &terminator) { - if (std::memchr(str, '\0', length) == nullptr) { + if (Fortran::runtime::memchr(str, '\0', length) == nullptr) { char *newCmd{(char *)AllocateMemoryOrCrash(terminator, length + 1)}; std::memcpy(newCmd, str, length); newCmd[length] = '\0'; diff --git a/flang/runtime/tools.h b/flang/runtime/tools.h index 392e3fc6c89136..d656e985d6b46d 100644 --- a/flang/runtime/tools.h +++ b/flang/runtime/tools.h @@ -430,7 +430,7 @@ template <> inline RT_API_ATTRS const char *FindCharacter( const char *data, char ch, std::size_t chars) { return reinterpret_cast( - std::memchr(data, static_cast(ch), chars)); + Fortran::runtime::memchr(data, static_cast(ch), chars)); } // Copy payload data from one allocated descriptor to another. diff --git a/flang/runtime/utf.cpp b/flang/runtime/utf.cpp index e9ccc2c04b6b07..8126c50c5ae447 100644 --- a/flang/runtime/utf.cpp +++ b/flang/runtime/utf.cpp @@ -11,7 +11,8 @@ namespace Fortran::runtime { // clang-format off -const std::uint8_t UTF8FirstByteTable[256]{ +RT_OFFLOAD_VAR_GROUP_BEGIN +const RT_CONST_VAR_ATTRS std::uint8_t UTF8FirstByteTable[256]{ /* 00 - 7F: 7 bit payload in single byte */ 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, @@ -37,10 +38,12 @@ const std::uint8_t UTF8FirstByteTable[256]{ /* FE: 32 bit payload */ 7, /* FF: invalid */ 0 }; +RT_OFFLOAD_VAR_GROUP_END // clang-format on +RT_OFFLOAD_API_GROUP_BEGIN // Non-minimal encodings are accepted. -Fortran::common::optional DecodeUTF8(const char *p0) { +RT_API_ATTRS Fortran::common::optional DecodeUTF8(const char *p0) { const std::uint8_t *p{reinterpret_cast(p0)}; std::size_t bytes{MeasureUTF8Bytes(*p0)}; if (bytes == 1) { @@ -61,7 +64,7 @@ Fortran::common::optional DecodeUTF8(const char *p0) { return Fortran::common::nullopt; } -std::size_t EncodeUTF8(char *p0, char32_t ucs) { +RT_API_ATTRS std::size_t EncodeUTF8(char *p0, char32_t ucs) { std::uint8_t *p{reinterpret_cast(p0)}; if (ucs <= 0x7f) { p[0] = ucs; @@ -107,5 +110,6 @@ std::size_t EncodeUTF8(char *p0, char32_t ucs) { return 7; } } +RT_OFFLOAD_API_GROUP_END } // namespace Fortran::runtime diff --git a/flang/runtime/utf.h b/flang/runtime/utf.h index 2b4e4f9a188758..29670d54b3eb6f 100644 --- a/flang/runtime/utf.h +++ b/flang/runtime/utf.h @@ -49,20 +49,22 @@ namespace Fortran::runtime { // Derive the length of a UTF-8 character encoding from its first byte. // A zero result signifies an invalid encoding. -extern const std::uint8_t UTF8FirstByteTable[256]; -static inline std::size_t MeasureUTF8Bytes(char first) { +RT_OFFLOAD_VAR_GROUP_BEGIN +extern const RT_CONST_VAR_ATTRS std::uint8_t UTF8FirstByteTable[256]; +static constexpr std::size_t maxUTF8Bytes{7}; +RT_OFFLOAD_VAR_GROUP_END + +static inline RT_API_ATTRS std::size_t MeasureUTF8Bytes(char first) { return UTF8FirstByteTable[static_cast(first)]; } -static constexpr std::size_t maxUTF8Bytes{7}; - // Ensure that all bytes are present in sequence in the input buffer // before calling; use MeasureUTF8Bytes(first byte) to count them. -Fortran::common::optional DecodeUTF8(const char *); +RT_API_ATTRS Fortran::common::optional DecodeUTF8(const char *); // Ensure that at least maxUTF8Bytes remain in the output // buffer before calling. -std::size_t EncodeUTF8(char *, char32_t); +RT_API_ATTRS std::size_t EncodeUTF8(char *, char32_t); } // namespace Fortran::runtime #endif // FORTRAN_RUNTIME_UTF_H_