This is an archive of the discontinued LLVM Phabricator instance.

clang/AMDGPU: Remove flat-address-space from feature map
ClosedPublic

Authored by arsenm on Nov 28 2022, 2:59 PM.

Download Raw Diff

Details

Reviewers

yaxunl
JonChesterfield
b-sumner
jdoerfert
Joe_Nash

Summary

This was only used for checking if is_shared/is_private were legal,
which we're not bothering to do anymore.

This is apparently visible to more than the target attribute (which
seems to silently ignore unrecognized features), so this has the
potential to break something (i.e. see the OpenMP test change)

Diff Detail

Unit TestsFailed

	Time	Test
	400 ms	x64 debian > Profile-x86_64.Linux::counter_promo_for.c
	360 ms	x64 debian > Profile-x86_64.Linux::counter_promo_while.c

Event Timeline

arsenm created this revision.Nov 28 2022, 2:59 PM

Herald added a project: Restricted Project. · View Herald TranscriptNov 28 2022, 2:59 PM

Herald added subscribers: kosarev, kerbowa, tpr and 3 others. · View Herald Transcript

arsenm requested review of this revision.Nov 28 2022, 2:59 PM

Herald added subscribers: sstefan1, wdng. · View Herald TranscriptNov 28 2022, 2:59 PM

arsenm added a parent revision: D138868: AMDGPU/clang: Remove target features from address space test builtins.Nov 28 2022, 2:59 PM

Harbormaster completed remote builds in B199877: Diff 478387.Nov 28 2022, 4:23 PM

The OpenMP part is fine.

Rebase

Harbormaster completed remote builds in B205185: Diff 485650.Dec 29 2022, 4:29 PM

The code looks fine, but as you say, the change visible in user code and could break something. Do you want to handle that somehow? Maybe wait for @b-sumner

clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp
17	Was the "dpp" attribute chosen arbitrarily?

In D138870#4020204, @Joe_Nash wrote:

The code looks fine, but as you say, the change visible in user code and could break something. Do you want to handle that somehow? Maybe wait for @b-sumner

OpenMP assumes flat pointers all over, so if someone was relying on this behavior it wasn't doing anything useful. They had dead code

clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp
17	Yes, also meaningfully varies

Rebase

In D138870#4020211, @arsenm wrote:

In D138870#4020204, @Joe_Nash wrote:

The code looks fine, but as you say, the change visible in user code and could break something. Do you want to handle that somehow? Maybe wait for @b-sumner

OpenMP assumes flat pointers all over, so if someone was relying on this behavior it wasn't doing anything useful. They had dead code

I can't verify the idea that users of the behavior had dead code, but assuming that's right LGTM.

This revision is now accepted and ready to land.Jan 5 2023, 7:16 AM

Harbormaster completed remote builds in B205900: Diff 486561.Jan 5 2023, 8:52 AM

81849497b42e1a865af8aff65ab768e56a301c87

Revision Contents

Path

Size

clang/

lib/

Basic/

Targets/

AMDGPU.cpp

4 lines

test/

CodeGenOpenCL/

amdgpu-features.cl

72 lines

OpenMP/

amdgcn-attributes.cpp

4 lines

metadirective_device_isa_codegen_amdgcn.cpp

2 lines

Diff 486561

clang/lib/Basic/Targets/AMDGPU.cpp

Show First 20 Lines • Show All 193 Lines • ▼ Show 20 Lines	case GK_GFX1100:
IsWave32Capable = true;		IsWave32Capable = true;
Features["ci-insts"] = true;		Features["ci-insts"] = true;
Features["dot1-insts"] = true;		Features["dot1-insts"] = true;
Features["dot5-insts"] = true;		Features["dot5-insts"] = true;
Features["dot6-insts"] = true;		Features["dot6-insts"] = true;
Features["dot7-insts"] = true;		Features["dot7-insts"] = true;
Features["dot8-insts"] = true;		Features["dot8-insts"] = true;
Features["dl-insts"] = true;		Features["dl-insts"] = true;
Features["flat-address-space"] = true;
Features["16-bit-insts"] = true;		Features["16-bit-insts"] = true;
Features["dpp"] = true;		Features["dpp"] = true;
Features["gfx8-insts"] = true;		Features["gfx8-insts"] = true;
Features["gfx9-insts"] = true;		Features["gfx9-insts"] = true;
Features["gfx10-insts"] = true;		Features["gfx10-insts"] = true;
Features["gfx10-3-insts"] = true;		Features["gfx10-3-insts"] = true;
Features["gfx11-insts"] = true;		Features["gfx11-insts"] = true;
break;		break;
case GK_GFX1036:		case GK_GFX1036:
case GK_GFX1035:		case GK_GFX1035:
case GK_GFX1034:		case GK_GFX1034:
case GK_GFX1033:		case GK_GFX1033:
case GK_GFX1032:		case GK_GFX1032:
case GK_GFX1031:		case GK_GFX1031:
case GK_GFX1030:		case GK_GFX1030:
IsWave32Capable = true;		IsWave32Capable = true;
Features["ci-insts"] = true;		Features["ci-insts"] = true;
Features["dot1-insts"] = true;		Features["dot1-insts"] = true;
Features["dot2-insts"] = true;		Features["dot2-insts"] = true;
Features["dot5-insts"] = true;		Features["dot5-insts"] = true;
Features["dot6-insts"] = true;		Features["dot6-insts"] = true;
Features["dot7-insts"] = true;		Features["dot7-insts"] = true;
Features["dl-insts"] = true;		Features["dl-insts"] = true;
Features["flat-address-space"] = true;
Features["16-bit-insts"] = true;		Features["16-bit-insts"] = true;
Features["dpp"] = true;		Features["dpp"] = true;
Features["gfx8-insts"] = true;		Features["gfx8-insts"] = true;
Features["gfx9-insts"] = true;		Features["gfx9-insts"] = true;
Features["gfx10-insts"] = true;		Features["gfx10-insts"] = true;
Features["gfx10-3-insts"] = true;		Features["gfx10-3-insts"] = true;
Features["s-memrealtime"] = true;		Features["s-memrealtime"] = true;
Features["s-memtime-inst"] = true;		Features["s-memtime-inst"] = true;
break;		break;
case GK_GFX1012:		case GK_GFX1012:
case GK_GFX1011:		case GK_GFX1011:
Features["dot1-insts"] = true;		Features["dot1-insts"] = true;
Features["dot2-insts"] = true;		Features["dot2-insts"] = true;
Features["dot5-insts"] = true;		Features["dot5-insts"] = true;
Features["dot6-insts"] = true;		Features["dot6-insts"] = true;
Features["dot7-insts"] = true;		Features["dot7-insts"] = true;
[[fallthrough]];		[[fallthrough]];
case GK_GFX1013:		case GK_GFX1013:
case GK_GFX1010:		case GK_GFX1010:
IsWave32Capable = true;		IsWave32Capable = true;
Features["dl-insts"] = true;		Features["dl-insts"] = true;
Features["ci-insts"] = true;		Features["ci-insts"] = true;
Features["flat-address-space"] = true;
Features["16-bit-insts"] = true;		Features["16-bit-insts"] = true;
Features["dpp"] = true;		Features["dpp"] = true;
Features["gfx8-insts"] = true;		Features["gfx8-insts"] = true;
Features["gfx9-insts"] = true;		Features["gfx9-insts"] = true;
Features["gfx10-insts"] = true;		Features["gfx10-insts"] = true;
Features["s-memrealtime"] = true;		Features["s-memrealtime"] = true;
Features["s-memtime-inst"] = true;		Features["s-memtime-inst"] = true;
break;		break;
Show All 36 Lines	case GK_GFX801:
[[fallthrough]];		[[fallthrough]];
case GK_GFX705:		case GK_GFX705:
case GK_GFX704:		case GK_GFX704:
case GK_GFX703:		case GK_GFX703:
case GK_GFX702:		case GK_GFX702:
case GK_GFX701:		case GK_GFX701:
case GK_GFX700:		case GK_GFX700:
Features["ci-insts"] = true;		Features["ci-insts"] = true;
Features["flat-address-space"] = true;
[[fallthrough]];		[[fallthrough]];
case GK_GFX602:		case GK_GFX602:
case GK_GFX601:		case GK_GFX601:
case GK_GFX600:		case GK_GFX600:
Features["s-memtime-inst"] = true;		Features["s-memtime-inst"] = true;
break;		break;
case GK_NONE:		case GK_NONE:
break;		break;
▲ Show 20 Lines • Show All 213 Lines • Show Last 20 Lines

clang/test/CodeGenOpenCL/amdgpu-features.cl

	Show First 20 Lines • Show All 49 Lines • ▼ Show 20 Lines

	// NOCPU-NOT: "target-features"			// NOCPU-NOT: "target-features"
	// NOCPU-WAVE32: "target-features"="+wavefrontsize32"			// NOCPU-WAVE32: "target-features"="+wavefrontsize32"
	// NOCPU-WAVE64: "target-features"="+wavefrontsize64"			// NOCPU-WAVE64: "target-features"="+wavefrontsize64"

	// GFX600: "target-features"="+s-memtime-inst,+wavefrontsize64"			// GFX600: "target-features"="+s-memtime-inst,+wavefrontsize64"
	// GFX601: "target-features"="+s-memtime-inst,+wavefrontsize64"			// GFX601: "target-features"="+s-memtime-inst,+wavefrontsize64"
	// GFX602: "target-features"="+s-memtime-inst,+wavefrontsize64"			// GFX602: "target-features"="+s-memtime-inst,+wavefrontsize64"
	// GFX700: "target-features"="+ci-insts,+flat-address-space,+s-memtime-inst,+wavefrontsize64"			// GFX700: "target-features"="+ci-insts,+s-memtime-inst,+wavefrontsize64"
	// GFX701: "target-features"="+ci-insts,+flat-address-space,+s-memtime-inst,+wavefrontsize64"			// GFX701: "target-features"="+ci-insts,+s-memtime-inst,+wavefrontsize64"
	// GFX702: "target-features"="+ci-insts,+flat-address-space,+s-memtime-inst,+wavefrontsize64"			// GFX702: "target-features"="+ci-insts,+s-memtime-inst,+wavefrontsize64"
	// GFX703: "target-features"="+ci-insts,+flat-address-space,+s-memtime-inst,+wavefrontsize64"			// GFX703: "target-features"="+ci-insts,+s-memtime-inst,+wavefrontsize64"
	// GFX704: "target-features"="+ci-insts,+flat-address-space,+s-memtime-inst,+wavefrontsize64"			// GFX704: "target-features"="+ci-insts,+s-memtime-inst,+wavefrontsize64"
	// GFX705: "target-features"="+ci-insts,+flat-address-space,+s-memtime-inst,+wavefrontsize64"			// GFX705: "target-features"="+ci-insts,+s-memtime-inst,+wavefrontsize64"
	// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX802: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX802: "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX803: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX803: "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX805: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX805: "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX810: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX810: "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX900: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX900: "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX902: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX902: "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot7-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX909: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX909: "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX90A: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX90A: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX90C: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX90C: "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX940: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-address-space,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"			// GFX940: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
	// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"			// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
	// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"			// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
	// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"			// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
	// GFX1013: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"			// GFX1013: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
	// GFX1030: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"			// GFX1030: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
	// GFX1031: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"			// GFX1031: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
	// GFX1032: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"			// GFX1032: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
	// GFX1033: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"			// GFX1033: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
	// GFX1034: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"			// GFX1034: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
	// GFX1035: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"			// GFX1035: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
	// GFX1036: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"			// GFX1036: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
	// GFX1100: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"			// GFX1100: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"
	// GFX1101: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"			// GFX1101: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"
	// GFX1102: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"			// GFX1102: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"
	// GFX1103: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"			// GFX1103: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"
	// GFX1103-W64: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize64"			// GFX1103-W64: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize64"

	kernel void test() {}			kernel void test() {}

clang/test/OpenMP/amdgcn-attributes.cpp

	Show All 27 Lines
	}			}

	int callable(int x) {			int callable(int x) {
	// ALL-LABEL: @_Z8callablei(i32 noundef %x) #1			// ALL-LABEL: @_Z8callablei(i32 noundef %x) #1
	return x + 1;			return x + 1;
	}			}

	// DEFAULT: attributes #0 = { convergent noinline norecurse nounwind optnone "kernel" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }			// DEFAULT: attributes #0 = { convergent noinline norecurse nounwind optnone "kernel" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
	// CPU: attributes #0 = { convergent noinline norecurse nounwind optnone "kernel" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx900" "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" "uniform-work-group-size"="true" }			// CPU: attributes #0 = { convergent noinline norecurse nounwind optnone "kernel" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx900" "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" "uniform-work-group-size"="true" }
	// NOIEEE: attributes #0 = { convergent noinline norecurse nounwind optnone "amdgpu-ieee"="false" "kernel" "no-nans-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }			// NOIEEE: attributes #0 = { convergent noinline norecurse nounwind optnone "amdgpu-ieee"="false" "kernel" "no-nans-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
	// UNSAFEATOMIC: attributes #0 = { convergent noinline norecurse nounwind optnone "amdgpu-unsafe-fp-atomics"="true" "kernel" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }			// UNSAFEATOMIC: attributes #0 = { convergent noinline norecurse nounwind optnone "amdgpu-unsafe-fp-atomics"="true" "kernel" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }

	// DEFAULT: attributes #1 = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }			// DEFAULT: attributes #1 = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
	// CPU: attributes #1 = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx900" "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" }			// CPU: attributes #1 = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx900" "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" }
	// NOIEEE: attributes #1 = { convergent mustprogress noinline nounwind optnone "amdgpu-ieee"="false" "no-nans-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }			// NOIEEE: attributes #1 = { convergent mustprogress noinline nounwind optnone "amdgpu-ieee"="false" "no-nans-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
	// UNSAFEATOMIC: attributes #1 = { convergent mustprogress noinline nounwind optnone "amdgpu-unsafe-fp-atomics"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }			// UNSAFEATOMIC: attributes #1 = { convergent mustprogress noinline nounwind optnone "amdgpu-unsafe-fp-atomics"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }

clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp

	// REQUIRES: amdgpu-registered-target			// REQUIRES: amdgpu-registered-target

	// RUN: %clang_cc1 -fopenmp -x c++ -w -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc			// RUN: %clang_cc1 -fopenmp -x c++ -w -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
	// RUN: %clang_cc1 -fopenmp -x c++ -w -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -target-cpu gfx906 -o - \| FileCheck %s			// RUN: %clang_cc1 -fopenmp -x c++ -w -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -target-cpu gfx906 -o - \| FileCheck %s
	// expected-no-diagnostics			// expected-no-diagnostics

	#ifndef HEADER			#ifndef HEADER
	#define HEADER			#define HEADER

	int amdgcn_device_isa_selected() {			int amdgcn_device_isa_selected() {
	int threadCount = 0;			int threadCount = 0;

	#pragma omp target map(tofrom \			#pragma omp target map(tofrom \
	: threadCount)			: threadCount)
	{			{
	#pragma omp metadirective \			#pragma omp metadirective \
	when(device = {isa("flat-address-space")} \			when(device = {isa("dpp")} \
				Joe_NashUnsubmitted Not Done Reply Inline Actions Was the "dpp" attribute chosen arbitrarily? Joe_Nash: Was the "dpp" attribute chosen arbitrarily?
				arsenmAuthorUnsubmitted Done Reply Inline Actions Yes, also meaningfully varies arsenm: Yes, also meaningfully varies
	: parallel) default(single)			: parallel) default(single)
	threadCount++;			threadCount++;
	}			}

	return threadCount;			return threadCount;
	}			}

	// CHECK: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}amdgcn_device_isa_selected			// CHECK: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}amdgcn_device_isa_selected
	Show All 28 Lines