[clang] [Clang][AMDGPU] Use unsigned data type for `__builtin_amdgcn_raw_buffer_store_*` (PR #99546)

[Shilei Tian via cfe-commits]

https://github.com/shiltian created https://github.com/llvm/llvm-project/pull/99546

None

>From e93ddb989d3360b1606247200ab15e21a1dc222c Mon Sep 17 00:00:00 2001
From: Shilei Tian <i at tianshilei.me>
Date: Thu, 18 Jul 2024 14:28:27 -0400
Subject: [PATCH] [Clang][AMDGPU] Use unsigned data type for
 `__builtin_amdgcn_raw_buffer_store_*`

---
 clang/include/clang/Basic/BuiltinsAMDGPU.def         | 12 ++++++------
 .../builtins-amdgcn-raw-buffer-store.cl              | 12 ++++++------
 .../builtins-amdgcn-raw-buffer-store-error.cl        | 12 ++++++------
 3 files changed, 18 insertions(+), 18 deletions(-)

diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index e62315eea277a..d9c1f7da66d78 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -149,12 +149,12 @@ BUILTIN(__builtin_amdgcn_mqsad_pk_u16_u8, "WUiWUiUiWUi", "nc")
 BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiWUiUiV4Ui", "nc")
 
 BUILTIN(__builtin_amdgcn_make_buffer_rsrc, "Qbv*sii", "nc")
-BUILTIN(__builtin_amdgcn_raw_buffer_store_b8, "vcQbiiIi", "n")
-BUILTIN(__builtin_amdgcn_raw_buffer_store_b16, "vsQbiiIi", "n")
-BUILTIN(__builtin_amdgcn_raw_buffer_store_b32, "viQbiiIi", "n")
-BUILTIN(__builtin_amdgcn_raw_buffer_store_b64, "vV2iQbiiIi", "n")
-BUILTIN(__builtin_amdgcn_raw_buffer_store_b96, "vV3iQbiiIi", "n")
-BUILTIN(__builtin_amdgcn_raw_buffer_store_b128, "vV4iQbiiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_b8, "vUcQbiiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_b16, "vUsQbiiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_b32, "vUiQbiiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_b64, "vV2UiQbiiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_b96, "vV3UiQbiiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_b128, "vV4UiQbiiIi", "n")
 
 //===----------------------------------------------------------------------===//
 // Ballot builtins.
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store.cl
index 37975d59730c5..361531381ac9e 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store.cl
@@ -2,12 +2,12 @@
 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s
 
-typedef char i8;
-typedef short i16;
-typedef int i32;
-typedef int i64 __attribute__((ext_vector_type(2)));
-typedef int i96 __attribute__((ext_vector_type(3)));
-typedef int i128 __attribute__((ext_vector_type(4)));
+typedef unsigned char i8;
+typedef unsigned short i16;
+typedef unsigned int i32;
+typedef unsigned int i64 __attribute__((ext_vector_type(2)));
+typedef unsigned int i96 __attribute__((ext_vector_type(3)));
+typedef unsigned int i128 __attribute__((ext_vector_type(4)));
 
 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b8(
 // CHECK-NEXT:  entry:
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-store-error.cl b/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-store-error.cl
index 356c031640483..3b78900420c9a 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-store-error.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-store-error.cl
@@ -3,12 +3,12 @@
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
-typedef char i8;
-typedef short i16;
-typedef int i32;
-typedef int i64 __attribute__((ext_vector_type(2)));
-typedef int i96 __attribute__((ext_vector_type(3)));
-typedef int i128 __attribute__((ext_vector_type(4)));
+typedef unsigned char i8;
+typedef unsigned short i16;
+typedef unsigned int i32;
+typedef unsigned int i64 __attribute__((ext_vector_type(2)));
+typedef unsigned int i96 __attribute__((ext_vector_type(3)));
+typedef unsigned int i128 __attribute__((ext_vector_type(4)));
 
 void test_amdgcn_raw_ptr_buffer_store_b8(i8 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset, int aux) {
   __builtin_amdgcn_raw_buffer_store_b8(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_buffer_store_b8' must be a constant integer}}