commit b72fce1ffd0a0de5b46b486c7030d54cc5d8c225
author Yaxun (Sam) Liu <yaxun.liu@amd.com> Fri Apr 03 09:24:58 2020 -0400
committer Yaxun (Sam) Liu <yaxun.liu@amd.com> Fri Apr 03 09:56:30 2020 -0400
tree f73631e65757812a0fbe80ebcd85f1f5ad25a478
parent 7521f3c2f068bcef9ef57d15e6bd5c6b0b619d4e

Fix __builtin_amdgcn_workgroup_size_x/y/z return type

https://reviews.llvm.org/D77390

clang/include/clang/Basic/BuiltinsAMDGPU.def

diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index e5b256c..b42c8a7 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -33,9 +33,9 @@
 BUILTIN(__builtin_amdgcn_workitem_id_y, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc")
 
-BUILTIN(__builtin_amdgcn_workgroup_size_x, "Ui", "nc")
-BUILTIN(__builtin_amdgcn_workgroup_size_y, "Ui", "nc")
-BUILTIN(__builtin_amdgcn_workgroup_size_z, "Ui", "nc")
+BUILTIN(__builtin_amdgcn_workgroup_size_x, "Us", "nc")
+BUILTIN(__builtin_amdgcn_workgroup_size_y, "Us", "nc")
+BUILTIN(__builtin_amdgcn_workgroup_size_z, "Us", "nc")
 
 BUILTIN(__builtin_amdgcn_mbcnt_hi, "UiUiUi", "nc")
 BUILTIN(__builtin_amdgcn_mbcnt_lo, "UiUiUi", "nc")

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 9d7916c..8f2f149 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -538,7 +538,7 @@
 void test_get_workgroup_size(int d, global int *out)
 {
 	switch (d) {
-	case 0: *out = __builtin_amdgcn_workgroup_size_x(); break;
+	case 0: *out = __builtin_amdgcn_workgroup_size_x() + 1; break;
 	case 1: *out = __builtin_amdgcn_workgroup_size_y(); break;
 	case 2: *out = __builtin_amdgcn_workgroup_size_z(); break;
 	default: *out = 0;