[clang] [Clang][AMDGPU] Use unsigned data type for `__builtin_amdgcn_raw_buffer_store_*` (PR #99546)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Jul 18 11:36:43 PDT 2024
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Shilei Tian (shiltian)
<details>
<summary>Changes</summary>
---
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
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}}
``````````
</details>
https://github.com/llvm/llvm-project/pull/99546
More information about the cfe-commits
mailing list