// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1250

typedef int    v4i   __attribute__((ext_vector_type(4)));
typedef int    v8i   __attribute__((ext_vector_type(8)));

// CHECK-GFX1250-LABEL: @test_amdgcn_tensor_load_to_lds(
// CHECK-GFX1250-NEXT:  entry:
// CHECK-GFX1250-NEXT:    tail call void @llvm.amdgcn.tensor.load.to.lds(<4 x i32> [[SG0:%.*]], <8 x i32> [[SG1:%.*]], <4 x i32> [[SG2:%.*]], <4 x i32> [[SG3:%.*]], i32 0)
// CHECK-GFX1250-NEXT:    ret void
//
void test_amdgcn_tensor_load_to_lds(v4i sg0, v8i sg1, v4i sg2, v4i sg3)
{
  __builtin_amdgcn_tensor_load_to_lds(sg0, sg1, sg2, sg3, 0);
}

// CHECK-GFX1250-LABEL: @test_amdgcn_tensor_load_to_lds_d2(
// CHECK-GFX1250-NEXT:  entry:
// CHECK-GFX1250-NEXT:    tail call void @llvm.amdgcn.tensor.load.to.lds.d2(<4 x i32> [[SG0:%.*]], <8 x i32> [[SG1:%.*]], i32 27)
// CHECK-GFX1250-NEXT:    ret void
//
void test_amdgcn_tensor_load_to_lds_d2(v4i sg0, v8i sg1)
{
  __builtin_amdgcn_tensor_load_to_lds_d2(sg0, sg1, 27);
}

// CHECK-GFX1250-LABEL: @test_amdgcn_tensor_store_from_lds(
// CHECK-GFX1250-NEXT:  entry:
// CHECK-GFX1250-NEXT:    tail call void @llvm.amdgcn.tensor.store.from.lds(<4 x i32> [[SG0:%.*]], <8 x i32> [[SG1:%.*]], <4 x i32> [[SG2:%.*]], <4 x i32> [[SG3:%.*]], i32 22)
// CHECK-GFX1250-NEXT:    ret void
//
void test_amdgcn_tensor_store_from_lds(v4i sg0, v8i sg1, v4i sg2, v4i sg3)
{
  __builtin_amdgcn_tensor_store_from_lds(sg0, sg1, sg2, sg3, 22);
}

// CHECK-GFX1250-LABEL: @test_amdgcn_tensor_store_from_lds_d2(
// CHECK-GFX1250-NEXT:  entry:
// CHECK-GFX1250-NEXT:    tail call void @llvm.amdgcn.tensor.store.from.lds.d2(<4 x i32> [[SG0:%.*]], <8 x i32> [[SG1:%.*]], i32 0)
// CHECK-GFX1250-NEXT:    ret void
//
void test_amdgcn_tensor_store_from_lds_d2(v4i sg0, v8i sg1)
{
  __builtin_amdgcn_tensor_store_from_lds_d2(sg0, sg1, 0);
}