clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl

// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s

#pragma OPENCL EXTENSION cl_khr_fp16 : enable

typedef int v4i32 __attribute__((ext_vector_type(4)));
typedef int v2i32 __attribute__((ext_vector_type(2)));
typedef int v3i32 __attribute__((ext_vector_type(3)));
typedef int v8i32 __attribute__((ext_vector_type(8)));
typedef int v16i32 __attribute__((ext_vector_type(16)));
typedef float v2f32 __attribute__((ext_vector_type(2)));
typedef float v3f32 __attribute__((ext_vector_type(3)));
typedef float v4f32 __attribute__((ext_vector_type(4)));
typedef float v8f32 __attribute__((ext_vector_type(8)));
typedef float v16f32 __attribute__((ext_vector_type(16)));
typedef char v2i8 __attribute__((ext_vector_type(2)));
typedef char v3i8 __attribute__((ext_vector_type(3)));
typedef char v4i8 __attribute__((ext_vector_type(4)));
typedef half v2f16 __attribute__((ext_vector_type(2)));
typedef half v3f16 __attribute__((ext_vector_type(3)));
typedef half v4f16 __attribute__((ext_vector_type(4)));

// CHECK-LABEL: @test_amdgcn_s_buffer_load_i32(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.ptr.s.buffer.load.i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8:![0-9]+]]
// CHECK-NEXT:    ret i32 [[TMP0]]
//
int test_amdgcn_s_buffer_load_i32(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_i32(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_i32_non_const_offset(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.ptr.s.buffer.load.i32(ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret i32 [[TMP0]]
//
int test_amdgcn_s_buffer_load_i32_non_const_offset(__amdgpu_buffer_rsrc_t rsrc, int offset) {
  return __builtin_amdgcn_s_buffer_load_i32(rsrc, offset, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v2i32(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.ptr.s.buffer.load.v2i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <2 x i32> [[TMP0]]
//
v2i32 test_amdgcn_s_buffer_load_v2i32(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v2i32(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v3i32(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <3 x i32> @llvm.amdgcn.ptr.s.buffer.load.v3i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <3 x i32> [[TMP0]]
//
v3i32 test_amdgcn_s_buffer_load_v3i32(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v3i32(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v4i32(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.ptr.s.buffer.load.v4i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <4 x i32> [[TMP0]]
//
v4i32 test_amdgcn_s_buffer_load_v4i32(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v4i32(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v8i32(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <8 x i32> @llvm.amdgcn.ptr.s.buffer.load.v8i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <8 x i32> [[TMP0]]
//
v8i32 test_amdgcn_s_buffer_load_v8i32(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v8i32(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v16i32(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <16 x i32> @llvm.amdgcn.ptr.s.buffer.load.v16i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <16 x i32> [[TMP0]]
//
v16i32 test_amdgcn_s_buffer_load_v16i32(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v16i32(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_f32(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call float @llvm.amdgcn.ptr.s.buffer.load.f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret float [[TMP0]]
//
float test_amdgcn_s_buffer_load_f32(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_f32(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v2f32(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <2 x float> @llvm.amdgcn.ptr.s.buffer.load.v2f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <2 x float> [[TMP0]]
//
v2f32 test_amdgcn_s_buffer_load_v2f32(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v2f32(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v3f32(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <3 x float> @llvm.amdgcn.ptr.s.buffer.load.v3f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <3 x float> [[TMP0]]
//
v3f32 test_amdgcn_s_buffer_load_v3f32(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v3f32(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v4f32(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <4 x float> @llvm.amdgcn.ptr.s.buffer.load.v4f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <4 x float> [[TMP0]]
//
v4f32 test_amdgcn_s_buffer_load_v4f32(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v4f32(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v8f32(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <8 x float> @llvm.amdgcn.ptr.s.buffer.load.v8f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <8 x float> [[TMP0]]
//
v8f32 test_amdgcn_s_buffer_load_v8f32(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v8f32(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v16f32(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <16 x float> @llvm.amdgcn.ptr.s.buffer.load.v16f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <16 x float> [[TMP0]]
//
v16f32 test_amdgcn_s_buffer_load_v16f32(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v16f32(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_i8(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.ptr.s.buffer.load.i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret i8 [[TMP0]]
//
char test_amdgcn_s_buffer_load_i8(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_i8(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_u8(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.ptr.s.buffer.load.i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret i8 [[TMP0]]
//
unsigned char test_amdgcn_s_buffer_load_u8(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_u8(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_i16(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.ptr.s.buffer.load.i16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret i16 [[TMP0]]
//
short test_amdgcn_s_buffer_load_i16(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_i16(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_u16(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.ptr.s.buffer.load.i16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret i16 [[TMP0]]
//
unsigned short test_amdgcn_s_buffer_load_u16(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_u16(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_f16(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call half @llvm.amdgcn.ptr.s.buffer.load.f16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret half [[TMP0]]
//
half test_amdgcn_s_buffer_load_f16(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_f16(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v2f16(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <2 x half> @llvm.amdgcn.ptr.s.buffer.load.v2f16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <2 x half> [[TMP0]]
//
v2f16 test_amdgcn_s_buffer_load_v2f16(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v2f16(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v3f16(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <3 x half> @llvm.amdgcn.ptr.s.buffer.load.v3f16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <3 x half> [[TMP0]]
//
v3f16 test_amdgcn_s_buffer_load_v3f16(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v3f16(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v4f16(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <4 x half> @llvm.amdgcn.ptr.s.buffer.load.v4f16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <4 x half> [[TMP0]]
//
v4f16 test_amdgcn_s_buffer_load_v4f16(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v4f16(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v2i8(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <2 x i8> @llvm.amdgcn.ptr.s.buffer.load.v2i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <2 x i8> [[TMP0]]
//
v2i8 test_amdgcn_s_buffer_load_v2i8(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v2i8(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v3i8(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <3 x i8> @llvm.amdgcn.ptr.s.buffer.load.v3i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <3 x i8> [[TMP0]]
//
v3i8 test_amdgcn_s_buffer_load_v3i8(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v3i8(rsrc, /*offset=*/0, /*aux=*/0);
}

// CHECK-LABEL: @test_amdgcn_s_buffer_load_v4i8(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[TMP0:%.*]] = tail call <4 x i8> @llvm.amdgcn.ptr.s.buffer.load.v4i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT:    ret <4 x i8> [[TMP0]]
//
v4i8 test_amdgcn_s_buffer_load_v4i8(__amdgpu_buffer_rsrc_t rsrc) {
  return __builtin_amdgcn_s_buffer_load_v4i8(rsrc, /*offset=*/0, /*aux=*/0);
}