clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store.cl

   1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
   2 // REQUIRES: amdgpu-registered-target
   3 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s
   4
   5 typedef unsigned char u8;
   6 typedef unsigned short u16;
   7 typedef unsigned int u32;
   8 typedef unsigned int v2u32 __attribute__((ext_vector_type(2)));
   9 typedef unsigned int v3u32 __attribute__((ext_vector_type(3)));
  10 typedef unsigned int v4u32 __attribute__((ext_vector_type(4)));
  11
  12 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b8(
  13 // CHECK-NEXT:  entry:
  14 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.i8(i8 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
  15 // CHECK-NEXT:    ret void
  16 //
  17 void test_amdgcn_raw_ptr_buffer_store_b8(u8 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
  18   __builtin_amdgcn_raw_buffer_store_b8(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
  19 }
  20
  21 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b16(
  22 // CHECK-NEXT:  entry:
  23 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.i16(i16 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
  24 // CHECK-NEXT:    ret void
  25 //
  26 void test_amdgcn_raw_ptr_buffer_store_b16(u16 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
  27   __builtin_amdgcn_raw_buffer_store_b16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
  28 }
  29
  30 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b32(
  31 // CHECK-NEXT:  entry:
  32 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
  33 // CHECK-NEXT:    ret void
  34 //
  35 void test_amdgcn_raw_ptr_buffer_store_b32(u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
  36   __builtin_amdgcn_raw_buffer_store_b32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
  37 }
  38
  39 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b64(
  40 // CHECK-NEXT:  entry:
  41 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2i32(<2 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
  42 // CHECK-NEXT:    ret void
  43 //
  44 void test_amdgcn_raw_ptr_buffer_store_b64(v2u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
  45   __builtin_amdgcn_raw_buffer_store_b64(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
  46 }
  47
  48 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b96(
  49 // CHECK-NEXT:  entry:
  50 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.v3i32(<3 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
  51 // CHECK-NEXT:    ret void
  52 //
  53 void test_amdgcn_raw_ptr_buffer_store_b96(v3u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
  54   __builtin_amdgcn_raw_buffer_store_b96(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
  55 }
  56
  57 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b128(
  58 // CHECK-NEXT:  entry:
  59 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4i32(<4 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
  60 // CHECK-NEXT:    ret void
  61 //
  62 void test_amdgcn_raw_ptr_buffer_store_b128(v4u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
  63   __builtin_amdgcn_raw_buffer_store_b128(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
  64 }
  65
  66 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b8_non_const_offset(
  67 // CHECK-NEXT:  entry:
  68 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.i8(i8 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)
  69 // CHECK-NEXT:    ret void
  70 //
  71 void test_amdgcn_raw_ptr_buffer_store_b8_non_const_offset(u8 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
  72   __builtin_amdgcn_raw_buffer_store_b8(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);
  73 }
  74
  75 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b16_non_const_offset(
  76 // CHECK-NEXT:  entry:
  77 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.i16(i16 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)
  78 // CHECK-NEXT:    ret void
  79 //
  80 void test_amdgcn_raw_ptr_buffer_store_b16_non_const_offset(u16 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
  81   __builtin_amdgcn_raw_buffer_store_b16(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);
  82 }
  83
  84 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b32_non_const_offset(
  85 // CHECK-NEXT:  entry:
  86 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)
  87 // CHECK-NEXT:    ret void
  88 //
  89 void test_amdgcn_raw_ptr_buffer_store_b32_non_const_offset(u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
  90   __builtin_amdgcn_raw_buffer_store_b32(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);
  91 }
  92
  93 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b64_non_const_offset(
  94 // CHECK-NEXT:  entry:
  95 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2i32(<2 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)
  96 // CHECK-NEXT:    ret void
  97 //
  98 void test_amdgcn_raw_ptr_buffer_store_b64_non_const_offset(v2u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
  99   __builtin_amdgcn_raw_buffer_store_b64(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);
 100 }
 101
 102 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b96_non_const_offset(
 103 // CHECK-NEXT:  entry:
 104 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.v3i32(<3 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)
 105 // CHECK-NEXT:    ret void
 106 //
 107 void test_amdgcn_raw_ptr_buffer_store_b96_non_const_offset(v3u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
 108   __builtin_amdgcn_raw_buffer_store_b96(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);
 109 }
 110
 111 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b128_non_const_offset(
 112 // CHECK-NEXT:  entry:
 113 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4i32(<4 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)
 114 // CHECK-NEXT:    ret void
 115 //
 116 void test_amdgcn_raw_ptr_buffer_store_b128_non_const_offset(v4u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
 117   __builtin_amdgcn_raw_buffer_store_b128(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);
 118 }
 119
 120 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b8_non_const_soffset(
 121 // CHECK-NEXT:  entry:
 122 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.i8(i8 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)
 123 // CHECK-NEXT:    ret void
 124 //
 125 void test_amdgcn_raw_ptr_buffer_store_b8_non_const_soffset(u8 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
 126   __builtin_amdgcn_raw_buffer_store_b8(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);
 127 }
 128
 129 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b16_non_const_soffset(
 130 // CHECK-NEXT:  entry:
 131 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.i16(i16 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)
 132 // CHECK-NEXT:    ret void
 133 //
 134 void test_amdgcn_raw_ptr_buffer_store_b16_non_const_soffset(u16 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
 135   __builtin_amdgcn_raw_buffer_store_b16(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);
 136 }
 137
 138 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b32_non_const_soffset(
 139 // CHECK-NEXT:  entry:
 140 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)
 141 // CHECK-NEXT:    ret void
 142 //
 143 void test_amdgcn_raw_ptr_buffer_store_b32_non_const_soffset(u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
 144   __builtin_amdgcn_raw_buffer_store_b32(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);
 145 }
 146
 147 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b64_non_const_soffset(
 148 // CHECK-NEXT:  entry:
 149 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2i32(<2 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)
 150 // CHECK-NEXT:    ret void
 151 //
 152 void test_amdgcn_raw_ptr_buffer_store_b64_non_const_soffset(v2u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
 153   __builtin_amdgcn_raw_buffer_store_b64(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);
 154 }
 155
 156 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b96_non_const_soffset(
 157 // CHECK-NEXT:  entry:
 158 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.v3i32(<3 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)
 159 // CHECK-NEXT:    ret void
 160 //
 161 void test_amdgcn_raw_ptr_buffer_store_b96_non_const_soffset(v3u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
 162   __builtin_amdgcn_raw_buffer_store_b96(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);
 163 }
 164
 165 // CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b128_non_const_soffset(
 166 // CHECK-NEXT:  entry:
 167 // CHECK-NEXT:    tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4i32(<4 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)
 168 // CHECK-NEXT:    ret void
 169 //
 170 void test_amdgcn_raw_ptr_buffer_store_b128_non_const_soffset(v4u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
 171   __builtin_amdgcn_raw_buffer_store_b128(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);
 172 }