[flang-commits] [flang] ed0624e - [flang][CUDA] Clarify host and device call diagnostics (#232850)
via flang-commits
flang-commits at lists.llvm.org
Thu Oct 22 08:54:23 PDT 2026
Author: Andre Kuhlenschmidt
Date: 2026-10-22T08:54:18-07:00
New Revision: ed0624e4bdb16060db08159ebce3e2402c123a84
URL: https://github.com/llvm/llvm-project/commit/ed0624e4bdb16060db08159ebce3e2402c123a84
DIFF: https://github.com/llvm/llvm-project/commit/ed0624e4bdb16060db08159ebce3e2402c123a84.diff
LOG: [flang][CUDA] Clarify host and device call diagnostics (#232850)
Report the callee's CUDA attribute and the calling context in CUDA
procedure-call diagnostics. Retain both the referenced generic name and
the selected specific procedure name when they differ, and report actual
attributes such as GRID_GLOBAL instead of inferring the callee kind from
the caller.
For example, a call to a HOST procedure from an unguarded HOST,DEVICE
procedure reports:
```text
'foo', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
```
Use lowercase `host` and `device` for individual calling contexts and
uppercase `HOST,DEVICE` for the combined context. Keep attribute names
uppercase. The message does not prescribe an ON_DEVICE() guard, since
some rejected calls occur in specification expressions where an IF guard
cannot be used.
This follows LLVM PR 228176. Call-validity checks and ON_DEVICE
recognition are unchanged. Tests cover selected generic specifics,
GRID_GLOBAL callees, module interfaces, guarded calls, and result-length
and dummy-bound expressions.
Validation: formatting and the full Flang checks pass.
Added:
flang/test/Semantics/CUDA/cuf-call-callee-kind.cuf
flang/test/Semantics/CUDA/cuf-call-diagnostics.cuf
Modified:
flang/lib/Semantics/check-cuda.cpp
flang/test/Semantics/CUDA/cuf-hostdevice-c-pointers.cuf
flang/test/Semantics/CUDA/cuf-hostdevice-call-context.cuf
flang/test/Semantics/CUDA/cuf-hostdevice-else-if-guard.cuf
flang/test/Semantics/CUDA/cuf-hostdevice-expression-context.cuf
flang/test/Semantics/CUDA/cuf-hostdevice-generic-diagnostics.cuf
flang/test/Semantics/CUDA/cuf-hostdevice-host-call.cuf
flang/test/Semantics/CUDA/cuf-hostdevice-inherited-spec.cuf
flang/test/Semantics/CUDA/cuf-hostdevice-module-calls.cuf
flang/test/Semantics/CUDA/cuf-hostdevice-return-guard.cuf
flang/test/Semantics/CUDA/cuf-hostdevice-stop-calls.cuf
flang/test/Semantics/CUDA/cuf02.cuf
flang/test/Semantics/CUDA/cuf09.cuf
Removed:
################################################################################
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index 581236cc725474..2aa6a911246186 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"
@@ -73,7 +74,7 @@ static const llvm::StringSet<> inlinePointerFunctions_ = {"c_associated_c_ptr",
"c_associated_c_funptr", "__builtin_c_ptr_eq", "__builtin_c_ptr_ne",
"__builtin_c_devptr_eq", "__builtin_c_devptr_ne", "__builtin_c_funloc"};
-enum class CallContext { Device, HostDevice, GuardedHostDevice };
+enum class CallContext { Host, Device, HostDevice, GuardedHostDevice };
// Match the BIND(C) intrinsic-module procedure intercepted by CUDA lowering,
// including calls through a USE rename. A user procedure named ON_DEVICE is
@@ -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,31 +189,36 @@ struct DeviceExprChecker
if (callContext_ == CallContext::GuardedHostDevice) {
return {};
}
- return CallError(x, false);
+ return CallError(x, calleeAttrs);
}
SemanticsContext &context_;
CallContext callContext_;
private:
- static Result CallError(
- const evaluate::ProcedureDesignator &proc, bool inHostCode) {
+ 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::Host ? "host"
+ : callContext_ == CallContext::Device ? "device"
+ : "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);
+ return parser::MessageFormattedText(
+ "'%s' (specific procedure '%s'), a procedure with the %s attribute, may not be called from %s code"_err_en_US,
+ *generic, specificName, procedureType, codeContext);
}
}
- 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());
+ return parser::MessageFormattedText(
+ "'%s', a procedure with the %s attribute, may not 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
new file mode 100644
index 00000000000000..da5512320793de
--- /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: 'grid_kernel', a procedure with the GRID_GLOBAL attribute, may not be called from device code
+ call grid_kernel<<<1,1>>>()
+ !ERROR: 'grid_generic' (specific procedure 'grid_kernel'), a procedure with the GRID_GLOBAL attribute, may not 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: 'grid_kernel', a procedure with the GRID_GLOBAL attribute, may not be called from device code
+ call grid_kernel<<<1,1>>>()
+ end subroutine
+
+ attributes(grid_global) subroutine grid_caller()
+ !ERROR: 'grid_kernel', a procedure with the GRID_GLOBAL attribute, may not be called from device code
+ call grid_kernel<<<1,1>>>()
+ end subroutine
+
+ attributes(host,device) subroutine both_caller()
+ !ERROR: 'grid_kernel', a procedure with the GRID_GLOBAL attribute, may not be called from HOST,DEVICE code
+ call grid_kernel<<<1,1>>>()
+ !ERROR: 'grid_generic' (specific procedure 'grid_kernel'), a procedure with the GRID_GLOBAL attribute, may not be called from HOST,DEVICE code
+ 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: 'renamed_kernel', a procedure with the GRID_GLOBAL attribute, may not 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 00000000000000..88ff56fbc36e9e
--- /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, may not 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, may not be called from HOST,DEVICE code
+ call k<<<1,1>>>()
+ !ERROR: 'mixed' (specific procedure 'mixed_host'), a procedure with the HOST attribute, may not be called from HOST,DEVICE code
+ n = mixed(n)
+ !ERROR: 'mixed' (specific procedure 'mixed_dev'), a procedure with the DEVICE attribute, may not 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, may not 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, may not 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 5e5750051e5da5..a18abbd0e26dca 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: 'host_value', a procedure with the HOST attribute, may not 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: 'c_associated_c_ptr' may not be called in device code
+ !ERROR: 'c_associated_c_ptr', a procedure with the HOST attribute, may not 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 e32476a67faa6c..3d49dcdd40381c 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: 'host_only', a procedure with the HOST attribute, may not 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: 'device_only' may not be called in host code
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
ordinary_branch = host_only(n)
else
- !ERROR: 'device_only' may not be called in host code
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_subroutine' may not be called in device code
+ !ERROR: 'host_subroutine', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
call host_subroutine(n)
- !ERROR: 'device_subroutine' may not be called in host code
+ !ERROR: 'device_subroutine', a procedure with the DEVICE attribute, may not be called from HOST,DEVICE code
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: 'host_generic' (specific procedure 'host_subroutine'), a procedure with the HOST attribute, may not be called from HOST,DEVICE code
call host_generic(n)
- !ERROR: 'device_generic' (specific procedure 'device_subroutine') may not be called in host code
+ !ERROR: 'device_generic' (specific procedure 'device_subroutine'), a procedure with the DEVICE attribute, may not 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: 'device_only' may not be called in host code
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, may not be called from HOST,DEVICE code
call device_only(n)
else
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'device_only' may not be called in host code
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, may not be called from HOST,DEVICE code
call device_only(n)
else
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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 c1ab90032cd61c..6b90fb25bf8efd 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: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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 00ca8f6502c6a4..810fb679805c6f 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: 'h', a procedure with the HOST attribute, may not 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: 'd' may not be called in host code
+ !ERROR: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
print *, h()
end subroutine
attributes(host,device) subroutine print_d()
- !ERROR: 'd' may not be called in host code
+ !ERROR: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
write(*,*) h()
end subroutine
attributes(host,device) subroutine write_item_d()
- !ERROR: 'd' may not be called in host code
+ !ERROR: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not 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: 'd' may not be called in host code
+ !ERROR: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
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: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
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: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not 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: 'd' may not be called in host code
+ !ERROR: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
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: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
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: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
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: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
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: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not 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: 'd' may not be called in host code
+ !ERROR: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
integer :: a(h())
a = 1
end subroutine
attributes(host,device) subroutine automatic_bound_d()
- !ERROR: 'd' may not be called in host code
+ !ERROR: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
character(h()) :: text
text = 'x'
end subroutine
attributes(host,device) subroutine character_length_d()
- !ERROR: 'd' may not be called in host code
+ !ERROR: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
integer :: a(h())
a = 1
end block
end subroutine
- !ERROR: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not 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: 'd' may not be called in host code
+ !ERROR: 'd', a procedure with the DEVICE attribute, may not be called from HOST,DEVICE code
integer :: a(d())
a = 1
end block
end subroutine
- !ERROR: 'd' may not be called in host code
+ !ERROR: 'd', a procedure with the DEVICE attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from 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: 'h', a procedure with the HOST attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not 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 599775338dd4d5..8b8385bcb37390 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: 'host_generic' (specific procedure 'host_integer'), a procedure with the HOST attribute, may not be called from HOST,DEVICE code
n = host_generic(n)
- !ERROR: 'host_generic' (specific procedure 'host_real') may not be called in device code
+ !ERROR: 'host_generic' (specific procedure 'host_real'), a procedure with the HOST attribute, may not be called from HOST,DEVICE code
x = host_generic(x)
- !ERROR: 'device_generic' (specific procedure 'device_integer') may not be called in host code
+ !ERROR: 'device_generic' (specific procedure 'device_integer'), a procedure with the DEVICE attribute, may not be called from HOST,DEVICE code
n = device_generic(n)
- !ERROR: 'host_sub_generic' (specific procedure 'host_sub') may not be called in device code
+ !ERROR: 'host_sub_generic' (specific procedure 'host_sub'), a procedure with the HOST attribute, may not be called from HOST,DEVICE code
call host_sub_generic(n)
- !ERROR: 'device_sub_generic' (specific procedure 'device_sub') may not be called in host code
+ !ERROR: 'device_sub_generic' (specific procedure 'device_sub'), a procedure with the DEVICE attribute, may not be called from HOST,DEVICE code
call device_sub_generic(n)
! Ordinary calls keep their single-name diagnostics.
- !ERROR: 'host_integer' may not be called in device code
+ !ERROR: 'host_integer', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
n = host_integer(n)
- !ERROR: 'device_integer' may not be called in host code
+ !ERROR: 'device_integer', a procedure with the DEVICE attribute, may not 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: '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, may not 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: 'renamed_host' (specific procedure 'host_integer') may not be called in device code
+ !ERROR: 'renamed_host' (specific procedure 'host_integer'), a procedure with the HOST attribute, may not be called from HOST,DEVICE code
n = renamed_host(n)
- !ERROR: 'another_host' (specific procedure 'host_integer') may not be called in device code
+ !ERROR: 'another_host' (specific procedure 'host_integer'), a procedure with the HOST attribute, may not be called from HOST,DEVICE code
n = another_host(n)
- !ERROR: 'renamed_device' (specific procedure 'device_integer') may not be called in host code
+ !ERROR: 'renamed_device' (specific procedure 'device_integer'), a procedure with the DEVICE attribute, may not be called from HOST,DEVICE code
n = renamed_device(n)
- !ERROR: 'renamed_sub' (specific procedure 'host_sub') may not be called in device code
+ !ERROR: 'renamed_sub' (specific procedure 'host_sub'), a procedure with the HOST attribute, may not be called from HOST,DEVICE code
call renamed_sub(n)
- !ERROR: 'renamed_host' (specific procedure 'host_integer') may not be called in device code
+ !ERROR: 'renamed_host' (specific procedure 'host_integer'), a procedure with the HOST attribute, may not 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: 'specific_alias' may not be called in device code
+ !ERROR: 'specific_alias', a procedure with the HOST attribute, may not 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: 'host_bound' (specific procedure 'host_integer') may not be called in device code
+ !ERROR: 'host_bound' (specific procedure 'host_integer'), a procedure with the HOST attribute, may not be called from HOST,DEVICE code
n = t%host_bound(n)
- !ERROR: 'device_bound' (specific procedure 'device_integer') may not be called in host code
+ !ERROR: 'device_bound' (specific procedure 'device_integer'), a procedure with the DEVICE attribute, may not be called from HOST,DEVICE code
n = t%device_bound(n)
- !ERROR: 'host_bound' (specific procedure 'host_integer') may not be called in device code
+ !ERROR: 'host_bound' (specific procedure 'host_integer'), a procedure with the HOST attribute, may not be called from HOST,DEVICE code
n = poly%host_bound(n)
- !ERROR: 'device_bound' (specific procedure 'device_integer') may not be called in host code
+ !ERROR: 'device_bound' (specific procedure 'device_integer'), a procedure with the DEVICE attribute, may not 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: 'atomicadd' (specific procedure 'atomicaddi') may not be called in host code
+ !ERROR: 'atomicadd' (specific procedure 'atomicaddi'), a procedure with the DEVICE attribute, may not 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: 'same_name' may not be called in device code
+ !ERROR: 'same_name', a procedure with the HOST attribute, may not 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 a86bcb5c7ac36a..316accc24f676b 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: 'hostfunc', a procedure with the HOST attribute, may not 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 2228ead80c7394..c1b78f16120c1d 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: 'extent', a procedure with the HOST attribute, may not 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 bbfb1dd655aeb3..f1db008032164a 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: 'host_only', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
caller = host_only(n)
- !ERROR: 'device_only' may not be called in host code
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, may not 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 d59f25787249ae..b9a84f8fefd249 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: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
call host_only()
else
- !ERROR: 'device_only' may not be called in host code
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, may not be called from HOST,DEVICE code
call device_only()
end if
- !ERROR: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'device_only' may not be called in host code
+ !ERROR: 'device_only', a procedure with the DEVICE attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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: 'device_value' may not be called in host code
+ !ERROR: 'device_value', a procedure with the DEVICE attribute, may not 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: 'host_only' may not be called in device code
+ !ERROR: 'host_only', a procedure with the HOST attribute, may not 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 a08448af736d5c..7a10fbb87e5139 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: 'd', a procedure with the DEVICE attribute, may not be called from HOST,DEVICE code
stop d()
end subroutine
attributes(host,device) subroutine host_code()
- !ERROR: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not be called from HOST,DEVICE code
stop h()
end subroutine
attributes(host,device) subroutine error_code()
- !ERROR: 'd' may not be called in host code
+ !ERROR: 'd', a procedure with the DEVICE attribute, may not be called from HOST,DEVICE code
error stop d()
end subroutine
attributes(host,device) subroutine quiet_device()
- !ERROR: 'device_quiet' may not be called in host code
+ !ERROR: 'device_quiet', a procedure with the DEVICE attribute, may not be called from HOST,DEVICE code
stop 1, quiet=device_quiet()
end subroutine
attributes(host,device) subroutine quiet_host()
- !ERROR: 'host_quiet' may not be called in device code
+ !ERROR: 'host_quiet', a procedure with the HOST attribute, may not 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: 'h' may not be called in device code
+ !ERROR: 'h', a procedure with the HOST attribute, may not 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 9a01a693f16e75..edeb0fd88a54a5 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: 'hostfunc', a procedure with the HOST attribute, may not 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 09f1abce098eb9..46a537988c6521 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: 'foo', a procedure with the HOST attribute, may not be called from device code
call foo
- !ERROR: 'bar' may not be called in device code
+ !ERROR: 'bar', a procedure with the HOST attribute, may not be called from device code
x = bar()
- !ERROR: 'ifunc' may not be called in device code
+ !ERROR: 'ifunc', a procedure with the HOST attribute, may not be called from device code
if (ifunc() /= 0) continue
- !ERROR: 'ifunc' may not be called in device code
+ !ERROR: 'ifunc', a procedure with the HOST attribute, may not be called from device code
if (ifunc() /= 0) then
- !ERROR: 'ifunc' may not be called in device code
+ !ERROR: 'ifunc', a procedure with the HOST attribute, may not 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: 'syncthreads' may not be called in device code
+ !ERROR: 'syncthreads', a procedure with the HOST attribute, may not 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