From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id B98D53858D37 for ; Mon, 21 Feb 2022 15:19:11 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org B98D53858D37 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: VE4f0W05xmtvL11Rj7688HkwD2DUcTRZDhf64j/vtoeYGiOJeX+O9cwDJEkm/z/BPqqCrsZS9i Wrzkhl9sez5VXSLkdAmc5mMaGjmq4PG8VzXr63L1LIbUBDVKBS605TbIfFPiOyWmo1dxEsRqzC JbTix7gDMvXMbLbwhX8IUAfcwIHHL8RcfuNBBk3boAWTV7Mpj6b1ipxIEw7GWvsnZbJQxLZP7R r9PwwyHV0s1TanuPnqm3VQnuPgm0t7VJ6Q56AJTIaa3KIYh26dzxNWM7zeZuwZczUVG6CbEdeO RCMAk3COs1ryXOMEmH9PW48w X-IronPort-AV: E=Sophos;i="5.88,386,1635235200"; d="scan'208";a="72094090" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 21 Feb 2022 07:19:11 -0800 IronPort-SDR: Mac7oo6IRIJgM0uDAec60Z1ZbgaF4HFcdM03sKoZJXlBXYP0iLoHxxJ1kWkKanJ2TM5Vo7RvtX vbOLqFrVKh+sA+7dXX+31T45wEEmEX3+pAiBp5gffYD2bBvuQu+7ImBX0fEBkE9INmcBb0md6S kgJvJ9zkzMvGPpXXx7Vuv2IXkNFAxRb9kxN84KloU4CKcoB8BKImhb7BHvr7/nl/V5te/Lac4b ugz9stlLGMbMBtRDhVdHsqe2woJjwQ/HJ27Hm2wbAM9omhg8aLnwi1W/rFZEYyjiMebneM/X5A pAw= Content-Type: multipart/mixed; boundary="------------lsx2S7f8WeAM3vY0W8DOy3i7" Message-ID: <65e1462e-d17f-2975-1401-358fe9c69e28@codesourcery.com> Date: Mon, 21 Feb 2022 23:18:57 +0800 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:91.0) Gecko/20100101 Thunderbird/91.6.0 Content-Language: en-US To: gcc-patches , Jakub Jelinek From: Chung-Lin Tang Subject: [PATCH, OpenMP, C/C++] Handle array reference base-pointers in array sections X-ClientProxiedBy: svr-orw-mbx-03.mgc.mentorg.com (147.34.90.203) To svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) X-Spam-Status: No, score=-10.5 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) 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-List-Received-Date: Mon, 21 Feb 2022 15:19:14 -0000 --------------lsx2S7f8WeAM3vY0W8DOy3i7 Content-Type: text/plain; charset="UTF-8"; format=flowed Content-Transfer-Encoding: 7bit Hi Jakub, as encountered in cases where a program constructs its own deep-copying for arrays-of-pointers, e.g: #pragma omp target enter data map(to:level->vectors[:N]) for (i = 0; i < N; i++) #pragma omp target enter data map(to:level->vectors[i][:N]) We need to treat the part of the array reference before the array section as a base-pointer (here 'level->vectors[i]'), providing pointer-attachment behavior. This patch adds this inside handle_omp_array_sections(), tracing the whole sequence of array dimensions, creating a whole base-pointer reference iteratively using build_array_ref(). The conditions are that each of the "absorbed" dimensions must be length==1, and the final reference must be of pointer-type (so that pointer attachment makes sense). There's also a little patch in gimplify_scan_omp_clauses(), to make sure the array-ref base-pointer goes down the right path. This case was encountered when working to make 534.hpgmgfv_t from SPEChpc 2021 properly compile. Tested without regressions on trunk. Okay to go in once stage1 opens? Thanks, Chung-Lin 2022-02-21 Chung-Lin Tang gcc/c/ChangeLog: * c-typeck.cc (handle_omp_array_sections): Add handling for creating array-reference base-pointer attachment clause. gcc/cp/ChangeLog: * semantics.cc (handle_omp_array_sections): Add handling for creating array-reference base-pointer attachment clause. gcc/ChangeLog: * gimplify.cc (gimplify_scan_omp_clauses): Add case for attach/detach map kind for ARRAY_REF of POINTER_TYPE. gcc/testsuite/ChangeLog: * c-c++-common/gomp/target-enter-data-1.c: Adjust testcase. libgomp/testsuite/ChangeLog: * libgomp.c-c++-common/ptr-attach-2.c: New test. --------------lsx2S7f8WeAM3vY0W8DOy3i7 Content-Type: text/plain; charset="UTF-8"; name="arrayref-bptr.patch" Content-Disposition: attachment; filename="arrayref-bptr.patch" Content-Transfer-Encoding: base64 ZGlmZiAtLWdpdCBhL2djYy9jL2MtdHlwZWNrLmNjIGIvZ2NjL2MvYy10eXBlY2suY2MKaW5k ZXggMzA3NWM4ODM1NDguLjQyNTdlMzczNTU3IDEwMDY0NAotLS0gYS9nY2MvYy9jLXR5cGVj ay5jYworKysgYi9nY2MvYy9jLXR5cGVjay5jYwpAQCAtMTM2NDksNiArMTM2NDksMTAgQEAg aGFuZGxlX29tcF9hcnJheV9zZWN0aW9ucyAodHJlZSBjLCBlbnVtIGNfb21wX3JlZ2lvbl90 eXBlIG9ydCkKICAgICAgIGlmIChpbnRfc2l6ZV9pbl9ieXRlcyAoVFJFRV9UWVBFIChmaXJz dCkpIDw9IDApCiAJbWF5YmVfemVyb19sZW4gPSB0cnVlOwogCisgICAgICBzdHJ1Y3QgZGlt IHsgdHJlZSBsb3dfYm91bmQsIGxlbmd0aDsgfTsKKyAgICAgIGF1dG9fdmVjPGRpbT4gZGlt cyAobnVtKTsKKyAgICAgIGRpbXMuc2FmZV9ncm93IChudW0pOworCiAgICAgICBmb3IgKGkg PSBudW0sIHQgPSBPTVBfQ0xBVVNFX0RFQ0wgKGMpOyBpID4gMDsKIAkgICB0ID0gVFJFRV9D SEFJTiAodCkpCiAJewpAQCAtMTM3NjMsNiArMTM3NjcsOSBAQCBoYW5kbGVfb21wX2FycmF5 X3NlY3Rpb25zICh0cmVlIGMsIGVudW0gY19vbXBfcmVnaW9uX3R5cGUgb3J0KQogCSAgICAg IGVsc2UKIAkJc2l6ZSA9IHNpemVfYmlub3AgKE1VTFRfRVhQUiwgc2l6ZSwgbCk7CiAJICAg IH0KKworCSAgZGltIGQgPSB7IGxvd19ib3VuZCwgbGVuZ3RoIH07CisJICBkaW1zW2ldID0g ZDsKIAl9CiAgICAgICBpZiAoc2lkZV9lZmZlY3RzKQogCXNpemUgPSBidWlsZDIgKENPTVBP VU5EX0VYUFIsIHNpemV0eXBlLCBzaWRlX2VmZmVjdHMsIHNpemUpOwpAQCAtMTM4MDIsNiAr MTM4MDksMjMgQEAgaGFuZGxlX29tcF9hcnJheV9zZWN0aW9ucyAodHJlZSBjLCBlbnVtIGNf b21wX3JlZ2lvbl90eXBlIG9ydCkKIAkgIE9NUF9DTEFVU0VfREVDTCAoYykgPSB0OwogCSAg cmV0dXJuIGZhbHNlOwogCX0KKworICAgICAgdHJlZSBhcmVmID0gdDsKKyAgICAgIGZvciAo aSA9IDA7IGkgPCBkaW1zLmxlbmd0aCAoKTsgaSsrKQorCXsKKwkgIGlmIChkaW1zW2ldLmxl bmd0aCAmJiBpbnRlZ2VyX29uZXAgKGRpbXNbaV0ubGVuZ3RoKSkKKwkgICAgeworCSAgICAg IHRyZWUgbGIgPSBkaW1zW2ldLmxvd19ib3VuZDsKKwkgICAgICBhcmVmID0gYnVpbGRfYXJy YXlfcmVmIChPTVBfQ0xBVVNFX0xPQ0FUSU9OIChjKSwgYXJlZiwgbGIpOworCSAgICB9CisJ ICBlbHNlCisJICAgIHsKKwkgICAgICBpZiAoVFJFRV9DT0RFIChUUkVFX1RZUEUgKGFyZWYp KSA9PSBQT0lOVEVSX1RZUEUpCisJCXQgPSBhcmVmOworCSAgICAgIGJyZWFrOworCSAgICB9 CisJfQorCiAgICAgICBmaXJzdCA9IGNfZnVsbHlfZm9sZCAoZmlyc3QsIGZhbHNlLCBOVUxM KTsKICAgICAgIE9NUF9DTEFVU0VfREVDTCAoYykgPSBmaXJzdDsKICAgICAgIGlmIChPTVBf Q0xBVVNFX0NPREUgKGMpID09IE9NUF9DTEFVU0VfSEFTX0RFVklDRV9BRERSKQpAQCAtMTM4 MzYsNyArMTM4NjAsOCBAQCBoYW5kbGVfb21wX2FycmF5X3NlY3Rpb25zICh0cmVlIGMsIGVu dW0gY19vbXBfcmVnaW9uX3R5cGUgb3J0KQogCSAgYnJlYWs7CiAJfQogICAgICAgdHJlZSBj MiA9IGJ1aWxkX29tcF9jbGF1c2UgKE9NUF9DTEFVU0VfTE9DQVRJT04gKGMpLCBPTVBfQ0xB VVNFX01BUCk7Ci0gICAgICBpZiAoVFJFRV9DT0RFICh0KSA9PSBDT01QT05FTlRfUkVGKQor ICAgICAgaWYgKFRSRUVfQ09ERSAodCkgPT0gQ09NUE9ORU5UX1JFRiB8fCBUUkVFX0NPREUg KHQpID09IEFSUkFZX1JFRgorCSAgfHwgVFJFRV9DT0RFICh0KSA9PSBJTkRJUkVDVF9SRUYp CiAJT01QX0NMQVVTRV9TRVRfTUFQX0tJTkQgKGMyLCBHT01QX01BUF9BVFRBQ0hfREVUQUNI KTsKICAgICAgIGVsc2UKIAlPTVBfQ0xBVVNFX1NFVF9NQVBfS0lORCAoYzIsIEdPTVBfTUFQ X0ZJUlNUUFJJVkFURV9QT0lOVEVSKTsKZGlmZiAtLWdpdCBhL2djYy9jcC9zZW1hbnRpY3Mu Y2MgYi9nY2MvY3Avc2VtYW50aWNzLmNjCmluZGV4IDBjYjE3YTZhOGFiLi42NDZmNDg4M2Q2 NiAxMDA2NDQKLS0tIGEvZ2NjL2NwL3NlbWFudGljcy5jYworKysgYi9nY2MvY3Avc2VtYW50 aWNzLmNjCkBAIC01NDk3LDYgKzU0OTcsMTAgQEAgaGFuZGxlX29tcF9hcnJheV9zZWN0aW9u cyAodHJlZSBjLCBlbnVtIGNfb21wX3JlZ2lvbl90eXBlIG9ydCkKICAgICAgIGlmIChwcm9j ZXNzaW5nX3RlbXBsYXRlX2RlY2wgJiYgbWF5YmVfemVyb19sZW4pCiAJcmV0dXJuIGZhbHNl OwogCisgICAgICBzdHJ1Y3QgZGltIHsgdHJlZSBsb3dfYm91bmQsIGxlbmd0aDsgfTsKKyAg ICAgIGF1dG9fdmVjPGRpbT4gZGltcyAobnVtKTsKKyAgICAgIGRpbXMuc2FmZV9ncm93IChu dW0pOworCiAgICAgICBmb3IgKGkgPSBudW0sIHQgPSBPTVBfQ0xBVVNFX0RFQ0wgKGMpOyBp ID4gMDsKIAkgICB0ID0gVFJFRV9DSEFJTiAodCkpCiAJewpAQCAtNTYwNCw2ICs1NjA4LDkg QEAgaGFuZGxlX29tcF9hcnJheV9zZWN0aW9ucyAodHJlZSBjLCBlbnVtIGNfb21wX3JlZ2lv bl90eXBlIG9ydCkKIAkgICAgICBlbHNlCiAJCXNpemUgPSBzaXplX2Jpbm9wIChNVUxUX0VY UFIsIHNpemUsIGwpOwogCSAgICB9CisKKwkgIGRpbSBkID0geyBsb3dfYm91bmQsIGxlbmd0 aCB9OworCSAgZGltc1tpXSA9IGQ7CiAJfQogICAgICAgaWYgKCFwcm9jZXNzaW5nX3RlbXBs YXRlX2RlY2wpCiAJewpAQCAtNTY0Nyw2ICs1NjU0LDI0IEBAIGhhbmRsZV9vbXBfYXJyYXlf c2VjdGlvbnMgKHRyZWUgYywgZW51bSBjX29tcF9yZWdpb25fdHlwZSBvcnQpCiAJICAgICAg T01QX0NMQVVTRV9ERUNMIChjKSA9IHQ7CiAJICAgICAgcmV0dXJuIGZhbHNlOwogCSAgICB9 CisKKwkgIHRyZWUgYXJlZiA9IHQ7CisJICBmb3IgKGkgPSAwOyBpIDwgZGltcy5sZW5ndGgg KCk7IGkrKykKKwkgICAgeworCSAgICAgIGlmIChkaW1zW2ldLmxlbmd0aCAmJiBpbnRlZ2Vy X29uZXAgKGRpbXNbaV0ubGVuZ3RoKSkKKwkJeworCQkgIHRyZWUgbGIgPSBkaW1zW2ldLmxv d19ib3VuZDsKKwkJICBhcmVmID0gY29udmVydF9mcm9tX3JlZmVyZW5jZSAoYXJlZik7CisJ CSAgYXJlZiA9IGJ1aWxkX2FycmF5X3JlZiAoT01QX0NMQVVTRV9MT0NBVElPTiAoYyksIGFy ZWYsIGxiKTsKKwkJfQorCSAgICAgIGVsc2UKKwkJeworCQkgIGlmIChUUkVFX0NPREUgKFRS RUVfVFlQRSAoYXJlZikpID09IFBPSU5URVJfVFlQRSkKKwkJICAgIHQgPSBhcmVmOworCQkg IGJyZWFrOworCQl9CisJICAgIH0KKwogCSAgT01QX0NMQVVTRV9ERUNMIChjKSA9IGZpcnN0 OwogCSAgaWYgKE9NUF9DTEFVU0VfQ09ERSAoYykgPT0gT01QX0NMQVVTRV9IQVNfREVWSUNF X0FERFIpCiAJICAgIHJldHVybiBmYWxzZTsKQEAgLTU2ODEsNyArNTcwNiw4IEBAIGhhbmRs ZV9vbXBfYXJyYXlfc2VjdGlvbnMgKHRyZWUgYywgZW51bSBjX29tcF9yZWdpb25fdHlwZSBv cnQpCiAJICBib29sIHJlZmVyZW5jZV9hbHdheXNfcG9pbnRlciA9IHRydWU7CiAJICB0cmVl IGMyID0gYnVpbGRfb21wX2NsYXVzZSAoT01QX0NMQVVTRV9MT0NBVElPTiAoYyksCiAJCQkJ ICAgICAgT01QX0NMQVVTRV9NQVApOwotCSAgaWYgKFRSRUVfQ09ERSAodCkgPT0gQ09NUE9O RU5UX1JFRikKKwkgIGlmIChUUkVFX0NPREUgKHQpID09IENPTVBPTkVOVF9SRUYgfHwgVFJF RV9DT0RFICh0KSA9PSBBUlJBWV9SRUYKKwkgICAgICB8fCAoVFJFRV9DT0RFICh0KSA9PSBJ TkRJUkVDVF9SRUYgJiYgIVJFRkVSRU5DRV9SRUZfUCAodCkpKQogCSAgICB7CiAJICAgICAg T01QX0NMQVVTRV9TRVRfTUFQX0tJTkQgKGMyLCBHT01QX01BUF9BVFRBQ0hfREVUQUNIKTsK IApkaWZmIC0tZ2l0IGEvZ2NjL2dpbXBsaWZ5LmNjIGIvZ2NjL2dpbXBsaWZ5LmNjCmluZGV4 IGY1NzBkYWEwMTVhLi43N2I5NWNkODAwMCAxMDA2NDQKLS0tIGEvZ2NjL2dpbXBsaWZ5LmNj CisrKyBiL2djYy9naW1wbGlmeS5jYwpAQCAtOTYyNiw3ICs5NjI2LDEwIEBAIGdpbXBsaWZ5 X3NjYW5fb21wX2NsYXVzZXMgKHRyZWUgKmxpc3RfcCwgZ2ltcGxlX3NlcSAqcHJlX3AsCiAJ CSAgIHx8IChjb21wb25lbnRfcmVmX3AKIAkJICAgICAgICYmIChJTkRJUkVDVF9SRUZfUCAo ZGVjbCkKIAkJCSAgIHx8IFRSRUVfQ09ERSAoZGVjbCkgPT0gTUVNX1JFRgotCQkJICAgfHwg VFJFRV9DT0RFIChkZWNsKSA9PSBBUlJBWV9SRUYpKSkKKwkJCSAgIHx8IFRSRUVfQ09ERSAo ZGVjbCkgPT0gQVJSQVlfUkVGKSkKKwkJICAgfHwgKFRSRUVfQ09ERSAoZGVjbCkgPT0gQVJS QVlfUkVGCisJCSAgICAgICAmJiBUUkVFX0NPREUgKFRSRUVfVFlQRSAoZGVjbCkpID09IFBP SU5URVJfVFlQRQorCQkgICAgICAgJiYgT01QX0NMQVVTRV9NQVBfS0lORCAoYykgPT0gR09N UF9NQVBfQVRUQUNIX0RFVEFDSCkpCiAJCSAgJiYgT01QX0NMQVVTRV9NQVBfS0lORCAoYykg IT0gR09NUF9NQVBfVE9fUFNFVAogCQkgICYmIE9NUF9DTEFVU0VfTUFQX0tJTkQgKGMpICE9 IEdPTVBfTUFQX0FUVEFDSAogCQkgICYmIE9NUF9DTEFVU0VfTUFQX0tJTkQgKGMpICE9IEdP TVBfTUFQX0RFVEFDSApkaWZmIC0tZ2l0IGEvZ2NjL3Rlc3RzdWl0ZS9jLWMrKy1jb21tb24v Z29tcC90YXJnZXQtZW50ZXItZGF0YS0xLmMgYi9nY2MvdGVzdHN1aXRlL2MtYysrLWNvbW1v bi9nb21wL3RhcmdldC1lbnRlci1kYXRhLTEuYwppbmRleCBjZTc2NmQyOWUyZC4uM2ExYjQ4 OGZhMWYgMTAwNjQ0Ci0tLSBhL2djYy90ZXN0c3VpdGUvYy1jKystY29tbW9uL2dvbXAvdGFy Z2V0LWVudGVyLWRhdGEtMS5jCisrKyBiL2djYy90ZXN0c3VpdGUvYy1jKystY29tbW9uL2dv bXAvdGFyZ2V0LWVudGVyLWRhdGEtMS5jCkBAIC0yMSw0ICsyMSw1IEBAIHZvaWQgZnVuYyAo c3RydWN0IGZvbyAqZiwgaW50IG4sIGludCBtKQogICAjcHJhZ21hIG9tcCB0YXJnZXQgZW50 ZXIgZGF0YSBtYXAgKHRvOiBmLT5iYXJzW25dLnZlY3RvcnNbOmYtPmJhcnNbbl0ubnVtX3Zl Y3RvcnNdKQogfQogCi0vKiB7IGRnLWZpbmFsIHsgc2Nhbi10cmVlLWR1bXAtdGltZXMgIm1h cFxcKHRvOlxcKl9cWzAtOVxdKyBcXFxbbGVuOiBfXFswLTlcXStcXFxdXFwpIG1hcFxcKGF0 dGFjaDpcW14tXF0rLT52ZWN0b3JzIFxcXFtiaWFzOiBcW15cXVxdK1xcXF1cXCkiIDMgImdp bXBsZSIgfSB9ICovCisvKiB7IGRnLWZpbmFsIHsgc2Nhbi10cmVlLWR1bXAtdGltZXMgIm1h cFxcKHRvOlxcKl9cWzAtOVxdKyBcXFxbbGVuOiBfXFswLTlcXStcXFxdXFwpIG1hcFxcKGF0 dGFjaDpcXCpfXFswLTlcXSsgXFxcW2JpYXM6IFxbXlxdXF0rXFxcXVxcKSIgMSAiZ2ltcGxl IiB9IH0gKi8KKy8qIHsgZGctZmluYWwgeyBzY2FuLXRyZWUtZHVtcC10aW1lcyAibWFwXFwo dG86XFwqX1xbMC05XF0rIFxcXFtsZW46IF9cWzAtOVxdK1xcXF1cXCkgbWFwXFwoYXR0YWNo OlxbXi1cXSstPnZlY3RvcnMgXFxcW2JpYXM6IFxbXlxdXF0rXFxcXVxcKSIgMiAiZ2ltcGxl IiB9IH0gKi8KZGlmZiAtLWdpdCBhL2xpYmdvbXAvdGVzdHN1aXRlL2xpYmdvbXAuYy1jKyst Y29tbW9uL3B0ci1hdHRhY2gtMi5jIGIvbGliZ29tcC90ZXN0c3VpdGUvbGliZ29tcC5jLWMr Ky1jb21tb24vcHRyLWF0dGFjaC0yLmMKbmV3IGZpbGUgbW9kZSAxMDA2NDQKaW5kZXggMDAw MDAwMDAwMDAuLjg4OWE0YTI1M2FlCi0tLSAvZGV2L251bGwKKysrIGIvbGliZ29tcC90ZXN0 c3VpdGUvbGliZ29tcC5jLWMrKy1jb21tb24vcHRyLWF0dGFjaC0yLmMKQEAgLTAsMCArMSw2 MCBAQAorI2luY2x1ZGUgPHN0ZGxpYi5oPgorCitzdHJ1Y3QgYmxrIHsgaW50IHgsIHk7IH07 CitzdHJ1Y3QgTAoreworICAjZGVmaW5lIE4gMTAKKyAgc3RydWN0IHsKKyAgICBpbnQgbnVt X2Jsb2Nrc1tOXTsKKyAgICBzdHJ1Y3QgYmxrICogYmxvY2tzW05dOworICB9IG07Cit9Owor Cit2b2lkIGZvbyAoc3RydWN0IEwgKmwpCit7CisgIGZvciAoaW50IGkgPSAwOyBpIDwgTjsg aSsrKQorICAgIHsKKyAgICAgIGwtPm0uYmxvY2tzW2ldID0gKHN0cnVjdCBibGsgKikgbWFs bG9jIChzaXplb2YgKHN0cnVjdCBibGspICogTik7CisgICAgICBsLT5tLm51bV9ibG9ja3Nb aV0gPSBOOworICAgIH0KKworICAjcHJhZ21hIG9tcCB0YXJnZXQgZW50ZXIgZGF0YSBtYXAo dG86bFs6MV0pCisgIGZvciAoaW50IGkgPSAwOyBpIDwgTjsgaSsrKQorICAgIHsKKyAgICAg ICNwcmFnbWEgb21wIHRhcmdldCBlbnRlciBkYXRhIG1hcCh0bzpsLT5tLmJsb2Nrc1tpXVs6 bC0+bS5udW1fYmxvY2tzW2ldXSkKKyAgICB9CisKKyAgI3ByYWdtYSBvbXAgdGFyZ2V0Cisg IHsKKyAgICBmb3IgKGludCBpID0gMDsgaSA8IE47IGkrKykKKyAgICAgIGZvciAoaW50IGog PSAwOyBqIDwgTjsgaisrKQorCXsKKwkgIGwtPm0uYmxvY2tzW2ldW2pdLnggPSBpICsgajsK KwkgIGwtPm0uYmxvY2tzW2ldW2pdLnkgPSBpICogajsKKwl9CisgIH0KKworICBmb3IgKGlu dCBpID0gMDsgaSA8IE47IGkrKykKKyAgICB7CisgICAgICAjcHJhZ21hIG9tcCB0YXJnZXQg ZXhpdCBkYXRhIG1hcChmcm9tOmwtPm0uYmxvY2tzW2ldWzpsLT5tLm51bV9ibG9ja3NbaV1d KQorICAgIH0KKyAgI3ByYWdtYSBvbXAgdGFyZ2V0IGV4aXQgZGF0YSBtYXAoZnJvbTpsWzox XSkKKworCisgIGZvciAoaW50IGkgPSAwOyBpIDwgTjsgaSsrKQorICAgIGZvciAoaW50IGog PSAwOyBqIDwgTjsgaisrKQorICAgICAgeworCWlmIChsLT5tLmJsb2Nrc1tpXVtqXS54ICE9 IGkgKyBqKQorCSAgYWJvcnQgKCk7CisJaWYgKGwtPm0uYmxvY2tzW2ldW2pdLnkgIT0gaSAq IGopCisJICBhYm9ydCAoKTsKKyAgICAgIH0KKworfQorCitpbnQgbWFpbiAodm9pZCkKK3sK KyAgc3RydWN0IEwgbDsKKyAgZm9vICgmbCk7CisgIHJldHVybiAwOworfQo= --------------lsx2S7f8WeAM3vY0W8DOy3i7--