[clang] b72fce1 - Fix __builtin_amdgcn_workgroup_size_x/y/z return type

[Yaxun Liu via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2020-04-03T09:56:30-04:00
New Revision: b72fce1ffd0a0de5b46b486c7030d54cc5d8c225

URL: https://github.com/llvm/llvm-project/commit/b72fce1ffd0a0de5b46b486c7030d54cc5d8c225
DIFF: https://github.com/llvm/llvm-project/commit/b72fce1ffd0a0de5b46b486c7030d54cc5d8c225.diff

LOG: Fix __builtin_amdgcn_workgroup_size_x/y/z return type

https://reviews.llvm.org/D77390

Added: 
    

Modified: 
    clang/include/clang/Basic/BuiltinsAMDGPU.def
    clang/test/CodeGenOpenCL/builtins-amdgcn.cl

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index e5b256c07a49..b42c8a77c4bc 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -33,9 +33,9 @@ BUILTIN(__builtin_amdgcn_workitem_id_x, "Ui", "nc")
 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")

diff  --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 9d7916c236c5..8f2f149103b3 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -538,7 +538,7 @@ void test_get_local_id(int d, global int *out)
 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;