A sticky CUDA error (e.g. an illegal memory access in a kernel) leaves
the primary context active but unusable, so CUFDeviceIsActive() reports
it as fine and the compiler-generated scope-exit frees abort a program
that ran to completion: 'cudaFree(p)' failed with
'cudaErrorIllegalAddress'.
Detect this by freeing a null pointer, a no-op that still reports the
sticky error. It runs only once the primary context is known active, so
it cannot lazily create one.
Resolve cuDevicePrimaryCtxGetState via the installed driver version, and
prefer running scope-exit cleanup when the probe cannot prove teardown,
so local device allocations are not leaked when the unversioned lookup
fails.
The CUDA device runtime probes primary-context state during scope-exit
cleanup via a one-time
cudaGetDriverEntryPoint(`cuDevicePrimaryCtxGetState`) lookup. On some
runtime/driver combinations (e.g. a CUDA runtime newer than the
installed driver) that lookup fails and leaves a sticky error in the
thread's last-error slot. Since the probe runs between user kernel
launches, the user's next `cudaGetLastError()` misattributes the
leftover error to the launch, causing a spurious `invalid argument`
failure.
Fix: snapshot the last-error on entry and consume an error only if the
slot started clean, so the probe never leaks its own error and never
discards a pre-existing one.
Reproducer: a local device variable forces the guarded scope-exit device
cleanup on every call, so the probe's one-time lookup runs between the
two launches; the second `cudaGetLastError()` then observes the stale
error.
```
module m
contains
attributes(global) subroutine kern(a)
real*8 :: a(:)
if (threadIdx%x <= size(a)) a(threadIdx%x) = 1.0d0
end subroutine
end module
module h
real*8, device, allocatable :: d(:)
contains
subroutine step()
use cudafor
use m
integer :: ierr
real*8, device :: s ! local device var -> per-call guarded cleanup
if (.not. allocated(d)) allocate(d(32))
call kern<<<1,32>>>(d)
ierr = cudaGetLastError()
if (ierr /= 0) then
print *, "Kernel Launch Error ierr=", ierr, trim(cudaGetErrorString(ierr))
stop 1
endif
print *, "step ok"
end subroutine
end module
program main
use h
call step()
call step()
print *, "done"
end program
```
Before the fix (on an affected runtime): prints "step ok" then
`Kernel Launch Error ierr= 1 invalid argument`.
After the fix: prints "step ok" / "step ok" / "done".
Example:
```fortran
program main
use cudafor
integer, allocatable, device :: b(:)
integer :: istat
allocate(b(1024))
b = 2
istat = cudaDeviceReset()
end program
```
Managed descriptors and main-program allocatable data could leak. After
`cudaDeviceReset()`, compiler-generated scope-exit cleanup could also
access
stale descriptors or recreate a CUDA context.
Fix: free device data and descriptors in the correct order, and avoid
unnecessary
managed temporary descriptors. Each routine’s CUDA scope-exit cleanup is
grouped
under one `CUFDeviceIsActive()` check using
`cuDevicePrimaryCtxGetState`.
Routines without CUDA cleanup—and individual runtime frees—incur no
query.
Two fixes when using CMake < 3.28 after #209482:
* In addition to `flang`, also recognize `fakeflang` to determine
whether any workaround for CMake not completely supporting Flang is
needed. `fakeflang` is used as a shim in #209482.
* For CMake < 3.24, CMake does not reognize Flang at all and tries other
methods to detect to the compiler, including the file `CMakeTestGNU.c`.
Make `fakeflang` recognize it as a probe attempt.
This should fix the
[openmp-offload-amdgpu-clang-flang](https://lab.llvm.org/buildbot/#/builders/67)
buildbot. The buildbot is currently upgrade to a newer CMake version,
and the minimum required version of CMake will become 3.31 soon.
The external shell will be deleted soon, so drop the option to enable
the use of it for testing flang/flang-rt.
Reviewers: petrhosek, clementval, ilovepi
Pull Request: https://github.com/llvm/llvm-project/pull/209624
`time-intrinsic.cpp` implements the timing intrinsics as SFINAE overload
sets (preferred/fallback per platform), so on any given platform some
overloads are intentionally never instantiated and trip
-Wunused-template. Mark all variants [[maybe_unused]].
`StoreElement` and `MakeArray` in `unittests/Runtime/tools.h` are static
function templates in a header, so any test TU that includes it without
using all of them trips the warning.
Part of #202945
Use CMake INTERFACE libraries to express dependencies and common build
options.
Using INTERFACE libraries was proposed here:
https://github.com/llvm/llvm-project/pull/171610#discussion_r3122187957
Since such libraries can also propagate build options use this instead
of `target_compile_options`. Unfortunately this does not mean we can get
rid of `flang_module_target` because it also needs to set the module
output directory.
---------
Co-authored-by: Tarun Prabhu <tarunprabhu@gmail.com>
MallocWrapper called std::aligned_alloc for over-aligned requests, but
that C11 function is only available on macOS 10.15 and newer. flang-rt
builds with a Darwin deployment target of 10.7 (set in
AddFlangRT.cmake), so the build failed under
-Werror=unguarded-availability-new.
Use posix_memalign instead, as it is available on all supported POSIX
targets.
The work queue system is also used in the assignment and could lead to
unresolved function when building non relocatable code for the gpu using
the CUDA thin I/O. Just add a guard to keep descriptor I/O out of the
CUDA thin I/O build.
…bytes
in https://github.com/llvm/llvm-project/pull/194969 the default array
alignment was changed to 64bytes for Fortran globals. In this PR the
malloc wrapper in the flang runtime is modified to accept an alignment
argument (0 default) and for cases above 16 byte (std::max_align_t) uses
`aligned_alloc` instead of malloc.
The issue of alignment was discussed in this RFC:
https://discourse.llvm.org/t/rfc-alignment-of-global-arrays/90397/13
Assisted-by: Opus 4.6
This patch to set OBJECT_MODE=any on AIX to have tools able to handle
both different bit mode objects.
This fixes the compare_iso_fortran_env_symbols.f90 failure on AIX. The
default of llvm-nm is 32-bit but the objects in the test is 64-bit.
Setting OBJECT_MODE=any allows the tool accepts either bit modes.
When compiling CUDA-Fortran, a use of the module __cuda_builtins is
added implicitly by the compiler. Add explicit use so CMake can consider
it in the built order. The other modules that are compiled with CUDA
enabled (`cudadevice.f90` and `cooperativegroups.f90`) transitively
depend on depend on __cuda_builtins as well.
Two build failures reported after #204260
* Unix Makefiles generator stops working: The cause is that the rules
for building each OBJECT library lands in its own Makefile, e.g.
`flang_rt.mod.fortran.builtins.dir/build.make` and
`libomp-mod.dir/build.make`. Trying to inject dependencies directly for
build rules in the other file does not work.
* `__ppc_types.f90` not tracked: Forgotten in #204260 due to being only
conditionally enabled for PowerPC targets.
The solution for both is to just remove the workaround for CMake not
recognizing modules uses declared using `intrinsic` which caused these
problems. This PR promotes the `use` constructs in the module sources to
normal dependencies that are not ignored by CMake.
The `intrinsic` modifier changes the search path to only look for such a
module in `-fintrinsic-modules-path`. The difference is only
theoretical:
1. When compiling the intrinsic modules, there are no other search paths
than for the intrinsics anyway
2. Users should strongly avoid naming their modules the same as an
intrinsic module
3. Most of our intrinsic sources don't even use the `intrinsic` modifier
in the first place (e.g.
1c9412432f/flang-rt/lib/runtime/__cuda_builtins.f90 (L13)).
That is, if there was a problem with lookup ambiguity, these would have
already been a problem for those.
4. The checksum of used modules is only added into the .mod file (the
only reason why modules would need to be transitively rebuilt) when a
`use` statement appears. Many intrinsic modules don't explicitly use the
modules that are added implicitly by the compiler, such as
`__fortran_builtins`, and hence do not even need to be rebuilt.
I started this PR with a tested-working fix
([71c5e7f0ffcf2ebd58812d2e2812b22663be1631](https://github.com/llvm/llvm-project/pull/205634/changes/BASE..71c5e7f0ffcf2ebd58812d2e2812b22663be1631))
to the workaround code. But given the points above, there seems to be no
benefit of keeping the workaround code handling CMake version < 4.5
around just for some intrinsic modules to keep the `intrinsic` modifier
with no practical effect. A non-workaround path using CMake's
`Fortran_BUILDING_IN(S)TRINSIC_MODULES` setting had to be added in any
case for Unix Makefile generators (reported issue 1.) because injecting
dependencies is not possible with Unix Makefiles as stated, but
`Fortran_BUILDING_IN(S)TRINSIC_MODULES` also does not work with the
Ninja generator. Avoiding `use, intrinsic` keeps just one configuration
path.
I tested that modules are properly updated with Ninja and Unix
Makefiles:
* Build from scratch
* Modifying `__fortran_builtins.f90` rebuilds its dependent modules
* Modifying `iso_c_binding.f90` rebuilds its dependent modules
(omp_lib.mod)
CMake currently intentionally ignores intrinsic dependencies. flang-rt
already had a workaround using target-level dependencies, but it does
not know about dependencies between .mod files created within the same
add_library that CMake ignores. As a result, as reported in #203549,
updating a .mod did not trigger rebuilding the .mod files that depend on
it. Specifically, .mod files store the checksum of used .mod files which
need to be updated and therefore require transitive rebuidling.
As mentioned, CMake already adds this file-level dependency itself for
non-intrinsic modules dependencies. In this PR we are injecting the
additional dependencies that CMake does not add via OBJECT_DEPENDS.
Three caveats:
1. Using OBJECT_DEPENDS for dependencies between modules of the same
OBJECT library makes Ninja complain about circular dependencies. To
avoid, split __fortran_builtins.f90 and __cuda_builtins.f90 into their
own OBJECT libraries each.
2. The dependency cannot be on the .mod files. Because of how CMake's
dependency mechanism works, the dependency scanning runs only after
ensuring that the OBJECT_DEPENDS files exist. So if they don't exist
yet, we get a missing dependency error. To avoid, we depend on the .o
file instead.
3. Generator expressions do not work in OBJECT_DEPENDS. That is, we
cannot use `$<TARGET_FILES:flang_rt.mod.fortran.builtins>`, but
fortunately we can use OBJECT_OUTPUTS to make the location predictable.
At some point in the future,
https://gitlab.kitware.com/cmake/cmake/-/issues/26803 should be resolved
and become the minimum required version to build LLVM, in which case
these workarounds can be removed.
A WRITE after ENDFILE with ERR= or IOSTAT= was crashing instead of
handling the error properly. Earlier, the program was crashing because
the error was triggered too early (before error handling was ready).
---------
Co-authored-by: Jay Satish Kumar Patel <kumarpat@hpe.com>
The representation of `TEAM_TYPE` in the `ISO_FORTRAN_ENV` module is
opaque to the compiler, but the size (and alignment) needs to match the
PRIF specification for `prif_team_type` to ensure ABI compatibility with
the multi-image runtime library.
Flang's old definition as a derived type containing only an
`integer(int64)` component was leading to an 8-byte `TEAM_TYPE`. However
PRIF specifies the `prif_team_type` component as a `pointer` to an
opaque scalar type, which flang compiles to a 40-byte representation.
This mismatch was leading to incorrect behavior at runtime for programs
compiled with `-fcoarray` and using `TEAM_TYPE`, where `TEAM_TYPE`
values returned by PRIF were being silently truncated by the compiler.
Change the declaration of `__builtin_team_type` to match the PRIF
specification, thus ensuring the representation of `TEAM_TYPE` has the
required size.
Add binary 'B' input edit descriptor support for logical values to flang
runtime EditLogicalInput(). The logical binary 'B' edit descriptor is
already supported in EditLogicalOutput().
Co-authored-by: John Otken john.otken@hpe.com
Summary:
This PR cuts `flang` out of the core runtimes configure dependency. We
will need the runtimes infrastructure to handle `flang` module
generation, but this dependency poisons the entire dependency stack.
`flang` and `mlir` are by far the heavily parts of the LLVM stack and
for distribution we want to only build it when absolutely necessary,
which as of now is only to install flang module files.
The approach here is to simply remove the `flang` target from the core
dependency tree, but intead configure it for the top-level `module`
targets which are part of `all`. To make this work we need to pass
COMPILER_WORKS and set up a dummy location so that configuration passes.
it's a little backdoor, but this is an extremely important quality of
life improvement for LLVM distribution support.
IEEE-754 binary16 ("half") maps cleanly to the C `_Float16` type
(ISO/IEC TS 18661-3), but flang previously rejected `REAL(KIND=2)` and
`COMPLEX(KIND=2)` in C-interoperable contexts. Make `REAL(KIND=2)` and
`COMPLEX(KIND=2)` into actual interoperable types.
`ISO_C_BINDING` now exports the gfortran-compatible named constants
`c_float16` and `c_float16_complex` (both value 2), the kind parameter
for the half-precision C interoperable types.
Assisted-by: AI
Treat '!' as a self-delimiting value separator when reading NAMELIST
input, so that "name=value!comment" is accepted without an intervening
blank, comma, slash, or end of record. This matches gfortran, ifx, and
classic nvfortran behavior on real-world namelist input files.
F2023 13.11.3.6 p.1 requires a value separator before a '!' comment
introducer in namelist input, so this is a documented extension. The
change does not affect '!' characters inside character literal
constants, which continue to be taken literally.
The extension was documented in flang/docs/Extensions.md.
Assisted-by: AI
Summary:
Ever since https://github.com/llvm/llvm-project/pull/171515 we now build
the moduel files in the runtime step. This is problematic because pretty
much all of the tests depend on them. The current dependency chain
doesn't work correctly because the dependencies are set much later.
This fix just adds a global property so that we can set these
out-of-order and sets that property to depend on all the runtime tests.
Flang now becomes the only test suite who depends on the runtimes like
this, every other project's test suite is hermetic. However, moving
pretty much every flang test to flang-rt isn't ideal. We can investigate
this in depth later but for now this fixes the build.
Reapplication of #137828 and #169638
Move building the .mod files from openmp/flang to openmp/flang-rt using
a shared mechanism. Motivations to do so are:
1. Most modules are target-dependent and need to be re-compiled for each
target separately, which is something the LLVM_ENABLE_RUNTIMES system
already does. Prime example is `iso_c_binding.mod` which encodes the
target's ABI. Constants such as `c_long_double` also have different
values).
Most other modules have `#ifdef`-enclosed code as well. For instance
this caused offload targets nvptx64-nvidia-cuda/amdgpu-amd-amdhsa to use
the modules files compiled for the host which may contrain uses of the
types REAL(10) or REAL(16) not available for nvptx/amdgpu.
2. CMake has support for Fortran that we should use. Among other things,
it automatically determines module dependencies so there is no need to
hardcode them in the CMakeLists.txt.
Currently the dependencies between builtin modules and `omp_lib.mod`
seem to be incorrect:
https://github.com/llvm/llvm-project/pull/174474#issuecomment-3880189581
3. It allows using Fortran itself to implement Flang-RT. Currently, only
`iso_fortran_env_impl.f90` emits object files that are needed by Fortran
applications (#89403). The workaround of #95388 could be reverted (PR
#169525).
If using Flang for cross-compilation or target-offloading, flang-rt must
now be compiled for each target not only for the library, but also to
get the target-specific module files. For instance in a bootstrapping
runtime build, this can be done by adding:
`-DLLVM_RUNTIME_TARGETS=default;nvptx64-nvidia-cuda;amdgpu-amd-amdhsa`.
Some new dependencies come into play:
* openmp depends on flang-rt for building `lib_omp.mod` and
`lib_omp_kinds.mod`. Currently, if flang-rt is not found then the
modules are not built.
* check-flang depends on flang-rt: If not found, the majority of tests
are disabled. If not building in a bootstrpping build, the location of
the module files can be pointed to using
`-DFLANG_INTRINSIC_MODULES_DIR=<path>`, e.g. in a flang-standalone
build. Alternatively, the test needing any of the intrinsic modules
could be marked with `REQUIRES: flangrt-modules`.
* check-flang depends on openmp: Not a change; tests requiring
`lib_omp.mod` and `lib_omp_kinds.mod` those are already marked with
`openmp_runtime`.
As intrinsic are now specific to the target, their location is moved
from `include/flang` to `<resource-dir>/finclude/flang/<triple>`. The
mechnism to compute the location have been moved from flang-rt
(previously used to compute the location of `libflang_rt.*.a`) to common
locations in `cmake/GetToolchainDirs.cmake` and
`runtimes/CMakeLists.txt` so they can be used by both, openmp and
flang-rt. Potentially the mechnism could also be shared by other
libraries such as compiler-rt.
`finclude` was chosen because `gfortran` uses it as well and avoids
misuse such as `#include <flang/iso_c_binding.mod>`. The search location
is now determined by `ToolChain` in the driver, instead of by the
frontend. Another subdirectory `flang` avoids accidental inclusion of
gfortran-modules which due to compression would result in
user-unfriendly errors. Now the driver adds `-fintrinsic-module-path`
for that location to the frontend call (Just like gfortran does).
Since the driver determines the location,
tests invoking `flang -fc1` and `bbc` must also be passed the location
by llvm-lit. This works like llvm-lit does for finding the include dirs
for Clang using `-print-file-name=...`.
Summary:
These are stored in the libc/shared and have a unified CMake helper to
find them. Likely these will be a more core dependency as LLVM uses them
for constexpr math, libcxx uses it, and compiler-rt will probably use
bits of it.
The original intention was to allow building flang-rt with a partial
checkout, but i don't think this is a reasonable use-case and I do not
think this exists in practice.
Summary:
We need the offload project's RPC thread to handle the IO requests
originating from the GPU. Previously we did the 'easy' solution and just
linked this handler directly into the offload proejct. This is not ideal
because it prevents people's ability to build and configure libraries
separately.
This PR inverts the dependency, flang-rt now conditionally enables
support using the existing RPC callback mechanism. The cost is that
every flang-rt program now pays the cost of a boolean compare, the
benefit is the libraries are now independent of each-other.
Upstream clang's commit 5f2bedca, PR #187860, enabled a warning that
fires whenever code calls a `<fenv.h>` / `<cfenv>` primitive without
enabling proper FP exception behavior. This caused warnings in compiling
flang-rt, which in turn caused certain buildsbots to fail with
`-Werror`.
This change enables FP trap behavior for parts of flang-rt that are
known to require it. The other parts will be addressed by a future
change.
Assisted-by: AI
Summary:
The test previously did not account for CMake overrides, so we just grab
the file that's actually generated. `sort -u` should handle the case
where there's both a .so and .a.
Summary:
Right now we have a very heavy dependency on `flang` for a single module
file. The flang build will handle creating the actual module file, we
just need to export the symbols. This PR ports the source file from the
flang .f90 to equivalent C++. This allows us to build `flang-rt`
standalone without bringing in the entire flang infrastructure, which is
by far the heaviest build in the entire LLVM project.
I verified the symbols and values against a previous build, so they are
identical.
Summary:
This PR simply takes the existing `.cpp` files for the heaviest
implementations and separates them logically, typically between real,
integer, and complex types. The existing `.cpp` file is turned into a
`.h` file and we create new `.cpp` files that *only* contain the old
portion that used `RTDEF`. This allows for far more build system
parallelism, and it also means that static library linking semantics
mean that if the user only uses integer routines the linker will not
include the unused complex / real routines in the final executable.
All around this is a good practice for runtime libraries. Verified that
all `_Fortran` entrypoint routines are still present, port was strictly
mechanical.
The result of all of this is that I can now build `flang-rt` in ~10s
with all threads instead of ~50s due to the most expensive files being
split into parallelizable chunks.
A check recently added to clang will now flag uses of floating-point
exception routines. There are some calls to these functions in the flang
runtime. Since the flang runtime is compiled with the freshly built
clang, it now shows these warnings.
Since the CI build uses -Werror, this breaks the flang-rt build even
though no changes to the flang-rt have been made:
```
flang-rt/lib/runtime/main.cpp:20:3: error:
'feclearexcept' used without enabling floating-point exception behavior;
use 'pragma STDC FENV_ACCESS ON' or compile with
'-ffp-exception-behavior=maytrap' [-Werror,-Wfenv-access]
20 | std::feclearexcept(FE_ALL_EXCEPT);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
flang-rt/lib/runtime/main.cpp:25:3: error:
'fesetround' used without enabling floating-point exception behavior; use
'pragma STDC FENV_ACCESS ON' or compile with
'-ffp-exception-behavior=maytrap' [-Werror,-Wfenv-access]
25 | std::fesetround(FE_TONEAREST);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~
2 errors generated.
```
This PR will add a flag to turn off the warning as a means to restore
the flang-rt build.
The test binary was run without setting LD_LIBRARY_PATH, causing
libflang_rt.runtime.so to not be found at runtime. Match the pattern
used other tests in the same directory.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Summary:
The previous code had a combinatorial explosion of functions by
templating on both the source and target types. This created around 170
instantiations. Instead we just template on the source type and then use
a simple runtime check. This should not affect performance in a
significant way, it introduces maybe a few branches in what is already a
non-trivial operation that I do not think justifies a two-minute compile
time.
The result is that this file goes from 120 seconds to 12 on my machine
and the resulting file goes from 7.2 MiB to 757 kiB. Functinally this
makes us instantiate 1/10th the functions.
Summary:
These files take an extraordinarily long amount of time to compile. This
is due to the combinatorial explosino of template parameters. This PR
takes a few choice factors and puts them as runtime components instead.
This should have no real impact on runtime with optimizations, but
brings the compilation for these from 60s each to 30s on my machine. A
future PR will split these up by component so we get more build-system
parallelism.
The `flang-rt :: Driver/safe-trampoline-gnustack.f90` test `FAIL`s on
Solaris. It checks for a `GNU_STACK` section, however that isn't used
and not emitted on Solaris, so this patch just skips the test.
Tested on `x86_64-pc-solaris2.11` and `x86_64-pc-linux-gnu`.
Information:
The "buffer_" is used to store data in memory, and then the entire
"buffer_" is written at once using the Write(..) function. Consider a
write statement inside the nested implied do loop, the buffer_ size
keeps increasing when the length exceeds the buffer size. To increase
the size, we need to reallocate memory by copying the entire buffer into
a new buffer. This process consumes more time.
Implementation:
By reducing the number of reallocations, the performance could be
improved. Initially, we increase the size linearly, i.e., +65536 for
every new allocation. Then, when we cross 1 MB of buffer size, we
increase the size geometrically, i.e., 2x. Later, for more than 64MB, we
do 1.5x.
Issue: https://github.com/llvm/llvm-project/issues/163945
