clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl

   1 // REQUIRES: amdgpu-registered-target
   2 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s
   3 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s
   4 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
   5 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s
   6
   7 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
   8
   9 typedef unsigned long ulong;
  10 typedef unsigned int  uint;
  11
  12 // CHECK-LABEL: @test_div_fixup_f16
  13 // CHECK: call half @llvm.amdgcn.div.fixup.f16
  14 void test_div_fixup_f16(global half* out, half a, half b, half c)
  15 {
  16   *out = __builtin_amdgcn_div_fixuph(a, b, c);
  17 }
  18
  19 // CHECK-LABEL: @test_rcp_f16
  20 // CHECK: call half @llvm.amdgcn.rcp.f16
  21 void test_rcp_f16(global half* out, half a)
  22 {
  23   *out = __builtin_amdgcn_rcph(a);
  24 }
  25
  26 // CHECK-LABEL: @test_sqrt_f16
  27 // CHECK: call half @llvm.amdgcn.sqrt.f16
  28 void test_sqrt_f16(global half* out, half a)
  29 {
  30   *out = __builtin_amdgcn_sqrth(a);
  31 }
  32
  33 // CHECK-LABEL: @test_rsq_f16
  34 // CHECK: call half @llvm.amdgcn.rsq.f16
  35 void test_rsq_f16(global half* out, half a)
  36 {
  37   *out = __builtin_amdgcn_rsqh(a);
  38 }
  39
  40 // CHECK-LABEL: @test_sin_f16
  41 // CHECK: call half @llvm.amdgcn.sin.f16
  42 void test_sin_f16(global half* out, half a)
  43 {
  44   *out = __builtin_amdgcn_sinh(a);
  45 }
  46
  47 // CHECK-LABEL: @test_cos_f16
  48 // CHECK: call half @llvm.amdgcn.cos.f16
  49 void test_cos_f16(global half* out, half a)
  50 {
  51   *out = __builtin_amdgcn_cosh(a);
  52 }
  53
  54 // CHECK-LABEL: @test_ldexp_f16
  55 // CHECK: call half @llvm.amdgcn.ldexp.f16
  56 void test_ldexp_f16(global half* out, half a, int b)
  57 {
  58   *out = __builtin_amdgcn_ldexph(a, b);
  59 }
  60
  61 // CHECK-LABEL: @test_frexp_mant_f16
  62 // CHECK: call half @llvm.amdgcn.frexp.mant.f16
  63 void test_frexp_mant_f16(global half* out, half a)
  64 {
  65   *out = __builtin_amdgcn_frexp_manth(a);
  66 }
  67
  68 // CHECK-LABEL: @test_frexp_exp_f16
  69 // CHECK: call i16 @llvm.amdgcn.frexp.exp.i16.f16
  70 void test_frexp_exp_f16(global short* out, half a)
  71 {
  72   *out = __builtin_amdgcn_frexp_exph(a);
  73 }
  74
  75 // CHECK-LABEL: @test_fract_f16
  76 // CHECK: call half @llvm.amdgcn.fract.f16
  77 void test_fract_f16(global half* out, half a)
  78 {
  79   *out = __builtin_amdgcn_fracth(a);
  80 }
  81
  82 // CHECK-LABEL: @test_class_f16
  83 // CHECK: call i1 @llvm.amdgcn.class.f16
  84 void test_class_f16(global half* out, half a, int b)
  85 {
  86   *out = __builtin_amdgcn_classh(a, b);
  87 }
  88
  89 // CHECK-LABEL: @test_s_memrealtime
  90 // CHECK: call i64 @llvm.amdgcn.s.memrealtime()
  91 void test_s_memrealtime(global ulong* out)
  92 {
  93   *out = __builtin_amdgcn_s_memrealtime();
  94 }
  95
  96 // CHECK-LABEL: @test_s_dcache_wb()
  97 // CHECK: call void @llvm.amdgcn.s.dcache.wb()
  98 void test_s_dcache_wb()
  99 {
 100   __builtin_amdgcn_s_dcache_wb();
 101 }
 102
 103 // CHECK-LABEL: @test_mov_dpp
 104 // CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 undef, i32 %src, i32 0, i32 0, i32 0, i1 false)
 105 void test_mov_dpp(global int* out, int src)
 106 {
 107   *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);
 108 }
 109
 110 // CHECK-LABEL: @test_update_dpp
 111 // CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 %arg1, i32 %arg2, i32 0, i32 0, i32 0, i1 false)
 112 void test_update_dpp(global int* out, int arg1, int arg2)
 113 {
 114   *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false);
 115 }
 116
 117 // CHECK-LABEL: @test_ds_fadd
 118 // CHECK: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
 119 void test_ds_faddf(local float *out, float src) {
 120   *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, false);
 121 }
 122
 123 // CHECK-LABEL: @test_ds_fmin
 124 // CHECK: call float @llvm.amdgcn.ds.fmin.f32(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
 125 void test_ds_fminf(local float *out, float src) {
 126   *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, false);
 127 }
 128
 129 // CHECK-LABEL: @test_ds_fmax
 130 // CHECK: call float @llvm.amdgcn.ds.fmax.f32(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
 131 void test_ds_fmaxf(local float *out, float src) {
 132   *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, false);
 133 }
 134
 135 // CHECK-LABEL: @test_s_memtime
 136 // CHECK: call i64 @llvm.amdgcn.s.memtime()
 137 void test_s_memtime(global ulong* out)
 138 {
 139   *out = __builtin_amdgcn_s_memtime();
 140 }
 141
 142 // CHECK-LABEL: @test_perm
 143 // CHECK: call i32 @llvm.amdgcn.perm(i32 %a, i32 %b, i32 %s)
 144 void test_perm(global uint* out, uint a, uint b, uint s)
 145 {
 146   *out = __builtin_amdgcn_perm(a, b, s);
 147 }
 148
 149 // CHECK-LABEL: @test_groupstaticsize
 150 // CHECK: call i32 @llvm.amdgcn.groupstaticsize()
 151 void test_groupstaticsize(global uint* out)
 152 {
 153   *out = __builtin_amdgcn_groupstaticsize();
 154 }