test/CodeGenOpenCL/builtins-amdgcn-vi.cl

   1 // REQUIRES: amdgpu-registered-target
   2 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s
   3
   4 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
   5
   6 typedef unsigned long ulong;
   7
   8 // CHECK-LABEL: @test_div_fixup_f16
   9 // CHECK: call half @llvm.amdgcn.div.fixup.f16
  10 void test_div_fixup_f16(global half* out, half a, half b, half c)
  11 {
  12   *out = __builtin_amdgcn_div_fixuph(a, b, c);
  13 }
  14
  15 // CHECK-LABEL: @test_rcp_f16
  16 // CHECK: call half @llvm.amdgcn.rcp.f16
  17 void test_rcp_f16(global half* out, half a)
  18 {
  19   *out = __builtin_amdgcn_rcph(a);
  20 }
  21
  22 // CHECK-LABEL: @test_rsq_f16
  23 // CHECK: call half @llvm.amdgcn.rsq.f16
  24 void test_rsq_f16(global half* out, half a)
  25 {
  26   *out = __builtin_amdgcn_rsqh(a);
  27 }
  28
  29 // CHECK-LABEL: @test_sin_f16
  30 // CHECK: call half @llvm.amdgcn.sin.f16
  31 void test_sin_f16(global half* out, half a)
  32 {
  33   *out = __builtin_amdgcn_sinh(a);
  34 }
  35
  36 // CHECK-LABEL: @test_cos_f16
  37 // CHECK: call half @llvm.amdgcn.cos.f16
  38 void test_cos_f16(global half* out, half a)
  39 {
  40   *out = __builtin_amdgcn_cosh(a);
  41 }
  42
  43 // CHECK-LABEL: @test_ldexp_f16
  44 // CHECK: call half @llvm.amdgcn.ldexp.f16
  45 void test_ldexp_f16(global half* out, half a, int b)
  46 {
  47   *out = __builtin_amdgcn_ldexph(a, b);
  48 }
  49
  50 // CHECK-LABEL: @test_frexp_mant_f16
  51 // CHECK: call half @llvm.amdgcn.frexp.mant.f16
  52 void test_frexp_mant_f16(global half* out, half a)
  53 {
  54   *out = __builtin_amdgcn_frexp_manth(a);
  55 }
  56
  57 // CHECK-LABEL: @test_frexp_exp_f16
  58 // CHECK: call i16 @llvm.amdgcn.frexp.exp.i16.f16
  59 void test_frexp_exp_f16(global short* out, half a)
  60 {
  61   *out = __builtin_amdgcn_frexp_exph(a);
  62 }
  63
  64 // CHECK-LABEL: @test_fract_f16
  65 // CHECK: call half @llvm.amdgcn.fract.f16
  66 void test_fract_f16(global half* out, half a)
  67 {
  68   *out = __builtin_amdgcn_fracth(a);
  69 }
  70
  71 // CHECK-LABEL: @test_class_f16
  72 // CHECK: call i1 @llvm.amdgcn.class.f16
  73 void test_class_f16(global half* out, half a, int b)
  74 {
  75   *out = __builtin_amdgcn_classh(a, b);
  76 }
  77
  78 // CHECK-LABEL: @test_s_memrealtime
  79 // CHECK: call i64 @llvm.amdgcn.s.memrealtime()
  80 void test_s_memrealtime(global ulong* out)
  81 {
  82   *out = __builtin_amdgcn_s_memrealtime();
  83 }
  84
  85 // CHECK-LABEL: @test_mov_dpp
  86 // CHECK: call i32 @llvm.amdgcn.mov.dpp.i32(i32 %src, i32 0, i32 0, i32 0, i1 false)
  87 void test_mov_dpp(global int* out, int src)
  88 {
  89   *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);
  90 }
  91
  92 // CHECK-LABEL: @test_ds_fadd
  93 // CHECK: call float @llvm.amdgcn.ds.fadd(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
  94 void test_ds_faddf(local float *out, float src) {
  95   *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, false);
  96 }
  97
  98 // CHECK-LABEL: @test_ds_fmin
  99 // CHECK: call float @llvm.amdgcn.ds.fmin(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
 100 void test_ds_fminf(local float *out, float src) {
 101   *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, false);
 102 }
 103
 104 // CHECK-LABEL: @test_ds_fmax
 105 // CHECK: call float @llvm.amdgcn.ds.fmax(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
 106 void test_ds_fmaxf(local float *out, float src) {
 107   *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, false);
 108 }