// 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 gfx1200 -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1200

typedef short  v4s   __attribute__((ext_vector_type(4)));
typedef  half  v4h   __attribute__((ext_vector_type(4)));
typedef __bf16 v4y   __attribute__((ext_vector_type(4)));

// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_b64_i32(
// CHECK-GFX1200-NEXT:  entry:
// CHECK-GFX1200-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.tr.b64.i32(ptr addrspace(1) [[INPTR:%.*]])
// CHECK-GFX1200-NEXT:    ret i32 [[TMP0]]
//
int test_amdgcn_global_load_tr_b64_i32(global int* inptr)
{
  return __builtin_amdgcn_global_load_tr_b64_i32(inptr);
}

// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_b128_v4i16(
// CHECK-GFX1200-NEXT:  entry:
// CHECK-GFX1200-NEXT:    [[TMP0:%.*]] = tail call <4 x i16> @llvm.amdgcn.global.load.tr.b128.v4i16(ptr addrspace(1) [[INPTR:%.*]])
// CHECK-GFX1200-NEXT:    ret <4 x i16> [[TMP0]]
//
v4s test_amdgcn_global_load_tr_b128_v4i16(global v4s* inptr)
{
  return __builtin_amdgcn_global_load_tr_b128_v4i16(inptr);
}

// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_b128_v4f16(
// CHECK-GFX1200-NEXT:  entry:
// CHECK-GFX1200-NEXT:    [[TMP0:%.*]] = tail call <4 x half> @llvm.amdgcn.global.load.tr.b128.v4f16(ptr addrspace(1) [[INPTR:%.*]])
// CHECK-GFX1200-NEXT:    ret <4 x half> [[TMP0]]
//
v4h test_amdgcn_global_load_tr_b128_v4f16(global v4h* inptr)
{
  return __builtin_amdgcn_global_load_tr_b128_v4f16(inptr);
}

// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_b128_v4bf16(
// CHECK-GFX1200-NEXT:  entry:
// CHECK-GFX1200-NEXT:    [[TMP0:%.*]] = tail call <4 x bfloat> @llvm.amdgcn.global.load.tr.b128.v4bf16(ptr addrspace(1) [[INPTR:%.*]])
// CHECK-GFX1200-NEXT:    ret <4 x bfloat> [[TMP0]]
//
v4y test_amdgcn_global_load_tr_b128_v4bf16(global v4y* inptr)
{
  return __builtin_amdgcn_global_load_tr_b128_v4bf16(inptr);
}