Differential D128158 Diff 445965 llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Changeset View

Standalone View

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

	Show First 20 Lines • Show All 231 Lines • ▼ Show 20 Lines
	// MASK = 0x0000 0040: VMEM write instructions may be scheduled across SCHED_BARRIER.			// MASK = 0x0000 0040: VMEM write instructions may be scheduled across SCHED_BARRIER.
	// MASK = 0x0000 0080: ALL DS instructions may be scheduled across SCHED_BARRIER.			// MASK = 0x0000 0080: ALL DS instructions may be scheduled across SCHED_BARRIER.
	// MASK = 0x0000 0100: ALL DS read instructions may be scheduled accoss SCHED_BARRIER.			// MASK = 0x0000 0100: ALL DS read instructions may be scheduled accoss SCHED_BARRIER.
	// MASK = 0x0000 0200: ALL DS write instructions may be scheduled across SCHED_BARRIER.			// MASK = 0x0000 0200: ALL DS write instructions may be scheduled across SCHED_BARRIER.
	def int_amdgcn_sched_barrier : ClangBuiltin<"__builtin_amdgcn_sched_barrier">,			def int_amdgcn_sched_barrier : ClangBuiltin<"__builtin_amdgcn_sched_barrier">,
	Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrConvergent,			Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrConvergent,
	IntrWillReturn]>;			IntrWillReturn]>;

				// The first parameter is a mask that determines the types of instructions that
				// you would like to synchronize around and add to a scheduling group. The
				// values of the mask are defined above for sched_barrier. These instructions
				// will be selected from the bottom up starting from the sched_group_barrier's
				// location during instruction scheduling. The second parameter is the number of
				// matching instructions that will be associated with this sched_group_barrier.
				// The third parameter is an identifier which is used to describe what other
				// sched_group_barriers should be synchronized with.
				def int_amdgcn_sched_group_barrier : ClangBuiltin<"__builtin_amdgcn_sched_group_barrier">,
				Intrinsic<[], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
				[ImmArg<ArgIndex<0>>, ImmArg<ArgIndex<1>>, ImmArg<ArgIndex<2>>, IntrNoMem, IntrHasSideEffects,
				IntrConvergent, IntrWillReturn]>;

	def int_amdgcn_s_waitcnt : ClangBuiltin<"__builtin_amdgcn_s_waitcnt">,			def int_amdgcn_s_waitcnt : ClangBuiltin<"__builtin_amdgcn_s_waitcnt">,
	Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]>;			Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]>;

	def int_amdgcn_div_scale : Intrinsic<			def int_amdgcn_div_scale : Intrinsic<
	// 1st parameter: Numerator			// 1st parameter: Numerator
	// 2nd parameter: Denominator			// 2nd parameter: Denominator
	// 3rd parameter: Select quotient. Must equal Numerator or Denominator.			// 3rd parameter: Select quotient. Must equal Numerator or Denominator.
	// (0 = Denominator, 1 = Numerator).			// (0 = Denominator, 1 = Numerator).
	▲ Show 20 Lines • Show All 2,115 Lines • Show Last 20 Lines