From patchwork Thu Jul 13 10:54:00 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 119771 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:a6b2:0:b0:3e4:2afc:c1 with SMTP id c18csp1736384vqm; Thu, 13 Jul 2023 03:54:58 -0700 (PDT) X-Google-Smtp-Source: APBJJlGX0PuUn3I117xw2WI1gvZLDSYHozmINn9hxdq4IJiDtzlsbM2Qy2ZZVnbxc7j5vbnFWDiV X-Received: by 2002:a17:906:f15:b0:988:6e75:6b3d with SMTP id z21-20020a1709060f1500b009886e756b3dmr1015419eji.33.1689245697573; Thu, 13 Jul 2023 03:54:57 -0700 (PDT) Received: from server2.sourceware.org (server2.sourceware.org. [2620:52:3:1:0:246e:9693:128c]) by mx.google.com with ESMTPS id k17-20020a17090627d100b0098886d50308si6485602ejc.490.2023.07.13.03.54.57 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Thu, 13 Jul 2023 03:54:57 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) client-ip=2620:52:3:1:0:246e:9693:128c; Authentication-Results: mx.google.com; dkim=pass header.i=@gcc.gnu.org header.s=default header.b=gnlWH6Fj; arc=fail (signature failed); spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org"; dmarc=pass (p=NONE sp=NONE dis=NONE) header.from=gnu.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 89B45385734E for ; Thu, 13 Jul 2023 10:54:55 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 89B45385734E DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1689245695; bh=QYesUyonszDoPo6plxaN42xLzOzcFzIq9KECuY5P0fI=; h=Date:Subject:To:Cc:References:In-Reply-To:List-Id: List-Unsubscribe:List-Archive:List-Post:List-Help:List-Subscribe: From:Reply-To:From; b=gnlWH6Fj/nf/Zg0d0ueF4IIa3bbwG92XGb77lCRa4VmiwpT3ebMZ9SdSBmhKTyjw+ 0U1BublqO+9ci03URVGTfUPQTsutrPIaxTP+ppRRgITuuyQdr2yil/o+ayJFs9pBjA VZJSzM4SUGO0JxQvRlcCUWEFsPMjjbmw+eaiHD7Q= X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from APC01-PSA-obe.outbound.protection.outlook.com (mail-psaapc01on2057.outbound.protection.outlook.com [40.107.255.57]) by sourceware.org (Postfix) with ESMTPS id CA3C43858C74 for ; Thu, 13 Jul 2023 10:54:10 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org CA3C43858C74 ARC-Seal: i=1; a=rsa-sha256; s=arcselector9901; d=microsoft.com; cv=none; b=l/5MG72S9a+W2ORxTnFT9rJQUR7tuCDXe1F1fTccZ9/OkXoQSU78AzT320zf3bmNGr6b5gUwrteNbwkB4JkfQEpRw5rbdyJ4lbvVUVV8NwSOpWcaxm2oL7m+XwaYADy8RgrKgbTMVS1ZQFq8iM3aJ1YQhgMgpJCfbpdayhWsagP6Ah9yPHEwmy6FF5n2T9xW7gJaTATiSCt8xiHuxApVzBIIOKjN8kWUDkh1pQBg4z9YEYRat2gPoxFa+/Llbh6AmIgSdn75GLg/XUYYfPqbbla7ixsqFrw67M73nzy3Jc3YLR/roWXqv0dKBR5lN3PLLdIs6Igp8CelYg4EiyfxEw== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=microsoft.com; s=arcselector9901; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-AntiSpam-MessageData-ChunkCount:X-MS-Exchange-AntiSpam-MessageData-0:X-MS-Exchange-AntiSpam-MessageData-1; bh=QYesUyonszDoPo6plxaN42xLzOzcFzIq9KECuY5P0fI=; b=XtFQOQ0odZV2TylFOS0yzL82cfMeNZr8A0fGgbPps4vQfbXS9giLQvm1/1B4dWQfc8JyNFFTTGiNc8d63T0hzrbsDs18tT8ILjC022vProw70+iQy7iTewy9IbNRsSWC/vwRQyWD8PfCgpohYW6xdjsvqnR0wrDxul3IVnCWdiOV4Ob7cHP8PUNVdRNl7K2Y3QMydML3dgREKDBgAcw7yMgrWOG6RgheSnxEipRO+O6NCYUrjqIvdpPdkP9M5bU0xFwBBQNX7ikdRl0UFveNxD6PkH9c1IP98EdbQQxLGqHfy2PNqETwUoY2IoNfPEDRpHbnqCYqDVI+lN4Eg6zucQ== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass smtp.mailfrom=siemens.com; dmarc=pass action=none header.from=siemens.com; dkim=pass header.d=siemens.com; arc=none Received: from SG2PR06MB5430.apcprd06.prod.outlook.com (2603:1096:4:1ba::14) by SEZPR06MB5479.apcprd06.prod.outlook.com (2603:1096:101:a2::13) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.6588.20; Thu, 13 Jul 2023 10:54:06 +0000 Received: from SG2PR06MB5430.apcprd06.prod.outlook.com ([fe80::e9fc:a1ea:cfed:90b8]) by SG2PR06MB5430.apcprd06.prod.outlook.com ([fe80::e9fc:a1ea:cfed:90b8%6]) with mapi id 15.20.6565.028; Thu, 13 Jul 2023 10:54:06 +0000 Message-ID: Date: Thu, 13 Jul 2023 18:54:00 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.15; rv:102.0) Gecko/20100101 Thunderbird/102.13.0 Subject: [PATCH, OpenACC 2.7, v2] Implement host_data must have use_device clause requirement Content-Language: en-US To: Thomas Schwinge , Chung-Lin Tang Cc: gcc-patches@gcc.gnu.org, Catherine Moore References: <3be2222f-48ae-12a1-a83b-415360e0a506@siemens.com> <87legjepn1.fsf@euler.schwinge.homeip.net> In-Reply-To: <87legjepn1.fsf@euler.schwinge.homeip.net> X-ClientProxiedBy: TYAPR01CA0200.jpnprd01.prod.outlook.com (2603:1096:404:29::20) To SG2PR06MB5430.apcprd06.prod.outlook.com (2603:1096:4:1ba::14) MIME-Version: 1.0 X-MS-PublicTrafficType: Email X-MS-TrafficTypeDiagnostic: SG2PR06MB5430:EE_|SEZPR06MB5479:EE_ X-MS-Office365-Filtering-Correlation-Id: 879d64fe-33d7-4dd3-1ee7-08db838f79e4 X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; X-Microsoft-Antispam-Message-Info: 781PUVEHXrto8nhRE6xBKlkMuINBJwJ2omkf5PwbHj0lKo1n/6aGOiHXfRSRK4h8oivK9UtzOs236sJFbVBD8sltG2UB7V+tNbW4NP7kB5bRLrvDjc+RJg6wB2lWa7k4f45PI0Q+7JTmv4Hh4bw/4EiyyMbyyVRJuizrCevdX+qhx9f2+GFq6oZVRXPM6msaJGFoGXPH0Vi7+sq4lpb7HzIC9VvZBhCnOC10N62XJMpPr9drjZTTUoJHB9aY7Nlf6VZif+CHJM8htCpLCED1YAp0+oZ6a6uDiy3JVZAxhBMy1aPSKojDoCsjSoyw91VpPijnSmqZmoR1uZF/WQGGoTz6HOwRY1zVLL/R6OhT9ZdbwO2kGGAMPss9Dx3XeyOY5pRGx0o2ngHMYOsoEsODxLv6dcoWatPEMO6+13HSms3m6Ag7+T2ubH+hnuUZVnD4SXeFaqGlj1oZKQZc09u7lzkspKFc2yDBDiErHpbSRe6nM3kWdTIl6QT6k85H7cujkk6zvBE+BdzX3k4KzHe5858ZYOxx9A+w+3MFK03h74OdzztIXukH4OTSisiiGlp4rojRB2YUQyPG9upHUZXLqUNYWhRDjA8IMIGs82tGMIlCGNlYY2OAn6jqdiNRzCK+M5I0mxva714exEcsMqbnQA== X-Forefront-Antispam-Report: CIP:255.255.255.255; CTRY:; LANG:en; SCL:1; SRV:; IPV:NLI; SFV:NSPM; H:SG2PR06MB5430.apcprd06.prod.outlook.com; PTR:; CAT:NONE; SFS:(13230028)(4636009)(396003)(39860400002)(136003)(376002)(366004)(346002)(451199021)(41300700001)(8676002)(6512007)(316002)(8936002)(66476007)(110136005)(33964004)(31686004)(66946007)(6486002)(66556008)(4326008)(478600001)(6506007)(53546011)(6666004)(83380400001)(84970400001)(38100700002)(82960400001)(31696002)(86362001)(66899021)(36756003)(2906002)(26005)(5660300002)(235185007)(2616005)(186003)(45980500001)(43740500002); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: =?utf-8?q?xp9+dQLc2ch9UOOP0r5Z0bSf/Cju?= =?utf-8?q?FTYWl6/pLAf2gVgFfSPWE3w6D9Ozq5bTqTolzQCm0SZRt1fn4xl+WKcu85P2T68Hs?= =?utf-8?q?EkIXEqv1HmFEMIWOBMOAvV+NggEdXic2tCQO/oHsbeFfkBZsoKLXtku0GBdzfM9hG?= =?utf-8?q?PPf7ErMAGh9cOT5GTCnSTyNt/yR+dODlKIddOP9o1Y2x1/nH3Zvsphw1kRxjv2Ziu?= =?utf-8?q?1DayGJGnqpeNfnj8DXrmSg1UQ5hh/129Unkb4fEyc+z7V6YDUt17Dn9Jr3TcmbxT6?= =?utf-8?q?uW9TsSp0r2RAtfIwTLDNNTyNcMGq9ZbA5VfU+TW+CPWlL7bjJU8EEi+fwl9BB3Jf5?= =?utf-8?q?yySf1Xy0jTzvS4pqUmxPZQwChyv45F3K5VDETDU6N4JsEDA56R7ziy/7H2dx3NE35?= =?utf-8?q?MqQAvTI/2Cweos4yMpFsfGpjiBK3lgKX6No0FCiqj/Y+0e+wPeAwmFlbRWPPn3SOE?= =?utf-8?q?C9MZRDByHGjDFWJ9Xhry0MUZ+/DUDZ580LSTIEN3LcIGPevrI2mLd3j6HOD0D4ept?= =?utf-8?q?z7w5yx55FfCo2U/RjQ2YE5XzgHCMmjdcbcGVhUfXF5fZFcFZuc9oIbBJ3OC3GPJJE?= =?utf-8?q?LBtozF5pgE8KDesPWp5mm39hZByv1VztnG/PX4Btoi+D+0Q7rhvv5fDaBTPa2Y24W?= =?utf-8?q?KpwaXttgCxleEhPO0YD604uvM7TT3a4t4suTtYT1GCwk3whmoEIw+irJmEajjgFFG?= =?utf-8?q?agxQphct3ajyjjgBWROSKkhyvMhjFydtszyBWT95Ho/Ca4zrjs8AIry8bk5URvSkZ?= =?utf-8?q?l1Y5h+P2dHCQ/vdjcwpp86wxSgGrqDtw8l6rKmS6Ll8JyxGDltVAVawygpNus8sRy?= =?utf-8?q?UTKPAZIXTpm9MTcy+lqaxVwB5MzER6KGM5nsPyPSKy2xQQWy9mPf/OSnX54g66s8w?= =?utf-8?q?dgTMrnm6L5P8teNGzgwAQx+459iYbqdukan9vThMAetst4UFbGa4OIrINQ18w4+lC?= =?utf-8?q?bRLG7I95SqwpXD1aIJms6XGVCDPI0bzx9xFe32+TT7BMdPn1/f/CmXfHYtyG4AbM+?= =?utf-8?q?LCpDbGtka3AFr2eZIuHkV0C1HpIKctKCOUWr8ZzTh9s5afxuRNOzaAtLAESgDv5oc?= =?utf-8?q?zpFH9RRWc7Tzqei5t7zvk9Knf2iXD8UL+68pvqrtFRByhaj7SeCV8LjcUiBVi2aRR?= =?utf-8?q?t/NDr7rEuKsjYey5KB0EJ0YI/3uAJHZ2Dsc4V5eTGbVB68JtOpfyvWUBNkzYT1SQ9?= =?utf-8?q?nR6+YsO5xl2zuKf6F7Qh1pI5HCtN8y+mpx+7CcWqM/CYSeVOZKvmgXhQZxNdB1YQF?= =?utf-8?q?RlVvx+hPi1FUDbZEDUbRlia9xz6ityqvby5GN5Q5AQOuRZ6Y4LUuqW4/hbwf1lln8?= =?utf-8?q?/53pW1wxqLZ9JOFsgI1fLizML4t7dM3uALhOlPEtKNvhuQdqjLOGl75Fk6etdRVo7?= =?utf-8?q?h6fK3RMsw5v5XJu9LJj64ANyETqo8pWVGXnfi4ws7ma2qiF2SVubSzc7lisLZz112?= =?utf-8?q?C/Pp85PLK96wCpdzH7O6ICGz3CA5K+zHV4o0knAKIuEgojj+7h9kxdtcRzUk4hbSn?= =?utf-8?q?kDnOGK0Cm+3Qa9/6N+ggoMsvKo+wbRyOcQ=3D=3D?= X-OriginatorOrg: siemens.com X-MS-Exchange-CrossTenant-Network-Message-Id: 879d64fe-33d7-4dd3-1ee7-08db838f79e4 X-MS-Exchange-CrossTenant-AuthSource: SG2PR06MB5430.apcprd06.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Internal X-MS-Exchange-CrossTenant-OriginalArrivalTime: 13 Jul 2023 10:54:05.4166 (UTC) X-MS-Exchange-CrossTenant-FromEntityHeader: Hosted X-MS-Exchange-CrossTenant-Id: 38ae3bcd-9579-4fd4-adda-b42e1495d55a X-MS-Exchange-CrossTenant-MailboxType: HOSTED X-MS-Exchange-CrossTenant-UserPrincipalName: YJ0JPuc8faj6cVn8fRAjouTU9lk/7ksmX1DdZec+535w0bePEUU4BGJpkG9A8mBc8dFd6Jwk0orsZxMlmdq3tVoR0vTcr9mXglcV0KxrSaw= X-MS-Exchange-Transport-CrossTenantHeadersStamped: SEZPR06MB5479 X-Spam-Status: No, score=-10.0 required=5.0 tests=BAYES_00, DKIMWL_WL_MED, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, FORGED_SPF_HELO, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, RCVD_IN_MSPIKE_H2, SPF_HELO_PASS, SPF_NONE, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , X-Patchwork-Original-From: Chung-Lin Tang via Gcc-patches From: Chung-Lin Tang Reply-To: Chung-Lin Tang Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1767966564356427374 X-GMAIL-MSGID: 1771302497033540880 On 2023/6/16 5:13 PM, Thomas Schwinge wrote: > OK with one small change, please -- unless there's a reason for doing it > this way: > >> --- a/gcc/fortran/trans-openmp.cc >> +++ b/gcc/fortran/trans-openmp.cc >> @@ -4677,6 +4677,12 @@ gfc_trans_oacc_construct (gfc_code *code) >> break; >> case EXEC_OACC_HOST_DATA: >> construct_code = OACC_HOST_DATA; >> + if (code->ext.omp_clauses->lists[OMP_LIST_USE_DEVICE] == NULL) >> + { >> + error_at (gfc_get_location (&code->loc), >> + "% construct requires % clause"); >> + return NULL_TREE; >> + } >> break; >> default: >> gcc_unreachable (); > The OpenMP "must contain at least one [...] clause" checks are done in > 'gcc/fortran/openmp.cc:resolve_omp_clauses'. For consistency (or, to let > 'gcc/fortran/trans-openmp.cc' continue to just deal with "directive > translation"), do similar for OpenACC 'host_data'? (..., and we later > accordingly adjust 'gcc/fortran/openmp.cc:gfc_match_oacc_update', too?) Hi Thomas, I've adjusted the Fortran implementation as you described. Yes, I agree this way more fits current Fortran FE conventions. I've re-tested the attached v2 patch, will commit later this week if no major objections. Thanks, Chung-Lin gcc/c/ChangeLog: * c-parser.cc (c_parser_oacc_host_data): Add checking requiring OpenACC host_data construct to have an use_device clause. gcc/cp/ChangeLog: * parser.cc (cp_parser_oacc_host_data): Add checking requiring OpenACC host_data construct to have an use_device clause. gcc/fortran/ChangeLog: * openmp.cc (resolve_omp_clauses): Add checking requiring OpenACC host_data construct to have an use_device clause. gcc/testsuite/ChangeLog: * c-c++-common/goacc/host_data-2.c: Adjust testcase. * gfortran.dg/goacc/host_data-error.f90: New testcase. * gfortran.dg/goacc/pr71704.f90: Adjust testcase. diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 24a6eb6e459..80920b31f83 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -18461,8 +18461,13 @@ c_parser_oacc_host_data (location_t loc, c_parser *parser, bool *if_p) tree stmt, clauses, block; clauses = c_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, - "#pragma acc host_data"); - + "#pragma acc host_data", false); + if (!omp_find_clause (clauses, OMP_CLAUSE_USE_DEVICE_PTR)) + { + error_at (loc, "% construct requires % clause"); + return error_mark_node; + } + clauses = c_finish_omp_clauses (clauses, C_ORT_ACC); block = c_begin_omp_parallel (); add_stmt (c_parser_omp_structured_block (parser, if_p)); stmt = c_finish_oacc_host_data (loc, clauses, block); diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 5e2b5cba57e..beb5b632e5e 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -45895,8 +45895,15 @@ cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) unsigned int save; clauses = cp_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, - "#pragma acc host_data", pragma_tok); - + "#pragma acc host_data", pragma_tok, + false); + if (!omp_find_clause (clauses, OMP_CLAUSE_USE_DEVICE_PTR)) + { + error_at (pragma_tok->location, + "% construct requires % clause"); + return error_mark_node; + } + clauses = finish_omp_clauses (clauses, C_ORT_ACC); block = begin_omp_parallel (); save = cp_parser_begin_omp_structured_block (parser); cp_parser_statement (parser, NULL_TREE, false, if_p); diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 8efc4b3ecfa..f7af02845de 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -8764,6 +8764,12 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, "% clause", &omp_clauses->detach->where); } + if (openacc + && code->op == EXEC_OACC_HOST_DATA + && omp_clauses->lists[OMP_LIST_USE_DEVICE] == NULL) + gfc_error ("% construct at %L requires % clause", + &code->loc); + if (omp_clauses->assume) gfc_resolve_omp_assumptions (omp_clauses->assume); } diff --git a/gcc/testsuite/c-c++-common/goacc/host_data-2.c b/gcc/testsuite/c-c++-common/goacc/host_data-2.c index b3093e575ff..862a764eb3a 100644 --- a/gcc/testsuite/c-c++-common/goacc/host_data-2.c +++ b/gcc/testsuite/c-c++-common/goacc/host_data-2.c @@ -8,7 +8,9 @@ void f (void) { int v2 = 3; -#pragma acc host_data copy(v2) /* { dg-error ".copy. is not valid for ..pragma acc host_data." } */ +#pragma acc host_data copy(v2) + /* { dg-error ".copy. is not valid for ..pragma acc host_data." "" { target *-*-* } .-1 } */ + /* { dg-error ".host_data. construct requires .use_device. clause" "" { target *-*-* } .-2 } */ ; #pragma acc host_data use_device(v2) @@ -20,6 +22,9 @@ f (void) /* { dg-error ".use_device_ptr. variable is neither a pointer nor an array" "" { target c } .-1 } */ /* { dg-error ".use_device_ptr. variable is neither a pointer, nor an array nor reference to pointer or array" "" { target c++ } .-2 } */ ; + +#pragma acc host_data /* { dg-error ".host_data. construct requires .use_device. clause" } */ + ; } diff --git a/gcc/testsuite/gfortran.dg/goacc/host_data-error.f90 b/gcc/testsuite/gfortran.dg/goacc/host_data-error.f90 new file mode 100644 index 00000000000..bd262989410 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/host_data-error.f90 @@ -0,0 +1,6 @@ +! { dg-do compile } + +subroutine foo () +!$acc host_data ! { dg-error "'host_data' construct at .1. requires 'use_device' clause" } +!$acc end host_data +end diff --git a/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 b/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 index 0235e85d42a..31724c8b046 100644 --- a/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 @@ -47,8 +47,9 @@ real function f8 () f8 = 1 end -real function f9 () -!$acc host_data +real function f9 (a) + integer a(:) +!$acc host_data use_device(a) !$acc end host_data f8 = 1 end