| 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 i8 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} |
|---|
| 56 | |
|---|