[Clang][AMDGPU] Use unsigned data type for __builtin_amdgcn_raw_buffer_store_*

[shiltian]

No description provided.

[shiltian]

• [Clang][AMDGPU] Use unsigned data type for __builtin_amdgcn_raw_buffer_store_* #99546 👈
• main

This stack of pull requests is managed by Graphite. Learn more about stacking.

Join @shiltian and the rest of your teammates on Graphite

[llvmbot]

@llvm/pr-subscribers-clang

@llvm/pr-subscribers-backend-amdgpu

Author: Shilei Tian (shiltian)

Changes

---

Full diff: https://github.com/llvm/llvm-project/pull/99546.diff

3 Files Affected:

• (modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+6-6)
• (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store.cl (+6-6)
• (modified) clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-store-error.cl (+6-6)

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}}