280d90d0fd
half and bfloat are common types for 16-bit elements. The support of them was original there and dropped due to some reasons. This work adds the support of the float types back.
49 lines
2.1 KiB
Common Lisp
49 lines
2.1 KiB
Common Lisp
// 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 +wavefrontsize32 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1200
|
|
|
|
typedef int v2i __attribute__((ext_vector_type(2)));
|
|
typedef short v8s __attribute__((ext_vector_type(8)));
|
|
typedef half v8h __attribute__((ext_vector_type(8)));
|
|
typedef __bf16 v8y __attribute__((ext_vector_type(8)));
|
|
|
|
// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_b64_v2i32(
|
|
// CHECK-GFX1200-NEXT: entry:
|
|
// CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.global.load.tr.b64.v2i32(ptr addrspace(1) [[INPTR:%.*]])
|
|
// CHECK-GFX1200-NEXT: ret <2 x i32> [[TMP0]]
|
|
//
|
|
v2i test_amdgcn_global_load_tr_b64_v2i32(global v2i* inptr)
|
|
{
|
|
return __builtin_amdgcn_global_load_tr_b64_v2i32(inptr);
|
|
}
|
|
|
|
// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_b128_v8i16(
|
|
// CHECK-GFX1200-NEXT: entry:
|
|
// CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call <8 x i16> @llvm.amdgcn.global.load.tr.b128.v8i16(ptr addrspace(1) [[INPTR:%.*]])
|
|
// CHECK-GFX1200-NEXT: ret <8 x i16> [[TMP0]]
|
|
//
|
|
v8s test_amdgcn_global_load_tr_b128_v8i16(global v8s* inptr)
|
|
{
|
|
return __builtin_amdgcn_global_load_tr_b128_v8i16(inptr);
|
|
}
|
|
|
|
// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_b128_v8f16(
|
|
// CHECK-GFX1200-NEXT: entry:
|
|
// CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call <8 x half> @llvm.amdgcn.global.load.tr.b128.v8f16(ptr addrspace(1) [[INPTR:%.*]])
|
|
// CHECK-GFX1200-NEXT: ret <8 x half> [[TMP0]]
|
|
//
|
|
v8h test_amdgcn_global_load_tr_b128_v8f16(global v8h* inptr)
|
|
{
|
|
return __builtin_amdgcn_global_load_tr_b128_v8f16(inptr);
|
|
}
|
|
|
|
// CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_b128_v8bf16(
|
|
// CHECK-GFX1200-NEXT: entry:
|
|
// CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call <8 x bfloat> @llvm.amdgcn.global.load.tr.b128.v8bf16(ptr addrspace(1) [[INPTR:%.*]])
|
|
// CHECK-GFX1200-NEXT: ret <8 x bfloat> [[TMP0]]
|
|
//
|
|
v8y test_amdgcn_global_load_tr_b128_v8bf16(global v8y* inptr)
|
|
{
|
|
return __builtin_amdgcn_global_load_tr_b128_v8bf16(inptr);
|
|
}
|