162 lines · cpp
1//===--------- SPIR.cpp - Emit LLVM Code for builtins ---------------------===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8//9// This contains code to emit Builtin calls as LLVM code.10//11//===----------------------------------------------------------------------===//12 13#include "CGHLSLRuntime.h"14#include "CodeGenFunction.h"15#include "clang/Basic/TargetBuiltins.h"16#include "llvm/IR/Intrinsics.h"17 18using namespace clang;19using namespace CodeGen;20using namespace llvm;21 22Value *CodeGenFunction::EmitSPIRVBuiltinExpr(unsigned BuiltinID,23 const CallExpr *E) {24 switch (BuiltinID) {25 case SPIRV::BI__builtin_spirv_distance: {26 Value *X = EmitScalarExpr(E->getArg(0));27 Value *Y = EmitScalarExpr(E->getArg(1));28 assert(E->getArg(0)->getType()->hasFloatingRepresentation() &&29 E->getArg(1)->getType()->hasFloatingRepresentation() &&30 "Distance operands must have a float representation");31 assert(E->getArg(0)->getType()->isVectorType() &&32 E->getArg(1)->getType()->isVectorType() &&33 "Distance operands must be a vector");34 return Builder.CreateIntrinsic(35 /*ReturnType=*/X->getType()->getScalarType(), Intrinsic::spv_distance,36 ArrayRef<Value *>{X, Y}, nullptr, "spv.distance");37 }38 case SPIRV::BI__builtin_spirv_length: {39 Value *X = EmitScalarExpr(E->getArg(0));40 assert(E->getArg(0)->getType()->hasFloatingRepresentation() &&41 "length operand must have a float representation");42 assert(E->getArg(0)->getType()->isVectorType() &&43 "length operand must be a vector");44 return Builder.CreateIntrinsic(45 /*ReturnType=*/X->getType()->getScalarType(), Intrinsic::spv_length,46 ArrayRef<Value *>{X}, nullptr, "spv.length");47 }48 case SPIRV::BI__builtin_spirv_reflect: {49 Value *I = EmitScalarExpr(E->getArg(0));50 Value *N = EmitScalarExpr(E->getArg(1));51 assert(E->getArg(0)->getType()->hasFloatingRepresentation() &&52 E->getArg(1)->getType()->hasFloatingRepresentation() &&53 "Reflect operands must have a float representation");54 assert(E->getArg(0)->getType()->isVectorType() &&55 E->getArg(1)->getType()->isVectorType() &&56 "Reflect operands must be a vector");57 return Builder.CreateIntrinsic(58 /*ReturnType=*/I->getType(), Intrinsic::spv_reflect,59 ArrayRef<Value *>{I, N}, nullptr, "spv.reflect");60 }61 case SPIRV::BI__builtin_spirv_refract: {62 Value *I = EmitScalarExpr(E->getArg(0));63 Value *N = EmitScalarExpr(E->getArg(1));64 Value *eta = EmitScalarExpr(E->getArg(2));65 assert(E->getArg(0)->getType()->hasFloatingRepresentation() &&66 E->getArg(1)->getType()->hasFloatingRepresentation() &&67 E->getArg(2)->getType()->isFloatingType() &&68 "refract operands must have a float representation");69 return Builder.CreateIntrinsic(70 /*ReturnType=*/I->getType(), Intrinsic::spv_refract,71 ArrayRef<Value *>{I, N, eta}, nullptr, "spv.refract");72 }73 case SPIRV::BI__builtin_spirv_smoothstep: {74 Value *Min = EmitScalarExpr(E->getArg(0));75 Value *Max = EmitScalarExpr(E->getArg(1));76 Value *X = EmitScalarExpr(E->getArg(2));77 assert(E->getArg(0)->getType()->hasFloatingRepresentation() &&78 E->getArg(1)->getType()->hasFloatingRepresentation() &&79 E->getArg(2)->getType()->hasFloatingRepresentation() &&80 "SmoothStep operands must have a float representation");81 return Builder.CreateIntrinsic(82 /*ReturnType=*/Min->getType(), Intrinsic::spv_smoothstep,83 ArrayRef<Value *>{Min, Max, X}, /*FMFSource=*/nullptr,84 "spv.smoothstep");85 }86 case SPIRV::BI__builtin_spirv_faceforward: {87 Value *N = EmitScalarExpr(E->getArg(0));88 Value *I = EmitScalarExpr(E->getArg(1));89 Value *Ng = EmitScalarExpr(E->getArg(2));90 assert(E->getArg(0)->getType()->hasFloatingRepresentation() &&91 E->getArg(1)->getType()->hasFloatingRepresentation() &&92 E->getArg(2)->getType()->hasFloatingRepresentation() &&93 "FaceForward operands must have a float representation");94 return Builder.CreateIntrinsic(95 /*ReturnType=*/N->getType(), Intrinsic::spv_faceforward,96 ArrayRef<Value *>{N, I, Ng}, /*FMFSource=*/nullptr, "spv.faceforward");97 }98 case SPIRV::BI__builtin_spirv_generic_cast_to_ptr_explicit: {99 Value *Ptr = EmitScalarExpr(E->getArg(0));100 assert(E->getArg(0)->getType()->hasPointerRepresentation() &&101 E->getArg(1)->getType()->hasIntegerRepresentation() &&102 "GenericCastToPtrExplicit takes a pointer and an int");103 llvm::Type *Res = getTypes().ConvertType(E->getType());104 assert(Res->isPointerTy() &&105 "GenericCastToPtrExplicit doesn't return a pointer");106 llvm::CallInst *Call = Builder.CreateIntrinsic(107 /*ReturnType=*/Res, Intrinsic::spv_generic_cast_to_ptr_explicit,108 ArrayRef<Value *>{Ptr}, nullptr, "spv.generic_cast");109 Call->addRetAttr(llvm::Attribute::AttrKind::NoUndef);110 return Call;111 }112 case SPIRV::BI__builtin_spirv_num_workgroups:113 return Builder.CreateIntrinsic(114 /*ReturnType=*/getTypes().ConvertType(E->getType()),115 Intrinsic::spv_num_workgroups,116 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,117 "spv.num.workgroups");118 case SPIRV::BI__builtin_spirv_workgroup_size:119 return Builder.CreateIntrinsic(120 /*ReturnType=*/getTypes().ConvertType(E->getType()),121 Intrinsic::spv_workgroup_size,122 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,123 "spv.workgroup.size");124 case SPIRV::BI__builtin_spirv_workgroup_id:125 return Builder.CreateIntrinsic(126 /*ReturnType=*/getTypes().ConvertType(E->getType()),127 Intrinsic::spv_group_id,128 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,129 "spv.group.id");130 case SPIRV::BI__builtin_spirv_local_invocation_id:131 return Builder.CreateIntrinsic(132 /*ReturnType=*/getTypes().ConvertType(E->getType()),133 Intrinsic::spv_thread_id_in_group,134 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,135 "spv.thread.id.in.group");136 case SPIRV::BI__builtin_spirv_global_invocation_id:137 return Builder.CreateIntrinsic(138 /*ReturnType=*/getTypes().ConvertType(E->getType()),139 Intrinsic::spv_thread_id,140 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,141 "spv.thread.id");142 case SPIRV::BI__builtin_spirv_global_size:143 return Builder.CreateIntrinsic(144 /*ReturnType=*/getTypes().ConvertType(E->getType()),145 Intrinsic::spv_global_size,146 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,147 "spv.num.workgroups");148 case SPIRV::BI__builtin_spirv_global_offset:149 return Builder.CreateIntrinsic(150 /*ReturnType=*/getTypes().ConvertType(E->getType()),151 Intrinsic::spv_global_offset,152 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,153 "spv.global.offset");154 case SPIRV::BI__builtin_spirv_fwidth:155 return Builder.CreateIntrinsic(156 /*ReturnType=*/getTypes().ConvertType(E->getType()),157 Intrinsic::spv_fwidth, ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))},158 nullptr, "spv.fwidth");159 }160 return nullptr;161}162