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-exception 6 // 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 18 using namespace clang; 19 using namespace CodeGen; 20 using namespace llvm; 21 22 Value *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_smoothstep: { 62 Value *Min = EmitScalarExpr(E->getArg(0)); 63 Value *Max = EmitScalarExpr(E->getArg(1)); 64 Value *X = EmitScalarExpr(E->getArg(2)); 65 assert(E->getArg(0)->getType()->hasFloatingRepresentation() && 66 E->getArg(1)->getType()->hasFloatingRepresentation() && 67 E->getArg(2)->getType()->hasFloatingRepresentation() && 68 "SmoothStep operands must have a float representation"); 69 return Builder.CreateIntrinsic( 70 /*ReturnType=*/Min->getType(), Intrinsic::spv_smoothstep, 71 ArrayRef<Value *>{Min, Max, X}, /*FMFSource=*/nullptr, 72 "spv.smoothstep"); 73 } 74 case SPIRV::BI__builtin_spirv_faceforward: { 75 Value *N = EmitScalarExpr(E->getArg(0)); 76 Value *I = EmitScalarExpr(E->getArg(1)); 77 Value *Ng = EmitScalarExpr(E->getArg(2)); 78 assert(E->getArg(0)->getType()->hasFloatingRepresentation() && 79 E->getArg(1)->getType()->hasFloatingRepresentation() && 80 E->getArg(2)->getType()->hasFloatingRepresentation() && 81 "FaceForward operands must have a float representation"); 82 return Builder.CreateIntrinsic( 83 /*ReturnType=*/N->getType(), Intrinsic::spv_faceforward, 84 ArrayRef<Value *>{N, I, Ng}, /*FMFSource=*/nullptr, "spv.faceforward"); 85 } 86 case SPIRV::BI__builtin_spirv_generic_cast_to_ptr_explicit: { 87 Value *Ptr = EmitScalarExpr(E->getArg(0)); 88 assert(E->getArg(0)->getType()->hasPointerRepresentation() && 89 E->getArg(1)->getType()->hasIntegerRepresentation() && 90 "GenericCastToPtrExplicit takes a pointer and an int"); 91 llvm::Type *Res = getTypes().ConvertType(E->getType()); 92 assert(Res->isPointerTy() && 93 "GenericCastToPtrExplicit doesn't return a pointer"); 94 llvm::CallInst *Call = Builder.CreateIntrinsic( 95 /*ReturnType=*/Res, Intrinsic::spv_generic_cast_to_ptr_explicit, 96 ArrayRef<Value *>{Ptr}, nullptr, "spv.generic_cast"); 97 Call->addRetAttr(llvm::Attribute::AttrKind::NoUndef); 98 return Call; 99 } 100 case SPIRV::BI__builtin_spirv_num_workgroups: 101 return Builder.CreateIntrinsic( 102 /*ReturnType=*/getTypes().ConvertType(E->getType()), 103 Intrinsic::spv_num_workgroups, 104 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr, 105 "spv.num.workgroups"); 106 case SPIRV::BI__builtin_spirv_workgroup_size: 107 return Builder.CreateIntrinsic( 108 /*ReturnType=*/getTypes().ConvertType(E->getType()), 109 Intrinsic::spv_workgroup_size, 110 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr, 111 "spv.workgroup.size"); 112 case SPIRV::BI__builtin_spirv_workgroup_id: 113 return Builder.CreateIntrinsic( 114 /*ReturnType=*/getTypes().ConvertType(E->getType()), 115 Intrinsic::spv_group_id, 116 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr, 117 "spv.group.id"); 118 case SPIRV::BI__builtin_spirv_local_invocation_id: 119 return Builder.CreateIntrinsic( 120 /*ReturnType=*/getTypes().ConvertType(E->getType()), 121 Intrinsic::spv_thread_id_in_group, 122 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr, 123 "spv.thread.id.in.group"); 124 case SPIRV::BI__builtin_spirv_global_invocation_id: 125 return Builder.CreateIntrinsic( 126 /*ReturnType=*/getTypes().ConvertType(E->getType()), 127 Intrinsic::spv_thread_id, 128 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr, 129 "spv.thread.id"); 130 case SPIRV::BI__builtin_spirv_global_size: 131 return Builder.CreateIntrinsic( 132 /*ReturnType=*/getTypes().ConvertType(E->getType()), 133 Intrinsic::spv_global_size, 134 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr, 135 "spv.num.workgroups"); 136 case SPIRV::BI__builtin_spirv_global_offset: 137 return Builder.CreateIntrinsic( 138 /*ReturnType=*/getTypes().ConvertType(E->getType()), 139 Intrinsic::spv_global_offset, 140 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr, 141 "spv.global.offset"); 142 } 143 return nullptr; 144 } 145