SPIR.cpp Source File (original) (raw)

24 switch (BuiltinID) {

25 case SPIRV::BI__builtin_spirv_distance: {

30 "Distance operands must have a float representation");

33 "Distance operands must be a vector");

34 return Builder.CreateIntrinsic(

35 X->getType()->getScalarType(), Intrinsic::spv_distance,

37 }

38 case SPIRV::BI__builtin_spirv_length: {

41 "length operand must have a float representation");

43 "length operand must be a vector");

44 return Builder.CreateIntrinsic(

45 X->getType()->getScalarType(), Intrinsic::spv_length,

47 }

48 case SPIRV::BI__builtin_spirv_reflect: {

53 "Reflect operands must have a float representation");

56 "Reflect operands must be a vector");

57 return Builder.CreateIntrinsic(

58 I->getType(), Intrinsic::spv_reflect,

60 }

61 case SPIRV::BI__builtin_spirv_refract: {

68 "refract operands must have a float representation");

69 return Builder.CreateIntrinsic(

70 I->getType(), Intrinsic::spv_refract,

72 }

73 case SPIRV::BI__builtin_spirv_smoothstep: {

80 "SmoothStep operands must have a float representation");

81 return Builder.CreateIntrinsic(

82 Min->getType(), Intrinsic::spv_smoothstep,

84 "spv.smoothstep");

85 }

86 case SPIRV::BI__builtin_spirv_faceforward: {

93 "FaceForward operands must have a float representation");

94 return Builder.CreateIntrinsic(

95 N->getType(), Intrinsic::spv_faceforward,

96 ArrayRef<Value *>{N, I, Ng}, nullptr, "spv.faceforward");

97 }

98 case SPIRV::BI__builtin_spirv_generic_cast_to_ptr_explicit: {

102 "GenericCastToPtrExplicit takes a pointer and an int");

104 assert(Res->isPointerTy() &&

105 "GenericCastToPtrExplicit doesn't return a pointer");

106 llvm::CallInst *Call = Builder.CreateIntrinsic(

107 Res, Intrinsic::spv_generic_cast_to_ptr_explicit,

109 Call->addRetAttr(llvm::Attribute::AttrKind::NoUndef);

111 }

112 case SPIRV::BI__builtin_spirv_num_workgroups:

113 return Builder.CreateIntrinsic(

115 Intrinsic::spv_num_workgroups,

117 "spv.num.workgroups");

118 case SPIRV::BI__builtin_spirv_workgroup_size:

119 return Builder.CreateIntrinsic(

121 Intrinsic::spv_workgroup_size,

123 "spv.workgroup.size");

124 case SPIRV::BI__builtin_spirv_workgroup_id:

125 return Builder.CreateIntrinsic(

127 Intrinsic::spv_group_id,

129 "spv.group.id");

130 case SPIRV::BI__builtin_spirv_local_invocation_id:

131 return Builder.CreateIntrinsic(

133 Intrinsic::spv_thread_id_in_group,

135 "spv.thread.id.in.group");

136 case SPIRV::BI__builtin_spirv_global_invocation_id:

137 return Builder.CreateIntrinsic(

139 Intrinsic::spv_thread_id,

141 "spv.thread.id");

142 case SPIRV::BI__builtin_spirv_global_size:

143 return Builder.CreateIntrinsic(

145 Intrinsic::spv_global_size,

147 "spv.num.workgroups");

148 case SPIRV::BI__builtin_spirv_global_offset:

149 return Builder.CreateIntrinsic(

151 Intrinsic::spv_global_offset,

153 "spv.global.offset");

154 case SPIRV::BI__builtin_spirv_ddx:

155 return Builder.CreateIntrinsic(

157 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr, "spv.ddx");

158 case SPIRV::BI__builtin_spirv_ddy:

159 return Builder.CreateIntrinsic(

161 ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr, "spv.ddy");

162 case SPIRV::BI__builtin_spirv_fwidth:

163 return Builder.CreateIntrinsic(

165 Intrinsic::spv_fwidth, ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))},

166 nullptr, "spv.fwidth");

167 }

168 return nullptr;

169}