From patchwork Sun Sep 1 05:41:13 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Prathamesh Kulkarni X-Patchwork-Id: 1979374 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; dkim=pass (2048-bit key; unprotected) header.d=Nvidia.com header.i=@Nvidia.com header.a=rsa-sha256 header.s=selector2 header.b=DCb1xBSn; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4WxLNr41L6z1yZ9 for ; Sun, 1 Sep 2024 15:41:46 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id A51883858280 for ; Sun, 1 Sep 2024 05:41:42 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from NAM11-DM6-obe.outbound.protection.outlook.com (mail-dm6nam11on20603.outbound.protection.outlook.com [IPv6:2a01:111:f403:2415::603]) by sourceware.org (Postfix) with ESMTPS id F043E3858D34 for ; Sun, 1 Sep 2024 05:41:18 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org F043E3858D34 Authentication-Results: sourceware.org; dmarc=fail (p=reject dis=none) header.from=nvidia.com Authentication-Results: sourceware.org; spf=fail smtp.mailfrom=nvidia.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org F043E3858D34 Authentication-Results: server2.sourceware.org; arc=pass smtp.remote-ip=2a01:111:f403:2415::603 ARC-Seal: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1725169280; cv=pass; b=wKJNoMCFJ96wPTieU73Ha7k1ibxO1UFX+gAXc4OJ0liKEwFAff4hQ4g1G0sx69Bz2xgBoRljOgsL2jTgXxqhXMmt/+9CoGRjq5OEJ3JrKzC/zwfOGRFpFb5H4tGfMJL2PWA6uQj00ngmw7ptTpyYG3WuNwwFCbyHzC1ryWO/YOI= ARC-Message-Signature: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1725169280; c=relaxed/simple; bh=6ap5uvVih2aiX9WwffXRPWjN1+7Kd7bjchF/paw7bHw=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=Xv3UZorYJWVv6VMTebAG20wGEPxeYBEMjkEWgQLGw1r168MokEC3Nh2jm+JHHBBo8MwlfPDDy+J/O9h43JSRVOxIqM69cWW5tDzOXX4wBzGWLpCtMdtTCqk3r/8etgHulAZ0/Wy9zsY5SqBPSAJq1wgvdNjHvfrrBZEFVxWMj9k= ARC-Authentication-Results: i=2; server2.sourceware.org ARC-Seal: i=1; a=rsa-sha256; s=arcselector10001; d=microsoft.com; cv=none; b=ZbjC2hwOG5mi5WsI4XqMU+Tcan6c/UjJKsB+CPvcLBbPJkpaJixYDn7NHOHXOKI18VnebsP0JaUS7MbkzOc326MW3KwqGUrZYIgG6Vo7TJcw+1xMLoLm6fH3jW5TLytBCkhv2QuOTjuyq85GkEa906z7m+YZeWOdKQg3ipMTI4G+3ZbvgQKr0hrkHXAVhud3UrgDf7gQpsPkuglRQ5Sol4JGEYA92OZq2drSH0j4dKqXgn91HGqB2kO0kykdbpjTDoPKfd1+3JNGIKMMcuOvwpsbYZNHJUmN6OUJ96WOFn0pRXENMrCzp5B3SVSuMWyTLppyOYRUmEXtB5SEpPED1w== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=microsoft.com; s=arcselector10001; 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=inRY9mz+BS3lN7tDcwNU8axRTckQucbF3JlSTIIGHpU=; b=XXSH6NejfLPEwZ7OvKALboOaLjLBhUPoBOagd/PyOLE4oKvjnBLFGIQ2AGMjV2jyR2j4TIsU5YSmAGsr8bRkt6mbP8edrc+NztObWkSeiuYamiMx45a9+x9h1n6LTWrMhlB+O7Cpb+dh11GgGDA+UnrHa3Vj8LbgLGz9JbLRuEfNExS79tL1f0EaXPvMFiVEtNJ3AyvN8edQFSks7Pth5dItdvQcwfyen042kEGjsISIcc3nty1DYM+qcXOtFJUfgVqVsm9V222TQYO8BU4IP6UpYohShzc4TfnjRjMtrVEtnenORGDpG+ev2QbPZCotKjBzgX9JYgVn/jWTYCpIdw== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass smtp.mailfrom=nvidia.com; dmarc=pass action=none header.from=nvidia.com; dkim=pass header.d=nvidia.com; arc=none DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=Nvidia.com; s=selector2; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-SenderADCheck; bh=inRY9mz+BS3lN7tDcwNU8axRTckQucbF3JlSTIIGHpU=; b=DCb1xBSnxzK1ej01eLo5F6xP/TWsMzd2SZMxuEMkcssfcI3KfD5fkN/lQPi3mZ7ftKzgOHsQs7Z5WeF773GkKHA+DMwXXSVzaTDuEuKMMYfq0myR2JaOdsfEoSq68rsQAy4SaLR0u0L5TUZ20yvshMN3IEtpkVZbDVJaf2e3kpD8dWI1MvA25lH5EHE/xDc1iZvVeprPbHa3MQAZhVFV8BdS/EKSJcBMxz25p4gx2tqTLX8KUDypDrRUg23l+yHtOX6FC856kjrtROzc45sb2eCS1vS2PVJeJwsJdsB4r6BSTu0z6LS7H+Bcvjysqn+25ET/W7CpZhT0O8XuU+ss5A== Received: from IA1PR12MB9031.namprd12.prod.outlook.com (2603:10b6:208:3f9::19) by IA0PR12MB8893.namprd12.prod.outlook.com (2603:10b6:208:484::19) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.7897.26; Sun, 1 Sep 2024 05:41:14 +0000 Received: from IA1PR12MB9031.namprd12.prod.outlook.com ([fe80::1fb7:5076:77b5:559c]) by IA1PR12MB9031.namprd12.prod.outlook.com ([fe80::1fb7:5076:77b5:559c%3]) with mapi id 15.20.7918.020; Sun, 1 Sep 2024 05:41:13 +0000 From: Prathamesh Kulkarni To: "gcc-patches@gcc.gnu.org" , "rguenther@suse.de" Subject: [gimplify.cc] Avoid ICE when passing VLA vector to accelerator Thread-Topic: [gimplify.cc] Avoid ICE when passing VLA vector to accelerator Thread-Index: Adr8MQWqQ2fcBtS5Qb6YiEEpvK1zWQ== Date: Sun, 1 Sep 2024 05:41:13 +0000 Message-ID: Accept-Language: en-US Content-Language: en-US X-MS-Has-Attach: yes X-MS-TNEF-Correlator: authentication-results: dkim=none (message not signed) header.d=none;dmarc=none action=none header.from=nvidia.com; x-ms-publictraffictype: Email x-ms-traffictypediagnostic: IA1PR12MB9031:EE_|IA0PR12MB8893:EE_ x-ms-office365-filtering-correlation-id: 87c7ac11-b884-4b0d-e0df-08dcca48b0df x-ms-exchange-senderadcheck: 1 x-ms-exchange-antispam-relay: 0 x-microsoft-antispam: BCL:0; ARA:13230040|376014|1800799024|366016|38070700018; x-microsoft-antispam-message-info: =?iso-8859-1?q?Y9LgwTFvjWhKD73Ell/Zq9pQwE?= =?iso-8859-1?q?f4NOQmp9TrYde2knGa3PeZLugxGjTjK/i95wJO2/JjSdx77nCH/XloTM2TX7?= =?iso-8859-1?q?mpZbS9ISLBwRJ33H/FirfcfL2zPBzy9e2WFshJY4Sh7NgT517/JpmfWvuvSL?= =?iso-8859-1?q?aYy3ApbpdzzN3N63ER640926K+cjsHXM7LPbFtz2oOTBPm8ZfUryW9dJeLNI?= =?iso-8859-1?q?om1vHyRq2XrpjfqcWGn6HtSt9WGECg4IzXGmw6nn8KwYcX+n8RWE4GmLq8Hl?= =?iso-8859-1?q?kUDWtRyNAv7+5fbt6eVy2iqxZyWuFSuYvsV3dNR6ZRK0e36UQ5jf7y+VnauF?= =?iso-8859-1?q?93NzfAhKr5D8nFPr8OvZSBzG0a/nUNPZXuW5SPegBHu5UFqjpKqIh7AvyatI?= =?iso-8859-1?q?qpriuWROcRBKL05BFTQuoecxxEnSerkmSvlEdBrAOuXdoOJBswnqiSvlMj/I?= =?iso-8859-1?q?1f4aBmDOlTkEv9xtaA5YMBBrx3uRcA7n8TIw/jIto2XJjiyLj7Uotib9RfW/?= =?iso-8859-1?q?UvJ8CkfHR7qFNJm7hcen4pZtuD1HBKtElf78vkO1auR+ggtVkg1TwfmZdefh?= =?iso-8859-1?q?YAfYgPMoCpdOt9SOUdgpRpsoQxtUYrtZcZelAxBNRYOeoEvSXFvsfqobGNWO?= =?iso-8859-1?q?LEJbhbY3V1FPv5yHo12OuMvy0iXh8w8Ph9scQTSKWs/p8wTP/mLF6E9xtb6g?= =?iso-8859-1?q?9gK2Xh1zBaQfUVJoTBwPZu2DO7S9MQDegE/5JqI7+H7gXHfHc57rC2MMA6Oi?= =?iso-8859-1?q?v+/PLNbNjUrpGCn/nQelz2Uuhtz8KfrgOSU8W/siM/GhQRfNE6qwdjw7WD/S?= =?iso-8859-1?q?zplbRvFQQT2HKVVc7Zu5c0vYSLwlxTb4tAQbaQYzrZ3N6irxSL047k3uW6lG?= =?iso-8859-1?q?uU7sDtuI5U/P9PFBgN/emAjPHfY0CxYoJHUYUxhax2OUhaVAhnyvlMJ9cXZh?= =?iso-8859-1?q?g4Tage67E17b+ZlYx2u8Hx2fLA48Ck/CjMvU1Y9nqD7hBA0U2jZkl826gPmb?= =?iso-8859-1?q?wENT82AbXtDfCTDJX0tb24+BIZXkyv8Cn5ti1RAo84WBCwRZbYwTrw9iqUTZ?= =?iso-8859-1?q?I2gEqlOD7QbOfUgRE7l2bONniVioYeTAXLvp5bDg2uzDAxcQR6tPao0Z+vEl?= =?iso-8859-1?q?Uudha7pnY16zTBmuCpG32l4+bjDz7oHnY8cC9vH3ky2X1LPouHiGaUKzkRt7?= =?iso-8859-1?q?2jSd7WWn8LS+AvYLto4841foE080R43RhKlDnWsZELQU80DmPn23pFv+Qg/4?= =?iso-8859-1?q?A3xX4bO8HZmn90fDeZ+/u1VsifbbDFhF+xdpHIK6yO/AiJx7aAdzsDbPRPvx?= =?iso-8859-1?q?sRzZjid8HsY+mKbbd9Qf6PmTtB72vR5e0WkpOEjc1ZyO3vg1Zyhq2eCcfTC4?= =?iso-8859-1?q?Xj5C9+hwv9JDN5gB+8wfuBmOYEMjd2qfOBY4TXKGTljlbOtMcMbS3h2cqnhO?= =?iso-8859-1?q?6pshlBTPWB3qN2XEaYLmtP6Q=3D=3D?= x-forefront-antispam-report: CIP:255.255.255.255; CTRY:; LANG:en; SCL:1; SRV:; IPV:NLI; SFV:NSPM; H:IA1PR12MB9031.namprd12.prod.outlook.com; PTR:; CAT:NONE; SFS:(13230040)(376014)(1800799024)(366016)(38070700018); DIR:OUT; SFP:1101; x-ms-exchange-antispam-messagedata-chunkcount: 1 x-ms-exchange-antispam-messagedata-0: =?iso-8859-1?q?AqCXkEyGIWgERZTpEMYlyfY?= =?iso-8859-1?q?jMTJPMoVjrajiazFdXvUyFz2R3L0DRHFiZy9vaZSc71mjek02bTvcW6BumZp?= =?iso-8859-1?q?Qj0Xmi/HWNjPC5iB0Ch6URv/SY0Dr1/jxk5p/9n+b6EL7aRIEfTs8xiTkAya?= =?iso-8859-1?q?BSvp/mvWtnT4r5zBU6M00GHJjEXLjCG1Kggy/9C4XlBil05Ewk2+qi20z5f4?= =?iso-8859-1?q?CGpQ+IEc9oHSg5w5oK7UgalA59Y1EQ3BGhXTL0W8Yv0nZ9L+jEwIE4nA0WZI?= =?iso-8859-1?q?xYawzfJGIkTcRt17L9l+o2h1h44OfVt7yIOHhovbWebnzE4JtgIiCyPFk4AH?= =?iso-8859-1?q?ef05cL66A9GeF4DNQjL5QLEYBN6NvnuVi2vNhKohRrBsKqTQl2caqFYbFzRc?= =?iso-8859-1?q?3D+oCPuZTymCT+U/EA2lPPh/l1u5dAqE+JoftOprjy9LbXQ5LQCz1ugstNh4?= =?iso-8859-1?q?fMsGEZ7EensqSLHjnURCz0aulwoINan4/a38Nkr5OZG1sHSBe/uqW7lUwI7j?= =?iso-8859-1?q?yZVeZFU8kYeOfcTg7Ibt2tSz09E+XAsXk+AY2E672wUgSuit1AQvLXVBOHOE?= =?iso-8859-1?q?PlEO1jXMuTAnS41T/P+ijxeLlGA/8KknVuHNBB7R0bfq5Z3lhoRchdokdUpa?= =?iso-8859-1?q?SSbCQ0ET1o3vd4x4dvdd7bK2yr/iRbyQKofR8RHFoAWwV/dkMOGee1zWWIId?= =?iso-8859-1?q?fwK6V4EcTiQkYbC3BXJV/lid7iUIprj9UueI3HCSBnzGq3BSb+cG9+eCCazh?= =?iso-8859-1?q?p1MaOncp9l71Fn/q6BAXIo2AOx1nypytckryrlzoQKeQIjgCTe9Dpr5RqML2?= =?iso-8859-1?q?5jTCTnoS1wzkziiglnJdPXjslBhf/y45ojQtliotw/sprLk5lz0jRDj5iHRG?= =?iso-8859-1?q?Tn1hB5IlzlPnQZofRVJgL0OvuzfgsWl+GC4TpaJzE/lauy/FIFAjS52obnoL?= =?iso-8859-1?q?QY+ybSXffADW+zTs3X9muU744pPl9J94U0C2gda42qVOumCmqapBy7lqQoXv?= =?iso-8859-1?q?KXyk7toL5mejf0/AYT5T+Hvhxk5J1u03mhyrnjA1aR+ByPu/gM4JVJasfBqg?= =?iso-8859-1?q?2Oe5dlp3bd1v2jGg3I1oW6nR+1uZdvMFoQ8h09+KPdRU/BiuDtoBzwPpekQU?= =?iso-8859-1?q?NlhQf6xJILgCoNdquKdaIHg6fEIIAyy5cZ/7fPpYN1eIxGFMtwZIgWZiUdbR?= =?iso-8859-1?q?L0gsdFk0qx9vmNIL/vVC7cT514rjlMp8UZfBEHgVCJ7WSz8IuoutVlrFnow0?= =?iso-8859-1?q?pYTf6ZjH5PODLkrCjcwGccM+ZRebdeY5WekHZIHTWqOSOfpl/rtJabj202at?= =?iso-8859-1?q?bSPLypt3rEtwigHqH5KAqYCs/jlmTigTlkP2mtC/xCOrL1MRFkx5dJ7hZLAL?= =?iso-8859-1?q?GIq2MkGXo4aYy4bWpfNq5v4Bcva3WMG8z+6gwE0Ut5wOJnYQykIxv8cp7Ixz?= =?iso-8859-1?q?VMrXjJyLps2/+GXfDbVGLY+sxXVRJ/z0DcZFJwmu323/YokXDw6xyEBmp8gW?= =?iso-8859-1?q?9muRVpQmT88GYQIn36IuXO4+ghxDrwG6ba2D0MJqggERZCbD8lyGHUaAmxmJ?= =?iso-8859-1?q?9tTIaGGuyKJT90EBjTjuJtzJztcCr/At1NpuJ22cXgcm6oDlII2RF9diolNC?= =?iso-8859-1?q?8zD7VpHKWbKk5i9nY?= MIME-Version: 1.0 X-OriginatorOrg: Nvidia.com X-MS-Exchange-CrossTenant-AuthAs: Internal X-MS-Exchange-CrossTenant-AuthSource: IA1PR12MB9031.namprd12.prod.outlook.com X-MS-Exchange-CrossTenant-Network-Message-Id: 87c7ac11-b884-4b0d-e0df-08dcca48b0df X-MS-Exchange-CrossTenant-originalarrivaltime: 01 Sep 2024 05:41:13.3583 (UTC) X-MS-Exchange-CrossTenant-fromentityheader: Hosted X-MS-Exchange-CrossTenant-id: 43083d15-7273-40c1-b7db-39efd9ccc17a X-MS-Exchange-CrossTenant-mailboxtype: HOSTED X-MS-Exchange-CrossTenant-userprincipalname: mjrH+aQvlHUyXIY/so5tzTjShDtZWgn9KtT/BhqaqlOqEUeFJs3KvOw3tXbu1aCoM571vEJmL6lDJuCUuIWK/w== X-MS-Exchange-Transport-CrossTenantHeadersStamped: IA0PR12MB8893 X-Spam-Status: No, score=-10.4 required=5.0 tests=BAYES_00, DKIMWL_WL_HIGH, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, FORGED_SPF_HELO, GIT_PATCH_0, SPF_HELO_PASS, SPF_NONE, TXREP, T_SCC_BODY_TEXT_LINE, WEIRD_PORT 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.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org Hi, For the following test: #include int main() {   svint32_t x;   #pragma omp target map(x)     x;   return 0; } compiling with -fopenmp -foffload=nvptx-none results in following ICE: t_sve.c: In function 'main': t_sve.c:6:11: internal compiler error: Segmentation fault     6 |   #pragma omp target map(x)       |           ^~~ 0x228ed13 internal_error(char const*, ...)         ../../gcc/gcc/diagnostic-global-context.cc:491 0xfcf68f crash_signal         ../../gcc/gcc/toplev.cc:321 0xc17d9c omp_add_variable         ../../gcc/gcc/gimplify.cc:7811 0xc17d9c omp_add_variable         ../../gcc/gcc/gimplify.cc:7752 0xc4176b gimplify_scan_omp_clauses         ../../gcc/gcc/gimplify.cc:12881 0xc46d53 gimplify_omp_workshare         ../../gcc/gcc/gimplify.cc:17139 0xc23383 gimplify_expr(tree_node**, gimple**, gimple**, bool (*)(tree_node*), int)         ../../gcc/gcc/gimplify.cc:18668 0xc27f53 gimplify_stmt(tree_node**, gimple**)         ../../gcc/gcc/gimplify.cc:7646 0xc24ef7 gimplify_statement_list         ../../gcc/gcc/gimplify.cc:2250 0xc24ef7 gimplify_expr(tree_node**, gimple**, gimple**, bool (*)(tree_node*), int)         ../../gcc/gcc/gimplify.cc:18565 0xc27f53 gimplify_stmt(tree_node**, gimple**)         ../../gcc/gcc/gimplify.cc:7646 0xc289d3 gimplify_bind_expr         ../../gcc/gcc/gimplify.cc:1642 0xc24b9b gimplify_expr(tree_node**, gimple**, gimple**, bool (*)(tree_node*), int)         ../../gcc/gcc/gimplify.cc:18315 0xc27f53 gimplify_stmt(tree_node**, gimple**)         ../../gcc/gcc/gimplify.cc:7646 0xc24ef7 gimplify_statement_list         ../../gcc/gcc/gimplify.cc:2250 0xc24ef7 gimplify_expr(tree_node**, gimple**, gimple**, bool (*)(tree_node*), int)         ../../gcc/gcc/gimplify.cc:18565 0xc27f53 gimplify_stmt(tree_node**, gimple**)         ../../gcc/gcc/gimplify.cc:7646 0xc2aadb gimplify_body(tree_node*, bool)         ../../gcc/gcc/gimplify.cc:19393 0xc2b05f gimplify_function_tree(tree_node*)         ../../gcc/gcc/gimplify.cc:19594 0xa0e47f cgraph_node::analyze()         ../../gcc/gcc/cgraphunit.cc:687 The attached patch fixes the issue by checking if variable is VLA vector, and emits an error in that case since no accel currently supports VLA vectors. Does the patch look OK ? Signed-off-by: Prathamesh Kulkarni Thanks, Prathamesh [gimplify.cc] Emit an error if VLA vector is passed to accelerator. gcc/ChangeLog: * gimplify.cc (omp_add_variable): Emit an error if VLA vector is passed to accelerator. Signed-off-by: Prathamesh Kulkarni Signed-off-by: Prathamesh Kulkarni Signed-off-by: Prathamesh Kulkarni Signed-off-by: Prathamesh Kulkarni diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 26a216e151d..fb7bd919b54 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -7789,6 +7789,11 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags) the parameters of the type. */ if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { + /* For now, bail out if a VLA vector is passed to accelerator. */ + if (VECTOR_TYPE_P (TREE_TYPE (decl)) + && POLY_INT_CST_P (DECL_SIZE (decl))) + fatal_error (input_location, + "passing VLA vector to accelerator not implemented"); /* Add the pointer replacement variable as PRIVATE if the variable replacement is private, else FIRSTPRIVATE since we'll need the address of the original variable either for SHARED, or for the