This is an archive of the discontinued LLVM Phabricator instance.

AMDGPU: Make amdgcn.s.get.waveid.in.workgroup inaccessiblememonly
ClosedPublic

Authored by arsenm on Jun 24 2019, 6:20 PM.

Reviewers

rampitec

Summary

This should probably be readnone, even though the instruction looks
like a load.

LGTM

This revision is now accepted and ready to land.Jun 24 2019, 7:31 PM

r364304

Path

Size

include/

llvm/

IR/

2 lines

	Show First 20 Lines • Show All 1,518 Lines • ▼ Show 20 Lines
	// the lanes to read from.			// the lanes to read from.
	def int_amdgcn_mov_dpp8 :			def int_amdgcn_mov_dpp8 :
	Intrinsic<[llvm_anyint_ty],			Intrinsic<[llvm_anyint_ty],
	[LLVMMatchType<0>, llvm_i32_ty],			[LLVMMatchType<0>, llvm_i32_ty],
	[IntrNoMem, IntrConvergent, ImmArg<1>]>;			[IntrNoMem, IntrConvergent, ImmArg<1>]>;

	def int_amdgcn_s_get_waveid_in_workgroup :			def int_amdgcn_s_get_waveid_in_workgroup :
	GCCBuiltin<"__builtin_amdgcn_s_get_waveid_in_workgroup">,			GCCBuiltin<"__builtin_amdgcn_s_get_waveid_in_workgroup">,
	Intrinsic<[llvm_i32_ty], [], [IntrReadMem]>;			Intrinsic<[llvm_i32_ty], [], [IntrReadMem, IntrInaccessibleMemOnly]>;

	//===----------------------------------------------------------------------===//			//===----------------------------------------------------------------------===//
	// Deep learning intrinsics.			// Deep learning intrinsics.
	//===----------------------------------------------------------------------===//			//===----------------------------------------------------------------------===//

	// f32 %r = llvm.amdgcn.fdot2(v2f16 %a, v2f16 %b, f32 %c, i1 %clamp)			// f32 %r = llvm.amdgcn.fdot2(v2f16 %a, v2f16 %b, f32 %c, i1 %clamp)
	// %r = %a[0] * %b[0] + %a[1] * %b[1] + %c			// %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
	def int_amdgcn_fdot2 :			def int_amdgcn_fdot2 :
	▲ Show 20 Lines • Show All 136 Lines • Show Last 20 Lines