This is an archive of the discontinued LLVM Phabricator instance.

Differential D77390

Fix __builtin_amdgcn_workgroup_size_x/y/z return type
ClosedPublic

Authored by yaxunl on Apr 3 2020, 6:27 AM.

Download Raw Diff

Details

Reviewers

b-sumner
arsenm

Diff Detail

Event Timeline

yaxunl created this revision.Apr 3 2020, 6:27 AM

Herald added subscribers: kerbowa, nhaehnle, wdng, jvesely. · View Herald TranscriptApr 3 2020, 6:27 AM

arsenm accepted this revision.Apr 3 2020, 6:38 AM

This revision is now accepted and ready to land.Apr 3 2020, 6:38 AM

LGTM

b72fce1ffd0a0de5b46b486c7030d54cc5d8c225

JonChesterfield added a subscriber: JonChesterfield.Apr 11 2020, 2:03 PM

JonChesterfield added inline comments.

clang/test/CodeGenOpenCL/builtins-amdgcn.cl
541	This looks unrelated to the return type change

yaxunl marked 2 inline comments as done.Apr 11 2020, 6:46 PM

yaxunl added inline comments.

clang/test/CodeGenOpenCL/builtins-amdgcn.cl
541	Without that change, the builtin return type is wrong and clang does not insert conversion for this "add" operation and causes assertion. Therefore this small change makes sure that error will not happen again.

Revision Contents

Path

Size

clang/

include/

clang/

Basic/

BuiltinsAMDGPU.def

6 lines

test/

CodeGenOpenCL/

builtins-amdgcn.cl

2 lines

Diff 254774

clang/include/clang/Basic/BuiltinsAMDGPU.def

Context not available.
	BUILTIN(__builtin_amdgcn_workitem_id_y, "Ui", "nc")			BUILTIN(__builtin_amdgcn_workitem_id_y, "Ui", "nc")
	BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc")			BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc")

	BUILTIN(__builtin_amdgcn_workgroup_size_x, "Ui", "nc")			BUILTIN(__builtin_amdgcn_workgroup_size_x, "Us", "nc")
	BUILTIN(__builtin_amdgcn_workgroup_size_y, "Ui", "nc")			BUILTIN(__builtin_amdgcn_workgroup_size_y, "Us", "nc")
	BUILTIN(__builtin_amdgcn_workgroup_size_z, "Ui", "nc")			BUILTIN(__builtin_amdgcn_workgroup_size_z, "Us", "nc")

	BUILTIN(__builtin_amdgcn_mbcnt_hi, "UiUiUi", "nc")			BUILTIN(__builtin_amdgcn_mbcnt_hi, "UiUiUi", "nc")
	BUILTIN(__builtin_amdgcn_mbcnt_lo, "UiUiUi", "nc")			BUILTIN(__builtin_amdgcn_mbcnt_lo, "UiUiUi", "nc")
Context not available.

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

Context not available.
	void test_get_workgroup_size(int d, global int *out)			void test_get_workgroup_size(int d, global int *out)
	{			{
	switch (d) {			switch (d) {
	case 0: *out = __builtin_amdgcn_workgroup_size_x(); break;			case 0: *out = __builtin_amdgcn_workgroup_size_x() + 1; break;
				JonChesterfieldUnsubmitted Not Done Reply Inline Actions This looks unrelated to the return type change JonChesterfield: This looks unrelated to the return type change
				yaxunlAuthorUnsubmitted Done Reply Inline Actions Without that change, the builtin return type is wrong and clang does not insert conversion for this "add" operation and causes assertion. Therefore this small change makes sure that error will not happen again. yaxunl: Without that change, the builtin return type is wrong and clang does not insert conversion for…
	case 1: *out = __builtin_amdgcn_workgroup_size_y(); break;			case 1: *out = __builtin_amdgcn_workgroup_size_y(); break;
	case 2: *out = __builtin_amdgcn_workgroup_size_z(); break;			case 2: *out = __builtin_amdgcn_workgroup_size_z(); break;
	default: *out = 0;			default: *out = 0;
Context not available.