<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/88451>88451</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[flang] Runtime crash when parsing cuda fortran program with scale function on device
</td>
</tr>
<tr>
<th>Labels</th>
<td>
flang
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
ImanHosseini
</td>
</tr>
</table>
<pre>
Runtime crash when running `bbc -emit-hlfir -fcuda f.cuf`:
```
! f.cuf
module mod1
implicit none
contains
attributes(global) subroutine sub1(adev)
real :: adev(10)
integer :: tid
tid = threadIdx%x
adev(tid + 1) = scale(real(tid), 2)
end subroutine sub1
subroutine host_sub()
real, device :: adev(10)
real :: ahost(10)
integer :: i
ahost = 0.0
adev = ahost
call sub1<<<1,10>>>(adev)
ahost = adev
print*, ahost
end subroutine host_sub
end module mod1
program test
use mod1
implicit none
call host_sub()
end program test
```
The stack dump:
```
0. Program arguments: bbc -emit-hlfir -fcuda f.cuf -o f.mlir
#0 0x000055dd8ee2bede llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/local/home/imanh/llvm_i/install/bin/bbc+0x4bd2ede)
#1 0x000055dd8ee2914b llvm::sys::RunSignalHandlers() (/local/home/imanh/llvm_i/install/bin/bbc+0x4bd014b)
#2 0x000055dd8ee29285 SignalHandler(int) Signals.cpp:0:0
#3 0x00007f21a5a42520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
#4 0x000055dd8cba84cb std::set<Fortran::common::Reference<Fortran::semantics::Symbol const>, Fortran::semantics::SymbolAddressCompare, std::allocator<Fortran::common::Reference<Fortran::semantics::Symbol const>>> Fortran::evaluate::CollectSymbols<Fortran::evaluate::Expr<Fortran::evaluate::SomeType>>(Fortran::evaluate::Expr<Fortran::evaluate::SomeType> const&) (/local/home/imanh/llvm_i/install/bin/bbc+0x294f4cb)
#5 0x000055dd8d3c8005 Fortran::semantics::CUDAChecker::Enter(Fortran::parser::AssignmentStmt const&) (/local/home/imanh/llvm_i/install/bin/bbc+0x316f005)
#6 0x000055dd8d0f339c std::enable_if<TupleTrait<Fortran::parser::AssignmentStmt>, void>::type Fortran::parser::Walk<Fortran::parser::AssignmentStmt, Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>>(Fortran::parser::AssignmentStmt const&, Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>&) (/local/home/imanh/llvm_i/install/bin/bbc+0x2e9a39c)
#7 0x000055dd8d112921 void Fortran::common::log2visit::Log2VisitHelper<0ul, 5ul, void, void Fortran::parser::Walk<Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>, Fortran::parser::ExecutableConstruct, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::FormatStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::EntryStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::DataStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::NamelistStmt, false>>, Fortran::parser::ErrorRecovery>(std::variant<Fortran::parser::ExecutableConstruct, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::FormatStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::EntryStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::DataStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::NamelistStmt, false>>, Fortran::parser::ErrorRecovery> const&, Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>&)::'lambda'(Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker> const&), std::variant<Fortran::parser::ExecutableConstruct, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::FormatStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::EntryStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::DataStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::NamelistStmt, false>>, Fortran::parser::ErrorRecovery> const&>(void Fortran::parser::Walk<Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>, Fortran::parser::ExecutableConstruct, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::FormatStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::EntryStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::DataStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::NamelistStmt, false>>, Fortran::parser::ErrorRecovery>(std::variant<Fortran::parser::ExecutableConstruct, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::FormatStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::EntryStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::DataStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::NamelistStmt, false>>, Fortran::parser::ErrorRecovery> const&, Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>&)::'lambda'(Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker> const&)&&, unsigned long, std::variant<Fortran::parser::ExecutableConstruct, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::FormatStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::EntryStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::DataStmt, false>>, Fortran::parser::Statement<Fortran::common::Indirection<Fortran::parser::NamelistStmt, false>>, Fortran::parser::ErrorRecovery> const&) semantics.cpp:0:0
#8 0x000055dd8d114dd5 std::enable_if<TupleTrait<Fortran::parser::SubroutineSubprogram>, void>::type Fortran::parser::Walk<Fortran::parser::SubroutineSubprogram, Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>>(Fortran::parser::SubroutineSubprogram const&, Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>&) (/local/home/imanh/llvm_i/install/bin/bbc+0x2ebbdd5)
#9 0x000055dd8d114fe4 void Fortran::parser::ForEachInTuple<0ul, void Fortran::parser::Walk<Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>, Fortran::parser::Statement<Fortran::parser::ContainsStmt>, std::__cxx11::list<Fortran::parser::ModuleSubprogram, std::allocator<Fortran::parser::ModuleSubprogram>>>(std::tuple<Fortran::parser::Statement<Fortran::parser::ContainsStmt>, std::__cxx11::list<Fortran::parser::ModuleSubprogram, std::allocator<Fortran::parser::ModuleSubprogram>>> const&, Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>&)::'lambda'(Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker> const&), std::tuple<Fortran::parser::Statement<Fortran::parser::ContainsStmt>, std::__cxx11::list<Fortran::parser::ModuleSubprogram, std::allocator<Fortran::parser::ModuleSubprogram>>>>(std::tuple<Fortran::parser::Statement<Fortran::parser::ContainsStmt>, std::__cxx11::list<Fortran::parser::ModuleSubprogram, std::allocator<Fortran::parser::ModuleSubprogram>>> const&, void Fortran::parser::Walk<Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>, Fortran::parser::Statement<Fortran::parser::ContainsStmt>, std::__cxx11::list<Fortran::parser::ModuleSubprogram, std::allocator<Fortran::parser::ModuleSubprogram>>>(std::tuple<Fortran::parser::Statement<Fortran::parser::ContainsStmt>, std::__cxx11::list<Fortran::parser::ModuleSubprogram, std::allocator<Fortran::parser::ModuleSubprogram>>> const&, Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>&)::'lambda'(Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker> const&)) (/local/home/imanh/llvm_i/install/bin/bbc+0x2ebbfe4)
#10 0x000055dd8d115057 void Fortran::parser::ForEachInTuple<0ul, void Fortran::parser::Walk<Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>, Fortran::parser::Statement<Fortran::parser::ModuleStmt>, Fortran::parser::SpecificationPart, std::optional<Fortran::parser::ModuleSubprogramPart>, Fortran::parser::Statement<Fortran::parser::EndModuleStmt>>(std::tuple<Fortran::parser::Statement<Fortran::parser::ModuleStmt>, Fortran::parser::SpecificationPart, std::optional<Fortran::parser::ModuleSubprogramPart>, Fortran::parser::Statement<Fortran::parser::EndModuleStmt>> const&, Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>&)::'lambda'(Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker> const&), std::tuple<Fortran::parser::Statement<Fortran::parser::ModuleStmt>, Fortran::parser::SpecificationPart, std::optional<Fortran::parser::ModuleSubprogramPart>, Fortran::parser::Statement<Fortran::parser::EndModuleStmt>>>(std::tuple<Fortran::parser::Statement<Fortran::parser::ModuleStmt>, Fortran::parser::SpecificationPart, std::optional<Fortran::parser::ModuleSubprogramPart>, Fortran::parser::Statement<Fortran::parser::EndModuleStmt>> const&, void Fortran::parser::Walk<Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>, Fortran::parser::Statement<Fortran::parser::ModuleStmt>, Fortran::parser::SpecificationPart, std::optional<Fortran::parser::ModuleSubprogramPart>, Fortran::parser::Statement<Fortran::parser::EndModuleStmt>>(std::tuple<Fortran::parser::Statement<Fortran::parser::ModuleStmt>, Fortran::parser::SpecificationPart, std::optional<Fortran::parser::ModuleSubprogramPart>, Fortran::parser::Statement<Fortran::parser::EndModuleStmt>> const&, Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker>&)::'lambda'(Fortran::semantics::SemanticsVisitor<Fortran::semantics::CUDAChecker> const&)) (/local/home/imanh/llvm_i/install/bin/bbc+0x2ebc057)
#11 0x000055dd8d11671f Fortran::semantics::Semantics::Perform() (/local/home/imanh/llvm_i/install/bin/bbc+0x2ebd71f)
#12 0x000055dd8b7c2209 main (/local/home/imanh/llvm_i/install/bin/bbc+0x1569209)
#13 0x00007f21a5a29d90 __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:58:16
#14 0x00007f21a5a29e40 call_init ./csu/../csu/libc-start.c:128:20
#15 0x00007f21a5a29e40 __libc_start_main ./csu/../csu/libc-start.c:379:5
#16 0x000055dd8b7b36c5 _start (/local/home/imanh/llvm_i/install/bin/bbc+0x155a6c5)
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzsW99z2y4S_2vwCxMPQr-sBz-kTjztzN1Np8ndPXoQrGyuCDyAUue_v0FSbMtx0vbitt9r1XFjAcuyPz67XoNhzsm1Bpij9B1Kbyas8Rtj5x9qpt8b50BqOSmNeJx_arSXNWBumdvgLxvQ2DZaS73GKCNlyfEV1NJfbVQlLb6qeCMYrqa8qVBGUHyNyA0i1-G5feG-TaOeqG3WRjQKcG1E1HVgjLGst0py6bE2Gg7d3GjPpHZ9D_PeyrLx4BCdrZUpmUK0wK4prWm81BAeI0RnTMADosWBUfhngSkcpIyvcUcwi8iBSmoPa7BPFF6K4XQvBUbxDfYbC0x8EDtE092QpOfaUtJ3OArChSmOMwWIzoIE3XhYli4wHQgJWjzTZT94NLAxzq9cUyI6O6tk4CzgQXJ4Rd1zRgmMz5OdGEee6B0mtpqSKXluknakpxkMcqZUp2a86F4RoouIoPi2f-1diV9asVugG9xaqT2i18EAnTIv2XZvwpYgDD6D5daatWU19nDMp3HfCt2g2zlPhcWe894HTde83wB2nvHPWDT19nlodU0y7azxsWfH7LqpQXsXfPRauOIrg6tpraTtBUY0JpjsCCEkTYWYAdASBGClHurO5-7RdQ8fg5Hvgmz3lvEA6wORZV9WxnkLrEY0C25oHVLg1gBLZXhA53JjakB0KWumN6FbPdQrGTq080wFilLq8LfkiL4ju6QUFAQcUIloHJ2IW0RJeU7cT42-k2vN1HumhQLrOl-8XSISJeVAInoqEZ2leLA2orPeHl23m_JtcC5p_-8ZxT2jvKIRS1lCU0qe5JUlosvdLFtlyZWSutldrXXTDfCpM9OsEy9MGQiXHAvHSzZLeImdF721wKN4sTTWW6a7Lm7q2vTPn6ACC5rDKZGDmmkveW_ru8e6NCpkbefb6F3gr5NfC2HBuYWpt8xCmLMXi6ngH2_sjxCufQ3lgwemGuahay2MUsB9N8-dch_S3u62z4QcUtyZGu4ft7BPbJfi1msUwu2tqKZFUiV8iOr0GDgi5jNC0le9uvjnzfViA_wz2F4b7VvoDyZtmXVPBNdtfRIS152v_QX1iaOsIiQd6JMN9CFVHBf8gDjQrFSwkhWKF_fNVsG9ZfJZbLwsfI_6ByNF6-hA4R-3gF-c_2-mPn8H_6-F1FPzX9LJM4Hzmq_O4fKb3PQjRboIrKFgccEHMMgHMIgiWtCodRt-MdEos6YPQYeu-Tezpq1K70Ftg7AL0rSlV9q9tRjo37_P_Zc14OLlxW93wBsfEL8IzrQNfw6wY_o7zzwEALyWjz9oIS1wL0PHy7heGluzPaYrptw-Nf4kCW61t4-_UoAb5tmvXP8frAYl3Xc74dZaYz8BNw9gH7uksU-gD8xK9ly6EXUj6i6Mup_5AdSNIJorVpeCIZqfflBecs3jCmhQD4_BNQbXzw6uNr2PNcwI-N8b8GMNM6JurGF-VA2T9So2uj2JElgZvR5LmzHmfm3MFXiP4fM78bOTXaJEiPQtm4V3-xOou6bsD4Euv2V4dpVfuHF4Tp7_q-3DshRiuItcnAKjguQrO31LY28Z33zQLVQO24V_4dr6xfA9Jlr0Z_RH29_7CFmt-G4XRV0jBO9rjP7ensAOIfv1o6BXORydI-85-d78v7_Wf0p1cWSwP8e5I6rH5DkmzzF5vjV5XqI6qiDZV0eIxhE5qY5SkuZ_aHXU4-sA9JfZbYHLSnIWvgh9ZNYPwGy2oZup78Fyy-QSStxqMdTjRySF39VSf0oiuXgV9rsCYoye_zV6xg-Dv6aTRjiPHwYXryo5SfNBVRmdVJVZHlXfZsOu-RFsZWx9kd9_UyhFHlUD-Qa__y5zTikpcM2kfuNaUZoVlBSDtU5-Ik4LURC8WilZ8pXzzPoVZ0qt2tWniC65axBdTsOje3QCtg7Rpd561f90_HTSdIPi63SG4usoO6yanK4KCWkvOqyklv50oe4xcL9quU95YEcDU0oOTNNzTAeqnNPiLPM4D6BPD7yzoUvKOOMp7ri-2Skpy_hhU_jkZkb3dyLmsSjigk1gHuURneVZkswmm3keZ6KKeD6rCgFxFqdVToo8i2OSR1nO4omcU0ITkkQRjQlJ8mnGaZJUNCE8hjRlFUoI1EyqaZB0aux6Ip1rYD6bJWk0UawE5dr7XpRWiuk1ohSlNxM7D_RXZbN2KCHh67U7cPDSq_aSWDcjvcFnLoSFBCf1GneXSbro299n-SL9prvthKtGt0c62Oiny0iTxqr5xvttG5B0iehyLf2mKafc1L3R-7errTX_Ae6DB4JeAa6tav8NAAD__zoCOPY">