87 lines · plain
1// REQUIRES: amdgpu-registered-target2// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu fiji -emit-llvm -o - %s | FileCheck -enable-var-scope --check-prefixes=CHECK %s3 4 5#pragma OPENCL EXTENSION cl_khr_fp64 : enable6 7typedef unsigned long ulong;8typedef unsigned int uint;9typedef unsigned short ushort;10typedef half __attribute__((ext_vector_type(2))) half2;11typedef short __attribute__((ext_vector_type(2))) short2;12typedef ushort __attribute__((ext_vector_type(2))) ushort2;13typedef uint __attribute__((ext_vector_type(4))) uint4;14 15// CHECK-LABEL: @test_lerp16// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.lerp17void test_lerp(global int* out, int a, int b, int c)18{19 *out = __builtin_amdgcn_lerp(a, b, c);20}21 22// CHECK-LABEL: @test_cubeid(23// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.cubeid(float %a, float %b, float %c)24void test_cubeid(global float* out, float a, float b, float c) {25 *out = __builtin_amdgcn_cubeid(a, b, c);26}27 28// CHECK-LABEL: @test_cubesc(29// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.cubesc(float %a, float %b, float %c)30void test_cubesc(global float* out, float a, float b, float c) {31 *out = __builtin_amdgcn_cubesc(a, b, c);32}33 34// CHECK-LABEL: @test_cubetc(35// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.cubetc(float %a, float %b, float %c)36void test_cubetc(global float* out, float a, float b, float c) {37 *out = __builtin_amdgcn_cubetc(a, b, c);38}39 40// CHECK-LABEL: @test_cubema(41// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.cubema(float %a, float %b, float %c)42void test_cubema(global float* out, float a, float b, float c) {43 *out = __builtin_amdgcn_cubema(a, b, c);44}45 46// CHECK-LABEL: @test_cvt_pknorm_i16(47// CHECK: tail call{{.*}} <2 x i16> @llvm.amdgcn.cvt.pknorm.i16(float %src0, float %src1)48kernel void test_cvt_pknorm_i16(global short2* out, float src0, float src1) {49 *out = __builtin_amdgcn_cvt_pknorm_i16(src0, src1);50}51 52// CHECK-LABEL: @test_cvt_pknorm_u16(53// CHECK: tail call{{.*}} <2 x i16> @llvm.amdgcn.cvt.pknorm.u16(float %src0, float %src1)54kernel void test_cvt_pknorm_u16(global ushort2* out, float src0, float src1) {55 *out = __builtin_amdgcn_cvt_pknorm_u16(src0, src1);56}57 58// CHECK-LABEL: @test_sad_u8(59// CHECK: tail call{{.*}} i32 @llvm.amdgcn.sad.u8(i32 %src0, i32 %src1, i32 %src2)60kernel void test_sad_u8(global uint* out, uint src0, uint src1, uint src2) {61 *out = __builtin_amdgcn_sad_u8(src0, src1, src2);62}63 64// CHECK-LABEL: test_msad_u8(65// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.msad.u8(i32 %src0, i32 %src1, i32 %src2)66kernel void test_msad_u8(global uint* out, uint src0, uint src1, uint src2) {67 *out = __builtin_amdgcn_msad_u8(src0, src1, src2);68}69 70// CHECK-LABEL: test_sad_hi_u8(71// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.sad.hi.u8(i32 %src0, i32 %src1, i32 %src2)72kernel void test_sad_hi_u8(global uint* out, uint src0, uint src1, uint src2) {73 *out = __builtin_amdgcn_sad_hi_u8(src0, src1, src2);74}75 76// CHECK-LABEL: @test_sad_u16(77// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.sad.u16(i32 %src0, i32 %src1, i32 %src2)78kernel void test_sad_u16(global uint* out, uint src0, uint src1, uint src2) {79 *out = __builtin_amdgcn_sad_u16(src0, src1, src2);80}81 82// CHECK-LABEL: @test_qsad_pk_u16_u8(83// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.qsad.pk.u16.u8(i64 %src0, i32 %src1, i64 %src2)84kernel void test_qsad_pk_u16_u8(global ulong* out, ulong src0, uint src1, ulong src2) {85 *out = __builtin_amdgcn_qsad_pk_u16_u8(src0, src1, src2);86}87