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
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: E->getArg(Arg: 0));
27 Value *Y = EmitScalarExpr(E: E->getArg(Arg: 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=*/RetTy: X->getType()->getScalarType(), ID: Intrinsic::spv_distance,
36 Args: ArrayRef<Value *>{X, Y}, FMFSource: nullptr, Name: "spv.distance");
37 }
38 case SPIRV::BI__builtin_spirv_length: {
39 Value *X = EmitScalarExpr(E: E->getArg(Arg: 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=*/RetTy: X->getType()->getScalarType(), ID: Intrinsic::spv_length,
46 Args: ArrayRef<Value *>{X}, FMFSource: nullptr, Name: "spv.length");
47 }
48 case SPIRV::BI__builtin_spirv_reflect: {
49 Value *I = EmitScalarExpr(E: E->getArg(Arg: 0));
50 Value *N = EmitScalarExpr(E: E->getArg(Arg: 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=*/RetTy: I->getType(), ID: Intrinsic::spv_reflect,
59 Args: ArrayRef<Value *>{I, N}, FMFSource: nullptr, Name: "spv.reflect");
60 }
61 case SPIRV::BI__builtin_spirv_refract: {
62 Value *I = EmitScalarExpr(E: E->getArg(Arg: 0));
63 Value *N = EmitScalarExpr(E: E->getArg(Arg: 1));
64 Value *eta = EmitScalarExpr(E: E->getArg(Arg: 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=*/RetTy: I->getType(), ID: Intrinsic::spv_refract,
71 Args: ArrayRef<Value *>{I, N, eta}, FMFSource: nullptr, Name: "spv.refract");
72 }
73 case SPIRV::BI__builtin_spirv_smoothstep: {
74 Value *Min = EmitScalarExpr(E: E->getArg(Arg: 0));
75 Value *Max = EmitScalarExpr(E: E->getArg(Arg: 1));
76 Value *X = EmitScalarExpr(E: E->getArg(Arg: 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=*/RetTy: Min->getType(), ID: Intrinsic::spv_smoothstep,
83 Args: ArrayRef<Value *>{Min, Max, X}, /*FMFSource=*/nullptr,
84 Name: "spv.smoothstep");
85 }
86 case SPIRV::BI__builtin_spirv_faceforward: {
87 Value *N = EmitScalarExpr(E: E->getArg(Arg: 0));
88 Value *I = EmitScalarExpr(E: E->getArg(Arg: 1));
89 Value *Ng = EmitScalarExpr(E: E->getArg(Arg: 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=*/RetTy: N->getType(), ID: Intrinsic::spv_faceforward,
96 Args: ArrayRef<Value *>{N, I, Ng}, /*FMFSource=*/nullptr, Name: "spv.faceforward");
97 }
98 case SPIRV::BI__builtin_spirv_generic_cast_to_ptr_explicit: {
99 Value *Ptr = EmitScalarExpr(E: E->getArg(Arg: 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(T: E->getType());
104 assert(Res->isPointerTy() &&
105 "GenericCastToPtrExplicit doesn't return a pointer");
106 llvm::CallInst *Call = Builder.CreateIntrinsic(
107 /*ReturnType=*/RetTy: Res, ID: Intrinsic::spv_generic_cast_to_ptr_explicit,
108 Args: ArrayRef<Value *>{Ptr}, FMFSource: nullptr, Name: "spv.generic_cast");
109 Call->addRetAttr(Kind: llvm::Attribute::AttrKind::NoUndef);
110 return Call;
111 }
112 case SPIRV::BI__builtin_spirv_subgroup_shuffle: {
113 Value *X = EmitScalarExpr(E: E->getArg(Arg: 0));
114 Value *Y = EmitScalarExpr(E: E->getArg(Arg: 1));
115 assert(E->getArg(1)->getType()->hasIntegerRepresentation());
116 return Builder.CreateIntrinsic(
117 /*ReturnType=*/RetTy: getTypes().ConvertType(T: E->getArg(Arg: 0)->getType()),
118 ID: Intrinsic::spv_wave_readlane, Args: {X, Y}, FMFSource: nullptr, Name: "spv.shuffle");
119 }
120 case SPIRV::BI__builtin_spirv_num_workgroups:
121 return Builder.CreateIntrinsic(
122 /*ReturnType=*/RetTy: getTypes().ConvertType(T: E->getType()),
123 ID: Intrinsic::spv_num_workgroups,
124 Args: ArrayRef<Value *>{EmitScalarExpr(E: E->getArg(Arg: 0))}, FMFSource: nullptr,
125 Name: "spv.num.workgroups");
126 case SPIRV::BI__builtin_spirv_workgroup_size:
127 return Builder.CreateIntrinsic(
128 /*ReturnType=*/RetTy: getTypes().ConvertType(T: E->getType()),
129 ID: Intrinsic::spv_workgroup_size,
130 Args: ArrayRef<Value *>{EmitScalarExpr(E: E->getArg(Arg: 0))}, FMFSource: nullptr,
131 Name: "spv.workgroup.size");
132 case SPIRV::BI__builtin_spirv_workgroup_id:
133 return Builder.CreateIntrinsic(
134 /*ReturnType=*/RetTy: getTypes().ConvertType(T: E->getType()),
135 ID: Intrinsic::spv_group_id,
136 Args: ArrayRef<Value *>{EmitScalarExpr(E: E->getArg(Arg: 0))}, FMFSource: nullptr,
137 Name: "spv.group.id");
138 case SPIRV::BI__builtin_spirv_local_invocation_id:
139 return Builder.CreateIntrinsic(
140 /*ReturnType=*/RetTy: getTypes().ConvertType(T: E->getType()),
141 ID: Intrinsic::spv_thread_id_in_group,
142 Args: ArrayRef<Value *>{EmitScalarExpr(E: E->getArg(Arg: 0))}, FMFSource: nullptr,
143 Name: "spv.thread.id.in.group");
144 case SPIRV::BI__builtin_spirv_global_invocation_id:
145 return Builder.CreateIntrinsic(
146 /*ReturnType=*/RetTy: getTypes().ConvertType(T: E->getType()),
147 ID: Intrinsic::spv_thread_id,
148 Args: ArrayRef<Value *>{EmitScalarExpr(E: E->getArg(Arg: 0))}, FMFSource: nullptr,
149 Name: "spv.thread.id");
150 case SPIRV::BI__builtin_spirv_global_size:
151 return Builder.CreateIntrinsic(
152 /*ReturnType=*/RetTy: getTypes().ConvertType(T: E->getType()),
153 ID: Intrinsic::spv_global_size,
154 Args: ArrayRef<Value *>{EmitScalarExpr(E: E->getArg(Arg: 0))}, FMFSource: nullptr,
155 Name: "spv.num.workgroups");
156 case SPIRV::BI__builtin_spirv_global_offset:
157 return Builder.CreateIntrinsic(
158 /*ReturnType=*/RetTy: getTypes().ConvertType(T: E->getType()),
159 ID: Intrinsic::spv_global_offset,
160 Args: ArrayRef<Value *>{EmitScalarExpr(E: E->getArg(Arg: 0))}, FMFSource: nullptr,
161 Name: "spv.global.offset");
162 case SPIRV::BI__builtin_spirv_ddx:
163 return Builder.CreateIntrinsic(
164 /*ReturnType=*/RetTy: getTypes().ConvertType(T: E->getType()), ID: Intrinsic::spv_ddx,
165 Args: ArrayRef<Value *>{EmitScalarExpr(E: E->getArg(Arg: 0))}, FMFSource: nullptr, Name: "spv.ddx");
166 case SPIRV::BI__builtin_spirv_ddy:
167 return Builder.CreateIntrinsic(
168 /*ReturnType=*/RetTy: getTypes().ConvertType(T: E->getType()), ID: Intrinsic::spv_ddy,
169 Args: ArrayRef<Value *>{EmitScalarExpr(E: E->getArg(Arg: 0))}, FMFSource: nullptr, Name: "spv.ddy");
170 case SPIRV::BI__builtin_spirv_fwidth:
171 return Builder.CreateIntrinsic(
172 /*ReturnType=*/RetTy: getTypes().ConvertType(T: E->getType()),
173 ID: Intrinsic::spv_fwidth, Args: ArrayRef<Value *>{EmitScalarExpr(E: E->getArg(Arg: 0))},
174 FMFSource: nullptr, Name: "spv.fwidth");
175 }
176 return nullptr;
177}
178