clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-gfx12.cl

   1 // REQUIRES: amdgpu-registered-target
   2
   3 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -emit-llvm -o - %s | FileCheck %s
   4
   5 typedef unsigned int uint;
   6
   7 // CHECK-LABEL: @builtins_amdgcn_dl_insts
   8 // CHECK: call float @llvm.amdgcn.dot4.f32.fp8.bf8(i32 %uiA, i32 %uiB, float %fC)
   9 // CHECK: call float @llvm.amdgcn.dot4.f32.bf8.fp8(i32 %uiA, i32 %uiB, float %fC)
  10 // CHECK: call float @llvm.amdgcn.dot4.f32.fp8.fp8(i32 %uiA, i32 %uiB, float %fC)
  11 // CHECK: call float @llvm.amdgcn.dot4.f32.bf8.bf8(i32 %uiA, i32 %uiB, float %fC)
  12
  13 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
  14 kernel void builtins_amdgcn_dl_insts_err(global float *fOut,
  15                                          uint uiA, uint uiB, float fC) {
  16   fOut[0] = __builtin_amdgcn_dot4_f32_fp8_bf8(uiA, uiB, fC);
  17   fOut[1] = __builtin_amdgcn_dot4_f32_bf8_fp8(uiA, uiB, fC);
  18   fOut[2] = __builtin_amdgcn_dot4_f32_fp8_fp8(uiA, uiB, fC);
  19   fOut[3] = __builtin_amdgcn_dot4_f32_bf8_bf8(uiA, uiB, fC);
  20 }