This is an archive of the discontinued LLVM Phabricator instance.

Differential D63421

AMDGPU: Make getreg intrinsic inaccessiblememonly
ClosedPublic

Authored by arsenm on Jun 17 2019, 5:31 AM.

Download Raw Diff

Details

Reviewers

nhaehnle
rampitec

Diff Detail

Event Timeline

arsenm created this revision.Jun 17 2019, 5:31 AM

Herald added subscribers: t-tye, tpr, dstuttard and 4 others. · View Herald TranscriptJun 17 2019, 5:31 AM

LGTM

This revision is now accepted and ready to land.Jun 17 2019, 9:39 AM

r363591

Revision Contents

Path

Size

include/

llvm/

IR/

IntrinsicsAMDGPU.td

2 lines

Diff 205047

include/llvm/IR/IntrinsicsAMDGPU.td

	Show First 20 Lines • Show All 1,165 Lines • ▼ Show 20 Lines
	def int_amdgcn_s_decperflevel :			def int_amdgcn_s_decperflevel :
	GCCBuiltin<"__builtin_amdgcn_s_decperflevel">,			GCCBuiltin<"__builtin_amdgcn_s_decperflevel">,
	Intrinsic<[], [llvm_i32_ty], [ImmArg<0>]> {			Intrinsic<[], [llvm_i32_ty], [ImmArg<0>]> {
	}			}

	def int_amdgcn_s_getreg :			def int_amdgcn_s_getreg :
	GCCBuiltin<"__builtin_amdgcn_s_getreg">,			GCCBuiltin<"__builtin_amdgcn_s_getreg">,
	Intrinsic<[llvm_i32_ty], [llvm_i32_ty],			Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
	[IntrReadMem, IntrSpeculatable, ImmArg<0>]			[IntrInaccessibleMemOnly, IntrReadMem, IntrSpeculatable, ImmArg<0>]
	>;			>;

	// int_amdgcn_s_getpc is provided to allow a specific style of position			// int_amdgcn_s_getpc is provided to allow a specific style of position
	// independent code to determine the high part of its address when it is			// independent code to determine the high part of its address when it is
	// known (through convention) that the code and any data of interest does			// known (through convention) that the code and any data of interest does
	// not cross a 4Gb address boundary. Use for any other purpose may not			// not cross a 4Gb address boundary. Use for any other purpose may not
	// produce the desired results as optimizations may cause code movement,			// produce the desired results as optimizations may cause code movement,
	// especially as we explicitly use IntrNoMem to allow optimizations.			// especially as we explicitly use IntrNoMem to allow optimizations.
	▲ Show 20 Lines • Show All 431 Lines • Show Last 20 Lines