diff --git a/flang/include/flang/Parser/parse-tree.h b/flang/include/flang/Parser/parse-tree.h index 6dd4f2492cf22..73888bb640e40 100644 --- a/flang/include/flang/Parser/parse-tree.h +++ b/flang/include/flang/Parser/parse-tree.h @@ -3274,13 +3274,13 @@ struct FunctionReference { // R1521 call-stmt -> CALL procedure-designator [ chevrons ] // [( [actual-arg-spec-list] )] // (CUDA) chevrons -> <<< * | scalar-expr, scalar-expr [, -// scalar-int-expr [, scalar-int-expr ] ] >>> +// scalar-expr [, scalar-int-expr ] ] >>> struct CallStmt { BOILERPLATE(CallStmt); WRAPPER_CLASS(StarOrExpr, std::optional); struct Chevrons { TUPLE_CLASS_BOILERPLATE(Chevrons); - std::tuple, + std::tuple, std::optional> t; }; diff --git a/flang/lib/Parser/program-parsers.cpp b/flang/lib/Parser/program-parsers.cpp index 92c0a64b39a9d..740dbbfab9db7 100644 --- a/flang/lib/Parser/program-parsers.cpp +++ b/flang/lib/Parser/program-parsers.cpp @@ -484,7 +484,7 @@ constexpr auto starOrExpr{ applyFunction(presentOptional, scalarExpr))}; TYPE_PARSER(extension( "<<<" >> construct(starOrExpr, ", " >> scalarExpr, - maybe("," >> scalarIntExpr), maybe("," >> scalarIntExpr)) / + maybe("," >> scalarExpr), maybe("," >> scalarIntExpr)) / ">>>")) constexpr auto actualArgSpecList{optionalList(actualArgSpec)}; TYPE_CONTEXT_PARSER("CALL statement"_en_US, diff --git a/flang/test/Lower/CUDA/cuda-kernel-calls.cuf b/flang/test/Lower/CUDA/cuda-kernel-calls.cuf index 71e594e4742ec..e0941f74072ba 100644 --- a/flang/test/Lower/CUDA/cuda-kernel-calls.cuf +++ b/flang/test/Lower/CUDA/cuda-kernel-calls.cuf @@ -16,6 +16,7 @@ contains subroutine host() real, device :: a integer(8) :: stream + integer(4) :: nbytes ! CHECK-LABEL: func.func @_QMtest_callPhost() ! CHECK: %[[A:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda, uniq_name = "_QMtest_callFhostEa"} : (!fir.ref) -> (!fir.ref, !fir.ref) @@ -57,6 +58,10 @@ contains call dev_kernel1<<<*,32,0,stream>>>(a) ! CHECK: cuf.kernel_launch @_QMtest_callPdev_kernel1<<<%c-1{{.*}}, %c1{{.*}}, %c1{{.*}}, %c32{{.*}}, %c1{{.*}}, %c1{{.*}}, %c0{{.*}}, %{{.*}} : !fir.ref>>>(%{{.*}}) : (!fir.ref) + call dev_kernel1<<<*, 32, 0.8 * nbytes>>>(a) +! CHECK: %[[MUL:.*]] = arith.mulf %{{.*}}, %{{.*}} fastmath : f32 +! CHECK: %[[BYTES:.*]] = fir.convert %[[MUL]] : (f32) -> i32 +! CHECK: cuf.kernel_launch @_QMtest_callPdev_kernel1<<<%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %[[BYTES]]>>>(%{{.*}}) : (!fir.ref) end end diff --git a/flang/test/Parser/cuf-sanity-common b/flang/test/Parser/cuf-sanity-common index 816e03bed7220..2348c2edf3b73 100644 --- a/flang/test/Parser/cuf-sanity-common +++ b/flang/test/Parser/cuf-sanity-common @@ -43,6 +43,7 @@ module m call globalsub<<<1, 2>>> call globalsub<<<1, 2, 3>>> call globalsub<<<1, 2, 3, 4>>> + call globalsub<<<1, 2, 0.9*10, 4>>> call globalsub<<<*,5>>> allocate(pa(32), pinned = isPinned) end subroutine diff --git a/flang/test/Parser/cuf-sanity-tree.CUF b/flang/test/Parser/cuf-sanity-tree.CUF index 83d7540b8dec5..b4d53f27cf395 100644 --- a/flang/test/Parser/cuf-sanity-tree.CUF +++ b/flang/test/Parser/cuf-sanity-tree.CUF @@ -178,7 +178,7 @@ include "cuf-sanity-common" !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1' !CHECK: | | | | | | Scalar -> Expr = '2_4' !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2' -!CHECK: | | | | | | Scalar -> Integer -> Expr = '3_4' +!CHECK: | | | | | | Scalar -> Expr = '3_4' !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '3' !CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub<<<1_4,2_4,3_4,4_4>>>()' !CHECK: | | | | | Call @@ -188,10 +188,27 @@ include "cuf-sanity-common" !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1' !CHECK: | | | | | | Scalar -> Expr = '2_4' !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2' -!CHECK: | | | | | | Scalar -> Integer -> Expr = '3_4' +!CHECK: | | | | | | Scalar -> Expr = '3_4' !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '3' !CHECK: | | | | | | Scalar -> Integer -> Expr = '4_4' !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '4' +!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub<<<1_4,2_4,9._4,4_4>>>()' +!CHECK: | | | | | Call +!CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub' +!CHECK: | | | | | Chevrons +!CHECK: | | | | | | StarOrExpr -> Scalar -> Expr = '1_4' +!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1' +!CHECK: | | | | | | Scalar -> Expr = '2_4' +!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2' +!CHECK: | | | | | | Scalar -> Expr = '9._4' +!CHECK: | | | | | | | Multiply +!CHECK: | | | | | | | | Expr = '8.9999997615814208984375e-1_4' +!CHECK: | | | | | | | | | LiteralConstant -> RealLiteralConstant +!CHECK: | | | | | | | | | | Real = '0.9' +!CHECK: | | | | | | | | Expr = '10_4' +!CHECK: | | | | | | | | | LiteralConstant -> IntLiteralConstant = '10' +!CHECK: | | | | | | Scalar -> Integer -> Expr = '4_4' +!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '4' !CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> AllocateStmt !CHECK: | | | | | Allocation !CHECK: | | | | | | AllocateObject = 'pa'