[flang-commits] [flang] [flang][CUDA] Clarify host and device call diagnostics (PR #232850)
Andre Kuhlenschmidt via flang-commits
flang-commits at lists.llvm.org
Fri Oct 16 16:58:51 PDT 2026
https://github.com/akuhlens updated https://github.com/llvm/llvm-project/pull/232850
>From 49a31b2b18963c39264c22c73bf57c6a5c82adbd Mon Sep 17 00:00:00 2001
From: Andre Kuhlenschmidt <akuhlenschmi at nvidia.com>
Date: Fri, 16 Oct 2026 09:50:57 -0700
Subject: [PATCH 1/6] [flang][CUDA] Clarify procedure call context diagnostics
---
flang/lib/Semantics/check-cuda.cpp | 30 +++++---
.../CUDA/cuf-hostdevice-c-pointers.cuf | 4 +-
.../CUDA/cuf-hostdevice-call-context.cuf | 26 +++----
.../CUDA/cuf-hostdevice-else-if-guard.cuf | 4 +-
.../cuf-hostdevice-expression-context.cuf | 70 +++++++++----------
.../cuf-hostdevice-generic-diagnostics.cuf | 40 +++++------
.../CUDA/cuf-hostdevice-host-call.cuf | 2 +-
.../CUDA/cuf-hostdevice-inherited-spec.cuf | 2 +-
.../CUDA/cuf-hostdevice-module-calls.cuf | 4 +-
.../CUDA/cuf-hostdevice-return-guard.cuf | 40 +++++------
.../CUDA/cuf-hostdevice-stop-calls.cuf | 12 ++--
flang/test/Semantics/CUDA/cuf02.cuf | 2 +-
flang/test/Semantics/CUDA/cuf09.cuf | 12 ++--
13 files changed, 129 insertions(+), 119 deletions(-)
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index 581236cc72547..28e3c68b2c32f 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -193,24 +193,34 @@ struct DeviceExprChecker
CallContext callContext_;
private:
- static Result CallError(
- const evaluate::ProcedureDesignator &proc, bool inHostCode) {
+ Result CallError(
+ const evaluate::ProcedureDesignator &proc, bool inHostCode) const {
+ const char *procedureType{inHostCode ? "DEVICE" : "HOST"};
+ const char *codeType{inHostCode ? "HOST" : "DEVICE"};
if (auto generic{proc.genericName()}) {
const Symbol *specific{proc.GetInterfaceSymbol()};
const std::string specificName{specific
? specific->GetUltimate().name().ToString()
: proc.GetName()};
if (*generic != specificName) {
- return parser::MessageFormattedText(inHostCode
- ? "'%s' (specific procedure '%s') may not be called in host code"_err_en_US
- : "'%s' (specific procedure '%s') may not be called in device code"_err_en_US,
- *generic, specificName);
+ if (callContext_ == CallContext::HostDevice) {
+ return parser::MessageFormattedText(
+ "The %s procedure '%s' (specific procedure '%s') may not be called in HOST,DEVICE code without an ON_DEVICE() guard"_err_en_US,
+ procedureType, *generic, specificName);
+ }
+ return parser::MessageFormattedText(
+ "The %s procedure '%s' (specific procedure '%s') may not be called in %s code"_err_en_US,
+ procedureType, *generic, specificName, codeType);
}
}
- return parser::MessageFormattedText(inHostCode
- ? "'%s' may not be called in host code"_err_en_US
- : "'%s' may not be called in device code"_err_en_US,
- proc.GetName());
+ if (callContext_ == CallContext::HostDevice) {
+ return parser::MessageFormattedText(
+ "The %s procedure '%s' may not be called in HOST,DEVICE code without an ON_DEVICE() guard"_err_en_US,
+ procedureType, proc.GetName());
+ }
+ return parser::MessageFormattedText(
+ "The %s procedure '%s' may not be called in %s code"_err_en_US,
+ procedureType, proc.GetName(), codeType);
}
};
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf
index 5e5750051e5da..df955b23835ec 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf
@@ -30,7 +30,7 @@ contains
integer, value :: i
end subroutine
attributes(host,device) subroutine pass_expression()
- !ERROR: 'host_value' may not be called in device code
+ !ERROR: The HOST procedure 'host_value' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call consume_value(host_value())
end subroutine
attributes(host,device) subroutine both(p, q, f, g, d, e, result)
@@ -70,7 +70,7 @@ contains
end function
attributes(host,device) subroutine caller(result)
logical :: result
- !ERROR: 'c_associated_c_ptr' may not be called in device code
+ !ERROR: The HOST procedure 'c_associated_c_ptr' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
result = c_associated_c_ptr()
end subroutine
end module
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf
index e32476a67faa6..6f36702eef6cf 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf
@@ -34,13 +34,13 @@ contains
attributes(host,device) integer function unguarded_host(n)
integer, value :: n
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
unguarded_host = host_only(n)
end function
attributes(host,device) integer function unguarded_device(n)
integer, value :: n
- !ERROR: 'device_only' may not be called in host code
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
unguarded_device = device_only(n)
end function
@@ -113,10 +113,10 @@ contains
attributes(host,device) integer function ordinary_branch(n)
integer, value :: n
if (n > 0) then
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
ordinary_branch = host_only(n)
else
- !ERROR: 'device_only' may not be called in host code
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
ordinary_branch = device_only(n)
end if
end function
@@ -190,7 +190,7 @@ contains
after_guard = host_only(n)
end if
! Both copies reach the statement after END IF.
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
after_guard = after_guard + host_only(n)
end function
@@ -224,17 +224,17 @@ contains
attributes(host,device) subroutine unguarded_subroutines(n)
integer, intent(inout) :: n
- !ERROR: 'host_subroutine' may not be called in device code
+ !ERROR: The HOST procedure 'host_subroutine' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_subroutine(n)
- !ERROR: 'device_subroutine' may not be called in host code
+ !ERROR: The DEVICE procedure 'device_subroutine' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call device_subroutine(n)
end subroutine
attributes(host,device) subroutine generic_subroutines(n)
integer, intent(inout) :: n
- !ERROR: 'host_generic' (specific procedure 'host_subroutine') may not be called in device code
+ !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_subroutine') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_generic(n)
- !ERROR: 'device_generic' (specific procedure 'device_subroutine') may not be called in host code
+ !ERROR: The DEVICE procedure 'device_generic' (specific procedure 'device_subroutine') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call device_generic(n)
end subroutine
@@ -272,10 +272,10 @@ contains
attributes(host,device) subroutine misleading_bindc_guard(n)
integer, intent(inout) :: n
if (on_device()) then
- !ERROR: 'device_only' may not be called in host code
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call device_only(n)
else
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only(n)
end if
end subroutine
@@ -325,10 +325,10 @@ contains
attributes(host,device) subroutine misleading_guard(n)
integer, intent(inout) :: n
if (on_device()) then
- !ERROR: 'device_only' may not be called in host code
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call device_only(n)
else
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only(n)
end if
end subroutine
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf
index c1ab90032cd61..6352d1533dbed 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf
@@ -18,7 +18,7 @@ contains
call host_only()
end if
! The guard does not apply after the entire IF.
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine later_else(x, y)
@@ -54,7 +54,7 @@ contains
call host_only()
else
! A nested guard in another arm does not guard this ELSE.
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end if
end subroutine
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf
index 00ca8f6502c6a..98b6af55888f7 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf
@@ -15,7 +15,7 @@ contains
end function
attributes(host,device) subroutine select_h()
integer :: x
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
select case(h())
case(1)
x = 1
@@ -23,7 +23,7 @@ contains
end subroutine
attributes(host,device) subroutine select_d()
integer :: x
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
select case(d())
case(1)
x = 1
@@ -37,34 +37,34 @@ contains
end select
end subroutine
attributes(host,device) subroutine print_h()
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
print *, h()
end subroutine
attributes(host,device) subroutine print_d()
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
print *, d()
end subroutine
attributes(host,device) subroutine print_b()
print *, b()
end subroutine
attributes(host,device) subroutine write_item_h()
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
write(*,*) h()
end subroutine
attributes(host,device) subroutine write_item_d()
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
write(*,*) d()
end subroutine
attributes(host,device) subroutine write_item_b()
write(*,*) b()
end subroutine
attributes(host,device) subroutine write_unit_h()
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
!WARNING: I/O statement might not be supported on device [-Wcuda-usage]
write(h(),*) 1
end subroutine
attributes(host,device) subroutine write_unit_d()
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
!WARNING: I/O statement might not be supported on device [-Wcuda-usage]
write(d(),*) 1
end subroutine
@@ -74,12 +74,12 @@ contains
end subroutine
attributes(host,device) subroutine write_implied_do_h()
integer :: i
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
write(*,*) (i, i=1,h())
end subroutine
attributes(host,device) subroutine write_implied_do_d()
integer :: i
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
write(*,*) (i, i=1,d())
end subroutine
attributes(host,device) subroutine write_implied_do_b()
@@ -88,12 +88,12 @@ contains
end subroutine
attributes(host,device) subroutine write_subscript_h()
integer :: x(4)
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
write(*,*) x(h())
end subroutine
attributes(host,device) subroutine write_subscript_d()
integer :: x(4)
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
write(*,*) x(d())
end subroutine
attributes(host,device) subroutine write_subscript_b()
@@ -102,13 +102,13 @@ contains
end subroutine
attributes(host,device) subroutine read_subscript_h()
integer :: x(4)
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
!WARNING: I/O statement might not be supported on device [-Wcuda-usage]
read(*,*) x(h())
end subroutine
attributes(host,device) subroutine read_subscript_d()
integer :: x(4)
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
!WARNING: I/O statement might not be supported on device [-Wcuda-usage]
read(*,*) x(d())
end subroutine
@@ -119,12 +119,12 @@ contains
end subroutine
attributes(host,device) subroutine allocate_bound_h()
integer, allocatable :: a(:)
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
allocate(a(h()))
end subroutine
attributes(host,device) subroutine allocate_bound_d()
integer, allocatable :: a(:)
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
allocate(a(d()))
end subroutine
attributes(host,device) subroutine allocate_bound_b()
@@ -133,12 +133,12 @@ contains
end subroutine
attributes(host,device) subroutine allocate_source_h()
integer, allocatable :: a(:)
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
allocate(a,source=[h()])
end subroutine
attributes(host,device) subroutine allocate_source_d()
integer, allocatable :: a(:)
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
allocate(a,source=[d()])
end subroutine
attributes(host,device) subroutine allocate_source_b()
@@ -147,12 +147,12 @@ contains
end subroutine
attributes(host,device) subroutine allocate_mold_h()
integer, allocatable :: a(:)
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
allocate(a,mold=[h()])
end subroutine
attributes(host,device) subroutine allocate_mold_d()
integer, allocatable :: a(:)
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
allocate(a,mold=[d()])
end subroutine
attributes(host,device) subroutine allocate_mold_b()
@@ -161,12 +161,12 @@ contains
end subroutine
attributes(host,device) subroutine allocate_length_h()
character(:), allocatable :: text
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
allocate(character(h()) :: text)
end subroutine
attributes(host,device) subroutine allocate_length_d()
character(:), allocatable :: text
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
allocate(character(d()) :: text)
end subroutine
attributes(host,device) subroutine allocate_length_b()
@@ -176,13 +176,13 @@ contains
attributes(host,device) subroutine deallocate_stat_h()
integer, allocatable :: a(:)
integer :: status(4)
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
deallocate(a,stat=status(h()))
end subroutine
attributes(host,device) subroutine deallocate_stat_d()
integer, allocatable :: a(:)
integer :: status(4)
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
deallocate(a,stat=status(d()))
end subroutine
attributes(host,device) subroutine deallocate_stat_b()
@@ -191,12 +191,12 @@ contains
deallocate(a,stat=status(b()))
end subroutine
attributes(host,device) subroutine automatic_bound_h()
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
integer :: a(h())
a = 1
end subroutine
attributes(host,device) subroutine automatic_bound_d()
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
integer :: a(d())
a = 1
end subroutine
@@ -205,12 +205,12 @@ contains
a = 1
end subroutine
attributes(host,device) subroutine character_length_h()
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
character(h()) :: text
text = 'x'
end subroutine
attributes(host,device) subroutine character_length_d()
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
character(d()) :: text
text = 'x'
end subroutine
@@ -285,23 +285,23 @@ contains
end subroutine
attributes(host,device) subroutine block_bound_h()
block
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
integer :: a(h())
a = 1
end block
end subroutine
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
attributes(host,device) character(h()) function header_h()
header_h = 'x'
end function
attributes(host,device) subroutine block_bound_d()
block
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
integer :: a(d())
a = 1
end block
end subroutine
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
attributes(host,device) character(d()) function header_d()
header_d = 'x'
end function
@@ -325,13 +325,13 @@ contains
end function
attributes(device) subroutine device_bound(n)
integer, value :: n
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in DEVICE code
real :: a(h(n))
a = 0.0
end subroutine
attributes(global) subroutine global_print(n)
integer, value :: n
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in DEVICE code
print *, h(n)
end subroutine
subroutine kernel_print(a, n)
@@ -339,7 +339,7 @@ contains
integer, device :: a(n)
!$cuf kernel do <<<*,*>>>
do i = 1, n
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in DEVICE code
print *, h(i)
a(i) = i
end do
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf
index 599775338dd4d..3ce3c69986eb3 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf
@@ -57,26 +57,26 @@ contains
attributes(host,device) subroutine unguarded(n, x)
integer :: n
real :: x
- !ERROR: 'host_generic' (specific procedure 'host_integer') may not be called in device code
+ !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = host_generic(n)
- !ERROR: 'host_generic' (specific procedure 'host_real') may not be called in device code
+ !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_real') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
x = host_generic(x)
- !ERROR: 'device_generic' (specific procedure 'device_integer') may not be called in host code
+ !ERROR: The DEVICE procedure 'device_generic' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = device_generic(n)
- !ERROR: 'host_sub_generic' (specific procedure 'host_sub') may not be called in device code
+ !ERROR: The HOST procedure 'host_sub_generic' (specific procedure 'host_sub') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_sub_generic(n)
- !ERROR: 'device_sub_generic' (specific procedure 'device_sub') may not be called in host code
+ !ERROR: The DEVICE procedure 'device_sub_generic' (specific procedure 'device_sub') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call device_sub_generic(n)
! Ordinary calls keep their single-name diagnostics.
- !ERROR: 'host_integer' may not be called in device code
+ !ERROR: The HOST procedure 'host_integer' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = host_integer(n)
- !ERROR: 'device_integer' may not be called in host code
+ !ERROR: The DEVICE procedure 'device_integer' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = device_integer(n)
n = both_generic(n)
end subroutine
attributes(device) subroutine device_caller(n)
integer :: n
- !ERROR: 'host_generic' (specific procedure 'host_integer') may not be called in device code
+ !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_integer') may not be called in DEVICE code
n = host_generic(n)
n = device_generic(n)
end subroutine
@@ -99,18 +99,18 @@ module generic_alias_callers
contains
attributes(host,device) subroutine aliases(n)
integer :: n
- !ERROR: 'renamed_host' (specific procedure 'host_integer') may not be called in device code
+ !ERROR: The HOST procedure 'renamed_host' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = renamed_host(n)
- !ERROR: 'another_host' (specific procedure 'host_integer') may not be called in device code
+ !ERROR: The HOST procedure 'another_host' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = another_host(n)
- !ERROR: 'renamed_device' (specific procedure 'device_integer') may not be called in host code
+ !ERROR: The DEVICE procedure 'renamed_device' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = renamed_device(n)
- !ERROR: 'renamed_sub' (specific procedure 'host_sub') may not be called in device code
+ !ERROR: The HOST procedure 'renamed_sub' (specific procedure 'host_sub') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call renamed_sub(n)
- !ERROR: 'renamed_host' (specific procedure 'host_integer') may not be called in device code
+ !ERROR: The HOST procedure 'renamed_host' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = both_generic(renamed_host(n))
! A USE alias of an ordinary specific is not a generic reference.
- !ERROR: 'specific_alias' may not be called in device code
+ !ERROR: The HOST procedure 'specific_alias' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = specific_alias(n)
end subroutine
end module
@@ -122,13 +122,13 @@ contains
type(bound_targets) :: t
class(bound_targets) :: poly
integer :: n
- !ERROR: 'host_bound' (specific procedure 'host_integer') may not be called in device code
+ !ERROR: The HOST procedure 'host_bound' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = t%host_bound(n)
- !ERROR: 'device_bound' (specific procedure 'device_integer') may not be called in host code
+ !ERROR: The DEVICE procedure 'device_bound' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = t%device_bound(n)
- !ERROR: 'host_bound' (specific procedure 'host_integer') may not be called in device code
+ !ERROR: The HOST procedure 'host_bound' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = poly%host_bound(n)
- !ERROR: 'device_bound' (specific procedure 'device_integer') may not be called in host code
+ !ERROR: The DEVICE procedure 'device_bound' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = poly%device_bound(n)
end subroutine
end module
@@ -139,7 +139,7 @@ contains
attributes(host,device) subroutine atomic_call(a, n, old)
integer, device :: a(*)
integer :: n, old
- !ERROR: 'atomicadd' (specific procedure 'atomicaddi') may not be called in host code
+ !ERROR: The DEVICE procedure 'atomicadd' (specific procedure 'atomicaddi') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
old = atomicadd(a(1), n)
end subroutine
end module
@@ -156,7 +156,7 @@ contains
attributes(host,device) subroutine caller(n)
integer :: n
! Do not repeat the name when generic and specific have the same spelling.
- !ERROR: 'same_name' may not be called in device code
+ !ERROR: The HOST procedure 'same_name' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
n = same_name(n)
end subroutine
end module
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-host-call.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-host-call.cuf
index a86bcb5c7ac36..ca18be94d7437 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-host-call.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-host-call.cuf
@@ -49,7 +49,7 @@ contains
end function
end interface
double precision, intent(inout) :: x
- !ERROR: 'hostfunc' may not be called in device code
+ !ERROR: The HOST procedure 'hostfunc' may not be called in DEVICE code
x = hostfunc()
end subroutine
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf
index 2228ead80c739..5bd310cc2babb 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf
@@ -132,7 +132,7 @@ contains
extent = 1
end function
attributes(host,device) subroutine caller(a)
- !ERROR: 'extent' may not be called in device code
+ !ERROR: The HOST procedure 'extent' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
integer :: a(extent())
a = 1
end subroutine
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf
index bbfb1dd655aeb..e3e503d966bbe 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf
@@ -39,9 +39,9 @@ contains
module procedure caller
caller = both(n)
call both_subroutine(n)
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
caller = host_only(n)
- !ERROR: 'device_only' may not be called in host code
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
caller = device_only(n)
if (on_device()) then
caller = device_only(n)
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf
index d59f25787249a..50065cad5d680 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf
@@ -69,7 +69,7 @@ contains
end if
end if
! The nested IF does not unconditionally return.
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine nested_else_return(choose)
@@ -82,7 +82,7 @@ contains
end if
end if
! The nested ELSE does not make the enclosing arm return unconditionally.
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine guarded_loop_return(choose)
@@ -93,7 +93,7 @@ contains
end do
end if
! Loops may not execute and are not traversed by the return scan.
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine conditional_block_guard(choose)
@@ -105,7 +105,7 @@ contains
call host_only()
end if
! The enclosing conditional can bypass the whole BLOCK.
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine return_before_last_statement()
@@ -134,7 +134,7 @@ contains
call device_only()
end if
! A nested conditional return does not eliminate fallthrough.
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine else_return()
@@ -152,7 +152,7 @@ contains
call host_only()
end if
! The outer IF can bypass the return guard.
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine loop_return(choose)
@@ -162,7 +162,7 @@ contains
call host_only()
end do
! The loop need not execute.
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine reset_between_branches(choose, other)
@@ -171,13 +171,13 @@ contains
if (on_device()) return
call host_only()
else if (other) then
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
else
- !ERROR: 'device_only' may not be called in host code
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call device_only()
end if
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine preserve_guarded_else_if(choose, other)
@@ -194,7 +194,7 @@ contains
! the host, but guard direction is not tracked yet.
call device_only()
end if
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine preserve_inherited_guard(choose)
@@ -215,7 +215,7 @@ contains
else
call both()
end if
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine ordinary_else_return(choose)
@@ -225,7 +225,7 @@ contains
else
return
end if
- !ERROR: 'device_only' may not be called in host code
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call device_only()
end subroutine
attributes(host,device) subroutine unrelated_arm_returns(choose)
@@ -236,7 +236,7 @@ contains
return
end if
! The returning arm does not test ON_DEVICE.
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine return_before_guard_arm(choose)
@@ -246,7 +246,7 @@ contains
else if (on_device()) then
call device_only()
end if
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine else_if_guard_else_returns(choose)
@@ -298,7 +298,7 @@ contains
attributes(host,device) subroutine ordinary_return(choose)
logical, intent(in) :: choose
if (choose) return
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine conditional_return(choose)
@@ -307,7 +307,7 @@ contains
if (choose) return
end if
! The device branch can fall through.
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
attributes(host,device) subroutine unrelated_label()
@@ -332,7 +332,7 @@ contains
attributes(host,device) subroutine else_if_bypass(choose)
logical, intent(in) :: choose
if (choose) then
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
else if (on_device()) then
return
@@ -343,7 +343,7 @@ contains
attributes(host,device) subroutine select_dispatch(result)
integer, intent(out) :: result
! The selector example from the review is already diagnosed.
- !ERROR: 'device_value' may not be called in host code
+ !ERROR: The DEVICE procedure 'device_value' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
select case (device_value())
case (1)
result = 10
@@ -365,7 +365,7 @@ contains
end subroutine
attributes(host,device) subroutine caller()
if (on_device()) return
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
call host_only()
end subroutine
end module
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf
index a08448af736d5..4b203890d7a61 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf
@@ -21,23 +21,23 @@ contains
attributes(host,device) subroutine caller()
! Review example: the STOP code expression runs on both targets.
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
stop d()
end subroutine
attributes(host,device) subroutine host_code()
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
stop h()
end subroutine
attributes(host,device) subroutine error_code()
- !ERROR: 'd' may not be called in host code
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
error stop d()
end subroutine
attributes(host,device) subroutine quiet_device()
- !ERROR: 'device_quiet' may not be called in host code
+ !ERROR: The DEVICE procedure 'device_quiet' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
stop 1, quiet=device_quiet()
end subroutine
attributes(host,device) subroutine quiet_host()
- !ERROR: 'host_quiet' may not be called in device code
+ !ERROR: The HOST procedure 'host_quiet' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
error stop 'failed', quiet=host_quiet()
end subroutine
attributes(host,device) subroutine both_code()
@@ -51,7 +51,7 @@ contains
end if
end subroutine
attributes(device) subroutine device_caller()
- !ERROR: 'h' may not be called in device code
+ !ERROR: The HOST procedure 'h' may not be called in DEVICE code
stop h()
end subroutine
end module
diff --git a/flang/test/Semantics/CUDA/cuf02.cuf b/flang/test/Semantics/CUDA/cuf02.cuf
index 9a01a693f16e7..3543b41feb687 100644
--- a/flang/test/Semantics/CUDA/cuf02.cuf
+++ b/flang/test/Semantics/CUDA/cuf02.cuf
@@ -40,7 +40,7 @@ module m
end
attributes(device) subroutine s6_bad
badstmt(x) = hostfunc(x)
- !ERROR: 'hostfunc' may not be called in device code
+ !ERROR: The HOST procedure 'hostfunc' may not be called in DEVICE code
x = badstmt(x)
end
!ERROR: A function may not have ATTRIBUTES(GLOBAL) or ATTRIBUTES(GRID_GLOBAL)
diff --git a/flang/test/Semantics/CUDA/cuf09.cuf b/flang/test/Semantics/CUDA/cuf09.cuf
index 09f1abce098eb..65eb06fd25b9b 100644
--- a/flang/test/Semantics/CUDA/cuf09.cuf
+++ b/flang/test/Semantics/CUDA/cuf09.cuf
@@ -219,15 +219,15 @@ program main
end do
!$cuf kernel do <<< 1, 2 >>>
do j = 1, 10
- !ERROR: 'foo' may not be called in device code
+ !ERROR: The HOST procedure 'foo' may not be called in DEVICE code
call foo
- !ERROR: 'bar' may not be called in device code
+ !ERROR: The HOST procedure 'bar' may not be called in DEVICE code
x = bar()
- !ERROR: 'ifunc' may not be called in device code
+ !ERROR: The HOST procedure 'ifunc' may not be called in DEVICE code
if (ifunc() /= 0) continue
- !ERROR: 'ifunc' may not be called in device code
+ !ERROR: The HOST procedure 'ifunc' may not be called in DEVICE code
if (ifunc() /= 0) then
- !ERROR: 'ifunc' may not be called in device code
+ !ERROR: The HOST procedure 'ifunc' may not be called in DEVICE code
else if (ifunc() /= 1) then
end if
end do
@@ -256,7 +256,7 @@ subroutine host1()
!$cuf kernel do(1) <<<*,32>>>
do i = 1, 32
a(i) = a(i) * 2.0
- !ERROR: 'syncthreads' may not be called in device code
+ !ERROR: The HOST procedure 'syncthreads' may not be called in DEVICE code
call syncthreads() ! missing explicit use cudadevice
a(i) = a(i) + a(j) - 34.0
end do
>From 7bcb9e20316dd78976fea6dc4577511842756165 Mon Sep 17 00:00:00 2001
From: Andre Kuhlenschmidt <akuhlenschmi at nvidia.com>
Date: Fri, 16 Oct 2026 11:14:21 -0700
Subject: [PATCH 2/6] [flang][CUDA] Unify call diagnostic formatting
---
flang/lib/Semantics/check-cuda.cpp | 23 ++++++++---------------
1 file changed, 8 insertions(+), 15 deletions(-)
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index 28e3c68b2c32f..ed5d4c5798600 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -196,31 +196,24 @@ struct DeviceExprChecker
Result CallError(
const evaluate::ProcedureDesignator &proc, bool inHostCode) const {
const char *procedureType{inHostCode ? "DEVICE" : "HOST"};
- const char *codeType{inHostCode ? "HOST" : "DEVICE"};
+ const char *codeContext{callContext_ == CallContext::HostDevice
+ ? "HOST,DEVICE code without an ON_DEVICE() guard"
+ : inHostCode ? "HOST code"
+ : "DEVICE code"};
if (auto generic{proc.genericName()}) {
const Symbol *specific{proc.GetInterfaceSymbol()};
const std::string specificName{specific
? specific->GetUltimate().name().ToString()
: proc.GetName()};
if (*generic != specificName) {
- if (callContext_ == CallContext::HostDevice) {
- return parser::MessageFormattedText(
- "The %s procedure '%s' (specific procedure '%s') may not be called in HOST,DEVICE code without an ON_DEVICE() guard"_err_en_US,
- procedureType, *generic, specificName);
- }
return parser::MessageFormattedText(
- "The %s procedure '%s' (specific procedure '%s') may not be called in %s code"_err_en_US,
- procedureType, *generic, specificName, codeType);
+ "The %s procedure '%s' (specific procedure '%s') may not be called in %s"_err_en_US,
+ procedureType, *generic, specificName, codeContext);
}
}
- if (callContext_ == CallContext::HostDevice) {
- return parser::MessageFormattedText(
- "The %s procedure '%s' may not be called in HOST,DEVICE code without an ON_DEVICE() guard"_err_en_US,
- procedureType, proc.GetName());
- }
return parser::MessageFormattedText(
- "The %s procedure '%s' may not be called in %s code"_err_en_US,
- procedureType, proc.GetName(), codeType);
+ "The %s procedure '%s' may not be called in %s"_err_en_US,
+ procedureType, proc.GetName(), codeContext);
}
};
>From b8d7445e2f998325ab83428733a730f3f57841d3 Mon Sep 17 00:00:00 2001
From: Andre Kuhlenschmidt <akuhlenschmi at nvidia.com>
Date: Fri, 16 Oct 2026 12:01:35 -0700
Subject: [PATCH 3/6] [flang][CUDA] Identify CUDADEVICE in guard diagnostics
---
flang/lib/Semantics/check-cuda.cpp | 3 +-
.../CUDA/cuf-hostdevice-c-pointers.cuf | 4 +-
.../CUDA/cuf-hostdevice-call-context.cuf | 26 ++++----
.../CUDA/cuf-hostdevice-else-if-guard.cuf | 4 +-
.../cuf-hostdevice-expression-context.cuf | 64 +++++++++----------
.../cuf-hostdevice-generic-diagnostics.cuf | 38 +++++------
.../CUDA/cuf-hostdevice-inherited-spec.cuf | 2 +-
.../CUDA/cuf-hostdevice-module-calls.cuf | 4 +-
.../CUDA/cuf-hostdevice-return-guard.cuf | 40 ++++++------
.../CUDA/cuf-hostdevice-stop-calls.cuf | 10 +--
10 files changed, 98 insertions(+), 97 deletions(-)
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index ed5d4c5798600..be0c4f21becaa 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -197,7 +197,8 @@ struct DeviceExprChecker
const evaluate::ProcedureDesignator &proc, bool inHostCode) const {
const char *procedureType{inHostCode ? "DEVICE" : "HOST"};
const char *codeContext{callContext_ == CallContext::HostDevice
- ? "HOST,DEVICE code without an ON_DEVICE() guard"
+ ? "HOST,DEVICE code without a guard using ON_DEVICE() from the "
+ "CUDADEVICE module"
: inHostCode ? "HOST code"
: "DEVICE code"};
if (auto generic{proc.genericName()}) {
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf
index df955b23835ec..9383ed014680b 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf
@@ -30,7 +30,7 @@ contains
integer, value :: i
end subroutine
attributes(host,device) subroutine pass_expression()
- !ERROR: The HOST procedure 'host_value' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_value' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call consume_value(host_value())
end subroutine
attributes(host,device) subroutine both(p, q, f, g, d, e, result)
@@ -70,7 +70,7 @@ contains
end function
attributes(host,device) subroutine caller(result)
logical :: result
- !ERROR: The HOST procedure 'c_associated_c_ptr' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'c_associated_c_ptr' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
result = c_associated_c_ptr()
end subroutine
end module
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf
index 6f36702eef6cf..73362eafd89e3 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf
@@ -34,13 +34,13 @@ contains
attributes(host,device) integer function unguarded_host(n)
integer, value :: n
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
unguarded_host = host_only(n)
end function
attributes(host,device) integer function unguarded_device(n)
integer, value :: n
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
unguarded_device = device_only(n)
end function
@@ -113,10 +113,10 @@ contains
attributes(host,device) integer function ordinary_branch(n)
integer, value :: n
if (n > 0) then
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
ordinary_branch = host_only(n)
else
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
ordinary_branch = device_only(n)
end if
end function
@@ -190,7 +190,7 @@ contains
after_guard = host_only(n)
end if
! Both copies reach the statement after END IF.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
after_guard = after_guard + host_only(n)
end function
@@ -224,17 +224,17 @@ contains
attributes(host,device) subroutine unguarded_subroutines(n)
integer, intent(inout) :: n
- !ERROR: The HOST procedure 'host_subroutine' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_subroutine' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_subroutine(n)
- !ERROR: The DEVICE procedure 'device_subroutine' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_subroutine' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call device_subroutine(n)
end subroutine
attributes(host,device) subroutine generic_subroutines(n)
integer, intent(inout) :: n
- !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_subroutine') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_subroutine') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_generic(n)
- !ERROR: The DEVICE procedure 'device_generic' (specific procedure 'device_subroutine') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_generic' (specific procedure 'device_subroutine') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call device_generic(n)
end subroutine
@@ -272,10 +272,10 @@ contains
attributes(host,device) subroutine misleading_bindc_guard(n)
integer, intent(inout) :: n
if (on_device()) then
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call device_only(n)
else
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only(n)
end if
end subroutine
@@ -325,10 +325,10 @@ contains
attributes(host,device) subroutine misleading_guard(n)
integer, intent(inout) :: n
if (on_device()) then
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call device_only(n)
else
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only(n)
end if
end subroutine
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf
index 6352d1533dbed..1465960a7f717 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf
@@ -18,7 +18,7 @@ contains
call host_only()
end if
! The guard does not apply after the entire IF.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine later_else(x, y)
@@ -54,7 +54,7 @@ contains
call host_only()
else
! A nested guard in another arm does not guard this ELSE.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end if
end subroutine
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf
index 98b6af55888f7..17d597f0c8381 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf
@@ -15,7 +15,7 @@ contains
end function
attributes(host,device) subroutine select_h()
integer :: x
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
select case(h())
case(1)
x = 1
@@ -23,7 +23,7 @@ contains
end subroutine
attributes(host,device) subroutine select_d()
integer :: x
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
select case(d())
case(1)
x = 1
@@ -37,34 +37,34 @@ contains
end select
end subroutine
attributes(host,device) subroutine print_h()
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
print *, h()
end subroutine
attributes(host,device) subroutine print_d()
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
print *, d()
end subroutine
attributes(host,device) subroutine print_b()
print *, b()
end subroutine
attributes(host,device) subroutine write_item_h()
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
write(*,*) h()
end subroutine
attributes(host,device) subroutine write_item_d()
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
write(*,*) d()
end subroutine
attributes(host,device) subroutine write_item_b()
write(*,*) b()
end subroutine
attributes(host,device) subroutine write_unit_h()
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
!WARNING: I/O statement might not be supported on device [-Wcuda-usage]
write(h(),*) 1
end subroutine
attributes(host,device) subroutine write_unit_d()
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
!WARNING: I/O statement might not be supported on device [-Wcuda-usage]
write(d(),*) 1
end subroutine
@@ -74,12 +74,12 @@ contains
end subroutine
attributes(host,device) subroutine write_implied_do_h()
integer :: i
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
write(*,*) (i, i=1,h())
end subroutine
attributes(host,device) subroutine write_implied_do_d()
integer :: i
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
write(*,*) (i, i=1,d())
end subroutine
attributes(host,device) subroutine write_implied_do_b()
@@ -88,12 +88,12 @@ contains
end subroutine
attributes(host,device) subroutine write_subscript_h()
integer :: x(4)
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
write(*,*) x(h())
end subroutine
attributes(host,device) subroutine write_subscript_d()
integer :: x(4)
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
write(*,*) x(d())
end subroutine
attributes(host,device) subroutine write_subscript_b()
@@ -102,13 +102,13 @@ contains
end subroutine
attributes(host,device) subroutine read_subscript_h()
integer :: x(4)
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
!WARNING: I/O statement might not be supported on device [-Wcuda-usage]
read(*,*) x(h())
end subroutine
attributes(host,device) subroutine read_subscript_d()
integer :: x(4)
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
!WARNING: I/O statement might not be supported on device [-Wcuda-usage]
read(*,*) x(d())
end subroutine
@@ -119,12 +119,12 @@ contains
end subroutine
attributes(host,device) subroutine allocate_bound_h()
integer, allocatable :: a(:)
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
allocate(a(h()))
end subroutine
attributes(host,device) subroutine allocate_bound_d()
integer, allocatable :: a(:)
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
allocate(a(d()))
end subroutine
attributes(host,device) subroutine allocate_bound_b()
@@ -133,12 +133,12 @@ contains
end subroutine
attributes(host,device) subroutine allocate_source_h()
integer, allocatable :: a(:)
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
allocate(a,source=[h()])
end subroutine
attributes(host,device) subroutine allocate_source_d()
integer, allocatable :: a(:)
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
allocate(a,source=[d()])
end subroutine
attributes(host,device) subroutine allocate_source_b()
@@ -147,12 +147,12 @@ contains
end subroutine
attributes(host,device) subroutine allocate_mold_h()
integer, allocatable :: a(:)
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
allocate(a,mold=[h()])
end subroutine
attributes(host,device) subroutine allocate_mold_d()
integer, allocatable :: a(:)
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
allocate(a,mold=[d()])
end subroutine
attributes(host,device) subroutine allocate_mold_b()
@@ -161,12 +161,12 @@ contains
end subroutine
attributes(host,device) subroutine allocate_length_h()
character(:), allocatable :: text
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
allocate(character(h()) :: text)
end subroutine
attributes(host,device) subroutine allocate_length_d()
character(:), allocatable :: text
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
allocate(character(d()) :: text)
end subroutine
attributes(host,device) subroutine allocate_length_b()
@@ -176,13 +176,13 @@ contains
attributes(host,device) subroutine deallocate_stat_h()
integer, allocatable :: a(:)
integer :: status(4)
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
deallocate(a,stat=status(h()))
end subroutine
attributes(host,device) subroutine deallocate_stat_d()
integer, allocatable :: a(:)
integer :: status(4)
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
deallocate(a,stat=status(d()))
end subroutine
attributes(host,device) subroutine deallocate_stat_b()
@@ -191,12 +191,12 @@ contains
deallocate(a,stat=status(b()))
end subroutine
attributes(host,device) subroutine automatic_bound_h()
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
integer :: a(h())
a = 1
end subroutine
attributes(host,device) subroutine automatic_bound_d()
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
integer :: a(d())
a = 1
end subroutine
@@ -205,12 +205,12 @@ contains
a = 1
end subroutine
attributes(host,device) subroutine character_length_h()
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
character(h()) :: text
text = 'x'
end subroutine
attributes(host,device) subroutine character_length_d()
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
character(d()) :: text
text = 'x'
end subroutine
@@ -285,23 +285,23 @@ contains
end subroutine
attributes(host,device) subroutine block_bound_h()
block
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
integer :: a(h())
a = 1
end block
end subroutine
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
attributes(host,device) character(h()) function header_h()
header_h = 'x'
end function
attributes(host,device) subroutine block_bound_d()
block
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
integer :: a(d())
a = 1
end block
end subroutine
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
attributes(host,device) character(d()) function header_d()
header_d = 'x'
end function
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf
index 3ce3c69986eb3..bba05a010f6f7 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf
@@ -57,20 +57,20 @@ contains
attributes(host,device) subroutine unguarded(n, x)
integer :: n
real :: x
- !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = host_generic(n)
- !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_real') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_real') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
x = host_generic(x)
- !ERROR: The DEVICE procedure 'device_generic' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_generic' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = device_generic(n)
- !ERROR: The HOST procedure 'host_sub_generic' (specific procedure 'host_sub') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_sub_generic' (specific procedure 'host_sub') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_sub_generic(n)
- !ERROR: The DEVICE procedure 'device_sub_generic' (specific procedure 'device_sub') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_sub_generic' (specific procedure 'device_sub') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call device_sub_generic(n)
! Ordinary calls keep their single-name diagnostics.
- !ERROR: The HOST procedure 'host_integer' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_integer' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = host_integer(n)
- !ERROR: The DEVICE procedure 'device_integer' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_integer' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = device_integer(n)
n = both_generic(n)
end subroutine
@@ -99,18 +99,18 @@ module generic_alias_callers
contains
attributes(host,device) subroutine aliases(n)
integer :: n
- !ERROR: The HOST procedure 'renamed_host' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'renamed_host' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = renamed_host(n)
- !ERROR: The HOST procedure 'another_host' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'another_host' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = another_host(n)
- !ERROR: The DEVICE procedure 'renamed_device' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'renamed_device' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = renamed_device(n)
- !ERROR: The HOST procedure 'renamed_sub' (specific procedure 'host_sub') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'renamed_sub' (specific procedure 'host_sub') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call renamed_sub(n)
- !ERROR: The HOST procedure 'renamed_host' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'renamed_host' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = both_generic(renamed_host(n))
! A USE alias of an ordinary specific is not a generic reference.
- !ERROR: The HOST procedure 'specific_alias' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'specific_alias' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = specific_alias(n)
end subroutine
end module
@@ -122,13 +122,13 @@ contains
type(bound_targets) :: t
class(bound_targets) :: poly
integer :: n
- !ERROR: The HOST procedure 'host_bound' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_bound' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = t%host_bound(n)
- !ERROR: The DEVICE procedure 'device_bound' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_bound' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = t%device_bound(n)
- !ERROR: The HOST procedure 'host_bound' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_bound' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = poly%host_bound(n)
- !ERROR: The DEVICE procedure 'device_bound' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_bound' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = poly%device_bound(n)
end subroutine
end module
@@ -139,7 +139,7 @@ contains
attributes(host,device) subroutine atomic_call(a, n, old)
integer, device :: a(*)
integer :: n, old
- !ERROR: The DEVICE procedure 'atomicadd' (specific procedure 'atomicaddi') may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'atomicadd' (specific procedure 'atomicaddi') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
old = atomicadd(a(1), n)
end subroutine
end module
@@ -156,7 +156,7 @@ contains
attributes(host,device) subroutine caller(n)
integer :: n
! Do not repeat the name when generic and specific have the same spelling.
- !ERROR: The HOST procedure 'same_name' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'same_name' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
n = same_name(n)
end subroutine
end module
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf
index 5bd310cc2babb..09d4cf4037fd1 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf
@@ -132,7 +132,7 @@ contains
extent = 1
end function
attributes(host,device) subroutine caller(a)
- !ERROR: The HOST procedure 'extent' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'extent' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
integer :: a(extent())
a = 1
end subroutine
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf
index e3e503d966bbe..4c06d5ef80c1d 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf
@@ -39,9 +39,9 @@ contains
module procedure caller
caller = both(n)
call both_subroutine(n)
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
caller = host_only(n)
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
caller = device_only(n)
if (on_device()) then
caller = device_only(n)
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf
index 50065cad5d680..64d3009f4f261 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf
@@ -69,7 +69,7 @@ contains
end if
end if
! The nested IF does not unconditionally return.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine nested_else_return(choose)
@@ -82,7 +82,7 @@ contains
end if
end if
! The nested ELSE does not make the enclosing arm return unconditionally.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine guarded_loop_return(choose)
@@ -93,7 +93,7 @@ contains
end do
end if
! Loops may not execute and are not traversed by the return scan.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine conditional_block_guard(choose)
@@ -105,7 +105,7 @@ contains
call host_only()
end if
! The enclosing conditional can bypass the whole BLOCK.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine return_before_last_statement()
@@ -134,7 +134,7 @@ contains
call device_only()
end if
! A nested conditional return does not eliminate fallthrough.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine else_return()
@@ -152,7 +152,7 @@ contains
call host_only()
end if
! The outer IF can bypass the return guard.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine loop_return(choose)
@@ -162,7 +162,7 @@ contains
call host_only()
end do
! The loop need not execute.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine reset_between_branches(choose, other)
@@ -171,13 +171,13 @@ contains
if (on_device()) return
call host_only()
else if (other) then
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
else
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call device_only()
end if
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine preserve_guarded_else_if(choose, other)
@@ -194,7 +194,7 @@ contains
! the host, but guard direction is not tracked yet.
call device_only()
end if
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine preserve_inherited_guard(choose)
@@ -215,7 +215,7 @@ contains
else
call both()
end if
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine ordinary_else_return(choose)
@@ -225,7 +225,7 @@ contains
else
return
end if
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call device_only()
end subroutine
attributes(host,device) subroutine unrelated_arm_returns(choose)
@@ -236,7 +236,7 @@ contains
return
end if
! The returning arm does not test ON_DEVICE.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine return_before_guard_arm(choose)
@@ -246,7 +246,7 @@ contains
else if (on_device()) then
call device_only()
end if
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine else_if_guard_else_returns(choose)
@@ -298,7 +298,7 @@ contains
attributes(host,device) subroutine ordinary_return(choose)
logical, intent(in) :: choose
if (choose) return
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine conditional_return(choose)
@@ -307,7 +307,7 @@ contains
if (choose) return
end if
! The device branch can fall through.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
attributes(host,device) subroutine unrelated_label()
@@ -332,7 +332,7 @@ contains
attributes(host,device) subroutine else_if_bypass(choose)
logical, intent(in) :: choose
if (choose) then
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
else if (on_device()) then
return
@@ -343,7 +343,7 @@ contains
attributes(host,device) subroutine select_dispatch(result)
integer, intent(out) :: result
! The selector example from the review is already diagnosed.
- !ERROR: The DEVICE procedure 'device_value' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_value' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
select case (device_value())
case (1)
result = 10
@@ -365,7 +365,7 @@ contains
end subroutine
attributes(host,device) subroutine caller()
if (on_device()) return
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
call host_only()
end subroutine
end module
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf
index 4b203890d7a61..def7ec0d9a827 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf
@@ -21,23 +21,23 @@ contains
attributes(host,device) subroutine caller()
! Review example: the STOP code expression runs on both targets.
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
stop d()
end subroutine
attributes(host,device) subroutine host_code()
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
stop h()
end subroutine
attributes(host,device) subroutine error_code()
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
error stop d()
end subroutine
attributes(host,device) subroutine quiet_device()
- !ERROR: The DEVICE procedure 'device_quiet' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The DEVICE procedure 'device_quiet' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
stop 1, quiet=device_quiet()
end subroutine
attributes(host,device) subroutine quiet_host()
- !ERROR: The HOST procedure 'host_quiet' may not be called in HOST,DEVICE code without an ON_DEVICE() guard
+ !ERROR: The HOST procedure 'host_quiet' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
error stop 'failed', quiet=host_quiet()
end subroutine
attributes(host,device) subroutine both_code()
>From 056c980d87b2c6dd33dff4d22b0ad5aaece75429 Mon Sep 17 00:00:00 2001
From: Andre Kuhlenschmidt <akuhlenschmi at nvidia.com>
Date: Fri, 16 Oct 2026 12:39:21 -0700
Subject: [PATCH 4/6] [flang][CUDA] Use callee attributes in call diagnostics
---
flang/lib/Semantics/check-cuda.cpp | 19 ++++---
.../Semantics/CUDA/cuf-call-callee-kind.cuf | 53 +++++++++++++++++++
2 files changed, 65 insertions(+), 7 deletions(-)
create mode 100644 flang/test/Semantics/CUDA/cuf-call-callee-kind.cuf
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index be0c4f21becaa..d761a57e99f0a 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -12,6 +12,7 @@
#include "flang/Evaluate/fold.h"
#include "flang/Evaluate/tools.h"
#include "flang/Evaluate/traverse.h"
+#include "flang/Parser/characters.h"
#include "flang/Parser/parse-tree-visitor.h"
#include "flang/Parser/parse-tree.h"
#include "flang/Parser/tools.h"
@@ -128,6 +129,7 @@ struct DeviceExprChecker
if (IsOnDevice(x)) {
return {};
}
+ common::CUDASubprogramAttrs calleeAttrs{common::CUDASubprogramAttrs::Host};
if (const Symbol * sym{x.GetInterfaceSymbol()}) {
const Symbol &ultimate{sym->GetUltimate()};
const auto *subp{ultimate.detailsIf<semantics::SubprogramDetails>()};
@@ -142,6 +144,7 @@ struct DeviceExprChecker
return (*this)(*stmtFunction);
}
if (auto attrs{subp->cudaSubprogramAttrs()}) {
+ calleeAttrs = *attrs;
if (*attrs == common::CUDASubprogramAttrs::HostDevice ||
*attrs == common::CUDASubprogramAttrs::Device) {
if (warpFunctions_.contains(sym->name().ToString()) &&
@@ -152,7 +155,7 @@ struct DeviceExprChecker
}
if (*attrs == common::CUDASubprogramAttrs::Device &&
callContext_ == CallContext::HostDevice) {
- return CallError(x, true);
+ return CallError(x, calleeAttrs);
}
return {};
}
@@ -186,21 +189,23 @@ struct DeviceExprChecker
if (callContext_ == CallContext::GuardedHostDevice) {
return {};
}
- return CallError(x, false);
+ return CallError(x, calleeAttrs);
}
SemanticsContext &context_;
CallContext callContext_;
private:
- Result CallError(
- const evaluate::ProcedureDesignator &proc, bool inHostCode) const {
- const char *procedureType{inHostCode ? "DEVICE" : "HOST"};
+ Result CallError(const evaluate::ProcedureDesignator &proc,
+ common::CUDASubprogramAttrs calleeAttrs) const {
+ const std::string procedureType{
+ calleeAttrs == common::CUDASubprogramAttrs::HostDevice
+ ? "HOST,DEVICE"
+ : parser::ToUpperCaseLetters(common::EnumToString(calleeAttrs))};
const char *codeContext{callContext_ == CallContext::HostDevice
? "HOST,DEVICE code without a guard using ON_DEVICE() from the "
"CUDADEVICE module"
- : inHostCode ? "HOST code"
- : "DEVICE code"};
+ : "DEVICE code"};
if (auto generic{proc.genericName()}) {
const Symbol *specific{proc.GetInterfaceSymbol()};
const std::string specificName{specific
diff --git a/flang/test/Semantics/CUDA/cuf-call-callee-kind.cuf b/flang/test/Semantics/CUDA/cuf-call-callee-kind.cuf
new file mode 100644
index 0000000000000..547c9bf3f831a
--- /dev/null
+++ b/flang/test/Semantics/CUDA/cuf-call-callee-kind.cuf
@@ -0,0 +1,53 @@
+! RUN: %python %S/../test_errors.py %s %flang_fc1
+
+module callee_kinds
+ interface grid_generic
+ module procedure grid_kernel
+ end interface
+contains
+ attributes(grid_global) subroutine grid_kernel()
+ end subroutine
+
+ attributes(global) subroutine global_kernel()
+ end subroutine
+
+ attributes(device) subroutine device_caller()
+ !ERROR: The GRID_GLOBAL procedure 'grid_kernel' may not be called in DEVICE code
+ call grid_kernel<<<1,1>>>()
+ !ERROR: The GRID_GLOBAL procedure 'grid_generic' (specific procedure 'grid_kernel') may not be called in DEVICE code
+ call grid_generic<<<1,1>>>()
+ !ERROR: not yet implemented: CUDA dynamic parallelism
+ call global_kernel<<<1,1>>>()
+ end subroutine
+
+ attributes(global) subroutine global_caller()
+ !ERROR: The GRID_GLOBAL procedure 'grid_kernel' may not be called in DEVICE code
+ call grid_kernel<<<1,1>>>()
+ end subroutine
+
+ attributes(grid_global) subroutine grid_caller()
+ !ERROR: The GRID_GLOBAL procedure 'grid_kernel' may not be called in DEVICE code
+ call grid_kernel<<<1,1>>>()
+ end subroutine
+
+ attributes(host,device) subroutine both_caller()
+ !ERROR: The GRID_GLOBAL procedure 'grid_kernel' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ call grid_kernel<<<1,1>>>()
+ !ERROR: The GRID_GLOBAL procedure 'grid_generic' (specific procedure 'grid_kernel') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ call grid_generic<<<1,1>>>()
+ end subroutine
+
+ subroutine host_caller()
+ call grid_kernel<<<1,1>>>()
+ call grid_generic<<<1,1>>>()
+ end subroutine
+end module
+
+module imported_callee_kind
+ use callee_kinds, only: renamed_kernel => grid_kernel
+contains
+ attributes(device) subroutine imported_caller()
+ !ERROR: The GRID_GLOBAL procedure 'renamed_kernel' may not be called in DEVICE code
+ call renamed_kernel<<<1,1>>>()
+ end subroutine
+end module
>From a0661464b927251518b47e9e6214bdefe4304d9e Mon Sep 17 00:00:00 2001
From: Andre Kuhlenschmidt <akuhlenschmi at nvidia.com>
Date: Fri, 16 Oct 2026 13:05:03 -0700
Subject: [PATCH 5/6] [flang][CUDA] Preserve explicit caller context in
diagnostics
---
flang/lib/Semantics/check-cuda.cpp | 14 ++++++++------
1 file changed, 8 insertions(+), 6 deletions(-)
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index d761a57e99f0a..bd949416c4fe4 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -155,7 +155,7 @@ struct DeviceExprChecker
}
if (*attrs == common::CUDASubprogramAttrs::Device &&
callContext_ == CallContext::HostDevice) {
- return CallError(x, calleeAttrs);
+ return CallError(x, calleeAttrs, /*inHostCode=*/true);
}
return {};
}
@@ -189,7 +189,7 @@ struct DeviceExprChecker
if (callContext_ == CallContext::GuardedHostDevice) {
return {};
}
- return CallError(x, calleeAttrs);
+ return CallError(x, calleeAttrs, /*inHostCode=*/false);
}
SemanticsContext &context_;
@@ -197,15 +197,17 @@ struct DeviceExprChecker
private:
Result CallError(const evaluate::ProcedureDesignator &proc,
- common::CUDASubprogramAttrs calleeAttrs) const {
+ common::CUDASubprogramAttrs calleeAttrs, bool inHostCode) const {
const std::string procedureType{
calleeAttrs == common::CUDASubprogramAttrs::HostDevice
? "HOST,DEVICE"
: parser::ToUpperCaseLetters(common::EnumToString(calleeAttrs))};
+ // clang-format off
const char *codeContext{callContext_ == CallContext::HostDevice
- ? "HOST,DEVICE code without a guard using ON_DEVICE() from the "
- "CUDADEVICE module"
- : "DEVICE code"};
+ ? "HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module"
+ : inHostCode ? "HOST code"
+ : "DEVICE code"};
+ // clang-format on
if (auto generic{proc.genericName()}) {
const Symbol *specific{proc.GetInterfaceSymbol()};
const std::string specificName{specific
>From 0c8d494ef8571ef07c27e2ec99a8f887a353a6c4 Mon Sep 17 00:00:00 2001
From: Andre Kuhlenschmidt <akuhlenschmi at nvidia.com>
Date: Fri, 16 Oct 2026 16:48:53 -0700
Subject: [PATCH 6/6] [flang][CUDA] Refine call diagnostic wording and cover
review examples
---
flang/lib/Semantics/check-cuda.cpp | 16 ++---
.../Semantics/CUDA/cuf-call-callee-kind.cuf | 14 ++--
.../Semantics/CUDA/cuf-call-diagnostics.cuf | 69 ++++++++++++++++++
.../CUDA/cuf-hostdevice-c-pointers.cuf | 4 +-
.../CUDA/cuf-hostdevice-call-context.cuf | 26 +++----
.../CUDA/cuf-hostdevice-else-if-guard.cuf | 4 +-
.../cuf-hostdevice-expression-context.cuf | 70 +++++++++----------
.../cuf-hostdevice-generic-diagnostics.cuf | 40 +++++------
.../CUDA/cuf-hostdevice-host-call.cuf | 2 +-
.../CUDA/cuf-hostdevice-inherited-spec.cuf | 2 +-
.../CUDA/cuf-hostdevice-module-calls.cuf | 4 +-
.../CUDA/cuf-hostdevice-return-guard.cuf | 40 +++++------
.../CUDA/cuf-hostdevice-stop-calls.cuf | 12 ++--
flang/test/Semantics/CUDA/cuf02.cuf | 2 +-
flang/test/Semantics/CUDA/cuf09.cuf | 12 ++--
15 files changed, 192 insertions(+), 125 deletions(-)
create mode 100644 flang/test/Semantics/CUDA/cuf-call-diagnostics.cuf
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index bd949416c4fe4..bedcbae0230fe 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -202,12 +202,10 @@ struct DeviceExprChecker
calleeAttrs == common::CUDASubprogramAttrs::HostDevice
? "HOST,DEVICE"
: parser::ToUpperCaseLetters(common::EnumToString(calleeAttrs))};
- // clang-format off
const char *codeContext{callContext_ == CallContext::HostDevice
- ? "HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module"
- : inHostCode ? "HOST code"
- : "DEVICE code"};
- // clang-format on
+ ? "host,device"
+ : inHostCode ? "host"
+ : "device"};
if (auto generic{proc.genericName()}) {
const Symbol *specific{proc.GetInterfaceSymbol()};
const std::string specificName{specific
@@ -215,13 +213,13 @@ struct DeviceExprChecker
: proc.GetName()};
if (*generic != specificName) {
return parser::MessageFormattedText(
- "The %s procedure '%s' (specific procedure '%s') may not be called in %s"_err_en_US,
- procedureType, *generic, specificName, codeContext);
+ "'%s' (specific procedure '%s'), a procedure with the %s attribute, can't be called from %s code."_err_en_US,
+ *generic, specificName, procedureType, codeContext);
}
}
return parser::MessageFormattedText(
- "The %s procedure '%s' may not be called in %s"_err_en_US,
- procedureType, proc.GetName(), codeContext);
+ "'%s', a procedure with the %s attribute, can't be called from %s code."_err_en_US,
+ proc.GetName(), procedureType, codeContext);
}
};
diff --git a/flang/test/Semantics/CUDA/cuf-call-callee-kind.cuf b/flang/test/Semantics/CUDA/cuf-call-callee-kind.cuf
index 547c9bf3f831a..c656d3d310b77 100644
--- a/flang/test/Semantics/CUDA/cuf-call-callee-kind.cuf
+++ b/flang/test/Semantics/CUDA/cuf-call-callee-kind.cuf
@@ -12,28 +12,28 @@ contains
end subroutine
attributes(device) subroutine device_caller()
- !ERROR: The GRID_GLOBAL procedure 'grid_kernel' may not be called in DEVICE code
+ !ERROR: 'grid_kernel', a procedure with the GRID_GLOBAL attribute, can't be called from device code.
call grid_kernel<<<1,1>>>()
- !ERROR: The GRID_GLOBAL procedure 'grid_generic' (specific procedure 'grid_kernel') may not be called in DEVICE code
+ !ERROR: 'grid_generic' (specific procedure 'grid_kernel'), a procedure with the GRID_GLOBAL attribute, can't be called from device code.
call grid_generic<<<1,1>>>()
!ERROR: not yet implemented: CUDA dynamic parallelism
call global_kernel<<<1,1>>>()
end subroutine
attributes(global) subroutine global_caller()
- !ERROR: The GRID_GLOBAL procedure 'grid_kernel' may not be called in DEVICE code
+ !ERROR: 'grid_kernel', a procedure with the GRID_GLOBAL attribute, can't be called from device code.
call grid_kernel<<<1,1>>>()
end subroutine
attributes(grid_global) subroutine grid_caller()
- !ERROR: The GRID_GLOBAL procedure 'grid_kernel' may not be called in DEVICE code
+ !ERROR: 'grid_kernel', a procedure with the GRID_GLOBAL attribute, can't be called from device code.
call grid_kernel<<<1,1>>>()
end subroutine
attributes(host,device) subroutine both_caller()
- !ERROR: The GRID_GLOBAL procedure 'grid_kernel' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'grid_kernel', a procedure with the GRID_GLOBAL attribute, can't be called from host,device code.
call grid_kernel<<<1,1>>>()
- !ERROR: The GRID_GLOBAL procedure 'grid_generic' (specific procedure 'grid_kernel') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'grid_generic' (specific procedure 'grid_kernel'), a procedure with the GRID_GLOBAL attribute, can't be called from host,device code.
call grid_generic<<<1,1>>>()
end subroutine
@@ -47,7 +47,7 @@ module imported_callee_kind
use callee_kinds, only: renamed_kernel => grid_kernel
contains
attributes(device) subroutine imported_caller()
- !ERROR: The GRID_GLOBAL procedure 'renamed_kernel' may not be called in DEVICE code
+ !ERROR: 'renamed_kernel', a procedure with the GRID_GLOBAL attribute, can't be called from device code.
call renamed_kernel<<<1,1>>>()
end subroutine
end module
diff --git a/flang/test/Semantics/CUDA/cuf-call-diagnostics.cuf b/flang/test/Semantics/CUDA/cuf-call-diagnostics.cuf
new file mode 100644
index 0000000000000..4de655af9e1d2
--- /dev/null
+++ b/flang/test/Semantics/CUDA/cuf-call-diagnostics.cuf
@@ -0,0 +1,69 @@
+! RUN: %python %S/../test_errors.py %s %flang_fc1
+
+! The attribute belongs to the selected specific procedure, including kernels.
+module mixed_callees
+ use cudadevice, only: on_device
+ interface mixed
+ module procedure mixed_host, mixed_dev
+ end interface
+contains
+ attributes(grid_global) subroutine k()
+ end subroutine
+
+ integer function mixed_host(n)
+ integer, value :: n
+ mixed_host = n
+ end function
+
+ attributes(device) real function mixed_dev(x)
+ real, value :: x
+ mixed_dev = x
+ end function
+
+ attributes(device) subroutine dev_caller()
+ !ERROR: 'k', a procedure with the GRID_GLOBAL attribute, can't be called from device code.
+ call k<<<1,1>>>()
+ end subroutine
+
+ attributes(host,device) subroutine hd_caller(n, x)
+ integer :: n
+ real :: x
+ !ERROR: 'k', a procedure with the GRID_GLOBAL attribute, can't be called from host,device code.
+ call k<<<1,1>>>()
+ !ERROR: 'mixed' (specific procedure 'mixed_host'), a procedure with the HOST attribute, can't be called from host,device code.
+ n = mixed(n)
+ !ERROR: 'mixed' (specific procedure 'mixed_dev'), a procedure with the DEVICE attribute, can't be called from host,device code.
+ x = mixed(x)
+ end subroutine
+
+ attributes(host,device) subroutine guarded_caller()
+ if (on_device()) then
+ call k<<<1,1>>>()
+ else
+ call k<<<1,1>>>()
+ end if
+ end subroutine
+end module
+
+! Calls in result lengths and dummy bounds cannot be enclosed in an IF guard.
+module specification_callees
+contains
+ pure integer function host_len(n)
+ integer, intent(in) :: n
+ host_len = n
+ end function
+
+ attributes(host,device) function f(n) result(r)
+ integer, value :: n
+ !ERROR: 'host_len', a procedure with the HOST attribute, can't be called from host,device code.
+ character(host_len(n)) :: r
+ r = ''
+ end function
+
+ attributes(host,device) subroutine s(n, a)
+ integer, value :: n
+ !ERROR: 'host_len', a procedure with the HOST attribute, can't be called from host,device code.
+ real :: a(host_len(n))
+ a = 0
+ end subroutine
+end module
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf
index 9383ed014680b..016af3222d635 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf
@@ -30,7 +30,7 @@ contains
integer, value :: i
end subroutine
attributes(host,device) subroutine pass_expression()
- !ERROR: The HOST procedure 'host_value' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_value', a procedure with the HOST attribute, can't be called from host,device code.
call consume_value(host_value())
end subroutine
attributes(host,device) subroutine both(p, q, f, g, d, e, result)
@@ -70,7 +70,7 @@ contains
end function
attributes(host,device) subroutine caller(result)
logical :: result
- !ERROR: The HOST procedure 'c_associated_c_ptr' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'c_associated_c_ptr', a procedure with the HOST attribute, can't be called from host,device code.
result = c_associated_c_ptr()
end subroutine
end module
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf
index 73362eafd89e3..8c8dada7becc8 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf
@@ -34,13 +34,13 @@ contains
attributes(host,device) integer function unguarded_host(n)
integer, value :: n
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
unguarded_host = host_only(n)
end function
attributes(host,device) integer function unguarded_device(n)
integer, value :: n
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, can't be called from host,device code.
unguarded_device = device_only(n)
end function
@@ -113,10 +113,10 @@ contains
attributes(host,device) integer function ordinary_branch(n)
integer, value :: n
if (n > 0) then
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
ordinary_branch = host_only(n)
else
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, can't be called from host,device code.
ordinary_branch = device_only(n)
end if
end function
@@ -190,7 +190,7 @@ contains
after_guard = host_only(n)
end if
! Both copies reach the statement after END IF.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
after_guard = after_guard + host_only(n)
end function
@@ -224,17 +224,17 @@ contains
attributes(host,device) subroutine unguarded_subroutines(n)
integer, intent(inout) :: n
- !ERROR: The HOST procedure 'host_subroutine' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_subroutine', a procedure with the HOST attribute, can't be called from host,device code.
call host_subroutine(n)
- !ERROR: The DEVICE procedure 'device_subroutine' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_subroutine', a procedure with the DEVICE attribute, can't be called from host,device code.
call device_subroutine(n)
end subroutine
attributes(host,device) subroutine generic_subroutines(n)
integer, intent(inout) :: n
- !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_subroutine') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_generic' (specific procedure 'host_subroutine'), a procedure with the HOST attribute, can't be called from host,device code.
call host_generic(n)
- !ERROR: The DEVICE procedure 'device_generic' (specific procedure 'device_subroutine') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_generic' (specific procedure 'device_subroutine'), a procedure with the DEVICE attribute, can't be called from host,device code.
call device_generic(n)
end subroutine
@@ -272,10 +272,10 @@ contains
attributes(host,device) subroutine misleading_bindc_guard(n)
integer, intent(inout) :: n
if (on_device()) then
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, can't be called from host,device code.
call device_only(n)
else
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only(n)
end if
end subroutine
@@ -325,10 +325,10 @@ contains
attributes(host,device) subroutine misleading_guard(n)
integer, intent(inout) :: n
if (on_device()) then
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, can't be called from host,device code.
call device_only(n)
else
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only(n)
end if
end subroutine
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf
index 1465960a7f717..6809e64034be0 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf
@@ -18,7 +18,7 @@ contains
call host_only()
end if
! The guard does not apply after the entire IF.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine later_else(x, y)
@@ -54,7 +54,7 @@ contains
call host_only()
else
! A nested guard in another arm does not guard this ELSE.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end if
end subroutine
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf
index 17d597f0c8381..e7d39d2411195 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf
@@ -15,7 +15,7 @@ contains
end function
attributes(host,device) subroutine select_h()
integer :: x
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
select case(h())
case(1)
x = 1
@@ -23,7 +23,7 @@ contains
end subroutine
attributes(host,device) subroutine select_d()
integer :: x
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
select case(d())
case(1)
x = 1
@@ -37,34 +37,34 @@ contains
end select
end subroutine
attributes(host,device) subroutine print_h()
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
print *, h()
end subroutine
attributes(host,device) subroutine print_d()
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
print *, d()
end subroutine
attributes(host,device) subroutine print_b()
print *, b()
end subroutine
attributes(host,device) subroutine write_item_h()
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
write(*,*) h()
end subroutine
attributes(host,device) subroutine write_item_d()
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
write(*,*) d()
end subroutine
attributes(host,device) subroutine write_item_b()
write(*,*) b()
end subroutine
attributes(host,device) subroutine write_unit_h()
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
!WARNING: I/O statement might not be supported on device [-Wcuda-usage]
write(h(),*) 1
end subroutine
attributes(host,device) subroutine write_unit_d()
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
!WARNING: I/O statement might not be supported on device [-Wcuda-usage]
write(d(),*) 1
end subroutine
@@ -74,12 +74,12 @@ contains
end subroutine
attributes(host,device) subroutine write_implied_do_h()
integer :: i
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
write(*,*) (i, i=1,h())
end subroutine
attributes(host,device) subroutine write_implied_do_d()
integer :: i
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
write(*,*) (i, i=1,d())
end subroutine
attributes(host,device) subroutine write_implied_do_b()
@@ -88,12 +88,12 @@ contains
end subroutine
attributes(host,device) subroutine write_subscript_h()
integer :: x(4)
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
write(*,*) x(h())
end subroutine
attributes(host,device) subroutine write_subscript_d()
integer :: x(4)
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
write(*,*) x(d())
end subroutine
attributes(host,device) subroutine write_subscript_b()
@@ -102,13 +102,13 @@ contains
end subroutine
attributes(host,device) subroutine read_subscript_h()
integer :: x(4)
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
!WARNING: I/O statement might not be supported on device [-Wcuda-usage]
read(*,*) x(h())
end subroutine
attributes(host,device) subroutine read_subscript_d()
integer :: x(4)
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
!WARNING: I/O statement might not be supported on device [-Wcuda-usage]
read(*,*) x(d())
end subroutine
@@ -119,12 +119,12 @@ contains
end subroutine
attributes(host,device) subroutine allocate_bound_h()
integer, allocatable :: a(:)
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
allocate(a(h()))
end subroutine
attributes(host,device) subroutine allocate_bound_d()
integer, allocatable :: a(:)
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
allocate(a(d()))
end subroutine
attributes(host,device) subroutine allocate_bound_b()
@@ -133,12 +133,12 @@ contains
end subroutine
attributes(host,device) subroutine allocate_source_h()
integer, allocatable :: a(:)
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
allocate(a,source=[h()])
end subroutine
attributes(host,device) subroutine allocate_source_d()
integer, allocatable :: a(:)
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
allocate(a,source=[d()])
end subroutine
attributes(host,device) subroutine allocate_source_b()
@@ -147,12 +147,12 @@ contains
end subroutine
attributes(host,device) subroutine allocate_mold_h()
integer, allocatable :: a(:)
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
allocate(a,mold=[h()])
end subroutine
attributes(host,device) subroutine allocate_mold_d()
integer, allocatable :: a(:)
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
allocate(a,mold=[d()])
end subroutine
attributes(host,device) subroutine allocate_mold_b()
@@ -161,12 +161,12 @@ contains
end subroutine
attributes(host,device) subroutine allocate_length_h()
character(:), allocatable :: text
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
allocate(character(h()) :: text)
end subroutine
attributes(host,device) subroutine allocate_length_d()
character(:), allocatable :: text
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
allocate(character(d()) :: text)
end subroutine
attributes(host,device) subroutine allocate_length_b()
@@ -176,13 +176,13 @@ contains
attributes(host,device) subroutine deallocate_stat_h()
integer, allocatable :: a(:)
integer :: status(4)
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
deallocate(a,stat=status(h()))
end subroutine
attributes(host,device) subroutine deallocate_stat_d()
integer, allocatable :: a(:)
integer :: status(4)
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
deallocate(a,stat=status(d()))
end subroutine
attributes(host,device) subroutine deallocate_stat_b()
@@ -191,12 +191,12 @@ contains
deallocate(a,stat=status(b()))
end subroutine
attributes(host,device) subroutine automatic_bound_h()
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
integer :: a(h())
a = 1
end subroutine
attributes(host,device) subroutine automatic_bound_d()
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
integer :: a(d())
a = 1
end subroutine
@@ -205,12 +205,12 @@ contains
a = 1
end subroutine
attributes(host,device) subroutine character_length_h()
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
character(h()) :: text
text = 'x'
end subroutine
attributes(host,device) subroutine character_length_d()
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
character(d()) :: text
text = 'x'
end subroutine
@@ -285,23 +285,23 @@ contains
end subroutine
attributes(host,device) subroutine block_bound_h()
block
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
integer :: a(h())
a = 1
end block
end subroutine
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
attributes(host,device) character(h()) function header_h()
header_h = 'x'
end function
attributes(host,device) subroutine block_bound_d()
block
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
integer :: a(d())
a = 1
end block
end subroutine
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
attributes(host,device) character(d()) function header_d()
header_d = 'x'
end function
@@ -325,13 +325,13 @@ contains
end function
attributes(device) subroutine device_bound(n)
integer, value :: n
- !ERROR: The HOST procedure 'h' may not be called in DEVICE code
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from device code.
real :: a(h(n))
a = 0.0
end subroutine
attributes(global) subroutine global_print(n)
integer, value :: n
- !ERROR: The HOST procedure 'h' may not be called in DEVICE code
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from device code.
print *, h(n)
end subroutine
subroutine kernel_print(a, n)
@@ -339,7 +339,7 @@ contains
integer, device :: a(n)
!$cuf kernel do <<<*,*>>>
do i = 1, n
- !ERROR: The HOST procedure 'h' may not be called in DEVICE code
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from device code.
print *, h(i)
a(i) = i
end do
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf
index bba05a010f6f7..4dc1b33c05524 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf
@@ -57,26 +57,26 @@ contains
attributes(host,device) subroutine unguarded(n, x)
integer :: n
real :: x
- !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_generic' (specific procedure 'host_integer'), a procedure with the HOST attribute, can't be called from host,device code.
n = host_generic(n)
- !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_real') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_generic' (specific procedure 'host_real'), a procedure with the HOST attribute, can't be called from host,device code.
x = host_generic(x)
- !ERROR: The DEVICE procedure 'device_generic' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_generic' (specific procedure 'device_integer'), a procedure with the DEVICE attribute, can't be called from host,device code.
n = device_generic(n)
- !ERROR: The HOST procedure 'host_sub_generic' (specific procedure 'host_sub') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_sub_generic' (specific procedure 'host_sub'), a procedure with the HOST attribute, can't be called from host,device code.
call host_sub_generic(n)
- !ERROR: The DEVICE procedure 'device_sub_generic' (specific procedure 'device_sub') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_sub_generic' (specific procedure 'device_sub'), a procedure with the DEVICE attribute, can't be called from host,device code.
call device_sub_generic(n)
! Ordinary calls keep their single-name diagnostics.
- !ERROR: The HOST procedure 'host_integer' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_integer', a procedure with the HOST attribute, can't be called from host,device code.
n = host_integer(n)
- !ERROR: The DEVICE procedure 'device_integer' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_integer', a procedure with the DEVICE attribute, can't be called from host,device code.
n = device_integer(n)
n = both_generic(n)
end subroutine
attributes(device) subroutine device_caller(n)
integer :: n
- !ERROR: The HOST procedure 'host_generic' (specific procedure 'host_integer') may not be called in DEVICE code
+ !ERROR: 'host_generic' (specific procedure 'host_integer'), a procedure with the HOST attribute, can't be called from device code.
n = host_generic(n)
n = device_generic(n)
end subroutine
@@ -99,18 +99,18 @@ module generic_alias_callers
contains
attributes(host,device) subroutine aliases(n)
integer :: n
- !ERROR: The HOST procedure 'renamed_host' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'renamed_host' (specific procedure 'host_integer'), a procedure with the HOST attribute, can't be called from host,device code.
n = renamed_host(n)
- !ERROR: The HOST procedure 'another_host' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'another_host' (specific procedure 'host_integer'), a procedure with the HOST attribute, can't be called from host,device code.
n = another_host(n)
- !ERROR: The DEVICE procedure 'renamed_device' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'renamed_device' (specific procedure 'device_integer'), a procedure with the DEVICE attribute, can't be called from host,device code.
n = renamed_device(n)
- !ERROR: The HOST procedure 'renamed_sub' (specific procedure 'host_sub') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'renamed_sub' (specific procedure 'host_sub'), a procedure with the HOST attribute, can't be called from host,device code.
call renamed_sub(n)
- !ERROR: The HOST procedure 'renamed_host' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'renamed_host' (specific procedure 'host_integer'), a procedure with the HOST attribute, can't be called from host,device code.
n = both_generic(renamed_host(n))
! A USE alias of an ordinary specific is not a generic reference.
- !ERROR: The HOST procedure 'specific_alias' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'specific_alias', a procedure with the HOST attribute, can't be called from host,device code.
n = specific_alias(n)
end subroutine
end module
@@ -122,13 +122,13 @@ contains
type(bound_targets) :: t
class(bound_targets) :: poly
integer :: n
- !ERROR: The HOST procedure 'host_bound' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_bound' (specific procedure 'host_integer'), a procedure with the HOST attribute, can't be called from host,device code.
n = t%host_bound(n)
- !ERROR: The DEVICE procedure 'device_bound' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_bound' (specific procedure 'device_integer'), a procedure with the DEVICE attribute, can't be called from host,device code.
n = t%device_bound(n)
- !ERROR: The HOST procedure 'host_bound' (specific procedure 'host_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_bound' (specific procedure 'host_integer'), a procedure with the HOST attribute, can't be called from host,device code.
n = poly%host_bound(n)
- !ERROR: The DEVICE procedure 'device_bound' (specific procedure 'device_integer') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_bound' (specific procedure 'device_integer'), a procedure with the DEVICE attribute, can't be called from host,device code.
n = poly%device_bound(n)
end subroutine
end module
@@ -139,7 +139,7 @@ contains
attributes(host,device) subroutine atomic_call(a, n, old)
integer, device :: a(*)
integer :: n, old
- !ERROR: The DEVICE procedure 'atomicadd' (specific procedure 'atomicaddi') may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'atomicadd' (specific procedure 'atomicaddi'), a procedure with the DEVICE attribute, can't be called from host,device code.
old = atomicadd(a(1), n)
end subroutine
end module
@@ -156,7 +156,7 @@ contains
attributes(host,device) subroutine caller(n)
integer :: n
! Do not repeat the name when generic and specific have the same spelling.
- !ERROR: The HOST procedure 'same_name' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'same_name', a procedure with the HOST attribute, can't be called from host,device code.
n = same_name(n)
end subroutine
end module
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-host-call.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-host-call.cuf
index ca18be94d7437..de9d5124f6bd5 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-host-call.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-host-call.cuf
@@ -49,7 +49,7 @@ contains
end function
end interface
double precision, intent(inout) :: x
- !ERROR: The HOST procedure 'hostfunc' may not be called in DEVICE code
+ !ERROR: 'hostfunc', a procedure with the HOST attribute, can't be called from device code.
x = hostfunc()
end subroutine
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf
index 09d4cf4037fd1..095a9adf0735f 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf
@@ -132,7 +132,7 @@ contains
extent = 1
end function
attributes(host,device) subroutine caller(a)
- !ERROR: The HOST procedure 'extent' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'extent', a procedure with the HOST attribute, can't be called from host,device code.
integer :: a(extent())
a = 1
end subroutine
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf
index 4c06d5ef80c1d..ca5b3a8495a42 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf
@@ -39,9 +39,9 @@ contains
module procedure caller
caller = both(n)
call both_subroutine(n)
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
caller = host_only(n)
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, can't be called from host,device code.
caller = device_only(n)
if (on_device()) then
caller = device_only(n)
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf
index 64d3009f4f261..0d7d1026ed9af 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf
@@ -69,7 +69,7 @@ contains
end if
end if
! The nested IF does not unconditionally return.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine nested_else_return(choose)
@@ -82,7 +82,7 @@ contains
end if
end if
! The nested ELSE does not make the enclosing arm return unconditionally.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine guarded_loop_return(choose)
@@ -93,7 +93,7 @@ contains
end do
end if
! Loops may not execute and are not traversed by the return scan.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine conditional_block_guard(choose)
@@ -105,7 +105,7 @@ contains
call host_only()
end if
! The enclosing conditional can bypass the whole BLOCK.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine return_before_last_statement()
@@ -134,7 +134,7 @@ contains
call device_only()
end if
! A nested conditional return does not eliminate fallthrough.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine else_return()
@@ -152,7 +152,7 @@ contains
call host_only()
end if
! The outer IF can bypass the return guard.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine loop_return(choose)
@@ -162,7 +162,7 @@ contains
call host_only()
end do
! The loop need not execute.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine reset_between_branches(choose, other)
@@ -171,13 +171,13 @@ contains
if (on_device()) return
call host_only()
else if (other) then
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
else
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, can't be called from host,device code.
call device_only()
end if
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine preserve_guarded_else_if(choose, other)
@@ -194,7 +194,7 @@ contains
! the host, but guard direction is not tracked yet.
call device_only()
end if
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine preserve_inherited_guard(choose)
@@ -215,7 +215,7 @@ contains
else
call both()
end if
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine ordinary_else_return(choose)
@@ -225,7 +225,7 @@ contains
else
return
end if
- !ERROR: The DEVICE procedure 'device_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, can't be called from host,device code.
call device_only()
end subroutine
attributes(host,device) subroutine unrelated_arm_returns(choose)
@@ -236,7 +236,7 @@ contains
return
end if
! The returning arm does not test ON_DEVICE.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine return_before_guard_arm(choose)
@@ -246,7 +246,7 @@ contains
else if (on_device()) then
call device_only()
end if
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine else_if_guard_else_returns(choose)
@@ -298,7 +298,7 @@ contains
attributes(host,device) subroutine ordinary_return(choose)
logical, intent(in) :: choose
if (choose) return
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine conditional_return(choose)
@@ -307,7 +307,7 @@ contains
if (choose) return
end if
! The device branch can fall through.
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
attributes(host,device) subroutine unrelated_label()
@@ -332,7 +332,7 @@ contains
attributes(host,device) subroutine else_if_bypass(choose)
logical, intent(in) :: choose
if (choose) then
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
else if (on_device()) then
return
@@ -343,7 +343,7 @@ contains
attributes(host,device) subroutine select_dispatch(result)
integer, intent(out) :: result
! The selector example from the review is already diagnosed.
- !ERROR: The DEVICE procedure 'device_value' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_value', a procedure with the DEVICE attribute, can't be called from host,device code.
select case (device_value())
case (1)
result = 10
@@ -365,7 +365,7 @@ contains
end subroutine
attributes(host,device) subroutine caller()
if (on_device()) return
- !ERROR: The HOST procedure 'host_only' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_only', a procedure with the HOST attribute, can't be called from host,device code.
call host_only()
end subroutine
end module
diff --git a/flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf b/flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf
index def7ec0d9a827..ea9c272b825c5 100644
--- a/flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf
+++ b/flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf
@@ -21,23 +21,23 @@ contains
attributes(host,device) subroutine caller()
! Review example: the STOP code expression runs on both targets.
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
stop d()
end subroutine
attributes(host,device) subroutine host_code()
- !ERROR: The HOST procedure 'h' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from host,device code.
stop h()
end subroutine
attributes(host,device) subroutine error_code()
- !ERROR: The DEVICE procedure 'd' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'd', a procedure with the DEVICE attribute, can't be called from host,device code.
error stop d()
end subroutine
attributes(host,device) subroutine quiet_device()
- !ERROR: The DEVICE procedure 'device_quiet' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'device_quiet', a procedure with the DEVICE attribute, can't be called from host,device code.
stop 1, quiet=device_quiet()
end subroutine
attributes(host,device) subroutine quiet_host()
- !ERROR: The HOST procedure 'host_quiet' may not be called in HOST,DEVICE code without a guard using ON_DEVICE() from the CUDADEVICE module
+ !ERROR: 'host_quiet', a procedure with the HOST attribute, can't be called from host,device code.
error stop 'failed', quiet=host_quiet()
end subroutine
attributes(host,device) subroutine both_code()
@@ -51,7 +51,7 @@ contains
end if
end subroutine
attributes(device) subroutine device_caller()
- !ERROR: The HOST procedure 'h' may not be called in DEVICE code
+ !ERROR: 'h', a procedure with the HOST attribute, can't be called from device code.
stop h()
end subroutine
end module
diff --git a/flang/test/Semantics/CUDA/cuf02.cuf b/flang/test/Semantics/CUDA/cuf02.cuf
index 3543b41feb687..0467fcdf09a2c 100644
--- a/flang/test/Semantics/CUDA/cuf02.cuf
+++ b/flang/test/Semantics/CUDA/cuf02.cuf
@@ -40,7 +40,7 @@ module m
end
attributes(device) subroutine s6_bad
badstmt(x) = hostfunc(x)
- !ERROR: The HOST procedure 'hostfunc' may not be called in DEVICE code
+ !ERROR: 'hostfunc', a procedure with the HOST attribute, can't be called from device code.
x = badstmt(x)
end
!ERROR: A function may not have ATTRIBUTES(GLOBAL) or ATTRIBUTES(GRID_GLOBAL)
diff --git a/flang/test/Semantics/CUDA/cuf09.cuf b/flang/test/Semantics/CUDA/cuf09.cuf
index 65eb06fd25b9b..8d04f821755e7 100644
--- a/flang/test/Semantics/CUDA/cuf09.cuf
+++ b/flang/test/Semantics/CUDA/cuf09.cuf
@@ -219,15 +219,15 @@ program main
end do
!$cuf kernel do <<< 1, 2 >>>
do j = 1, 10
- !ERROR: The HOST procedure 'foo' may not be called in DEVICE code
+ !ERROR: 'foo', a procedure with the HOST attribute, can't be called from device code.
call foo
- !ERROR: The HOST procedure 'bar' may not be called in DEVICE code
+ !ERROR: 'bar', a procedure with the HOST attribute, can't be called from device code.
x = bar()
- !ERROR: The HOST procedure 'ifunc' may not be called in DEVICE code
+ !ERROR: 'ifunc', a procedure with the HOST attribute, can't be called from device code.
if (ifunc() /= 0) continue
- !ERROR: The HOST procedure 'ifunc' may not be called in DEVICE code
+ !ERROR: 'ifunc', a procedure with the HOST attribute, can't be called from device code.
if (ifunc() /= 0) then
- !ERROR: The HOST procedure 'ifunc' may not be called in DEVICE code
+ !ERROR: 'ifunc', a procedure with the HOST attribute, can't be called from device code.
else if (ifunc() /= 1) then
end if
end do
@@ -256,7 +256,7 @@ subroutine host1()
!$cuf kernel do(1) <<<*,32>>>
do i = 1, 32
a(i) = a(i) * 2.0
- !ERROR: The HOST procedure 'syncthreads' may not be called in DEVICE code
+ !ERROR: 'syncthreads', a procedure with the HOST attribute, can't be called from device code.
call syncthreads() ! missing explicit use cudadevice
a(i) = a(i) + a(j) - 34.0
end do
More information about the flang-commits
mailing list