// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa -cl-std=CL1.2 \
// RUN:     -emit-llvm -o - | FileCheck %s

// CHECK-LABEL: @test_builtin_amdgcn_cvt_off_f32_i4_ui(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-NEXT:    store i32 [[N:%.*]], ptr addrspace(5) [[N_ADDR]], align 4
// CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspace(5) [[N_ADDR]], align 4
// CHECK-NEXT:    [[TMP1:%.*]] = call float @llvm.amdgcn.cvt.off.f32.i4(i32 [[TMP0]])
// CHECK-NEXT:    ret float [[TMP1]]
//
float test_builtin_amdgcn_cvt_off_f32_i4_ui(unsigned n) {
    return __builtin_amdgcn_cvt_off_f32_i4(n);
}

// CHECK-LABEL: @test_builtin_amdgcn_cvt_off_f32_i4_i(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-NEXT:    store i32 [[N:%.*]], ptr addrspace(5) [[N_ADDR]], align 4
// CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspace(5) [[N_ADDR]], align 4
// CHECK-NEXT:    [[TMP1:%.*]] = call float @llvm.amdgcn.cvt.off.f32.i4(i32 [[TMP0]])
// CHECK-NEXT:    ret float [[TMP1]]
//
float test_builtin_amdgcn_cvt_off_f32_i4_i(int n) {
    return __builtin_amdgcn_cvt_off_f32_i4(n);
}