clang/test/CodeGenOpenCL/builtins-r600.cl

   1 // REQUIRES: amdgpu-registered-target
   2 // RUN: %clang_cc1 -triple r600-unknown-unknown -target-cpu cypress -S -emit-llvm -o - %s | FileCheck %s
   3
   4 // CHECK-LABEL: @test_recipsqrt_ieee_f32
   5 // CHECK: call float @llvm.r600.recipsqrt.ieee.f32
   6 void test_recipsqrt_ieee_f32(global float* out, float a)
   7 {
   8   *out = __builtin_r600_recipsqrt_ieeef(a);
   9 }
  10
  11 #if cl_khr_fp64
  12 // XCHECK-LABEL: @test_recipsqrt_ieee_f64
  13 // XCHECK: call double @llvm.r600.recipsqrt.ieee.f64
  14 void test_recipsqrt_ieee_f64(global double* out, double a)
  15 {
  16   *out = __builtin_r600_recipsqrt_ieee(a);
  17 }
  18 #endif
  19
  20 // CHECK-LABEL: @test_implicitarg_ptr
  21 // CHECK: call ptr addrspace(7) @llvm.r600.implicitarg.ptr()
  22 void test_implicitarg_ptr(__attribute__((address_space(7))) unsigned char ** out)
  23 {
  24   *out = __builtin_r600_implicitarg_ptr();
  25 }
  26
  27 // CHECK-LABEL: @test_get_group_id(
  28 // CHECK: tail call i32 @llvm.r600.read.tgid.x()
  29 // CHECK: tail call i32 @llvm.r600.read.tgid.y()
  30 // CHECK: tail call i32 @llvm.r600.read.tgid.z()
  31 void test_get_group_id(int d, global int *out)
  32 {
  33         switch (d) {
  34         case 0: *out = __builtin_r600_read_tgid_x(); break;
  35         case 1: *out = __builtin_r600_read_tgid_y(); break;
  36         case 2: *out = __builtin_r600_read_tgid_z(); break;
  37         default: *out = 0;
  38         }
  39 }
  40
  41 // CHECK-LABEL: @test_get_local_id(
  42 // CHECK: tail call i32 @llvm.r600.read.tidig.x(), !range [[WI_RANGE:![0-9]*]]
  43 // CHECK: tail call i32 @llvm.r600.read.tidig.y(), !range [[WI_RANGE]]
  44 // CHECK: tail call i32 @llvm.r600.read.tidig.z(), !range [[WI_RANGE]]
  45 void test_get_local_id(int d, global int *out)
  46 {
  47         switch (d) {
  48         case 0: *out = __builtin_r600_read_tidig_x(); break;
  49         case 1: *out = __builtin_r600_read_tidig_y(); break;
  50         case 2: *out = __builtin_r600_read_tidig_z(); break;
  51         default: *out = 0;
  52         }
  53 }
  54
  55 // CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}