This is an archive of the discontinued LLVM Phabricator instance.

Differential D67338

AMDGPU: Add immarg tto llvm.amdgcn.init.exec.from.input
ClosedPublic

Authored by arsenm on Sep 8 2019, 5:53 PM.

Download Raw Diff

Details

Reviewers

mareko
nhaehnle

Summary

As far as I can tell this has to be a constant.

Diff Detail

Event Timeline

arsenm created this revision.Sep 8 2019, 5:53 PM

Herald added subscribers: t-tye, tpr, dstuttard and 4 others. · View Herald TranscriptSep 8 2019, 5:53 PM

Yes.

This revision is now accepted and ready to land.Sep 12 2019, 1:56 PM

r371793

Revision Contents

Path

Size

include/

llvm/

IR/

IntrinsicsAMDGPU.td

2 lines

Diff 219282

include/llvm/IR/IntrinsicsAMDGPU.td

	Show First 20 Lines • Show All 179 Lines • ▼ Show 20 Lines		def int_amdgcn_init_exec : Intrinsic<[],
	[IntrConvergent, ImmArg<0>]>;			[IntrConvergent, ImmArg<0>]>;

	// Set EXEC according to a thread count packed in an SGPR input:			// Set EXEC according to a thread count packed in an SGPR input:
	// thread_count = (input >> bitoffset) & 0x7f;			// thread_count = (input >> bitoffset) & 0x7f;
	// This is always moved to the beginning of the basic block.			// This is always moved to the beginning of the basic block.
	def int_amdgcn_init_exec_from_input : Intrinsic<[],			def int_amdgcn_init_exec_from_input : Intrinsic<[],
	[llvm_i32_ty, // 32-bit SGPR input			[llvm_i32_ty, // 32-bit SGPR input
	llvm_i32_ty], // bit offset of the thread count			llvm_i32_ty], // bit offset of the thread count
	[IntrConvergent]>;			[IntrConvergent, ImmArg<1>]>;

	def int_amdgcn_wavefrontsize :			def int_amdgcn_wavefrontsize :
	GCCBuiltin<"__builtin_amdgcn_wavefrontsize">,			GCCBuiltin<"__builtin_amdgcn_wavefrontsize">,
	Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;			Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;


	//===----------------------------------------------------------------------===//			//===----------------------------------------------------------------------===//
	// Instruction Intrinsics			// Instruction Intrinsics
	▲ Show 20 Lines • Show All 1,638 Lines • Show Last 20 Lines