Fix __builtin_amdgcn_workgroup_size_x/y/z return type

https://reviews.llvm.org/D77390

diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index e5b256c..b42c8a7 100644 (file)
--- 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 9d7916c..8f2f149 100644 (file)
--- 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;