This is an archive of the discontinued LLVM Phabricator instance.

AMDGPU: Add Clang Builtin for v_lerp_u8
ClosedPublic

Authored by wdng on Jul 14 2016, 1:21 PM.

Download Raw Diff

Details

Reviewers

• tstellarAMD
arsenm

Commits

rGea41f356bbeb: AMDGPU: Add Clang Builtin for v_lerp_u8
rC275577: AMDGPU: Add Clang Builtin for v_lerp_u8
rL275577: AMDGPU: Add Clang Builtin for v_lerp_u8

Summary

Add Clang Builtin for v_lerp_u8

Diff Detail

Repository: rL LLVM

Event Timeline

wdng updated this revision to Diff 64032.Jul 14 2016, 1:21 PM

wdng retitled this revision from to AMDGPU: Add Clang Builtin for v_lerp_u8 .

wdng updated this object.

wdng added reviewers: arsenm, • tstellarAMD.

wdng set the repository for this revision to rL LLVM.

wdng added a project: Restricted Project.

Herald added a subscriber: kzhuravl. · View Herald TranscriptJul 14 2016, 1:21 PM

Fixed an indentation issue.

LGTM

This revision is now accepted and ready to land.Jul 14 2016, 3:50 PM

Closed by commit rL275577: AMDGPU: Add Clang Builtin for v_lerp_u8 (authored by wdng). · Explain WhyJul 15 2016, 9:50 AM

This revision was automatically updated to reflect the committed changes.

Revision Contents

Path

Size

include/

clang/

Basic/

BuiltinsAMDGPU.def

1 line

lib/

CodeGen/

CGBuiltin.cpp

2 lines

test/

CodeGenOpenCL/

builtins-amdgcn.cl

7 lines

Diff 64032

include/clang/Basic/BuiltinsAMDGPU.def

Context not available.
	BUILTIN(__builtin_amdgcn_frexp_expf, "if", "nc")	BUILTIN(__builtin_amdgcn_frexp_expf, "if", "nc")
	BUILTIN(__builtin_amdgcn_fract, "dd", "nc")	BUILTIN(__builtin_amdgcn_fract, "dd", "nc")
	BUILTIN(__builtin_amdgcn_fractf, "ff", "nc")	BUILTIN(__builtin_amdgcn_fractf, "ff", "nc")
		BUILTIN(__builtin_amdgcn_lerp, "UiUiUiUi", "nc")
	BUILTIN(__builtin_amdgcn_class, "bdi", "nc")	BUILTIN(__builtin_amdgcn_class, "bdi", "nc")
	BUILTIN(__builtin_amdgcn_classf, "bfi", "nc")	BUILTIN(__builtin_amdgcn_classf, "bfi", "nc")
	BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc")	BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc")
Context not available.

lib/CodeGen/CGBuiltin.cpp

Context not available.
	case AMDGPU::BI__builtin_amdgcn_fract:	case AMDGPU::BI__builtin_amdgcn_fract:
	case AMDGPU::BI__builtin_amdgcn_fractf:	case AMDGPU::BI__builtin_amdgcn_fractf:
	return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_fract);	return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_fract);
		case AMDGPU::BI__builtin_amdgcn_lerp:
		return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_lerp);
	case AMDGPU::BI__builtin_amdgcn_class:	case AMDGPU::BI__builtin_amdgcn_class:
	case AMDGPU::BI__builtin_amdgcn_classf:	case AMDGPU::BI__builtin_amdgcn_classf:
	return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);	return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);
Context not available.

test/CodeGenOpenCL/builtins-amdgcn.cl

Context not available.
	*out = __builtin_amdgcn_fract(a);	*out = __builtin_amdgcn_fract(a);
	}	}

		// CHECK-LABEL: @test_lerp
		// CHECK: call i32 @llvm.amdgcn.lerp
		void test_lerp(global int* out, int a, int b, int c)
		{
		*out = __builtin_amdgcn_lerp(a, b, c);
		}

	// CHECK-LABEL: @test_class_f32	// CHECK-LABEL: @test_class_f32
	// CHECK: call i1 @llvm.amdgcn.class.f32	// CHECK: call i1 @llvm.amdgcn.class.f32
	void test_class_f32(global float* out, float a, int b)	void test_class_f32(global float* out, float a, int b)
Context not available.