clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-w64.cl

   1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
   2 // REQUIRES: amdgpu-registered-target
   3 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1200
   4
   5 typedef short  v4s   __attribute__((ext_vector_type(4)));
   6 typedef  half  v4h   __attribute__((ext_vector_type(4)));
   7 typedef __bf16 v4y   __attribute__((ext_vector_type(4)));
   8
   9 // CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_b64_i32(
  10 // CHECK-GFX1200-NEXT:  entry:
  11 // CHECK-GFX1200-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.tr.b64.i32(ptr addrspace(1) [[INPTR:%.*]])
  12 // CHECK-GFX1200-NEXT:    ret i32 [[TMP0]]
  13 //
  14 int test_amdgcn_global_load_tr_b64_i32(global int* inptr)
  15 {
  16   return __builtin_amdgcn_global_load_tr_b64_i32(inptr);
  17 }
  18
  19 // CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_b128_v4i16(
  20 // CHECK-GFX1200-NEXT:  entry:
  21 // CHECK-GFX1200-NEXT:    [[TMP0:%.*]] = tail call <4 x i16> @llvm.amdgcn.global.load.tr.b128.v4i16(ptr addrspace(1) [[INPTR:%.*]])
  22 // CHECK-GFX1200-NEXT:    ret <4 x i16> [[TMP0]]
  23 //
  24 v4s test_amdgcn_global_load_tr_b128_v4i16(global v4s* inptr)
  25 {
  26   return __builtin_amdgcn_global_load_tr_b128_v4i16(inptr);
  27 }
  28
  29 // CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_b128_v4f16(
  30 // CHECK-GFX1200-NEXT:  entry:
  31 // CHECK-GFX1200-NEXT:    [[TMP0:%.*]] = tail call <4 x half> @llvm.amdgcn.global.load.tr.b128.v4f16(ptr addrspace(1) [[INPTR:%.*]])
  32 // CHECK-GFX1200-NEXT:    ret <4 x half> [[TMP0]]
  33 //
  34 v4h test_amdgcn_global_load_tr_b128_v4f16(global v4h* inptr)
  35 {
  36   return __builtin_amdgcn_global_load_tr_b128_v4f16(inptr);
  37 }
  38
  39 // CHECK-GFX1200-LABEL: @test_amdgcn_global_load_tr_b128_v4bf16(
  40 // CHECK-GFX1200-NEXT:  entry:
  41 // CHECK-GFX1200-NEXT:    [[TMP0:%.*]] = tail call <4 x bfloat> @llvm.amdgcn.global.load.tr.b128.v4bf16(ptr addrspace(1) [[INPTR:%.*]])
  42 // CHECK-GFX1200-NEXT:    ret <4 x bfloat> [[TMP0]]
  43 //
  44 v4y test_amdgcn_global_load_tr_b128_v4bf16(global v4y* inptr)
  45 {
  46   return __builtin_amdgcn_global_load_tr_b128_v4bf16(inptr);
  47 }