| // 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 -S -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1200 |
| |
| typedef int v2i __attribute__((ext_vector_type(2))); |
| typedef half v8h __attribute__((ext_vector_type(8))); |
| typedef short v8s __attribute__((ext_vector_type(8))); |
| |
| // Wave32 |
| |
| // |
| // amdgcn_global_load_tr |
| // |
| |
| // CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_v2i32( |
| // CHECK-GFX1200-NEXT: entry: |
| // CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.global.load.tr.v2i32(ptr addrspace(1) [[INPTR:%.*]]) |
| // CHECK-GFX1200-NEXT: ret <2 x i32> [[TMP0]] |
| // |
| v2i test_amdgcn_global_load_tr_v2i32(global v2i* inptr) |
| { |
| return __builtin_amdgcn_global_load_tr_v2i32(inptr); |
| } |
| |
| // |
| // amdgcn_global_load_tr |
| // |
| |
| // CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_v8i16( |
| // CHECK-GFX1200-NEXT: entry: |
| // CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call <8 x i16> @llvm.amdgcn.global.load.tr.v8i16(ptr addrspace(1) [[INPTR:%.*]]) |
| // CHECK-GFX1200-NEXT: ret <8 x i16> [[TMP0]] |
| // |
| v8s test_amdgcn_global_load_tr_v8i16(global v8s* inptr) |
| { |
| return __builtin_amdgcn_global_load_tr_v8i16(inptr); |
| } |
| |
| // CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_v8f16( |
| // CHECK-GFX1200-NEXT: entry: |
| // CHECK-GFX1200-NEXT: [[TMP0:%.*]] = tail call <8 x half> @llvm.amdgcn.global.load.tr.v8f16(ptr addrspace(1) [[INPTR:%.*]]) |
| // CHECK-GFX1200-NEXT: ret <8 x half> [[TMP0]] |
| // |
| v8h test_amdgcn_global_load_tr_v8f16(global v8h* inptr) |
| { |
| return __builtin_amdgcn_global_load_tr_v8f16(inptr); |
| } |
| |
