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
EmitSPIRVBuiltinExpr(unsigned BuiltinID,const CallExpr * E)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