From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from mga14.intel.com (mga14.intel.com [192.55.52.115]) by sourceware.org (Postfix) with ESMTPS id 76ADE3858D1E for ; Wed, 19 Apr 2023 06:46:48 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 76ADE3858D1E Authentication-Results: sourceware.org; dmarc=pass (p=none dis=none) header.from=intel.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=intel.com DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/simple; d=intel.com; i=@intel.com; q=dns/txt; s=Intel; t=1681886808; x=1713422808; h=from:to:cc:subject:date:message-id:references: in-reply-to:content-transfer-encoding:mime-version; bh=qOnQlcXo69/fewmvaTJkTVY4bipwugIoOtHQ5Wr+Kjg=; b=UUWQoLRIEfmWxfGQHIUAlih1JHerfguD5rujr0hYI9x3Q61Khm9XZbWo Rz/WpG/+Hz9vkiNBM5FqUY3bTZIYGZfZ9pXdduVKKHHi0pjIggcrOBUme 7HckiKMGTfzO3nDrOh1l/NvI+gxagVAknCJvzc5g1VPwzf+oKRyJJAW/W i8Fm9q4TwOeyrSo3XZBpAfoFkduS8tRPtP7iQL0VBPH6/6Jg44AcVRCay 1WuTPY7sNOyCeJnP3bBejhcmU0nsrWtHcJiHmKsH9S3XJXtSdQiC7trPR 4SOTfSwziKuFVevQ3WNlW/IIe1PmiVV2oCi1sJv2yFurPUOsCIibSpw+n g==; X-IronPort-AV: E=McAfee;i="6600,9927,10684"; a="345369866" X-IronPort-AV: E=Sophos;i="5.99,208,1677571200"; d="scan'208";a="345369866" Received: from orsmga002.jf.intel.com ([10.7.209.21]) by fmsmga103.fm.intel.com with ESMTP/TLS/ECDHE-RSA-AES256-GCM-SHA384; 18 Apr 2023 23:46:47 -0700 X-ExtLoop1: 1 X-IronPort-AV: E=McAfee;i="6600,9927,10684"; a="691382644" X-IronPort-AV: E=Sophos;i="5.99,208,1677571200"; d="scan'208";a="691382644" Received: from orsmsx602.amr.corp.intel.com ([10.22.229.15]) by orsmga002.jf.intel.com with ESMTP; 18 Apr 2023 23:46:47 -0700 Received: from orsmsx601.amr.corp.intel.com (10.22.229.14) by ORSMSX602.amr.corp.intel.com (10.22.229.15) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_128_GCM_SHA256) id 15.1.2507.23; Tue, 18 Apr 2023 23:46:46 -0700 Received: from ORSEDG601.ED.cps.intel.com (10.7.248.6) by orsmsx601.amr.corp.intel.com (10.22.229.14) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_128_GCM_SHA256) id 15.1.2507.23 via Frontend Transport; Tue, 18 Apr 2023 23:46:46 -0700 Received: from NAM10-DM6-obe.outbound.protection.outlook.com (104.47.58.101) by edgegateway.intel.com (134.134.137.102) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.1.2507.23; Tue, 18 Apr 2023 23:46:46 -0700 ARC-Seal: i=1; a=rsa-sha256; s=arcselector9901; d=microsoft.com; cv=none; b=g1MNcr7+Ma6G9hN2As9xjKP6edSteCfcgi89m3sUOksBspSdkLga9VadmeevfT3XrzMApn4RQm3E3PscoB4b4xwS2AdkULwQG2pNjyQc2w+vPcAfLE0aI8Cs77OGMY2qFUs3pbTZ9inxNCmA8NaMEBwTaVbDiX9yH4EkONBvD495MpCUqbLmNPbfJvKR+Pg4K7+bSENfo82mlgbK1+fUHbp5K0i8ftEgZT2Rs+zrcYxyAk4procSUHIvUXXxD1ztusU0H13gFNbnDbtD81gN426m88Y4sCQkMmBdYC54GN1dTtl36A5M0+sF6RONesCxTGrh9mONVymoi5TiZzWkQw== 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=COMKYNErKL4VTukiYMd1cncl9MctHLGe76TtQCMcKK0=; b=m5KxN7n7jSf4OtGbFOi4E9aYnpYjV5o0UbeR9v16K6/vGg9XcP4QHp5wGXnnukJGwD+7VTbMvOLNftoBIvj6LClPG213asbdsE8WYO3sUcbn90imMOir7lbdtKhUOTSictHdwYJiBaI3Z65Hew4obVHLHIqhKXeNmzg8f/5JnYztcGpPZ7MsSoeVV9myJvitgb/HA711zFeSH2EAzQ662LOKDfkdFkelIgyMubabamXv5Bn1MxOd5EUfVghGTH7jc6wavKsvdpQuuhA+iuJZpDpbwhClQy9qsHwYb2EJrkBA3Ge998J2Xv+oq+dhRTG1WVXx+jMoHs7iPuLTfS9Hbg== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass smtp.mailfrom=intel.com; dmarc=pass action=none header.from=intel.com; dkim=pass header.d=intel.com; arc=none Received: from MW5PR11MB5908.namprd11.prod.outlook.com (2603:10b6:303:194::10) by MW3PR11MB4572.namprd11.prod.outlook.com (2603:10b6:303:5e::14) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.6319.21; Wed, 19 Apr 2023 06:46:44 +0000 Received: from MW5PR11MB5908.namprd11.prod.outlook.com ([fe80::bbc5:f013:1f53:10a9]) by MW5PR11MB5908.namprd11.prod.outlook.com ([fe80::bbc5:f013:1f53:10a9%6]) with mapi id 15.20.6298.045; Wed, 19 Apr 2023 06:46:44 +0000 From: "Li, Pan2" To: Richard Biener CC: "gcc-patches@gcc.gnu.org" , "juzhe.zhong@rivai.ai" , "kito.cheng@sifive.com" , "richard.sandiford@arm.com" , "Wang, Yanzhang" Subject: RE: [PATCH v2] RISC-V: Allow Vector IOR(V1, NOT V1) optimization Thread-Topic: [PATCH v2] RISC-V: Allow Vector IOR(V1, NOT V1) optimization Thread-Index: AQHZcdVuFQDFsqXx0Um5zVl6ZEYotq8yL+aAgAAA38A= Date: Wed, 19 Apr 2023 06:46:44 +0000 Message-ID: References: <20230417145025.2291874-1-pan2.li@intel.com> <20230418090855.3012513-1-pan2.li@intel.com> In-Reply-To: Accept-Language: en-US Content-Language: en-US X-MS-Has-Attach: X-MS-TNEF-Correlator: authentication-results: dkim=none (message not signed) header.d=none;dmarc=none action=none header.from=intel.com; x-ms-publictraffictype: Email x-ms-traffictypediagnostic: MW5PR11MB5908:EE_|MW3PR11MB4572:EE_ x-ms-office365-filtering-correlation-id: 6d95fd7a-b589-4ded-dc5f-08db40a1d71c x-ms-exchange-senderadcheck: 1 x-ms-exchange-antispam-relay: 0 x-microsoft-antispam: BCL:0; x-microsoft-antispam-message-info: anKGewXIHSeqkftTz9PYbYi+TPq4Fu0cvjdgjgP1ToSBy5/l4AZwDs2yw7i/f2dsjNQpuuwv0S7F/2IInPB1J/CtHCt4xSRh+0CfmdnFE6b3y5SqmzI9Zb0PlBj1H84eJsWZRoA2nfZ8gbRbcoXK5R2ti03eo4JHNc3/VZVlI6ON0Rx3NXfxvdn35Nm8ETzINlRzRghp/0PlIpekKm1abk9yQ8cpRO6OXqmaT3tykEefcSq68R8YPaqeqPUjnscM4A7LZK224xKfrIvcu3m678A6df0PpeH8WFZjenZc3Dk7mL7YeEWCClKPcPJGpRt91SG0u/uK1Yk+44Z9K6wQh5BK5QGGIWB268v3rcWR8qBQKVTfUa6BXxDkOREmJ+r3OOaK9gS45FJvuOElqNDpbrUiJe2v908Ro+ka2NLXVJIaCoCeN0wJRNqXWyW6FU+O33c1H0ECoT3RP8T+l+THpYfofcp2CuOnW8q6dvb/w8yCB5Kmr57ppymvhHdbAf+BRGyEqbBh3cSjvV/VWp9haGcJ9Lpv2nRP+y4nnldsNUVqb7rjqpU6KvlBHmlXVeJOp3j/9+twP3P0ggiCg31jdurLe6+G9BK1lhJTn/5aRKk= x-forefront-antispam-report: CIP:255.255.255.255;CTRY:;LANG:en;SCL:1;SRV:;IPV:NLI;SFV:NSPM;H:MW5PR11MB5908.namprd11.prod.outlook.com;PTR:;CAT:NONE;SFS:(13230028)(366004)(396003)(136003)(346002)(376002)(39860400002)(451199021)(38100700002)(55016003)(2906002)(8676002)(8936002)(52536014)(5660300002)(33656002)(71200400001)(86362001)(186003)(7696005)(107886003)(4326008)(64756008)(26005)(66946007)(66476007)(66446008)(9686003)(66556008)(76116006)(6916009)(41300700001)(53546011)(6506007)(83380400001)(316002)(54906003)(478600001)(122000001)(82960400001)(38070700005)(84970400001);DIR:OUT;SFP:1102; x-ms-exchange-antispam-messagedata-chunkcount: 1 x-ms-exchange-antispam-messagedata-0: =?us-ascii?Q?cipJbw8eEHIASovpqc3zA3K+olMqwSHBhDHP3lXnASN8ERiKL5TQZVvsF1ef?= =?us-ascii?Q?4twQvkIYlX4mYCRTYrsEvB3FmdU4HHa6wB0Z/QI7Los1rfU40x7K8nUX1QLx?= =?us-ascii?Q?I4qV7XeUj1jZY2a0WvmO8bJpAFHlFOjYidk4VH6z41GOwXep/yLbmP++3xOL?= =?us-ascii?Q?Tx1/UPhQICpWMmHJUNcW7FwlUS0BOiqUYOT+LK8FiR4JVEhP/ZOn2hdktqUV?= =?us-ascii?Q?9EitHccT8EmdFWRmOkrABXLcLXdt9FXwUvVJvxqlsRUK8BeiZQrEGS5ty+G7?= =?us-ascii?Q?5zH5F63ShQgjhxuXbbD9aTd1eTzkQRbyiHlszdycycaWHMBbI31PQN8QhjNV?= =?us-ascii?Q?aAxE1nY9uv4fRTDPM9hCkgYXU//YzMC47VdqzV00TIIhuh2e7gfbRYRuGHtw?= =?us-ascii?Q?2QWnhMOc3p5zxRnpkqdc0Erf2AT064sXfcyva68vLtxlllkV+2DexVZ28pe+?= =?us-ascii?Q?63VMTwb5Kk0ZaaButnS6IomsdI08BqlOyC2Rh0p1hFNjEVO4Tl+VPZlZtV+H?= =?us-ascii?Q?WDDXNm2I2ICD3UaD9hKoXb6WDwuCWQySWlQgPP3xlvU27WJ1/wpixwMD/OTB?= =?us-ascii?Q?waRDg7BqWh+gxtFd1lQYUJ7l1ROfXZO0L47jFHWBiLRveFDF/Hae9Gh8mFjU?= =?us-ascii?Q?tGhCcNvzEFMJ8LT/5s/W7wGFXiU4tkgPP1BrMV+JsKVLQqzJc6ajeOQpyClX?= =?us-ascii?Q?Iw38HLYeGq6dcn11b10bKYJmE1Y9tYsI1KK53Aq7QTrPR8vho9xvx0kk2NJJ?= =?us-ascii?Q?KeZh+PsSon2ULCpQDjTMkQr8gewz7/hDJ//NXedxH8XdZXgrpZdtEiTbHQ+t?= =?us-ascii?Q?1WAkZuTZ3RVcOU7KdI3bpzKOf/tR9sQkAjpNeSfiCpiCMQHSYIvgMjZWOB2t?= =?us-ascii?Q?ADGUM/r6nH/1cVkndUkYY3QUfKEZctQcmBfT4FYLQ0OzpR5fGYo63/UDvlNX?= =?us-ascii?Q?UHlf9VIEFP6Q8xkjRJf0Mf+C4LbgmYATXH7JJ8kxhPNuaCIadRYi7n07Qvvw?= =?us-ascii?Q?xht1w8OxTjCJgHeWyV+1E2ZgyWDVn0TivM69h3Ra3dLsJNczwK9jvY32slPA?= =?us-ascii?Q?gfPtzcNEkmZiRvy71cLwbyew7TskNqbChHuCvgUmonCIX3OQ1XO/Vr70EBTj?= =?us-ascii?Q?bfgib6a5NTexdY5PJKZJ49lSR9sfrpvKhwcxwvb+pnNmnezg9ZY8WvwL69uB?= =?us-ascii?Q?BN7iNxXgseIemT1VYk6FJh5tQkaw94qf5N98f85KfXS3TgFcxXaYfrB+GKep?= =?us-ascii?Q?y87eAdYS3+KxE7o3QNEdwl0MCno31BOy28YjArCBJjbxEDXGYMSkxwAf1t39?= =?us-ascii?Q?VN6pJ9sNGkj3a/RuymYtPJAldNHcpuBgaipyWo/YoxmfPDOv7GcoCTB8ogI+?= =?us-ascii?Q?TOa6RHHXOB1ihzcAwU2oO3TFaP2ZEj0Tj1O7QgbOCLgesd2Z8tDKgH62qwr3?= =?us-ascii?Q?Jze0eYQ+1DhY56BvAuNmptfFPlU3oe1EAAg06gDplqeoADna5DQK2ZACjd2E?= =?us-ascii?Q?Vzwl6Uy+zsC21M/Ah59gBr3AExZiJF/fAKb14CPQWQsn+4DJ6y1go2sXilfn?= =?us-ascii?Q?Byek+vPAOqwWDbpVgM8=3D?= Content-Type: text/plain; charset="us-ascii" Content-Transfer-Encoding: quoted-printable MIME-Version: 1.0 X-MS-Exchange-CrossTenant-AuthAs: Internal X-MS-Exchange-CrossTenant-AuthSource: MW5PR11MB5908.namprd11.prod.outlook.com X-MS-Exchange-CrossTenant-Network-Message-Id: 6d95fd7a-b589-4ded-dc5f-08db40a1d71c X-MS-Exchange-CrossTenant-originalarrivaltime: 19 Apr 2023 06:46:44.5413 (UTC) X-MS-Exchange-CrossTenant-fromentityheader: Hosted X-MS-Exchange-CrossTenant-id: 46c98d88-e344-4ed4-8496-4ed7712e255d X-MS-Exchange-CrossTenant-mailboxtype: HOSTED X-MS-Exchange-CrossTenant-userprincipalname: Je/Xcs5BpCgQ16uD5z4sYuLvKPVaMjWVVTX3usw82FUaRR7xUtZLVZDKJ4ZylHJrc9FHr03/xGj94RwaxMp/dg== X-MS-Exchange-Transport-CrossTenantHeadersStamped: MW3PR11MB4572 X-OriginatorOrg: intel.com X-Spam-Status: No, score=-12.6 required=5.0 tests=BAYES_00,DKIMWL_WL_HIGH,DKIM_SIGNED,DKIM_VALID,DKIM_VALID_AU,DKIM_VALID_EF,GIT_PATCH_0,KAM_SHORT,SPF_HELO_NONE,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 List-Id: Oh, I see. The message need to be re-generated. Thank you for pointing out,= will update ASPA. Pan -----Original Message----- From: Richard Biener =20 Sent: Wednesday, April 19, 2023 2:40 PM To: Li, Pan2 Cc: gcc-patches@gcc.gnu.org; juzhe.zhong@rivai.ai; kito.cheng@sifive.com; r= ichard.sandiford@arm.com; Wang, Yanzhang Subject: Re: [PATCH v2] RISC-V: Allow Vector IOR(V1, NOT V1) optimization On Tue, 18 Apr 2023, pan2.li@intel.com wrote: > From: Pan Li >=20 > This patch add the optimization for the vector IOR(V1, NOT V1). Assume=20 > we have below sample code. >=20 > vbool32_t test_shortcut_for_riscv_vmorn_case_5(vbool32_t v1, size_t=20 > vl) { > return __riscv_vmorn_mm_b32(v1, v1, vl); } >=20 > Before this patch: > vsetvli a5,zero,e8,mf4,ta,ma > vlm.v v24,0(a1) > vsetvli zero,a2,e8,mf4,ta,ma > vmorn.mm v24,v24,v24 > vsetvli a5,zero,e8,mf4,ta,ma > vsm.v v24,0(a0) > ret >=20 > After this patch: > vsetvli zero,a2,e8,mf4,ta,ma > vmset.m v24 > vsetvli a5,zero,e8,mf4,ta,ma > vsm.v v24,0(a0) > ret >=20 > Or in RTL's perspective, > from: > (ior:VNx2BI (reg/v:VNx2BI 137 [ v1 ]) (not:VNx2BI (reg/v:VNx2BI 137 [=20 > v1 ]))) > to: > (const_vector:VNx2BI repeat [ (const_int 1 [0x1]) ]) >=20 > The similar optimization like VMANDN has enabled already. There should=20 > be no difference execpt the operator when compare the VMORN and VMANDN=20 > for such kind of optimization. The patch allows the VECTOR_BOOL=20 > IOR(V1, NOT V1) simplification besides the existing SCALAR_INT mode. >=20 > gcc/ChangeLog: >=20 > * simplify-rtx.cc (simplify_context::simplify_binary_operation_1): This needs some text > gcc/testsuite/ChangeLog: >=20 > * gcc.target/riscv/rvv/base/mask_insn_shortcut.c: Likewise. OK with that fixed. > * gcc.target/riscv/simplify_ior_optimization.c: New test. >=20 > Signed-off-by: Pan Li > --- > gcc/simplify-rtx.cc | 4 +- > .../riscv/rvv/base/mask_insn_shortcut.c | 3 +- > .../riscv/simplify_ior_optimization.c | 50 +++++++++++++++++++ > 3 files changed, 53 insertions(+), 4 deletions(-) create mode 100644=20 > gcc/testsuite/gcc.target/riscv/simplify_ior_optimization.c >=20 > diff --git a/gcc/simplify-rtx.cc b/gcc/simplify-rtx.cc index=20 > ee75079917f..3bc9b2f55ea 100644 > --- a/gcc/simplify-rtx.cc > +++ b/gcc/simplify-rtx.cc > @@ -3332,8 +3332,8 @@ simplify_context::simplify_binary_operation_1 (rtx_= code code, > if (((GET_CODE (op0) =3D=3D NOT && rtx_equal_p (XEXP (op0, 0), op1= )) > || (GET_CODE (op1) =3D=3D NOT && rtx_equal_p (XEXP (op1, 0), op0))) > && ! side_effects_p (op0) > - && SCALAR_INT_MODE_P (mode)) > - return constm1_rtx; > + && GET_MODE_CLASS (mode) !=3D MODE_CC) > + return CONSTM1_RTX (mode); > =20 > /* (ior A C) is C if all bits of A that might be nonzero are on in= C. */ > if (CONST_INT_P (op1) > diff --git=20 > a/gcc/testsuite/gcc.target/riscv/rvv/base/mask_insn_shortcut.c=20 > b/gcc/testsuite/gcc.target/riscv/rvv/base/mask_insn_shortcut.c > index 83cc4a1b5a5..57d0241675a 100644 > --- a/gcc/testsuite/gcc.target/riscv/rvv/base/mask_insn_shortcut.c > +++ b/gcc/testsuite/gcc.target/riscv/rvv/base/mask_insn_shortcut.c > @@ -233,9 +233,8 @@ vbool64_t=20 > test_shortcut_for_riscv_vmxnor_case_6(vbool64_t v1, size_t vl) { > /* { dg-final { scan-assembler-not {vmxor\.mm\s+v[0-9]+,\s*v[0-9]+} }=20 > } */ > /* { dg-final { scan-assembler-not {vmor\.mm\s+v[0-9]+,\s*v[0-9]+} }=20 > } */ > /* { dg-final { scan-assembler-not {vmnor\.mm\s+v[0-9]+,\s*v[0-9]+} }=20 > } */ > -/* { dg-final { scan-assembler-times=20 > {vmorn\.mm\s+v[0-9]+,\s*v[0-9]+,\s*v[0-9]+} 7 } } */ > /* { dg-final { scan-assembler-not {vmxnor\.mm\s+v[0-9]+,\s*v[0-9]+}=20 > } } */ > /* { dg-final { scan-assembler-times {vmclr\.m\s+v[0-9]+} 14 } } */ > -/* { dg-final { scan-assembler-times {vmset\.m\s+v[0-9]+} 7 } } */ > +/* { dg-final { scan-assembler-times {vmset\.m\s+v[0-9]+} 14 } } */ > /* { dg-final { scan-assembler-times {vmmv\.m\s+v[0-9]+,\s*v[0-9]+}=20 > 14 } } */ > /* { dg-final { scan-assembler-times {vmnot\.m\s+v[0-9]+,\s*v[0-9]+}=20 > 14 } } */ diff --git=20 > a/gcc/testsuite/gcc.target/riscv/simplify_ior_optimization.c=20 > b/gcc/testsuite/gcc.target/riscv/simplify_ior_optimization.c > new file mode 100644 > index 00000000000..ec3bd0baf03 > --- /dev/null > +++ b/gcc/testsuite/gcc.target/riscv/simplify_ior_optimization.c > @@ -0,0 +1,50 @@ > +/* { dg-do compile } */ > +/* { dg-options "-march=3Drv64gc -mabi=3Dlp64 -O2" } */ > + > +#include > + > +uint8_t test_simplify_ior_scalar_case_0 (uint8_t a) { > + return a | ~a; > +} > + > +uint16_t test_simplify_ior_scalar_case_1 (uint16_t a) { > + return a | ~a; > +} > + > +uint32_t test_simplify_ior_scalar_case_2 (uint32_t a) { > + return a | ~a; > +} > + > +uint64_t test_simplify_ior_scalar_case_3 (uint64_t a) { > + return a | ~a; > +} > + > +int8_t test_simplify_ior_scalar_case_4 (int8_t a) { > + return a | ~a; > +} > + > +int16_t test_simplify_ior_scalar_case_5 (int16_t a) { > + return a | ~a; > +} > + > +int32_t test_simplify_ior_scalar_case_6 (int32_t a) { > + return a | ~a; > +} > + > +int64_t test_simplify_ior_scalar_case_7 (int64_t a) { > + return a | ~a; > +} > + > +/* { dg-final { scan-assembler-times {li\s+a[0-9]+,\s*-1} 6 } } */ > +/* { dg-final { scan-assembler-times {li\s+a[0-9]+,\s*255} 1 } } */ > +/* { dg-final { scan-assembler-times {li\s+a[0-9]+,\s*65536} 1 } } */ > +/* { dg-final { scan-assembler-not {or\s+a[0-9]+} } } */ > +/* { dg-final { scan-assembler-not {not\s+a[0-9]+} } } */ >=20 -- Richard Biener SUSE Software Solutions Germany GmbH, Frankenstrasse 146, 90461 Nuernberg, = Germany; GF: Ivo Totev, Andrew Myers, Andrew McDonald, Boudien Moerman; HRB= 36809 (AG Nuernberg)